CUDA racecheck、共享内存阵列和 cudaDeviceSynchronize()

Question

我最近发现了 cuda-memcheck 的 racecheck 工具，可用于 CUDA 5.0(cuda-memcheck --tool racecheck，请参阅 NVIDIA doc)。该工具可以检测 CUDA 内核中共享内存的竞争条件。

在 Debug模式下，这个工具没有检测到任何东西，这显然是正常的。但是，在 Release模式 (-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(Data*) ========= Write Thread (0, 0, 0) at 0x00000048 in ....h:22:void kernel_test3(Data*)
========= Current Value : 13, Incoming Value : 0

1. 让我感到惊讶的第一件事是线程 ID。当我第一次遇到错误时，每个 block 包含 32 个线程(id 0 到 31)。那么为什么线程id 32会出现问题呢？我什至在 threadIdx.x 上添加了额外的检查，但这并没有改变什么。
2. 我使用共享内存作为临时缓冲区，每个线程处理自己的多维数组参数，例如__shared__ float arr[SIZE_1][SIZE_2][NB_THREADS_PER_BLOCK] .我真的不明白怎么会有竞争条件，因为每个线程都处理自己的共享内存部分。
3. 将网格大小从 64 个 block 减少到 32 个 block 似乎可以解决问题(每个 block 有 32 个线程)。我不明白为什么。

为了了解发生了什么，我测试了一些更简单的内核。让我向您展示一个导致此类错误的内核示例。基本上，这个内核使用 SIZE_X*SIZE_Y*NTHREADS*sizeof(float) B 的共享内存，每个 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 个 block ，32 个线程 => 不会导致任何明显的 racecheck 错误。
• kernel_test<32><<<64, 32>>>(); : 64 个 block ，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

那么我在这里缺少什么？我在共享内存上做错了吗？ (我还是初学者)

**更新**

问题似乎来自cudaDeviceSynchronize()什么时候NBLOCKS > 32 .为什么会这样？

Answer 1

对于初学者来说，cudaDeviceSynchronize() 不是原因；你的内核是原因，但它是一个异步调用，所以错误是在你调用 cudaDeviceSynchronize() 时捕获的。

至于内核，您的共享内存大小为 SIZE_X*SIZE_Y*NTHREADS(在示例中转换为每个 block 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-/