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