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
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24
25 #include "procs.h"
26 #include "harness/conversions.h"
27 #include "harness/typeWrappers.h"
28
29 #define DECLARE_S2V_IDENT_KERNEL(srctype,dsttype,size) \
30 "__kernel void test_conversion(__global " srctype " *sourceValues, __global " dsttype #size " *destValues )\n" \
31 "{\n" \
32 " int tid = get_global_id(0);\n" \
33 " " srctype " src = sourceValues[tid];\n" \
34 "\n" \
35 " destValues[tid] = (" dsttype #size ")src;\n" \
36 "\n" \
37 "}\n"
38
39 #define DECLARE_S2V_IDENT_KERNELS(srctype,dsttype) \
40 { \
41 DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,2), \
42 DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,4), \
43 DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,8), \
44 DECLARE_S2V_IDENT_KERNEL(srctype,#dsttype,16) \
45 }
46
47 #define DECLARE_EMPTY { NULL, NULL, NULL, NULL, NULL }
48
49 /* Note: the next four arrays all must match in order and size to the ExplicitTypes enum in conversions.h!!! */
50
51 #define DECLARE_S2V_IDENT_KERNELS_SET(srctype) \
52 { \
53 DECLARE_S2V_IDENT_KERNELS(#srctype,bool), \
54 DECLARE_S2V_IDENT_KERNELS(#srctype,char), \
55 DECLARE_S2V_IDENT_KERNELS(#srctype,uchar), \
56 DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned char), \
57 DECLARE_S2V_IDENT_KERNELS(#srctype,short), \
58 DECLARE_S2V_IDENT_KERNELS(#srctype,ushort), \
59 DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned short), \
60 DECLARE_S2V_IDENT_KERNELS(#srctype,int), \
61 DECLARE_S2V_IDENT_KERNELS(#srctype,uint), \
62 DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned int), \
63 DECLARE_S2V_IDENT_KERNELS(#srctype,long), \
64 DECLARE_S2V_IDENT_KERNELS(#srctype,ulong), \
65 DECLARE_S2V_IDENT_KERNELS(#srctype,unsigned long), \
66 DECLARE_S2V_IDENT_KERNELS(#srctype,float), \
67 DECLARE_EMPTY \
68 }
69
70 #define DECLARE_EMPTY_SET \
71 { \
72 DECLARE_EMPTY, \
73 DECLARE_EMPTY, \
74 DECLARE_EMPTY, \
75 DECLARE_EMPTY, \
76 DECLARE_EMPTY, \
77 DECLARE_EMPTY, \
78 DECLARE_EMPTY, \
79 DECLARE_EMPTY, \
80 DECLARE_EMPTY, \
81 DECLARE_EMPTY, \
82 DECLARE_EMPTY, \
83 DECLARE_EMPTY, \
84 DECLARE_EMPTY, \
85 DECLARE_EMPTY, \
86 DECLARE_EMPTY \
87 }
88
89
90 /* The overall array */
91 const char * kernel_explicit_s2v_set[kNumExplicitTypes][kNumExplicitTypes][5] = {
92 DECLARE_S2V_IDENT_KERNELS_SET(bool),
93 DECLARE_S2V_IDENT_KERNELS_SET(char),
94 DECLARE_S2V_IDENT_KERNELS_SET(uchar),
95 DECLARE_S2V_IDENT_KERNELS_SET(unsigned char),
96 DECLARE_S2V_IDENT_KERNELS_SET(short),
97 DECLARE_S2V_IDENT_KERNELS_SET(ushort),
98 DECLARE_S2V_IDENT_KERNELS_SET(unsigned short),
99 DECLARE_S2V_IDENT_KERNELS_SET(int),
100 DECLARE_S2V_IDENT_KERNELS_SET(uint),
101 DECLARE_S2V_IDENT_KERNELS_SET(unsigned int),
102 DECLARE_S2V_IDENT_KERNELS_SET(long),
103 DECLARE_S2V_IDENT_KERNELS_SET(ulong),
104 DECLARE_S2V_IDENT_KERNELS_SET(unsigned long),
105 DECLARE_S2V_IDENT_KERNELS_SET(float),
106 DECLARE_EMPTY_SET
107 };
108
test_explicit_s2v_function(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * programSrc,ExplicitType srcType,unsigned int count,ExplicitType destType,unsigned int vecSize,void * inputData)109 int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char *programSrc,
110 ExplicitType srcType, unsigned int count, ExplicitType destType, unsigned int vecSize, void *inputData )
111 {
112 clProgramWrapper program;
113 clKernelWrapper kernel;
114 int error;
115 clMemWrapper streams[2];
116 void *outData;
117 unsigned char convertedData[ 8 ]; /* Max type size is 8 bytes */
118 size_t threadSize[3], groupSize[3];
119 unsigned int i, s;
120 unsigned char *inPtr, *outPtr;
121 size_t paramSize, destTypeSize;
122
123 const char* finalProgramSrc[2] = {
124 "", // optional pragma
125 programSrc
126 };
127
128 if (srcType == kDouble || destType == kDouble) {
129 finalProgramSrc[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
130 }
131
132
133 if( programSrc == NULL )
134 return 0;
135
136 paramSize = get_explicit_type_size( srcType );
137 destTypeSize = get_explicit_type_size( destType );
138
139 size_t destStride = destTypeSize * vecSize;
140
141 outData = malloc( destStride * count );
142
143 if( create_single_kernel_helper( context, &program, &kernel, 2, finalProgramSrc, "test_conversion" ) )
144 {
145 log_info( "****** %s%s *******\n", finalProgramSrc[0], finalProgramSrc[1] );
146 return -1;
147 }
148
149 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
150 paramSize * count, inputData, &error);
151 test_error( error, "clCreateBuffer failed");
152 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, destStride * count,
153 NULL, &error);
154 test_error( error, "clCreateBuffer failed");
155
156 /* Set the arguments */
157 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0] );
158 test_error( error, "Unable to set indexed kernel arguments" );
159 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1] );
160 test_error( error, "Unable to set indexed kernel arguments" );
161
162 /* Run the kernel */
163 threadSize[0] = count;
164
165 error = get_max_common_work_group_size( context, kernel, threadSize[0], &groupSize[0] );
166 test_error( error, "Unable to get work group size to use" );
167
168 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threadSize, groupSize, 0, NULL, NULL );
169 test_error( error, "Unable to execute test kernel" );
170
171 /* Now verify the results. Each value should have been duplicated four times, and we should be able to just
172 do a memcpy instead of relying on the actual type of data */
173 error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, destStride * count, outData, 0, NULL, NULL );
174 test_error( error, "Unable to read output values!" );
175
176 inPtr = (unsigned char *)inputData;
177 outPtr = (unsigned char *)outData;
178
179 for( i = 0; i < count; i++ )
180 {
181 /* Convert the input data element to our output data type to compare against */
182 convert_explicit_value( (void *)inPtr, (void *)convertedData, srcType, false, kDefaultRoundingType, destType );
183
184 /* Now compare every element of the vector */
185 for( s = 0; s < vecSize; s++ )
186 {
187 if( memcmp( convertedData, outPtr + destTypeSize * s, destTypeSize ) != 0 )
188 {
189 unsigned int *p = (unsigned int *)outPtr;
190 log_error( "ERROR: Output value %d:%d does not validate for size %d:%d!\n", i, s, vecSize, (int)destTypeSize );
191 log_error( " Input: 0x%0*x\n", (int)( paramSize * 2 ), *(unsigned int *)inPtr & ( 0xffffffff >> ( 32 - paramSize * 8 ) ) );
192 log_error( " Actual: 0x%08x 0x%08x 0x%08x 0x%08x\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
193 return -1;
194 }
195 }
196 inPtr += paramSize;
197 outPtr += destStride;
198 }
199
200 free( outData );
201
202 return 0;
203 }
204
test_explicit_s2v_function_set(cl_device_id deviceID,cl_context context,cl_command_queue queue,ExplicitType srcType,unsigned int count,void * inputData)205 int test_explicit_s2v_function_set(cl_device_id deviceID, cl_context context, cl_command_queue queue, ExplicitType srcType,
206 unsigned int count, void *inputData )
207 {
208 unsigned int sizes[] = { 2, 4, 8, 16, 0 };
209 int i, dstType, failed = 0;
210
211
212 for( dstType = kBool; dstType < kNumExplicitTypes; dstType++ )
213 {
214 if( dstType == kDouble && !is_extension_available( deviceID, "cl_khr_fp64" ) )
215 continue;
216
217 if (( dstType == kLong || dstType == kULong ) && !gHasLong )
218 continue;
219
220 for( i = 0; sizes[i] != 0; i++ )
221 {
222 if( dstType != srcType )
223 continue;
224 if( strchr( get_explicit_type_name( (ExplicitType)srcType ), ' ' ) != NULL ||
225 strchr( get_explicit_type_name( (ExplicitType)dstType ), ' ' ) != NULL )
226 continue;
227
228 if( test_explicit_s2v_function( deviceID, context, queue, kernel_explicit_s2v_set[ srcType ][ dstType ][ i ],
229 srcType, count, (ExplicitType)dstType, sizes[ i ], inputData ) != 0 )
230 {
231 log_error( "ERROR: Explicit cast of scalar %s to vector %s%d FAILED; skipping other %s vector tests\n",
232 get_explicit_type_name(srcType), get_explicit_type_name((ExplicitType)dstType), sizes[i], get_explicit_type_name((ExplicitType)dstType) );
233 failed = -1;
234 break;
235 }
236 }
237 }
238
239 return failed;
240 }
241
test_explicit_s2v_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)242 int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
243 {
244 char data[128];
245 RandomSeed seed(gRandomSeed);
246
247 generate_random_data( kChar, 128, seed, data );
248
249 return test_explicit_s2v_function_set( deviceID, context, queue, kChar, 128, data );
250 }
251
test_explicit_s2v_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)252 int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
253 {
254 unsigned char data[128];
255 RandomSeed seed(gRandomSeed);
256
257 generate_random_data( kUChar, 128, seed, data );
258
259 if( test_explicit_s2v_function_set( deviceID, context, queue, kUChar, 128, data ) != 0 )
260 return -1;
261 if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedChar, 128, data ) != 0 )
262 return -1;
263 return 0;
264 }
265
test_explicit_s2v_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)266 int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
267 {
268 short data[128];
269 RandomSeed seed(gRandomSeed);
270
271 generate_random_data( kShort, 128, seed, data );
272
273 if( test_explicit_s2v_function_set( deviceID, context, queue, kShort, 128, data ) != 0 )
274 return -1;
275 return 0;
276 }
277
test_explicit_s2v_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)278 int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
279 {
280 unsigned short data[128];
281 RandomSeed seed(gRandomSeed);
282
283 generate_random_data( kUShort, 128, seed, data );
284
285 if( test_explicit_s2v_function_set( deviceID, context, queue, kUShort, 128, data ) != 0 )
286 return -1;
287 if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedShort, 128, data ) != 0 )
288 return -1;
289 return 0;
290 }
291
test_explicit_s2v_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)292 int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
293 {
294 int data[128];
295 RandomSeed seed(gRandomSeed);
296
297 generate_random_data( kInt, 128, seed, data );
298
299 if( test_explicit_s2v_function_set( deviceID, context, queue, kInt, 128, data ) != 0 )
300 return -1;
301 return 0;
302 }
303
test_explicit_s2v_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)304 int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
305 {
306 unsigned int data[128];
307 RandomSeed seed(gRandomSeed);
308
309 generate_random_data( kUInt, 128, seed, data );
310
311 if( test_explicit_s2v_function_set( deviceID, context, queue, kUInt, 128, data ) != 0 )
312 return -1;
313 if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedInt, 128, data ) != 0 )
314 return -1;
315 return 0;
316 }
317
test_explicit_s2v_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)318 int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
319 {
320 cl_long data[128];
321 RandomSeed seed(gRandomSeed);
322
323 generate_random_data( kLong, 128, seed, data );
324
325 if( test_explicit_s2v_function_set( deviceID, context, queue, kLong, 128, data ) != 0 )
326 return -1;
327 return 0;
328 }
329
test_explicit_s2v_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)330 int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
331 {
332 cl_ulong data[128];
333 RandomSeed seed(gRandomSeed);
334
335 generate_random_data( kULong, 128, seed, data );
336
337 if( test_explicit_s2v_function_set( deviceID, context, queue, kULong, 128, data ) != 0 )
338 return -1;
339 if( test_explicit_s2v_function_set( deviceID, context, queue, kUnsignedLong, 128, data ) != 0 )
340 return -1;
341 return 0;
342 }
343
test_explicit_s2v_float(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)344 int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
345 {
346 float data[128];
347 RandomSeed seed(gRandomSeed);
348
349 generate_random_data( kFloat, 128, seed, data );
350
351 if( test_explicit_s2v_function_set( deviceID, context, queue, kFloat, 128, data ) != 0 )
352 return -1;
353 return 0;
354 }
355
356
test_explicit_s2v_double(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)357 int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
358 {
359 double data[128];
360 RandomSeed seed(gRandomSeed);
361
362 if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) {
363 log_info("Extension cl_khr_fp64 not supported. Skipping test.\n");
364 return 0;
365 }
366
367 generate_random_data( kDouble, 128, seed, data );
368
369 if( test_explicit_s2v_function_set( deviceID, context, queue, kDouble, 128, data ) != 0 )
370 return -1;
371 return 0;
372 }
373
374
375