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 <climits>
22 #include <cstring>
23
BuildKernel(const char * name,int vectorSize,cl_uint kernel_count,cl_kernel * k,cl_program * p,bool relaxedMode)24 static int BuildKernel(const char *name, int vectorSize, cl_uint kernel_count,
25 cl_kernel *k, cl_program *p, bool relaxedMode)
26 {
27 const char *c[] = { "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
28 "__kernel void math_kernel",
29 sizeNames[vectorSize],
30 "( __global double",
31 sizeNames[vectorSize],
32 "* out, __global double",
33 sizeNames[vectorSize],
34 "* in1, __global int",
35 sizeNames[vectorSize],
36 "* in2 )\n"
37 "{\n"
38 " size_t i = get_global_id(0);\n"
39 " out[i] = ",
40 name,
41 "( in1[i], in2[i] );\n"
42 "}\n" };
43
44 const char *c3[] = {
45 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
46 "__kernel void math_kernel",
47 sizeNames[vectorSize],
48 "( __global double* out, __global double* in, __global int* in2)\n"
49 "{\n"
50 " size_t i = get_global_id(0);\n"
51 " if( i + 1 < get_global_size(0) )\n"
52 " {\n"
53 " double3 d0 = vload3( 0, in + 3 * i );\n"
54 " int3 i0 = vload3( 0, in2 + 3 * i );\n"
55 " d0 = ",
56 name,
57 "( d0, i0 );\n"
58 " vstore3( d0, 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 " double3 d0;\n"
66 " int3 i0;\n"
67 " switch( parity )\n"
68 " {\n"
69 " case 1:\n"
70 " d0 = (double3)( in[3*i], NAN, NAN ); \n"
71 " i0 = (int3)( in2[3*i], 0xdead, 0xdead ); \n"
72 " break;\n"
73 " case 0:\n"
74 " d0 = (double3)( in[3*i], in[3*i+1], NAN ); \n"
75 " i0 = (int3)( in2[3*i], in2[3*i+1], 0xdead ); \n"
76 " break;\n"
77 " }\n"
78 " d0 = ",
79 name,
80 "( d0, i0 );\n"
81 " switch( parity )\n"
82 " {\n"
83 " case 0:\n"
84 " out[3*i+1] = d0.y; \n"
85 " // fall through\n"
86 " case 1:\n"
87 " out[3*i] = d0.x; \n"
88 " break;\n"
89 " }\n"
90 " }\n"
91 "}\n"
92 };
93
94 const char **kern = c;
95 size_t kernSize = sizeof(c) / sizeof(c[0]);
96
97 if (sizeValues[vectorSize] == 3)
98 {
99 kern = c3;
100 kernSize = sizeof(c3) / sizeof(c3[0]);
101 }
102
103 char testName[32];
104 snprintf(testName, sizeof(testName) - 1, "math_kernel%s",
105 sizeNames[vectorSize]);
106
107 return MakeKernels(kern, (cl_uint)kernSize, testName, kernel_count, k, p,
108 relaxedMode);
109 }
110
111 typedef struct BuildKernelInfo
112 {
113 cl_uint offset; // the first vector size to build
114 cl_uint kernel_count;
115 cl_kernel **kernels;
116 cl_program *programs;
117 const char *nameInCode;
118 bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
119 } BuildKernelInfo;
120
BuildKernelFn(cl_uint job_id,cl_uint thread_id UNUSED,void * p)121 static cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
122 {
123 BuildKernelInfo *info = (BuildKernelInfo *)p;
124 cl_uint i = info->offset + job_id;
125 return BuildKernel(info->nameInCode, i, info->kernel_count,
126 info->kernels[i], info->programs + i, info->relaxedMode);
127 }
128
129 // Thread specific data for a worker thread
130 typedef struct ThreadInfo
131 {
132 cl_mem inBuf; // input buffer for the thread
133 cl_mem inBuf2; // input buffer for the thread
134 cl_mem outBuf[VECTOR_SIZE_COUNT]; // output buffers for the thread
135 float maxError; // max error value. Init to 0.
136 double
137 maxErrorValue; // position of the max error value (param 1). Init to 0.
138 cl_int maxErrorValue2; // position of the max error value (param 2). Init
139 // to 0.
140 MTdata d;
141 cl_command_queue tQueue; // per thread command queue to improve performance
142 } ThreadInfo;
143
144 typedef struct TestInfo
145 {
146 size_t subBufferSize; // Size of the sub-buffer in elements
147 const Func *f; // A pointer to the function info
148 cl_program programs[VECTOR_SIZE_COUNT]; // programs for various vector sizes
149 cl_kernel
150 *k[VECTOR_SIZE_COUNT]; // arrays of thread-specific kernels for each
151 // worker thread: k[vector_size][thread_id]
152 ThreadInfo *
153 tinfo; // An array of thread specific information for each worker thread
154 cl_uint threadCount; // Number of worker threads
155 cl_uint jobCount; // Number of jobs
156 cl_uint step; // step between each chunk and the next.
157 cl_uint scale; // stride between individual test values
158 float ulps; // max_allowed ulps
159 int ftz; // non-zero if running in flush to zero mode
160
161 // no special values
162 } TestInfo;
163
164 // A table of more difficult cases to get right
165 static const double specialValues[] = {
166 -NAN,
167 -INFINITY,
168 -DBL_MAX,
169 MAKE_HEX_DOUBLE(-0x1.0000000000001p64, -0x10000000000001LL, 12),
170 MAKE_HEX_DOUBLE(-0x1.0p64, -0x1LL, 64),
171 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp63, -0x1fffffffffffffLL, 11),
172 MAKE_HEX_DOUBLE(-0x1.0000000000001p63, -0x10000000000001LL, 11),
173 MAKE_HEX_DOUBLE(-0x1.0p63, -0x1LL, 63),
174 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp62, -0x1fffffffffffffLL, 10),
175 MAKE_HEX_DOUBLE(-0x1.000002p32, -0x1000002LL, 8),
176 MAKE_HEX_DOUBLE(-0x1.0p32, -0x1LL, 32),
177 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp31, -0x1fffffffffffffLL, -21),
178 MAKE_HEX_DOUBLE(-0x1.0000000000001p31, -0x10000000000001LL, -21),
179 MAKE_HEX_DOUBLE(-0x1.0p31, -0x1LL, 31),
180 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp30, -0x1fffffffffffffLL, -22),
181 -1000.0,
182 -100.0,
183 -4.0,
184 -3.5,
185 -3.0,
186 MAKE_HEX_DOUBLE(-0x1.8000000000001p1, -0x18000000000001LL, -51),
187 -2.5,
188 MAKE_HEX_DOUBLE(-0x1.7ffffffffffffp1, -0x17ffffffffffffLL, -51),
189 -2.0,
190 MAKE_HEX_DOUBLE(-0x1.8000000000001p0, -0x18000000000001LL, -52),
191 -1.5,
192 MAKE_HEX_DOUBLE(-0x1.7ffffffffffffp0, -0x17ffffffffffffLL, -52),
193 MAKE_HEX_DOUBLE(-0x1.0000000000001p0, -0x10000000000001LL, -52),
194 -1.0,
195 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp-1, -0x1fffffffffffffLL, -53),
196 MAKE_HEX_DOUBLE(-0x1.0000000000001p-1, -0x10000000000001LL, -53),
197 -0.5,
198 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp-2, -0x1fffffffffffffLL, -54),
199 MAKE_HEX_DOUBLE(-0x1.0000000000001p-2, -0x10000000000001LL, -54),
200 -0.25,
201 MAKE_HEX_DOUBLE(-0x1.fffffffffffffp-3, -0x1fffffffffffffLL, -55),
202 MAKE_HEX_DOUBLE(-0x1.0000000000001p-1022, -0x10000000000001LL, -1074),
203 -DBL_MIN,
204 MAKE_HEX_DOUBLE(-0x0.fffffffffffffp-1022, -0x0fffffffffffffLL, -1074),
205 MAKE_HEX_DOUBLE(-0x0.0000000000fffp-1022, -0x00000000000fffLL, -1074),
206 MAKE_HEX_DOUBLE(-0x0.00000000000fep-1022, -0x000000000000feLL, -1074),
207 MAKE_HEX_DOUBLE(-0x0.000000000000ep-1022, -0x0000000000000eLL, -1074),
208 MAKE_HEX_DOUBLE(-0x0.000000000000cp-1022, -0x0000000000000cLL, -1074),
209 MAKE_HEX_DOUBLE(-0x0.000000000000ap-1022, -0x0000000000000aLL, -1074),
210 MAKE_HEX_DOUBLE(-0x0.0000000000008p-1022, -0x00000000000008LL, -1074),
211 MAKE_HEX_DOUBLE(-0x0.0000000000007p-1022, -0x00000000000007LL, -1074),
212 MAKE_HEX_DOUBLE(-0x0.0000000000006p-1022, -0x00000000000006LL, -1074),
213 MAKE_HEX_DOUBLE(-0x0.0000000000005p-1022, -0x00000000000005LL, -1074),
214 MAKE_HEX_DOUBLE(-0x0.0000000000004p-1022, -0x00000000000004LL, -1074),
215 MAKE_HEX_DOUBLE(-0x0.0000000000003p-1022, -0x00000000000003LL, -1074),
216 MAKE_HEX_DOUBLE(-0x0.0000000000002p-1022, -0x00000000000002LL, -1074),
217 MAKE_HEX_DOUBLE(-0x0.0000000000001p-1022, -0x00000000000001LL, -1074),
218 -0.0,
219
220 +NAN,
221 +INFINITY,
222 +DBL_MAX,
223 MAKE_HEX_DOUBLE(+0x1.0000000000001p64, +0x10000000000001LL, 12),
224 MAKE_HEX_DOUBLE(+0x1.0p64, +0x1LL, 64),
225 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp63, +0x1fffffffffffffLL, 11),
226 MAKE_HEX_DOUBLE(+0x1.0000000000001p63, +0x10000000000001LL, 11),
227 MAKE_HEX_DOUBLE(+0x1.0p63, +0x1LL, 63),
228 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp62, +0x1fffffffffffffLL, 10),
229 MAKE_HEX_DOUBLE(+0x1.000002p32, +0x1000002LL, 8),
230 MAKE_HEX_DOUBLE(+0x1.0p32, +0x1LL, 32),
231 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp31, +0x1fffffffffffffLL, -21),
232 MAKE_HEX_DOUBLE(+0x1.0000000000001p31, +0x10000000000001LL, -21),
233 MAKE_HEX_DOUBLE(+0x1.0p31, +0x1LL, 31),
234 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp30, +0x1fffffffffffffLL, -22),
235 +1000.0,
236 +100.0,
237 +4.0,
238 +3.5,
239 +3.0,
240 MAKE_HEX_DOUBLE(+0x1.8000000000001p1, +0x18000000000001LL, -51),
241 +2.5,
242 MAKE_HEX_DOUBLE(+0x1.7ffffffffffffp1, +0x17ffffffffffffLL, -51),
243 +2.0,
244 MAKE_HEX_DOUBLE(+0x1.8000000000001p0, +0x18000000000001LL, -52),
245 +1.5,
246 MAKE_HEX_DOUBLE(+0x1.7ffffffffffffp0, +0x17ffffffffffffLL, -52),
247 MAKE_HEX_DOUBLE(-0x1.0000000000001p0, -0x10000000000001LL, -52),
248 +1.0,
249 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp-1, +0x1fffffffffffffLL, -53),
250 MAKE_HEX_DOUBLE(+0x1.0000000000001p-1, +0x10000000000001LL, -53),
251 +0.5,
252 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp-2, +0x1fffffffffffffLL, -54),
253 MAKE_HEX_DOUBLE(+0x1.0000000000001p-2, +0x10000000000001LL, -54),
254 +0.25,
255 MAKE_HEX_DOUBLE(+0x1.fffffffffffffp-3, +0x1fffffffffffffLL, -55),
256 MAKE_HEX_DOUBLE(+0x1.0000000000001p-1022, +0x10000000000001LL, -1074),
257 +DBL_MIN,
258 MAKE_HEX_DOUBLE(+0x0.fffffffffffffp-1022, +0x0fffffffffffffLL, -1074),
259 MAKE_HEX_DOUBLE(+0x0.0000000000fffp-1022, +0x00000000000fffLL, -1074),
260 MAKE_HEX_DOUBLE(+0x0.00000000000fep-1022, +0x000000000000feLL, -1074),
261 MAKE_HEX_DOUBLE(+0x0.000000000000ep-1022, +0x0000000000000eLL, -1074),
262 MAKE_HEX_DOUBLE(+0x0.000000000000cp-1022, +0x0000000000000cLL, -1074),
263 MAKE_HEX_DOUBLE(+0x0.000000000000ap-1022, +0x0000000000000aLL, -1074),
264 MAKE_HEX_DOUBLE(+0x0.0000000000008p-1022, +0x00000000000008LL, -1074),
265 MAKE_HEX_DOUBLE(+0x0.0000000000007p-1022, +0x00000000000007LL, -1074),
266 MAKE_HEX_DOUBLE(+0x0.0000000000006p-1022, +0x00000000000006LL, -1074),
267 MAKE_HEX_DOUBLE(+0x0.0000000000005p-1022, +0x00000000000005LL, -1074),
268 MAKE_HEX_DOUBLE(+0x0.0000000000004p-1022, +0x00000000000004LL, -1074),
269 MAKE_HEX_DOUBLE(+0x0.0000000000003p-1022, +0x00000000000003LL, -1074),
270 MAKE_HEX_DOUBLE(+0x0.0000000000002p-1022, +0x00000000000002LL, -1074),
271 MAKE_HEX_DOUBLE(+0x0.0000000000001p-1022, +0x00000000000001LL, -1074),
272 +0.0,
273 };
274
275 static size_t specialValuesCount =
276 sizeof(specialValues) / sizeof(specialValues[0]);
277
278 static const int specialValuesInt[] = {
279 0, 1, 2, 3, 1022, 1023, 1024, INT_MIN,
280 INT_MAX, -1, -2, -3, -1022, -1023, -11024, -INT_MAX,
281 };
282 static constexpr size_t specialValuesIntCount =
283 sizeof(specialValuesInt) / sizeof(specialValuesInt[0]);
284
285 static cl_int Test(cl_uint job_id, cl_uint thread_id, void *data);
286
TestFunc_Double_Double_Int(const Func * f,MTdata d,bool relaxedMode)287 int TestFunc_Double_Double_Int(const Func *f, MTdata d, bool relaxedMode)
288 {
289 TestInfo test_info;
290 cl_int error;
291 float maxError = 0.0f;
292 double maxErrorVal = 0.0;
293 cl_int maxErrorVal2 = 0;
294
295 logFunctionInfo(f->name, sizeof(cl_double), relaxedMode);
296
297 // Init test_info
298 memset(&test_info, 0, sizeof(test_info));
299 test_info.threadCount = GetThreadCount();
300 test_info.subBufferSize = BUFFER_SIZE
301 / (sizeof(cl_double) * RoundUpToNextPowerOfTwo(test_info.threadCount));
302 test_info.scale = getTestScale(sizeof(cl_double));
303
304 test_info.step = (cl_uint)test_info.subBufferSize * test_info.scale;
305 if (test_info.step / test_info.subBufferSize != test_info.scale)
306 {
307 // there was overflow
308 test_info.jobCount = 1;
309 }
310 else
311 {
312 test_info.jobCount = (cl_uint)((1ULL << 32) / test_info.step);
313 }
314
315 test_info.f = f;
316 test_info.ulps = f->double_ulps;
317 test_info.ftz = f->ftz || gForceFTZ;
318
319 // cl_kernels aren't thread safe, so we make one for each vector size for
320 // every thread
321 for (auto i = gMinVectorSizeIndex; i < gMaxVectorSizeIndex; i++)
322 {
323 size_t array_size = test_info.threadCount * sizeof(cl_kernel);
324 test_info.k[i] = (cl_kernel *)malloc(array_size);
325 if (NULL == test_info.k[i])
326 {
327 vlog_error("Error: Unable to allocate storage for kernels!\n");
328 error = CL_OUT_OF_HOST_MEMORY;
329 goto exit;
330 }
331 memset(test_info.k[i], 0, array_size);
332 }
333 test_info.tinfo =
334 (ThreadInfo *)malloc(test_info.threadCount * sizeof(*test_info.tinfo));
335 if (NULL == test_info.tinfo)
336 {
337 vlog_error(
338 "Error: Unable to allocate storage for thread specific data.\n");
339 error = CL_OUT_OF_HOST_MEMORY;
340 goto exit;
341 }
342 memset(test_info.tinfo, 0,
343 test_info.threadCount * sizeof(*test_info.tinfo));
344 for (cl_uint i = 0; i < test_info.threadCount; i++)
345 {
346 cl_buffer_region region = {
347 i * test_info.subBufferSize * sizeof(cl_double),
348 test_info.subBufferSize * sizeof(cl_double)
349 };
350 test_info.tinfo[i].inBuf =
351 clCreateSubBuffer(gInBuffer, CL_MEM_READ_ONLY,
352 CL_BUFFER_CREATE_TYPE_REGION, ®ion, &error);
353 if (error || NULL == test_info.tinfo[i].inBuf)
354 {
355 vlog_error("Error: Unable to create sub-buffer of gInBuffer for "
356 "region {%zd, %zd}\n",
357 region.origin, region.size);
358 goto exit;
359 }
360 cl_buffer_region region2 = { i * test_info.subBufferSize
361 * sizeof(cl_int),
362 test_info.subBufferSize * sizeof(cl_int) };
363 test_info.tinfo[i].inBuf2 =
364 clCreateSubBuffer(gInBuffer2, CL_MEM_READ_ONLY,
365 CL_BUFFER_CREATE_TYPE_REGION, ®ion2, &error);
366 if (error || NULL == test_info.tinfo[i].inBuf2)
367 {
368 vlog_error("Error: Unable to create sub-buffer of gInBuffer2 for "
369 "region {%zd, %zd}\n",
370 region.origin, region.size);
371 goto exit;
372 }
373
374 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
375 {
376 test_info.tinfo[i].outBuf[j] = clCreateSubBuffer(
377 gOutBuffer[j], CL_MEM_WRITE_ONLY, CL_BUFFER_CREATE_TYPE_REGION,
378 ®ion, &error);
379 if (error || NULL == test_info.tinfo[i].outBuf[j])
380 {
381 vlog_error("Error: Unable to create sub-buffer of "
382 "gOutBuffer[%d] for region {%zd, %zd}\n",
383 (int)j, region.origin, region.size);
384 goto exit;
385 }
386 }
387 test_info.tinfo[i].tQueue =
388 clCreateCommandQueue(gContext, gDevice, 0, &error);
389 if (NULL == test_info.tinfo[i].tQueue || error)
390 {
391 vlog_error("clCreateCommandQueue failed. (%d)\n", error);
392 goto exit;
393 }
394
395 test_info.tinfo[i].d = init_genrand(genrand_int32(d));
396 }
397
398 // Init the kernels
399 {
400 BuildKernelInfo build_info = {
401 gMinVectorSizeIndex, test_info.threadCount, test_info.k,
402 test_info.programs, f->nameInCode, relaxedMode
403 };
404 if ((error = ThreadPool_Do(BuildKernelFn,
405 gMaxVectorSizeIndex - gMinVectorSizeIndex,
406 &build_info)))
407 goto exit;
408 }
409
410 // Run the kernels
411 if (!gSkipCorrectnessTesting)
412 {
413 error = ThreadPool_Do(Test, test_info.jobCount, &test_info);
414
415 // Accumulate the arithmetic errors
416 for (cl_uint i = 0; i < test_info.threadCount; i++)
417 {
418 if (test_info.tinfo[i].maxError > maxError)
419 {
420 maxError = test_info.tinfo[i].maxError;
421 maxErrorVal = test_info.tinfo[i].maxErrorValue;
422 maxErrorVal2 = test_info.tinfo[i].maxErrorValue2;
423 }
424 }
425
426 if (error) goto exit;
427
428 if (gWimpyMode)
429 vlog("Wimp pass");
430 else
431 vlog("passed");
432
433 vlog("\t%8.2f @ {%a, %d}", maxError, maxErrorVal, maxErrorVal2);
434 }
435
436 vlog("\n");
437
438 exit:
439 // Release
440 for (auto i = gMinVectorSizeIndex; i < gMaxVectorSizeIndex; i++)
441 {
442 clReleaseProgram(test_info.programs[i]);
443 if (test_info.k[i])
444 {
445 for (cl_uint j = 0; j < test_info.threadCount; j++)
446 clReleaseKernel(test_info.k[i][j]);
447
448 free(test_info.k[i]);
449 }
450 }
451 if (test_info.tinfo)
452 {
453 for (cl_uint i = 0; i < test_info.threadCount; i++)
454 {
455 free_mtdata(test_info.tinfo[i].d);
456 clReleaseMemObject(test_info.tinfo[i].inBuf);
457 clReleaseMemObject(test_info.tinfo[i].inBuf2);
458 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
459 clReleaseMemObject(test_info.tinfo[i].outBuf[j]);
460 clReleaseCommandQueue(test_info.tinfo[i].tQueue);
461 }
462
463 free(test_info.tinfo);
464 }
465
466 return error;
467 }
468
Test(cl_uint job_id,cl_uint thread_id,void * data)469 static cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
470 {
471 const TestInfo *job = (const TestInfo *)data;
472 size_t buffer_elements = job->subBufferSize;
473 size_t buffer_size = buffer_elements * sizeof(cl_double);
474 cl_uint base = job_id * (cl_uint)job->step;
475 ThreadInfo *tinfo = job->tinfo + thread_id;
476 float ulps = job->ulps;
477 dptr func = job->f->dfunc;
478 int ftz = job->ftz;
479 MTdata d = tinfo->d;
480 cl_int error;
481 const char *name = job->f->name;
482 cl_ulong *t;
483 cl_double *r;
484 cl_double *s;
485 cl_int *s2;
486
487 Force64BitFPUPrecision();
488
489 // start the map of the output arrays
490 cl_event e[VECTOR_SIZE_COUNT];
491 cl_ulong *out[VECTOR_SIZE_COUNT];
492 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
493 {
494 out[j] = (cl_ulong *)clEnqueueMapBuffer(
495 tinfo->tQueue, tinfo->outBuf[j], CL_FALSE, CL_MAP_WRITE, 0,
496 buffer_size, 0, NULL, e + j, &error);
497 if (error || NULL == out[j])
498 {
499 vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
500 error);
501 return error;
502 }
503 }
504
505 // Get that moving
506 if ((error = clFlush(tinfo->tQueue))) vlog("clFlush failed\n");
507
508 // Init input array
509 cl_ulong *p = (cl_ulong *)gIn + thread_id * buffer_elements;
510 cl_int *p2 = (cl_int *)gIn2 + thread_id * buffer_elements;
511 size_t idx = 0;
512 int totalSpecialValueCount = specialValuesCount * specialValuesIntCount;
513 int lastSpecialJobIndex = (totalSpecialValueCount - 1) / buffer_elements;
514
515 if (job_id <= (cl_uint)lastSpecialJobIndex)
516 { // test edge cases
517 cl_double *fp = (cl_double *)p;
518 cl_int *ip2 = (cl_int *)p2;
519 uint32_t x, y;
520
521 x = (job_id * buffer_elements) % specialValuesCount;
522 y = (job_id * buffer_elements) / specialValuesCount;
523
524 for (; idx < buffer_elements; idx++)
525 {
526 fp[idx] = specialValues[x];
527 ip2[idx] = specialValuesInt[y];
528 if (++x >= specialValuesCount)
529 {
530 x = 0;
531 y++;
532 if (y >= specialValuesIntCount) break;
533 }
534 }
535 }
536
537 // Init any remaining values.
538 for (; idx < buffer_elements; idx++)
539 {
540 p[idx] = DoubleFromUInt32(genrand_int32(d));
541 p2[idx] = genrand_int32(d);
542 }
543
544 if ((error = clEnqueueWriteBuffer(tinfo->tQueue, tinfo->inBuf, CL_FALSE, 0,
545 buffer_size, p, 0, NULL, NULL)))
546 {
547 vlog_error("Error: clEnqueueWriteBuffer failed! err: %d\n", error);
548 goto exit;
549 }
550
551 if ((error = clEnqueueWriteBuffer(tinfo->tQueue, tinfo->inBuf2, CL_FALSE, 0,
552 buffer_size / 2, p2, 0, NULL, NULL)))
553 {
554 vlog_error("Error: clEnqueueWriteBuffer failed! err: %d\n", error);
555 goto exit;
556 }
557
558 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
559 {
560 // Wait for the map to finish
561 if ((error = clWaitForEvents(1, e + j)))
562 {
563 vlog_error("Error: clWaitForEvents failed! err: %d\n", error);
564 goto exit;
565 }
566 if ((error = clReleaseEvent(e[j])))
567 {
568 vlog_error("Error: clReleaseEvent failed! err: %d\n", error);
569 goto exit;
570 }
571
572 // Fill the result buffer with garbage, so that old results don't carry
573 // over
574 uint32_t pattern = 0xffffdead;
575 memset_pattern4(out[j], &pattern, buffer_size);
576 if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
577 out[j], 0, NULL, NULL)))
578 {
579 vlog_error("Error: clEnqueueMapBuffer failed! err: %d\n", error);
580 goto exit;
581 }
582
583 // run the kernel
584 size_t vectorCount =
585 (buffer_elements + sizeValues[j] - 1) / sizeValues[j];
586 cl_kernel kernel = job->k[j][thread_id]; // each worker thread has its
587 // own copy of the cl_kernel
588 cl_program program = job->programs[j];
589
590 if ((error = clSetKernelArg(kernel, 0, sizeof(tinfo->outBuf[j]),
591 &tinfo->outBuf[j])))
592 {
593 LogBuildError(program);
594 return error;
595 }
596 if ((error = clSetKernelArg(kernel, 1, sizeof(tinfo->inBuf),
597 &tinfo->inBuf)))
598 {
599 LogBuildError(program);
600 return error;
601 }
602 if ((error = clSetKernelArg(kernel, 2, sizeof(tinfo->inBuf2),
603 &tinfo->inBuf2)))
604 {
605 LogBuildError(program);
606 return error;
607 }
608
609 if ((error = clEnqueueNDRangeKernel(tinfo->tQueue, kernel, 1, NULL,
610 &vectorCount, NULL, 0, NULL, NULL)))
611 {
612 vlog_error("FAILED -- could not execute kernel\n");
613 goto exit;
614 }
615 }
616
617 // Get that moving
618 if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 2 failed\n");
619
620 if (gSkipCorrectnessTesting) return CL_SUCCESS;
621
622 // Calculate the correctly rounded reference result
623 r = (cl_double *)gOut_Ref + thread_id * buffer_elements;
624 s = (cl_double *)gIn + thread_id * buffer_elements;
625 s2 = (cl_int *)gIn2 + thread_id * buffer_elements;
626 for (size_t j = 0; j < buffer_elements; j++)
627 r[j] = (cl_double)func.f_fi(s[j], s2[j]);
628
629 // Read the data back -- no need to wait for the first N-1 buffers but wait
630 // for the last buffer. This is an in order queue.
631 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
632 {
633 cl_bool blocking = (j + 1 < gMaxVectorSizeIndex) ? CL_FALSE : CL_TRUE;
634 out[j] = (cl_ulong *)clEnqueueMapBuffer(
635 tinfo->tQueue, tinfo->outBuf[j], blocking, CL_MAP_READ, 0,
636 buffer_size, 0, NULL, NULL, &error);
637 if (error || NULL == out[j])
638 {
639 vlog_error("Error: clEnqueueMapBuffer %d failed! err: %d\n", j,
640 error);
641 goto exit;
642 }
643 }
644
645 // Verify data
646 t = (cl_ulong *)r;
647 for (size_t j = 0; j < buffer_elements; j++)
648 {
649 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
650 {
651 cl_ulong *q = out[k];
652
653 // If we aren't getting the correctly rounded result
654 if (t[j] != q[j])
655 {
656 cl_double test = ((cl_double *)q)[j];
657 long double correct = func.f_fi(s[j], s2[j]);
658 float err = Bruteforce_Ulp_Error_Double(test, correct);
659 int fail = !(fabsf(err) <= ulps);
660
661 if (fail && ftz)
662 {
663 // retry per section 6.5.3.2
664 if (IsDoubleResultSubnormal(correct, ulps))
665 {
666 fail = fail && (test != 0.0f);
667 if (!fail) err = 0.0f;
668 }
669
670 // retry per section 6.5.3.3
671 if (IsDoubleSubnormal(s[j]))
672 {
673 long double correct2 = func.f_fi(0.0, s2[j]);
674 long double correct3 = func.f_fi(-0.0, s2[j]);
675 float err2 =
676 Bruteforce_Ulp_Error_Double(test, correct2);
677 float err3 =
678 Bruteforce_Ulp_Error_Double(test, correct3);
679 fail = fail
680 && ((!(fabsf(err2) <= ulps))
681 && (!(fabsf(err3) <= ulps)));
682 if (fabsf(err2) < fabsf(err)) err = err2;
683 if (fabsf(err3) < fabsf(err)) err = err3;
684
685 // retry per section 6.5.3.4
686 if (IsDoubleResultSubnormal(correct2, ulps)
687 || IsDoubleResultSubnormal(correct3, ulps))
688 {
689 fail = fail && (test != 0.0f);
690 if (!fail) err = 0.0f;
691 }
692 }
693 }
694
695 if (fabsf(err) > tinfo->maxError)
696 {
697 tinfo->maxError = fabsf(err);
698 tinfo->maxErrorValue = s[j];
699 tinfo->maxErrorValue2 = s2[j];
700 }
701 if (fail)
702 {
703 vlog_error("\nERROR: %s%s: %f ulp error at {%.13la, %d}: "
704 "*%.13la vs. %.13la\n",
705 name, sizeNames[k], err, s[j], s2[j], r[j],
706 test);
707 error = -1;
708 goto exit;
709 }
710 }
711 }
712 }
713
714 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
715 {
716 if ((error = clEnqueueUnmapMemObject(tinfo->tQueue, tinfo->outBuf[j],
717 out[j], 0, NULL, NULL)))
718 {
719 vlog_error("Error: clEnqueueUnmapMemObject %d failed 2! err: %d\n",
720 j, error);
721 return error;
722 }
723 }
724
725 if ((error = clFlush(tinfo->tQueue))) vlog("clFlush 3 failed\n");
726
727
728 if (0 == (base & 0x0fffffff))
729 {
730 if (gVerboseBruteForce)
731 {
732 vlog("base:%14u step:%10u scale:%10u buf_elements:%10zd ulps:%5.3f "
733 "ThreadCount:%2u\n",
734 base, job->step, job->scale, buffer_elements, job->ulps,
735 job->threadCount);
736 }
737 else
738 {
739 vlog(".");
740 }
741 fflush(stdout);
742 }
743
744 exit:
745 return error;
746 }
747