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
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
|
// 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 "../test_precomp.hpp"
#include <opencv2/core/ocl.hpp>
namespace opencv_test { namespace {
static void testOpenCLKernel(cv::ocl::Kernel& k)
{
ASSERT_FALSE(k.empty());
cv::UMat src(cv::Size(4096, 2048), CV_8UC1, cv::Scalar::all(100));
cv::UMat dst(src.size(), CV_8UC1);
size_t globalSize[2] = {(size_t)src.cols, (size_t)src.rows};
size_t localSize[2] = {8, 8};
int64 kernel_time = k.args(
cv::ocl::KernelArg::ReadOnlyNoSize(src), // size is not used (similar to 'dst' size)
cv::ocl::KernelArg::WriteOnly(dst),
(int)5
).runProfiling(2, globalSize, localSize);
ASSERT_GE(kernel_time, (int64)0);
std::cout << "Kernel time: " << (kernel_time * 1e-6) << " ms" << std::endl;
cv::Mat res, reference(src.size(), CV_8UC1, cv::Scalar::all(105));
dst.copyTo(res);
EXPECT_EQ(0, cvtest::norm(reference, res, cv::NORM_INF));
}
TEST(OpenCL, support_binary_programs)
{
cv::ocl::Context ctx = cv::ocl::Context::getDefault();
if (!ctx.ptr())
{
throw cvtest::SkipTestException("OpenCL is not available");
}
cv::ocl::Device device = cv::ocl::Device::getDefault();
if (!device.compilerAvailable())
{
throw cvtest::SkipTestException("OpenCL compiler is not available");
}
std::vector<char> program_binary_code;
cv::String module_name; // empty to disable OpenCL cache
{ // Generate program binary from OpenCL C source
static const char* opencl_kernel_src =
"__kernel void test_kernel(__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"
" int c)\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 src_idx = y * src_step + x + src_offset;\n"
" int dst_idx = y * dst_step + x + dst_offset;\n"
" dst[dst_idx] = src[src_idx] + c;\n"
" }\n"
"}\n";
cv::ocl::ProgramSource src(module_name, "simple", opencl_kernel_src, "");
cv::String errmsg;
cv::ocl::Program program(src, "", errmsg);
ASSERT_TRUE(program.ptr() != NULL);
cv::ocl::Kernel k("test_kernel", program);
EXPECT_FALSE(k.empty());
program.getBinary(program_binary_code);
std::cout << "Program binary size: " << program_binary_code.size() << " bytes" << std::endl;
}
cv::ocl::Kernel k;
{ // Load program from binary (without sources)
ASSERT_FALSE(program_binary_code.empty());
cv::ocl::ProgramSource src = cv::ocl::ProgramSource::fromBinary(module_name, "simple_binary", (uchar*)&program_binary_code[0], program_binary_code.size(), "");
cv::String errmsg;
cv::ocl::Program program(src, "", errmsg);
ASSERT_TRUE(program.ptr() != NULL);
k.create("test_kernel", program);
}
testOpenCLKernel(k);
}
TEST(OpenCL, support_SPIR_programs)
{
cv::ocl::Context ctx = cv::ocl::Context::getDefault();
if (!ctx.ptr())
{
throw cvtest::SkipTestException("OpenCL is not available");
}
cv::ocl::Device device = cv::ocl::Device::getDefault();
if (!device.isExtensionSupported("cl_khr_spir"))
{
throw cvtest::SkipTestException("'cl_khr_spir' extension is not supported by OpenCL device");
}
std::vector<char> program_binary_code;
cv::String fname = cv::format("test_kernel.spir%d", device.addressBits());
std::string full_path = cvtest::findDataFile(std::string("opencl/") + fname);
{
std::fstream f(full_path.c_str(), std::ios::in|std::ios::binary);
ASSERT_TRUE(f.is_open());
size_t pos = (size_t)f.tellg();
f.seekg(0, std::fstream::end);
size_t fileSize = (size_t)f.tellg();
std::cout << "Program SPIR size: " << fileSize << " bytes" << std::endl;
f.seekg(pos, std::fstream::beg);
program_binary_code.resize(fileSize);
f.read(&program_binary_code[0], fileSize);
ASSERT_FALSE(f.fail());
}
cv::String module_name; // empty to disable OpenCL cache
cv::ocl::Kernel k;
{ // Load program from SPIR format
ASSERT_FALSE(program_binary_code.empty());
cv::ocl::ProgramSource src = cv::ocl::ProgramSource::fromSPIR(module_name, "simple_spir", (uchar*)&program_binary_code[0], program_binary_code.size(), "");
cv::String errmsg;
cv::ocl::Program program(src, "", errmsg);
if (program.ptr() == NULL && device.isAMD())
{
// https://community.amd.com/t5/opencl/spir-support-in-new-drivers-lost/td-p/170165
throw cvtest::SkipTestException("Bypass AMD OpenCL runtime bug: 'cl_khr_spir' extension is declared, but it doesn't really work");
}
ASSERT_TRUE(program.ptr() != NULL);
k.create("test_kernel", program);
}
testOpenCLKernel(k);
}
TEST(OpenCL, image2Dcount_regression_19334)
{
cv::ocl::Context ctx = cv::ocl::Context::getDefault();
if (!ctx.ptr())
{
throw cvtest::SkipTestException("OpenCL is not available");
}
cv::ocl::Device device = cv::ocl::Device::getDefault();
if (!device.compilerAvailable())
{
throw cvtest::SkipTestException("OpenCL compiler is not available");
}
std::string module_name; // empty to disable OpenCL cache
static const char* opencl_kernel_src =
"__kernel void test_kernel(int a,\n"
" __global const uchar* src0, int src0_step, int src0_offset, int src0_rows, int src0_cols,\n"
" __global const uchar* src1, int src1_step, int src1_offset, int src1_rows, int src1_cols,\n"
" __global const uchar* src2, int src2_step, int src2_offset, int src2_rows, int src2_cols,\n"
" __read_only image2d_t image)\n"
"{\n"
"}";
cv::ocl::ProgramSource src(module_name, "test_opencl_image_arg", opencl_kernel_src, "");
cv::String errmsg;
cv::ocl::Program program(src, "", errmsg);
ASSERT_TRUE(program.ptr() != NULL);
cv::ocl::Kernel k("test_kernel", program);
ASSERT_FALSE(k.empty());
std::vector<UMat> images(4);
for (size_t i = 0; i < images.size(); ++i)
images[i] = UMat(10, 10, CV_8UC1);
cv::ocl::Image2D image;
try
{
cv::ocl::Image2D image_(images.back());
image = image_;
}
catch (const cv::Exception&)
{
throw cvtest::SkipTestException("OpenCL images are not supported");
}
int nargs = 0;
int a = 0;
nargs = k.set(nargs, a);
ASSERT_EQ(1, nargs);
nargs = k.set(nargs, images[0]);
ASSERT_EQ(6, nargs);
nargs = k.set(nargs, images[1]);
ASSERT_EQ(11, nargs);
nargs = k.set(nargs, images[2]);
ASSERT_EQ(16, nargs);
// do not throw (issue of #19334)
ASSERT_NO_THROW(nargs = k.set(nargs, image));
ASSERT_EQ(17, nargs);
// allow to replace image argument if kernel is not running
UMat image2(10, 10, CV_8UC1);
ASSERT_NO_THROW(nargs = k.set(16, cv::ocl::Image2D(image2)));
ASSERT_EQ(17, nargs);
}
TEST(OpenCL, move_construct_assign)
{
cv::ocl::Context ctx1 = cv::ocl::Context::getDefault();
if (!ctx1.ptr())
{
throw cvtest::SkipTestException("OpenCL is not available");
}
void* const ctx_ptr = ctx1.ptr();
cv::ocl::Context ctx2(std::move(ctx1));
ASSERT_EQ(ctx1.ptr(), nullptr);
ASSERT_EQ(ctx2.ptr(), ctx_ptr);
cv::ocl::Context ctx3 = std::move(ctx2);
ASSERT_EQ(ctx2.ptr(), nullptr);
ASSERT_EQ(ctx3.ptr(), ctx_ptr);
cv::ocl::Platform pl1 = cv::ocl::Platform::getDefault();
void* const pl_ptr = pl1.ptr();
cv::ocl::Platform pl2(std::move(pl1));
ASSERT_EQ(pl1.ptr(), nullptr);
ASSERT_EQ(pl2.ptr(), pl_ptr);
cv::ocl::Platform pl3 = std::move(pl2);
ASSERT_EQ(pl2.ptr(), nullptr);
ASSERT_EQ(pl3.ptr(), pl_ptr);
std::vector<cv::ocl::PlatformInfo> platformInfos;
cv::ocl::getPlatfomsInfo(platformInfos);
const cv::String pi_name = platformInfos[0].name();
cv::ocl::PlatformInfo pinfo2(std::move(platformInfos[0]));
ASSERT_EQ(platformInfos[0].name(), cv::String());
ASSERT_EQ(pinfo2.name(), pi_name);
cv::ocl::PlatformInfo pinfo3 = std::move(pinfo2);
ASSERT_EQ(pinfo2.name(), cv::String());
ASSERT_EQ(pinfo3.name(), pi_name);
cv::ocl::Queue q1 = cv::ocl::Queue::getDefault();
void* const q_ptr = q1.ptr();
cv::ocl::Queue q2(std::move(q1));
ASSERT_EQ(q1.ptr(), nullptr);
ASSERT_EQ(q2.ptr(), q_ptr);
cv::ocl::Queue q3 = std::move(q2);
ASSERT_EQ(q2.ptr(), nullptr);
ASSERT_EQ(q3.ptr(), q_ptr);
cv::ocl::Device d1 = cv::ocl::Device::getDefault();
if (!d1.compilerAvailable())
{
throw cvtest::SkipTestException("OpenCL compiler is not available");
}
void* const d_ptr = d1.ptr();
cv::ocl::Device d2(std::move(d1));
ASSERT_EQ(d1.ptr(), nullptr);
ASSERT_EQ(d2.ptr(), d_ptr);
cv::ocl::Device d3 = std::move(d2);
ASSERT_EQ(d2.ptr(), nullptr);
ASSERT_EQ(d3.ptr(), d_ptr);
if (d3.imageSupport()) {
cv::UMat umat1 = cv::UMat::ones(640, 480, CV_32FC1);
cv::ocl::Image2D img1(umat1);
void *const img_ptr = img1.ptr();
cv::ocl::Image2D img2(std::move(img1));
ASSERT_EQ(img1.ptr(), nullptr);
ASSERT_EQ(img2.ptr(), img_ptr);
cv::ocl::Image2D img3 = std::move(img2);
ASSERT_EQ(img2.ptr(), nullptr);
ASSERT_EQ(img3.ptr(), img_ptr);
}
static const char* opencl_kernel_src =
"__kernel void test_kernel(__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"
" int c)\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 src_idx = y * src_step + x + src_offset;\n"
" int dst_idx = y * dst_step + x + dst_offset;\n"
" dst[dst_idx] = src[src_idx] + c;\n"
" }\n"
"}\n";
cv::String module_name; // empty to disable OpenCL cache
cv::ocl::ProgramSource ps1(module_name, "move_construct_assign", opencl_kernel_src, "");
cv::ocl::ProgramSource::Impl* const ps_ptr = ps1.getImpl();
cv::ocl::ProgramSource ps2(std::move(ps1));
ASSERT_EQ(ps1.getImpl(), nullptr);
ASSERT_EQ(ps2.getImpl(), ps_ptr);
cv::ocl::ProgramSource ps3 = std::move(ps2);
ASSERT_EQ(ps2.getImpl(), nullptr);
ASSERT_EQ(ps3.getImpl(), ps_ptr);
cv::String errmsg;
cv::ocl::Program prog1(ps3, "", errmsg);
void* const prog_ptr = prog1.ptr();
ASSERT_NE(prog_ptr, nullptr);
cv::ocl::Program prog2(std::move(prog1));
ASSERT_EQ(prog1.ptr(), nullptr);
ASSERT_EQ(prog2.ptr(), prog_ptr);
cv::ocl::Program prog3 = std::move(prog2);
ASSERT_EQ(prog2.ptr(), nullptr);
ASSERT_EQ(prog3.ptr(), prog_ptr);
cv::ocl::Kernel k1("test_kernel", prog3);
void* const k_ptr = k1.ptr();
ASSERT_NE(k_ptr, nullptr);
cv::ocl::Kernel k2(std::move(k1));
ASSERT_EQ(k1.ptr(), nullptr);
ASSERT_EQ(k2.ptr(), k_ptr);
cv::ocl::Kernel k3 = std::move(k2);
ASSERT_EQ(k2.ptr(), nullptr);
ASSERT_EQ(k3.ptr(), k_ptr);
testOpenCLKernel(k3);
}
}} // namespace
|
