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
17 #include "function_list.h"
18 #include "test_functions.h"
19 #include "utility.h"
20
21 #include <cstring>
22
BuildKernel(const char * name,int vectorSize,cl_kernel * k,cl_program * p,bool relaxedMode)23 static int BuildKernel(const char *name, int vectorSize, cl_kernel *k,
24 cl_program *p, bool relaxedMode)
25 {
26 const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
27 "__kernel void math_kernel",
28 sizeNames[vectorSize],
29 "( __global double",
30 sizeNames[vectorSize],
31 "* out, __global ulong",
32 sizeNames[vectorSize],
33 "* in )\n"
34 "{\n"
35 " size_t i = get_global_id(0);\n"
36 " out[i] = ",
37 name,
38 "( in[i] );\n"
39 "}\n" };
40
41 const char *c3[] = {
42 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
43 "__kernel void math_kernel",
44 sizeNames[vectorSize],
45 "( __global double* out, __global ulong* in )\n"
46 "{\n"
47 " size_t i = get_global_id(0);\n"
48 " if( i + 1 < get_global_size(0) )\n"
49 " {\n"
50 " ulong3 u0 = vload3( 0, in + 3 * i );\n"
51 " double3 f0 = ",
52 name,
53 "( u0 );\n"
54 " vstore3( f0, 0, out + 3*i );\n"
55 " }\n"
56 " else\n"
57 " {\n"
58 " size_t parity = i & 1; // Figure out how many elements are "
59 "left over after BUFFER_SIZE % (3*sizeof(float)). Assume power of two "
60 "buffer size \n"
61 " ulong3 u0;\n"
62 " switch( parity )\n"
63 " {\n"
64 " case 1:\n"
65 " u0 = (ulong3)( in[3*i], 0xdeaddeaddeaddeadUL, "
66 "0xdeaddeaddeaddeadUL ); \n"
67 " break;\n"
68 " case 0:\n"
69 " u0 = (ulong3)( in[3*i], in[3*i+1], "
70 "0xdeaddeaddeaddeadUL ); \n"
71 " break;\n"
72 " }\n"
73 " double3 f0 = ",
74 name,
75 "( u0 );\n"
76 " switch( parity )\n"
77 " {\n"
78 " case 0:\n"
79 " out[3*i+1] = f0.y; \n"
80 " // fall through\n"
81 " case 1:\n"
82 " out[3*i] = f0.x; \n"
83 " break;\n"
84 " }\n"
85 " }\n"
86 "}\n"
87 };
88
89 const char **kern = c;
90 size_t kernSize = sizeof(c) / sizeof(c[0]);
91
92 if (sizeValues[vectorSize] == 3)
93 {
94 kern = c3;
95 kernSize = sizeof(c3) / sizeof(c3[0]);
96 }
97
98 char testName[32];
99 snprintf(testName, sizeof(testName) - 1, "math_kernel%s",
100 sizeNames[vectorSize]);
101
102 return MakeKernel(kern, (cl_uint)kernSize, testName, k, p, relaxedMode);
103 }
104
105 typedef struct BuildKernelInfo
106 {
107 cl_uint offset; // the first vector size to build
108 cl_kernel *kernels;
109 cl_program *programs;
110 const char *nameInCode;
111 bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
112 } BuildKernelInfo;
113
BuildKernelFn(cl_uint job_id,cl_uint thread_id UNUSED,void * p)114 static cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
115 {
116 BuildKernelInfo *info = (BuildKernelInfo *)p;
117 cl_uint i = info->offset + job_id;
118 return BuildKernel(info->nameInCode, i, info->kernels + i,
119 info->programs + i, info->relaxedMode);
120 }
121
random64(MTdata d)122 static cl_ulong random64(MTdata d)
123 {
124 return (cl_ulong)genrand_int32(d) | ((cl_ulong)genrand_int32(d) << 32);
125 }
126
TestFunc_Double_ULong(const Func * f,MTdata d,bool relaxedMode)127 int TestFunc_Double_ULong(const Func *f, MTdata d, bool relaxedMode)
128 {
129 int error;
130 cl_program programs[VECTOR_SIZE_COUNT];
131 cl_kernel kernels[VECTOR_SIZE_COUNT];
132 float maxError = 0.0f;
133 int ftz = f->ftz || gForceFTZ;
134 double maxErrorVal = 0.0f;
135 uint64_t step = getTestStep(sizeof(cl_double), BUFFER_SIZE);
136
137 logFunctionInfo(f->name, sizeof(cl_double), relaxedMode);
138
139 Force64BitFPUPrecision();
140
141 // Init the kernels
142 {
143 BuildKernelInfo build_info = { gMinVectorSizeIndex, kernels, programs,
144 f->nameInCode, relaxedMode };
145 if ((error = ThreadPool_Do(BuildKernelFn,
146 gMaxVectorSizeIndex - gMinVectorSizeIndex,
147 &build_info)))
148 return error;
149 }
150
151 for (uint64_t i = 0; i < (1ULL << 32); i += step)
152 {
153 // Init input array
154 cl_ulong *p = (cl_ulong *)gIn;
155 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_ulong); j++)
156 p[j] = random64(d);
157
158 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
159 BUFFER_SIZE, gIn, 0, NULL, NULL)))
160 {
161 vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
162 return error;
163 }
164
165 // write garbage into output arrays
166 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
167 {
168 uint32_t pattern = 0xffffdead;
169 memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
170 if ((error =
171 clEnqueueWriteBuffer(gQueue, gOutBuffer[j], CL_FALSE, 0,
172 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
173 {
174 vlog_error("\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
175 error, j);
176 goto exit;
177 }
178 }
179
180 // Run the kernels
181 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
182 {
183 size_t vectorSize = sizeValues[j] * sizeof(cl_double);
184 size_t localCount = (BUFFER_SIZE + vectorSize - 1) / vectorSize;
185 if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
186 &gOutBuffer[j])))
187 {
188 LogBuildError(programs[j]);
189 goto exit;
190 }
191 if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
192 &gInBuffer)))
193 {
194 LogBuildError(programs[j]);
195 goto exit;
196 }
197
198 if ((error =
199 clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
200 &localCount, NULL, 0, NULL, NULL)))
201 {
202 vlog_error("FAILED -- could not execute kernel\n");
203 goto exit;
204 }
205 }
206
207 // Get that moving
208 if ((error = clFlush(gQueue))) vlog("clFlush failed\n");
209
210 // Calculate the correctly rounded reference result
211 double *r = (double *)gOut_Ref;
212 cl_ulong *s = (cl_ulong *)gIn;
213 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
214 r[j] = (double)f->dfunc.f_u(s[j]);
215
216 // Read the data back
217 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
218 {
219 if ((error =
220 clEnqueueReadBuffer(gQueue, gOutBuffer[j], CL_TRUE, 0,
221 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
222 {
223 vlog_error("ReadArray failed %d\n", error);
224 goto exit;
225 }
226 }
227
228 if (gSkipCorrectnessTesting) break;
229
230 // Verify data
231 uint64_t *t = (uint64_t *)gOut_Ref;
232 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
233 {
234 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
235 {
236 uint64_t *q = (uint64_t *)(gOut[k]);
237
238 // If we aren't getting the correctly rounded result
239 if (t[j] != q[j])
240 {
241 double test = ((double *)q)[j];
242 long double correct = f->dfunc.f_u(s[j]);
243 float err = Bruteforce_Ulp_Error_Double(test, correct);
244 int fail = !(fabsf(err) <= f->double_ulps);
245
246 if (fail)
247 {
248 if (ftz)
249 {
250 // retry per section 6.5.3.2
251 if (IsDoubleResultSubnormal(correct,
252 f->double_ulps))
253 {
254 fail = fail && (test != 0.0);
255 if (!fail) err = 0.0f;
256 }
257 }
258 }
259 if (fabsf(err) > maxError)
260 {
261 maxError = fabsf(err);
262 maxErrorVal = s[j];
263 }
264 if (fail)
265 {
266 vlog_error("\n%s%sD: %f ulp error at 0x%16.16llx: "
267 "*%.13la vs. %.13la\n",
268 f->name, sizeNames[k], err,
269 ((uint64_t *)gIn)[j],
270 ((double *)gOut_Ref)[j], test);
271 error = -1;
272 goto exit;
273 }
274 }
275 }
276 }
277
278 if (0 == (i & 0x0fffffff))
279 {
280 if (gVerboseBruteForce)
281 {
282 vlog("base:%14u step:%10zu bufferSize:%10zd \n", i, step,
283 BUFFER_SIZE);
284 }
285 else
286 {
287 vlog(".");
288 }
289 fflush(stdout);
290 }
291 }
292
293 if (!gSkipCorrectnessTesting)
294 {
295 if (gWimpyMode)
296 vlog("Wimp pass");
297 else
298 vlog("passed");
299
300 vlog("\t%8.2f @ %a", maxError, maxErrorVal);
301 }
302
303 vlog("\n");
304
305 exit:
306 // Release
307 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
308 {
309 clReleaseKernel(kernels[k]);
310 clReleaseProgram(programs[k]);
311 }
312
313 return error;
314 }
315