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/conversions.h"
18 #include "harness/typeWrappers.h"
19 #include "harness/testHarness.h"
20
21 // clang-format off
22
23 const char *anyAllTestKernelPattern =
24 "%s\n" // optional pragma
25 "%s\n" // optional pragma
26 "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n"
27 "{\n"
28 " int tid = get_global_id(0);\n"
29 " destValues[tid] = %s( sourceA[tid] );\n"
30 "\n"
31 "}\n";
32
33 const char *anyAllTestKernelPatternVload =
34 "%s\n" // optional pragma
35 "%s\n" // optional pragma
36 "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n"
37 "{\n"
38 " int tid = get_global_id(0);\n"
39 " destValues[tid] = %s(vload3(tid, (__global %s *)sourceA));\n" // ugh, almost
40 "\n"
41 "}\n";
42
43 // clang-format on
44
45 #define TEST_SIZE 512
46
47 typedef int (*anyAllVerifyFn)( ExplicitType vecType, unsigned int vecSize, void *inData );
48
test_any_all_kernel(cl_context context,cl_command_queue queue,const char * fnName,ExplicitType vecType,unsigned int vecSize,anyAllVerifyFn verifyFn,MTdata d)49 int test_any_all_kernel(cl_context context, cl_command_queue queue,
50 const char *fnName, ExplicitType vecType,
51 unsigned int vecSize, anyAllVerifyFn verifyFn,
52 MTdata d )
53 {
54 clProgramWrapper program;
55 clKernelWrapper kernel;
56 clMemWrapper streams[2];
57 cl_long inDataA[TEST_SIZE * 16], clearData[TEST_SIZE * 16];
58 int outData[TEST_SIZE];
59 int error, i;
60 size_t threads[1], localThreads[1];
61 char kernelSource[10240];
62 char *programPtr;
63 char sizeName[4];
64
65
66 /* Create the source */
67 if( g_vector_aligns[vecSize] == 1 ) {
68 sizeName[ 0 ] = 0;
69 } else {
70 sprintf( sizeName, "%d", vecSize );
71 }
72 log_info("Testing any/all on %s%s\n",
73 get_explicit_type_name( vecType ), sizeName);
74 if(DENSE_PACK_VECS && vecSize == 3) {
75 // anyAllTestKernelPatternVload
76 sprintf(
77 kernelSource, anyAllTestKernelPatternVload,
78 vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
79 : "",
80 vecType == kHalf ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable"
81 : "",
82 get_explicit_type_name(vecType), sizeName, fnName,
83 get_explicit_type_name(vecType));
84 } else {
85 sprintf(
86 kernelSource, anyAllTestKernelPattern,
87 vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
88 : "",
89 vecType == kHalf ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable"
90 : "",
91 get_explicit_type_name(vecType), sizeName, fnName);
92 }
93 /* Create kernels */
94 programPtr = kernelSource;
95 if( create_single_kernel_helper( context, &program, &kernel, 1,
96 (const char **)&programPtr,
97 "sample_test" ) )
98 {
99 return -1;
100 }
101
102 /* Generate some streams */
103 generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataA );
104 memset( clearData, 0, sizeof( clearData ) );
105
106 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
107 get_explicit_type_size(vecType)
108 * g_vector_aligns[vecSize] * TEST_SIZE,
109 &inDataA, &error);
110 if( streams[0] == NULL )
111 {
112 print_error( error, "Creating input array A failed!\n");
113 return -1;
114 }
115 streams[1] =
116 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
117 sizeof(cl_int) * g_vector_aligns[vecSize] * TEST_SIZE,
118 clearData, &error);
119 if( streams[1] == NULL )
120 {
121 print_error( error, "Creating output array failed!\n");
122 return -1;
123 }
124
125 /* Assign streams and execute */
126 error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
127 test_error( error, "Unable to set indexed kernel arguments" );
128 error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
129 test_error( error, "Unable to set indexed kernel arguments" );
130
131 /* Run the kernel */
132 threads[0] = TEST_SIZE;
133
134 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
135 test_error( error, "Unable to get work group size to use" );
136
137 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
138 test_error( error, "Unable to execute test kernel" );
139
140 /* Now get the results */
141 error = clEnqueueReadBuffer( queue, streams[1], true, 0, sizeof( int ) * TEST_SIZE, outData, 0, NULL, NULL );
142 test_error( error, "Unable to read output array!" );
143
144 /* And verify! */
145 for( i = 0; i < TEST_SIZE; i++ )
146 {
147 int expected = verifyFn( vecType, vecSize, (char *)inDataA + i * get_explicit_type_size( vecType ) * g_vector_aligns[vecSize] );
148 if( expected != outData[ i ] )
149 {
150 unsigned int *ptr = (unsigned int *)( (char *)inDataA + i * get_explicit_type_size( vecType ) * g_vector_aligns[vecSize] );
151 log_error( "ERROR: Data sample %d does not validate! Expected (%d), got (%d), source 0x%08x\n",
152 i, expected, outData[i], *ptr );
153 return -1;
154 }
155 }
156
157 return 0;
158 }
159
anyVerifyFn(ExplicitType vecType,unsigned int vecSize,void * inData)160 int anyVerifyFn( ExplicitType vecType, unsigned int vecSize, void *inData )
161 {
162 unsigned int i;
163 switch( vecType )
164 {
165 case kChar:
166 {
167 char sum = 0;
168 char *tData = (char *)inData;
169 for( i = 0; i < vecSize; i++ )
170 sum |= tData[ i ] & 0x80;
171 return (sum != 0) ? 1 : 0;
172 }
173 case kShort:
174 {
175 short sum = 0;
176 short *tData = (short *)inData;
177 for( i = 0; i < vecSize; i++ )
178 sum |= tData[ i ] & 0x8000;
179 return (sum != 0);
180 }
181 case kInt:
182 {
183 cl_int sum = 0;
184 cl_int *tData = (cl_int *)inData;
185 for( i = 0; i < vecSize; i++ )
186 sum |= tData[ i ] & (cl_int)0x80000000L;
187 return (sum != 0);
188 }
189 case kLong:
190 {
191 cl_long sum = 0;
192 cl_long *tData = (cl_long *)inData;
193 for( i = 0; i < vecSize; i++ )
194 sum |= tData[ i ] & 0x8000000000000000LL;
195 return (sum != 0);
196 }
197 default:
198 return 0;
199 }
200 }
201
test_relational_any(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)202 int test_relational_any(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
203 {
204 ExplicitType vecType[] = { kChar, kShort, kInt, kLong };
205 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
206 unsigned int index, typeIndex;
207 int retVal = 0;
208 RandomSeed seed(gRandomSeed );
209
210 for( typeIndex = 0; typeIndex < 4; typeIndex++ )
211 {
212 if (vecType[typeIndex] == kLong && !gHasLong)
213 continue;
214
215 for( index = 0; vecSizes[ index ] != 0; index++ )
216 {
217 // Test!
218 if( test_any_all_kernel(context, queue, "any", vecType[ typeIndex ], vecSizes[ index ], anyVerifyFn, seed ) != 0 )
219 {
220 log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] );
221 retVal = -1;
222 }
223 }
224 }
225
226 return retVal;
227 }
228
allVerifyFn(ExplicitType vecType,unsigned int vecSize,void * inData)229 int allVerifyFn( ExplicitType vecType, unsigned int vecSize, void *inData )
230 {
231 unsigned int i;
232 switch( vecType )
233 {
234 case kChar:
235 {
236 char sum = 0x80;
237 char *tData = (char *)inData;
238 for( i = 0; i < vecSize; i++ )
239 sum &= tData[ i ] & 0x80;
240 return (sum != 0) ? 1 : 0;
241 }
242 case kShort:
243 {
244 short sum = 0x8000;
245 short *tData = (short *)inData;
246 for( i = 0; i < vecSize; i++ )
247 sum &= tData[ i ] & 0x8000;
248 return (sum != 0);
249 }
250 case kInt:
251 {
252 cl_int sum = 0x80000000L;
253 cl_int *tData = (cl_int *)inData;
254 for( i = 0; i < vecSize; i++ )
255 sum &= tData[ i ] & (cl_int)0x80000000L;
256 return (sum != 0);
257 }
258 case kLong:
259 {
260 cl_long sum = 0x8000000000000000LL;
261 cl_long *tData = (cl_long *)inData;
262 for( i = 0; i < vecSize; i++ )
263 sum &= tData[ i ] & 0x8000000000000000LL;
264 return (sum != 0);
265 }
266 default:
267 return 0;
268 }
269 }
270
test_relational_all(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)271 int test_relational_all(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
272 {
273 ExplicitType vecType[] = { kChar, kShort, kInt, kLong };
274 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
275 unsigned int index, typeIndex;
276 int retVal = 0;
277 RandomSeed seed(gRandomSeed );
278
279
280 for( typeIndex = 0; typeIndex < 4; typeIndex++ )
281 {
282 if (vecType[typeIndex] == kLong && !gHasLong)
283 continue;
284
285 for( index = 0; vecSizes[ index ] != 0; index++ )
286 {
287 // Test!
288 if( test_any_all_kernel(context, queue, "all", vecType[ typeIndex ], vecSizes[ index ], allVerifyFn, seed ) != 0 )
289 {
290 log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] );
291 retVal = -1;
292 }
293 }
294 }
295
296 return retVal;
297 }
298
299 // clang-format off
300
301 const char *selectTestKernelPattern =
302 "%s\n" // optional pragma
303 "%s\n" // optional pragma
304 "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n"
305 "{\n"
306 " int tid = get_global_id(0);\n"
307 " destValues[tid] = %s( sourceA[tid], sourceB[tid], sourceC[tid] );\n"
308 "\n"
309 "}\n";
310
311
312 const char *selectTestKernelPatternVload =
313 "%s\n" // optional pragma
314 "%s\n" // optional pragma
315 "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n"
316 "{\n"
317 " int tid = get_global_id(0);\n"
318 " %s%s tmp = %s( vload3(tid, (__global %s *)sourceA), vload3(tid, (__global %s *)sourceB), vload3(tid, (__global %s *)sourceC) );\n"
319 " vstore3(tmp, tid, (__global %s *)destValues);\n"
320 "\n"
321 "}\n";
322
323 // clang-format on
324
325 typedef void (*selectVerifyFn)( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData );
326
test_select_kernel(cl_context context,cl_command_queue queue,const char * fnName,ExplicitType vecType,unsigned int vecSize,ExplicitType testVecType,selectVerifyFn verifyFn,MTdata d)327 int test_select_kernel(cl_context context, cl_command_queue queue, const char *fnName,
328 ExplicitType vecType, unsigned int vecSize, ExplicitType testVecType, selectVerifyFn verifyFn, MTdata d )
329 {
330 clProgramWrapper program;
331 clKernelWrapper kernel;
332 clMemWrapper streams[4];
333 cl_long inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ], inDataC[ TEST_SIZE * 16 ];
334 cl_long outData[TEST_SIZE * 16], expected[16];
335 int error, i;
336 size_t threads[1], localThreads[1];
337 char kernelSource[10240];
338 char *programPtr;
339 char sizeName[4], outSizeName[4];
340 unsigned int outVecSize;
341
342
343 /* Create the source */
344 if( vecSize == 1 )
345 sizeName[ 0 ] = 0;
346 else
347 sprintf( sizeName, "%d", vecSize );
348
349 outVecSize = vecSize;
350
351 if( outVecSize == 1 )
352 outSizeName[ 0 ] = 0;
353 else
354 sprintf( outSizeName, "%d", outVecSize );
355
356 if(DENSE_PACK_VECS && vecSize == 3) {
357 // anyAllTestKernelPatternVload
358 sprintf(kernelSource, selectTestKernelPatternVload,
359 (vecType == kDouble || testVecType == kDouble)
360 ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
361 : "",
362 (vecType == kHalf || testVecType == kHalf)
363 ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable"
364 : "",
365 get_explicit_type_name(vecType), sizeName,
366 get_explicit_type_name(vecType), sizeName,
367 get_explicit_type_name(testVecType), sizeName,
368 get_explicit_type_name(vecType), outSizeName,
369 get_explicit_type_name(vecType), sizeName, fnName,
370 get_explicit_type_name(vecType),
371 get_explicit_type_name(vecType),
372 get_explicit_type_name(vecType),
373 get_explicit_type_name(testVecType));
374 } else {
375 sprintf(kernelSource, selectTestKernelPattern,
376 (vecType == kDouble || testVecType == kDouble)
377 ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
378 : "",
379 (vecType == kHalf || testVecType == kHalf)
380 ? "#pragma OPENCL EXTENSION cl_khr_fp16 : enable"
381 : "",
382 get_explicit_type_name(vecType), sizeName,
383 get_explicit_type_name(vecType), sizeName,
384 get_explicit_type_name(testVecType), sizeName,
385 get_explicit_type_name(vecType), outSizeName, fnName);
386 }
387
388 /* Create kernels */
389 programPtr = kernelSource;
390 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
391 {
392 return -1;
393 }
394
395 /* Generate some streams */
396 generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataA );
397 generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataB );
398 generate_random_data( testVecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataC );
399
400 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
401 get_explicit_type_size(vecType)
402 * g_vector_aligns[vecSize] * TEST_SIZE,
403 &inDataA, &error);
404 if( streams[0] == NULL )
405 {
406 print_error( error, "Creating input array A failed!\n");
407 return -1;
408 }
409 streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
410 get_explicit_type_size(vecType)
411 * g_vector_aligns[vecSize] * TEST_SIZE,
412 &inDataB, &error);
413 if( streams[1] == NULL )
414 {
415 print_error( error, "Creating input array A failed!\n");
416 return -1;
417 }
418 streams[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
419 get_explicit_type_size(testVecType)
420 * g_vector_aligns[vecSize] * TEST_SIZE,
421 &inDataC, &error);
422 if( streams[2] == NULL )
423 {
424 print_error( error, "Creating input array A failed!\n");
425 return -1;
426 }
427 streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, get_explicit_type_size( vecType ) * g_vector_aligns[outVecSize] * TEST_SIZE, NULL, &error);
428 if( streams[3] == NULL )
429 {
430 print_error( error, "Creating output array failed!\n");
431 return -1;
432 }
433
434 /* Assign streams and execute */
435 error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
436 test_error( error, "Unable to set indexed kernel arguments" );
437 error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
438 test_error( error, "Unable to set indexed kernel arguments" );
439 error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] );
440 test_error( error, "Unable to set indexed kernel arguments" );
441 error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] );
442 test_error( error, "Unable to set indexed kernel arguments" );
443
444 /* Run the kernel */
445 threads[0] = TEST_SIZE;
446
447 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
448 test_error( error, "Unable to get work group size to use" );
449
450 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
451 test_error( error, "Unable to execute test kernel" );
452
453 /* Now get the results */
454 error = clEnqueueReadBuffer( queue, streams[3], true, 0, get_explicit_type_size( vecType ) * TEST_SIZE * g_vector_aligns[outVecSize], outData, 0, NULL, NULL );
455 test_error( error, "Unable to read output array!" );
456
457 /* And verify! */
458 for( i = 0; i < (int)(TEST_SIZE * g_vector_aligns[vecSize]); i++ )
459 {
460 if(i%g_vector_aligns[vecSize] >= (int) vecSize) {
461 continue;
462 }
463 verifyFn( vecType, testVecType, vecSize, (char *)inDataA + i * get_explicit_type_size( vecType ),
464 (char *)inDataB + i * get_explicit_type_size( vecType ),
465 (char *)inDataC + i * get_explicit_type_size( testVecType ),
466 expected);
467
468 char *outPtr = (char *)outData;
469 outPtr += ( i / g_vector_aligns[vecSize] ) * get_explicit_type_size( vecType ) * g_vector_aligns[outVecSize];
470 outPtr += ( i % g_vector_aligns[vecSize] ) * get_explicit_type_size( vecType );
471 if( memcmp( expected, outPtr, get_explicit_type_size( vecType ) ) != 0 )
472 {
473 log_error( "ERROR: Data sample %d:%d does not validate! Expected (0x%08x), got (0x%08x) from (0x%08x) and (0x%08x) with test (0x%08x)\n",
474 i / g_vector_aligns[vecSize],
475 i % g_vector_aligns[vecSize],
476 *( (int *)expected ),
477 *( (int *)( (char *)outData +
478 i * get_explicit_type_size( vecType
479 ) ) ),
480 *( (int *)( (char *)inDataA +
481 i * get_explicit_type_size( vecType
482 ) ) ),
483 *( (int *)( (char *)inDataB +
484 i * get_explicit_type_size( vecType
485 ) ) ),
486 *( (int *)( (char *)inDataC +
487 i*get_explicit_type_size( testVecType
488 ) ) ) );
489 int j;
490 log_error( "inA: " );
491 unsigned char *a = (unsigned char *)( (char *)inDataA + i * get_explicit_type_size( vecType ) );
492 unsigned char *b = (unsigned char *)( (char *)inDataB + i * get_explicit_type_size( vecType ) );
493 unsigned char *c = (unsigned char *)( (char *)inDataC + i * get_explicit_type_size( testVecType ) );
494 unsigned char *e = (unsigned char *)( expected );
495 unsigned char *g = (unsigned char *)( (char *)outData + i * get_explicit_type_size( vecType ) );
496 for( j = 0; j < 16; j++ )
497 log_error( "0x%02x ", a[ j ] );
498 log_error( "\ninB: " );
499 for( j = 0; j < 16; j++ )
500 log_error( "0x%02x ", b[ j ] );
501 log_error( "\ninC: " );
502 for( j = 0; j < 16; j++ )
503 log_error( "0x%02x ", c[ j ] );
504 log_error( "\nexp: " );
505 for( j = 0; j < 16; j++ )
506 log_error( "0x%02x ", e[ j ] );
507 log_error( "\ngot: " );
508 for( j = 0; j < 16; j++ )
509 log_error( "0x%02x ", g[ j ] );
510 return -1;
511 }
512 }
513
514 return 0;
515 }
516
bitselect_verify_fn(ExplicitType vecType,ExplicitType testVecType,unsigned int vecSize,void * inDataA,void * inDataB,void * inDataTest,void * outData)517 void bitselect_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData )
518 {
519 char *inA = (char *)inDataA, *inB = (char *)inDataB, *inT = (char *)inDataTest, *out = (char *)outData;
520 size_t i, numBytes = get_explicit_type_size( vecType );
521
522 // Type is meaningless, this is all bitwise!
523 for( i = 0; i < numBytes; i++ )
524 {
525 out[ i ] = ( inA[ i ] & ~inT[ i ] ) | ( inB[ i ] & inT[ i ] );
526 }
527 }
528
test_relational_bitselect(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)529 int test_relational_bitselect(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
530 {
531 constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort,
532 kInt, kUInt, kLong, kULong,
533 kHalf, kFloat, kDouble };
534 constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType);
535 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
536 unsigned int index, typeIndex;
537 int retVal = 0;
538 RandomSeed seed( gRandomSeed );
539
540
541 for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++)
542 {
543 if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong)
544 continue;
545
546 if (vecType[typeIndex] == kDouble)
547 {
548 if(!is_extension_available(device, "cl_khr_fp64"))
549 {
550 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
551 continue;
552 }
553 else
554 log_info("Testing doubles.\n");
555 }
556
557 if (vecType[typeIndex] == kHalf)
558 {
559 if (!is_extension_available(device, "cl_khr_fp16"))
560 {
561 log_info("Extension cl_khr_fp16 not supported; skipping half "
562 "tests.\n");
563 continue;
564 }
565 else
566 log_info("Testing halfs.\n");
567 }
568
569 for( index = 0; vecSizes[ index ] != 0; index++ )
570 {
571 // Test!
572 if( test_select_kernel(context, queue, "bitselect", vecType[ typeIndex ], vecSizes[ index ], vecType[typeIndex], bitselect_verify_fn, seed ) != 0 )
573 {
574 log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] );
575 retVal = -1;
576 }
577 }
578 }
579
580 return retVal;
581 }
582
select_signed_verify_fn(ExplicitType vecType,ExplicitType testVecType,unsigned int vecSize,void * inDataA,void * inDataB,void * inDataTest,void * outData)583 void select_signed_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData )
584 {
585 bool yep = false;
586 if (vecSize == 1) {
587 switch( testVecType )
588 {
589 case kChar:
590 yep = *( (char *)inDataTest ) ? true : false;
591 break;
592 case kShort:
593 yep = *( (short *)inDataTest ) ? true : false;
594 break;
595 case kInt:
596 yep = *( (int *)inDataTest ) ? true : false;
597 break;
598 case kLong:
599 yep = *( (cl_long *)inDataTest ) ? true : false;
600 break;
601 default:
602 // Should never get here
603 return;
604 }
605 }
606 else {
607 switch( testVecType )
608 {
609 case kChar:
610 yep = *( (char *)inDataTest ) & 0x80 ? true : false;
611 break;
612 case kShort:
613 yep = *( (short *)inDataTest ) & 0x8000 ? true : false;
614 break;
615 case kInt:
616 yep = *( (int *)inDataTest ) & 0x80000000L ? true : false;
617 break;
618 case kLong:
619 yep = *( (cl_long *)inDataTest ) & 0x8000000000000000LL ? true : false;
620 break;
621 default:
622 // Should never get here
623 return;
624 }
625 }
626 memcpy( outData, ( yep ) ? inDataB : inDataA, get_explicit_type_size( vecType ) );
627 }
628
test_relational_select_signed(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)629 int test_relational_select_signed(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
630 {
631 constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort,
632 kInt, kUInt, kLong, kULong,
633 kHalf, kFloat, kDouble };
634 constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType);
635
636 ExplicitType testVecType[] = { kChar, kShort, kInt, kLong, kNumExplicitTypes };
637 unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
638 unsigned int index, typeIndex, testTypeIndex;
639 int retVal = 0;
640 RandomSeed seed( gRandomSeed );
641
642 for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++)
643 {
644 if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong)
645 continue;
646
647 if (vecType[typeIndex] == kDouble) {
648 if(!is_extension_available(device, "cl_khr_fp64")) {
649 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
650 continue;
651 } else {
652 log_info("Testing doubles.\n");
653 }
654 }
655 if (vecType[typeIndex] == kHalf)
656 {
657 if (!is_extension_available(device, "cl_khr_fp16"))
658 {
659 log_info("Extension cl_khr_fp16 not supported; skipping half "
660 "tests.\n");
661 continue;
662 }
663 else
664 {
665 log_info("Testing halfs.\n");
666 }
667 }
668 for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ )
669 {
670 if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] )
671 continue;
672
673 for( index = 0; vecSizes[ index ] != 0; index++ )
674 {
675 // Test!
676 if( test_select_kernel(context, queue, "select", vecType[ typeIndex ], vecSizes[ index ], testVecType[ testTypeIndex ], select_signed_verify_fn, seed ) != 0 )
677 {
678 log_error( " Vector %s%d, test vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ],
679 get_explicit_type_name( testVecType[ testTypeIndex ] ), vecSizes[ index ] );
680 retVal = -1;
681 }
682 }
683 }
684 }
685
686 return retVal;
687 }
688
select_unsigned_verify_fn(ExplicitType vecType,ExplicitType testVecType,unsigned int vecSize,void * inDataA,void * inDataB,void * inDataTest,void * outData)689 void select_unsigned_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData )
690 {
691 bool yep = false;
692 if (vecSize == 1) {
693 switch( testVecType )
694 {
695 case kUChar:
696 yep = *( (unsigned char *)inDataTest ) ? true : false;
697 break;
698 case kUShort:
699 yep = *( (unsigned short *)inDataTest ) ? true : false;
700 break;
701 case kUInt:
702 yep = *( (unsigned int *)inDataTest ) ? true : false;
703 break;
704 case kULong:
705 yep = *( (cl_ulong *)inDataTest ) ? true : false;
706 break;
707 default:
708 // Should never get here
709 return;
710 }
711 }
712 else {
713 switch( testVecType )
714 {
715 case kUChar:
716 yep = *( (unsigned char *)inDataTest ) & 0x80 ? true : false;
717 break;
718 case kUShort:
719 yep = *( (unsigned short *)inDataTest ) & 0x8000 ? true : false;
720 break;
721 case kUInt:
722 yep = *( (unsigned int *)inDataTest ) & 0x80000000L ? true : false;
723 break;
724 case kULong:
725 yep = *( (cl_ulong *)inDataTest ) & 0x8000000000000000LL ? true : false;
726 break;
727 default:
728 // Should never get here
729 return;
730 }
731 }
732 memcpy( outData, ( yep ) ? inDataB : inDataA, get_explicit_type_size( vecType ) );
733 }
734
test_relational_select_unsigned(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)735 int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
736 {
737 constexpr ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort,
738 kInt, kUInt, kLong, kULong,
739 kHalf, kFloat, kDouble };
740 constexpr auto vecTypeSize = sizeof(vecType) / sizeof(ExplicitType);
741
742 ExplicitType testVecType[] = { kUChar, kUShort, kUInt, kULong, kNumExplicitTypes };
743 unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
744 unsigned int index, typeIndex, testTypeIndex;
745 int retVal = 0;
746 RandomSeed seed(gRandomSeed);
747
748
749 for (typeIndex = 0; typeIndex < vecTypeSize; typeIndex++)
750 {
751 if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong)
752 continue;
753
754 if (vecType[typeIndex] == kDouble) {
755 if(!is_extension_available(device, "cl_khr_fp64")) {
756 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
757 continue;
758 } else {
759 log_info("Testing doubles.\n");
760 }
761 }
762 if (vecType[typeIndex] == kHalf)
763 {
764 if (!is_extension_available(device, "cl_khr_fp16"))
765 {
766 log_info("Extension cl_khr_fp16 not supported; skipping half "
767 "tests.\n");
768 continue;
769 }
770 else
771 {
772 log_info("Testing halfs.\n");
773 }
774 }
775 for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ )
776 {
777 if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] )
778 continue;
779
780 for( index = 0; vecSizes[ index ] != 0; index++ )
781 {
782 // Test!
783 if( test_select_kernel(context, queue, "select", vecType[ typeIndex ], vecSizes[ index ], testVecType[ testTypeIndex ], select_unsigned_verify_fn, seed ) != 0 )
784 {
785 log_error( " Vector %s%d, test vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ],
786 get_explicit_type_name( testVecType[ testTypeIndex ] ), vecSizes[ index ] );
787 retVal = -1;
788 }
789 }
790 }
791 }
792
793 return retVal;
794 }
795