GPUプログラミング

上の(1)と(2)は、どちらもスレッドの半分にA、残りにBを実行させる。ワープの効率の点で正しい説明はどれか(1ワープ32スレッドとする)。

if (threadIdx.x % 2 == 0) A(); else B();   // (1)
if ((threadIdx.x / 32) % 2 == 0) A(); else B();   // (2)
  1. A(1)の方が、分岐がそろうので効率がよい
  2. B(1)はすべてのワープの中で分岐が分かれるが、(2)はワープごとにそろうので、(2)の方が効率がよい
  3. Cどちらも効率はまったく同じである
答えと解説を見る

正解:B. (1)はすべてのワープの中で分岐が分かれるが、(2)はワープごとにそろうので、(2)の方が効率がよい

(1)では偶数と奇数のスレッドで分かれるので、どのワープの中にもA側とB側のスレッドが混ざり、すべてのワープがAとBを順に実行します。(2)ではスレッドの番号を32で割った値で分けるので、1つのワープの32本はすべて同じ側になり、ワープごとにAかBの一方だけを実行します。同じ仕事でも、条件をワープの単位でそろえるとダイバージェンスがなくなります。

  • A. (1)は隣り合うスレッドで条件が変わります。
  • B. 正しい。分岐の条件をワープの単位でそろえるのが工夫です。
  • C. ダイバージェンスの有無で大きく違います。
この分野を解く 分野ページへ

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

  1. GPUのメモリ合体アクセス(コアレッシング)の説明として正しいものはどれか。
  2. 各スレッドの番号をiとして、上のように32要素おきに配列を読むと、a[i]を読む場合に比べて遅くなる。主な理由はどれか。
  3. 粒子の位置x、y、zを持つデータをGPUで処理するとき、メモリ合体アクセスの点で有利な並べ方はどれか。
  4. CUDAの共有メモリ(シェアードメモリ)の特徴として正しいものはどれか。
  5. CUDAの__syncthreads()の働きとして正しいものはどれか。