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