synchronization - CUDA/OpenCL 中的现实死锁示例

Question

对于我正在编写的教程，我正在寻找一个“现实”且简单的死锁示例，该示例是由于对 SIMT/SIMD 的无知而导致的。

我想出了这个片段，这似乎是一个很好的例子。

我们将不胜感激。

…
int x = threadID / 2;
if (threadID > x) {
    value[threadID] = 42;
    barrier();
    }
else {
    value2[threadID/2] = 13
    barrier();
}
result = value[threadID/2] + value2[threadID/2];

我知道，它既不是正确的 CUDA C 也不是 OpenCL C。

Answer 1

一个简单的死锁实际上很容易被新手 CUDA 程序员捕捉到，当一个人试图为单个线程实现一个临界区时，它最终应该由所有线程执行。它或多或少是这样的:

__global__ kernel() {
  __shared__ int semaphore;
  semaphore=0;
  __syncthreads();
  while (true) {
    int prev=atomicCAS(&semaphore,0,1);
    if (prev==0) {
      //critical section
      semaphore=0;
      break;
    }
  }
}

atomicCAS 指令确保 exaclty 一个线程获得 0 分配给 prev，而所有其他线程获得 1。当一个线程完成其临界区时，它将信号量设置回 0，以便其他线程有机会进入临界区。

问题是，当 1 个线程获得 prev=0 时，属于同一个 SIMD 单元的 31 个线程获得值 1。在 if 语句中，CUDA 调度程序将单个线程置于暂停状态(将其屏蔽)并且让其他 31 线程继续工作。在正常情况下，这是一个很好的策略，但在这种特殊情况下，您最终会得到 1 个永远不会执行的临界区线程和 31 个等待无穷大的线程。死锁。

还要注意，break 的存在将控制流引导到 while 循环之外。如果你省略了 break 指令，并且在 if-block 之后还有一些应该由所有线程执行的操作，它实际上可能有助于调度程序避免死锁。

关于您在问题中给出的示例:在 CUDA 中，明确禁止将 __syncthreads() 放在 SIMD 发散代码中。编译器不会捕捉到它，但手册说“未定义的行为”。实际上，在 Fermi 之前的设备上，所有 __syncthreads() 都被视为相同的障碍。有了这个假设，您的代码实际上会终止而不会出现错误。不应该不依赖这种行为。