GPUプログラミング

粒子の位置x、y、zを持つデータをGPUで処理するとき、メモリ合体アクセスの点で有利な並べ方はどれか。

  1. A粒子ごとにx、y、zを並べた構造体の配列にする
  2. Bすべての値をランダムな順番に並べる
  3. Cxだけの配列、yだけの配列、zだけの配列に分ける(構造体の配列ではなく、配列の構造体)
答えと解説を見る

正解:C. xだけの配列、yだけの配列、zだけの配列に分ける(構造体の配列ではなく、配列の構造体)

構造体の配列(AoS)では、粒子0のx・y・z、粒子1のx・y・z…と並ぶので、各スレッドがxを読むと番地が12バイトずつとびとびになります。配列の構造体(SoA)にしてxだけを続けて並べると、隣り合うスレッドが隣り合う番地のxを読むので合体アクセスになります。CPUのSIMDでも同じ理由でSoAが有利なことが多いです。

  • A. 隣のスレッドのxが12バイトずつ離れてしまいます。
  • B. 合体アクセスにならなくなります。
  • C. 正しい。同じ成分を読むスレッドの番地が続きます。
この分野を解く 分野ページへ

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

  1. CUDAの共有メモリ(シェアードメモリ)の特徴として正しいものはどれか。
  2. CUDAの__syncthreads()の働きとして正しいものはどれか。
  3. CUDAで、__syncthreads()を「if (threadIdx.x < 16)」の中に書き、一部のスレッドだけが通るようにした。起こりうることとして正しいものはどれか。
  4. CUDAの共有メモリのバンク競合の説明として正しいものはどれか。
  5. 32×32のfloatの表を共有メモリに置き、ワープの32本で列の方向(縦)に読むとバンク競合が起きる。よく使われる対策はどれか。