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 const char *binary_fn_code_pattern =
26 "%s\n" /* optional pragma */
27 "__kernel void test_fn(__global %s%s *x, __global %s%s *y, __global %s%s *dst)\n"
28 "{\n"
29 " int tid = get_global_id(0);\n"
30 "\n"
31 " dst[tid] = %s(x[tid], y[tid]);\n"
32 "}\n";
33
34 const char *binary_fn_code_pattern_v3 =
35 "%s\n" /* optional pragma */
36 "__kernel void test_fn(__global %s *x, __global %s *y, __global %s *dst)\n"
37 "{\n"
38 " int tid = get_global_id(0);\n"
39 "\n"
40 " vstore3(%s(vload3(tid,x), vload3(tid,y) ), tid, dst);\n"
41 "}\n";
42
43 const char *binary_fn_code_pattern_v3_scalar =
44 "%s\n" /* optional pragma */
45 "__kernel void test_fn(__global %s *x, __global %s *y, __global %s *dst)\n"
46 "{\n"
47 " int tid = get_global_id(0);\n"
48 "\n"
49 " vstore3(%s(vload3(tid,x), y[tid] ), tid, dst);\n"
50 "}\n";
51
test_binary_fn(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems,const char * fnName,bool vectorSecondParam,binary_verify_float_fn floatVerifyFn,binary_verify_double_fn doubleVerifyFn)52 int test_binary_fn( cl_device_id device, cl_context context, cl_command_queue queue, int n_elems,
53 const char *fnName, bool vectorSecondParam,
54 binary_verify_float_fn floatVerifyFn, binary_verify_double_fn doubleVerifyFn )
55 {
56 cl_mem streams[6];
57 cl_float *input_ptr[2], *output_ptr;
58 cl_double *input_ptr_double[2], *output_ptr_double=NULL;
59 cl_program *program;
60 cl_kernel *kernel;
61 size_t threads[1];
62 int num_elements;
63 int err;
64 int i, j;
65 MTdata d;
66
67 program = (cl_program*)malloc(sizeof(cl_program)*kTotalVecCount*2);
68 kernel = (cl_kernel*)malloc(sizeof(cl_kernel)*kTotalVecCount*2);
69
70 num_elements = n_elems * (1 << (kTotalVecCount-1));
71
72 int test_double = 0;
73 if(is_extension_available( device, "cl_khr_fp64" ))
74 {
75 log_info("Testing doubles.\n");
76 test_double = 1;
77 }
78
79 for( i = 0; i < 2; i++ )
80 {
81 input_ptr[i] = (cl_float*)malloc(sizeof(cl_float) * num_elements);
82 if (test_double) input_ptr_double[i] = (cl_double*)malloc(sizeof(cl_double) * num_elements);
83 }
84 output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
85 if (test_double) output_ptr_double = (cl_double*)malloc(sizeof(cl_double) * num_elements);
86
87 for( i = 0; i < 3; i++ )
88 {
89 streams[i] =
90 clCreateBuffer(context, CL_MEM_READ_WRITE,
91 sizeof(cl_float) * num_elements, NULL, &err);
92 test_error( err, "clCreateBuffer failed");
93 }
94
95 if (test_double)
96 for( i = 3; i < 6; i++ )
97 {
98 streams[i] =
99 clCreateBuffer(context, CL_MEM_READ_WRITE,
100 sizeof(cl_double) * num_elements, NULL, &err);
101 test_error(err, "clCreateBuffer failed");
102 }
103
104 d = init_genrand( gRandomSeed );
105 for( j = 0; j < num_elements; j++ )
106 {
107 input_ptr[0][j] = get_random_float(-0x20000000, 0x20000000, d);
108 input_ptr[1][j] = get_random_float(-0x20000000, 0x20000000, d);
109 if (test_double)
110 {
111 input_ptr_double[0][j] = get_random_double(-0x20000000, 0x20000000, d);
112 input_ptr_double[1][j] = get_random_double(-0x20000000, 0x20000000, d);
113 }
114 }
115 free_mtdata(d); d = NULL;
116
117 for( i = 0; i < 2; i++ )
118 {
119 err = clEnqueueWriteBuffer( queue, streams[ i ], CL_TRUE, 0, sizeof( cl_float ) * num_elements, input_ptr[ i ], 0, NULL, NULL );
120 test_error( err, "Unable to write input buffer" );
121
122 if (test_double)
123 {
124 err = clEnqueueWriteBuffer( queue, streams[ 3 + i ], CL_TRUE, 0, sizeof( cl_double ) * num_elements, input_ptr_double[ i ], 0, NULL, NULL );
125 test_error( err, "Unable to write input buffer" );
126 }
127 }
128
129 for( i = 0; i < kTotalVecCount; i++ )
130 {
131 char programSrc[ 10240 ];
132 char vecSizeNames[][ 3 ] = { "", "2", "4", "8", "16", "3" };
133
134 if(i >= kVectorSizeCount) {
135 // do vec3 print
136
137 if(vectorSecondParam) {
138 sprintf( programSrc,binary_fn_code_pattern_v3, "", "float", "float", "float", fnName );
139 } else {
140 sprintf( programSrc,binary_fn_code_pattern_v3_scalar, "", "float", "float", "float", fnName );
141 }
142 } else {
143 // do regular
144 sprintf( programSrc, binary_fn_code_pattern, "", "float", vecSizeNames[ i ], "float", vectorSecondParam ? vecSizeNames[ i ] : "", "float", vecSizeNames[ i ], fnName );
145 }
146 const char *ptr = programSrc;
147 err = create_single_kernel_helper( context, &program[ i ], &kernel[ i ], 1, &ptr, "test_fn" );
148 test_error( err, "Unable to create kernel" );
149
150 if (test_double)
151 {
152 if(i >= kVectorSizeCount) {
153 if(vectorSecondParam) {
154 sprintf( programSrc, binary_fn_code_pattern_v3, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable",
155 "double", "double", "double", fnName );
156 } else {
157
158 sprintf( programSrc, binary_fn_code_pattern_v3_scalar, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable",
159 "double", "double", "double", fnName );
160 }
161 } else {
162 sprintf( programSrc, binary_fn_code_pattern, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable",
163 "double", vecSizeNames[ i ], "double", vectorSecondParam ? vecSizeNames[ i ] : "", "double", vecSizeNames[ i ], fnName );
164 }
165 ptr = programSrc;
166 err = create_single_kernel_helper( context, &program[ kTotalVecCount + i ], &kernel[ kTotalVecCount + i ], 1, &ptr, "test_fn" );
167 test_error( err, "Unable to create kernel" );
168 }
169 }
170
171 for( i = 0; i < kTotalVecCount; i++ )
172 {
173 for( j = 0; j < 3; j++ )
174 {
175 err = clSetKernelArg( kernel[ i ], j, sizeof( streams[ j ] ), &streams[ j ] );
176 test_error( err, "Unable to set kernel argument" );
177 }
178
179 threads[0] = (size_t)n_elems;
180
181 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
182 test_error( err, "Unable to execute kernel" );
183
184 err = clEnqueueReadBuffer( queue, streams[2], true, 0, sizeof(cl_float)*num_elements, (void *)output_ptr, 0, NULL, NULL );
185 test_error( err, "Unable to read results" );
186
187
188
189 if( floatVerifyFn( input_ptr[0], input_ptr[1], output_ptr, n_elems, ((g_arrVecSizes[i])) ) )
190 {
191 log_error(" float%d%s test failed\n", ((g_arrVecSizes[i])), vectorSecondParam ? "" : ", float");
192 err = -1;
193 }
194 else
195 {
196 log_info(" float%d%s test passed\n", ((g_arrVecSizes[i])), vectorSecondParam ? "" : ", float");
197 err = 0;
198 }
199
200 if (err)
201 break;
202 }
203
204 if (test_double)
205 {
206 for( i = 0; i < kTotalVecCount; i++ )
207 {
208 for( j = 0; j < 3; j++ )
209 {
210 err = clSetKernelArg( kernel[ kTotalVecCount + i ], j, sizeof( streams[ 3 + j ] ), &streams[ 3 + j ] );
211 test_error( err, "Unable to set kernel argument" );
212 }
213
214 threads[0] = (size_t)n_elems;
215
216 err = clEnqueueNDRangeKernel( queue, kernel[kTotalVecCount + i], 1, NULL, threads, NULL, 0, NULL, NULL );
217 test_error( err, "Unable to execute kernel" );
218
219 err = clEnqueueReadBuffer( queue, streams[5], CL_TRUE, 0, sizeof(cl_double)*num_elements, (void *)output_ptr_double, 0, NULL, NULL );
220 test_error( err, "Unable to read results" );
221
222 if( doubleVerifyFn( input_ptr_double[0], input_ptr_double[1], output_ptr_double, n_elems, ((g_arrVecSizes[i]))))
223 {
224 log_error(" double%d%s test failed\n", ((g_arrVecSizes[i])), vectorSecondParam ? "" : ", double");
225 err = -1;
226 }
227 else
228 {
229 log_info(" double%d%s test passed\n", ((g_arrVecSizes[i])), vectorSecondParam ? "" : ", double");
230 err = 0;
231 }
232
233 if (err)
234 break;
235 }
236 }
237
238
239 for( i = 0; i < ((test_double) ? 6 : 3); i++ )
240 {
241 clReleaseMemObject(streams[i]);
242 }
243 for (i=0; i < ((test_double) ? kTotalVecCount * 2 : kTotalVecCount) ; i++)
244 {
245 clReleaseKernel(kernel[i]);
246 clReleaseProgram(program[i]);
247 }
248 free(input_ptr[0]);
249 free(input_ptr[1]);
250 free(output_ptr);
251 free(program);
252 free(kernel);
253
254 if (test_double)
255 {
256 free(input_ptr_double[0]);
257 free(input_ptr_double[1]);
258 free(output_ptr_double);
259 }
260
261 return err;
262 }
263
264
265