CUDA/OpenCL 中的现实死锁示例

对于我正在编写的教程，我正在寻找一个“现实”且简单的示例，说明由于不了解 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。

新手 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;
    }
  }
}

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

问题是，当 1 个线程获取 prev=0 时，属于同一 SIMD 单元的 31 个线程获取值 1。在 if 语句处，CUDA 调度程序将该单个线程置于保留状态（将其屏蔽），并让其他 31 个线程处于等待状态。 - 线程继续工作。在正常情况下，这是一个很好的策略，但在这种特殊情况下，您最终会得到 1 个从未执行的临界区线程和 31 个无限等待的线程。僵局。

另请注意，存在break这导致控制流外部while环形。如果省略break指令并在if块之后添加一些应该由所有线程执行的操作，它实际上可以帮助调度程序避免死锁。

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