• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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