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