原子函数做直方图统计

加速计算 CUDA
1

请问用原子函数做统计直方图，这么做为什么得到的结果是错的呢__global__ void Histkernel(floatdev_Img,int width,int height,intdev_Hist)
{
int tid_x=threadIdx.x+blockIdx.xblockDim.x;
int tid_y=threadIdx.y+blockIdx.yblockDim.y;
int thread_x=threadIdx.x;
int thread_y=threadIdx.y;

__shared__ int s_Hist[256];

if((tid_x<width)&&(tid_y<height))
{
	atomicAdd(&s_Hist[(int)dev_Img[tid_y*width+tid_x]],1);	
}
__syncthreads();

atomicAdd(&dev_Hist[thread_x+thread_y*16],s_Hist[thread_x+thread_y*16]);
__syncthreads();

}

得到的结果是错的

2

LZ您好：

错和对是相对的，为了鉴别何为错，需要知道何为对。
所以请您补充以下各点信息：

1：您要解决的问题的描述。

2：您打算解决这个问题的规划和设计。

3：您的kernel的各部分的实现意图。

4：您的kernel的各个参数的含义。

5：您的kernel的调用写法，以及您host端各个和kernel相关的参数的情况。

6：您的计算结果的情况，以及这和您的预计有何偏差。

希望LZ按照上述6点梳理一下思路和结果，并予以反馈，如此不仅能使LZ对自己的问题有进一步的认识，而且可以使得读帖者高效准确地了解您的情况，并与您讨论。

以及这并非是故意为难LZ，这是常规的分析问题的方式和方法。

大致如此，请LZ积极反馈。

祝您编码顺利~

3

以及，我补充一下：

您这里：
atomicAdd(&s_Hist[(int)dev_Img[tid_y*width+tid_x]],1);
需要确保您的直接(int)后的dev_img中的每个元素都在0-255范围内（您是从float cast过来的，这个不一定能保证）。否则将出错。

其二，
您没有初始化贵缓冲区s_dist，里面的内容不一定保证自动为0的。请初始化他们。

其三，
您还需要一个单独的kernel, 在运行这个kernel之前，对dev_Hist进行清零的(或者cudaMemset)

这是补充ICE之外的您可能有的问题。

4

谢谢版主提醒

5

谢谢版主的建议

6

正如版主说的
1.我的dev_img存放的是一幅灰度图像的像素值，数值的范围是0到255，经由float强制转化成int型，会超出0到255这个范围吗
2.我的dev_Hist是由kernel函数统计的直方图结果，在host端再有显存写入内存，清0的方法是直接
cudaMemset（dev_Hist，0，255*sizeof（int））
3.版主说再添加一个kernel函数吗，我的程序是先算出每一个block的小直方图统计，再用一个原子函数把各个小block的小直方图统计成一个大的，还有必要在写一个kernel函数吗

7

LZ您好：

1：因为float的表达范围远大于0~255，所以横扫斑竹提示您这一点。以及如果您的数据能够先验地保证这一点，可以放心使用。

2：（按照横扫斑竹顺序）您需要对您每个block申请的shared memory中的s_Hist[256]这个数组进行初始化，因为手册并不保证您直接申请的shared memory里面的数组里的数据是初始化过的，您需要在kerenl里面用线程对这个数组写一遍零，以保证后续操作的正确性。

3：（按照横扫斑竹顺序）您需要对您global memory中的dev_Hist数组进行清零（初始化），因为cudaMalloc申请的device存储空间同样不保证自动进行初始化。您可以使用runtime API的cudaMemset()函数在host端申请这一空间以后，立即予以清零；或者在您的这些统计的kernel执行前，先写一个对这段global memory清零的kernel，并加以执行。

4：（6# LZ的3：）如3：所示，横扫斑竹说的是另外的事情。

大致如此，上述内容基本上是重复说明3#横扫斑竹的内容，并无添加新的内容。

以及，上述内容仅代表在有限资料下，人肉编译和DEBUG看出的部分内容，并不保证找出了LZ的所有问题，也不保证LZ的完整的程序能够正确执行。

祝您好运~

8

谢谢版主的回答