はじめに

以前書いたUMatの内部処理の続編ということで，ここではOpenCVのUMat内部でどのようにOpenCLの処理が行われているかについてもう少し詳細に紹介します．

対象読者

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

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

UMat(OpenCL)の処理の流れ

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


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

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

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

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

OpenCLカーネルのコンパイル

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

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

OpenCLカーネルの格納場所

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

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

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"

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

OpenCLプログラムの引数指定

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

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プログラムを実行しています．

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プログラムを実行するようになっています．

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上で以下のコメントを頂きました．

@dandelion1124 OpenCLのカーネルは一度コンパイルするとキャッシュされますので、OpenCV独自でもないようです。ただしまだキャッシュの寿命についてはよくわかっていませんが
— まお（松岡洋） (@kuronekodaisuki) 2016年5月4日

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

サンプルコード

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

# 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
Windows 10 Pro（64bit）
Visual Studio 2013 Update5
NVIDIA CUDA Toolkit 7.5