Skip to content

Instantly share code, notes, and snippets.

@atinfinity
Last active September 14, 2023 07:26
Show Gist options
  • Save atinfinity/8c25c8fb1b3708aa0944 to your computer and use it in GitHub Desktop.
Save atinfinity/8c25c8fb1b3708aa0944 to your computer and use it in GitHub Desktop.
sample code to execute custom OpenCL kernel on OpenCV 2.4.9
#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";
cv::Mat mat_src = cv::imread("lena.jpg", cv::IMREAD_GRAYSCALE);
cv::Mat mat_dst;
if(mat_src.empty())
{
std::cerr << "Failed to open image file." << std::endl;
return -1;
}
unsigned int channels = mat_src.channels();
unsigned int depth = mat_src.depth();
cv::ocl::oclMat ocl_src(mat_src);
cv::ocl::oclMat ocl_dst(mat_src.size(), mat_src.type());
cv::ocl::ProgramSource program("negaposi", KernelSource);
std::size_t globalThreads[3]={ocl_src.rows * ocl_src.step, 1, 1};
std::vector<std::pair<size_t , const void *> > args;
args.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_src.data ));
args.push_back( std::make_pair( sizeof(cl_mem), (void *) &ocl_dst.data ));
cv::ocl::openCLExecuteKernelInterop(cv::ocl::Context::getContext(),
program, "negaposi", globalThreads, NULL, args, channels, depth, NULL);
ocl_dst.download(mat_dst);
cv::namedWindow("mat_src");
cv::namedWindow("mat_dst");
cv::imshow("mat_src", mat_src);
cv::imshow("mat_dst", mat_dst);
cv::waitKey(0);
cv::destroyAllWindows();
return 0;
}
@ashokbugude
Copy link

Hi,
Can I pls know the source for 'negaposi' file ?
Thanks

@jnodev
Copy link

jnodev commented Dec 5, 2016

Hi @ashokbugude
if I am not mistaken "negaposi" is just a name and the source code is in the KernelSource string.
Cheers

@waqarmumtaz369
Copy link

Yes. @jnodev is right. "negaposi" is the name of kernel function because one kernel source can consist of more than one kernel functions.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment