CUDA日記(Bank Conflictの実験)

計測したい処理を1メガ回繰り返し計算しているので消費クロック数が86000420みたいになる。だから端数は気にせずに86Mと表記することにした。

Bank Conflictが起きるのを確認したくてこういうコードにした。スレッド数は4。
v += i;を繰り返しているあたりのコードから、for文自体のコストは60Mだとわかるので、v += smem[x * 2];のコストが16M、smem[x] += smem[0];のコストが36Mと倍以上であることがわかる。

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;
    __shared__ int smem[16];

    if (x == 0) timer[bid] = clock();
    int v = 0;
    for(int i=0; i < LOOP_SIZE; i++){
      //smem[x] += smem[0]; // 96M
      smem[x * 2] += smem[0]; // 80M
      //smem[x] += smem[x + 4]; // 88M
      //smem[x * 2] += smem[x * 2 + 1]; // 80M 
      //v += smem[x * 2]; // 76M
      //v += smem[x]; // 80M
      //v += i; // 80M
      //v += i; v += i; // 100M
      //v += i; v += i; v += i;// 120M
      //v += i + smem[x * 2]; // 80M
    }
    output[x] = v; //smem[x];
    __syncthreads();
    if (x == 0) timer[bid+gridDim.x] = clock();
}

sharedメモリーがBankConflictさえ起こさなければレジスタ並みに速いって書いてあったけど、たしかにスレッドが4つしかない以上iはレジスタに載っているはず。そのiをお足すのより速いとか...あ、違うか。これsmem[x * 2]がループの間ずっと定数だからコンパイラが賢くループの外で取得して定数を足すコードに書き換えているのか。for文の中身を空にしたらfor文ごと削除してくれたりするし、定数に置き換えるくらいのことはやってそうだ。

      //v += smem[x * 2]; // 76M
      //v += smem[x]; // 80M
      //v += i; // 80M

メモリアクセスの方法を変えて節約できるコストはv += i;の5分の1程度だから、reductionや配列を単純にn倍するような「計算の密度がめちゃくちゃ低い処理」でもない限り気にしないでいいんじゃないかと思った。