我正准备为 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<<>>();
我希望每个部分直方图的值为 255。但是,这仅适用于前几个 block (低 blockIdx.x
/blockIdx.y
) .其他 block 的值差异很大。最后一个 block (blockIdx.y == 31
) 的值为 239 或 240。
我无法解释这种行为。它是一个常量 for 循环,毕竟它运行了 255 次。每个线程访问共享内存的不同部分,因此不应存在竞争条件。
谁能解释这种行为?
您没有初始化共享内存。它不会自动为您初始化为零。
在这行代码之后:
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.
Không đúng.
Tôi là một lập trình viên xuất sắc, rất giỏi!