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