【CUDA】ストリームを使いこなして高速化しよう

GPUで処理をさせるとき、たった一つだけのカーネルを実行させて終わり、という状況ばかりではないはずです。

例えば、

連続で複数カーネルを実行させる
複数スレッド（CPU側）から同時にカーネルを実行させる
GPUのカーネルや転送と並行してCPUで処理をさせる

といったことをしたいケースもあるでしょう。
このように複数処理を効率的に行いたい場合、ストリーム(Stream)の利用は必須です。

この記事では、そんなストリームの使い方や特性を解説していきます。

ストリーム(Stream) とは
動作の特徴
ストリーム関連のAPI
APIごとの同期・非同期
コードとタイムライン

ストリーム(Stream) とは

ストリームはGPU処理のスケジュール管理の単位、または処理キューのようなものです。

なぜ、このようなものが必要なのでしょうか。
GPUのスケジューリングの特性から解説していきます。

データ転送やカーネル実行をCPUからGPUに命令したときですが、
CPUはGPUの処理完了を待たずに、次の動作に移る非同期動作ができます。
また、CPUの複数スレッドから任意のタイミングでGPUに処理を命令することができます。
この時、命令された処理はGPU内で自動的にスケジューリングされます。

ここで、命令した処理が終わったかどうかを把握したかったり、関連するカーネルや転送の処理順序は固定したくなります。

そんな時に利用するのがストリームです。

動作の特徴

ストリームとはGPU内の処理キューのようなもので、複数生成でき、データ転送やカーネルごとに一つづつ指定します。
ストリームを指定されたデータ転送とカーネルは、以下のような動きをします。

同じストリームのデータ転送やカーネルどうしは、命令された順に実行される。
別のストリームのデータ転送とカーネルどうしは、オーバーラップして処理時間が短縮されることがある。

このような動作をします。
また、CPU側はストリームごとに処理完了を待機できたりします。

ストリーム関連のAPI

以下のAPIでストリームの生成をしたり、状態を確認します。

API名	解説
cudaStreamCreate	ストリームを生成する
cudaStreamDestroy	ストリームを破棄する
cudaStreamSynchronize	指定ストリームの処理完了を待機
cudaStreamQuery	指定ストリームの処理完了を確認

APIごとの同期・非同期

ストリームを利用するとき、同期・非同期APIを意識して使用するようにしましょう。
複数の処理を命令したときに、そのたび同期していては意味がありません。

APIごとの同期・非同期動作の違いは以下の通りです。

同期API	非同期API	解説
cudaMemset	cudaMemsetAsync	デバイスメモリに値をセットする
cudaMemcpy	cudaMemcpyAsync	ホスト・デバイスメモリ間でデータを転送する
カーネル (ストリーム指定無し)	カーネル (ストリーム指定有り)
cudaMalloc/cudaFree	–	デバイスメモリの確保・開放

コードとタイムライン

コーディング例

ストリームを利用したときのコード例を以下に示します。
ホストメモリやデバイスメモリの確保については、冗長になるので省略しています。

ざっと眺めていただくと判りますが、転送にcudaMemcpyAsyncを使っていたり、転送とカーネルにストリームを指定しています。

...

	//ストリーム生成
	cudaStream_t stream1;
	cudaStream_t stream2;
	cudaStream_t stream3;
	cudaStreamCreate(&stream1);
	cudaStreamCreate(&stream2);
	cudaStreamCreate(&stream3);

	//ホスト->デバイスへ数値を転送
	cudaMemcpyAsync(pDevSrc1, pHostSrc1, size, cudaMemcpyHostToDevice, stream1);
	cudaMemcpyAsync(pDevSrc2, pHostSrc2, size, cudaMemcpyHostToDevice, stream2);
	cudaMemcpyAsync(pDevSrc3, pHostSrc3, size, cudaMemcpyHostToDevice, stream3);

	//カーネル
	dim3 block(128, 1, 1);
	dim3 grid(128, 1, 1);
	kernel_stream <<<grid, block, 0, stream1 >>> (pDevSrc1);
	kernel_stream <<<grid, block, 0, stream2 >>> (pDevSrc2);
	kernel_stream <<<grid, block, 0, stream3 >>> (pDevSrc3);

	//デバイス->ホストへ結果を転送
	cudaMemcpyAsync(pHostResult1, pDevSrc1, size, cudaMemcpyDeviceToHost, stream1);
	cudaMemcpyAsync(pHostResult2, pDevSrc2, size, cudaMemcpyDeviceToHost, stream2);
	cudaMemcpyAsync(pHostResult3, pDevSrc3, size, cudaMemcpyDeviceToHost, stream3);

	//転送を待機
	cudaStreamSynchronize(stream1);
	cudaStreamSynchronize(stream2);
	cudaStreamSynchronize(stream3);

	//ストリーム破棄
	cudaStreamDestroy(stream1);
	cudaStreamDestroy(stream2);
	cudaStreamDestroy(stream3);

...

...

//ストリーム生成

cudaStream_t stream1;

cudaStream_t stream2;

cudaStream_t stream3;

cudaStreamCreate(&stream1);

cudaStreamCreate(&stream2);

cudaStreamCreate(&stream3);

//ホスト->デバイスへ数値を転送

cudaMemcpyAsync(pDevSrc1, pHostSrc1, size, cudaMemcpyHostToDevice, stream1);

cudaMemcpyAsync(pDevSrc2, pHostSrc2, size, cudaMemcpyHostToDevice, stream2);

cudaMemcpyAsync(pDevSrc3, pHostSrc3, size, cudaMemcpyHostToDevice, stream3);

//カーネル

dim3 block(128, 1, 1);

dim3 grid(128, 1, 1);

kernel_stream <<<grid, block, 0, stream1 >>> (pDevSrc1);

kernel_stream <<<grid, block, 0, stream2 >>> (pDevSrc2);

kernel_stream <<<grid, block, 0, stream3 >>> (pDevSrc3);

//デバイス->ホストへ結果を転送

cudaMemcpyAsync(pHostResult1, pDevSrc1, size, cudaMemcpyDeviceToHost, stream1);

cudaMemcpyAsync(pHostResult2, pDevSrc2, size, cudaMemcpyDeviceToHost, stream2);

cudaMemcpyAsync(pHostResult3, pDevSrc3, size, cudaMemcpyDeviceToHost, stream3);

//転送を待機

cudaStreamSynchronize(stream1);

cudaStreamSynchronize(stream2);

cudaStreamSynchronize(stream3);

//ストリーム破棄

cudaStreamDestroy(stream1);

cudaStreamDestroy(stream2);

cudaStreamDestroy(stream3);

...

ストリームの生成や破棄は処理毎にやるのは無駄なので、実際はどこか処理外でやるようにしたほうがよいでしょう。
また、このようなシンプルな処理では、「cudaStreamSynchronize」ではなく、「cudaThreadSynchronize」などでGPUの処理完了を待機してもよいです。

タイムライン

実行したときのタイムラインは以下の通りです。
各ストリームの転送やカーネルがオーバーラップしていることが判ります。
オーバーラップした分だけ、トータルの処理時間を短縮することができます。

注意

・ホストメモリはページロックされている必要があります。

・昔のGeForceでは、「ホスト->デバイス」と「デバイス->ホスト」はオーバーラップしない場合があります。