我最近发现了CUDA 5.0中的cuda-memcheck --tool racecheck工具，它包含了racecheck工具，详情请参见NVIDIA doc。该工具可以检测CUDA内核中共享内存的竞争条件。
在调试模式下，该工具不会检测到任何问题，这似乎是正常的。然而，在发布模式（-O3）下，根据问题的参数，我会得到错误信息。
以下是一个错误示例（第22行共享内存初始化，第119行赋值）：
========= ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (35, 0, 0) :
========= Write Thread (32, 0, 0) at 0x00000890 in ....h:119:void kernel_test3<float, unsigned int=4, unsigned int=32, unsigned int=64>(Data<float, unsigned int=4, unsigned int=32, unsigned int=64>*)
========= Write Thread (0, 0, 0) at 0x00000048 in ....h:22:void kernel_test3<float, unsigned int=4, unsigned int=32, unsigned int=64>(Data<float, unsigned int=4, unsigned int=32, unsigned int=64>*)
========= Current Value : 13, Incoming Value : 0
- 让我惊讶的第一件事是线程ID。当我第一次遇到错误时，每个块包含32个线程（ID从0到31）。那么为什么线程ID 32会有问题呢？我甚至在
threadIdx.x上添加了额外的检查，但这并没有改变什么。 - 我使用共享内存作为临时缓冲区，每个线程处理多维数组的自己的参数，例如
__shared__ float arr[SIZE_1][SIZE_2][NB_THREADS_PER_BLOCK]。我真的不明白为什么会有任何竞争条件，因为每个线程都处理自己的共享内存部分。 - 将网格大小从64个块减少到32个块似乎解决了这个问题（每个块有32个线程）。我不明白为什么。
为了理解发生了什么，我测试了一些更简单的内核。
让我展示一个创建这种错误的内核的例子。基本上，这个内核使用SIZE_X*SIZE_Y*NTHREADS*sizeof(float)字节的共享内存，而我可以每个SM使用48KB的共享内存。
test.cu
template <unsigned int NTHREADS>
__global__ void kernel_test()
{
const int SIZE_X = 4;
const int SIZE_Y = 4;
__shared__ float tmp[SIZE_X][SIZE_Y][NTHREADS];
for (unsigned int i = 0; i < SIZE_X; i++)
for (unsigned int j = 0; j < SIZE_Y; j++)
tmp[i][j][threadIdx.x] = threadIdx.x;
}
int main()
{
const unsigned int NTHREADS = 32;
//kernel_test<NTHREADS><<<32, NTHREADS>>>(); // ---> works fine
kernel_test<NTHREADS><<<64, NTHREADS>>>();
cudaDeviceSynchronize(); // ---> gives racecheck errors if NBLOCKS > 32
}
编译:
nvcc test.cu --ptxas-options=-v -o test
如果我们运行内核：
cuda-memcheck --tool racecheck test
kernel_test<32><<<32, 32>>>();：32个块，每个块32个线程 => 不会导致任何明显的竞争检查错误。kernel_test<32><<<64, 32>>>();：64个块，每个块32个线程 => 导致WAW危害（threadId.x= 32？！）和错误。========= ERROR: Potential WAW hazard detected at __shared__ 0x6 in block (57, 0, 0) : ========= Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void) ========= Write Thread (1, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void) ========= Current Value : 0, Incoming Value : 128 ========= INFO:(Identical data being written) Potential WAW hazard detected at __shared__ 0x0 in block (47, 0, 0) : ========= Write Thread (32, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void) ========= Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void) ========= Current Value : 0, Incoming Value : 0
那么我在这里缺少什么？我在共享内存方面做错了什么吗？（我在这方面还是个初学者）
更新：
当NBLOCKS > 32时，问题似乎来自于cudaDeviceSynchronize()。为什么会发生这种情况？
cudaDeviceSynchronize()对于某些块的原因。 - BenC