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 "harness/compat.h"
17 #include "harness/typeWrappers.h"
18
19 #include <assert.h>
20 #include <stdio.h>
21 #include <time.h>
22 #include <string.h>
23
24 #include <cinttypes>
25 #include <vector>
26
27 #if ! defined( _WIN32)
28 #if defined(__APPLE__)
29 #include <sys/sysctl.h>
30 #endif
31 #endif
32 #include <limits.h>
33 #include "test_select.h"
34
35 #include "harness/testHarness.h"
36 #include "harness/kernelHelpers.h"
37 #include "harness/mt19937.h"
38 #include "harness/parseParameters.h"
39
40
41 //-----------------------------------------
42 // Static functions
43 //-----------------------------------------
44
45 // initialize src1 and src2 buffer with values based on stype
46 static void initSrcBuffer(void* src1, Type stype, MTdata);
47
48 // initialize the valued used to compare with in the select with
49 // vlaues [start, count)
50 static void initCmpBuffer(void *cmp, Type cmptype, uint64_t start,
51 const size_t count);
52
53 // make a program that uses select for the given stype (src/dest type),
54 // ctype (comparison type), veclen (vector length)
55 static cl_program makeSelectProgram(cl_kernel *kernel_ptr, cl_context context,
56 Type stype, Type ctype,
57 const size_t veclen);
58
59 // Creates and execute the select test for the given device, context,
60 // stype (source/dest type), cmptype (comparison type), using max_tg_size
61 // number of threads. It runs test for all the different vector lengths
62 // for the given stype and cmptype.
63 static int doTest(cl_command_queue queue, cl_context context,
64 Type stype, Type cmptype, cl_device_id device);
65
66
67 static void printUsage( void );
68
69 //-----------------------------------------
70 // Definitions and initializations
71 //-----------------------------------------
72
73 // Define the buffer size that we want to block our test with
74 #define BUFFER_SIZE (1024*1024)
75 #define KPAGESIZE 4096
76
77 #define test_error_count(errCode, msg) \
78 { \
79 auto errCodeResult = errCode; \
80 if (errCodeResult != CL_SUCCESS) \
81 { \
82 gFailCount++; \
83 print_error(errCodeResult, msg); \
84 return errCode; \
85 } \
86 }
87
88 // When we indicate non wimpy mode, the types that are 32 bits value will
89 // test their entire range and 64 bits test will test the 32 bit
90 // range. Otherwise, we test a subset of the range
91 // [-min_short, min_short]
92 static bool s_wimpy_mode = false;
93 static int s_wimpy_reduction_factor = 256;
94
95 //-----------------------------------------
96 // Static helper functions
97 //-----------------------------------------
98
99 // calculates log2 for a 32 bit number
int_log2(size_t value)100 int int_log2(size_t value) {
101 if( 0 == value )
102 return INT_MIN;
103
104 #if defined( __GNUC__ )
105 return (unsigned) (8*sizeof(size_t) - 1UL - __builtin_clzl(value));
106 #else
107 int result = -1;
108 while(value)
109 {
110 result++;
111 value >>= 1;
112 }
113 return result;
114 #endif
115 }
116
117
initSrcBuffer(void * src1,Type stype,MTdata d)118 static void initSrcBuffer(void* src1, Type stype, MTdata d)
119 {
120 unsigned int* s1 = (unsigned int *)src1;
121 size_t i;
122
123 for ( i=0 ; i < BUFFER_SIZE/sizeof(cl_int); i++)
124 s1[i] = genrand_int32(d);
125 }
126
initCmpBuffer(void * cmp,Type cmptype,uint64_t start,const size_t count)127 static void initCmpBuffer(void *cmp, Type cmptype, uint64_t start,
128 const size_t count)
129
130 {
131 assert(cmptype != kfloat);
132 switch (type_size[cmptype]) {
133 case 1: {
134 uint8_t* ub = (uint8_t *)cmp;
135 for (size_t i = 0; i < count; ++i) ub[i] = (uint8_t)start++;
136 break;
137 }
138 case 2: {
139 uint16_t* us = (uint16_t *)cmp;
140 for (size_t i = 0; i < count; ++i) us[i] = (uint16_t)start++;
141 break;
142 }
143 case 4: {
144 if (!s_wimpy_mode) {
145 uint32_t* ui = (uint32_t *)cmp;
146 for (size_t i = 0; i < count; ++i) ui[i] = (uint32_t)start++;
147 }
148 else {
149 // The short test doesn't iterate over the entire 32 bit space so
150 // we alternate between positive and negative values
151 int32_t* ui = (int32_t *)cmp;
152 int32_t neg_start = (int32_t)start * -1;
153 for (size_t i = 0; i < count; i++)
154 {
155 ++start;
156 --neg_start;
157 ui[i] = (int32_t)((i % 2) ? start : neg_start);
158 }
159 }
160 break;
161 }
162 case 8: {
163 // We don't iterate over the entire space of 64 bit so for the
164 // selects, we want to test positive and negative values
165 int64_t* ll = (int64_t *)cmp;
166 int64_t neg_start = (int64_t)start * -1;
167 for (size_t i = 0; i < count; i++)
168 {
169 ++start;
170 --neg_start;
171 ll[i] = (int64_t)((i % 2) ? start : neg_start);
172 }
173 break;
174 }
175 default:
176 log_error("invalid cmptype %s\n",type_name[cmptype]);
177 } // end switch
178 }
179
180 // Make the various incarnations of the program we want to run
181 // stype: source and destination type for the select
182 // ctype: compare type
makeSelectProgram(cl_kernel * kernel_ptr,const cl_context context,Type srctype,Type cmptype,const size_t vec_len)183 static cl_program makeSelectProgram(cl_kernel *kernel_ptr,
184 const cl_context context, Type srctype,
185 Type cmptype, const size_t vec_len)
186 {
187 char testname[256];
188 char stypename[32];
189 char ctypename[32];
190 char extension[128] = "";
191 int err = 0;
192
193 const char *source[] = {
194 extension,
195 "__kernel void ", testname,
196 "(__global ", stypename, " *dest, __global ", stypename, " *src1,\n __global ",
197 stypename, " *src2, __global ", ctypename, " *cmp)\n",
198 "{\n"
199 " size_t tid = get_global_id(0);\n"
200 " if( tid < get_global_size(0) )\n"
201 " dest[tid] = select(src1[tid], src2[tid], cmp[tid]);\n"
202 "}\n"
203 };
204
205
206 const char *sourceV3[] = {
207 extension,
208 "__kernel void ", testname,
209 "(__global ", stypename, " *dest, __global ", stypename, " *src1,\n __global ",
210 stypename, " *src2, __global ", ctypename, " *cmp)\n",
211 "{\n"
212 " size_t tid = get_global_id(0);\n"
213 " size_t size = get_global_size(0);\n"
214 " if( tid + 1 < size ) // can't run off the end\n"
215 " vstore3( select( vload3(tid, src1), vload3(tid, src2), vload3(tid, cmp)), tid, dest );\n"
216 " else if(tid + 1 == size)\n"
217 " {\n"
218 // If the size is odd, then we have odd * 3 elements, which is an odd number of scalars in the array
219 // If the size is even, then we have even * 3 elements, which is an even number of scalars in the array
220 // 3 will never divide evenly into a power of two sized buffer, so the last vec3 will overhang by 1 or 2.
221 // The only even number x in power_of_two < x <= power_of_two+2 is power_of_two+2.
222 // The only odd number x in power_of_two < x <= power_of_two+2 is power_of_two+1.
223 // Therefore, odd sizes overhang the end of the array by 1, and even sizes overhang by 2.
224 " size_t leftovers = 1 + (size & 1);\n"
225 " ", stypename, "3 a, b; \n"
226 " ", ctypename, "3 c;\n"
227 " switch( leftovers ) \n"
228 " {\n"
229 " case 2:\n"
230 " a.y = src1[3*tid+1];\n"
231 " b.y = src2[3*tid+1];\n"
232 " c.y = cmp[3*tid+1];\n"
233 " // fall through \n"
234 " case 1:\n"
235 " a.x = src1[3*tid];\n"
236 " b.x = src2[3*tid];\n"
237 " c.x = cmp[3*tid];\n"
238 " break;\n"
239 " }\n"
240 " a = select( a, b, c );\n"
241 " switch( leftovers ) \n"
242 " {\n"
243 " case 2:\n"
244 " dest[3*tid+1] = a.y;\n"
245 " // fall through \n"
246 " case 1:\n"
247 " dest[3*tid] = a.x;\n"
248 " break;\n"
249 " }\n"
250 " }\n"
251 "}\n"
252 };
253
254 if (srctype == kdouble)
255 strcpy( extension, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" );
256
257 if (srctype == khalf)
258 strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n");
259
260 // create type name and testname
261 switch( vec_len )
262 {
263 case 1:
264 strncpy(stypename, type_name[srctype], sizeof(stypename));
265 strncpy(ctypename, type_name[cmptype], sizeof(ctypename));
266 snprintf(testname, sizeof(testname), "select_%s_%s", stypename, ctypename );
267 log_info("Building %s(%s, %s, %s)\n", testname, stypename, stypename, ctypename);
268 break;
269 case 3:
270 strncpy(stypename, type_name[srctype], sizeof(stypename));
271 strncpy(ctypename, type_name[cmptype], sizeof(ctypename));
272 snprintf(testname, sizeof(testname), "select_%s3_%s3", stypename, ctypename );
273 log_info("Building %s(%s3, %s3, %s3)\n", testname, stypename, stypename, ctypename);
274 break;
275 case 2:
276 case 4:
277 case 8:
278 case 16:
279 snprintf(stypename,sizeof(stypename), "%s%d", type_name[srctype],(int)vec_len);
280 snprintf(ctypename,sizeof(ctypename), "%s%d", type_name[cmptype],(int)vec_len);
281 snprintf(testname, sizeof(testname), "select_%s_%s", stypename, ctypename );
282 log_info("Building %s(%s, %s, %s)\n", testname, stypename, stypename, ctypename);
283 break;
284 default:
285 log_error( "Unkown vector type. Aborting...\n" );
286 exit(-1);
287 break;
288 }
289
290 /*
291 int j;
292 for( j = 0; j < sizeof( source ) / sizeof( source[0] ); j++ )
293 log_info( "%s", source[j] );
294 */
295
296 // create program
297 cl_program program;
298 const char **psrc = vec_len == 3 ? sourceV3 : source;
299 size_t src_size = vec_len == 3 ? ARRAY_SIZE(sourceV3) : ARRAY_SIZE(source);
300
301 if (create_single_kernel_helper(context, &program, kernel_ptr, src_size,
302 psrc, testname))
303 {
304 log_error("Failed to build program (%d)\n", err);
305 return NULL;
306 }
307
308 return program;
309 }
310
311 #define VECTOR_SIZE_COUNT 6
312
doTest(cl_command_queue queue,cl_context context,Type stype,Type cmptype,cl_device_id device)313 static int doTest(cl_command_queue queue, cl_context context, Type stype, Type cmptype, cl_device_id device)
314 {
315 int err = CL_SUCCESS;
316 MTdataHolder d(gRandomSeed);
317 const size_t element_count[VECTOR_SIZE_COUNT] = { 1, 2, 3, 4, 8, 16 };
318 clMemWrapper src1, src2, cmp, dest;
319
320 cl_ulong blocks = type_size[stype] * 0x100000000ULL / BUFFER_SIZE;
321 const size_t block_elements = BUFFER_SIZE / type_size[stype];
322 size_t step = s_wimpy_mode ? s_wimpy_reduction_factor : 1;
323 cl_ulong cmp_stride = block_elements * step;
324
325 // It is more efficient to create the tests all at once since we
326 // use the same test data on each of the vector sizes
327 clProgramWrapper programs[VECTOR_SIZE_COUNT];
328 clKernelWrapper kernels[VECTOR_SIZE_COUNT];
329
330 if (stype == kdouble && !is_extension_available(device, "cl_khr_fp64"))
331 {
332 log_info("Skipping double because cl_khr_fp64 extension is not supported.\n");
333 return 0;
334 }
335
336 if (stype == khalf && !is_extension_available(device, "cl_khr_fp16"))
337 {
338 log_info(
339 "Skipping half because cl_khr_fp16 extension is not supported.\n");
340 return 0;
341 }
342
343 if (gIsEmbedded)
344 {
345 if (( stype == klong || stype == kulong ) && ! is_extension_available( device, "cles_khr_int64" ))
346 {
347 log_info("Long types unsupported, skipping.");
348 return 0;
349 }
350
351 if (( cmptype == klong || cmptype == kulong ) && ! is_extension_available( device, "cles_khr_int64" ))
352 {
353 log_info("Long types unsupported, skipping.");
354 return 0;
355 }
356 }
357
358 src1 = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
359 test_error_count(err, "Error: could not allocate src1 buffer\n");
360 src2 = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
361 test_error_count(err, "Error: could not allocate src2 buffer\n");
362 cmp = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
363 test_error_count(err, "Error: could not allocate cmp buffer\n");
364 dest = clCreateBuffer( context, CL_MEM_WRITE_ONLY, BUFFER_SIZE, NULL, &err );
365 test_error_count(err, "Error: could not allocate dest buffer\n");
366
367 programs[0] = makeSelectProgram(&kernels[0], context, stype, cmptype,
368 element_count[0]);
369 programs[1] = makeSelectProgram(&kernels[1], context, stype, cmptype,
370 element_count[1]);
371 programs[2] = makeSelectProgram(&kernels[2], context, stype, cmptype,
372 element_count[2]);
373 programs[3] = makeSelectProgram(&kernels[3], context, stype, cmptype,
374 element_count[3]);
375 programs[4] = makeSelectProgram(&kernels[4], context, stype, cmptype,
376 element_count[4]);
377 programs[5] = makeSelectProgram(&kernels[5], context, stype, cmptype,
378 element_count[5]);
379
380 for (size_t vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
381 {
382 if (!programs[vecsize] || !kernels[vecsize])
383 {
384 return -1;
385 }
386
387 err = clSetKernelArg(kernels[vecsize], 0, sizeof dest, &dest);
388 test_error_count(err, "Error: Cannot set kernel arg dest!\n");
389 err = clSetKernelArg(kernels[vecsize], 1, sizeof src1, &src1);
390 test_error_count(err, "Error: Cannot set kernel arg dest!\n");
391 err = clSetKernelArg(kernels[vecsize], 2, sizeof src2, &src2);
392 test_error_count(err, "Error: Cannot set kernel arg dest!\n");
393 err = clSetKernelArg(kernels[vecsize], 3, sizeof cmp, &cmp);
394 test_error_count(err, "Error: Cannot set kernel arg dest!\n");
395 }
396
397 std::vector<char> ref(BUFFER_SIZE);
398 std::vector<char> sref(BUFFER_SIZE);
399 std::vector<char> src1_host(BUFFER_SIZE);
400 std::vector<char> src2_host(BUFFER_SIZE);
401 std::vector<char> cmp_host(BUFFER_SIZE);
402 std::vector<char> dest_host(BUFFER_SIZE);
403
404 // We block the test as we are running over the range of compare values
405 // "block the test" means "break the test into blocks"
406 if( type_size[stype] == 4 )
407 cmp_stride = block_elements * step * (0x100000000ULL / 0x100000000ULL);
408 if( type_size[stype] == 8 )
409 cmp_stride = block_elements * step * (0xffffffffffffffffULL / 0x100000000ULL + 1);
410
411 log_info("Testing...");
412 uint64_t i;
413
414 initSrcBuffer(src1_host.data(), stype, d);
415 initSrcBuffer(src2_host.data(), stype, d);
416 for (i=0; i < blocks; i+=step)
417 {
418 initCmpBuffer(cmp_host.data(), cmptype, i * cmp_stride, block_elements);
419
420 err = clEnqueueWriteBuffer(queue, src1, CL_FALSE, 0, BUFFER_SIZE,
421 src1_host.data(), 0, NULL, NULL);
422 test_error_count(err, "Error: Could not write src1");
423
424 err = clEnqueueWriteBuffer(queue, src2, CL_FALSE, 0, BUFFER_SIZE,
425 src2_host.data(), 0, NULL, NULL);
426 test_error_count(err, "Error: Could not write src2");
427
428 err = clEnqueueWriteBuffer(queue, cmp, CL_FALSE, 0, BUFFER_SIZE,
429 cmp_host.data(), 0, NULL, NULL);
430 test_error_count(err, "Error: Could not write cmp");
431
432 Select sfunc = (cmptype == ctype[stype][0]) ? vrefSelects[stype][0]
433 : vrefSelects[stype][1];
434 (*sfunc)(ref.data(), src1_host.data(), src2_host.data(),
435 cmp_host.data(), block_elements);
436
437 sfunc = (cmptype == ctype[stype][0]) ? refSelects[stype][0]
438 : refSelects[stype][1];
439 (*sfunc)(sref.data(), src1_host.data(), src2_host.data(),
440 cmp_host.data(), block_elements);
441
442 for (int vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
443 {
444 size_t vector_size = element_count[vecsize] * type_size[stype];
445 size_t vector_count = (BUFFER_SIZE + vector_size - 1) / vector_size;
446
447 const cl_int pattern = -1;
448 err = clEnqueueFillBuffer(queue, dest, &pattern, sizeof(cl_int), 0,
449 BUFFER_SIZE, 0, nullptr, nullptr);
450 test_error_count(err, "clEnqueueFillBuffer failed");
451
452
453 err = clEnqueueNDRangeKernel(queue, kernels[vecsize], 1, NULL, &vector_count, NULL, 0, NULL, NULL);
454 test_error_count(err, "clEnqueueNDRangeKernel failed errcode\n");
455
456 err = clEnqueueReadBuffer(queue, dest, CL_TRUE, 0, BUFFER_SIZE,
457 dest_host.data(), 0, NULL, NULL);
458 test_error_count(
459 err, "Error: Reading buffer from dest to dest_host failed\n");
460
461 if ((*checkResults[stype])(dest_host.data(),
462 vecsize == 0 ? sref.data() : ref.data(),
463 block_elements, element_count[vecsize])
464 != 0)
465 {
466 log_error("vec_size:%d indx: 0x%16.16" PRIx64 "\n",
467 (int)element_count[vecsize], i);
468 return TEST_FAIL;
469 }
470 } // for vecsize
471 } // for i
472
473 if (!s_wimpy_mode)
474 log_info(" Passed\n\n");
475 else
476 log_info(" Wimpy Passed\n\n");
477
478 return err;
479 }
480
test_select_uchar_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)481 int test_select_uchar_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
482 {
483 return doTest(queue, context, kuchar, kuchar, deviceID);
484 }
test_select_uchar_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)485 int test_select_uchar_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
486 {
487 return doTest(queue, context, kuchar, kchar, deviceID);
488 }
test_select_char_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)489 int test_select_char_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
490 {
491 return doTest(queue, context, kchar, kuchar, deviceID);
492 }
test_select_char_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)493 int test_select_char_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
494 {
495 return doTest(queue, context, kchar, kchar, deviceID);
496 }
test_select_ushort_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)497 int test_select_ushort_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
498 {
499 return doTest(queue, context, kushort, kushort, deviceID);
500 }
test_select_ushort_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)501 int test_select_ushort_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
502 {
503 return doTest(queue, context, kushort, kshort, deviceID);
504 }
test_select_short_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)505 int test_select_short_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
506 {
507 return doTest(queue, context, kshort, kushort, deviceID);
508 }
test_select_short_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)509 int test_select_short_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
510 {
511 return doTest(queue, context, kshort, kshort, deviceID);
512 }
test_select_half_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)513 int test_select_half_ushort(cl_device_id deviceID, cl_context context,
514 cl_command_queue queue, int num_elements)
515 {
516 return doTest(queue, context, khalf, kushort, deviceID);
517 }
test_select_half_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)518 int test_select_half_short(cl_device_id deviceID, cl_context context,
519 cl_command_queue queue, int num_elements)
520 {
521 return doTest(queue, context, khalf, kshort, deviceID);
522 }
test_select_uint_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)523 int test_select_uint_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
524 {
525 return doTest(queue, context, kuint, kuint, deviceID);
526 }
test_select_uint_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)527 int test_select_uint_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
528 {
529 return doTest(queue, context, kuint, kint, deviceID);
530 }
test_select_int_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)531 int test_select_int_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
532 {
533 return doTest(queue, context, kint, kuint, deviceID);
534 }
test_select_int_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)535 int test_select_int_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
536 {
537 return doTest(queue, context, kint, kint, deviceID);
538 }
test_select_float_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)539 int test_select_float_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
540 {
541 return doTest(queue, context, kfloat, kuint, deviceID);
542 }
test_select_float_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)543 int test_select_float_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
544 {
545 return doTest(queue, context, kfloat, kint, deviceID);
546 }
test_select_ulong_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)547 int test_select_ulong_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
548 {
549 return doTest(queue, context, kulong, kulong, deviceID);
550 }
test_select_ulong_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)551 int test_select_ulong_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
552 {
553 return doTest(queue, context, kulong, klong, deviceID);
554 }
test_select_long_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)555 int test_select_long_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
556 {
557 return doTest(queue, context, klong, kulong, deviceID);
558 }
test_select_long_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)559 int test_select_long_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
560 {
561 return doTest(queue, context, klong, klong, deviceID);
562 }
test_select_double_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)563 int test_select_double_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
564 {
565 return doTest(queue, context, kdouble, kulong, deviceID);
566 }
test_select_double_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)567 int test_select_double_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
568 {
569 return doTest(queue, context, kdouble, klong, deviceID);
570 }
571
572 test_definition test_list[] = {
573 ADD_TEST(select_uchar_uchar), ADD_TEST(select_uchar_char),
574 ADD_TEST(select_char_uchar), ADD_TEST(select_char_char),
575 ADD_TEST(select_ushort_ushort), ADD_TEST(select_ushort_short),
576 ADD_TEST(select_short_ushort), ADD_TEST(select_short_short),
577 ADD_TEST(select_half_ushort), ADD_TEST(select_half_short),
578 ADD_TEST(select_uint_uint), ADD_TEST(select_uint_int),
579 ADD_TEST(select_int_uint), ADD_TEST(select_int_int),
580 ADD_TEST(select_float_uint), ADD_TEST(select_float_int),
581 ADD_TEST(select_ulong_ulong), ADD_TEST(select_ulong_long),
582 ADD_TEST(select_long_ulong), ADD_TEST(select_long_long),
583 ADD_TEST(select_double_ulong), ADD_TEST(select_double_long),
584 };
585
586 const int test_num = ARRAY_SIZE( test_list );
587
main(int argc,const char * argv[])588 int main(int argc, const char* argv[])
589 {
590 test_start();
591
592 argc = parseCustomParam(argc, argv);
593 if (argc == -1)
594 {
595 return EXIT_FAILURE;
596 }
597
598 const char ** argList = (const char **)calloc( argc, sizeof( char*) );
599
600 if( NULL == argList )
601 {
602 log_error( "Failed to allocate memory for argList array.\n" );
603 return 1;
604 }
605
606 argList[0] = argv[0];
607 size_t argCount = 1;
608
609 for( int i = 1; i < argc; ++i )
610 {
611 const char *arg = argv[i];
612 if (arg == NULL)
613 break;
614
615 if (arg[0] == '-')
616 {
617 arg++;
618 while(*arg != '\0')
619 {
620 switch(*arg) {
621 case 'h':
622 printUsage();
623 return 0;
624 case 'w':
625 s_wimpy_mode = true;
626 break;
627 case '[':
628 parseWimpyReductionFactor(arg, s_wimpy_reduction_factor);
629 break;
630 default:
631 break;
632 }
633 arg++;
634 }
635 }
636 else
637 {
638 argList[argCount] = arg;
639 argCount++;
640 }
641 }
642
643 if (getenv("CL_WIMPY_MODE")) {
644 s_wimpy_mode = true;
645 }
646
647 if (s_wimpy_mode) {
648 log_info("\n");
649 log_info("*** WARNING: Testing in Wimpy mode! ***\n");
650 log_info("*** Wimpy mode is not sufficient to verify correctness. ***\n");
651 log_info("*** It gives warm fuzzy feelings and then nevers calls. ***\n\n");
652 log_info("*** Wimpy Reduction Factor: %-27u ***\n\n", s_wimpy_reduction_factor);
653 }
654
655 int err = runTestHarness(argCount, argList, test_num, test_list, false, 0);
656
657 free( argList );
658
659 return err;
660 }
661
printUsage(void)662 static void printUsage( void )
663 {
664 log_info("test_select: [-w] <optional: test_names> \n");
665 log_info("\tdefault is to run the full test on the default device\n");
666 log_info("\t-w run in wimpy mode (smoke test)\n");
667 log_info("\t-[2^n] Set wimpy reduction factor, recommended range of n is 1-12, default factor(%u)\n", s_wimpy_reduction_factor);
668 log_info("\n");
669 log_info("Test names:\n");
670 for( int i = 0; i < test_num; i++ )
671 {
672 log_info( "\t%s\n", test_list[i].name );
673 }
674 }
675