条件同步线程和死锁（或不）

Question

跟进 Q 至：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 支持提前退出这种方式，因为硬件为每个块维护一个活动线程计数，这个计数用于屏障同步而不是初始线程计数为块。

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

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

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

重要提示：屏障计数是按 Warp 更新的（请参阅 here），而不是按线程更新。因此，您可能会遇到这样的情况，例如，只有少数（或零个）线程提前返回。这意味着屏障计数不会减少。但是，只要每个 warp 中至少有一个线程到达屏障，它就不会死锁。

因此，一般来说，您需要谨慎使用障碍。但具体来说，像这样的（简单）提前退出模式确实有效。

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

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

Iteration Idx==18: 1 个活动 warp，因此屏障退出计数为 32。来自 warp 0 的 18 个线程到达屏障，计数从 0 增加到 32。满足屏障并释放 warp 0。

等等……