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 
20 #define TEST_SIZE 512
21 
22 const char *equivTestKernelPattern_float =
23 "__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
24 "{\n"
25 "    int  tid = get_global_id(0);\n"
26 "    destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
27 "    destValuesB[tid] = sourceA[tid] %s sourceB[tid];\n"
28 "\n"
29 "}\n";
30 
31 const char *equivTestKernelPatternLessGreater_float =
32 "__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
33 "{\n"
34 "    int  tid = get_global_id(0);\n"
35 "    destValues[tid] = %s( sourceA[tid], sourceB[tid] );\n"
36 "    destValuesB[tid] = (sourceA[tid] < sourceB[tid]) | (sourceA[tid] > sourceB[tid]);\n"
37 "\n"
38 "}\n";
39 
40 
41 const char *equivTestKernelPattern_float3 =
42 "__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
43 "{\n"
44 "    int  tid = get_global_id(0);\n"
45 "    float3 sampA = vload3(tid, (__global float *)sourceA);\n"
46 "    float3 sampB = vload3(tid, (__global float *)sourceB);\n"
47 "    vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n"
48 "    vstore3(( sampA %s sampB ), tid, (__global int *)destValuesB);\n"
49 "\n"
50 "}\n";
51 
52 const char *equivTestKernelPatternLessGreater_float3 =
53 "__kernel void sample_test(__global float%s *sourceA, __global float%s *sourceB, __global int%s *destValues, __global int%s *destValuesB)\n"
54 "{\n"
55 "    int  tid = get_global_id(0);\n"
56 "    float3 sampA = vload3(tid, (__global float *)sourceA);\n"
57 "    float3 sampB = vload3(tid, (__global float *)sourceB);\n"
58 "    vstore3(%s( sampA, sampB ), tid, (__global int *)destValues);\n"
59 "    vstore3(( sampA < sampB ) | (sampA > sampB), tid, (__global int *)destValuesB);\n"
60 "\n"
61 "}\n";
62 
63 typedef bool (*equivVerifyFn)( float inDataA, float inDataB );
64 extern int gInfNanSupport;
65 
IsFloatInfinity(float x)66 int IsFloatInfinity(float x)
67 {
68     return isinf(x);
69 }
70 
IsFloatNaN(float x)71 int IsFloatNaN(float x)
72 {
73     return isnan(x);
74 }
75 
verify_equiv_values_float(unsigned int vecSize,float * inDataA,float * inDataB,int * outData,equivVerifyFn verifyFn)76 void verify_equiv_values_float( unsigned int vecSize, float *inDataA, float *inDataB, int *outData, equivVerifyFn verifyFn )
77 {
78     unsigned int i;
79     int trueResult;
80     bool result;
81 
82     trueResult = ( vecSize == 1 ) ? 1 : -1;
83     for( i = 0; i < vecSize; i++ )
84     {
85         result = verifyFn( inDataA[ i ], inDataB[ i ] );
86         outData[ i ] = result ? trueResult : 0;
87     }
88 }
89 
generate_equiv_test_data_float(float * outData,unsigned int vecSize,bool alpha,MTdata d)90 void generate_equiv_test_data_float( float *outData, unsigned int vecSize, bool alpha, MTdata d )
91 {
92     unsigned int i;
93 
94     generate_random_data( kFloat, vecSize * TEST_SIZE, d, outData );
95 
96     // Fill the first few vectors with NAN in each vector element (or the second set if we're alpha, so we can test either case)
97     if( alpha )
98         outData += vecSize * vecSize;
99     for( i = 0; i < vecSize; i++ )
100     {
101         outData[ 0 ] = NAN;
102         outData += vecSize + 1;
103     }
104     // Make sure the third set is filled regardless, to test the case where both have NANs
105     if( !alpha )
106         outData += vecSize * vecSize;
107     for( i = 0; i < vecSize; i++ )
108     {
109         outData[ 0 ] = NAN;
110         outData += vecSize + 1;
111     }
112 }
113 
test_equiv_kernel_float(cl_context context,cl_command_queue queue,const char * fnName,const char * opName,unsigned int vecSize,equivVerifyFn verifyFn,MTdata d)114 int test_equiv_kernel_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName,
115                        unsigned int vecSize, equivVerifyFn verifyFn, MTdata d )
116 {
117     clProgramWrapper program;
118     clKernelWrapper kernel;
119     clMemWrapper streams[4];
120     float inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ];
121     int outData[TEST_SIZE * 16], expected[16];
122     int error, i, j;
123     size_t threads[1], localThreads[1];
124     char kernelSource[10240];
125     char *programPtr;
126     char sizeName[4];
127 
128 
129     /* Create the source */
130     if( vecSize == 1 )
131         sizeName[ 0 ] = 0;
132     else
133         sprintf( sizeName, "%d", vecSize );
134 
135 
136     if(DENSE_PACK_VECS && vecSize == 3) {
137   if (strcmp(fnName, "islessgreater")) {
138             sprintf( kernelSource, equivTestKernelPattern_float3, sizeName, sizeName, sizeName, sizeName, fnName, opName );
139         } else {
140             sprintf( kernelSource, equivTestKernelPatternLessGreater_float3, sizeName, sizeName, sizeName, sizeName, fnName );
141         }
142     } else {
143         if (strcmp(fnName, "islessgreater")) {
144           sprintf( kernelSource, equivTestKernelPattern_float, sizeName, sizeName, sizeName, sizeName, fnName, opName );
145   } else {
146     sprintf( kernelSource, equivTestKernelPatternLessGreater_float, sizeName, sizeName, sizeName, sizeName, fnName );
147   }
148     }
149 
150     /* Create kernels */
151     programPtr = kernelSource;
152     if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
153     {
154         return -1;
155     }
156 
157     /* Generate some streams */
158     generate_equiv_test_data_float( inDataA, vecSize, true, d );
159     generate_equiv_test_data_float( inDataB, vecSize, false, d );
160 
161     streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
162                                 sizeof(cl_float) * vecSize * TEST_SIZE,
163                                 &inDataA, &error);
164     if( streams[0] == NULL )
165     {
166         print_error( error, "Creating input array A failed!\n");
167         return -1;
168     }
169     streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
170                                 sizeof(cl_float) * vecSize * TEST_SIZE,
171                                 &inDataB, &error);
172     if( streams[1] == NULL )
173     {
174         print_error( error, "Creating input array A failed!\n");
175         return -1;
176     }
177     streams[2] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error);
178     if( streams[2] == NULL )
179     {
180         print_error( error, "Creating output array failed!\n");
181         return -1;
182     }
183   streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( cl_int ) * vecSize * TEST_SIZE, NULL, &error);
184     if( streams[3] == NULL )
185     {
186         print_error( error, "Creating output array failed!\n");
187         return -1;
188     }
189 
190 
191     /* Assign streams and execute */
192     error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
193     test_error( error, "Unable to set indexed kernel arguments" );
194     error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
195     test_error( error, "Unable to set indexed kernel arguments" );
196     error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] );
197     test_error( error, "Unable to set indexed kernel arguments" );
198   error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] );
199     test_error( error, "Unable to set indexed kernel arguments" );
200 
201 
202     /* Run the kernel */
203     threads[0] = TEST_SIZE;
204 
205     error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
206     test_error( error, "Unable to get work group size to use" );
207 
208     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
209     test_error( error, "Unable to execute test kernel" );
210 
211   /* Now get the results */
212     error = clEnqueueReadBuffer( queue, streams[2], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL );
213     test_error( error, "Unable to read output array!" );
214 
215   /* And verify! */
216   for( i = 0; i < TEST_SIZE; i++ )
217   {
218         verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn);
219 
220         for( j = 0; j < (int)vecSize; j++ )
221         {
222             if( expected[ j ] != outData[ i * vecSize + j ] )
223             {
224                 log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n",
225                   i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] );
226                 return -1;
227             }
228         }
229   }
230 
231   /* Now get the results */
232     error = clEnqueueReadBuffer( queue, streams[3], true, 0, sizeof( int ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL );
233     test_error( error, "Unable to read output array!" );
234 
235   /* And verify! */
236     int fail = 0;
237     for( i = 0; i < TEST_SIZE; i++ )
238     {
239         verify_equiv_values_float( vecSize, &inDataA[ i * vecSize ], &inDataB[ i * vecSize ], expected, verifyFn);
240 
241         for( j = 0; j < (int)vecSize; j++ )
242         {
243             if( expected[ j ] != outData[ i * vecSize + j ] )
244             {
245                 if (gInfNanSupport == 0)
246                 {
247                     if (IsFloatNaN(inDataA[i*vecSize + j]) || IsFloatNaN (inDataB[i*vecSize + j]))
248                     {
249                         fail = 0;
250                     }
251                     else
252                         fail = 1;
253                 }
254                 if (fail)
255                 {
256                     log_error( "ERROR: Data sample %d:%d at size %d does not validate! Expected %d, got %d, source %f,%f\n",
257                       i, j, vecSize, expected[ j ], outData[ i * vecSize + j ], inDataA[i*vecSize + j], inDataB[i*vecSize + j] );
258                     return -1;
259                 }
260             }
261         }
262   }
263 
264   return 0;
265 }
266 
test_equiv_kernel_set_float(cl_context context,cl_command_queue queue,const char * fnName,const char * opName,equivVerifyFn verifyFn,MTdata d)267 int test_equiv_kernel_set_float(cl_context context, cl_command_queue queue, const char *fnName, const char *opName, equivVerifyFn verifyFn, MTdata d )
268 {
269     unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
270     unsigned int index;
271     int retVal = 0;
272 
273     for( index = 0; vecSizes[ index ] != 0; index++ )
274     {
275         // Test!
276         if( test_equiv_kernel_float(context, queue, fnName, opName, vecSizes[ index ], verifyFn, d ) != 0 )
277         {
278             log_error( "   Vector float%d FAILED\n", vecSizes[ index ] );
279             retVal = -1;
280         }
281     }
282 
283     return retVal;
284 }
285 
isequal_verify_fn_float(float valueA,float valueB)286 bool isequal_verify_fn_float( float valueA, float valueB )
287 {
288     return valueA == valueB;
289 }
290 
test_relational_isequal_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)291 int test_relational_isequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
292 {
293     RandomSeed seed( gRandomSeed );
294     return test_equiv_kernel_set_float( context, queue, "isequal", "==", isequal_verify_fn_float, seed );
295 }
296 
isnotequal_verify_fn_float(float valueA,float valueB)297 bool isnotequal_verify_fn_float( float valueA, float valueB )
298 {
299     return valueA != valueB;
300 }
301 
test_relational_isnotequal_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)302 int test_relational_isnotequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
303 {
304     RandomSeed seed( gRandomSeed );
305     return test_equiv_kernel_set_float( context, queue, "isnotequal", "!=", isnotequal_verify_fn_float, seed );
306 }
307 
isgreater_verify_fn_float(float valueA,float valueB)308 bool isgreater_verify_fn_float( float valueA, float valueB )
309 {
310     return valueA > valueB;
311 }
312 
test_relational_isgreater_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)313 int test_relational_isgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
314 {
315     RandomSeed seed( gRandomSeed );
316     return test_equiv_kernel_set_float( context, queue, "isgreater", ">", isgreater_verify_fn_float, seed );
317 }
318 
isgreaterequal_verify_fn_float(float valueA,float valueB)319 bool isgreaterequal_verify_fn_float( float valueA, float valueB )
320 {
321     return valueA >= valueB;
322 }
323 
test_relational_isgreaterequal_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)324 int test_relational_isgreaterequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
325 {
326     RandomSeed seed( gRandomSeed );
327     return test_equiv_kernel_set_float( context, queue, "isgreaterequal", ">=", isgreaterequal_verify_fn_float, seed );
328 }
329 
isless_verify_fn_float(float valueA,float valueB)330 bool isless_verify_fn_float( float valueA, float valueB )
331 {
332     return valueA < valueB;
333 }
334 
test_relational_isless_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)335 int test_relational_isless_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
336 {
337     RandomSeed seed( gRandomSeed );
338     return test_equiv_kernel_set_float( context, queue, "isless", "<", isless_verify_fn_float, seed );
339 }
340 
islessequal_verify_fn_float(float valueA,float valueB)341 bool islessequal_verify_fn_float( float valueA, float valueB )
342 {
343     return valueA <= valueB;
344 }
345 
test_relational_islessequal_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)346 int test_relational_islessequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
347 {
348     RandomSeed seed( gRandomSeed );
349     return test_equiv_kernel_set_float( context, queue, "islessequal", "<=", islessequal_verify_fn_float, seed );
350 }
351 
islessgreater_verify_fn_float(float valueA,float valueB)352 bool islessgreater_verify_fn_float( float valueA, float valueB )
353 {
354     return ( valueA < valueB ) || ( valueA > valueB );
355 }
356 
test_relational_islessgreater_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)357 int test_relational_islessgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
358 {
359     RandomSeed seed( gRandomSeed );
360     return test_equiv_kernel_set_float( context, queue, "islessgreater", "<>", islessgreater_verify_fn_float, seed );
361 }
362 
363 
364