为什么 CUDA 同步点不能防止竞争条件？

Question

我们运行我们的代码 cuda-memcheck --tool racecheck <executable>。我们收到以下内存危险错误。

========= Race reported between Read access at 0x00004098 CUDA.cu:123:KernelFunction()
=========     and Write access at 0x00005058 in CUDA.cu:146:KernelFunction() [529996 hazards]

这是代码。它声称第 123 行 value = sharedMemory0[sharedMemoryIndex]; 与第 146 行 sharedMemory0[sharedIndex0] = sharedMemory1[sharedIndex1]; 处于竞争状态。我们有

// Synchronization Point 1 
__syncthreads(); 
__threadfence_block();

两行之间。难道所有的线程不应该在那个时候同步并且所有以前的内存 read/writes 在那个时候完成吗？在开始第二个 j 循环之前，所有线程和内存访问都应该在第一个 j 循环之后完成。所以在我们看来，同步点 1 应该隔离两个 j 循环并防止竞争条件，但该工具说这不是真的。

为什么该工具报告竞争条件？关于我们可以做些什么来防止它的任何见解？

我们还看到了对一种工具的引用，该工具可能能够报告执行轨迹以更轻松地查看竞争条件。我们可以使用什么工具和选项来获取跟踪以更清楚地了解竞争条件存在的原因？

   for (i = 0; i < COUNT0; i++) {
       // Synchronization Point 0
       __syncthreads();
       __threadfence_block();
       for (j = 0; j < COUNT1; j++) {
          index = j*blockDim.x + threadIdx.x;
          if (index < THREAD_COUNT0) {
             for (k = 0; k < COUNT2; k++)
                sharedMemoryIndex = function0(index);
                value = sharedMemory0[sharedMemoryIndex];
             }
          }         
       }

       // Synchronization Point 1
       __syncthreads();
       __threadfence_block();
       for (j = 0; j < COUNT2; j++) {
          index = j*blockDim.x + threadIdx.x;
          if (index < THREAD_COUNT1) {
            sharedIndex0 = function1(index);
            sharedIndex1 = function2(index);
            sharedMemory0[sharedIndex0] = sharedMemory1[sharedIndex1];
          }
       }
    }

我们还运行 Synchcheck 工具，cuda-memcheck --tool synccheck <executable> 它在同步点 1 上报告了以下错误。这两个错误之间可能有很强的相关性，但不是很相关cuda-memcheck 指南中有很多关于不同代码的同步是什么、它为什么不好以及如何修复它的文档。

有意见吗？

========= Barrier error detected. Encountered barrier with divergent threads in block
=========     at 0x00004ad8 in CUDA.cu:139:KernelFunction()
=========     by thread (0,0,0) in block (8,0,0)

Answer 1

没有足够的数据来准确查明您的问题。然而，最后的错误信息非常关键：

Barrier error detected. Encountered barrier with divergent threads in block

似乎您的一个线程在一个块中达到了某种障碍而另一个没有，因为它在一个未被采用的分支中。请注意，发散分支不仅出现在 if 条件下，而且如果它们的循环条件在块中的线程之间不同，也会出现在循环中。

当某些线程错过 __syncthreads() 时，可能会发生奇怪的事情。在实践中，这通常意味着那些线程停止在 不同的 __syncthreads() 并且系统认为一切都同步，而实际上并非如此。这可能会导致您描述的赛车情况。

所以 - 找到你的分歧 __syncthreads() - 这很可能是你遇到问题的原因。问题可能出在之前您包含的代码段。

还有：

是i局部变量（非共享）吗？
对于一个块中的所有线程来说 COUNT0 是否相同？

Answer 2

无论用于执行计算的线程数如何，此代码都会给出相同的结果。我们运行只有单线程的代码然后运行多线程的代码。根据定义，单线程版本不可能运行进入竞争状态。然而，单线程版本给出了与多线程版本相同的结果。 cuda-memcheck --tool racecheck 报告了多线程版本上的许多竞争违规。如果实际发生了竞争冲突，多线程结果将不会与单线程结果完全匹配。因此，cuda-memcheck 一定是错误的，并且有处理复杂循环结构的错误。 cuda-memcheck 能够在简单的循环结构中找到竞争条件，只是在这个复杂的结构中找不到。

为什么 CUDA 同步点不能防止竞争条件？

Why doesn't CUDA synchronization point prevent race condition?

c++

cuda

race-condition