Skip to content

Commit 1625ffa

Browse files
alalekalalek
committed
ocl: opencl_custom_kernel.cpp example
1 parent 15b909e commit 1625ffa

File tree

1 file changed

+160
-0
lines changed

1 file changed

+160
-0
lines changed

‎samples/tapi/opencl_custom_kernel.cpp

Lines changed: 160 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,160 @@
1+
#include "opencv2/core.hpp"
2+
#include "opencv2/core/ocl.hpp"
3+
#include "opencv2/highgui.hpp"
4+
#include "opencv2/imgcodecs.hpp"
5+
#include "opencv2/imgproc.hpp"
6+
7+
#include <iostream>
8+
9+
using namespace std;
10+
using namespace cv;
11+
12+
static const char* opencl_kernel_src =
13+
"__kernel void magnutude_filter_8u(\n"
14+
" __global const uchar* src, int src_step, int src_offset,\n"
15+
" __global uchar* dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,\n"
16+
" float scale)\n"
17+
"{\n"
18+
" int x = get_global_id(0);\n"
19+
" int y = get_global_id(1);\n"
20+
" if (x < dst_cols && y < dst_rows)\n"
21+
" {\n"
22+
" int dst_idx = y * dst_step + x + dst_offset;\n"
23+
" if (x > 0 && x < dst_cols - 1 && y > 0 && y < dst_rows - 2)\n"
24+
" {\n"
25+
" int src_idx = y * src_step + x + src_offset;\n"
26+
" int dx = (int)src[src_idx]*2 - src[src_idx - 1] - src[src_idx + 1];\n"
27+
" int dy = (int)src[src_idx]*2 - src[src_idx - 1*src_step] - src[src_idx + 1*src_step];\n"
28+
" dst[dst_idx] = convert_uchar_sat(sqrt((float)(dx*dx + dy*dy)) * scale);\n"
29+
" }\n"
30+
" else\n"
31+
" {\n"
32+
" dst[dst_idx] = 0;\n"
33+
" }\n"
34+
" }\n"
35+
"}\n";
36+
37+
int main(int argc, char** argv)
38+
{
39+
const char* keys =
40+
"{ i input | | specify input image }"
41+
"{ h help | | print help message }";
42+
43+
cv::CommandLineParser args(argc, argv, keys);
44+
if (args.has("help"))
45+
{
46+
cout << "Usage : " << argv[0] << " [options]" << endl;
47+
cout << "Available options:" << endl;
48+
args.printMessage();
49+
return EXIT_SUCCESS;
50+
}
51+
52+
cv::ocl::Context ctx = cv::ocl::Context::getDefault();
53+
if (!ctx.ptr())
54+
{
55+
cerr << "OpenCL is not available" << endl;
56+
return 1;
57+
}
58+
cv::ocl::Device device = cv::ocl::Device::getDefault();
59+
if (!device.compilerAvailable())
60+
{
61+
cerr << "OpenCL compiler is not available" << endl;
62+
return 1;
63+
}
64+
65+
66+
UMat src;
67+
{
68+
string image_file = args.get<string>("i");
69+
if (!image_file.empty())
70+
{
71+
Mat image = imread(image_file);
72+
if (image.empty())
73+
{
74+
cout << "error read image: " << image_file << endl;
75+
return 1;
76+
}
77+
cvtColor(image, src, COLOR_BGR2GRAY);
78+
}
79+
else
80+
{
81+
Mat frame(cv::Size(640, 480), CV_8U, Scalar::all(128));
82+
Point p(frame.cols / 2, frame.rows / 2);
83+
line(frame, Point(0, frame.rows / 2), Point(frame.cols, frame.rows / 2), 1);
84+
circle(frame, p, 200, Scalar(32, 32, 32), 8, LINE_AA);
85+
string str = "OpenCL";
86+
int baseLine = 0;
87+
Size box = getTextSize(str, FONT_HERSHEY_COMPLEX, 2, 5, &baseLine);
88+
putText(frame, str, Point((frame.cols - box.width) / 2, (frame.rows - box.height) / 2 + baseLine),
89+
FONT_HERSHEY_COMPLEX, 2, Scalar(255, 255, 255), 5, LINE_AA);
90+
frame.copyTo(src);
91+
}
92+
}
93+
94+
95+
cv::String module_name; // empty to disable OpenCL cache
96+
97+
{
98+
cout << "OpenCL program source: " << endl;
99+
cout << "======================================================================================================" << endl;
100+
cout << opencl_kernel_src << endl;
101+
cout << "======================================================================================================" << endl;
102+
//! [Define OpenCL program source]
103+
cv::ocl::ProgramSource source(module_name, "simple", opencl_kernel_src, "");
104+
//! [Define OpenCL program source]
105+
106+
//! [Compile/build OpenCL for current OpenCL device]
107+
cv::String errmsg;
108+
cv::ocl::Program program(source, "", errmsg);
109+
if (program.ptr() == NULL)
110+
{
111+
cerr << "Can't compile OpenCL program:" << endl << errmsg << endl;
112+
return 1;
113+
}
114+
//! [Compile/build OpenCL for current OpenCL device]
115+
116+
if (!errmsg.empty())
117+
{
118+
cout << "OpenCL program build log:" << endl << errmsg << endl;
119+
}
120+
121+
//! [Get OpenCL kernel by name]
122+
cv::ocl::Kernel k("magnutude_filter_8u", program);
123+
if (k.empty())
124+
{
125+
cerr << "Can't get OpenCL kernel" << endl;
126+
return 1;
127+
}
128+
//! [Get OpenCL kernel by name]
129+
130+
UMat result(src.size(), CV_8UC1);
131+
132+
//! [Define kernel parameters and run]
133+
size_t globalSize[2] = {(size_t)src.cols, (size_t)src.rows};
134+
size_t localSize[2] = {8, 8};
135+
bool executionResult = k
136+
.args(
137+
cv::ocl::KernelArg::ReadOnlyNoSize(src), // size is not used (similar to 'dst' size)
138+
cv::ocl::KernelArg::WriteOnly(result),
139+
(float)2.0
140+
)
141+
.run(2, globalSize, localSize, true);
142+
if (!executionResult)
143+
{
144+
cerr << "OpenCL kernel launch failed" << endl;
145+
return 1;
146+
}
147+
//! [Define kernel parameters and run]
148+
149+
imshow("Source", src);
150+
imshow("Result", result);
151+
152+
for (;;)
153+
{
154+
int key = waitKey();
155+
if (key == 27/*ESC*/ || key == 'q' || key == 'Q')
156+
break;
157+
}
158+
}
159+
return 0;
160+
}

0 commit comments

Comments
 (0)