現在地 >> メニュー >> CUDA >> CUDA::基本編03
INDEX:CUDA::基本編02 << CUDA::基本編03 >> CUDA::基本編04

配列操作の基本


配列データを用意し、それをGPU側で処理をする場合
  • GPUにデータを転送
  • GPUで処理(global関数へ渡す)
  • 計算結果を戻す
となる。

[1]GPU―GPU間のデータ転送


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への転送
	... ...


[2]スレッドId、ブロックIdなど


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

<スレッドId>


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

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

とすると、
  • threadIdx.x → 0〜49
  • threadIdx.y → 0
となる。

<ブロックId>


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 → 50
blockDim.y → 1
となる。

<グリッド幅>


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

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

とすれば、
gridDim.x → 2
gridDim.y → 1
となる。


[3]データのアクセス方法

<1次元グリッド、1次元スレッドのブロック>


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

などのようにa、bの部分のみ変わる場合
スレッド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);

		... ...

<2次元グリッド、1次元スレッドのブロック>


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

などのようにa、b、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);

		... ...


<2次元グリッド、2次元スレッドのブロック>


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

などのようにa、b、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次元配列風にアクセスしてもよいが、すなおに2つに分けた方がわかりやすい(と思う)。

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



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

この広告は60日間更新がないwikiに表示されております。

目次

― その他 ―

Wiki内検索

計測中...(07.10.8〜)

Save The World






▲よろしければ広告のクリックもお願いします


▲ランキングに参加しました

管理人/副管理人のみ編集できます