sample code to execute custom OpenCL kernel on OpenCV 2.4.9

oclMat_custom_kernel2.cpp

#include <opencv2/core/core.hpp>
#include <opencv2/highgui/highgui.hpp>
#include <opencv2/ocl/ocl.hpp>

// cl_mem構造体を参照するためにインクルード
#if defined __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif

#include <iostream>

int main(int argc, const char** argv)
{
    cv::ocl::DevicesInfo devInfo;
    int res = cv::ocl::getOpenCLDevices(devInfo);

    if(res == 0)
    {
        std::cerr << "There is no OPENCL Here !" << std::endl;
        return -1;
    }
    else
    {
        for(unsigned int i = 0 ; i < devInfo.size() ; ++i)
        {
            std::cout << "Device : " << devInfo[i]->deviceName << " is present" << std::endl;
        }
    }

    cv::ocl::setDevice(devInfo[0]);        // select device to use
    std::cout << CV_VERSION_EPOCH << "." << CV_VERSION_MAJOR << "." << CV_VERSION_MINOR << std::endl;

    // グレースケール用、カラー用のカーネル
    const char *KernelSource = "\n" \
    "__kernel void negaposi_C1_D0(               \n" \
    "   __global uchar* input,                   \n" \
    "   __global uchar* output)                  \n" \
    "{                                           \n" \
    "   int i = get_global_id(0);                \n" \
    "   output[i] = 255 - input[i];              \n" \
    "}                                           \n" \
    "__kernel void negaposi_C3_D0(               \n" \
    "   __global uchar* input,                   \n" \
    "   __global uchar* output)                  \n" \
    "{                                           \n" \
    "   int i = get_global_id(0);                \n" \
    "   output[i] = 255 - input[i];              \n" \
    "}\n";

    cv::ocl::ProgramSource program("negaposi", KernelSource);

    // グレースケール用の処理
    cv::Mat mat_src_gray = cv::imread("lena.jpg", cv::IMREAD_GRAYSCALE);
    cv::Mat mat_dst_gray;
    if(mat_src_gray.empty())
    {
        std::cerr << "Failed to open image file." << std::endl;
        return -1; 
    }

    cv::ocl::oclMat ocl_src_gray(mat_src_gray);
    cv::ocl::oclMat ocl_dst_gray(mat_src_gray.size(), mat_src_gray.type());
    std::size_t globalThreads_gray[3]={ocl_src_gray.rows * ocl_src_gray.step, 1, 1};
    std::vector<std::pair<size_t , const void *> > args_gray;
    args_gray.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_src_gray.data ));
    args_gray.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_dst_gray.data ));
    cv::ocl::openCLExecuteKernelInterop(cv::ocl::Context::getContext(), 
        program, "negaposi", globalThreads_gray, NULL, args_gray, mat_src_gray.channels(), mat_src_gray.depth(), NULL);
    ocl_dst_gray.download(mat_dst_gray);

    // カラー用の処理
    cv::Mat mat_src_color = cv::imread("lena.jpg", cv::IMREAD_COLOR);
    cv::Mat mat_dst_color;
    if(mat_src_color.empty())
    {
        std::cerr << "Failed to open image file." << std::endl;
        return -1; 
    }

    cv::ocl::oclMat ocl_src_color(mat_src_color);
    cv::ocl::oclMat ocl_dst_color(mat_src_color.size(), mat_src_color.type());
    std::size_t globalThreads_color[3]={ocl_src_color.rows * ocl_src_color.step, 1, 1};
    std::vector<std::pair<size_t , const void *> > args_color;
    args_color.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_src_color.data ));
    args_color.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_dst_color.data ));
    cv::ocl::openCLExecuteKernelInterop(cv::ocl::Context::getContext(), 
        program, "negaposi", globalThreads_color, NULL, args_color, mat_src_color.channels(), mat_src_color.depth(), NULL);
    ocl_dst_color.download(mat_dst_color);

    cv::namedWindow("mat_dst_gray");
    cv::namedWindow("mat_dst_color");
    cv::imshow("mat_dst_gray", mat_dst_gray);
    cv::imshow("mat_dst_color", mat_dst_color);
    cv::waitKey(0);
    cv::destroyAllWindows();

    return 0;
}