OpenCV 3.0.0での独自カーネルOpenCL

  • 30
    Like
  • 0
    Comment
More than 1 year has 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に置換されるようになっています。

おわりに

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