最小缩减cuda不起作用

我写了一个代码来找到最小化缩减。但结果始终为零。我不知道是什么问题。请帮帮我。最小缩减cuda不起作用

这里是内核代码：我修改了Nvidia的总和缩减代码。

#include <limits.h> 

#define NumThread 128 
#define NumBlock 32 

__global__ void min_reduce(int* In, int* Out, int n){ 
    __shared__ int sdata[NumThread]; 
    unsigned int i = blockIdx.x * NumThread + threadIdx.x; 
    unsigned int tid = threadIdx.x; 
    unsigned int gridSize = NumBlock * NumThread; 
    int myMin = INT_MAX; 

    while (i < n){ 
    if(In[i] < myMin) 
    myMin = In[i]; 
    i += gridSize; 
    } 
    sdata[tid] = myMin; 
    __syncthreads(); 

    if (NumThread >= 1024){ 
    if (tid < 512) 
    if(sdata[tid] > sdata[tid + 512]) sdata[tid] = sdata[tid + 512]; 
    __syncthreads(); 
    } 
    if (NumThread >= 512){ 
    if(sdata[tid] > sdata[tid + 256]) sdata[tid] = sdata[tid + 256]; 
    __syncthreads(); 
    } 
    if (NumThread >= 256){ 
    if(sdata[tid] > sdata[tid + 128] && sdata[tid + 128] !=0) sdata[tid] = sdata[tid + 128]; 
    __syncthreads(); 
    } 
    if (NumThread >= 128){ 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    } 
    //the following practice is deprecated 
    if (tid < 32){ 
    volatile int *smem = sdata; 
    if (NumThread >= 64) if(smem[tid] > smem[tid + 32]) smem[tid] = smem[tid+32]; 
    if (NumThread >= 32) if(smem[tid] > smem[tid + 16]) smem[tid] = smem[tid+16]; 
    if (NumThread >= 16) if(smem[tid] > smem[tid + 8]) smem[tid] = smem[tid+8]; 
    if (NumThread >= 8) if(smem[tid] > smem[tid + 4]) smem[tid] = smem[tid+4]; 
    if (NumThread >= 4) if(smem[tid] > smem[tid + 2]) smem[tid] = smem[tid+2]; 
    if (NumThread >= 2) if(smem[tid] > smem[tid + 1])  smem[tid] = smem[tid+1]; 
    } 
    if (tid == 0) 
    if(sdata[0] < sdata[1]) Out[blockIdx.x] = sdata[0]; 
    else Out[blockIdx.x] = sdata[1];  
}

在这里，这是我的主要代码：

#include <stdio.h> 
#include <stdlib.h> 

#include "min_reduction.cu" 

int main(int argc, char* argv[]){ 
    unsigned int length = 1048576; 
    int i, Size, min; 
    int *a, *out, *gpuA, *gpuOut; 

    cudaSetDevice(0); 
    Size = length * sizeof(int); 
    a = (int*)malloc(Size); 
    out = (int*)malloc(NumBlock*sizeof(int)); 
    for(i=0;i<length;i++) a[i] = (i + 10); 

    cudaMalloc((void**)&gpuA,Size); 
    cudaMalloc((void**)&gpuOut,NumBlock*sizeof(int)); 
    cudaMemcpy(gpuA,a,Size,cudaMemcpyHostToDevice); 
    min_reduce<<<NumBlock,NumThread>>>(gpuA,gpuOut,length); 
    cudaDeviceSynchronize(); 
    cudaMemcpy(out,gpuOut,NumBlock*sizeof(int),cudaMemcpyDeviceToHost); 

    min = out[0]; 
    for(i=1;i<NumBlock;i++) if(min < out[i]) min = out[i]; 
    return 0; 
}

来源

2016-02-12 Hamid_UMB

帮你一个忙，并用[适当的错误检查]包围所有你的CUDA调用（http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-错误 - 使用最CUDA的运行时API）。你的代码有不正确的同步，'__syncthreads'应该被执行[由块中的所有线程]（http://stackoverflow.com/questions/6666382/can-i-use-syncthreads-after-having-dropped-threads）。此外，在这里使用'volatile'是不正确的，歪曲同步编程被破坏**，请不要使用它。 –

我不知道我的一切，@HubertApplebaum表示同意，但我可以用proper cuda error checking的建议达成一致。正如你在代码中提到的，warp同步编程可以被认为是已弃用，但我不能支持它是中断（还）的说法。但我不想为此争论;这不是你的问题的核心。

另一个有用的调试建议将遵循步骤here编译您的代码与-lineinfo并运行您的代码与cuda-memcheck。如果你这样做，你会看到很多这样的报道：

========= Invalid __shared__ read of size 4 
=========  at 0x000001e0 in /home/bob/misc/t1074.cu:39:min_reduce(int*, int*, int) 
=========  by thread (64,0,0) in block (24,0,0) 
=========  Address 0x00000200 is out of bounds 
=========  Saved host backtrace up to driver entry point at kernel launch time 
=========  Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x15859d] 
=========  Host Frame:./t1074 [0x16dc1] 
=========  Host Frame:./t1074 [0x315d3] 
=========  Host Frame:./t1074 [0x28f5] 
=========  Host Frame:./t1074 [0x2623] 
=========  Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21d65] 
=========  Host Frame:./t1074 [0x271d]

这将表明双方是在你的代码的主要问题是，你是不正确索引到你的__shared__存储器阵列以及具体线路代码在哪里发生。整齐！（对我而言，这是第39行，但在您的情况下可能会有不同的行）。如果再钻入那行，你将要学习的这部分代码：

#define NumThread 128 
    ... 
    __shared__ int sdata[NumThread]; 
    ... 
    if (NumThread >= 128){ 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; //line 39 in my case 
    __syncthreads(); 
    }

您已经定义NumThread在128，并且有静态分配的，许多int数量的共享存储阵列。一切都很好。 if语句中的代码怎么样？ if条件将得到满足，这意味着块中的所有128个线程将执行该if语句的主体。但是，您正在从共享内存中读取sdata[tid + 64]，而对于tid大于63（即每个块中的线程数量的一半）的线程，这会生成一个大于127的共享内存索引（这是超出范围即非法）。

（你已经显示了这具体代码）的修补程序相当简单，只需添加另一种，如果测试：

if (NumThread >= 128){ 
    if (tid < 64) 
     if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    }

如果你作出这样的修改你的代码，并重新运行cuda-memcheck测试，你会看到所有运行时报告的错误都消失了。好极了！

但是代码仍然没有产生正确的答案。你在这里做另一个错误：

for(i=1;i<NumBlock;i++) if(min < out[i]) min = out[i];

如果你想找到最低值，并仔细想想这个逻辑，你会意识到你应该这样做：

for(i=1;i<NumBlock;i++) if(min > out[i]) min = out[i]; 
           ^
           | 
           greater than

有了这两个变化，你的代码产生正确的结果对我来说：

$ cat t1074.cu 
#include <stdio.h> 
#include <stdlib.h> 


#include <limits.h> 

#define NumThread 128 
#define NumBlock 32 

__global__ void min_reduce(int* In, int* Out, int n){ 
    __shared__ int sdata[NumThread]; 
    unsigned int i = blockIdx.x * NumThread + threadIdx.x; 
    unsigned int tid = threadIdx.x; 
    unsigned int gridSize = NumBlock * NumThread; 
    int myMin = INT_MAX; 

    while (i < n){ 
    if(In[i] < myMin) 
    myMin = In[i]; 
    i += gridSize; 
    } 
    sdata[tid] = myMin; 
    __syncthreads(); 

    if (NumThread >= 1024){ 
    if (tid < 512) 
    if(sdata[tid] > sdata[tid + 512]) sdata[tid] = sdata[tid + 512]; 
    __syncthreads(); 
    } 
    if (NumThread >= 512){ 
    if(sdata[tid] > sdata[tid + 256]) sdata[tid] = sdata[tid + 256]; 
    __syncthreads(); 
    } 
    if (NumThread >= 256){ 
    if(sdata[tid] > sdata[tid + 128] && sdata[tid + 128] !=0) sdata[tid] = sdata[tid + 128]; 
    __syncthreads(); 
    } 
    if (NumThread >= 128){ 
    if (tid < 64) 
    if(sdata[tid] > sdata[tid + 64]) sdata[tid] = sdata[tid + 64]; 
    __syncthreads(); 
    } 
    //the following practice is deprecated 
    if (tid < 32){ 
    volatile int *smem = sdata; 
    if (NumThread >= 64) if(smem[tid] > smem[tid + 32]) smem[tid] = smem[tid+32]; 
    if (NumThread >= 32) if(smem[tid] > smem[tid + 16]) smem[tid] = smem[tid+16]; 
    if (NumThread >= 16) if(smem[tid] > smem[tid + 8]) smem[tid] = smem[tid+8]; 
    if (NumThread >= 8) if(smem[tid] > smem[tid + 4]) smem[tid] = smem[tid+4]; 
    if (NumThread >= 4) if(smem[tid] > smem[tid + 2]) smem[tid] = smem[tid+2]; 
    if (NumThread >= 2) if(smem[tid] > smem[tid + 1])  smem[tid] = smem[tid+1]; 
    } 
    if (tid == 0) 
    if(sdata[0] < sdata[1]) Out[blockIdx.x] = sdata[0]; 
    else Out[blockIdx.x] = sdata[1]; 
} 

int main(int argc, char* argv[]){ 
    unsigned int length = 1048576; 
    int i, Size, min; 
    int *a, *out, *gpuA, *gpuOut; 

    cudaSetDevice(0); 
    Size = length * sizeof(int); 
    a = (int*)malloc(Size); 
    out = (int*)malloc(NumBlock*sizeof(int)); 
    for(i=0;i<length;i++) a[i] = (i + 10); 
    a[10]=5; 
    cudaMalloc((void**)&gpuA,Size); 
    cudaMalloc((void**)&gpuOut,NumBlock*sizeof(int)); 
    cudaMemcpy(gpuA,a,Size,cudaMemcpyHostToDevice); 
    min_reduce<<<NumBlock,NumThread>>>(gpuA,gpuOut,length); 
    cudaDeviceSynchronize(); 
    cudaMemcpy(out,gpuOut,NumBlock*sizeof(int),cudaMemcpyDeviceToHost); 

    min = out[0]; 
    for(i=1;i<NumBlock;i++) if(min > out[i]) min = out[i]; 
    printf("min = %d\n", min); 
    return 0; 
} 
$ nvcc -o t1074 t1074.cu 
$ cuda-memcheck ./t1074 
========= CUDA-MEMCHECK 
min = 5 
========= ERROR SUMMARY: 0 errors 
$

注意，你已经在1024个线程情况下，如果检查，你可能想要为512和256线程案例添加一个适当的if-check，就像我为上面的128个线程案例添加它一样。

来源

2016-02-12 05:16:30

谢谢你的回答。代码工作正常，但是当我改变输入数组a，并添加这一行a [10] = 5;初始化后。代码，找不到最小值？你能告诉我我的代码有什么问题吗？ –

我加了一个[10] = 5;初始化到我在我的答案张贴的代码后，它似乎工作正常。（我已经在我的答案中更新了完整的测试案例以证明这一点）。也许您应该针对您的新代码提出一个新问题，而这个新问题不起作用。您是否添加了适当的cuda错误检查并使用'cuda-memcheck'运行您的代码？ –

谢谢......问题已解决 –

最小缩减cuda不起作用

回答

相关问题