我最近发现了cuda memcheck的racecoheck工具，该工具在cuda 5.0中可用(cuda-memcheck --tool racecheck，请参阅NVIDIA文档)。该工具可以检测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的共享内存
测试.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
那么我在这里缺少了什么？我在共享记忆方面做错了什么吗？(我还是个初学者)
更新：
问题似乎来自CCD_ 9而不是CCD_ 10。为什么会发生这种情况？
对于初学者来说，cudaDeviceSynchronize()不是原因；内核是原因，但它是一个异步调用，因此在调用cudaDeviceSynchronize()时会发现错误。
至于内核，您的共享内存大小为size_X*size_Y*NTHREADS(在本例中转换为每个块512个元素)。在嵌套循环中，使用[i*blockDim.x*SIZE_Y+j*blockDim.x+threadIdx.x]对其进行索引——这就是问题所在。
更具体地说，您的i和j值的范围为[0，4)，您的threadIdx.x的范围为[0]，32)，并且您的SIZE_{x|Y}值为4。当blockDim.x为64时，循环中使用的最大索引将为991(从3*64*4+3*64+31)。当您的blockDim.x为32时，您的最大索引将为511。
根据您的代码，每当您的NBLOCKS超过NTHREADS 时，您都会出现错误
注意：我最初将此发布到https://devtalk.nvidia.com/default/topic/527292/cuda-programming-and-performance/cuda-racecheck-shared-memory-array-and-cudadevicesynchronize-/
这显然是NVIDIA Linux驱动程序中的一个错误。该错误在313.18发布后消失。