条件syncthreads和死锁(或不)

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中至少有一个线程到达屏障,它就不会死锁.

所以一般来说,你需要仔细使用障碍.但具体而言,这样的(简单)早期退出模式确实有效.

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

Iteration Idx == 36:2活动warp,因此屏障退出计数为64. warp 0的所有线程都到达屏障,递增计数从0到32.来自warp 1的4个线程到达屏障,将计数从32递增到64,并且warp 0和1从屏障释放.阅读上面的链接,了解发生这种情况的原因.

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

等等...