1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #ifndef TEST_CONFORMANCE_CLCPP_API_TEST_CTORS_HPP
17 #define TEST_CONFORMANCE_CLCPP_API_TEST_CTORS_HPP
18
19 #include <vector>
20 #include <limits>
21 #include <algorithm>
22 #include <numeric>
23
24 #include "../common.hpp"
25
26 // TEST 1
27 // Verify that constructors are executed before any kernel is executed.
28 // Verify that when present, multiple constructors are executed. The order between
29 // constructors is undefined, but they should all execute.
30
31 // -----------------------------------------------------------------------------------
32 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
33 // -----------------------------------------------------------------------------------
34 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
35 const char * kernel_test_ctors_executed =
36 "__kernel void test_ctors_executed(global uint *output)\n"
37 "{\n"
38 " ulong gid = get_global_id(0);\n"
39 " output[gid] = 0;\n"
40 "}\n"
41 ;
42 const char * kernel_test_ctors_executed_multiple_ctors =
43 "__kernel void test_ctors_executed_multiple_ctors(global uint *output)\n"
44 "{\n"
45 " ulong gid = get_global_id(0);\n"
46 " output[gid] = 0;\n"
47 "}\n"
48 ;
49 #else
50 const char * kernel_test_ctors_executed =
51 "#include <opencl_memory>\n"
52 "#include <opencl_work_item>\n"
53 "using namespace cl;\n"
54 "struct ctor_test_class {\n"
55 // non-trivial ctor
56 " ctor_test_class(int y) { x = y;};\n"
57 " int x;\n"
58 "};\n"
59 // global scope program variable
60 "ctor_test_class global_var(int(0xbeefbeef));\n"
61 "__kernel void test_ctors_executed(global_ptr<uint[]> output)\n"
62 "{\n"
63 " ulong gid = get_global_id(0);\n"
64 " int result = 0;\n"
65 " if(global_var.x != int(0xbeefbeef)) result = 1;\n"
66 " output[gid] = result;\n"
67 "}\n"
68 ;
69 const char * kernel_test_ctors_executed_multiple_ctors =
70 "#include <opencl_memory>\n"
71 "#include <opencl_work_item>\n"
72 "#include <opencl_limits>\n"
73 "using namespace cl;\n"
74 "template<class T>\n"
75 "struct ctor_test_class {\n"
76 // non-trivial ctor
77 " ctor_test_class(T y) { x = y;};\n"
78 " T x;\n"
79 "};\n"
80 // global scope program variables
81 "ctor_test_class<int> global_var0(int(0xbeefbeef));\n"
82 "ctor_test_class<uint> global_var1(uint(0xbeefbeefU));\n"
83 "ctor_test_class<float> global_var2(float(FLT_MAX));\n"
84 "__kernel void test_ctors_executed_multiple_ctors(global_ptr<uint[]> output)\n"
85 "{\n"
86 " ulong gid = get_global_id(0);\n"
87 " int result = 0;\n"
88 " if(global_var0.x != int(0xbeefbeef)) result = 1;\n"
89 " if(global_var1.x != uint(0xbeefbeefU)) result = 1;\n"
90 " if(global_var2.x != float(FLT_MAX)) result = 1;\n"
91 " output[gid] = result;\n"
92 "}\n"
93 ;
94 #endif
95
test_ctors_execution(cl_device_id device,cl_context context,cl_command_queue queue,int count,std::string kernel_name,const char * kernel_source)96 int test_ctors_execution(cl_device_id device,
97 cl_context context,
98 cl_command_queue queue,
99 int count,
100 std::string kernel_name,
101 const char * kernel_source)
102 {
103 int error = CL_SUCCESS;
104
105 cl_mem output_buffer;
106 cl_program program;
107 cl_kernel kernel;
108
109 size_t dim = 1;
110 size_t work_size[1];
111 // -----------------------------------------------------------------------------------
112 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
113 // -----------------------------------------------------------------------------------
114 // Only OpenCL C++ to SPIR-V compilation
115 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
116 error = create_opencl_kernel(context, &program, &kernel, kernel_source, kernel_name);
117 RETURN_ON_ERROR(error)
118 return error;
119 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
120 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
121 error = create_opencl_kernel(context, &program, &kernel, kernel_source, kernel_name, "", false);
122 RETURN_ON_ERROR(error)
123 // Normal run
124 #else
125 error = create_opencl_kernel(context, &program, &kernel, kernel_source, kernel_name);
126 RETURN_ON_ERROR(error)
127 #endif
128
129 // host vector, size == count, output[0...count-1] == 1
130 std::vector<cl_uint> output(count, cl_uint(1));
131 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
132 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
133
134 error = clEnqueueWriteBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_uint) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);
135 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
136
137 error = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
138 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
139
140 work_size[0] = output.size();
141 error = clEnqueueNDRangeKernel(queue, kernel, dim, NULL, work_size, NULL, 0, NULL, NULL);
142 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
143
144 error = clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_uint) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);
145 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
146
147 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
148 if(sum != 0)
149 {
150 error = -1;
151 CHECK_ERROR_MSG(error, "Test %s failed.", kernel_name.c_str());
152 }
153
154 clReleaseMemObject(output_buffer);
155 clReleaseKernel(kernel);
156 clReleaseProgram(program);
157 return error;
158 }
159
AUTO_TEST_CASE(test_global_scope_ctors_executed)160 AUTO_TEST_CASE(test_global_scope_ctors_executed)
161 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
162 {
163 int error = CL_SUCCESS;
164 int local_error = CL_SUCCESS;
165
166 local_error = test_ctors_execution(
167 device, context, queue, count,
168 "test_ctors_executed", kernel_test_ctors_executed
169 );
170 CHECK_ERROR(local_error);
171 error |= local_error;
172
173 local_error = test_ctors_execution(
174 device, context, queue, count,
175 "test_ctors_executed_multiple_ctors", kernel_test_ctors_executed_multiple_ctors
176 );
177 CHECK_ERROR(local_error);
178 error |= local_error;
179
180 if(error != CL_SUCCESS)
181 {
182 return -1;
183 }
184 return error;
185 }
186
187 // TEST 2
188 // Verify that constructors are only executed once when multiple kernels from a program are executed.
189
190 // How: The first kernel (test_ctors_executed_once_set) is run once. It changes values of program scope
191 // variables, then the second kernel is run multiple times, each time verifying that global variables
192 // have correct values (the second kernel should observe the values assigned by the first kernel, not
193 // by the constructors).
194
195 // -----------------------------------------------------------------------------------
196 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
197 // -----------------------------------------------------------------------------------
198 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
199 const char * program_test_ctors_executed_once =
200 "__kernel void test_ctors_executed_once_set()\n"
201 "{\n"
202 "}\n"
203 "__kernel void test_ctors_executed_once_read(global uint *output)\n"
204 "{\n"
205 " ulong gid = get_global_id(0);\n"
206 " output[gid] = 0;\n"
207 "}\n"
208 ;
209 #else
210 const char * program_test_ctors_executed_once =
211 "#include <opencl_memory>\n"
212 "#include <opencl_work_item>\n"
213 "using namespace cl;\n"
214 // struct template
215 "template<class T>\n"
216 "struct ctor_test_class {\n"
217 // non-trivial ctor
218 " ctor_test_class(T y) { x = y;};\n"
219 " T x;\n"
220 "};\n"
221 // global scope program variables
222 "ctor_test_class<int> global_var0(int(0));\n"
223 "ctor_test_class<uint> global_var1(uint(0));\n"
224
225 "__kernel void test_ctors_executed_once_set()\n"
226 "{\n"
227 " ulong gid = get_global_id(0);\n"
228 " if(gid == 0) {\n"
229 " global_var0.x = int(0xbeefbeef);\n"
230 " global_var1.x = uint(0xbeefbeefU);\n"
231 " }\n"
232 "}\n\n"
233
234 "__kernel void test_ctors_executed_once_read(global_ptr<uint[]> output)\n"
235 "{\n"
236 " ulong gid = get_global_id(0);\n"
237 " int result = 0;\n"
238 " if(global_var0.x != int(0xbeefbeef)) result = 1;\n"
239 " if(global_var1.x != uint(0xbeefbeefU)) result = 1;\n"
240 " output[gid] = result;\n"
241 "}\n"
242 ;
243 #endif
244
AUTO_TEST_CASE(test_global_scope_ctors_executed_once)245 AUTO_TEST_CASE(test_global_scope_ctors_executed_once)
246 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
247 {
248 int error = CL_SUCCESS;
249
250 cl_mem output_buffer;
251 cl_program program;
252 cl_kernel kernel_set_global_vars;
253 cl_kernel kernel_read_global_vars;
254
255 size_t dim = 1;
256 size_t work_size[1];
257 // -----------------------------------------------------------------------------------
258 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
259 // -----------------------------------------------------------------------------------
260 // Only OpenCL C++ to SPIR-V compilation
261 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
262 error = create_opencl_kernel(
263 context, &program, &kernel_set_global_vars,
264 program_test_ctors_executed_once, "test_ctors_executed_once_set"
265 );
266 RETURN_ON_ERROR(error)
267 return error;
268 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
269 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
270 error = create_opencl_kernel(
271 context, &program, &kernel_set_global_vars,
272 program_test_ctors_executed_once, "test_ctors_executed_once_set", "", false
273 );
274 RETURN_ON_ERROR(error)
275 // Get the second kernel
276 kernel_read_global_vars = clCreateKernel(program, "test_ctors_executed_once_read", &error);
277 RETURN_ON_CL_ERROR(error, "clCreateKernel");
278 // Normal run
279 #else
280 error = create_opencl_kernel(
281 context, &program, &kernel_set_global_vars,
282 program_test_ctors_executed_once, "test_ctors_executed_once_set"
283 );
284 RETURN_ON_ERROR(error)
285 // Get the second kernel
286 kernel_read_global_vars = clCreateKernel(program, "test_ctors_executed_once_read", &error);
287 RETURN_ON_CL_ERROR(error, "clCreateKernel");
288 #endif
289
290 // Execute kernel_set_global_vars
291
292 work_size[0] = count;
293 error = clEnqueueNDRangeKernel(queue, kernel_set_global_vars, dim, NULL, work_size, NULL, 0, NULL, NULL);
294 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
295
296 // Execute kernel_read_global_vars 4 times, each time we check if
297 // global variables have correct values.
298
299 // host vector, size == count, output[0...count-1] == 1
300 std::vector<cl_uint> output(count, cl_uint(1));
301 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
302 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
303
304 for(size_t i = 0; i < 4; i++)
305 {
306 std::fill(output.begin(), output.end(), cl_uint(1));
307 error = clEnqueueWriteBuffer(
308 queue, output_buffer, CL_TRUE,
309 0, sizeof(cl_uint) * output.size(),
310 static_cast<void *>(output.data()),
311 0, NULL, NULL
312 );
313 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
314
315 error = clSetKernelArg(kernel_read_global_vars, 0, sizeof(output_buffer), &output_buffer);
316 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
317
318 work_size[0] = output.size();
319 error = clEnqueueNDRangeKernel(
320 queue, kernel_read_global_vars,
321 dim, NULL, work_size, NULL,
322 0, NULL, NULL
323 );
324 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
325
326 error = clEnqueueReadBuffer(
327 queue, output_buffer, CL_TRUE,
328 0, sizeof(cl_uint) * output.size(),
329 static_cast<void *>(output.data()),
330 0, NULL, NULL
331 );
332 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
333
334 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
335 if(sum != 0)
336 {
337 error = -1;
338 CHECK_ERROR_MSG(error, "Test test_ctors_executed_onces failed.");
339 }
340 }
341
342 clReleaseMemObject(output_buffer);
343 clReleaseKernel(kernel_set_global_vars);
344 clReleaseKernel(kernel_read_global_vars);
345 clReleaseProgram(program);
346 return error;
347 }
348
349 // TEST3
350 // Verify that when constructor is executed, the ND-range used is (1,1,1).
351
352 // -----------------------------------------------------------------------------------
353 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
354 // -----------------------------------------------------------------------------------
355 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
356 const char * program_test_ctors_ndrange =
357 "__kernel void test_ctors_ndrange(global int *output)\n"
358 "{\n"
359 " ulong gid = get_global_id(0);\n"
360 " output[gid] = 0;\n"
361 "}\n"
362 ;
363 #else
364 const char * program_test_ctors_ndrange =
365 "#include <opencl_memory>\n"
366 "#include <opencl_work_item>\n"
367 "using namespace cl;\n"
368 // struct
369 "struct ctor_test_class {\n"
370 // non-trivial ctor
371 " ctor_test_class() {\n"
372 " x = get_global_size(0);\n"
373 " y = get_global_size(1);\n"
374 " z = get_global_size(2);\n"
375 " };\n"
376 " ulong x;\n"
377 " ulong y;\n"
378 " ulong z;\n"
379 // return true if the ND-range used when ctor was exectured was
380 // (1, 1, 1); otherwise - false
381 " bool check() { return (x == 1) && (y == 1) && (z == 1);}"
382 "};\n"
383 // global scope program variables
384 "ctor_test_class global_var0;\n"
385 "ctor_test_class global_var1;\n"
386
387 "__kernel void test_ctors_ndrange(global_ptr<uint[]> output)\n"
388 "{\n"
389 " ulong gid = get_global_id(0);\n"
390 " int result = 0;\n"
391 " if(!global_var0.check()) result = 1;\n"
392 " if(!global_var1.check()) result = 1;\n"
393 " output[gid] = result;\n"
394 "}\n"
395 ;
396 #endif
397
AUTO_TEST_CASE(test_global_scope_ctors_ndrange)398 AUTO_TEST_CASE(test_global_scope_ctors_ndrange)
399 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
400 {
401 int error = CL_SUCCESS;
402
403 cl_mem output_buffer;
404 cl_program program;
405 cl_kernel kernel;
406
407 size_t dim = 1;
408 size_t work_size[1];
409 // -----------------------------------------------------------------------------------
410 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
411 // -----------------------------------------------------------------------------------
412 // Only OpenCL C++ to SPIR-V compilation
413 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
414 error = create_opencl_kernel(
415 context, &program, &kernel,
416 program_test_ctors_ndrange, "test_ctors_ndrange"
417 );
418 RETURN_ON_ERROR(error)
419 return error;
420 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
421 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
422 error = create_opencl_kernel(
423 context, &program, &kernel,
424 program_test_ctors_ndrange, "test_ctors_ndrange", "", false
425 );
426 RETURN_ON_ERROR(error)
427 // Normal run
428 #else
429 error = create_opencl_kernel(
430 context, &program, &kernel,
431 program_test_ctors_ndrange, "test_ctors_ndrange"
432 );
433 RETURN_ON_ERROR(error)
434 #endif
435
436 // host vector, size == count, output[0...count-1] == 1
437 std::vector<cl_uint> output(count, cl_uint(1));
438 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
439 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
440
441 error = clEnqueueWriteBuffer(
442 queue, output_buffer, CL_TRUE,
443 0, sizeof(cl_uint) * output.size(),
444 static_cast<void *>(output.data()),
445 0, NULL, NULL
446 );
447 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
448
449 error = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
450 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
451
452 work_size[0] = output.size();
453 error = clEnqueueNDRangeKernel(
454 queue, kernel,
455 dim, NULL, work_size, NULL,
456 0, NULL, NULL
457 );
458 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
459
460 error = clEnqueueReadBuffer(
461 queue, output_buffer, CL_TRUE,
462 0, sizeof(cl_uint) * output.size(),
463 static_cast<void *>(output.data()),
464 0, NULL, NULL
465 );
466 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
467
468 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
469 if(sum != 0)
470 {
471 error = -1;
472 CHECK_ERROR_MSG(error, "Test test_ctors_executed_ndrange failed.");
473 }
474
475 clReleaseMemObject(output_buffer);
476 clReleaseKernel(kernel);
477 clReleaseProgram(program);
478 return error;
479 }
480
481 #endif // TEST_CONFORMANCE_CLCPP_API_TEST_CTORS_HPP
482