0

我即将为 CUDA 设备编写直方图内核。它基于NVIDIA 的论文。

这个想法是每个线程计算某个部分的部分直方图（在我的情况下是卷）并将其写入共享内存块。但是，我遇到了算法的一个奇怪问题，并将内核剥离到重要部分：

__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。但是，这仅适用于前几个块（低blockIdx.x/ blockIdx.y）。其他块的值变化很大。最后一个块 ( blockIdx.y == 31) 的值为 239 或 240。

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

谁能解释这种行为？

4

1 回答 1

3

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

在这行代码之后：

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

添加这个：

partialHistograms[seqTid] = 0;

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

每个线程访问共享内存的不同部分，因此不应该存在竞争条件。

是不正确的。

于 2013-11-01T12:08:06.003 回答