xref: /aosp_15_r20/external/OpenCL-CTS/test_conformance/select/test_select.cpp (revision 6467f958c7de8070b317fc65bcb0f6472e388d82)
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