データの転送

2011年1月8日

ファイル

01-copy.tar.gz

内容

GPU に処理を行わせるには、ホストからデバイスへデータを転送する必要があります。ここではそのデータ転送について見てみます。

ベースの C プログラム

copy.cu

...
int main(void)
{
	float *a;
	float *b;
	float *c;
	
	...

	a = (float *)malloc(N*sizeof(float));
	b = (float *)malloc(N*sizeof(float));
	c = (float *)malloc(N*sizeof(float));

	...

		memcpy(b, a, N*sizeof(float));

		memcpy(c, b, N*sizeof(float));

	...

	free(a);
	free(b);
	free(c);
	
	return 0;
}

上記は CUDA を使わない単純なデータのコピーの例です。CUDA ではソースコードの拡張子を .cu とし、nvcc というコマンドでコンパイルします。

$ nvcc -o copy copy.cu

copy.cu はただの C のプログラムですが、CUDA は C/C++ の拡張なので nvcc でもコンパイルできます。

CUDA 用のプログラム

では、これを CUDA 用に書き直します。

copy2.cu

	...
	float *a;
	float *b;
	float *c;
	
	...
	
	a = (float *)malloc(N*sizeof(float));
	cudaMalloc(&b, N*sizeof(float));
	c = (float *)malloc(N*sizeof(float));

	...

		cudaMemcpy(b, a, N*sizeof(float), cudaMemcpyHostToDevice);

		cudaMemcpy(c, b, N*sizeof(float), cudaMemcpyDeviceToHost);

	...

	free(a);
	cudaFree(b);
	free(c);
	...

b をデバイスのメモリとして確保し、a -> b (ホストからデバイスへ)、b -> c (デバイスからホストへ) とデータを転送しています。

デバイスのメモリの確保は cudaMalloc()、その解放は cudaFree() で行います。デバイスのメモリは通常のポインタで管理できます。ただし、デバイスのメモリにはデバイス側のコードからしかアクセスできません (だから転送が必要)。ホストのポインタかデバイスのポインタかは区別がつかないので、ポインタの名前できちんと区別をつけておく必要があります。

データの転送は cudaMemcpy() で行います。ホストからデバイスへの転送には cudaMemcpyHostToDevice、デバイスからホストへの転送には cudaMemcpyDeviceToHost を指定します。

ところで、ここでは float を使っていますが、そのままでは double を使うことができません。double を使う場合は、つぎのようにコンパイラへそのことを伝える必要があります。

$ nvcc -arch sm_13 -o copy copy.cu

オプション "-arc sm_13" は、倍精度浮動小数点数に対応した Compute Capability 1.3 以上のハードウェアを想定しますよ、という意味です。

ページロックの使用

ホストのメモリを確保するときに "ページロック" という機能を使うと、ホスト - デバイス間のデータ転送性能が上がります。そのためには malloc()/free() の代わりに cudaMallocHost()/cudaFreeHost() を使います。

copy3.cu

	...
	cudaMallocHost(&a, N*sizeof(float));
	cudaMalloc(&b, N*sizeof(float));
	cudaMallocHost(&c, N*sizeof(float));

	...

	cudaFreeHost(a);
	cudaFree(b);
	cudaFreeHost(c);
	...

速度の比較

上で見た 3 つのプログラムについて実行速度を比較します。単純に転送するだけでは一瞬で処理が終わってしまうので、数秒の時間がかかるように転送を繰り返すようにしています。

ホスト - デバイス間転送にはかなりの時間を要しています。このことから GPU 計算においてはデータ転送がネックになりそうなことがわかります。また、ページロックの使用で少しは転送が速くなることもわかります。

CUDA プログラミング 目次へ