ばぐだし

CUDA エクササイズ３は配列の反転。

単一ブロック、複数ブロック、複数ブロック高速版があった。

単一ブロック

// Part 1 of 1: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int in = threadIdx.x;
    int out = blockDim.x - 1 - threadIdx.x;
    d_out[out] = d_in[in];
}

複数ブロック

// Part3: implement the kernel
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int inOffset  = blockDim.x * blockIdx.x;
    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
    int in  = inOffset + threadIdx.x;
    int out = outOffset + (blockDim.x - 1 - threadIdx.x);
    d_out[out] = d_in[in];
}

複数ブロック高速版

// Part 2 of 2: implement the fast kernel using shared memory
__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    extern __shared__ int s_data[];

    int inOffset  = blockDim.x * blockIdx.x;
    int in  = inOffset + threadIdx.x;

    // Load one element per thread from device memory and store it
    // *in reversed order* into temporary shared memory
    s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];

    // Block until all threads in the block have written their data to shared mem
    __syncthreads();

    // write the data from shared memory in forward order,
    // but to the reversed block offset as before

    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);

    int out = outOffset + threadIdx.x;
    d_out[out] = s_data[threadIdx.x];
}

不思議なのは高速版で、sharedメモリ（ブロック内でのみアクセス）を介することで高速化しているらしい。
普通に見たらsharedメモリ挟む分速度が下がりそうだけど、ブロック単位にメモリアクセスを分ければ高速になるのだろうか。ぐぬぬ。

まぁこれで一通りのエクササイズが終了しました。
次はCUBLASについて調べようと思います。