Edited at

OpenCV 3.0.0での独自カーネルOpenCL

More than 3 years have passed since last update.


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の状態でビルド済みらしいので

#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関数を見ると、


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が提供されている。

これを用いると、例えば

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などのそれぞれのモジュールに分散して収められている。

また、ヘッダファイルも変わっているので注意。

//#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型のグレースケール画像をネガポジ反転するコードです。


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();
}



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をカーネル引数として指定している部分です。

    kernel.args(cv::ocl::KernelArg::ReadOnlyNoSize(umat_src), cv::ocl::KernelArg::ReadWrite(umat_dst));

ここで、カーネルの定義を見ると

__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な配列になっていないことがあるため必要になります。

さらに

int src_index = mad24(y, src_step, x + src_offset);

この部分のアクセス方法は、対象とするUMatのデータ型がCV_8U(1バイト)なので大丈夫なのですが、そうでないといろいろ問題が生じます。

詳しくは、最後に書いた"image2D型の渡し方"での例を見てください。

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

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です。

以下、コードを変更のあるところのみ一部抜粋。


setTo.cpp

    ...

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


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型の渡し方


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));
...



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 dst.png

ポイント:image2D引数の渡し方と受け取り方

UMatをcv::ocl::Image2Dのコンストラクタに渡して

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*型のポインタでアドレスを計算してから、対象とする型のポインタに直しています。

   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オプションで

cv::String buildopt = cv::format("-D dstT=%s", cv::ocl::typeToStr(umat_dst.depth()));

とすることで、dstTがfloatに置換されるようになっています。


おわりに

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