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。

- Framester

---

这个例子看起来过于复杂，不够简单易懂。我会在条件语句中只使用 get_local_id(0) > constant，并用注释 /* do some stuff */ 和 /* do another stuff */ 替换 "业务代码"（赋值）。 尽管如此，我认为 StackOverflow 不是讨论的最佳场所，它更适合提问和回答。 - Radim Vansa

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指令确保只有一个线程被分配0，其他所有线程都被分配1。当这个线程完成其临界区时，它将信号量设置回0，以便其他线程有机会进入临界区。

问题是，虽然一个线程得到了prev=0，但属于同一SIMD单元的31个线程却得到了1的值。在if语句中，CUDA调度程序使单个线程暂停（屏蔽掉），让其他31个线程继续执行它们的工作。通常情况下，这是一个不错的策略，但在这种特殊情况下，你最终得到了一个永远不会被执行的1个临界区线程和31个无限等待的线程。死锁。

还要注意break存在，它使控制流从while循环之外跳出。如果你省略了break指令，并且在if块之后有一些需要所有线程执行的其他操作，它实际上可能帮助调度程序避免死锁。

关于你在问题中给出的示例：在CUDA中，禁止将__syncthreads()放在分歧代码中。编译器不会捕捉到错误，但手册中提到了“未定义的行为”。在实践中，在Fermi之前的设备上，所有__syncthreads()都被视为相同的屏障。有了这个假设，你的代码实际上会在没有错误的情况下终止。但是不应该依赖于这种行为。

- CygnusX1