cuda - 条件同步线程和死锁(或不)

Question

后续问题:EarlyExit和 DroppedThreads

根据上面的链接，下面的代码应该死锁。
请解释为什么这不会死锁。 (费米上的 Cuda 5)

__device__ int add[144];
__device__ int result;

add<<<1,96>>>();  // the calling 

__global__ void add() {
 for(idx=72>>1; idx>0; idx>>=1) {
  if(thrdIdx < idx) 
   add[thrdIdx]+= add[thrdIdx+idx];
  else
   return;
  __syncthreads();
 }

 if(thrdIdx == 0)
  result= add[0];
}

Answer 1

从技术上讲，这是一个定义不明确的程序。

大多数(但不是全部)(例如 G80 不支持)，NVIDIA GPU 以这种方式支持提前退出，因为硬件为每个 block 维护一个事件线程计数，并且该计数用于屏障同步而不是初始线程计数对于 block 。

因此，当到达代码中的__syncthreads()时，硬件不会等待任何已经返回的线程，程序运行不会出现死锁。

这种风格更常见的用法是:

__global__ void foo(int n, ...) {
  int idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx >= n) return;
  ... // do some computation with remaining threads
}

重要提示:屏障计数按扭曲更新(请参阅 here )，而不是按线程更新。因此，您可能会遇到这样的情况:只有少数(或零个)线程提前返回。这意味着屏障计数不会减少。不过，只要每个 warp 中至少有一个线程到达屏障，就不会发生死锁。

所以一般来说，你需要谨慎使用屏障。但具体来说，像这样的(简单)提前退出模式确实有效。

编辑:针对您的具体情况。

迭代 Idx==36:2 个事件扭曲，因此屏障退出计数为 64。来自扭曲 0 的所有线程都到达屏障，计数从 0 增加到 32。来自扭曲 1 的 4 个线程到达屏障，计数从 32 增加到 64，并且扭曲 0 和 1 从屏障中释放。阅读上面的链接以了解为什么会发生这种情况。

迭代 Idx==18:1 个事件 warp，因此屏障退出计数为 32。来自 warp 0 的 18 个线程到达屏障，计数从 0 递增到 32。满足屏障并释放 warp 0。

等等...