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 <climits>
24 #include <cstring>
25
26 namespace {
27
BuildKernelFn(cl_uint job_id,cl_uint thread_id UNUSED,void * p)28 cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
29 {
30 BuildKernelInfo &info = *(BuildKernelInfo *)p;
31 auto generator = [](const std::string &kernel_name, const char *builtin,
32 cl_uint vector_size_index) {
33 return GetUnaryKernel(kernel_name, builtin, ParameterType::Double,
34 ParameterType::Int, ParameterType::Double,
35 vector_size_index);
36 };
37 return BuildKernels(info, job_id, generator);
38 }
39
abs_cl_long(cl_long i)40 cl_ulong abs_cl_long(cl_long i)
41 {
42 cl_long mask = i >> 63;
43 return (i ^ mask) - mask;
44 }
45
46 } // anonymous namespace
47
TestFunc_DoubleI_Double(const Func * f,MTdata d,bool relaxedMode)48 int TestFunc_DoubleI_Double(const Func *f, MTdata d, bool relaxedMode)
49 {
50 int error;
51 Programs programs;
52 const unsigned thread_id = 0; // Test is currently not multithreaded.
53 KernelMatrix kernels;
54 float maxError = 0.0f;
55 int64_t maxError2 = 0;
56 int ftz = f->ftz || gForceFTZ;
57 double maxErrorVal = 0.0f;
58 double maxErrorVal2 = 0.0f;
59 cl_ulong maxiError = f->double_ulps == INFINITY ? CL_ULONG_MAX : 0;
60 uint64_t step = getTestStep(sizeof(cl_double), BUFFER_SIZE);
61 int scale =
62 (int)((1ULL << 32) / (16 * BUFFER_SIZE / sizeof(cl_double)) + 1);
63
64 logFunctionInfo(f->name, sizeof(cl_double), relaxedMode);
65
66 Force64BitFPUPrecision();
67
68 // Init the kernels
69 BuildKernelInfo build_info{ 1, kernels, programs, f->nameInCode,
70 relaxedMode };
71 if ((error = ThreadPool_Do(BuildKernelFn,
72 gMaxVectorSizeIndex - gMinVectorSizeIndex,
73 &build_info)))
74 return error;
75
76 for (uint64_t i = 0; i < (1ULL << 32); i += step)
77 {
78 // Init input array
79 double *p = (double *)gIn;
80 if (gWimpyMode)
81 {
82 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
83 p[j] = DoubleFromUInt32((uint32_t)i + j * scale);
84 }
85 else
86 {
87 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
88 p[j] = DoubleFromUInt32((uint32_t)i + j);
89 }
90 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
91 BUFFER_SIZE, gIn, 0, NULL, NULL)))
92 {
93 vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
94 return error;
95 }
96
97 // Write garbage into output arrays
98 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
99 {
100 uint32_t pattern = 0xffffdead;
101 if (gHostFill)
102 {
103 memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
104 if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer[j],
105 CL_FALSE, 0, BUFFER_SIZE,
106 gOut[j], 0, NULL, NULL)))
107 {
108 vlog_error(
109 "\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
110 error, j);
111 return error;
112 }
113
114 memset_pattern4(gOut2[j], &pattern, BUFFER_SIZE);
115 if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer2[j],
116 CL_FALSE, 0, BUFFER_SIZE,
117 gOut2[j], 0, NULL, NULL)))
118 {
119 vlog_error(
120 "\n*** Error %d in clEnqueueWriteBuffer2b(%d) ***\n",
121 error, j);
122 return error;
123 }
124 }
125 else
126 {
127 if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer[j],
128 &pattern, sizeof(pattern), 0,
129 BUFFER_SIZE, 0, NULL, NULL)))
130 {
131 vlog_error("Error: clEnqueueFillBuffer 1 failed! err: %d\n",
132 error);
133 return error;
134 }
135
136 if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer2[j],
137 &pattern, sizeof(pattern), 0,
138 BUFFER_SIZE, 0, NULL, NULL)))
139 {
140 vlog_error("Error: clEnqueueFillBuffer 2 failed! err: %d\n",
141 error);
142 return error;
143 }
144 }
145 }
146
147 // Run the kernels
148 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
149 {
150 size_t vectorSize = sizeValues[j] * sizeof(cl_double);
151 size_t localCount = (BUFFER_SIZE + vectorSize - 1) / vectorSize;
152 if ((error = clSetKernelArg(kernels[j][thread_id], 0,
153 sizeof(gOutBuffer[j]), &gOutBuffer[j])))
154 {
155 LogBuildError(programs[j]);
156 return error;
157 }
158 if ((error =
159 clSetKernelArg(kernels[j][thread_id], 1,
160 sizeof(gOutBuffer2[j]), &gOutBuffer2[j])))
161 {
162 LogBuildError(programs[j]);
163 return error;
164 }
165 if ((error = clSetKernelArg(kernels[j][thread_id], 2,
166 sizeof(gInBuffer), &gInBuffer)))
167 {
168 LogBuildError(programs[j]);
169 return error;
170 }
171
172 if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j][thread_id],
173 1, NULL, &localCount, NULL, 0,
174 NULL, NULL)))
175 {
176 vlog_error("FAILED -- could not execute kernel\n");
177 return error;
178 }
179 }
180
181 // Get that moving
182 if ((error = clFlush(gQueue))) vlog("clFlush failed\n");
183
184 // Calculate the correctly rounded reference result
185 double *r = (double *)gOut_Ref;
186 int *r2 = (int *)gOut_Ref2;
187 double *s = (double *)gIn;
188 for (size_t j = 0; j < BUFFER_SIZE / sizeof(double); j++)
189 r[j] = (double)f->dfunc.f_fpI(s[j], r2 + j);
190
191 // Read the data back
192 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
193 {
194 if ((error =
195 clEnqueueReadBuffer(gQueue, gOutBuffer[j], CL_TRUE, 0,
196 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
197 {
198 vlog_error("ReadArray failed %d\n", error);
199 return error;
200 }
201 if ((error =
202 clEnqueueReadBuffer(gQueue, gOutBuffer2[j], CL_TRUE, 0,
203 BUFFER_SIZE, gOut2[j], 0, NULL, NULL)))
204 {
205 vlog_error("ReadArray2 failed %d\n", error);
206 return error;
207 }
208 }
209
210 if (gSkipCorrectnessTesting) break;
211
212 // Verify data
213 uint64_t *t = (uint64_t *)gOut_Ref;
214 int32_t *t2 = (int32_t *)gOut_Ref2;
215 for (size_t j = 0; j < BUFFER_SIZE / sizeof(double); j++)
216 {
217 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
218 {
219 uint64_t *q = (uint64_t *)(gOut[k]);
220 int32_t *q2 = (int32_t *)(gOut2[k]);
221
222 // If we aren't getting the correctly rounded result
223 if (t[j] != q[j] || t2[j] != q2[j])
224 {
225 double test = ((double *)q)[j];
226 int correct2 = INT_MIN;
227 long double correct = f->dfunc.f_fpI(s[j], &correct2);
228 float err = Bruteforce_Ulp_Error_Double(test, correct);
229 cl_long iErr = (long long)q2[j] - (long long)correct2;
230 int fail = !(fabsf(err) <= f->double_ulps
231 && abs_cl_long(iErr) <= maxiError);
232 if (ftz || relaxedMode)
233 {
234 // retry per section 6.5.3.2
235 if (IsDoubleResultSubnormal(correct, f->double_ulps))
236 {
237 fail = fail && !(test == 0.0f && iErr == 0);
238 if (!fail) err = 0.0f;
239 }
240
241 // retry per section 6.5.3.3
242 if (IsDoubleSubnormal(s[j]))
243 {
244 int correct5, correct6;
245 long double correct3 =
246 f->dfunc.f_fpI(0.0, &correct5);
247 long double correct4 =
248 f->dfunc.f_fpI(-0.0, &correct6);
249 float err2 =
250 Bruteforce_Ulp_Error_Double(test, correct3);
251 float err3 =
252 Bruteforce_Ulp_Error_Double(test, correct4);
253 cl_long iErr2 =
254 (long long)q2[j] - (long long)correct5;
255 cl_long iErr3 =
256 (long long)q2[j] - (long long)correct6;
257
258 // Did +0 work?
259 if (fabsf(err2) <= f->double_ulps
260 && abs_cl_long(iErr2) <= maxiError)
261 {
262 err = err2;
263 iErr = iErr2;
264 fail = 0;
265 }
266 // Did -0 work?
267 else if (fabsf(err3) <= f->double_ulps
268 && abs_cl_long(iErr3) <= maxiError)
269 {
270 err = err3;
271 iErr = iErr3;
272 fail = 0;
273 }
274
275 // retry per section 6.5.3.4
276 if (fail
277 && (IsDoubleResultSubnormal(correct2,
278 f->double_ulps)
279 || IsDoubleResultSubnormal(correct3,
280 f->double_ulps)))
281 {
282 fail = fail
283 && !(test == 0.0f
284 && (abs_cl_long(iErr2) <= maxiError
285 || abs_cl_long(iErr3)
286 <= maxiError));
287 if (!fail)
288 {
289 err = 0.0f;
290 iErr = 0;
291 }
292 }
293 }
294 }
295 if (fabsf(err) > maxError)
296 {
297 maxError = fabsf(err);
298 maxErrorVal = s[j];
299 }
300 if (llabs(iErr) > maxError2)
301 {
302 maxError2 = llabs(iErr);
303 maxErrorVal2 = s[j];
304 }
305
306 if (fail)
307 {
308 vlog_error("\nERROR: %sD%s: {%f, %d} ulp error at "
309 "%.13la: *{%.13la, %d} vs. {%.13la, %d}\n",
310 f->name, sizeNames[k], err, (int)iErr,
311 ((double *)gIn)[j], ((double *)gOut_Ref)[j],
312 ((int *)gOut_Ref2)[j], test, q2[j]);
313 return -1;
314 }
315 }
316 }
317 }
318
319 if (0 == (i & 0x0fffffff))
320 {
321 if (gVerboseBruteForce)
322 {
323 vlog("base:%14" PRIu64 " step:%10" PRIu64
324 " bufferSize:%10d \n",
325 i, step, BUFFER_SIZE);
326 }
327 else
328 {
329 vlog(".");
330 }
331 fflush(stdout);
332 }
333 }
334
335 if (!gSkipCorrectnessTesting)
336 {
337 if (gWimpyMode)
338 vlog("Wimp pass");
339 else
340 vlog("passed");
341
342 vlog("\t{%8.2f, %" PRId64 "} @ {%a, %a}", maxError, maxError2,
343 maxErrorVal, maxErrorVal2);
344 }
345
346 vlog("\n");
347
348 return CL_SUCCESS;
349 }
350