GPUプログラミング

各スレッドの番号をiとして、上のように32要素おきに配列を読むと、a[i]を読む場合に比べて遅くなる。主な理由はどれか。

float x = a[i * 32];
  1. A隣のスレッドの番地が128バイト離れるので合体アクセスにならず、スレッドごとに別のメモリのまとまりを読むことになるから
  2. B掛け算の計算に時間がかかるから
  3. C32要素おきに読むと、値が正しく読めないから
答えと解説を見る

正解:A. 隣のスレッドの番地が128バイト離れるので合体アクセスにならず、スレッドごとに別のメモリのまとまりを読むことになるから

floatは4バイトなので、a[i*32]では隣のスレッドの番地が128バイト離れます。ワープの32本は32個の別々のまとまりを読み、それぞれのまとまりから4バイトしか使いません。a[i]なら32本で128バイトの1つのまとまりで済むので、運ぶデータの量に何十倍もの差が出ます。2次元の配列では、スレッドの番号を内側(続いた番地)の添字に対応させるのが基本です。

  • A. 正しい。読んだまとまりの大部分が使われません。
  • B. 掛け算の時間は、メモリの待ちに比べて小さいです。
  • C. 値は正しく読めますが、遅くなります。
この分野を解く 分野ページへ

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

  1. 粒子の位置x、y、zを持つデータをGPUで処理するとき、メモリ合体アクセスの点で有利な並べ方はどれか。
  2. CUDAの共有メモリ(シェアードメモリ)の特徴として正しいものはどれか。
  3. CUDAの__syncthreads()の働きとして正しいものはどれか。
  4. CUDAで、__syncthreads()を「if (threadIdx.x < 16)」の中に書き、一部のスレッドだけが通るようにした。起こりうることとして正しいものはどれか。
  5. CUDAの共有メモリのバンク競合の説明として正しいものはどれか。