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 <string.h>
20 #include <limits.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24 #include "procs.h"
25
26
verify_abs_char(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)27 static int verify_abs_char( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
28 {
29 const cl_char *inA = (const cl_char*) p;
30 const cl_uchar *outptr = (const cl_uchar*) q;
31 size_t i;
32 for( i = 0; i < n; i++ )
33 {
34 cl_uchar r = inA[i];
35 if( inA[i] < 0 )
36 r = -inA[i];
37 if( r != outptr[i] )
38 { log_info( "%ld) Failure for abs( (char%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
39 }
40 return 0;
41 }
42
43
verify_abs_short(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)44 static int verify_abs_short( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
45 {
46 const cl_short *inA = (const cl_short*) p;
47 const cl_ushort *outptr = (const cl_ushort*) q;
48 size_t i;
49 for( i = 0; i < n; i++ )
50 {
51 cl_ushort r = inA[i];
52 if( inA[i] < 0 )
53 r = -inA[i];
54 if( r != outptr[i] )
55 { log_info( "%ld) Failure for abs( (short%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
56 }
57 return 0;
58 }
59
verify_abs_int(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)60 static int verify_abs_int( const void *p, const void *q, size_t n, const char *sizeName , size_t vecSize)
61 {
62 const cl_int *inA = (const cl_int*) p;
63 const cl_uint *outptr = (const cl_uint*) q;
64 size_t i;
65 for( i = 0; i < n; i++ )
66 {
67 cl_uint r = inA[i];
68 if( inA[i] < 0 )
69 r = -inA[i];
70 if( r != outptr[i] )
71 { log_info( "%ld) Failure for abs( (int%s) 0x%2.2x) = *0x%8.8x vs 0x%8.8x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
72 }
73 return 0;
74 }
75
verify_abs_long(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)76 static int verify_abs_long( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
77 {
78 const cl_long *inA = (const cl_long*) p;
79 const cl_ulong *outptr = (const cl_ulong*) q;
80 size_t i;
81 for( i = 0; i < n; i++ )
82 {
83 cl_ulong r = inA[i];
84 if( inA[i] < 0 )
85 r = -inA[i];
86 if( r != outptr[i] )
87 { log_info( "%ld) Failure for abs( (long%s) 0x%16.16llx) = *0x%16.16llx vs 0x%16.16llx\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
88 }
89 return 0;
90 }
91
92
93
verify_abs_uchar(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)94 static int verify_abs_uchar( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
95 {
96 const cl_uchar *inA = (const cl_uchar*) p;
97 const cl_uchar *outptr = (const cl_uchar*) q;
98 size_t i;
99 for( i = 0; i < n; i++ )
100 {
101 cl_uchar r = inA[i];
102 if( r != outptr[i] )
103 { log_info( "%ld) Failure for abs( (uchar%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
104 }
105 return 0;
106 }
107
108
verify_abs_ushort(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)109 static int verify_abs_ushort( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
110 {
111 const cl_ushort *inA = (const cl_ushort*) p;
112 const cl_ushort *outptr = (const cl_ushort*) q;
113 size_t i;
114 for( i = 0; i < n; i++ )
115 {
116 cl_ushort r = inA[i];
117 if( r != outptr[i] )
118 { log_info( "%ld) Failure for abs( (short%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
119 }
120 return 0;
121 }
122
verify_abs_uint(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)123 static int verify_abs_uint( const void *p, const void *q, size_t n, const char *sizeName , size_t vecSize)
124 {
125 const cl_uint *inA = (const cl_uint*) p;
126 const cl_uint *outptr = (const cl_uint*) q;
127 size_t i;
128 for( i = 0; i < n; i++ )
129 {
130 cl_uint r = inA[i];
131 if( r != outptr[i] )
132 { log_info( "%ld) Failure for abs( (int%s) 0x%2.2x) = *0x%8.8x vs 0x%8.8x\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
133 }
134 return 0;
135 }
136
verify_abs_ulong(const void * p,const void * q,size_t n,const char * sizeName,size_t vecSize)137 static int verify_abs_ulong( const void *p, const void *q, size_t n, const char *sizeName, size_t vecSize )
138 {
139 const cl_ulong *inA = (const cl_ulong*) p;
140 const cl_ulong *outptr = (const cl_ulong*) q;
141 size_t i;
142 for( i = 0; i < n; i++ )
143 {
144 cl_ulong r = inA[i];
145 if( r != outptr[i] )
146 { log_info( "%ld) Failure for abs( (long%s) 0x%16.16llx) = *0x%16.16llx vs 0x%16.16llx\n", i, sizeName, inA[i],r, outptr[i] ); return -1; }
147 }
148 return 0;
149 }
150
151
152 typedef int (*verifyFunc)( const void *, const void *, size_t n, const char *sizeName, size_t vecSize );
153 static const verifyFunc verify[] = {
154 verify_abs_char, verify_abs_short, verify_abs_int, verify_abs_long,
155 verify_abs_uchar, verify_abs_ushort, verify_abs_uint, verify_abs_ulong
156 };
157
158 static const char *test_str_names[] = { "char", "short", "int", "long" ,
159 "uchar", "ushort", "uint", "ulong"};
160 static const char *test_ustr_names[] = { "uchar", "ushort", "uint", "ulong" ,
161 "uchar", "ushort", "uint", "ulong"};
162 static const int vector_sizes[] = {1, 2, 3, 4, 8, 16};
163 static const char *vector_size_names[] = { "", "2", "3", "4", "8", "16" };
164 static const char *vector_size_names_io_types[] = { "", "2", "", "4", "8", "16" };
165 static const size_t kSizes[9] = { 1, 2, 4, 8, 1, 2, 4, 8 };
166
167 static const char * source_loads[] = {
168 "srcA[tid]",
169 "vload3(tid, srcA)"
170 };
171
172 static const char * dest_stores[] = {
173 " dst[tid] = tmp;\n",
174 " vstore3(tmp, tid, dst);\n"
175 };
176
test_integer_abs(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)177 int test_integer_abs(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
178 {
179 cl_int *input_ptr, *output_ptr, *p;
180 int err;
181 int i;
182 cl_uint vectorSizeIdx;
183 cl_uint type;
184 MTdata d;
185 int fail_count = 0;
186
187 size_t length = sizeof(cl_int) * 4 * n_elems;
188
189 input_ptr = (cl_int*)malloc(length);
190 output_ptr = (cl_int*)malloc(length);
191
192 p = input_ptr;
193 d = init_genrand( gRandomSeed );
194 for (i=0; i<n_elems * 4; i++)
195 p[i] = genrand_int32(d);
196 free_mtdata(d); d = NULL;
197
198 for( type = 0; type < sizeof( test_str_names ) / sizeof( test_str_names[0] ); type++ )
199 {
200 //embedded devices don't support long/ulong so skip over
201 if (! gHasLong && strstr(test_str_names[type],"long"))
202 {
203 log_info( "WARNING: 64 bit integers are not supported on this device. Skipping %s\n", test_str_names[type] );
204 continue;
205 }
206
207 verifyFunc f = verify[ type ];
208
209 size_t elementCount = length / kSizes[type];
210 cl_mem streams[2];
211
212 log_info( "%s", test_str_names[type] );
213 fflush( stdout );
214
215 // Set up data streams for the type
216 streams[0] = clCreateBuffer(context, 0, length, NULL, NULL);
217 if (!streams[0])
218 {
219 log_error("clCreateBuffer failed\n");
220 return -1;
221 }
222 streams[1] = clCreateBuffer(context, 0, length, NULL, NULL);
223 if (!streams[1])
224 {
225 log_error("clCreateBuffer failed\n");
226 return -1;
227 }
228
229 err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, length, input_ptr, 0, NULL, NULL);
230 if (err != CL_SUCCESS)
231 {
232 log_error("clEnqueueWriteBuffer failed\n");
233 return -1;
234 }
235
236
237
238 for( vectorSizeIdx = 0; vectorSizeIdx < sizeof( vector_size_names ) / sizeof( vector_size_names[0] ); vectorSizeIdx++ )
239 {
240 cl_program program = NULL;
241 cl_kernel kernel = NULL;
242
243 const char *source[] = {
244 "__kernel void test_abs_",
245 test_str_names[type],
246 vector_size_names[vectorSizeIdx],
247 "(__global ", test_str_names[type],
248 vector_size_names_io_types[vectorSizeIdx],
249 " *srcA, __global ", test_ustr_names[type],
250 vector_size_names_io_types[vectorSizeIdx],
251 " *dst)\n"
252 "{\n"
253 " int tid = get_global_id(0);\n"
254 "\n"
255 " ", test_ustr_names[type], vector_size_names[vectorSizeIdx],
256 " tmp = abs(", source_loads[!!(vector_sizes[vectorSizeIdx]==3)], ");\n",
257 dest_stores[!!(vector_sizes[vectorSizeIdx]==3)],
258 "}\n"
259 };
260
261 char kernelName[128];
262 snprintf( kernelName, sizeof( kernelName ), "test_abs_%s%s", test_str_names[type], vector_size_names[vectorSizeIdx] );
263 err = create_single_kernel_helper(context, &program, &kernel, sizeof( source ) / sizeof( source[0] ), source, kernelName );
264 if (err)
265 return -1;
266
267 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
268 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
269 if (err != CL_SUCCESS)
270 {
271 log_error("clSetKernelArgs failed\n");
272 return -1;
273 }
274
275 //Wipe the output buffer clean
276 uint32_t pattern = 0xdeadbeef;
277 memset_pattern4( output_ptr, &pattern, length );
278 err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
279 if (err != CL_SUCCESS)
280 {
281 log_error("clEnqueueWriteBuffer failed\n");
282 return -1;
283 }
284
285 size_t size = elementCount / ((vector_sizes[vectorSizeIdx]));
286 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL, 0, NULL, NULL);
287 if (err != CL_SUCCESS)
288 {
289 log_error("clEnqueueNDRangeKernel failed\n");
290 return -1;
291 }
292
293 err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
294 if (err != CL_SUCCESS)
295 {
296 log_error("clEnqueueReadBuffer failed\n");
297 return -1;
298 }
299
300 char *inP = (char *)input_ptr;
301 char *outP = (char *)output_ptr;
302
303 for( size_t e = 0; e < size; e++ )
304 {
305 if( f( inP, outP, (vector_sizes[vectorSizeIdx]), vector_size_names[vectorSizeIdx], vector_sizes[vectorSizeIdx] ) ) {
306 ++fail_count; break; // return -1;
307 }
308 inP += kSizes[type] * (vector_sizes[vectorSizeIdx] );
309 outP += kSizes[type] * (vector_sizes[vectorSizeIdx]);
310 }
311
312 clReleaseKernel( kernel );
313 clReleaseProgram( program );
314 log_info( "." );
315 fflush( stdout );
316 }
317
318 clReleaseMemObject( streams[0] );
319 clReleaseMemObject( streams[1] );
320 log_info( "done\n" );
321 }
322
323 if(fail_count) {
324 log_info("Failed on %d types\n", fail_count);
325 return -1;
326 }
327
328 free(input_ptr);
329 free(output_ptr);
330
331 return err;
332 }
333
334
335