CUDAカーネルでブロック同期を実行する可能性のある方法

Question

最近、カーネルが起動する前にグローバルメモリの初期化が必要なコードを作成しています。私は各カーネルが起動する前にcudaMemsetをやっていました。しかし、このカーネルを何千回も呼び出す必要がある場合、オーバーヘッドは無視できません。だから私は最終的にすべての初期化作業が完了したかどうかを判断するためにグローバルメモリを使用するという考え方を思いつきます。しかし、私はすぐにアクティブなブロック内のいくつかのスレッドがループを行っているとき、次のブロックは起動し続けず、結果としてデッドループになります。

int i = blockIdx.x * blockDim.x + threadIdx.x; 
if (i < (n + n)) { 
     data[i] = 0; 
}//working. 
__syncthreads();//sync 
if (threadIdx.x == 0) { 
     atomicAdd((unsigned *)&flag, 1);//voting  
     while (flag < gridDim.x); //waiting 
    } 
} 
__syncthreads(); 
//do something with data 

手動で現在のブロックをスリープ状態にしてカーネルを起動させる方法はありますか？または私の初期化の問題のためのより良い解決策はありますか？

Answer 1

CUDAでブロック同期を試みるべきではありません。これにより、後のブロックの起動（以前のブロックはリソースをあきらめないため）と同期ポイントのデッドロックを防ぐことができます。

作業が完了するまでブロックをスリープ状態にするのではなく、を実行して、現在実行中のブロックに作業を移動してください。Programming Guideには、カーネルの最後のブロックで追加作業を行うためのworked example at the end of it's memory fence sectionがあります。これを使用して、次のブロックのグローバルメモリ変数を準備することができます。

追加のcudaMemcpy（）または追加のカーネル起動を実行する必要がないという利点は、ブロックごとの余分なアトミックメモリアクセスと各ブロック内の同期化とを比較する必要があります。だから、ある時点でグリッドあたりのブロック数が増えると、余分なcudaMemcpy（）を実行するほうが安くなります。