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 - 线程继续他们的工作.在正常情况下,这是一个很好的策略,但在这种特殊情况下,你最终会得到一个永不执行的关键部分线程和31个等待无穷大的线程.僵局.

还要注意,其存在break导致控制流在while循环之外.如果你省略break指令并在if-block之后有一些应该由所有线程执行的操作,那么它实际上可以帮助调度程序避免死锁.

关于你在问题中给出的例子:在CUDA中,明确禁止__syncthreads()输入SIMD分歧代码.编译器不会捕获它,但手册中说的是"未定义的行为".实际上,在前费米器件上,所有器件都__syncthreads()被视为相同的障碍.有了这个假设,你的代码实际上会终止而不会出错.每个人都应该不依赖于这种行为虽然.