常量 for 循环产生错误的结果

Question

我正准备为 CUDA 设备编写直方图内核。它基于 NVIDIA's paper .

这个想法是每个线程计算某个部分(在我的例子中是体积)的部分直方图并将其写入共享内存块。然而，我遇到了一个奇怪的算法问题，并将内核剥离到重要部分:

__global__ void calcHist64()
{
    extern __shared__ unsigned char partialHistograms[];

    //a unique sequential thread id within this block, used to determine the memory in which to write the partial histogram
    unsigned int seqTid = threadIdx.x + threadIdx.y * blockDim.x;       
#pragma unroll
    for(int i = 0; i < 255; ++i)
    {
        //increment the thread's partial histogram value
        partialHistograms[seqTid]++; 
    }
    //each partial histogram should now be 255
    //Output the value for every thread in a certain block
    if(blockIdx.x == 0 && blockIdx.y == 31)
        printf("Partial[%i][%i]: %i\n", threadIdx.x, threadIdx.y, partialHistograms[partialHistRoot]);  
}

内核通过以下方式调用:

int sharedMemory = 4096;
dim blocks(32, 32, 1);
dim3 threadsPerBlock(8,8,1);
calcHist64<<<blocks, threadsPerBlock, sharedMemory>>>();

我希望每个部分直方图的值为 255。但是，这仅适用于前几个 block (低 blockIdx.x/blockIdx.y) .其他 block 的值差异很大。最后一个 block (blockIdx.y == 31) 的值为 239 或 240。

我无法解释这种行为。它是一个常量 for 循环，毕竟它运行了 255 次。每个线程访问共享内存的不同部分，因此不应存在竞争条件。

谁能解释这种行为？

Answer 1

您没有初始化共享内存。它不会自动为您初始化为零。

在这行代码之后:

unsigned int seqTid = threadIdx.x + threadIdx.y * blockDim.x;  

添加这个:

partialHistograms[seqTid] = 0;

此外，在您的代码中，您没有定义 partialHistRoot。我假设 partialHistRoot == seqTid。如果不是这种情况，那么你有一个竞争条件和你的声明

Each thread accesses different parts of shared memory so there should be no race conditions.

不正确。