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_DTORS_HPP
17 #define TEST_CONFORMANCE_CLCPP_API_TEST_DTORS_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 destructor is executed.
28
29 // How: destructor of struct dtor_test_class has a side effect: zeroing buffer. If values
30 // in buffer are not zeros after releasing program, destructor was not executed.
31
32 // -----------------------------------------------------------------------------------
33 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
34 // -----------------------------------------------------------------------------------
35 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
36 const char * program_test_dtor_is_executed =
37 "__kernel void test_dtor_is_executed(global uint *output)\n"
38 "{\n"
39 " ulong gid = get_global_id(0);\n"
40 " output[gid] = 0;\n"
41 "}\n"
42 ;
43 #else
44 const char * program_test_dtor_is_executed =
45 "#include <opencl_memory>\n"
46 "#include <opencl_work_item>\n"
47 "using namespace cl;\n"
48 // struct
49 "struct dtor_test_class {\n"
50 // non-trivial dtor
51 // set all values in buffer to 0
52 " ~dtor_test_class() {\n"
53 " for(ulong i = 0; i < size; i++)\n"
54 " buffer[i] = 0;\n"
55 " };\n"
56 " global_ptr<uint[]> buffer;\n"
57 " ulong size;\n"
58 "};\n"
59 // global scope program variable
60 "dtor_test_class global_var;\n"
61
62 // values in output __MUST BE__ greater than 0 for the test to work
63 // correctly
64 "__kernel void test_dtor_is_executed(global_ptr<uint[]> output)\n"
65 "{\n"
66 " ulong gid = get_global_id(0);\n"
67 // set buffer and size in global var
68 " if(gid == 0){\n"
69 " global_var.buffer = output;\n"
70 " global_var.size = get_global_size(0);\n"
71 " }\n"
72 "}\n"
73 ;
74 #endif
75
AUTO_TEST_CASE(test_global_scope_dtor_is_executed)76 AUTO_TEST_CASE(test_global_scope_dtor_is_executed)
77 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
78 {
79 int error = CL_SUCCESS;
80
81 cl_mem output_buffer;
82 cl_program program;
83 cl_kernel kernel;
84
85 size_t dim = 1;
86 size_t work_size[1];
87 // -----------------------------------------------------------------------------------
88 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
89 // -----------------------------------------------------------------------------------
90 // Only OpenCL C++ to SPIR-V compilation
91 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
92 error = create_opencl_kernel(
93 context, &program, &kernel,
94 program_test_dtor_is_executed, "test_dtor_is_executed"
95 );
96 RETURN_ON_ERROR(error)
97 return error;
98 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
99 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
100 error = create_opencl_kernel(
101 context, &program, &kernel,
102 program_test_dtor_is_executed, "test_dtor_is_executed", "", false
103 );
104 RETURN_ON_ERROR(error)
105 // Normal run
106 #else
107 error = create_opencl_kernel(
108 context, &program, &kernel,
109 program_test_dtor_is_executed, "test_dtor_is_executed"
110 );
111 RETURN_ON_ERROR(error)
112 #endif
113
114 // host vector, size == count, output[0...count-1] == 0xbeefbeef (3203383023)
115 // values in output __MUST BE__ greater than 0 for the test to work correctly
116 std::vector<cl_uint> output(count, cl_uint(0xbeefbeef));
117 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
118 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
119
120 error = clEnqueueWriteBuffer(
121 queue, output_buffer, CL_TRUE,
122 0, sizeof(cl_uint) * output.size(),
123 static_cast<void *>(output.data()),
124 0, NULL, NULL
125 );
126 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
127
128 error = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
129 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
130
131 work_size[0] = output.size();
132 error = clEnqueueNDRangeKernel(
133 queue, kernel,
134 dim, NULL, work_size, NULL,
135 0, NULL, NULL
136 );
137 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
138
139 // Release kernel and program
140 // Dtor should be called now
141 error = clReleaseKernel(kernel);
142 RETURN_ON_CL_ERROR(error, "clReleaseKernel")
143 error = clReleaseProgram(program);
144 RETURN_ON_CL_ERROR(error, "clReleaseProgram")
145
146 // Finish
147 error = clFinish(queue);
148 RETURN_ON_CL_ERROR(error, "clFinish")
149
150 // Read output buffer
151 error = clEnqueueReadBuffer(
152 queue, output_buffer, CL_TRUE,
153 0, sizeof(cl_uint) * output.size(),
154 static_cast<void *>(output.data()),
155 0, NULL, NULL
156 );
157 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
158
159 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
160 if(sum != 0)
161 {
162 error = -1;
163 CHECK_ERROR_MSG(error, "Test test_dtor_is_executed failed.");
164 }
165
166 clReleaseMemObject(output_buffer);
167 return error;
168 }
169
170 // TEST 2
171 // Verify that multiple destructors, if present, are executed. Order between multiple
172 // destructors is undefined.
173 // Verify that each destructor is executed only once.
174
175 // How:
176 // 0) dtor_test_class struct has a global pointer to a buffer, it's set by
177 // test_dtors_executed_once kernel.
178 // 1) Destructors have a side effect: each dtor writes to its part of the buffer. If all
179 // dtors are executed, all values in that buffer should be changed.
180 // 2) The first time destructors are executed, they set their parts of the buffer to zero.
181 // Next time to 1, next time to 2 etc. Since dtors should be executed only once, all
182 // values in that buffer should be equal to zero.
183
184 // -----------------------------------------------------------------------------------
185 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
186 // -----------------------------------------------------------------------------------
187 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
188 const char * program_test_dtors_executed_once =
189 "__kernel void test_dtors_executed_once(global uint *output)\n"
190 "{\n"
191 " ulong gid = get_global_id(0);\n"
192 " output[gid] = 0;\n"
193 "}\n"
194 ;
195 #else
196 const char * program_test_dtors_executed_once =
197 "#include <opencl_memory>\n"
198 "#include <opencl_work_item>\n"
199 "using namespace cl;\n"
200 // struct
201 "struct dtor_test_class {\n"
202 // non-trivial dtor
203 // Set all values in range [start; end - 1] in buffer to counter.
204 // If dtor is executed only once (correct), all values in range
205 // [start; end - 1] in buffer should be equal to zero after releasing
206 // the program
207 " ~dtor_test_class() {\n"
208 " for(ulong i = start; i < end; i++){\n"
209 " buffer[i] = counter;\n"
210 " };\n"
211 " counter++;\n"
212 " };\n"
213 " global_ptr<uint[]> buffer;\n"
214 " ulong start;\n"
215 " ulong end;\n"
216 " ulong counter;\n"
217 "};\n"
218 // global scope program variables
219 "dtor_test_class global_var0;\n"
220 "dtor_test_class global_var1;\n"
221 "dtor_test_class global_var2;\n"
222 "dtor_test_class global_var3;\n"
223
224 // values in output __MUST BE__ greater than 0 for the test to work correctly
225 "__kernel void test_dtors_executed_once(global_ptr<uint[]> output)\n"
226 "{\n"
227 " ulong gid = get_global_id(0);\n"
228 // set buffer and size in global var
229 " if(gid == 0){\n"
230 " ulong end = get_global_size(0) / 4;"
231 // global_var0
232 " global_var0.buffer = output;\n"
233 " global_var0.start = 0;\n"
234 " global_var0.end = end;\n"
235 " global_var0.counter = 0;\n"
236 // global_var1
237 " global_var1.buffer = output;\n"
238 " global_var1.start = end;\n"
239 " end += get_global_size(0) / 4;\n"
240 " global_var1.end = end;\n"
241 " global_var1.counter = 0;\n"
242 // global_var2
243 " global_var2.buffer = output;\n"
244 " global_var2.start = end;\n"
245 " end += get_global_size(0) / 4;\n"
246 " global_var2.end = end;\n"
247 " global_var2.counter = 0;\n"
248 // global_var3
249 " global_var3.buffer = output;\n"
250 " global_var3.start = end;\n"
251 " global_var3.end = get_global_size(0);\n"
252 " global_var3.counter = 0;\n"
253 " }\n"
254 "}\n"
255 ;
256 #endif
257
AUTO_TEST_CASE(test_global_scope_dtors_executed_once)258 AUTO_TEST_CASE(test_global_scope_dtors_executed_once)
259 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
260 {
261 int error = CL_SUCCESS;
262
263 cl_mem output_buffer;
264 cl_program program;
265 cl_kernel kernel;
266
267 size_t dim = 1;
268 size_t work_size[1];
269 // -----------------------------------------------------------------------------------
270 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
271 // -----------------------------------------------------------------------------------
272 // Only OpenCL C++ to SPIR-V compilation
273 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
274 error = create_opencl_kernel(
275 context, &program, &kernel,
276 program_test_dtors_executed_once, "test_dtors_executed_once"
277 );
278 RETURN_ON_ERROR(error)
279 return error;
280 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
281 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
282 error = create_opencl_kernel(
283 context, &program, &kernel,
284 program_test_dtors_executed_once, "test_dtors_executed_once", "", false
285 );
286 RETURN_ON_ERROR(error)
287 // Normal run
288 #else
289 error = create_opencl_kernel(
290 context, &program, &kernel,
291 program_test_dtors_executed_once, "test_dtors_executed_once"
292 );
293 RETURN_ON_ERROR(error)
294 #endif
295
296 // host vector, size == count, output[0...count-1] == 0xbeefbeef (3203383023)
297 // values in output __MUST BE__ greater than 0 for the test to work correctly
298 cl_uint init_value = cl_uint(0xbeefbeef);
299 std::vector<cl_uint> output(count, init_value);
300 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
301 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
302
303 error = clEnqueueWriteBuffer(
304 queue, output_buffer, CL_TRUE,
305 0, sizeof(cl_uint) * output.size(),
306 static_cast<void *>(output.data()),
307 0, NULL, NULL
308 );
309 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
310
311 error = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
312 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
313
314 work_size[0] = output.size();
315 error = clEnqueueNDRangeKernel(
316 queue, kernel,
317 dim, NULL, work_size, NULL,
318 0, NULL, NULL
319 );
320 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
321
322
323 // Increments the program reference count. Twice
324 error = clRetainProgram(program);
325 RETURN_ON_CL_ERROR(error, "clRetainProgram")
326 error = clRetainProgram(program);
327 RETURN_ON_CL_ERROR(error, "clRetainProgram")
328
329 // Should just decrement the program reference count.
330 error = clReleaseProgram(program);
331 RETURN_ON_CL_ERROR(error, "clReleaseProgram")
332 error = clFinish(queue);
333 RETURN_ON_CL_ERROR(error, "clFinish")
334
335 // Should just decrement the program reference count.
336 error = clReleaseProgram(program);
337 RETURN_ON_CL_ERROR(error, "clReleaseProgram")
338 error = clFinish(queue);
339 RETURN_ON_CL_ERROR(error, "clFinish")
340
341 #ifndef USE_OPENCLC_KERNELS
342 // At this point global scope variables should not be destroyed,
343 // values in output buffer should not be modified.
344
345 // Read output buffer
346 error = clEnqueueReadBuffer(
347 queue, output_buffer, CL_TRUE,
348 0, sizeof(cl_uint) * output.size(),
349 static_cast<void *>(output.data()),
350 0, NULL, NULL
351 );
352 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
353 for(auto& i : output)
354 {
355 if(i != init_value)
356 {
357 log_error("ERROR: Test test_global_scope_dtors_executed_once failed.");
358 log_error("\tDestructors were executed prematurely.\n");
359 RETURN_ON_ERROR(-1)
360 }
361 }
362 #endif
363
364 // Release kernel and program, destructors should be called now
365 error = clReleaseKernel(kernel);
366 RETURN_ON_CL_ERROR(error, "clReleaseKernel")
367 error = clReleaseProgram(program);
368 RETURN_ON_CL_ERROR(error, "clReleaseProgram")
369
370 // Finish
371 error = clFinish(queue);
372 RETURN_ON_CL_ERROR(error, "clFinish")
373
374 // Read output buffer
375 error = clEnqueueReadBuffer(
376 queue, output_buffer, CL_TRUE,
377 0, sizeof(cl_uint) * output.size(),
378 static_cast<void *>(output.data()),
379 0, NULL, NULL
380 );
381 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
382
383 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
384 if(sum != 0)
385 {
386 log_error("ERROR: Test test_global_scope_dtors_executed_once failed.");
387 // Maybe some dtors were not run?
388 for(auto& i : output)
389 {
390 if(i == init_value)
391 {
392 log_error("\tSome dtors were not executed.");
393 break;
394 }
395 }
396 log_error("\n");
397 RETURN_ON_ERROR(-1)
398 }
399
400 // Clean
401 clReleaseMemObject(output_buffer);
402 return error;
403 }
404
405 // TEST3
406 // Verify that ND-range during destructor execution is set to (1,1,1)
407
408 // -----------------------------------------------------------------------------------
409 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
410 // -----------------------------------------------------------------------------------
411 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
412 const char * program_test_dtor_ndrange =
413 "__kernel void test_dtor_ndrange(global uint *output)\n"
414 "{\n"
415 " ulong gid = get_global_id(0);\n"
416 " output[gid] = 0;\n"
417 "}\n"
418 ;
419 #else
420 const char * program_test_dtor_ndrange =
421 "#include <opencl_memory>\n"
422 "#include <opencl_work_item>\n"
423 "using namespace cl;\n"
424 // struct
425 "struct dtor_test_class {\n"
426 // non-trivial dtor
427 // set all values in buffer to 0 only if ND-range is (1, 1, 1)
428 " ~dtor_test_class() {\n"
429 " if(check()){\n"
430 " for(ulong i = 0; i < size; i++)\n"
431 " buffer[i] = 0;\n"
432 " }\n"
433 " };\n"
434 // return true if the ND-range is (1, 1, 1); otherwise - false
435 " bool check() {\n"
436 " return (get_global_size(0) == 1)"
437 " && (get_global_size(1) == 1)"
438 " && (get_global_size(2) == 1);\n"
439 " }"
440 " ulong size;\n"
441 " global_ptr<uint[]> buffer;\n"
442 "};\n"
443 // global scope program variable
444 "dtor_test_class global_var;\n"
445
446 // values in output __MUST BE__ greater than 0 for the test to work correctly
447 "__kernel void test_dtor_ndrange(global_ptr<uint[]> output)\n"
448 "{\n"
449 " ulong gid = get_global_id(0);\n"
450 // set buffer and size in global var
451 " if(gid == 0){\n"
452 " global_var.buffer = output;\n"
453 " global_var.size = get_global_size(0);\n"
454 " }\n"
455 "}\n"
456 ;
457 #endif
458
AUTO_TEST_CASE(test_global_scope_dtor_ndrange)459 AUTO_TEST_CASE(test_global_scope_dtor_ndrange)
460 (cl_device_id device, cl_context context, cl_command_queue queue, int count)
461 {
462 int error = CL_SUCCESS;
463
464 cl_mem output_buffer;
465 cl_program program;
466 cl_kernel kernel;
467
468 size_t dim = 1;
469 size_t work_size[1];
470 // -----------------------------------------------------------------------------------
471 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
472 // -----------------------------------------------------------------------------------
473 // Only OpenCL C++ to SPIR-V compilation
474 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
475 error = create_opencl_kernel(
476 context, &program, &kernel,
477 program_test_dtor_ndrange, "test_dtor_ndrange"
478 );
479 RETURN_ON_ERROR(error)
480 return error;
481 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
482 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
483 error = create_opencl_kernel(
484 context, &program, &kernel,
485 program_test_dtor_ndrange, "test_dtor_ndrange", "", false
486 );
487 RETURN_ON_ERROR(error)
488 // Normal run
489 #else
490 error = create_opencl_kernel(
491 context, &program, &kernel,
492 program_test_dtor_ndrange, "test_dtor_ndrange"
493 );
494 RETURN_ON_ERROR(error)
495 #endif
496
497 // host vector, size == count, output[0...count-1] == 0xbeefbeef (3203383023)
498 // values in output __MUST BE__ greater than 0 for the test to work correctly
499 std::vector<cl_uint> output(count, cl_uint(0xbeefbeef));
500 output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
501 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
502
503 error = clEnqueueWriteBuffer(
504 queue, output_buffer, CL_TRUE,
505 0, sizeof(cl_uint) * output.size(),
506 static_cast<void *>(output.data()),
507 0, NULL, NULL
508 );
509 RETURN_ON_CL_ERROR(error, "clEnqueueWriteBuffer")
510
511 error = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
512 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
513
514 work_size[0] = output.size();
515 error = clEnqueueNDRangeKernel(
516 queue, kernel,
517 dim, NULL, work_size, NULL,
518 0, NULL, NULL
519 );
520 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
521
522 // Release kernel and program
523 // Dtor should be called now
524 error = clReleaseKernel(kernel);
525 RETURN_ON_CL_ERROR(error, "clReleaseKernel")
526 error = clReleaseProgram(program);
527 RETURN_ON_CL_ERROR(error, "clReleaseProgram")
528
529 // Finish
530 error = clFinish(queue);
531 RETURN_ON_CL_ERROR(error, "clFinish")
532
533 // Read output buffer
534 error = clEnqueueReadBuffer(
535 queue, output_buffer, CL_TRUE,
536 0, sizeof(cl_uint) * output.size(),
537 static_cast<void *>(output.data()),
538 0, NULL, NULL
539 );
540 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
541
542 size_t sum = std::accumulate(output.begin(), output.end(), size_t(0));
543 if(sum != 0)
544 {
545 error = -1;
546 CHECK_ERROR_MSG(error, "Test test_dtor_ndrange failed.");
547 }
548
549 clReleaseMemObject(output_buffer);
550 return error;
551 }
552
553 #endif // TEST_CONFORMANCE_CLCPP_API_TEST_DTORS_HPP
554