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