ばぐだし

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について調べようと思います。

続いてエクササイズ２、カーネルの起動いきまーす。

ソース一部抜粋

// Part 3 of 5: implement the kernel
__global__ void myFirstKernel(int *d_a)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    d_a[idx] = 1000*blockIdx.x + threadIdx.x;
}

    // Part 2 of 5: configure and launch kernel
    dim3 dimGrid(numBlocks);
    dim3 dimBlock(numThreadsPerBlock);
    myFirstKernel<<<dimGrid,dimBlock>>>(d_a);

    // block until the device has completed
    cudaThreadSynchronize();

※以下私なりに丸めたちょっと怪しい表現を使います。正しい表現は本家ページで！
CUDAでは「関数名<<<グリッド内のブロックの数、１ブロックあたりのスレッドの数>>>(引数) 」で並列カーネル(デバイス上で実行される機能)を呼び出す。
CUDAで書いた関数(この場合myFirstKernel)が1つのスレッドに当たる。1度に実行できるのは1つのグリッドだけ。従って、一度に実行されるスレッド数(myFirstKernel)は、グリッド内のブロックの数(numBlocks)×１ブロックあたりのスレッドの数(numThreadsPerBlock)となる。
ちなみに、グリッド内のブロックの数、１ブロックあたりのスレッドの数は、2次元での表現も可能(ブロックあたりスレッド数は3次元まで可能)で、画像屋さんにはとても助かる。

また、カーネル内では、以下の変数が自動的に定義されており、アクセスができる。
dim3 gridDim;
ブロックのグリッドの次元（最大で2次元）
dim3 blockDim;
スレッドのブロックの次元
dim3 blockIdx;
グリッド内のブロックのインデックス
dim3 threadIdx;
ブロック内のスレッドのインデックス
てなわけでスレッドのインデックスとしては、

int idx = blockDim.x * blockIdx.x + threadIdx.x;

が1次元の並列カーネル呼び出しでは定番となる。

以上、次もエクササイズ！

CUDAってなんなのさ

CUDAをかじろうかと思う。
CUDAってのはGPGPUプログラムのための開発言語、開発環境のことで、まぁ詳しくはこちらを参照ってことで。
GPGPUに関してはこれから出るOpenCLの方が便利そうですし普及しそうだけど、そのベースになるCUDAも勉強しといて損はないかなと。
CUDAは以前サンプルを読んだりもしていたのですが、かなり断片的にしか分かっていないので今回は改めてお勉強。
基本的に本家本元NVIDIAのCUDAのページがとても親切に教えてくれるので、それを追う形でいこうかと。

CUDAのしくみ

アーキテクチャはこんなん。

NVIDIAHPより
ここでは実装に重きを置いて書くことにするので、基本的なCUDAの仕組みはここの「CUDA ベーシックプログラミングガイド」で学んでくださいな。というか私も今はこのページをなぞるだけ。

環境構築

以前よりSDKのバージョンも上がってるし、動かないサンプルがあったので開発環境構築から。
グラボ：GeForce9600GT
OS：Vista 32bit
ここから
1. 最新CUDAドライバ(185.85)
2. CUDAツールキット一式(2.2)
3. CUDA SDKコードサンプル(2.21)
をダウンロード・インストール。

CUDAトレーニング コース

nVidiaのページの学習コースに乗っかる。第1巻： CUDAプログラミング入門 (日本語版)ってとこ。
・エクササイズ1: ホストとデバイス間でのコピー
ホストはCPUで、デバイスはGPU。ホスト-デバイスはGlobal Memoryを通じてデータをやりとりする。大容量だけど遅いから使用は最小限に抑える。

課題部分のソースを抜粋

    // Part 1 of 5: allocate device memory
    size_t memSize = dimA*sizeof(float);
    cudaMalloc((void**)&d_a, memSize);
    cudaMalloc((void**)&d_b, memSize);

    // Part 2 of 5: host to device memory copy
    cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);

    // Part 3 of 5: device to device memory copy
    cudaMemcpy(d_b, d_a, memSize, cudaMemcpyDeviceToDevice);

    // Part 4 of 5: device to host copy
    cudaMemcpy(h_a, d_b, memSize, cudaMemcpyDeviceToHost);

    // Part 5 of 5: free device memory pointers d_a and d_b
    cudaFree(d_a);
    cudaFree(d_b);

とこんな風に、ホストとデバイスのメモリを意識して使い分ける。
おそらく、対象データをホストからデバイスにコピー→デバイス内で演算→結果データをデバイスからホストにコピーの流れになる。

今回は以上。次もエクササイズを追っていきます。