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
25 #ifndef M_PI
26 #define M_PI 3.14159265358979323846264338327950288
27 #endif
28
29 #define CLAMP_KERNEL( type ) \
30 const char *clamp_##type##_kernel_code = \
31 EMIT_PRAGMA_DIRECTIVE \
32 "__kernel void test_clamp(__global " #type " *x, __global " #type " *minval, __global " #type " *maxval, __global " #type " *dst)\n" \
33 "{\n" \
34 " int tid = get_global_id(0);\n" \
35 "\n" \
36 " dst[tid] = clamp(x[tid], minval[tid], maxval[tid]);\n" \
37 "}\n";
38
39 #define CLAMP_KERNEL_V( type, size) \
40 const char *clamp_##type##size##_kernel_code = \
41 EMIT_PRAGMA_DIRECTIVE \
42 "__kernel void test_clamp(__global " #type #size " *x, __global " #type #size " *minval, __global " #type #size " *maxval, __global " #type #size " *dst)\n" \
43 "{\n" \
44 " int tid = get_global_id(0);\n" \
45 "\n" \
46 " dst[tid] = clamp(x[tid], minval[tid], maxval[tid]);\n" \
47 "}\n";
48
49 #define CLAMP_KERNEL_V3( type, size) \
50 const char *clamp_##type##size##_kernel_code = \
51 EMIT_PRAGMA_DIRECTIVE \
52 "__kernel void test_clamp(__global " #type " *x, __global " #type " *minval, __global " #type " *maxval, __global " #type " *dst)\n" \
53 "{\n" \
54 " int tid = get_global_id(0);\n" \
55 "\n" \
56 " vstore3(clamp(vload3(tid, x), vload3(tid,minval), vload3(tid,maxval)), tid, dst);\n" \
57 "}\n";
58
59 #define EMIT_PRAGMA_DIRECTIVE " "
60 CLAMP_KERNEL( float )
61 CLAMP_KERNEL_V( float, 2 )
62 CLAMP_KERNEL_V( float, 4 )
63 CLAMP_KERNEL_V( float, 8 )
64 CLAMP_KERNEL_V( float, 16 )
65 CLAMP_KERNEL_V3( float, 3)
66 #undef EMIT_PRAGMA_DIRECTIVE
67
68 #define EMIT_PRAGMA_DIRECTIVE "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"
69 CLAMP_KERNEL( double )
70 CLAMP_KERNEL_V( double, 2 )
71 CLAMP_KERNEL_V( double, 4 )
72 CLAMP_KERNEL_V( double, 8 )
73 CLAMP_KERNEL_V( double, 16 )
74 CLAMP_KERNEL_V3( double, 3 )
75 #undef EMIT_PRAGMA_DIRECTIVE
76
77 const char *clamp_float_codes[] = { clamp_float_kernel_code, clamp_float2_kernel_code, clamp_float4_kernel_code, clamp_float8_kernel_code, clamp_float16_kernel_code, clamp_float3_kernel_code };
78 const char *clamp_double_codes[] = { clamp_double_kernel_code, clamp_double2_kernel_code, clamp_double4_kernel_code, clamp_double8_kernel_code, clamp_double16_kernel_code, clamp_double3_kernel_code };
79
verify_clamp(float * x,float * minval,float * maxval,float * outptr,int n)80 static int verify_clamp(float *x, float *minval, float *maxval, float *outptr, int n)
81 {
82 float t;
83 int i;
84
85 for (i=0; i<n; i++)
86 {
87 t = fminf( fmaxf( x[ i ], minval[ i ] ), maxval[ i ] );
88 if (t != outptr[i])
89 {
90 log_error( "%d) verification error: clamp( %a, %a, %a) = *%a vs. %a\n", i, x[i], minval[i], maxval[i], t, outptr[i] );
91 return -1;
92 }
93 }
94
95 return 0;
96 }
97
verify_clamp_double(double * x,double * minval,double * maxval,double * outptr,int n)98 static int verify_clamp_double(double *x, double *minval, double *maxval, double *outptr, int n)
99 {
100 double t;
101 int i;
102
103 for (i=0; i<n; i++)
104 {
105 t = fmin( fmax( x[ i ], minval[ i ] ), maxval[ i ] );
106 if (t != outptr[i])
107 {
108 log_error( "%d) verification error: clamp( %a, %a, %a) = *%a vs. %a\n", i, x[i], minval[i], maxval[i], t, outptr[i] );
109 return -1;
110 }
111 }
112
113 return 0;
114 }
115
116 int
test_clamp(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)117 test_clamp(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
118 {
119 cl_mem streams[8];
120 cl_float *input_ptr[3], *output_ptr;
121 cl_double *input_ptr_double[3], *output_ptr_double = NULL;
122 cl_program *program;
123 cl_kernel *kernel;
124 size_t threads[1];
125 int num_elements;
126 int err;
127 int i, j;
128 MTdata d;
129
130 program = (cl_program*)malloc(sizeof(cl_program)*kTotalVecCount*2);
131 kernel = (cl_kernel*)malloc(sizeof(cl_kernel)*kTotalVecCount*2);
132
133 num_elements = n_elems * (1 << (kVectorSizeCount-1));
134
135 int test_double = 0;
136 if(is_extension_available( device, "cl_khr_fp64" )) {
137 log_info("Testing doubles.\n");
138 test_double = 1;
139 }
140
141
142 // why does this go from 0 to 2?? -- Oh, I see, there are four function
143 // arguments to the function, and 3 of them are inputs?
144 for( i = 0; i < 3; i++ )
145 {
146 input_ptr[i] = (cl_float*)malloc(sizeof(cl_float) * num_elements);
147 if (test_double) input_ptr_double[i] = (cl_double*)malloc(sizeof(cl_double) * num_elements);
148 }
149 output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
150 if (test_double) output_ptr_double = (cl_double*)malloc(sizeof(cl_double) * num_elements);
151
152 // why does this go from 0 to 3?
153 for( i = 0; i < 4; i++ )
154 {
155 streams[i] =
156 clCreateBuffer(context, CL_MEM_READ_WRITE,
157 sizeof(cl_float) * num_elements, NULL, NULL);
158 if (!streams[0])
159 {
160 log_error("clCreateBuffer failed\n");
161 return -1;
162 }
163 }
164 if (test_double)
165 for( i = 4; i < 8; i++ )
166 {
167 streams[i] =
168 clCreateBuffer(context, CL_MEM_READ_WRITE,
169 sizeof(cl_double) * num_elements, NULL, NULL);
170 if (!streams[0])
171 {
172 log_error("clCreateBuffer failed\n");
173 return -1;
174 }
175 }
176
177 d = init_genrand( gRandomSeed );
178 for( j = 0; j < num_elements; j++ )
179 {
180 input_ptr[0][j] = get_random_float(-0x20000000, 0x20000000, d);
181 input_ptr[1][j] = get_random_float(-0x20000000, 0x20000000, d);
182 input_ptr[2][j] = get_random_float(input_ptr[1][j], 0x20000000, d);
183
184 if (test_double) {
185 input_ptr_double[0][j] = get_random_double(-0x20000000, 0x20000000, d);
186 input_ptr_double[1][j] = get_random_double(-0x20000000, 0x20000000, d);
187 input_ptr_double[2][j] = get_random_double(input_ptr_double[1][j], 0x20000000, d);
188 }
189 }
190 free_mtdata(d); d = NULL;
191
192 for( i = 0; i < 3; i++ )
193 {
194 err = clEnqueueWriteBuffer( queue, streams[ i ], CL_TRUE, 0, sizeof( cl_float ) * num_elements, input_ptr[ i ], 0, NULL, NULL );
195 test_error( err, "Unable to write input buffer" );
196
197 if (test_double) {
198 err = clEnqueueWriteBuffer( queue, streams[ 4 + i ], CL_TRUE, 0, sizeof( cl_double ) * num_elements, input_ptr_double[ i ], 0, NULL, NULL );
199 test_error( err, "Unable to write input buffer" );
200 }
201 }
202
203 for( i = 0; i < kTotalVecCount; i++ )
204 {
205 err = create_single_kernel_helper( context, &program[ i ], &kernel[ i ], 1, &clamp_float_codes[ i ], "test_clamp" );
206 test_error( err, "Unable to create kernel" );
207
208 log_info("Just made a program for float, i=%d, size=%d, in slot %d\n", i, g_arrVecSizes[i], i);
209 fflush(stdout);
210
211 if (test_double) {
212 err = create_single_kernel_helper( context, &program[ kTotalVecCount + i ], &kernel[ kTotalVecCount + i ], 1, &clamp_double_codes[ i ], "test_clamp" );
213 log_info("Just made a program for double, i=%d, size=%d, in slot %d\n", i, g_arrVecSizes[i], kTotalVecCount+i);
214 fflush(stdout);
215 test_error( err, "Unable to create kernel" );
216 }
217 }
218
219 for( i = 0; i < kTotalVecCount; i++ )
220 {
221 for( j = 0; j < 4; j++ )
222 {
223 err = clSetKernelArg( kernel[ i ], j, sizeof( streams[ j ] ), &streams[ j ] );
224 test_error( err, "Unable to set kernel argument" );
225 }
226
227 threads[0] = (size_t)n_elems;
228
229 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
230 test_error( err, "Unable to execute kernel" );
231
232 err = clEnqueueReadBuffer( queue, streams[3], true, 0, sizeof(cl_float)*num_elements, (void *)output_ptr, 0, NULL, NULL );
233 test_error( err, "Unable to read results" );
234
235 if (verify_clamp(input_ptr[0], input_ptr[1], input_ptr[2], output_ptr, n_elems*((g_arrVecSizes[i]))))
236 {
237 log_error("CLAMP float%d test failed\n", ((g_arrVecSizes[i])));
238 err = -1;
239 }
240 else
241 {
242 log_info("CLAMP float%d test passed\n", ((g_arrVecSizes[i])));
243 err = 0;
244 }
245
246
247
248 if (err)
249 break;
250 }
251
252 // If the device supports double precision then test that
253 if (test_double)
254 {
255 for( ; i < 2*kTotalVecCount; i++ )
256 {
257
258 log_info("Start of test_double loop, i is %d\n", i);
259 for( j = 0; j < 4; j++ )
260 {
261 err = clSetKernelArg( kernel[i], j, sizeof( streams[j+4] ), &streams[j+4] );
262 test_error( err, "Unable to set kernel argument" );
263 }
264
265 threads[0] = (size_t)n_elems;
266
267 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
268 test_error( err, "Unable to execute kernel" );
269
270 err = clEnqueueReadBuffer( queue, streams[7], CL_TRUE, 0, sizeof(cl_double)*num_elements, (void *)output_ptr_double, 0, NULL, NULL );
271 test_error( err, "Unable to read results" );
272
273 if (verify_clamp_double(input_ptr_double[0], input_ptr_double[1], input_ptr_double[2], output_ptr_double, n_elems*g_arrVecSizes[(i-kTotalVecCount)]))
274 {
275 log_error("CLAMP double%d test failed\n", g_arrVecSizes[(i-kTotalVecCount)]);
276 err = -1;
277 }
278 else
279 {
280 log_info("CLAMP double%d test passed\n", g_arrVecSizes[(i-kTotalVecCount)]);
281 err = 0;
282 }
283
284 if (err)
285 break;
286 }
287 }
288
289
290 for( i = 0; i < ((test_double) ? 8 : 4); i++ )
291 {
292 clReleaseMemObject(streams[i]);
293 }
294 for (i=0; i < ((test_double) ? kTotalVecCount * 2-1 : kTotalVecCount); i++)
295 {
296 clReleaseKernel(kernel[i]);
297 clReleaseProgram(program[i]);
298 }
299 free(input_ptr[0]);
300 free(input_ptr[1]);
301 free(input_ptr[2]);
302 free(output_ptr);
303 free(program);
304 free(kernel);
305 if (test_double) {
306 free(input_ptr_double[0]);
307 free(input_ptr_double[1]);
308 free(input_ptr_double[2]);
309 free(output_ptr_double);
310 }
311
312 return err;
313 }
314
315
316