• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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