CUDA日記(2の累乗でない配列のparallel reduction実験)

要素数が72個の配列の合計を取る計算は128まで0で埋めてしまった上で普通の2つずつのreductionをする方が速いのか、それとも素因数分解して一部3とかでreductionした方が速いのか、という話。結論から言うと最初に3で割るのが128まで埋めるタイプより1割程度速かった。

42個の場合も試してみたが580が551に減るだけの5%程度の高速化にしかならなかった。42の場合、素因数に7が入っているから微妙なんだな。

いま見直していて気づいたんだけども、zeroで埋めるのを128までではなくて、いっそ72 * 2の144個使うことにすれば__syncthreads()がいらなくなるな。あとsmem[x] = x;のあとに__suncthreads()が必要だな。

#define N 72
#define NUM_BLOCKS 1
#define CEIL 128
#define TYPE 0

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[CEIL];

    if (x == 0) timer[bid] = clock();
    for(int i=0; i < LOOP_SIZE; i++){
#if TYPE == 0
      // 684M
      // zero reset
      smem[CEIL - N + x] = 0;
      __syncthreads();

      // target values
      smem[x] = x;

      // reduction
      if(x < 64){
      	smem[x] += smem[x + 64];
      }
      __syncthreads();
      if(x < 32){
        smem[x] += smem[x + 32];
        smem[x] += smem[x + 16];
        smem[x] += smem[x + 8];
        smem[x] += smem[x + 4];
        smem[x] += smem[x + 2];
        smem[x] += smem[x + 1];
      }
#elif TYPE == 1
      // 632M
      // target values
      smem[x] = x;

      // reduction
      if(x < 36){
        smem[x] += smem[x + 36];
      }
      __syncthreads();
      if(x < 18){
        smem[x] += smem[x + 18];
	smem[x] += smem[x + 9];
	smem[x] += smem[x + 3] + smem[x + 6];
	smem[x] += smem[x + 1] + smem[x + 2];
      }
#else
      // 624M
      // target values
      smem[x] = x;

      // reduction
      if(x < 24){
        smem[x] += smem[x + 24] + smem[x + 48];
      }
      __syncthreads();
      if(x < 8){
        smem[x] += smem[x + 8] + smem[x + 16];
	smem[x] += smem[x + 4];
	smem[x] += smem[x + 2];
	smem[x] += smem[x + 1];
      }
#endif
      if(x == 0) output[0] = smem[0];
    }
    __syncthreads();
    if (x == 0) timer[bid+gridDim.x] = clock();
}