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 #include "testBase.h"
17 #include "harness/imageHelpers.h"
18 #include "harness/propertyHelpers.h"
19 #include <stdlib.h>
20 #include <ctype.h>
21 #include <algorithm>
22 
test_get_platform_info(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)23 int test_get_platform_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
24 {
25     cl_platform_id platform;
26     cl_int error;
27     char buffer[ 16384 ];
28     size_t length;
29 
30     // Get the platform to use
31     error = clGetPlatformIDs(1, &platform, NULL);
32     test_error( error, "Unable to get platform" );
33 
34     // Platform profile should either be FULL_PROFILE or EMBEDDED_PROFILE
35     error = clGetPlatformInfo(platform,  CL_PLATFORM_PROFILE, sizeof( buffer ), buffer, &length );
36     test_error( error, "Unable to get platform profile string" );
37 
38     log_info("Returned CL_PLATFORM_PROFILE %s.\n", buffer);
39 
40     if( strcmp( buffer, "FULL_PROFILE" ) != 0 && strcmp( buffer, "EMBEDDED_PROFILE" ) != 0 )
41     {
42         log_error( "ERROR: Returned platform profile string is not a valid string by OpenCL 1.2! (Returned: %s)\n", buffer );
43         return -1;
44     }
45     if( strlen( buffer )+1 != length )
46     {
47         log_error( "ERROR: Returned length of profile string is incorrect (actual length: %d, returned length: %d)\n",
48                   (int)strlen( buffer )+1, (int)length );
49         return -1;
50     }
51 
52     // Check just length return
53     error = clGetPlatformInfo(platform,  CL_PLATFORM_PROFILE, 0, NULL, &length );
54     test_error( error, "Unable to get platform profile length" );
55     if( strlen( (char *)buffer )+1 != length )
56     {
57         log_error( "ERROR: Returned length of profile string is incorrect (actual length: %d, returned length: %d)\n",
58                   (int)strlen( (char *)buffer )+1, (int)length );
59         return -1;
60     }
61 
62 
63     // Platform version should fit the regex "OpenCL *[0-9]+\.[0-9]+"
64     error = clGetPlatformInfo(platform,  CL_PLATFORM_VERSION, sizeof( buffer ), buffer, &length );
65     test_error( error, "Unable to get platform version string" );
66 
67     log_info("Returned CL_PLATFORM_VERSION %s.\n", buffer);
68 
69     if( memcmp( buffer, "OpenCL ", strlen( "OpenCL " ) ) != 0 )
70     {
71         log_error( "ERROR: Initial part of platform version string does not match required format! (returned: %s)\n", (char *)buffer );
72         return -1;
73     }
74     char *p1 = (char *)buffer + strlen( "OpenCL " );
75     while( *p1 == ' ' )
76         p1++;
77     char *p2 = p1;
78     while( isdigit( *p2 ) )
79         p2++;
80     if( *p2 != '.' )
81     {
82         log_error( "ERROR: Numeric part of platform version string does not match required format! (returned: %s)\n", (char *)buffer );
83         return -1;
84     }
85     char *p3 = p2 + 1;
86     while( isdigit( *p3 ) )
87         p3++;
88     if( *p3 != ' ' )
89     {
90         log_error( "ERROR: space expected after minor version number! (returned: %s)\n", (char *)buffer );
91         return -1;
92     }
93     *p2 = ' '; // Put in a space for atoi below.
94     p2++;
95 
96     // make sure it is null terminated
97     for( ; p3 != buffer + length; p3++ )
98         if( *p3 == '\0' )
99             break;
100     if( p3 == buffer + length )
101     {
102         log_error( "ERROR: platform version string is not NUL terminated!\n" );
103         return -1;
104     }
105 
106     int major = atoi( p1 );
107     int minor = atoi( p2 );
108     int minor_revision = 2;
109     if( major * 10 + minor < 10 + minor_revision )
110     {
111         log_error( "ERROR: OpenCL profile version returned is less than 1.%d!\n", minor_revision );
112         return -1;
113     }
114 
115     // Sanity checks on the returned values
116     if( length != strlen( (char *)buffer ) + 1)
117     {
118         log_error( "ERROR: Returned length of version string does not match actual length (actual: %d, returned: %d)\n", (int)strlen( (char *)buffer )+1, (int)length );
119         return -1;
120     }
121 
122     // Check just length
123     error = clGetPlatformInfo(platform,  CL_PLATFORM_VERSION, 0, NULL, &length );
124     test_error( error, "Unable to get platform version length" );
125     if( length != strlen( (char *)buffer )+1 )
126     {
127         log_error( "ERROR: Returned length of version string does not match actual length (actual: %d, returned: %d)\n", (int)strlen( buffer )+1, (int)length );
128         return -1;
129     }
130 
131     return 0;
132 }
133 
134 template <typename T>
sampler_param_test(cl_sampler sampler,cl_sampler_info param_name,T expected,const char * name)135 int sampler_param_test(cl_sampler sampler, cl_sampler_info param_name,
136                        T expected, const char *name)
137 {
138     size_t size;
139     T val;
140     int error = clGetSamplerInfo(sampler, param_name, sizeof(val), &val, &size);
141     test_error(error, "Unable to get sampler info");
142     if (val != expected)
143     {
144         test_fail("ERROR: Sampler %s did not validate!\n", name);
145     }
146     if (size != sizeof(val))
147     {
148         test_fail("ERROR: Returned size of sampler %s does not validate! "
149                   "(expected %d, got %d)\n",
150                   name, (int)sizeof(val), (int)size);
151     }
152     return 0;
153 }
154 
155 static cl_int normalized_coord_values[] = { CL_TRUE, CL_FALSE };
156 static cl_addressing_mode addressing_mode_values[] = {
157     CL_ADDRESS_NONE, CL_ADDRESS_CLAMP_TO_EDGE, CL_ADDRESS_CLAMP,
158     CL_ADDRESS_REPEAT, CL_ADDRESS_MIRRORED_REPEAT
159 };
160 static cl_filter_mode filter_mode_values[] = { CL_FILTER_NEAREST,
161                                                CL_FILTER_LINEAR };
162 
test_sampler_params(cl_device_id deviceID,cl_context context,bool is_compatibility,int norm_coord_num,int addr_mod_num,int filt_mod_num)163 int test_sampler_params(cl_device_id deviceID, cl_context context,
164                         bool is_compatibility, int norm_coord_num,
165                         int addr_mod_num, int filt_mod_num)
166 {
167     cl_uint refCount;
168     size_t size;
169     int error;
170 
171     clSamplerWrapper sampler;
172     cl_sampler_properties properties[] = {
173         CL_SAMPLER_NORMALIZED_COORDS,
174         normalized_coord_values[norm_coord_num],
175         CL_SAMPLER_ADDRESSING_MODE,
176         addressing_mode_values[addr_mod_num],
177         CL_SAMPLER_FILTER_MODE,
178         filter_mode_values[filt_mod_num],
179         0
180     };
181 
182     if (is_compatibility)
183     {
184         sampler =
185             clCreateSampler(context, normalized_coord_values[norm_coord_num],
186                             addressing_mode_values[addr_mod_num],
187                             filter_mode_values[filt_mod_num], &error);
188         test_error(error, "Unable to create sampler to test with");
189     }
190     else
191     {
192         sampler = clCreateSamplerWithProperties(context, properties, &error);
193         test_error(error, "Unable to create sampler to test with");
194     }
195 
196     error = clGetSamplerInfo(sampler, CL_SAMPLER_REFERENCE_COUNT,
197                              sizeof(refCount), &refCount, &size);
198     test_error(error, "Unable to get sampler ref count");
199     test_assert_error(size == sizeof(refCount),
200                       "Returned size of sampler refcount does not validate!\n");
201 
202     error = sampler_param_test(sampler, CL_SAMPLER_CONTEXT, context, "context");
203     test_error(error, "param checking failed");
204 
205     error = sampler_param_test(sampler, CL_SAMPLER_ADDRESSING_MODE,
206                                addressing_mode_values[addr_mod_num],
207                                "addressing mode");
208     test_error(error, "param checking failed");
209 
210     error = sampler_param_test(sampler, CL_SAMPLER_FILTER_MODE,
211                                filter_mode_values[filt_mod_num], "filter mode");
212     test_error(error, "param checking failed");
213 
214     error = sampler_param_test(sampler, CL_SAMPLER_NORMALIZED_COORDS,
215                                normalized_coord_values[norm_coord_num],
216                                "normalized coords");
217     test_error(error, "param checking failed");
218 
219     Version version = get_device_cl_version(deviceID);
220     if (version >= Version(3, 0))
221     {
222         std::vector<cl_sampler_properties> test_properties(
223             properties, properties + ARRAY_SIZE(properties));
224 
225         std::vector<cl_sampler_properties> check_properties;
226         size_t set_size;
227 
228         error = clGetSamplerInfo(sampler, CL_SAMPLER_PROPERTIES, 0, NULL,
229                                  &set_size);
230         test_error(
231             error,
232             "clGetSamplerInfo failed asking for CL_SAMPLER_PROPERTIES size.");
233 
234         if (is_compatibility)
235         {
236             if (set_size != 0)
237             {
238                 log_error(
239                     "ERROR: CL_SAMPLER_PROPERTIES size is %d, expected 0\n",
240                     set_size);
241                 return TEST_FAIL;
242             }
243         }
244         else
245         {
246             if (set_size
247                 != test_properties.size() * sizeof(cl_sampler_properties))
248             {
249                 log_error(
250                     "ERROR: CL_SAMPLER_PROPERTIES size is %d, expected %d.\n",
251                     set_size,
252                     test_properties.size() * sizeof(cl_sampler_properties));
253                 return TEST_FAIL;
254             }
255 
256             cl_uint number_of_props = set_size / sizeof(cl_sampler_properties);
257             check_properties.resize(number_of_props);
258             error = clGetSamplerInfo(sampler, CL_SAMPLER_PROPERTIES, set_size,
259                                      check_properties.data(), 0);
260             test_error(
261                 error,
262                 "clGetSamplerInfo failed asking for CL_SAMPLER_PROPERTIES.");
263 
264             error = compareProperties(check_properties, test_properties);
265             test_error(error, "checkProperties mismatch.");
266         }
267     }
268     return 0;
269 }
270 
get_sampler_info_params(cl_device_id deviceID,cl_context context,bool is_compatibility)271 int get_sampler_info_params(cl_device_id deviceID, cl_context context,
272                             bool is_compatibility)
273 {
274     for (int norm_coord_num = 0;
275          norm_coord_num < ARRAY_SIZE(normalized_coord_values); norm_coord_num++)
276     {
277         for (int addr_mod_num = 0;
278              addr_mod_num < ARRAY_SIZE(addressing_mode_values); addr_mod_num++)
279         {
280             if ((normalized_coord_values[norm_coord_num] == CL_FALSE)
281                 && ((addressing_mode_values[addr_mod_num] == CL_ADDRESS_REPEAT)
282                     || (addressing_mode_values[addr_mod_num]
283                         == CL_ADDRESS_MIRRORED_REPEAT)))
284             {
285                 continue;
286             }
287             for (int filt_mod_num = 0;
288                  filt_mod_num < ARRAY_SIZE(filter_mode_values); filt_mod_num++)
289             {
290                 int err = test_sampler_params(deviceID, context,
291                                               is_compatibility, norm_coord_num,
292                                               addr_mod_num, filt_mod_num);
293                 test_error(err, "testing clGetSamplerInfo params failed");
294             }
295         }
296     }
297     return 0;
298 }
test_get_sampler_info(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)299 int test_get_sampler_info(cl_device_id deviceID, cl_context context,
300                           cl_command_queue queue, int num_elements)
301 {
302     int error;
303     PASSIVE_REQUIRE_IMAGE_SUPPORT(deviceID)
304 
305     error = get_sampler_info_params(deviceID, context, false);
306     test_error(error, "Test Failed");
307 
308     return 0;
309 }
310 
test_get_sampler_info_compatibility(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)311 int test_get_sampler_info_compatibility(cl_device_id deviceID,
312                                         cl_context context,
313                                         cl_command_queue queue,
314                                         int num_elements)
315 {
316     int error;
317     PASSIVE_REQUIRE_IMAGE_SUPPORT(deviceID)
318 
319     error = get_sampler_info_params(deviceID, context, true);
320     test_error(error, "Test Failed");
321 
322     return 0;
323 }
324 
325 template <typename T>
command_queue_param_test(cl_command_queue queue,cl_command_queue_info param_name,T expected,const char * name)326 int command_queue_param_test(cl_command_queue queue,
327                              cl_command_queue_info param_name, T expected,
328                              const char *name)
329 {
330     size_t size;
331     T val;
332     int error =
333         clGetCommandQueueInfo(queue, param_name, sizeof(val), &val, &size);
334     test_error(error, "Unable to get command queue info");
335     if (val != expected)
336     {
337         test_fail("ERROR: Command queue %s did not validate!\n", name);
338     }
339     if (size != sizeof(val))
340     {
341         test_fail("ERROR: Returned size of command queue %s does not validate! "
342                   "(expected %d, got %d)\n",
343                   name, (int)sizeof(val), (int)size);
344     }
345     return 0;
346 }
347 
348 #define MIN_NUM_COMMAND_QUEUE_PROPERTIES 2
349 #define OOO_NUM_COMMAND_QUEUE_PROPERTIES 4
350 static cl_command_queue_properties property_options[] = {
351     0,
352 
353     CL_QUEUE_PROFILING_ENABLE,
354 
355     CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
356 
357     CL_QUEUE_PROFILING_ENABLE | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
358 
359     CL_QUEUE_ON_DEVICE | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
360 
361     CL_QUEUE_PROFILING_ENABLE | CL_QUEUE_ON_DEVICE
362         | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
363 
364     CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT
365         | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
366 
367     CL_QUEUE_PROFILING_ENABLE | CL_QUEUE_ON_DEVICE | CL_QUEUE_ON_DEVICE_DEFAULT
368         | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
369 };
370 
check_get_command_queue_info_params(cl_device_id deviceID,cl_context context,bool is_compatibility)371 int check_get_command_queue_info_params(cl_device_id deviceID,
372                                         cl_context context,
373                                         bool is_compatibility)
374 {
375     int error;
376     size_t size;
377 
378     cl_queue_properties host_queue_props, device_queue_props;
379     cl_queue_properties queue_props[] = { CL_QUEUE_PROPERTIES, 0, 0 };
380 
381     clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_HOST_PROPERTIES,
382                     sizeof(host_queue_props), &host_queue_props, NULL);
383     log_info("CL_DEVICE_QUEUE_ON_HOST_PROPERTIES is %d\n",
384              (int)host_queue_props);
385     clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES,
386                     sizeof(device_queue_props), &device_queue_props, NULL);
387     log_info("CL_DEVICE_QUEUE_ON_HOST_PROPERTIES is %d\n",
388              (int)device_queue_props);
389 
390     auto version = get_device_cl_version(deviceID);
391 
392     // Are on device queues supported
393     bool on_device_supported =
394         (version >= Version(2, 0) && version < Version(3, 0))
395         || (version >= Version(3, 0) && device_queue_props != 0);
396 
397     int num_test_options = MIN_NUM_COMMAND_QUEUE_PROPERTIES;
398     if (host_queue_props & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
399     {
400         // Test out-of-order queues properties if supported
401         num_test_options = OOO_NUM_COMMAND_QUEUE_PROPERTIES;
402     }
403     if (on_device_supported && !is_compatibility)
404     {
405         // Test queue on device if supported (in this case out-of-order must
406         // also be supported)
407         num_test_options = ARRAY_SIZE(property_options);
408     }
409 
410     for (int i = 0; i < num_test_options; i++)
411     {
412         queue_props[1] = property_options[i];
413         clCommandQueueWrapper queue;
414 
415         if (is_compatibility)
416         {
417             queue =
418                 clCreateCommandQueue(context, deviceID, queue_props[1], &error);
419             test_error(error, "Unable to create command queue to test with");
420         }
421         else
422         {
423             queue = clCreateCommandQueueWithProperties(context, deviceID,
424                                                        &queue_props[0], &error);
425             test_error(error, "Unable to create command queue to test with");
426         }
427 
428         cl_uint refCount;
429         error = clGetCommandQueueInfo(queue, CL_QUEUE_REFERENCE_COUNT,
430                                       sizeof(refCount), &refCount, &size);
431         test_error(error, "Unable to get command queue reference count");
432         test_assert_error(size == sizeof(refCount),
433                           "Returned size of command queue reference count does "
434                           "not validate!\n");
435 
436         error = command_queue_param_test(queue, CL_QUEUE_CONTEXT, context,
437                                          "context");
438         test_error(error, "param checking failed");
439 
440         error = command_queue_param_test(queue, CL_QUEUE_DEVICE, deviceID,
441                                          "deviceID");
442         test_error(error, "param checking failed");
443 
444         error = command_queue_param_test(queue, CL_QUEUE_PROPERTIES,
445                                          queue_props[1], "properties");
446         test_error(error, "param checking failed");
447     }
448     return 0;
449 }
test_get_command_queue_info(cl_device_id deviceID,cl_context context,cl_command_queue ignoreQueue,int num_elements)450 int test_get_command_queue_info(cl_device_id deviceID, cl_context context,
451                                 cl_command_queue ignoreQueue, int num_elements)
452 {
453     int error = check_get_command_queue_info_params(deviceID, context, false);
454     test_error(error, "Test Failed");
455     return 0;
456 }
457 
test_get_command_queue_info_compatibility(cl_device_id deviceID,cl_context context,cl_command_queue ignoreQueue,int num_elements)458 int test_get_command_queue_info_compatibility(cl_device_id deviceID,
459                                               cl_context context,
460                                               cl_command_queue ignoreQueue,
461                                               int num_elements)
462 {
463     int error = check_get_command_queue_info_params(deviceID, context, true);
464     test_error(error, "Test Failed");
465     return 0;
466 }
467 
test_get_context_info(cl_device_id deviceID,cl_context context,cl_command_queue ignoreQueue,int num_elements)468 int test_get_context_info(cl_device_id deviceID, cl_context context, cl_command_queue ignoreQueue, int num_elements)
469 {
470     int error;
471     size_t size;
472     cl_context_properties props;
473 
474     error = clGetContextInfo( context, CL_CONTEXT_PROPERTIES, sizeof( props ), &props, &size );
475     test_error( error, "Unable to get context props" );
476 
477     if (size == 0) {
478         // Valid size
479         return 0;
480     } else if (size == sizeof(cl_context_properties)) {
481         // Data must be NULL
482         if (props != 0) {
483             log_error("ERROR: Returned properties is no NULL.\n");
484             return -1;
485         }
486         // Valid data and size
487         return 0;
488     }
489     // Size was not 0 or 1
490     log_error( "ERROR: Returned size of context props is not valid! (expected 0 or %d, got %d)\n",
491               (int)sizeof(cl_context_properties), (int)size );
492     return -1;
493 }
494 
495 #define TEST_MEM_OBJECT_PARAM( mem, paramName, val, expected, name, type, cast )    \
496 error = clGetMemObjectInfo( mem, paramName, sizeof( val ), &val, &size );        \
497 test_error( error, "Unable to get mem object " name );                            \
498 if( val != expected )                                                                \
499 {                                                                                    \
500 log_error( "ERROR: Mem object " name " did not validate! (expected " type ", got " type ")\n", (cast)(expected), (cast)val );    \
501 return -1;                                                                        \
502 }            \
503 if( size != sizeof( val ) )                \
504 {                                        \
505 log_error( "ERROR: Returned size of mem object " name " does not validate! (expected %d, got %d)\n", (int)sizeof( val ), (int)size );    \
506 return -1;    \
507 }
508 
mem_obj_destructor_callback(cl_mem,void * data)509 void CL_CALLBACK mem_obj_destructor_callback( cl_mem, void *data )
510 {
511     free( data );
512 }
513 
514 // All possible combinations of valid cl_mem_flags.
515 static cl_mem_flags all_flags[16] = {
516   0,
517   CL_MEM_READ_WRITE,
518   CL_MEM_READ_ONLY,
519   CL_MEM_WRITE_ONLY,
520   CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
521   CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
522   CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR,
523   CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
524   CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,
525   CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR,
526   CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR,
527   CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR,
528   CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR | CL_MEM_COPY_HOST_PTR,
529   CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR,
530   CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
531   CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
532 };
533 
534 #define TEST_DEVICE_PARAM( device, paramName, val, name, type, cast )    \
535 error = clGetDeviceInfo( device, paramName, sizeof( val ), &val, &size );        \
536 test_error( error, "Unable to get device " name );                            \
537 if( size != sizeof( val ) )                \
538 {                                        \
539 log_error( "ERROR: Returned size of device " name " does not validate! (expected %d, got %d)\n", (int)sizeof( val ), (int)size );    \
540 return -1;    \
541 }                \
542 log_info( "\tReported device " name " : " type "\n", (cast)val );
543 
544 #define TEST_DEVICE_PARAM_MEM( device, paramName, val, name, type, div )    \
545 error = clGetDeviceInfo( device, paramName, sizeof( val ), &val, &size );        \
546 test_error( error, "Unable to get device " name );                            \
547 if( size != sizeof( val ) )                \
548 {                                        \
549 log_error( "ERROR: Returned size of device " name " does not validate! (expected %d, got %d)\n", (int)sizeof( val ), (int)size );    \
550 return -1;    \
551 }                \
552 log_info( "\tReported device " name " : " type "\n", (int)( val / div ) );
553 
test_get_device_info(cl_device_id deviceID,cl_context context,cl_command_queue ignoreQueue,int num_elements)554 int test_get_device_info(cl_device_id deviceID, cl_context context, cl_command_queue ignoreQueue, int num_elements)
555 {
556     int error;
557     size_t size;
558 
559     cl_uint vendorID;
560     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_VENDOR_ID, vendorID, "vendor ID", "0x%08x", int )
561 
562     char extensions[ 10240 ];
563     error = clGetDeviceInfo( deviceID, CL_DEVICE_EXTENSIONS, sizeof( extensions ), &extensions, &size );
564     test_error( error, "Unable to get device extensions" );
565     if( size != strlen( extensions ) + 1 )
566     {
567         log_error( "ERROR: Returned size of device extensions does not validate! (expected %d, got %d)\n", (int)( strlen( extensions ) + 1 ), (int)size );
568         return -1;
569     }
570     log_info( "\tReported device extensions: %s \n", extensions );
571 
572     cl_uint preferred;
573     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, preferred, "preferred vector char width", "%d", int )
574     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, preferred, "preferred vector short width", "%d", int )
575     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, preferred, "preferred vector int width", "%d", int )
576     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, preferred, "preferred vector long width", "%d", int )
577     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, preferred, "preferred vector float width", "%d", int )
578     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, preferred, "preferred vector double width", "%d", int )
579 
580     // Note that even if cl_khr_fp64, the preferred width for double can be non-zero.  For example, vendors
581     // extensions can support double but may not support cl_khr_fp64, which implies math library support.
582 
583     cl_uint baseAddrAlign;
584     TEST_DEVICE_PARAM(deviceID, CL_DEVICE_MEM_BASE_ADDR_ALIGN, baseAddrAlign,
585                       "base address alignment", "%d bits", int)
586 
587     cl_uint maxDataAlign;
588     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, maxDataAlign, "min data type alignment", "%d bytes", int )
589 
590     cl_device_mem_cache_type cacheType;
591     error = clGetDeviceInfo( deviceID, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof( cacheType ), &cacheType, &size );
592     test_error( error, "Unable to get device global mem cache type" );
593     if( size != sizeof( cacheType ) )
594     {
595         log_error( "ERROR: Returned size of device global mem cache type does not validate! (expected %d, got %d)\n", (int)sizeof( cacheType ), (int)size );
596         return -1;
597     }
598     const char *cacheTypeName = ( cacheType == CL_NONE ) ? "CL_NONE" : ( cacheType == CL_READ_ONLY_CACHE ) ? "CL_READ_ONLY_CACHE" : ( cacheType == CL_READ_WRITE_CACHE ) ? "CL_READ_WRITE_CACHE" : "<unknown>";
599     log_info( "\tReported device global mem cache type: %s \n", cacheTypeName );
600 
601     cl_uint cachelineSize;
602     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, cachelineSize, "global mem cacheline size", "%d bytes", int )
603 
604     cl_ulong cacheSize;
605     TEST_DEVICE_PARAM_MEM( deviceID, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, cacheSize, "global mem cache size", "%d KB", 1024 )
606 
607     cl_ulong memSize;
608     TEST_DEVICE_PARAM_MEM( deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, memSize, "global mem size", "%d MB", ( 1024 * 1024 ) )
609 
610     cl_device_local_mem_type localMemType;
611     error = clGetDeviceInfo( deviceID, CL_DEVICE_LOCAL_MEM_TYPE, sizeof( localMemType ), &localMemType, &size );
612     test_error( error, "Unable to get device local mem type" );
613     if( size != sizeof( cacheType ) )
614     {
615         log_error( "ERROR: Returned size of device local mem type does not validate! (expected %d, got %d)\n", (int)sizeof( localMemType ), (int)size );
616         return -1;
617     }
618     const char *localMemTypeName = ( localMemType == CL_LOCAL ) ? "CL_LOCAL" : ( cacheType == CL_GLOBAL ) ? "CL_GLOBAL" : "<unknown>";
619     log_info( "\tReported device local mem type: %s \n", localMemTypeName );
620 
621 
622     cl_bool errSupport;
623     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_ERROR_CORRECTION_SUPPORT, errSupport, "error correction support", "%d", int )
624 
625     size_t timerResolution;
626     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_PROFILING_TIMER_RESOLUTION, timerResolution, "profiling timer resolution", "%ld nanoseconds", long )
627 
628     cl_bool endian;
629     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_ENDIAN_LITTLE, endian, "little endian flag", "%d", int )
630 
631     cl_bool avail;
632     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_AVAILABLE, avail, "available flag", "%d", int )
633 
634     cl_bool compilerAvail;
635     TEST_DEVICE_PARAM( deviceID, CL_DEVICE_COMPILER_AVAILABLE, compilerAvail, "compiler available flag", "%d", int )
636 
637     char profile[ 1024 ];
638     error = clGetDeviceInfo( deviceID, CL_DEVICE_PROFILE, sizeof( profile ), &profile, &size );
639     test_error( error, "Unable to get device profile" );
640     if( size != strlen( profile ) + 1 )
641     {
642         log_error( "ERROR: Returned size of device profile does not validate! (expected %d, got %d)\n", (int)( strlen( profile ) + 1 ), (int)size );
643         return -1;
644     }
645     if( strcmp( profile, "FULL_PROFILE" ) != 0 && strcmp( profile, "EMBEDDED_PROFILE" ) != 0 )
646     {
647         log_error( "ERROR: Returned profile of device not FULL or EMBEDDED as required by OpenCL 1.2! (Returned %s)\n", profile );
648         return -1;
649     }
650     log_info( "\tReported device profile: %s \n", profile );
651 
652 
653     return 0;
654 }
655 
656 
657 
658 
659 static const char *sample_compile_size[2] = {
660     "__kernel void sample_test(__global int *src, __global int *dst)\n"
661     "{\n"
662     "    int  tid = get_global_id(0);\n"
663     "     dst[tid] = src[tid];\n"
664     "\n"
665     "}\n",
666     "__kernel __attribute__((reqd_work_group_size(%d,%d,%d))) void sample_test(__global int *src, __global int *dst)\n"
667     "{\n"
668     "    int  tid = get_global_id(0);\n"
669     "     dst[tid] = src[tid];\n"
670     "\n"
671     "}\n" };
672 
test_kernel_required_group_size(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)673 int test_kernel_required_group_size(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
674 {
675     int error;
676     size_t realSize;
677     size_t kernel_max_workgroup_size;
678     size_t global[] = {64,14,10};
679     size_t local[] = {0,0,0};
680 
681     cl_uint max_dimensions;
682 
683     error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(max_dimensions), &max_dimensions, NULL);
684     test_error(error,  "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS");
685     log_info("Device reported CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = %d.\n", (int)max_dimensions);
686 
687     {
688         clProgramWrapper program;
689         clKernelWrapper kernel;
690 
691         error = create_single_kernel_helper( context, &program, &kernel, 1, &sample_compile_size[ 0 ], "sample_test" );
692         if( error != 0 )
693             return error;
694 
695         error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_WORK_GROUP_SIZE, sizeof(kernel_max_workgroup_size), &kernel_max_workgroup_size, NULL);
696         test_error( error, "clGetKernelWorkGroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE");
697         log_info("The CL_KERNEL_WORK_GROUP_SIZE for the kernel is %d.\n", (int)kernel_max_workgroup_size);
698 
699         size_t size[ 3 ];
700         error = clGetKernelWorkGroupInfo( kernel, deviceID, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof( size ), size, &realSize );
701         test_error( error, "Unable to get work group info" );
702 
703         if( size[ 0 ] != 0 || size[ 1 ] != 0 || size[ 2 ] != 0 )
704         {
705             log_error( "ERROR: Nonzero compile work group size returned for nonspecified size! (returned %d,%d,%d)\n", (int)size[0], (int)size[1], (int)size[2] );
706             return -1;
707         }
708 
709         if( realSize != sizeof( size ) )
710         {
711             log_error( "ERROR: Returned size of compile work group size not valid! (Expected %d, got %d)\n", (int)sizeof( size ), (int)realSize );
712             return -1;
713         }
714 
715         // Determine some local dimensions to use for the test.
716         if (max_dimensions == 1) {
717             error = get_max_common_work_group_size(context, kernel, global[0], &local[0]);
718             test_error( error, "get_max_common_work_group_size failed");
719             log_info("For global dimension %d, kernel will require local dimension %d.\n", (int)global[0], (int)local[0]);
720         } else if (max_dimensions == 2) {
721             error = get_max_common_2D_work_group_size(context, kernel, global, local);
722             test_error( error, "get_max_common_2D_work_group_size failed");
723             log_info("For global dimension %d x %d, kernel will require local dimension %d x %d.\n", (int)global[0], (int)global[1], (int)local[0], (int)local[1]);
724         } else {
725             error = get_max_common_3D_work_group_size(context, kernel, global, local);
726             test_error( error, "get_max_common_3D_work_group_size failed");
727             log_info("For global dimension %d x %d x %d, kernel will require local dimension %d x %d x %d.\n",
728                      (int)global[0], (int)global[1], (int)global[2], (int)local[0], (int)local[1], (int)local[2]);
729         }
730     }
731 
732 
733     {
734         clProgramWrapper program;
735         clKernelWrapper kernel;
736         clMemWrapper in, out;
737         //char source[1024];
738         char *source = (char*)malloc(1024);
739         source[0] = '\0';
740 
741         sprintf(source, sample_compile_size[1], local[0], local[1], local[2]);
742 
743         error = create_single_kernel_helper( context, &program, &kernel, 1, (const char**)&source, "sample_test" );
744         if( error != 0 )
745             return error;
746 
747         size_t size[ 3 ];
748         error = clGetKernelWorkGroupInfo( kernel, deviceID, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof( size ), size, &realSize );
749         test_error( error, "Unable to get work group info" );
750 
751         if( size[ 0 ] != local[0] || size[ 1 ] != local[1] || size[ 2 ] != local[2] )
752         {
753             log_error( "ERROR: Incorrect compile work group size returned for specified size! (returned %d,%d,%d, expected %d,%d,%d)\n",
754                       (int)size[0], (int)size[1], (int)size[2], (int)local[0], (int)local[1], (int)local[2]);
755             return -1;
756         }
757 
758         // Verify that the kernel will only execute with that size.
759         in = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_int)*global[0], NULL, &error);
760         test_error(error, "clCreateBuffer failed");
761         out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_int)*global[0], NULL, &error);
762         test_error(error, "clCreateBuffer failed");
763 
764         error = clSetKernelArg(kernel, 0, sizeof(in), &in);
765         test_error(error, "clSetKernelArg failed");
766         error = clSetKernelArg(kernel, 1, sizeof(out), &out);
767         test_error(error, "clSetKernelArg failed");
768 
769         error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, local, 0, NULL, NULL);
770         test_error(error, "clEnqueueNDRangeKernel failed");
771 
772         error = clFinish(queue);
773         test_error(error, "clFinish failed");
774 
775         log_info("kernel_required_group_size may report spurious ERRORS in the conformance log.\n");
776 
777         local[0]++;
778         error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, local, 0, NULL, NULL);
779         if (error != CL_INVALID_WORK_GROUP_SIZE) {
780             log_error("Incorrect error returned for executing a kernel with the wrong required local work group size. (used %d,%d,%d, required %d,%d,%d)\n",
781                       (int)local[0], (int)local[1], (int)local[2], (int)local[0]-1, (int)local[1], (int)local[2] );
782             print_error(error, "Expected: CL_INVALID_WORK_GROUP_SIZE.");
783             return -1;
784         }
785 
786         error = clFinish(queue);
787         test_error(error, "clFinish failed");
788 
789         if (max_dimensions == 1) {
790             free(source);
791             return 0;
792         }
793 
794         local[0]--; local[1]++;
795         error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, local, 0, NULL, NULL);
796         if (error != CL_INVALID_WORK_GROUP_SIZE) {
797             log_error("Incorrect error returned for executing a kernel with the wrong required local work group size. (used %d,%d,%d, required %d,%d,%d)\n",
798                       (int)local[0], (int)local[1], (int)local[2], (int)local[0]-1, (int)local[1], (int)local[2]);
799             print_error(error, "Expected: CL_INVALID_WORK_GROUP_SIZE.");
800             return -1;
801         }
802 
803         error = clFinish(queue);
804         test_error(error, "clFinish failed");
805 
806         if (max_dimensions == 2) {
807             return 0;
808             free(source);
809         }
810 
811         local[1]--; local[2]++;
812         error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, local, 0, NULL, NULL);
813         if (error != CL_INVALID_WORK_GROUP_SIZE) {
814             log_error("Incorrect error returned for executing a kernel with the wrong required local work group size. (used %d,%d,%d, required %d,%d,%d)\n",
815                       (int)local[0], (int)local[1], (int)local[2], (int)local[0]-1, (int)local[1], (int)local[2]);
816             print_error(error, "Expected: CL_INVALID_WORK_GROUP_SIZE.");
817             return -1;
818         }
819 
820         error = clFinish(queue);
821         test_error(error, "clFinish failed");
822         free(source);
823     }
824 
825     return 0;
826 }
827 
828 
829