GPUプログラミング

CUDAで、カーネルの途中で、ほかのすべてのブロックの計算が終わるのを待ってから次の段階に進みたい。一般的な方法はどれか。

  1. A__syncthreads()を呼ぶ
  2. Bカーネルをその地点で2つに分け、1つ目のカーネルが終わってから2つ目を起動する
  3. Cループで、ほかのブロックが書いたフラグを待ち続ける
答えと解説を見る

正解:B. カーネルをその地点で2つに分け、1つ目のカーネルが終わってから2つ目を起動する

ブロックは独立に、どの順番で実行されてもよい決まりなので、全部のブロックが同時にGPUに載っているとは限りません。待っている相手のブロックがまだ始まっていないと、フラグを待つループは永久に終わりません。そこで、同じストリームで起動したカーネルは前のカーネルが終わってから始まる性質を使い、カーネルを分けることで全体の同期を取るのが一般的です。

  • A. __syncthreads()は同じブロックの中だけの同期です。
  • B. 正しい。カーネルの境目が全体の同期になります。
  • C. 全部のブロックが同時に載っている保証がないので、止まることがあります。
この分野を解く 分野ページへ

GPUプログラミングのほかの問題

  1. GPUで画素の値のヒストグラムを作るとき、全スレッドがグローバルメモリの256個のビンにatomicAddで直接足すと遅いことがある。改善の方法として正しいものはどれか。
  2. GPUで多数の値の総和を求めるとき、木の形のリダクション(隣どうしを足して半分にすることを繰り返す)を使う理由として正しいものはどれか。
  3. 1024個の値を木の形のリダクションで合計するとき、必要な段の数はいくつか。
  4. 単体のGPUを使うプログラムで、CPU側のメモリとGPU側のメモリについて一般的に正しいものはどれか。
  5. 配列の要素を1回だけ足し算するような軽い処理のためにデータをGPUへ送ると、CPUで計算するより遅くなることがある。主な理由はどれか。