|
ここでは二つのベクトルの和の計算（C=A+B）を例にとってCUDAプログラミング例を示す．まず，CPUで実行される通常のC言語でのサンプルコードを以下に示す．
|
// main.c
#include <stdio.h>
#define NVEC (10000)
int main(){
int a[NVEC], b[NVEC], c[NVEC], i; // declare all variables
for (i = 0; i < NVEC; i++) a[i] = i * 2; // initialize A
for (i = 0; i < NVEC; i++) b[i] = NVEC - i; // initialize B
for (i = 0; i < NVEC; i++) c[i] = a[i] + b[i]; // calculate C = A + B
return 0;
}
|
|
これと同じことをGPUで動かすためにCUDAでこれから書くわけだが，ここではC=A+Bを計算する部分を他のファイルに分離する．ここでは，gpu_addvecという関数をCUDAで書くことにする．
|
// followings are written in main.c
#include <stdio.h>
#define NVEC (10000)
#define GPUFLAG (1)
extern void gpu_addvec(int*, int*, int*);
int main(){
int a[NVEC], b[NVEC], c[NVEC], i; // declare all variables
for (i = 0; i < NVEC; i++) a[i] = i * 2; // initialize A
for (i = 0; i < NVEC; i++) b[i] = NVEC - i; // initialize B
if (GPUFLAG == 1){
gpu_addvec(c,a,b);
}else{
for (i = 0; i < NVEC; i++) c[i] = a[i] + b[i]; // calculate C = A + B
}
return 0;
}
// followings are written in addvec.cu
#include <stdio.h>
#define NVEC (10000)
extern "C" void gpu_addvec(int* c, int* a, int* b){
}
|
|
これからgpu_addvecの中身書いていく．まず必要なのは，ビデオメモリへの情報の転送である．ここでは，ビデオメモリ上に置く配列を宣言，アロケートし，メインメモリ-ビデオメモリ間の情報の転送を行い，最後にビデオメモリ上に置く配列を解放している．また，cutil.hというヘッダをインクルードしている．このファイルは
/usr/local/cuda/bin/nvcc --compile addvec.cu -I$CUDA_SDK_PATH/common/inc
という手続きでオブジェクトファイルに変換される．（64 bit Mac OS Xなどでは-m64を付ける必要あり）$CUDA_SDK_PATHはSDKがインストールされている場所のパスであり，環境変数に書いてもよい．また，/usr/local/cudaにtoolkitがインストールされているものとしている．そして実行ファイルは
gcc main.c addvec.o -L/usr/local/cuda/lib -lcuda -lcudart -lm -lstdc++
という手続きで生成される．
|
// followings are written in addvec.cu
#include <stdio.h>
#include <cutil.h>
#define NVEC (10000)
extern "C" void gpu_addvec(int* c, int* a, int* b){
int *a_dev, *b_dev, *c_dev; // declare arrays for GPU computation
unsigned int memsize = sizeof(int) *NVEC; // memory size of each array
cudaMalloc((void**)&a_dev,memsize); // allocate arrays on video memory (global memory)
cudaMalloc((void**)&b_dev,memsize);
cudaMalloc((void**)&c_dev,memsize);
cudaMemcpy(a_dev,a,memsize,cudaMemcpyHostToDevice); // transfer data to video memory
cudaMemcpy(b_dev,b,memsize,cudaMemcpyHostToDevice);
cudaMemcpy(c,c_dev,memsize,cudaMemcpyDeviceToHost); // transfer data from video memory
cudaFree(a_dev);
cudaFree(b_dev);
cudaFree(c_dev);
}
|
|
次に，GPU上で走る各スレッドの実行する関数（カーネル）を考える．gpu_addvecにはそれを制御する命令を書かなければならないのだが，先にカーネルそのものの方を示す．__global__という修飾子はカーネルの関数につけられるものである．indexという変数の中身は次に触れるが，ここでは各スレッドごとに異なる値を持つものである．
|
// followings are written in addvec.cu
#define BLOCKSIZE (100)
__global__ void gpu_addvec_kernel(int *c_dev, int *a_dev, int* b_dev){
int index = threadIdx.x + blockIdx.x * BLOCKSIZE;
c_dev[index] = a_dev[index] + b_dev[index];
}
|
|
最後に，gpu_addvec_kernelの制御をgpu_addvecから行えば完成である．gpu_addvec_kernelの呼び出しの際に<<<grids,threads>>>というのが付いているが，これはスレッドの個数を指定するものである．CUDAではスレッドをブロックという単位で複数に束ねて実行するため，1ブロックあたりのスレッド数（ここではBLOCKSIZE）とブロックの総数（ここではNVEC/BLOCKSIZE）という情報が必要になる．gpu_addvec_kernelを呼び出す直前でこれらを指定している．また，gpu_addvec_kernel内でthreadIdx.xとblockIdx.xという変数を用いているが，これはそれぞれブロック内でのスレッドのIDおよびブロックのIDを示している．
|
// followings are written in addvec.cu
#include <stdio.h>
#define NVEC (10000)
extern "C" void gpu_addvec(int* c, int* a, int* b){
int *a_dev, *b_dev, *c_dev; // declare arrays for GPU computation
unsigned int memsize = sizeof(int) *NVEC; // memory size of each array
cudaMalloc((void**)&a_dev,memsize); // allocate arrays on video memory (global memory)
cudaMalloc((void**)&b_dev,memsize);
cudaMalloc((void**)&c_dev,memsize);
cudaMemcpy(a_dev,a,memsize,cudaMemcpyHostToDevice); // transfer data to video memory
cudaMemcpy(b_dev,b,memsize,cudaMemcpyHostToDevice);
dim3 threads(BLOCKSIZE);
dim3 grids(NVEC/BLOCKSIZE);
gpu_addvec_kernel<<<grids,threads>>>(c_dev,a_dev,b_dev);
cudaMemcpy(c,c_dev,memsize,cudaMemcpyDeviceToHost); // transfer data from video memory
cudaFree(a_dev);
cudaFree(b_dev);
cudaFree(c_dev);
}
|
|
