CUDA/OpenCL 中的现实死锁示例答案

【问题标题】：Realistic deadlock example in CUDA/OpenCLCUDA/OpenCL 中的现实死锁示例
【发布时间】：2011-09-19 14:04:12
【问题描述】：

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

我想出了这个sn-p，这似乎是一个很好的例子。

我们将不胜感激。

…
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。

【问题讨论】：

作为一个例子似乎太复杂了，而对于一个“现实”的例子来说却相当简单。我会在条件中只使用get_local_id(0) > constant，并将“业务代码”（赋值）替换为cmets /* do some stuff */ 和/* do another stuff */。不过，我认为 StackOverflow 不是讨论的最佳场所，它是提问和回答的场所。

标签： synchronization cuda parallel-processing opencl simd

【解决方案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 指令确保一个线程获得 0 分配给 prev，而所有其他线程获得 1。当一个线程完成其临界区时，它将信号量设置回 0，以便其他线程有机会进入关键部分。

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

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

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

【讨论】：