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
}

重要な注意: バリア カウントは、スレッドごとではなく、ワープごとに更新されます (こちらを参照)。そのため、たとえば、少数 (またはゼロ) のスレッドのみが早期に返される場合があります。これは、バリア カウントがデクリメントされないことを意味します。ただし、各ワープから少なくとも 1 つのスレッドがバリアに到達する限り、デッドロックは発生しません。

したがって、一般に、バリアは慎重に使用する必要があります。しかし具体的には、このような (単純な) 早期終了パターンは機能します。

編集：あなたの特定のケースのために。

反復 Idx==36: 2 つのアクティブなワープがあるため、バリアの終了カウントは 64 です。ワープ 0 からのすべてのスレッドがバリアに到達し、カウントが 0 から 32 に増加します。ワープ 1 からの 4 つのスレッドがバリアに到達し、カウントが 32 から 64 に増加し、ワープ 0 と1体が結界から解放される。上記のリンクを読んで、なぜこれが起こるのかを理解してください。

反復 Idx==18: 1 つのアクティブなワープなので、バリア出口カウントは 32 です。ワープ 0 からの 18 スレッドがバリアに到達し、カウントが 0 から 32 に増加します。バリアが満たされ、ワー​​プ 0 が解放されます。

等...

于 2013-03-01T02:41:28.143 に答える