
GPUで処理をさせるとき、たった一つだけのカーネルを実行させて終わり、という状況ばかりではないはずです。
例えば、
- 連続で複数カーネルを実行させる
- 複数スレッド(CPU側)から同時にカーネルを実行させる
- GPUのカーネルや転送と並行してCPUで処理をさせる
といったことをしたいケースもあるでしょう。
このように複数処理を効率的に行いたい場合、ストリーム(Stream)の利用は必須です。
この記事では、そんなストリームの使い方や特性を解説していきます。
ストリーム(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を使っていたり、転送とカーネルにストリームを指定しています。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 |
... //ストリーム生成 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では、「ホスト->デバイス」と「デバイス->ホスト」はオーバーラップしない場合があります。
