histogram - 如何制作CUDA直方图内核？

Question

我正在为图片上的直方图编写一个 CUDA 内核，但我不知道如何从内核返回一个数组，并且当其他线程读取它时数组会发生变化。有什么可能的解决方案吗？

__global__ void Hist(
    TColor *dst, //input image
    int imageW,
    int imageH,
 int*data
){
    const int ix = blockDim.x * blockIdx.x + threadIdx.x;
    const int iy = blockDim.y * blockIdx.y + threadIdx.y;

if(ix < imageW && iy < imageH)
{
  int pixel = get_red(dst[imageW * (iy) + (ix)]);
                  //this assign specific RED value of image to pixel

  data[pixel] ++; // ?? problem statement ...
 }
}

@para d_dst:输入图像 TColor 等于 float4。

@para data:直方图大小[255]的数组

extern "C" void
cuda_Hist(TColor *d_dst, int imageW, int imageH,int* data) 
{
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(iDivUp(imageW, BLOCKDIM_X), iDivUp(imageH, BLOCKDIM_Y));
  Hist<<<grid, threads>>>(d_dst, imageW, imageH, data);
}

Answer 1

您看过 SDK 示例了吗？ “直方图”样本在 CUDA SDK 中可用(目前 NVIDIA 上的版本为 3.0 developer site，版本 3.1 beta 可用于 registered developers)。

示例文档很好地解释了如何处理求和，要么在 GPU 上使用全局内存原子，要么分别收集每个 block 的结果，然后单独进行归约(在主机或 GPU 上)。