UMatの内部処理（OpenCL編）

# はじめに
以前書いた[UMatの内部処理](http://qiita.com/dandelion1124/items/fad18a164108158585e3)の続編ということで，ここではOpenCVのUMat内部でどのようにOpenCLの処理が行われているかについてもう少し詳細に紹介します．

# 対象読者
本記事の対象読者は以下の通りです．

- UMatの内部処理をデバッグする方<br>or
- UMatのOpenCL実装をしようと思っている方

# UMat(OpenCL)の処理の流れ
ここではcv::thresholdを例に挙げてUMatの内部処理(OpenCL)を説明します．
cv::thresholdの入力がUMat型のインスタンスの場合，以下のような流れで処理が行われます．

<pre>
cv::threshold
→ocl_threshold
 →OpenCLカーネルのコンパイル
 →OpenCLプログラムの引数指定
 →OpenCLプログラムの実行
</pre>

ということでOpenCLを使ったことがある人ならおなじみの以下の処理がUMat(OpenCL)の中で実行されています．

- OpenCLカーネルのコンパイル
- OpenCLプログラムの引数指定
- OpenCLプログラムの実行

以降，これらの処理内容について述べていきます．

## OpenCLカーネルのコンパイル
ocl::KernelのコンストラクタでOpenCLカーネルのコンパイルを行います．
cv::thresholdの入力がUMat型のインスタンスであれば以下の処理でOpenCLカーネルのコンパイルを行います．

```cpp:
ocl::Kernel k("threshold", ocl::imgproc::threshold_oclsrc,
              format("-D %s -D T=%s -D T1=%s -D STRIDE_SIZE=%d%s", thresholdMap[thresh_type],
                     ocl::typeToStr(ktype), ocl::typeToStr(depth), stride_size,
                     doubleSupport ? " -D DOUBLE_SUPPORT" : ""));
```

ocl::Kernelクラスのコンストラクタ引数は以下の通りです．

| 引数    | 意味|
|:-------|:------------|
|第1引数 |OpenCLカーネル名<br>（上記の例では**threshold**） |
|第2引数 |OpenCLカーネルソースが格納された変数<br>（上記の例では**ocl::imgproc::threshold_oclsrc**） |
|第3引数 |カーネルのビルドオプション<br>（上記の例では**type**，**depth**，**strideサイズ**をビルドオプションで渡している） |
|第4引数 |エラーメッセージ<br>（上記の例では指定していない） |

### OpenCLカーネルの格納場所
前述の通りcv::thresholdのOpenCLカーネルは**ocl::imgproc::threshold_oclsrc**の中に文字列として格納されているわけですが，OpenCLカーネルの実体はOpenCVのどこで定義されているかというと，OpenCVビルド時に生成されるopencl_kernels_imgproc.cppにて定義されています．

**threshold_oclsrc**のOpenCLカーネルを抜粋したものを以下に示します．

```cpp:
const struct ProgramEntry threshold={"threshold",
"#ifdef DOUBLE_SUPPORT\n"
"#ifdef cl_amd_fp64\n"
"#pragma OPENCL EXTENSION cl_amd_fp64:enable\n"
"#elif defined (cl_khr_fp64)\n"
"#pragma OPENCL EXTENSION cl_khr_fp64:enable\n"
"#endif\n"
"#endif\n"
"__kernel void threshold(__global const uchar * srcptr, int src_step, int src_offset,\n"
"__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,\n"
"T1 thresh, T1 max_val, T1 min_val)\n"
"{\n"
"int gx = get_global_id(0);\n"
"int gy = get_global_id(1) * STRIDE_SIZE;\n"
"if (gx < cols)\n"
"{\n"
"int src_index = mad24(gy, src_step, mad24(gx, (int)sizeof(T), src_offset));\n"
"int dst_index = mad24(gy, dst_step, mad24(gx, (int)sizeof(T), dst_offset));\n"
```

<br>**引用符や\nがあって読みにくいんだけど・・・**と思った方は安心してください．
<code>modules\imgproc\src\opencl\threshold.cl</code>に元のOpenCLコードがあります．そのため，OpenCVでOpenCLを使ってどのように実装しているかを知りたい場合は同梱の*.clファイルを読むとよいでしょう．
※OpenCVビルド時にこれらの.clファイルがopencl_kernels_imgproc.cpp等にまとめられるという作りになっています．

## OpenCLプログラムの引数指定
ocl::Kernelクラスのargsメソッドを呼ぶことでOpenCLプログラムの引数を指定しています．
ocl_thresholdでは以下の処理によってOpenCLカーネルの引数を設定します．

```cpp:
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst, cn, kercn),
       ocl::KernelArg::Constant(Mat(1, 1, depth, Scalar::all(thresh))),
       ocl::KernelArg::Constant(Mat(1, 1, depth, Scalar::all(maxval))),
       ocl::KernelArg::Constant(Mat(1, 1, depth, Scalar::all(min_val))));
```

## OpenCLプログラムの実行
ocl::Kernelクラスのrunメソッドを呼ぶことでOpenCLプログラムを実行しています．
ocl_thresholdでは以下の処理によってOpenCLプログラムを実行しています．

```cpp:
size_t globalsize[2] = { (size_t)dst.cols * cn / kercn, (size_t)dst.rows };
globalsize[1] = (globalsize[1] + stride_size - 1) / stride_size;
return k.run(2, globalsize, NULL, false);
```

runメソッドの引数は以下の通りです．

| 引数    |意味|
|:-------|:------------|
|第1引数 |次元数|
|第2引数 |global_work_size|
|第3引数 |local_work_size|
|第4引数 |カーネル実行待ちを行うか否かを指定するブール値 |

# 補足
## OpenCLプログラムのキャッシュ
OpenCVのUMatではOpenCLプログラムをキャッシュする仕組みが備わっています．
具体的にはocl::ProgramクラスのgetProgメソッドでキャッシュ処理が動いています．そして，一度コンパイルしたOpenCLプログラムはmapに格納されており，ハッシュキーで検索してヒットしたOpenCLプログラムを実行するようになっています．

```cpp:
Program getProg(const ProgramSource& src,
                const String& buildflags, String& errmsg)
{
    String prefix = Program::getPrefix(buildflags);
    HashKey k(src.hash(), crc64((const uchar*)prefix.c_str(), prefix.size()));
    phash_t::iterator it = phash.find(k);
    if( it != phash.end() )
        return it->second;
    //String filename = format("%08x%08x_%08x%08x.clb2",
    Program prog(src, buildflags, errmsg);
    if(prog.ptr())
        phash.insert(std::pair<HashKey,Program>(k, prog));
    return prog;
}
```

このgetProgメソッドはocl::Kernelクラスのコンストラクタで呼ばれるようになっていることから，一度コンパイルしたOpenCLプログラムに関してはOpenCLカーネルのコンパイルが行われません．

### 追記（2016/5/5）

twitter上で以下のコメントを頂きました．

<blockquote class="twitter-tweet" data-conversation="none" data-lang="ja"><p lang="ja" dir="ltr"><a href="https://twitter.com/dandelion1124">@dandelion1124</a> OpenCLのカーネルは一度コンパイルするとキャッシュされますので、OpenCV独自でもないようです。ただしまだキャッシュの寿命についてはよくわかっていませんが</p>&mdash; まお（松岡洋） (@kuronekodaisuki) <a href="https://twitter.com/kuronekodaisuki/status/727907955446665217">2016年5月4日</a></blockquote>
<script async src="//platform.twitter.com/widgets.js" charset="utf-8"></script>

ご指摘のようにOpenCVで独自にキャッシュ機構を実装している背景は気になるところですね・・・．

# サンプルコード
今回用いたサンプルコードは以下の通りです．

```cpp:
#include <opencv2/core.hpp>
#include <opencv2/imgproc.hpp>
#include <opencv2/imgcodecs.hpp>
#include <opencv2/highgui.hpp>

#include <iostream>

int main(int argc, char *argv[])
{
    cv::Mat src = cv::imread("lena.jpg", cv::IMREAD_GRAYSCALE);
    if (src.empty())
    {
        std::cerr << "Failed to open image file." << std::endl;
        return -1;
    }

    cv::UMat u_src, u_bin;
    src.copyTo(u_src);

    cv::threshold(u_src, u_bin, 100, 255, cv::THRESH_BINARY);

    cv::namedWindow("threshold");
    cv::imshow("threshold", u_bin);
    cv::waitKey(0);
    cv::destroyAllWindows();

    return 0;
}
```

# おわりに
この記事ではOpenCVのUMat内部でどのようにOpenCLの処理が行われているかを紹介しました．

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

* [OpenCV 3.1.0](https://github.com/Itseez/opencv/releases/tag/3.1.0)
* Windows 10 Pro（64bit）
* Visual Studio 2013 Update5
* NVIDIA CUDA Toolkit 7.5