OpenGL de プログラミング

配列データを用意し、それをGPU側で処理をする場合

• GPUにデータを転送
• GPUで処理(global関数へ渡す)
• 計算結果を戻す

となる。

cudaMemcpy()関数を使う事でできる。

[例]

---

	float* x; //CPU用
	... ...
	float* gpu_x;//GPU用
	... ...
	CUDA_SAFE_CALL( cudaMemcpy( gpu_x, x, DataSize , cudaMemcpyHostToDevice)); //Host -> GPUのデータコピー
	... //--- 何らかの計算 ---//  ...
	CUDA_SAFE_CALL( cudaMemcpy( x, gpu_x,DataSize, cudaMemcpyDeviceToHost) );//GPU -> Hostへの転送
	... ...

---

__global__関数内で配列要素にアクセスする時は、「スレッドID」や「ブロックID」を用いる事ができる。

threadIdx型。「threadIdx.x」「threadIdx.y」などでスレッドのIDが取得できる
例えば、

---

dim3 threads(50,1,1);
 ... ...
GpuFunction<<<... ,threads, ...>>>(...);

---

とすると、

• threadIdx.x → 0〜49
• threadIdx.y → 0

となる。

blockIdx型。threadIdx型と同様。例えば、

---

dim3 grid(2,1,1);
... ...
GpuFunction<<<grid,... ...>>>(... ...);

---

とすると

• blockIdx.x = 0 か 1
• blockIdx.y = 0

となる。

blockDim型。1ブロックの大きさがわかる。例えば、

---

dim3 threads(50,1,1);
 ... ...
GpuFunction<<<... ,threads, ...>>>(...);

---

とすると、
blockDim.x　→　５０
blockDim.y　→　１
となる。

gridDim型。グリッドの大きさがわかる。例えば、

---

dim3 grid(2,1,1);
... ...
GpuFunction<<<grid,... ...>>>(... ...);

---

とすれば、
gridDim.x　→　２
gridDim.y　→　１
となる。

dim3 threads(a,1,1);
dim3 grid(b,1,1);

などのようにa、ｂの部分のみ変わる場合

スレッドId = blockIdx.x * blockDim.x + threadIdx.x

[例]：ブロック数(4×1)、グリッド数(25×1)個の場合

---

//----------- GPU上の計算 --------------------//
__global__ void GpuFunction(float *input)
{
	const int TidX = blockIdx.x * blockDim.x + threadIdx.x;  //スレッドIdを取得

	input[TidX] += 100;  //要素に100を加算
	__syncthreads(); //同期をとる
}
		... ...

	float* gpu_x;
	dim3 threads(25,1,1); //Db
	dim3 grid(4,1,1); //Dg
		... ...
	GpuFunction<<<grid,threads>>>(gpu_x);

		... ...

---

dim3 threads(a,1,1);
dim3 grid(b,c,1);

などのようにa、ｂ、cの部分のみ変わる場合

スレッドId = (blockIdx.y * gridDim.x + blockIdx.x) * blockDim.x + threadIdx.x

[例]：ブロック数(2×2)、グリッド数(25×1)個の場合

---

//----------- GPU上の計算 --------------------//
__global__ void GpuFunction(float *input)
{
	const int TidX = (blockIdx.y * gridDim.x + blockIdx.x) * blockDim.x + threadIdx.x;  //スレッドIdを取得

	input[TidX] += 100;  //要素に100を加算
	__syncthreads(); //同期をとる
}
		... ...

	float* gpu_x;
	dim3 threads(25,1,1); //Db
	dim3 grid(2,2,1); //Dg
		... ...
	GpuFunction<<<grid,threads>>>(gpu_x);

		... ...

---

dim3 threads(a,b,1);
dim3 grid(c,d,1);

などのようにa、ｂ、c、dのが変わる場合

---

スレッドId
 = threadIdx.x+blockDim.x*threadIdx.y
 + (blockIdx.x*blockDim.x*blockDim.y)
 + (blockIdx.y*blockDim.x*blockDim.y*gridDim.x)

---

[例]：ブロック数(2×2)、グリッド数(5×5)個の場合

---

//----------- GPU上の計算 --------------------//
__global__ void GpuFunction(float *input)
{
	const int TidX 
	= threadIdx.x+blockDim.x*threadIdx.y
	+(blockIdx.x*blockDim.x*blockDim.y)
	+(blockIdx.y*blockDim.x*blockDim.y*gridDim.x);  


	input[TidX] += 100;  //要素に100を加算
	__syncthreads(); //同期をとる
}
		... ...

	float* gpu_x;
	dim3 threads(5,5,1); //Db
	dim3 grid(2,2,1); //Dg
		... ...
	GpuFunction<<<grid,threads>>>(gpu_x);

		... ...

---

※メモ
1次元配列風にアクセスしてもよいが、すなおに２つに分けた方がわかりやすい(と思う)。

---

... ...
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
... ...

---

サンプルコード >> CUDA::基本編03::まとめコード