GPUプログラミング
上のようなループ(グリッドストライドループ)でカーネルを書く利点として正しいものはどれか。
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < n;
i += blockDim.x * gridDim.x)
y[i] = a * x[i] + y[i];
- Aループを使うと、メモリ合体アクセスにならなくなる
- B起動するスレッドの数が要素の数より少なくても、各スレッドが全スレッド数おきに複数の要素を処理して、すべての要素を扱える
- Cスレッドが1本しか動かなくなる
答えと解説を見る
正解:B. 起動するスレッドの数が要素の数より少なくても、各スレッドが全スレッド数おきに複数の要素を処理して、すべての要素を扱える
全スレッドの数(blockDim.x×gridDim.x)だけ番号を進めながら繰り返すので、スレッドの数がいくつでも全要素を漏れなく処理できます。GPUのSMの数に合わせたほどよい数のブロックで起動でき、スレッドの作成の手間も減ります。各回のループでは隣り合うスレッドが隣り合う要素を扱うので、メモリ合体アクセスも保たれます。
- A. 各回のループで隣のスレッドは隣の要素を読むので、合体アクセスになります。
- B. 正しい。グリッドの大きさを要素の数と切り離せます。
- C. すべてのスレッドが動きます。