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