GPUプログラミング

上のCUDAのカーネルで「if (i < n)」を入れている理由として正しいものはどれか。

__global__ void add(float *c, const float *a,
                    const float *b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        c[i] = a[i] + b[i];
}
  1. Aスレッドの総数が要素数より多いとき、余ったスレッドが配列の範囲の外を読み書きしないようにするため
  2. B足し算の結果がnを超えないようにするため
  3. CGPUの速さを2倍にするため
答えと解説を見る

正解:A. スレッドの総数が要素数より多いとき、余ったスレッドが配列の範囲の外を読み書きしないようにするため

スレッドはブロック単位で起動するので、要素数がブロックの大きさの倍数でないと、スレッドの総数が要素数を上回ります。番号を確かめずに書き込むと、確保した配列の外のメモリを壊したり、不正なアクセスでエラーになったりします。「i < n」の確認は、GPUのカーネルを書くときのほぼ決まった形です。

  • A. 正しい。ブロックの数を切り上げたために生じる余りに対処します。
  • B. 値の大きさではなく、番号の範囲を確かめています。
  • C. 速さのためではなく、正しさのためです。
この分野を解く 分野ページへ

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

  1. NVIDIAのGPUで、32本のスレッドを1つの束にして、同じ命令を一斉に実行する単位を何と呼ぶか。
  2. GPUのSIMT(単一命令・複数スレッド)の方式と、CPUのSIMD命令の違いとして正しいものはどれか。
  3. 1つのワープの中で、条件分岐によってスレッドの半分がif側、残りの半分がelse側に分かれた(ワープダイバージェンス)。一般にどう実行されるか。
  4. 上の(1)と(2)は、どちらもスレッドの半分にA、残りにBを実行させる。ワープの効率の点で正しい説明はどれか(1ワープ32スレッドとする)。
  5. GPUのメモリ合体アクセス(コアレッシング)の説明として正しいものはどれか。