いくつかの質問がまとめられているので、個別に解決しようと思います。
SM ごとに 1 ブロック
nVidia 自身のフォーラムでしばらく前にこれを尋ねたところ、これが起こっていないことを示す結果が得られました。どうやら、ブロックの数が SM の数と等しい場合、ブロック スケジューラは SM ごとにブロックを割り当てません。
暗黙の同期
いいえ。まず第一に、各ブロックに独自の SM があることを保証することはできません (上記を参照)。次に、すべてのブロックが同時にグローバル ストアにアクセスすることはできません。それらが同期的に実行されると、最初のメモリ読み取り/書き込みの時点でこの同期性が失われます。
ブロック同期
良いニュースがあります。はい、できます。CUDA C プログラミング ガイドのセクション B.11 で説明されているアトミック命令を使用して、バリアを作成できます。NGPU で同時に実行されるブロックがあるとします。
__device__ int barrier = N;
__global__ void mykernel ( ) {
/* Do whatever it is that this block does. */
...
/* Make sure all threads in this block are actually here. */
__syncthreads();
/* Once we're done, decrease the value of the barrier. */
if ( threadIdx.x == 0 )
atomicSub( &barrier , 1 );
/* Now wait for the barrier to be zero. */
if ( threadIdx.x == 0 )
while ( atomicCAS( &barrier , 0 , 0 ) != 0 );
/* Make sure everybody has waited for the barrier. */
__syncthreads();
/* Carry on with whatever else you wanted to do. */
...
}
命令はアトミックにatomicSub(p,i)計算*p -= iされ、ブロック内の 0 番目のスレッドによってのみ呼び出されます。つまり、デクリメントをbarrier1 回だけ行います。この命令atomicCAS(p,c,v)は*p = viffを設定*p == cし、 の古い値を返します*p。barrierこの部分は、到達するまで0、つまりすべてのブロックが通過するまでループします。
__synchtreads()ブロック内のスレッドは厳密なロックステップで実行されないため、この部分を の呼び出しでラップする必要があり、すべてのスレッドが 0 番目のスレッドを待機するように強制する必要があることに注意してください。
カーネルを複数回呼び出す場合は、 にbarrier戻す必要があることを覚えておいてくださいN。
アップデート
jHackTheRipperの回答とCicadaのコメントへの返信として、GPU で同時にスケジュールできるよりも多くのブロックを開始しようとするべきではないことを指摘する必要がありました。これは多くの要因によって制限されるため、CUDA Occupancy Calculatorを使用して、カーネルとデバイスの最大ブロック数を見つける必要があります。
ただし、元の質問から判断すると、SM の数だけブロックが開始されているため、この点は意味がありません。