cudaモジュールのcuda::Streamを使ってみよう

この記事は[OpenCV Advent Calendar 2016](http://qiita.com/advent-calendar/2016/opencv)の4日目の記事です．

# はじめに
この記事ではOpenCVのcudaモジュールで提供されている**cuda::Stream**とその使い方を紹介します．
今回紹介する**cuda::Stream**ですが，もしかしたらOpenCVドキュメントで見たことある方いらっしゃるかもしれませんね．

![opencv_cuda_stream.png](https://qiita-image-store.s3.amazonaws.com/0/31158/30baf4e7-f79b-6660-d7c1-ea93391f5869.png)

以降は基本的にOpenCV 3.x系を前提に説明します（最後の方で2.4系についても軽く触れます）．

# cuda::Streamクラスとは
CUDAのStreamはGPU上の処理を管理するキューのことで**カーネル実行やメモリ転送の並列性，実行順序**を制御するために用いられます．CUDAのStreamに関しては以下の記事を参照ください．

* [ストリームを用いたコンカレントカーネルプログラミングと最適化](http://on-demand.gputechconf.com/gtc/2014/jp/sessions/4004.pdf)
* [How to Overlap Data Transfers in CUDA C/C++](https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/)
* [GPU Pro Tip: CUDA 7 Streams Simplify Concurrency](https://devblogs.nvidia.com/parallelforall/gpu-pro-tip-cuda-7-streams-simplify-concurrency/)

また，OpenCVのcudaモジュールではこのStreamを簡単に利用するために**cuda::Stream**クラスを提供しています．

## **cuda::Stream**クラスが提供するAPI
冒頭でも紹介した**cuda::Stream**クラスは以下のAPIを提供しています．
各APIの詳細については[公式ドキュメント](http://docs.opencv.org/3.1.0/d9/df3/classcv_1_1cuda_1_1Stream.html)を参照ください．

| API                | 機能        |
|:-------------------|:------------|
| waitForCompletion  | Streamの一連の処理が終わるまで現在のCPUスレッドをブロッキングする |
| waitEvent          | Streamのイベントを待機する<br>※イベントについては[公式ドキュメント](http://docs.opencv.org/3.1.0/d5/d38/classcv_1_1cuda_1_1Event.html)を参照ください |
| queryIfComplete    | 現在のstream queueの処理が終わっているかどうかを調べる |
| enqueueHostCallback| Streamのキューイング処理完了後にホスト側で実行されるコールバック関数を設定する |


また，enqueueHostCallbackの使い方は[公式テストコード](https://github.com/opencv/opencv/blob/3.1.0/modules/cudaarithm/test/test_stream.cpp)が参考になるでしょう．

# **cuda::Stream**クラスの使い方
OpenCVのcudaモジュールAPIで明示的に**cuda::Stream**を使うのは非常に簡単です．
基本的な実装手順は以下の通りです．

1. **cuda::Stream**クラスのインスタンスを生成する
1. cudaモジュールのAPIに**cuda::Stream**クラスのインスタンスを渡す

一言で言うと**「cuda::StreamクラスのインタンスをcudaモジュールのAPIに明示的にセットする」**だけでOKです．

## サンプルコード
おそらく実際のコードを読むのがわかりやすいと思うのでサンプルコードを以下に示します．

```cpp:
// (1) cuda::Streamクラスのインスタンスを生成する
cv::cuda::Stream stream;

// (2) cudaモジュールのAPIにcuda::Streamクラスのインスタンスを渡す
cv::cuda::cvtColor(d_img, d_gray, cv::COLOR_BGR2GRAY, 0, stream);
```

# **cuda::Stream**を使うメリット
**cuda::Stream**を使うことでデータ転送とカーネル実行をオーバーラップさせることができ，データ転送時間を隠蔽できるようになります．と言ってもピンと来ないと思うのでサンプルコードを用いながら説明します．

## 明示的に**cuda::Stream**をセットしない場合
OpenCVのcudaモジュールAPIに**cuda::Stream**クラスのインスタンスを明示的にセットしない場合，
**Default Stream**を使って処理が行われます．ここではサンプルコードと[NVIDIA Visual Profiler](https://developer.nvidia.com/nvidia-visual-profiler)で取得したタイムラインを示します．

### サンプルコード
```cpp:
cv::Mat src(cv::Size(3840, 2160), CV_8UC3, cv::Scalar(0));
cv::cuda::HostMem dst1, dst2, dst3;
cv::cuda::GpuMat d_src;
cv::cuda::GpuMat d_gray1, d_gray2, d_gray3;
cv::cuda::GpuMat d_dst1, d_dst2, d_dst3;

// デバイスに転送
d_src.upload(src);

// GpuMatを使った処理(Default Stream)
cv::cuda::cvtColor(d_src, d_gray1, cv::COLOR_BGR2GRAY, 0);
cv::cuda::threshold(d_gray1, d_dst1, 200, 255, cv::THRESH_BINARY);

//ホストへ転送(Default Stream)
d_dst1.download(dst1);

// GpuMatを使った処理(Default Stream)
cv::cuda::cvtColor(d_src, d_gray2, cv::COLOR_BGR2GRAY, 0);
cv::cuda::threshold(d_gray2, d_dst2, 200, 255, cv::THRESH_BINARY);

//ホストに転送(Default Stream)
d_dst2.download(dst2);

// GpuMatを使った処理(Default Stream)
cv::cuda::cvtColor(d_src, d_gray3, cv::COLOR_BGR2GRAY, 0);
cv::cuda::threshold(d_gray3, d_dst3, 200, 255, cv::THRESH_BINARY);

//ホストに転送(Default Stream)
d_dst3.download(dst3);
```

### タイムライン
サンプルコードを実行したときのタイムラインは以下の通りです．

![timeline_default_stream.png](https://qiita-image-store.s3.amazonaws.com/0/31158/70aca3c3-6b9b-46f1-a35a-737e520cc86c.png)

cudaモジュールAPIに**cuda::Stream**クラスのインスタンスを明示的にセットしないと
**Default Stream**が使われていてデータ転送やカーネル実行が並列に行われてないことがわかります．

## 明示的に**cuda::Stream**をセットする場合
ここでは**cuda::Stream**を使った場合のサンプルコードとタイムラインを示します．

### サンプルコード
```cpp:
cv::Mat src(cv::Size(3840, 2160), CV_8UC3, cv::Scalar(0));
cv::cuda::HostMem dst1, dst2, dst3;
cv::cuda::GpuMat d_src;
cv::cuda::GpuMat d_gray1, d_gray2, d_gray3;
cv::cuda::GpuMat d_dst1, d_dst2, d_dst3;

cv::cuda::Stream stream[3];

// デバイスに転送
d_src.upload(src);

// GpuMatを使った処理(cuda::Stream)
cv::cuda::cvtColor(d_src, d_gray1, cv::COLOR_BGR2GRAY, 0, stream[0]);
cv::cuda::threshold(d_gray1, d_dst1, 200, 255, cv::THRESH_BINARY, stream[0]);

//ホストに転送(cuda::Stream)
d_dst1.download(dst1, stream[0]);

// GpuMatを使った処理(cuda::Stream)
cv::cuda::cvtColor(d_src, d_gray2, cv::COLOR_BGR2GRAY, 0, stream[1]);
cv::cuda::threshold(d_gray2, d_dst2, 200, 255, cv::THRESH_BINARY, stream[1]);

//ホストに転送(cuda::Stream)
d_dst2.download(dst2, stream[1]);

// GpuMatを使った処理(cuda::Stream)
cv::cuda::cvtColor(d_src, d_gray3, cv::COLOR_BGR2GRAY, 0, stream[2]);
cv::cuda::threshold(d_gray3, d_dst3, 200, 255, cv::THRESH_BINARY, stream[2]);

//ホストに転送(cuda::Stream)
d_dst3.download(dst3, stream[2]);
```

### タイムライン
サンプルコードを実行したときのタイムラインは以下の通りです．

![timeline_multi_stream1.png](https://qiita-image-store.s3.amazonaws.com/0/31158/1e21c85e-ee6b-c6b8-cbc3-0112e6fd2e9a.png)

このサンプルコードでは複数の**cuda::Stream**を使ってデータ転送やカーネル実行に割り当てることでデータ転送とカーネル実行をオーバーラップさせることができていることがわかります．

また，サンプルコードは割愛しますが，今回用いたGeForce GTX 1060はCopy Engineを2つ持っているので，処理によっては以下のように**データ転送（ホスト→デバイス）**と**データ転送（デバイス→ホスト）**をオーバーラップさせることもできます．

![timeline_multi_stream2.png](https://qiita-image-store.s3.amazonaws.com/0/31158/d8d348df-c067-0a1f-fd65-a12e657a3f3b.png)


# **cuda::Stream**クラスを使った実装方法
ここまではOpenCVのcudaモジュールのAPIで**cuda::Stream**を明示的に使うお話でした．
一方でOpenCVと連携した自作CUDAカーネルで**cuda::Stream**を利用したいケースもあると思います．

基本的な実装手順は以下の通りです．

1. 関数の引数に**cuda::Stream**クラスの変数を追加する
1. <code>cuda&#058;&#058;StreamAccessor&#058;&#058;getStream</code>メソッド（[公式ドキュメント](http://docs.opencv.org/3.1.0/d6/df1/structcv_1_1cuda_1_1StreamAccessor.html#a5830d99635f08b462a429771bfad71a9)）を使ってCUDAのStreamを取得する
1. CUDAカーネル起動時に2.で取得したStreamを指定する

## サンプルコード
こちらもおそらく実際のコードを読むのがわかりやすいと思うのでサンプルコードを以下に示します．

```cpp:
// (1) 関数の引数にcuda::Streamクラスの変数を追加する
void hoge_func(cv::cuda::GpuMat& img, cv::cuda::Stream& stream_)
{
    // (2) CUDAのStreamを取得する
    cudaStream_t stream = cv::cuda::StreamAccessor::getStream(stream_);

    const dim3 threads(32, 8);
    const dim3 grid(divUp(img.cols, threads.x), divUp(img.rows, threads.y));

    // (3) CUDAカーネル起動時に(2)で取得したStreamを指定する
    hoge_kernel<<<grid, threads, 0, stream>>>(img);

    cudaGetLastError();
    cudaDeviceSynchronize();
}
```

前に述べたように

1. 関数の引数に**cuda::Stream**クラスの変数を追加する
1. <code>cuda&#058;&#058;StreamAccessor&#058;&#058;getStream</code>メソッド（[公式ドキュメント](http://docs.opencv.org/3.1.0/d6/df1/structcv_1_1cuda_1_1StreamAccessor.html#a5830d99635f08b462a429771bfad71a9)）を使ってCUDAのStreamを取得する
1. CUDAカーネル起動時に2.で取得したStreamを指定する

という流れそのままなので簡単ですね！

# OpenCV 2.4系は？
冒頭で述べたようにここまでOpenCV 3.x系を前提とした話になっていました．ここまで聞いて「じゃあ，**cuda::Stream**クラスってOpenCV 2.4系じゃ使えないの？」と疑問が生じるかもしれませんが安心してください．
OpenCV 2.4系だと[**gpu::Stream**クラス](http://docs.opencv.org/2.4.13/modules/gpu/doc/data_structures.html#gpu-stream)がその役割を担っています．

ただし，GpuMatクラスのupload，downloadメソッドの扱いが若干異っていて，OpenCV 2.4系だと<code>gpu&#058;&#058;Stream::enqueueUpload</code>，<code>gpu&#058;&#058;Stream::enqueueDownload</code>メソッドを使う必要があります．

各APIの詳細は以下のURLを参照ください。
* <http://docs.opencv.org/2.4.13/modules/gpu/doc/data_structures.html#gpu-stream-enqueueupload>
* <http://docs.opencv.org/2.4.13/modules/gpu/doc/data_structures.html#gpu-stream-enqueuedownload>

ということで，2.4系だと具体的にどういう実装になるかイメージしやすいよう
3.x系と対比しながらサンプルコードを見てみましょう．

## サンプルコード（OpenCV 3.x系）

```cpp:
// cuda::Streamクラスのインスタンスを生成する
cv::cuda::Stream stream;

// uploadメソッドにcuda::Streamクラスのインスタンスを渡す
d_src.upload(src, stream);

// GpuMatを使った処理

// downloadメソッドにcuda::Streamクラスのインスタンスを渡す
d_dst.download(dst, stream);

```

## サンプルコード（OpenCV 2.4系）

```cpp:
// gpu::Streamクラスのインスタンスを生成する
cv::gpu::Stream stream;

// enqueueUploadメソッドを使う
stream.enqueueUpload(src, d_src);

// GpuMatを使った処理

// enqueueDownloadメソッドを使う
stream.enqueueDownload(d_dst, dst);
```

ということで，2.4系でもそこまで大きな違いがなく，Streamが使えることがわかると思います．

# おわりに
この記事ではOpenCVのcudaモジュールで提供されている**cuda::Stream**とその使い方を紹介しました．
cudaモジュールを使う機会があれば必要に応じて**cuda::Stream**を使ってみてください！

# 備考
筆者は以下の環境で動作確認しました．

* CPU：Intel Core i7-6700HQ
* メモリ：64GB
* GPU：NVIDIA GeForce GTX 1060 / 6GB
* Ubuntu 16.04 LTS(64bit)
* [OpenCV 3.1.0](https://github.com/Itseez/opencv/releases/tag/3.1.0)
* CUDA 8.0
* gcc 5.4.0

# 参考URL
* [How to Overlap Data Transfers in CUDA C/C++](https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/)
* [GPU Pro Tip: CUDA 7 Streams Simplify Concurrency](https://devblogs.nvidia.com/parallelforall/gpu-pro-tip-cuda-7-streams-simplify-concurrency/)
* [ストリームを用いたコンカレントカーネルプログラミングと最適化](http://on-demand.gputechconf.com/gtc/2014/jp/sessions/4004.pdf)
* [streamで速くする(1) ～イントロ](http://fareastprogramming.hatenadiary.jp/entry/2016/10/21/220352)
* [streamで速くする(2)](http://fareastprogramming.hatenadiary.jp/entry/2016/10/21/221314)
* [streamで速くする(3) ～からくり](http://fareastprogramming.hatenadiary.jp/entry/2016/10/21/225438)
* [Getting started with OpenCV on GPUs](http://on-demand.gputechconf.com/gtc/2013/webinar/gtc-express-itseez-opencv-webinar.pdf)
* [Computer Vision on GPU with OpenCV](http://www.gipsa-lab.grenoble-inp.fr/summerschool/gpu2011/fichiers/6_Obukhov_OpenCV_rev5.pdf)
* [Hyper-Qの限界までStreamを並列動作させてみた on JetsonTX1](http://wazalabo.com/hyper-q-on-jetsontx1.html)
* [OpenCVのgpu::countNonZeroをgpu::Streamで高速化する](http://phst.hateblo.jp/entry/2016/10/05/140102)