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
BuildKernel(const char * name,int vectorSize,cl_kernel * k,cl_program * p,bool relaxedMode)27 int BuildKernel(const char *name, int vectorSize, cl_kernel *k, cl_program *p,
28 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 int",
34 sizeNames[vectorSize],
35 "* out, __global double",
36 sizeNames[vectorSize],
37 "* in )\n"
38 "{\n"
39 " size_t i = get_global_id(0);\n"
40 " out[i] = ",
41 name,
42 "( in[i] );\n"
43 "}\n" };
44
45 const char *c3[] = {
46 "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n",
47 "__kernel void math_kernel",
48 sizeNames[vectorSize],
49 "( __global int* out, __global double* in)\n"
50 "{\n"
51 " size_t i = get_global_id(0);\n"
52 " if( i + 1 < get_global_size(0) )\n"
53 " {\n"
54 " double3 f0 = vload3( 0, in + 3 * i );\n"
55 " int3 i0 = ",
56 name,
57 "( f0 );\n"
58 " vstore3( i0, 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 f0;\n"
66 " switch( parity )\n"
67 " {\n"
68 " case 1:\n"
69 " f0 = (double3)( in[3*i], NAN, NAN ); \n"
70 " break;\n"
71 " case 0:\n"
72 " f0 = (double3)( in[3*i], in[3*i+1], NAN ); \n"
73 " break;\n"
74 " }\n"
75 " int3 i0 = ",
76 name,
77 "( f0 );\n"
78 " switch( parity )\n"
79 " {\n"
80 " case 0:\n"
81 " out[3*i+1] = i0.y; \n"
82 " // fall through\n"
83 " case 1:\n"
84 " out[3*i] = i0.x; \n"
85 " break;\n"
86 " }\n"
87 " }\n"
88 "}\n"
89 };
90
91 const char **kern = c;
92 size_t kernSize = sizeof(c) / sizeof(c[0]);
93
94 if (sizeValues[vectorSize] == 3)
95 {
96 kern = c3;
97 kernSize = sizeof(c3) / sizeof(c3[0]);
98 }
99
100 char testName[32];
101 snprintf(testName, sizeof(testName) - 1, "math_kernel%s",
102 sizeNames[vectorSize]);
103
104 return MakeKernel(kern, (cl_uint)kernSize, testName, k, p, relaxedMode);
105 }
106
107 struct BuildKernelInfo2
108 {
109 cl_kernel *kernels;
110 Programs &programs;
111 const char *nameInCode;
112 bool relaxedMode; // Whether to build with -cl-fast-relaxed-math.
113 };
114
BuildKernelFn(cl_uint job_id,cl_uint thread_id UNUSED,void * p)115 cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
116 {
117 BuildKernelInfo2 *info = (BuildKernelInfo2 *)p;
118 cl_uint vectorSize = gMinVectorSizeIndex + job_id;
119 return BuildKernel(info->nameInCode, vectorSize, info->kernels + vectorSize,
120 &(info->programs[vectorSize]), info->relaxedMode);
121 }
122
123 } // anonymous namespace
124
TestFunc_Int_Double(const Func * f,MTdata d,bool relaxedMode)125 int TestFunc_Int_Double(const Func *f, MTdata d, bool relaxedMode)
126 {
127 int error;
128 Programs programs;
129 cl_kernel kernels[VECTOR_SIZE_COUNT];
130 int ftz = f->ftz || gForceFTZ;
131 uint64_t step = getTestStep(sizeof(cl_double), BUFFER_SIZE);
132 int scale =
133 (int)((1ULL << 32) / (16 * BUFFER_SIZE / sizeof(cl_double)) + 1);
134
135 logFunctionInfo(f->name, sizeof(cl_double), relaxedMode);
136
137 // This test is not using ThreadPool so we need to disable FTZ here
138 // for reference computations
139 FPU_mode_type oldMode;
140 DisableFTZ(&oldMode);
141
142 Force64BitFPUPrecision();
143
144 // Init the kernels
145 {
146 BuildKernelInfo2 build_info{ kernels, programs, f->nameInCode,
147 relaxedMode };
148 if ((error = ThreadPool_Do(BuildKernelFn,
149 gMaxVectorSizeIndex - gMinVectorSizeIndex,
150 &build_info)))
151 return error;
152 }
153
154 for (uint64_t i = 0; i < (1ULL << 32); i += step)
155 {
156 // Init input array
157 double *p = (double *)gIn;
158 if (gWimpyMode)
159 {
160 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
161 p[j] = DoubleFromUInt32((uint32_t)i + j * scale);
162 }
163 else
164 {
165 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
166 p[j] = DoubleFromUInt32((uint32_t)i + j);
167 }
168
169 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
170 BUFFER_SIZE, gIn, 0, NULL, NULL)))
171 {
172 vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
173 return error;
174 }
175
176 // write garbage into output arrays
177 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
178 {
179 uint32_t pattern = 0xffffdead;
180 memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
181 if ((error =
182 clEnqueueWriteBuffer(gQueue, gOutBuffer[j], CL_FALSE, 0,
183 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
184 {
185 vlog_error("\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
186 error, j);
187 goto exit;
188 }
189 }
190
191 // Run the kernels
192 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
193 {
194 size_t vectorSize = sizeValues[j] * sizeof(cl_double);
195 size_t localCount = (BUFFER_SIZE + vectorSize - 1)
196 / vectorSize; // BUFFER_SIZE / vectorSize rounded up
197 if ((error = clSetKernelArg(kernels[j], 0, sizeof(gOutBuffer[j]),
198 &gOutBuffer[j])))
199 {
200 LogBuildError(programs[j]);
201 goto exit;
202 }
203 if ((error = clSetKernelArg(kernels[j], 1, sizeof(gInBuffer),
204 &gInBuffer)))
205 {
206 LogBuildError(programs[j]);
207 goto exit;
208 }
209
210 if ((error =
211 clEnqueueNDRangeKernel(gQueue, kernels[j], 1, NULL,
212 &localCount, NULL, 0, NULL, NULL)))
213 {
214 vlog_error("FAILED -- could not execute kernel\n");
215 goto exit;
216 }
217 }
218
219 // Get that moving
220 if ((error = clFlush(gQueue))) vlog("clFlush failed\n");
221
222 // Calculate the correctly rounded reference result
223 int *r = (int *)gOut_Ref;
224 double *s = (double *)gIn;
225 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
226 r[j] = f->dfunc.i_f(s[j]);
227
228 // Read the data back
229 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
230 {
231 if ((error =
232 clEnqueueReadBuffer(gQueue, gOutBuffer[j], CL_TRUE, 0,
233 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
234 {
235 vlog_error("ReadArray failed %d\n", error);
236 goto exit;
237 }
238 }
239
240 if (gSkipCorrectnessTesting) break;
241
242 // Verify data
243 uint32_t *t = (uint32_t *)gOut_Ref;
244 for (size_t j = 0; j < BUFFER_SIZE / sizeof(cl_double); j++)
245 {
246 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
247 {
248 uint32_t *q = (uint32_t *)(gOut[k]);
249 // If we aren't getting the correctly rounded result
250 if (t[j] != q[j])
251 {
252 if ((ftz || relaxedMode) && IsDoubleSubnormal(s[j]))
253 {
254 unsigned int correct0 = f->dfunc.i_f(0.0);
255 unsigned int correct1 = f->dfunc.i_f(-0.0);
256 if (q[j] == correct0 || q[j] == correct1) continue;
257 }
258
259 uint32_t err = t[j] - q[j];
260 if (q[j] > t[j]) err = q[j] - t[j];
261 vlog_error(
262 "\nERROR: %sD%s: %d ulp error at %.13la: *%d vs. %d\n",
263 f->name, sizeNames[k], err, ((double *)gIn)[j], t[j],
264 q[j]);
265 error = -1;
266 goto exit;
267 }
268 }
269 }
270
271 if (0 == (i & 0x0fffffff))
272 {
273 if (gVerboseBruteForce)
274 {
275 vlog("base:%14" PRIu64 " step:%10" PRIu64
276 " bufferSize:%10d \n",
277 i, step, BUFFER_SIZE);
278 }
279 else
280 {
281 vlog(".");
282 }
283
284 fflush(stdout);
285 }
286 }
287
288 if (!gSkipCorrectnessTesting)
289 {
290 if (gWimpyMode)
291 vlog("Wimp pass");
292 else
293 vlog("passed");
294 }
295
296 vlog("\n");
297
298 exit:
299 RestoreFPState(&oldMode);
300 // Release
301 for (auto k = gMinVectorSizeIndex; k < gMaxVectorSizeIndex; k++)
302 {
303 clReleaseKernel(kernels[k]);
304 }
305
306 return error;
307 }
308