Cuda编程直方图

Question

我想运行一个cuda程序,但我是初学者.我必须为直方图编写一个程序.但是用水桶.根据maxValue(示例中为40),该数字将添加到相应的存储桶中.如果我们有4个桶:

histo:| 1 | 10 | 30 | 39 | 32 | 2 | 4 | 5 | 1 |

0-9(第1桶)

10-19(第2桶)

20-29(第3桶)

30-39(第4桶)

我的GPU具有计算能力1.1.

我试图做一个像块一样的临时temp [],每个线程都在他的临时表上添加他的值:

__global__ void histo_kernel_optimized5( unsigned char *buffer, long size,
                               unsigned int *histo )
{
     extern __shared__ unsigned int temp[];
     temp[threadIdx.x] = 0;
     __syncthreads();

     int i = threadIdx.x + blockIdx.x * blockDim.x;
     int offset = blockDim.x * gridDim.x;
     int bucketID;
     while (i < size)
     {
              bucketID = array[i]/Bwidth;
              atomicAdd( &temp[bucketID], 1);
              i += offset;
     }
     __syncthreads();


    atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );
}

histo_kernel_optimized <<<array_size/buckets, buckets,buckets*sizeof(unsigned int)>>>(buffer,SIZE, histogram)

但编译器sais:指令'{atom,red} .shared'需要.目标sm_12或更高

我也试过为每个创建的线程都有一个临时表:

__global__ void histo_kernel_optimized5( unsigned char *buffer, long size,
                               unsigned int *histo )
{
    unsigned int temp[buckets];
     int j;
    for (j=0;j<buckets;j++){
        temp[j]=0;
    }

    int bucketID;

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int offset = blockDim.x * gridDim.x;
    while (i < size)
    {
        bucketID = array[i]/Bwidth;
        temp[bucketID]++;
        i += offset;
    }


    for (j=0;j<buckets;j++){
        histo[j] += temp[j];    
    }
 }

但编译器不允许我这样做,因为它需要一个常量来创建临时表.但问题是说,动态地为命令行提供了桶.

还有另一种方法吗？我不知道怎么做.我很迷惑.

Answer 1

使用原子时,启动较少的块将减少争用(从而提高性能),因为它不必在较少的块之间进行协调.启动更少的块,并使每个块循环遍历更多输入元素.

for (unsigned tid = blockIdx.x*blockDim.x+threadIdx.x; 
              tid < size; tid += gridDim.x*blockDim.x) {
    unsigned char value = array[tid]; // borrowing notation from another answer here
    int bin = value % buckets;
    atomicAdd(&histo[bin],1);
}