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