Go to Qiita Advent Calendar 2024 Top
LoginSignup
40
39

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

More than 5 years have passed since last update.

@tackson5

OpenCV 3.0.0での独自カーネルOpenCL

Last updated at Posted at 2015-02-20

#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に置換されるようになっています。

おわりに

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

40
39
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
40
39

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?