Skip to content

Instantly share code, notes, and snippets.

@tomoaki0705

tomoaki0705/CMakeLists.txt

Created April 13, 2018 08:58
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save tomoaki0705/bee46d7a35fa220459c1c76b243759bb to your computer and use it in GitHub Desktop.
Save tomoaki0705/bee46d7a35fa220459c1c76b243759bb to your computer and use it in GitHub Desktop.
Download ZIP
sample program for OpenCL
Raw
CMakeLists.txt
# CMake minimun version
cmake_minimum_required(VERSION 2.8)
# executable name
set(TARGET_NAME OpenCVTemplate)
# solution name
project(${TARGET_NAME})
# executable
add_executable(${TARGET_NAME} main.cpp)
# find OpenCV
find_package(OpenCV REQUIRED)
if(MSVC)
add_definitions("/wd4265 /wd4266 /wd4350 /wd4365 /wd4435 /wd4514 /wd4625 /wd4626 /wd4640 /wd4668 /wd4710 /wd4819 /wd4820 /wd4946")
endif(MSVC)
# Set path and libraries OpenCV
if(OpenCV_FOUND)
# include path
include_directories(${OpenCV_INCLUDE_DIRS})
# link library path
target_link_libraries(${TARGET_NAME} ${OpenCV_LIBS})
endif(OpenCV_FOUND)
Raw
main.cpp
#include <opencv2/opencv.hpp>
#include <opencv2/core/ocl.hpp>
#include <iostream>
cv::String windowName = "OpenCV image";
const char sourceNegaPosi[] =
"__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]; "
"}; ";
const char sourceTest00[] =
"__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 && y < dst_rows) "
" { "
" int src_index = mad24(y, src_step, x + src_offset); "
" int dst_index = mad24(y, dst_step, x + dst_offset); "
" dst[dst_index] = mad24(x, y, 1); "
" }; "
"}; ";
const char sourceTest01[] =
"__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 && y < dst_rows) "
" { "
" int src_index = mad24(y, src_step, x + src_offset); "
" int dst_index = mad24(y, dst_step, 4*(x + dst_offset));"
" *(__global float*)(dst + dst_index) = (float)(x + y); "
" }; "
"}; ";
const char sourceTest02[] =
"__kernel void negaposi( \n"
"__global uchar* src, \n"
"int src_step, int src_offset, int src_rows, int src_cols, \n"
"__global uchar* dst, \n"
"int dst_step, int dst_offset, int dst_rows, int dst_cols, \n"
" float ifx, float ify) \n"
"{ \n"
" int dx = get_global_id(0); \n"
" int dy = get_global_id(1); \n"
" if (dx < dst_cols && dy < dst_rows) \n"
" { \n"
" float sx = ((dx + 0.5f) * ifx - 0.5f), sy = ((dy + 0.5f) * ify - 0.5f);\n"
" int x = floor(sx), y = floor(sy); \n"
" float u = sx - x, v = sy - y; \n"
" if (x<0) x = 0, u = 0; \n"
" if (x >= src_cols) x = src_cols - 1, u = 0; \n"
" if (y<0) y = 0, v = 0; \n"
" if (y >= src_rows) y = src_rows - 1, v = 0; \n"
" int y_ = min(y + 1, src_rows - 1); \n"
" int x_ = min(x + 1, src_cols - 1); \n"
" int dst_index = mad24(dy, dst_step, 4*(dx + dst_offset)); \n"
" *(__global int*)(dst + dst_index) = (int)(x); \n"
" } \n"
"} \n";
const char sourceResize00[] =
"__kernel void resizeLN(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols, "
" __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, "
" float ifx, float ify) "
"{ "
" int dx = get_global_id(0); "
" int dy = get_global_id(1); "
" if (dx < dst_cols && dy < dst_rows) "
" { "
" float sx = ((dx + 0.5f) * ifx - 0.5f), sy = ((dy + 0.5f) * ify - 0.5f); "
" int x = floor(sx), y = floor(sy); "
" float u = sx - x, v = sy - y; "
" if (x<0) x = 0, u = 0; "
" if (x >= src_cols) x = src_cols - 1, u = 0; "
" if (y<0) y = 0, v = 0; "
" if (y >= src_rows) y = src_rows - 1, v = 0; "
" int y_ = min(y + 1, src_rows - 1); "
" int x_ = min(x + 1, src_cols - 1); "
" "
" u = u * (1 << 11); "
" v = v * (1 << 11); "
" int U = rint(u); "
" int V = rint(v); "
" int U1 = rint((1 << 11) - u); "
" int V1 = rint((1 << 11) - v); "
" int4 data0 = convert_int4(*(__global const uchar4 *)(srcptr + mad24(y, src_step, mad24(x, (int)sizeof(uchar4), src_offset)))); "
" int4 data1 = convert_int4(*(__global const uchar4 *)(srcptr + mad24(y, src_step, mad24(x_, (int)sizeof(uchar4), src_offset)))); "
" int4 data2 = convert_int4(*(__global const uchar4 *)(srcptr + mad24(y_, src_step, mad24(x, (int)sizeof(uchar4), src_offset)))); "
" int4 data3 = convert_int4(*(__global const uchar4 *)(srcptr + mad24(y_, src_step, mad24(x_, (int)sizeof(uchar4), src_offset))));"
" "
" int4 val = "
" mul24((int4)mul24(U1, V1), data0) "
" mul24((int4)mul24(U, V1), data1) "
" mul24((int4)mul24(U1, V), data2) "
" mul24((int4)mul24(U, V), data3); "
" "
" uchar4 uval = convert_uchar4_sat((val + (1 << ((11 << 1) - 1))) >> (11 << 1)); "
" "
" *(__global uchar4 *)(dstptr + mad24(dy, dst_step, mad24(dx, (int)sizeof(uchar4), dst_offset))) = uval; "
" } "
"} ";
int main(int argc, char**argv)
{
cv::CommandLineParser parser(argc, argv, "{@input||}{help h|false|show this message}{type t|0|kenel number}");
if (parser.has("@input") && parser.get<bool>("help") == false)
{
if (!cv::ocl::haveOpenCL())
{
std::cout << "OpenCL is not avaiable..." << std::endl;
return 0;
}
cv::String filename = parser.get<cv::String>("@input");
std::cout << filename << std::endl;
cv::UMat image = cv::imread(filename, cv::IMREAD_GRAYSCALE).getUMat(cv::ACCESS_READ, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
if (image.empty())
{
std::cerr << "Failed to open " << filename << std::endl;
return 1;
}
cv::imshow(windowName, image);
cv::ocl::Context context;
if (!context.create(cv::ocl::Device::TYPE_GPU))
{
std::cout << "Failed creating the context..." << std::endl;
return 2;
}
std::cout << context.ndevices() << " GPU devices are detected." << std::endl;
for (int i = 0; i < context.ndevices(); i++)
{
cv::ocl::Device device = context.device(i);
std::cout << "name : " << device.name() << std::endl;
std::cout << "available : " << device.available() << std::endl;
std::cout << "imageSupport : " << device.imageSupport() << std::endl;
std::cout << "OpenCL_C_Version : " << device.OpenCL_C_Version() << std::endl;
std::cout << std::endl;
}
// Select the first device
cv::ocl::Device(context.device(0));
bool success = false;
cv::UMat umat_dst;
switch (parser.get<int>("type"))
{
default:
{
// Transfer Mat data to the device
cv::UMat umat_src = image;
umat_dst = cv::UMat(umat_src.size(), umat_src.type(), cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
// Read the OpenCL kernel code
std::string kernelSource(sourceNegaPosi);
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] = { umat_src.cols, umat_src.rows, 1 };
//size_t localThreads[3] = { 16, 16, 1 };
success = kernel.run(3, globalThreads, NULL, true);
}
break;
case 0:
{
// Transfer Mat data to the device
cv::UMat umat_src = image;
umat_dst = cv::UMat(umat_src.size(), umat_src.type(), cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
// Read the OpenCL kernel code
std::string kernelSource(sourceTest00);
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] = { umat_src.cols, umat_src.rows, 1 };
//size_t localThreads[3] = { 16, 16, 1 };
success = kernel.run(3, globalThreads, NULL, true);
}
break;
case 1:
{
// Transfer Mat data to the device
cv::UMat umat_src = image;
umat_dst = cv::UMat(cv::Size(umat_src.cols*4, umat_src.rows), umat_src.type(), cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
// Read the OpenCL kernel code
std::string kernelSource(sourceTest01);
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] = { umat_src.cols, umat_src.rows, 1 };
//size_t localThreads[3] = { 16, 16, 1 };
success = kernel.run(3, globalThreads, NULL, true);
}
break;
case 2:
{
// Transfer Mat data to the device
cv::UMat umat_src = image;
umat_dst = cv::UMat(cv::Size(umat_src.cols * 4, umat_src.rows), umat_src.type(), cv::ACCESS_WRITE, cv::USAGE_ALLOCATE_DEVICE_MEMORY);
// Read the OpenCL kernel code
std::string kernelSource(sourceTest02);
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::ReadOnly(umat_src), cv::ocl::KernelArg::ReadWrite(umat_dst), (float)1.5, (float)2.0);
size_t globalThreads[3] = { umat_src.cols, umat_src.rows, 1 };
//size_t localThreads[3] = { 16, 16, 1 };
success = kernel.run(3, globalThreads, NULL, true);
}
break;
}
if (!success){
std::cout << "Failed running the kernel..." << std::endl;
return 3;
}
// Download the dst data from the device (?)
cv::Mat mat_dst = umat_dst.getMat(cv::ACCESS_READ);
cv::imshow("src", image);
cv::imshow("dst", mat_dst);
cv::waitKey();
}
else
{
parser.printMessage();
}
return 0;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment