0

我想运行一个 cuda 程序，但我是初学者。我必须为直方图编写一个程序。但是有桶。根据 maxValue（示例中的 40），该数字将被添加到相应的存储桶中。如果我们有 4 个桶：

历史：| 1 | 10 | 30 | 39 | 32 | 2 | 4 | 5 | 1 |

0-9（第一个桶）

10-19（第二桶）

20-29（第三桶）

30- 39（第四桶）

我的 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)

但是编译器说：指令 '{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];    
    }
 }

但是编译器不允许我这样做，因为它需要一个常量来创建临时表。但问题是，桶是为命令行动态给出的。

还有另一种方法吗？我不知道该怎么做。我很困惑。

4

3 回答 3

8

使用原子时，启动较少的块将减少争用（并因此提高性能），因为它不必在较少的块之间进行协调。启动更少的块，并让每个块在更多的输入元素上循环。

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);
}
于 2013-04-04T16:06:58.273 回答
4

使用原子操作很容易实现直方图。我不知道您为什么要编写如此复杂的内核。并行化操作的动机是利用算法的并行性。无需遍历内核内的整个直方图。这是一个示例 CUDA 内核和包装函数，用于计算具有指定数量的 bin 的数组的直方图。我认为它不能针对 Compute 1.1 设备进行进一步优化。但是对于 Compute 1.2，可以使用共享内存。

__global__ void kernel_getHist(unsigned char* array, long size, unsigned int* histo, int buckets)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;

    if(tid>=size)   return;

    unsigned char value = array[tid];

    int bin = value % buckets;

    atomicAdd(&histo[bin],1);
}

void getHist(unsigned char* array, long size, unsigned int* histo,int buckets)
{
    unsigned char* dArray;
    cudaMalloc(&dArray,size);
    cudaMemcpy(dArray,array,size,cudaMemcpyHostToDevice);

    unsigned int* dHist;
    cudaMalloc(&dHist,buckets * sizeof(int));
    cudaMemset(dHist,0,buckets * sizeof(int));

    dim3 block(32);
    dim3 grid((size + block.x - 1)/block.x);

    kernel_getHist<<<grid,block>>>(dArray,size,dHist,buckets);

    cudaMemcpy(histo,dHist,buckets * sizeof(int),cudaMemcpyDeviceToHost);

    cudaFree(dArray);
    cudaFree(dHist);
}
于 2013-04-03T11:08:06.137 回答
0

有一种针对没有原子操作的设备的解决方案，并展示了一种将片上内存冲突最小化的方法，Podlozhnyuk在 CUDA 的直方图计算中提出了对扭曲的细分

代码位于 CUDASamples\3_Imaging\histogram （来自 CUDA Samples）

于 2016-05-29T16:25:54.090 回答