GPUプログラミング
CUDAで、カーネルの途中で、ほかのすべてのブロックの計算が終わるのを待ってから次の段階に進みたい。一般的な方法はどれか。
- A__syncthreads()を呼ぶ
- Bカーネルをその地点で2つに分け、1つ目のカーネルが終わってから2つ目を起動する
- Cループで、ほかのブロックが書いたフラグを待ち続ける
答えと解説を見る
正解:B. カーネルをその地点で2つに分け、1つ目のカーネルが終わってから2つ目を起動する
ブロックは独立に、どの順番で実行されてもよい決まりなので、全部のブロックが同時にGPUに載っているとは限りません。待っている相手のブロックがまだ始まっていないと、フラグを待つループは永久に終わりません。そこで、同じストリームで起動したカーネルは前のカーネルが終わってから始まる性質を使い、カーネルを分けることで全体の同期を取るのが一般的です。
- A. __syncthreads()は同じブロックの中だけの同期です。
- B. 正しい。カーネルの境目が全体の同期になります。
- C. 全部のブロックが同時に載っている保証がないので、止まることがあります。