忘年会から朝帰りしたCUDA日記2

昨日書いていた「2の累乗でない要素数Nの配列のparallel reduction」について。ちなみにやりたい内容はリストの合計を取るだけなので素朴な方法として「2の累乗に足りない部分を0で埋める」が考えられる。あとif文で分岐する方法と、素因数分解して約数でreductionする方法も考えられる。Nが2とか3しか素因数に持たない場合は3番目の方法でいいと思うが、7とかある場合にどうするのが一番早いのか。

を考えたいけどやっぱり眠いので少し寝る。カフェインを摂取したので30〜60分後にタイマーをセットすればちょうど血中カフェイン濃度が最高になっているはずだ。

• • • • -

CUDAのコードがどういうインストラクションの列になって、どういう風に処理されていくか見れるといいんだけどな。

ドキュメントを読んだ感じ、下のコードでAとBが全く別の処理だとCPUなら1秒で済むけどGPUだと2秒かかるのだと思っている。試そう。

if(x & 1){ // 奇数のとき
  // 1秒かかる処理A
}else{
  // 1秒かかる処理B
}

• • • • -

確認のためのコードを書いた。スレッドは64個。ブロックが1個。差分は1行なのでdiffスタイルで書いた。

const int LOOP_SIZE = 1000000;
__global__ static void kernel(const int* input, int* output, clock_t * timer) {
    const int x = threadIdx.x;
    const int bid = blockIdx.x;
    if (x == 0) timer[bid] = clock();

    int v = 0;

-    if(x & 1){
+    if(x < 32){
      for(int i=0; i < LOOP_SIZE; i++){
        v += i;
      }
    }else{
      for(int i=LOOP_SIZE; i > 0; i--){
        v -= i;
      }
    }
    output[x] = v;
    __syncthreads();
    if (x == 0) timer[bid+gridDim.x] = clock();
}

(x & 1)の結果
time = 148000468
time = 148000468
time = 148000468

(x < 32)の結果
time = 86000420
time = 86000404
time = 86000420

倍近い性能差がでるという予想はあっていたけど、148と86がきれいに2倍になっていない理由がよくわからない。なぜだろう。