我最近发现了CUDA 5.0 中可用的cuda-memcheck竞赛检查工具（参见NVIDIA 文档）。该工具可以检测 CUDA 内核中共享内存的竞争条件。cuda-memcheck --tool racecheck
在调试模式下，这个工具没有检测到任何东西，这显然是正常的。但是，在发布模式 (-O3) 中，我会根据问题的参数得到错误。
这是一个错误示例（第 22 行共享内存的初始化，第 119 行的赋值）：
========= 错误：在块 (35, 0, 0) 中的共享0x0处检测到潜在的 WAW 危害：========== 在 0x00000890 处写入线程 (32, 0, 0)。 ...h:119:void kernel_test3(Data*) ========= 在 ....h:22:void kernel_test3(Data*)
==的 0x00000048 处写入线程 (0, 0, 0) ======= 当前值：13，传入值：0
- 让我惊讶的第一件事是线程ID。当我第一次遇到错误时，每个块包含 32 个线程（ids 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)B 的共享内存，我每个 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？！）和错误。
========= 错误：在块 (57, 0, 0) 中的共享0x6处检测到潜在的 WAW 危害：
========== 在 0x00000048 处写入线程 (0, 0, 0)。 ...h:403:void kernel_test(void)
========= 在 ....h:403:void kernel_test(void)
====的 0x00000048 处写入线程 (1, 0, 0) ===== 当前值：0，传入值：128========= 信息：（正在写入相同的数据）在块（47、0、0）中的共享0x0 处检测到潜在的 WAW 危害：
========= 写入线程（32、0、 0) at 0x00000048 in ....h:403:void kernel_test(void)
========= 在 ....h:403:void kernel_test( void)
========= 当前值：0，传入值：0
那么我在这里错过了什么？我在共享内存方面做错了吗？（我仍然是这个的初学者）
** 更新 **
问题似乎来自cudaDeviceSynchronize()when NBLOCKS > 32。为什么会这样？