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