OpenCV 3.0.0での独自カーネルOpenCL

#OpenCVとOpenCL
OpenCVではデフォルトでOpenCLが組み込まれていて、手軽にGPU等のデバイスでの並列化が実現できる。
OpenCV-CLを使えば、OpenCVのMat型とスムーズに連携でき、さらにOpenCVをインストールするだけでOpenCLの環境も整うので、コードの公開時に外部依存の少ないコードが書けそうな点に魅力を感じている。

ここでは特に、OpenCV 3.0における、独自カーネルによるOpenCLの実行方法について解説。
ただし、現状、公式のドキュメントがなく、ソースコードから類推しただけなので、間違いがあったら指摘してほしい。

# おさらい
## OpenCV 2.4でOpenCL
OpenCV 2.4では、cv::Matの代わりにcv::ocl::oclMatを使うと、フィルターなどのいくつかの処理がOpenCLで実行できるようになる。
詳しくは @dandelion1124 さんの解説を参考にしてほしい。
http://www.slideshare.net/YasuhiroYoshimura/gpgpu2opencvopencloclmat

基本的には、公式のOpenCV 2.4.7以降のバイナリであれば、OpenCLがONの状態でビルド済みらしいので

```cpp
#include <opencv2/ocl/ocl.hpp>
#pragma comment (lib, "opencv_ocl***.lib")
```
のように、oclモジュールをインクルード＆リンクしてやれば、（他にCUDAなどをインストールする必要なく）動作するようです。
ただし、*staticlib版のOpenCVを使用していると、プログラム終了時に異常終了する場合があります（OpenCV 2.4.10、Windows OS、Visual Studio2013で確認）*。

## OpenCV 2.4で独自カーネルOpenCL
予め用意された以外の処理を行いたい場合は、openCLExecuteKernelInterop関数を利用。
詳しくは @dandelion1124 さんの解説やコードを参考にしてほしい。
http://www.slideshare.net/YasuhiroYoshimura/cv-dandelion1124-20150125
https://gist.github.com/atinfinity/8c25c8fb1b3708aa0944

**OpenCV 2.4の問題点**
image2d_t型をカーネルに渡す方法がユーザーに提供されていないこと。
image2d_t は sampler_tと併せることで、GPU のテクスチャユニットを利用した高速な線形補完サンプリング（実数値座標での画素値取得）が実現できる。

ただ、OpenCV 2.4.10の /sources/modules/ocl/src/tv1flow.cpp のwarpBackward関数を見ると、

```C++:tv1flow.cpp
void ocl_tvl1flow::warpBackward(..., const oclMat &I1, ..)
{
    ...
    cl_mem I1_tex;
    ...
    I1_tex = bindTexture(I1);
    ...
    args.push_back( make_pair( sizeof(cl_mem), (void*)&I1_tex));
    ...
    openCLExecuteKernel(...);
    releaseTexture(I1_tex);
}
```

とあるので、bindTextureが呼べれば良さそうだけれど、この関数は opencv2/ocl/private/util.hppに書いてあって、ユーザー側には提供されていない模様。
この問題は、OpenCV 3.0で解決する。

## OpenCV 3.0でOpenCL
OpenCV 3.0では、cv::ocl::oclMatの代わりに、cv::UMatが提供されている。
これを用いると、例えば

```cpp-objdump
cv::bilateralFilter(umat_src, umat_dst, 10, 10, 10);
```

と呼び出した時に、普通にCPUで実行するか、OpenCLでデバイス上で実行するかをcv::ocl::setUseOpenCL(true/false)だけで切り替えできる。
詳しくは @dandelion1124 さんの解説を参考にしてほしい。
http://www.slideshare.net/YasuhiroYoshimura/gpgpu-dandelion1124-201301130

なお、OpenCV 3.0では、oclモジュール（opencv_ocl***.lib）が廃止されて、imgprocなどのそれぞれのモジュールに分散して収められている。
また、ヘッダファイルも変わっているので注意。

```cpp
//#include <opencv2/ocl/ocl.hpp>  // OpenCV 2.4
#include <opencv2/core/ocl.hpp> // OpenCV 3.0 ではこっち
```

#本題：OpenCV 3.0で独自カーネルOpenCL
ようやく本題ですが、OpenCV 3.0ではOpenCLまわりのインターフェイスがかなり整理されて、使いやすくなっています。
OpenCV 3.0.0beta版の現状では、まだ未実装な部分もあるため、複数デバイスの検出などがまだできないようですが、とりあえず動かすことはできました。

##環境
* Windows 8.1
* Visual Studio 2013
* OpenCV 3.0.0 beta（公式配布のバイナリをそのまま使用）
* GeForce GTX Titan Black（ドライバーバージョン：340.62）

##UMatを使った独自カーネルコード
CV_8U型のグレースケール画像をネガポジ反転するコードです。

```cpp:negaposi.cpp
#include <iostream>
#include <fstream>
#include <string>
#include <iterator>
#include <opencv2/opencv.hpp>
#include <opencv2/core/ocl.hpp>

using namespace std;

void main()
{
	if (!cv::ocl::haveOpenCL())
	{
		cout << "OpenCL is not avaiable..." << endl;
		return;
	}
	cv::ocl::Context context;
	if (!context.create(cv::ocl::Device::TYPE_GPU))
	{
		cout << "Failed creating the context..." << endl;
		return;
	}

	// In OpenCV 3.0.0 beta, only a single device is detected.
	cout << context.ndevices() << " GPU devices are detected." << endl;
	for (int i = 0; i < context.ndevices(); i++)
	{
		cv::ocl::Device device = context.device(i);
		cout << "name                 : " << device.name() << endl;
		cout << "available            : " << device.available() << endl;
		cout << "imageSupport         : " << device.imageSupport() << endl;
		cout << "OpenCL_C_Version     : " << device.OpenCL_C_Version() << endl;
		cout << endl;
	}

	// Select the first device
	cv::ocl::Device(context.device(0));

	// Transfer Mat data to the device
	cv::Mat mat_src = cv::imread("Lena.png", cv::IMREAD_GRAYSCALE);
	cv::UMat umat_src = mat_src.getUMat(cv::ACCESS_READ, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
	cv::UMat umat_dst(mat_src.size(), mat_src.type(), cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);

	// Read the OpenCL kernel code
	std::ifstream ifs("negaposi.cl");
	if (ifs.fail()) return;
	std::string kernelSource((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
	cv::ocl::ProgramSource programSource(kernelSource);

	// Compile the kernel code
	cv::String errmsg;
	cv::String buildopt = ""; // By setting "-D xxx=yyy ", we can replace xxx with yyy in the kernel
	cv::ocl::Program program = context.getProg(programSource, buildopt, errmsg);

	cv::ocl::Kernel kernel("negaposi", program);
	kernel.args(cv::ocl::KernelArg::ReadOnlyNoSize(umat_src), cv::ocl::KernelArg::ReadWrite(umat_dst));

	size_t globalThreads[3] = { mat_src.cols, mat_src.rows, 1 };
	//size_t localThreads[3] = { 16, 16, 1 };
	bool success = kernel.run(3, globalThreads, NULL, true);
	if (!success){
		cout << "Failed running the kernel..." << endl;
		return;
	}

	// Download the dst data from the device (?)
	cv::Mat mat_dst = umat_dst.getMat(cv::ACCESS_READ);

	cv::imshow("src", mat_src);
	cv::imshow("dst", mat_dst);
	cv::waitKey();
}
```

```cpp:negaposi.cl

__kernel void negaposi(
   __global uchar* src,
   int src_step, int src_offset,
   __global uchar* dst,
   int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
   int x = get_global_id(0);
   int y = get_global_id(1);
   if (x >= dst_cols) return;
   int src_index = mad24(y, src_step, x + src_offset);
   int dst_index = mad24(y, dst_step, x + dst_offset);
   dst[dst_index] = 255 - src[src_index];
};
```
**コメント**

* 現状のOpenCV 3.0.0betaでは、デバイスは1つしか列挙されないようです。（2つ目以降に検出されたデバイスは無視している）
* それゆえか、使用するデバイスの選択方法がないっぽい？（2.4ではcv::ocl::setDevice()が使えた）→ cv::ocl::Device(context.device(index)) で指定できると @dandelion1124 さんから情報いただきましたのでコードを修正しました。
* buildopt で、"-D xxx=yyy "を列挙するとカーネルコード内のxxxをyyyで置換してくれるようです。定数パラメータを数値に置き換える時などに使えます。
* MatからUMatにアップロードするときや、UMatからMatに落とすときのACCESSフラグやUSAGEフラグの設定は、これで正しいのかはよくわかってないです。あと、実は最後のところでUMatでそのまま cv::imshow("dst", umat_dst); としても動きます。
* カーネル内 mad24は mad24(x,y,z) = x*y + z という命令です。
* ローカルスレッドは、今回はNULLで省略していますが、適切に設定するだけかなり速くなります。

**ポイント：カーネル引数の渡し方と受け取り方**
自分がハマったポイントのメモ。
以下は、UMatをカーネル引数として指定している部分です。

```cpp
	kernel.args(cv::ocl::KernelArg::ReadOnlyNoSize(umat_src), cv::ocl::KernelArg::ReadWrite(umat_dst));
```
ここで、カーネルの定義を見ると

```cpp
__kernel void negaposi(
   __global uchar* src,
   int src_step, int src_offset,
   __global uchar* dst,
   int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
```
となっています。
つまり、KernelArgの作り方によって受け取る引数が変わるようで、まとめると以下のようになります。

* ReadOnly(): data_type *data, int step, int offset, int rows, int cols
* ReadOnlyNoSize(): data_type *data, int step, int offset
* PtrReadOnly(): data_type *data

stepやoffsetの情報は、データのアライメントの関係で、UMatがdenseな配列になっていないことがあるため必要になります。

さらに

```cpp
int src_index = mad24(y, src_step, x + src_offset);
```
この部分のアクセス方法は、対象とするUMatのデータ型がCV_8U（1バイト）なので大丈夫なのですが、そうでないといろいろ問題が生じます。
詳しくは、最後に書いた"image2D型の渡し方"での例を見てください。

また、Kernelに引数をセットする方法は、Kernel.args()以外にも

```cpp
int nArgs = 0;
nArgs = kernel.set(nArgs, kernelArg1);
nArgs = kernel.set(nArgs, kernelArg2);
nArgs = kernel.set(nArgs, kernelArg3);
```
と一つずつセットする方法もあります。

##int/float型の渡し方
UMat以外の通常のintやfloatやucharなどを渡す場合は、もっと簡単です。
そのまま kernel.args にintやfloat型の変数を渡せばOKです。
以下、コードを変更のあるところのみ一部抜粋。

```cpp:setTo.cpp
	...
	uchar value = 255;
	cv::ocl::Kernel kernel("setTo", program);
	kernel.args(value, cv::ocl::KernelArg::ReadWrite(umat_dst));
	...
```

```cpp:setTo.cl
__kernel void setTo( uchar value,
   __global uchar* dst,
   int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
   int x = get_global_id(0);
   int y = get_global_id(1);
   if (x >= dst_cols) return;
   int dst_index = mad24(y, dst_step, x + dst_offset);
   dst[dst_index] = value;
}
```

##image2D型の渡し方

```cpp:shift.cpp
	...
	// Transfer Mat data to the device
	cv::Mat mat_src = cv::imread("Lena.png", cv::IMREAD_GRAYSCALE);
	mat_src.convertTo(mat_src, CV_32F, 1.0 / 255);
	cv::UMat umat_src = mat_src.getUMat(cv::ACCESS_READ, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
	cv::UMat umat_dst(mat_src.size(), CV_32F, cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);

	std::ifstream ifs("shift.cl");
	if (ifs.fail()) return;
	std::string kernelSource((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
	cv::ocl::ProgramSource programSource(kernelSource);

	// Compile the kernel code
	cv::String errmsg;
	cv::String buildopt = cv::format("-D dstT=%s", cv::ocl::typeToStr(umat_dst.depth())); // "-D dstT=float"
	cv::ocl::Program program = context.getProg(programSource, buildopt, errmsg);

	cv::ocl::Image2D image(umat_src);
	float shift_x = 100.5;
	float shift_y = -50.0;
	cv::ocl::Kernel kernel("shift", program);
	kernel.args(image, shift_x, shift_y, cv::ocl::KernelArg::ReadWrite(umat_dst));
	...
```

```cpp:shift.cl
__constant sampler_t samplerLN = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
__kernel void shift(
   __global const image2d_t src,
   float shift_x,
   float shift_y,
   __global uchar* dst,
   int dst_step, int dst_offset, int dst_rows, int dst_cols)
{
   int x = get_global_id(0);
   int y = get_global_id(1);
   if (x >= dst_cols) return;
   int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset));
   __global dstT *dstf = (__global dstT *)(dst + dst_index);
   float2 coord = (float2)((float)x+0.5f+shift_x, (float)y+0.5f+shift_y);
   dstf[0] = (dstT)read_imagef(src, samplerLN, coord).x;
}
```
実行すると、画像領域外の境界を補完しながら、シフトしてくれます。
補完してくれるのは、sampler_t にCLK_ADDRESS_CLAMP_TO_EDGEオプションがあるからです。
![src.png](https://qiita-image-store.s3.amazonaws.com/0/69982/904aca4b-dae0-d155-fd89-3b50da811d89.png)　![dst.png](https://qiita-image-store.s3.amazonaws.com/0/69982/eddee61d-a96b-a2cb-5801-fd3d2af36631.png)


**ポイント：image2D引数の渡し方と受け取り方**
UMatをcv::ocl::Image2Dのコンストラクタに渡して

```cpp:
cv::ocl::Image2D image(umat_src);
```
そのまま kernel.args に image を渡せばOKです。カーネル側も image2d_t image で良い。

それから、カーネル内で read_imagef するときのcoordは、整数座標に0.5を足すことで画素の中心をサンプリングしてくれます。
そうでないと隣接画素との線形補完が発生します。
CLK_FILTER_LINEARをCLK_FILTER_NEARESTにすると、線形補完なしで最近傍画素の値を取得します。

また、read_imagefは画像が1チャンネルでも常に4チャンネル分の値をfloat4として返すので、.xを付けて最初のチャンネルの値を取得する。

**ポイント：float型のUMatのアクセス**
ちょっとめんどくさいですが、いったん uchar*型のポインタでアドレスを計算してから、対象とする型のポインタに直しています。

```cpp
   int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT), dst_offset));
   __global dstT *dstf = (__global dstT *)(dst + dst_index);
   dstf[0] = ...;
```
ちなみに dstTは、buildオプションで

```cpp
cv::String buildopt = cv::format("-D dstT=%s", cv::ocl::typeToStr(umat_dst.depth()));
```
とすることで、dstTがfloatに置換されるようになっています。

# おわりに
まだいろいろ分からないところなどあるので、間違いや新情報などあったら教えてください