【CUDA】グリッド，ブロック，スレッド

CUDA勉強中
CUDA by Exampleからのメモ．

#define N 1024
__global__ void add(int *a, int *b, int *c){
	//int tid = threadIdx.x;
	if(tid < N){
		c[tid] = a[tid] + b[tid];  
	}
}

といったベクトルの加算をデバイスにさせたいときに、どうブロックやスレッドを分ければよいか。
手元の環境（QuadroK5000）だと、最大ブロック数65535、最大スレッド/ブロック数は1024

長さN:1024のベクトルを計算する場合

add<<<N,1>>>(dev_a, dev_b, dev_c);

とブロックを増やして1ブロックあたり1スレッドでも、

add<<<1,N>>>(dev_a, dev_b, dev_c);

として、1ブロックNスレッドでも動く。

長さN:1024*100のベクトルを計算する場合

この場合、スレッドとブロックを組み合わせる必要がある。
128スレッド/ブロックにするとすると、

add<<<N/128,128>>>(dev_a, dev_b, dev_c);

としがちだが、これだとN<128のときにスレッドが起動しない。
こういう場合は、以下のように記述する

add<<<(N+127)/128,128>>>(dev_a, dev_b, dev_c);

ちなみにこのときカーネルは、

__global__ void add(int *a, int *b, int *c){
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	if(tid < N){
	　　　c[tid] = a[tid] + b[tid];  
	}

}

と書き換える

長さN:1024*100000のベクトルを計算する場合

この長さになると、同時起動できる最大スレッド数を超えるので、
ひとつのスレッドに複数回計算させる。

カーネルを以下のように書き換えて、

__global__ void add(int *a, int *b, int *c){
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	while(tid < N){
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;
	}

とする。
この場合カーネルの呼び出しは、

add<<<128, 128>>>(dev_a, dev_b, dev_c);

でいい。
パフォーマンスについてはどうなるんだろう

単純に並列処理させたい総スレッド数だけを指定するのではなく、わざわざブロック数を指定するのは、
GPU内部が複数のStreaming Multiprocessor(SM)に分割されているため。

CUDAはブロックをSMへ割り当てる単位にしている。
たとえば、ブロック数を1に指定してしまうと、1つのSMしか動作しない。

では、ブロック数=SMとすれば解決、かというとそうではない。
ブロックあたりのスレッド数には上限がある。