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