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::Float,
34 ParameterType::Int, ParameterType::Float,
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_FloatI_Float(const Func * f,MTdata d,bool relaxedMode)48 int TestFunc_FloatI_Float(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 || 0 == (CL_FP_DENORM & gFloatCapabilities);
57 float maxErrorVal = 0.0f;
58 float maxErrorVal2 = 0.0f;
59 uint64_t step = getTestStep(sizeof(float), BUFFER_SIZE);
60 int scale = (int)((1ULL << 32) / (16 * BUFFER_SIZE / sizeof(float)) + 1);
61 cl_ulong maxiError;
62
63 logFunctionInfo(f->name, sizeof(cl_float), relaxedMode);
64
65 float float_ulps;
66 if (gIsEmbedded)
67 float_ulps = f->float_embedded_ulps;
68 else
69 float_ulps = f->float_ulps;
70
71 maxiError = float_ulps == INFINITY ? CL_ULONG_MAX : 0;
72
73 // Init the kernels
74 BuildKernelInfo build_info{ 1, kernels, programs, f->nameInCode,
75 relaxedMode };
76 if ((error = ThreadPool_Do(BuildKernelFn,
77 gMaxVectorSizeIndex - gMinVectorSizeIndex,
78 &build_info)))
79 return error;
80
81 for (uint64_t i = 0; i < (1ULL << 32); i += step)
82 {
83 // Init input array
84 uint32_t *p = (uint32_t *)gIn;
85 if (gWimpyMode)
86 {
87 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
88 p[j] = (uint32_t)i + j * scale;
89 }
90 else
91 {
92 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
93 p[j] = (uint32_t)i + j;
94 }
95 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
96 BUFFER_SIZE, gIn, 0, NULL, NULL)))
97 {
98 vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
99 return error;
100 }
101
102 // Write garbage into output arrays
103 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
104 {
105 uint32_t pattern = 0xffffdead;
106 if (gHostFill)
107 {
108 memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
109 if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer[j],
110 CL_FALSE, 0, BUFFER_SIZE,
111 gOut[j], 0, NULL, NULL)))
112 {
113 vlog_error(
114 "\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
115 error, j);
116 return error;
117 }
118
119 memset_pattern4(gOut2[j], &pattern, BUFFER_SIZE);
120 if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer2[j],
121 CL_FALSE, 0, BUFFER_SIZE,
122 gOut2[j], 0, NULL, NULL)))
123 {
124 vlog_error(
125 "\n*** Error %d in clEnqueueWriteBuffer2b(%d) ***\n",
126 error, j);
127 return error;
128 }
129 }
130 else
131 {
132 if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer[j],
133 &pattern, sizeof(pattern), 0,
134 BUFFER_SIZE, 0, NULL, NULL)))
135 {
136 vlog_error("Error: clEnqueueFillBuffer 1 failed! err: %d\n",
137 error);
138 return error;
139 }
140
141 if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer2[j],
142 &pattern, sizeof(pattern), 0,
143 BUFFER_SIZE, 0, NULL, NULL)))
144 {
145 vlog_error("Error: clEnqueueFillBuffer 2 failed! err: %d\n",
146 error);
147 return error;
148 }
149 }
150 }
151
152 // Run the kernels
153 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
154 {
155 size_t vectorSize = sizeValues[j] * sizeof(cl_float);
156 size_t localCount = (BUFFER_SIZE + vectorSize - 1) / vectorSize;
157 if ((error = clSetKernelArg(kernels[j][thread_id], 0,
158 sizeof(gOutBuffer[j]), &gOutBuffer[j])))
159 {
160 LogBuildError(programs[j]);
161 return error;
162 }
163 if ((error =
164 clSetKernelArg(kernels[j][thread_id], 1,
165 sizeof(gOutBuffer2[j]), &gOutBuffer2[j])))
166 {
167 LogBuildError(programs[j]);
168 return error;
169 }
170 if ((error = clSetKernelArg(kernels[j][thread_id], 2,
171 sizeof(gInBuffer), &gInBuffer)))
172 {
173 LogBuildError(programs[j]);
174 return error;
175 }
176
177 if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j][thread_id],
178 1, NULL, &localCount, NULL, 0,
179 NULL, NULL)))
180 {
181 vlog_error("FAILED -- could not execute kernel\n");
182 return error;
183 }
184 }
185
186 // Get that moving
187 if ((error = clFlush(gQueue))) vlog("clFlush failed\n");
188
189 // Calculate the correctly rounded reference result
190 float *r = (float *)gOut_Ref;
191 int *r2 = (int *)gOut_Ref2;
192 float *s = (float *)gIn;
193 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
194 r[j] = (float)f->func.f_fpI(s[j], r2 + j);
195
196 // Read the data back
197 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
198 {
199 if ((error =
200 clEnqueueReadBuffer(gQueue, gOutBuffer[j], CL_TRUE, 0,
201 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
202 {
203 vlog_error("ReadArray failed %d\n", error);
204 return error;
205 }
206 if ((error =
207 clEnqueueReadBuffer(gQueue, gOutBuffer2[j], CL_TRUE, 0,
208 BUFFER_SIZE, gOut2[j], 0, NULL, NULL)))
209 {
210 vlog_error("ReadArray2 failed %d\n", error);
211 return error;
212 }
213 }
214
215 if (gSkipCorrectnessTesting) break;
216
217 // Verify data
218 uint32_t *t = (uint32_t *)gOut_Ref;
219 int32_t *t2 = (int32_t *)gOut_Ref2;
220 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
221 {
222 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
223 {
224 uint32_t *q = (uint32_t *)(gOut[k]);
225 int32_t *q2 = (int32_t *)(gOut2[k]);
226
227 // If we aren't getting the correctly rounded result
228 if (t[j] != q[j] || t2[j] != q2[j])
229 {
230 float test = ((float *)q)[j];
231 int correct2 = INT_MIN;
232 double correct = f->func.f_fpI(s[j], &correct2);
233 float err = Ulp_Error(test, correct);
234 cl_long iErr = (int64_t)q2[j] - (int64_t)correct2;
235 int fail = !(fabsf(err) <= float_ulps
236 && abs_cl_long(iErr) <= maxiError);
237 if (ftz || relaxedMode)
238 {
239 // retry per section 6.5.3.2
240 if (IsFloatResultSubnormal(correct, float_ulps))
241 {
242 fail = fail && !(test == 0.0f && iErr == 0);
243 if (!fail) err = 0.0f;
244 }
245
246 // retry per section 6.5.3.3
247 if (IsFloatSubnormal(s[j]))
248 {
249 int correct5, correct6;
250 double correct3 = f->func.f_fpI(0.0, &correct5);
251 double correct4 = f->func.f_fpI(-0.0, &correct6);
252 float err2 = Ulp_Error(test, correct3);
253 float err3 = Ulp_Error(test, correct4);
254 cl_long iErr2 =
255 (long long)q2[j] - (long long)correct5;
256 cl_long iErr3 =
257 (long long)q2[j] - (long long)correct6;
258
259 // Did +0 work?
260 if (fabsf(err2) <= float_ulps
261 && abs_cl_long(iErr2) <= maxiError)
262 {
263 err = err2;
264 iErr = iErr2;
265 fail = 0;
266 }
267 // Did -0 work?
268 else if (fabsf(err3) <= float_ulps
269 && abs_cl_long(iErr3) <= maxiError)
270 {
271 err = err3;
272 iErr = iErr3;
273 fail = 0;
274 }
275
276 // retry per section 6.5.3.4
277 if (fail
278 && (IsFloatResultSubnormal(correct2, float_ulps)
279 || IsFloatResultSubnormal(correct3,
280 float_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: %s%s: {%f, %d} ulp error at %a: "
309 "*{%a, %d} vs. {%a, %d}\n",
310 f->name, sizeNames[k], err, (int)iErr,
311 ((float *)gIn)[j], ((float *)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