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