CUDAのスレッドは2次元で持つことができるので、その使い方の解説をします。
以下に紹介する処理で、スレッドを2次元で利用するコーディングの例を示します。
処理内容
別の記事でも紹介しましたが、以下のような処理を作成します。
・ 8000×8000のint型 二次元配列で、ある要素を中心として、その周辺3x3の平均値を求める。
・結果を別の二次元配列に保存する。
・この処理を全要素に行う。

コーディング例
カーネルを呼び出す側の関数になります。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 |
//処理関数 void Process_MeanGPU() { int sizex = 8192; int sizey = 8192; int length = sizex * sizey; size_t size = sizeof(int) * length; //ホストメモリのポインタ int* pHostSrc; //数値 int* pHostResult; //結果 //デバイスメモリのポインタ int* pDevSrc; //数値 int* pDevResult; //結果 //ホストメモリの確保 cudaMallocHost(&pHostSrc, size); cudaMallocHost(&pHostResult, size); //デバイスメモリの確保 cudaMalloc(&pDevSrc, size); cudaMalloc(&pDevResult, size); //ホストメモリに適当な値を設定 unsigned int value = 0 ; for (int n = 0; n < length; n++) { pHostSrc[n] = value++ % 1000; } //ホスト->デバイスへ数値を転送 cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice); //カーネル dim3 block(128, 1, 1); dim3 grid((sizex + 128 - 1) / 128, sizey, 1); kernel_mean <<<grid, block >>> (pDevSrc, pDevResult, sizex, sizey); //デバイス->ホストへ結果を転送 cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost); //デバイスメモリの開放 cudaFree(pDevSrc); cudaFree(pDevResult); //ホストメモリの開放 cudaFreeHost(pHostSrc); cudaFreeHost(pHostResult); } |
こちらはカーネルです。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 |
//平均を取得するカーネル __global__ void kernel_mean(int* pSrc, int *pResult, int sizex, int sizey) { int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex int idy = blockDim.y * blockIdx.y + threadIdx.y; //自分のスレッドyのindex if (idx < 1 || idy < 1 || idx >= sizex - 1 || idy >= sizey - 1) { return; //配列の左右上下端は無視 } unsigned int sum = 0; for (int y = 0; y < 3; y++) { for (int x = 0; x < 3; x++) { sum += pSrc[(idy + y - 1) * sizex + (idx + x - 1)]; } } pResult[idy * sizex + idx] = sum / 9; return; } |
int idx = blockDim.x * blockIdx.x + threadIdx.x;
int idy = blockDim.y * blockIdx.y + threadIdx.y;
の記述で、自分のスレッドのx,yインデックスを計算しています。
おまじないのようなものなので覚えておきましょう。
