6

跟进 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];
}
4

1 回答 1

9

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

大多数但不是全部（例如 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 中至少有一个线程到达屏障，它就不会死锁。

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

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

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

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

ETC...

于 2013-03-01T02:41:28.143 回答