CUDA 直方图遇到非法内存访问 (77）

CUDA Histogram an illegal memory access was encountered (77)

本文关键字：访问 内存 直方图 遇到 非法 CUDA      更新时间：2023-10-16

所以这是我几乎完整的代码：第一个内核是正常的全局直方图，工作正常。但是我收到错误"遇到非法内存访问（77）"在计算shared_histogram后的最终memcpy。我不知道代码有什么问题。似乎共享直方图确实改变了d_hist2的大小。我还检查了bin_count是否更改了，但没有。那么我的shared_histog内核是错误的还是我在memCpy上做错了？注意：W * H * NC是我输入图像的大小

__global__ void histog( int *img, int *hist, int bin_count, int n)
{
int x = threadIdx.x + blockDim.x *blockIdx.x;
if(x>=n)   return;
    unsigned char value = img[x];
int bin = value % bin_count;
atomicAdd(&hist[bin],1);
}
__global__ void shared_histog( int *img, int *hist, int n)
{
int x = threadIdx.x + blockDim.x *blockIdx.x;
int indx = threadIdx.x;
if(x>n)   return;
__shared__ int shHist[256];
if (indx < 256)
shHist[indx] =0;
__syncthreads();
unsigned char value = img[x];
__syncthreads();
atomicAdd( (int*)&shHist[value], 1);
__syncthreads();
atomicAdd( (int*)&(hist[indx]), shHist[indx] );
}

int main(int argc, char **argv)
{
cudaDeviceSynchronize();  CUDA_CHECK;

int *imgval  = new int[(size_t)w*h*nc];
for (int i =0; i<w*h*nc; i++)
    imgval[i] = (imgIn[i])*256 + 1;

int bin_count = 256;
int *Histogram  = new int[bin_count];
int *Histogram2  = new int[bin_count];
for (int i =0; i <bin_count; i++)
    Histogram2[i] = 0;
Timer timer; timer.start();
for (int i =0; i <bin_count; i++)
    Histogram[i] = 0;
for (int i =0; i<w*h*nc; i++)
    Histogram[(imgval[i])]++;
showHistogram256("CPU_Histo", Histogram, 100 + w + 40, 100);

timer.end();  float t = timer.get();  // elapsed time in seconds
cout << "CPU time: " << t*1000 << " ms" << endl;

int *d_img = NULL;
int nbytes = w * h * nc * sizeof(int);
cudaMalloc(&d_img, nbytes); CUDA_CHECK;
cudaMemcpy(d_img, imgval, nbytes, cudaMemcpyHostToDevice); CUDA_CHECK;
int *d_hist = NULL;
cudaMalloc(&d_hist, bin_count * sizeof(int)); CUDA_CHECK;
cudaMemset(d_hist, 0, bin_count * sizeof(int)); CUDA_CHECK;
int *d_hist2 = NULL;
cudaMalloc(&d_hist2, bin_count * sizeof(int)); CUDA_CHECK;
cudaMemset(d_hist2, 0, bin_count * sizeof(int)); CUDA_CHECK;
dim3 block = dim3(1024,1,1);
dim3 grid = dim3 ((w*h*nc+block.x-1)/block.x, 1, 1);
Timer timer2; timer2.start();
histog <<<grid, block>>> (d_img, d_hist, bin_count, nbytes);    CUDA_CHECK;
    timer2.end();  float t2 = timer2.get();  // elapsed time in seconds
    cout << "GPU time: " << t2*1000 << " ms" << endl;
cudaMemcpy(Histogram, d_hist,bin_count * sizeof(int), cudaMemcpyDeviceToHost); CUDA_CHECK;
showHistogram256("GPU_Histo", Histogram, 100 + w + 40, 100 + h/2 + 10);

Timer timer3; timer3.start();
shared_histog <<<grid, block>>> (d_img, d_hist2, nbytes); CUDA_CHECK;
timer3.end();  float t3 = timer3.get();  // elapsed time in seconds
    cout << "Shared time: " << t3*1000 << " ms" << endl;
*

错误来了 *

cudaMemcpy(Histogram2, d_hist2, 256 * sizeof(int), cudaMemcpyDeviceToHost);  CUDA_CHECK;
showHistogram256("GPU_Histo_Shared", Histogram2, 100 + w + 40, 100 + h +10);

return 0;
}

你在条件语句后使用 __syncthreads()：

if(x>n)   return;

这可能会阻止块中的所有线程访问它。 这不是正确的用法：

条件代码中允许使用 __syncthreads（），但前提是条件在整个线程块中的计算方式相同，否则代码执行可能会挂起或产生意外的副作用。

但它可能与非法内存访问无关。

您正在启动此内核，每个块有 1024 个线程：

dim3 block = dim3(1024,1,1);

这意味着在内核中，您的indx变量：

int indx = threadIdx.x;

将根据线程从 0..1023 开始，这意味着此行：

atomicAdd( (int*)&(hist[indx]), shHist[indx] );
                        ^^^^           ^^^^

将尝试索引到indx值大于 255 的线程的 hist 和 shHist 越界，因为 hist 和 shHist 都只分配了 256 个元素。

您可以通过添加条件语句来解决此问题：

if (indx < 256) 
  atomicAdd( (int*)&(hist[indx]), shHist[indx] );

如果使用-lineinfo编译并使用cuda-memcheck，您实际上可以cuda-memcheck查明生成越界访问的源代码行。

相关文章：