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 <cstring>
23
24 namespace {
25
BuildKernelFn(cl_uint job_id,cl_uint thread_id UNUSED,void * p)26 cl_int BuildKernelFn(cl_uint job_id, cl_uint thread_id UNUSED, void *p)
27 {
28 BuildKernelInfo &info = *(BuildKernelInfo *)p;
29 auto generator = [](const std::string &kernel_name, const char *builtin,
30 cl_uint vector_size_index) {
31 return GetTernaryKernel(kernel_name, builtin, ParameterType::Float,
32 ParameterType::Float, ParameterType::Float,
33 ParameterType::Float, vector_size_index);
34 };
35 return BuildKernels(info, job_id, generator);
36 }
37
38 } // anonymous namespace
39
TestFunc_mad_Float(const Func * f,MTdata d,bool relaxedMode)40 int TestFunc_mad_Float(const Func *f, MTdata d, bool relaxedMode)
41 {
42 int error;
43
44 logFunctionInfo(f->name, sizeof(cl_float), relaxedMode);
45
46 Programs programs;
47 const unsigned thread_id = 0; // Test is currently not multithreaded.
48 KernelMatrix kernels;
49 float maxError = 0.0f;
50 float maxErrorVal = 0.0f;
51 float maxErrorVal2 = 0.0f;
52 float maxErrorVal3 = 0.0f;
53 uint64_t step = getTestStep(sizeof(float), BUFFER_SIZE);
54
55 // Init the kernels
56 BuildKernelInfo build_info{ 1, kernels, programs, f->nameInCode,
57 relaxedMode };
58 if ((error = ThreadPool_Do(BuildKernelFn,
59 gMaxVectorSizeIndex - gMinVectorSizeIndex,
60 &build_info)))
61 return error;
62
63 for (uint64_t i = 0; i < (1ULL << 32); i += step)
64 {
65 // Init input array
66 cl_uint *p = (cl_uint *)gIn;
67 cl_uint *p2 = (cl_uint *)gIn2;
68 cl_uint *p3 = (cl_uint *)gIn3;
69 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
70 {
71 p[j] = genrand_int32(d);
72 p2[j] = genrand_int32(d);
73 p3[j] = genrand_int32(d);
74 }
75
76 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0,
77 BUFFER_SIZE, gIn, 0, NULL, NULL)))
78 {
79 vlog_error("\n*** Error %d in clEnqueueWriteBuffer ***\n", error);
80 return error;
81 }
82
83 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer2, CL_FALSE, 0,
84 BUFFER_SIZE, gIn2, 0, NULL, NULL)))
85 {
86 vlog_error("\n*** Error %d in clEnqueueWriteBuffer2 ***\n", error);
87 return error;
88 }
89
90 if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer3, CL_FALSE, 0,
91 BUFFER_SIZE, gIn3, 0, NULL, NULL)))
92 {
93 vlog_error("\n*** Error %d in clEnqueueWriteBuffer3 ***\n", error);
94 return error;
95 }
96
97 // Write garbage into output arrays
98 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
99 {
100 uint32_t pattern = 0xffffdead;
101 if (gHostFill)
102 {
103 memset_pattern4(gOut[j], &pattern, BUFFER_SIZE);
104 if ((error = clEnqueueWriteBuffer(gQueue, gOutBuffer[j],
105 CL_FALSE, 0, BUFFER_SIZE,
106 gOut[j], 0, NULL, NULL)))
107 {
108 vlog_error(
109 "\n*** Error %d in clEnqueueWriteBuffer2(%d) ***\n",
110 error, j);
111 return error;
112 }
113 }
114 else
115 {
116 if ((error = clEnqueueFillBuffer(gQueue, gOutBuffer[j],
117 &pattern, sizeof(pattern), 0,
118 BUFFER_SIZE, 0, NULL, NULL)))
119 {
120 vlog_error("Error: clEnqueueFillBuffer failed! err: %d\n",
121 error);
122 return error;
123 }
124 }
125 }
126
127 // Run the kernels
128 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
129 {
130 size_t vectorSize = sizeof(cl_float) * sizeValues[j];
131 size_t localCount = (BUFFER_SIZE + vectorSize - 1)
132 / vectorSize; // BUFFER_SIZE / vectorSize rounded up
133 if ((error = clSetKernelArg(kernels[j][thread_id], 0,
134 sizeof(gOutBuffer[j]), &gOutBuffer[j])))
135 {
136 LogBuildError(programs[j]);
137 return error;
138 }
139 if ((error = clSetKernelArg(kernels[j][thread_id], 1,
140 sizeof(gInBuffer), &gInBuffer)))
141 {
142 LogBuildError(programs[j]);
143 return error;
144 }
145 if ((error = clSetKernelArg(kernels[j][thread_id], 2,
146 sizeof(gInBuffer2), &gInBuffer2)))
147 {
148 LogBuildError(programs[j]);
149 return error;
150 }
151 if ((error = clSetKernelArg(kernels[j][thread_id], 3,
152 sizeof(gInBuffer3), &gInBuffer3)))
153 {
154 LogBuildError(programs[j]);
155 return error;
156 }
157
158 if ((error = clEnqueueNDRangeKernel(gQueue, kernels[j][thread_id],
159 1, NULL, &localCount, NULL, 0,
160 NULL, NULL)))
161 {
162 vlog_error("FAILED -- could not execute kernel\n");
163 return error;
164 }
165 }
166
167 // Get that moving
168 if ((error = clFlush(gQueue))) vlog("clFlush failed\n");
169
170 // Calculate the correctly rounded reference result
171 float *r = (float *)gOut_Ref;
172 float *s = (float *)gIn;
173 float *s2 = (float *)gIn2;
174 float *s3 = (float *)gIn3;
175 for (size_t j = 0; j < BUFFER_SIZE / sizeof(float); j++)
176 r[j] = (float)f->func.f_fff(s[j], s2[j], s3[j]);
177
178 // Read the data back
179 for (auto j = gMinVectorSizeIndex; j < gMaxVectorSizeIndex; j++)
180 {
181 if ((error =
182 clEnqueueReadBuffer(gQueue, gOutBuffer[j], CL_TRUE, 0,
183 BUFFER_SIZE, gOut[j], 0, NULL, NULL)))
184 {
185 vlog_error("ReadArray failed %d\n", error);
186 return error;
187 }
188 }
189
190 if (gSkipCorrectnessTesting) break;
191
192 // Verify data -- No verification possible.
193 // MAD is a random number generator.
194 if (0 == (i & 0x0fffffff))
195 {
196 vlog(".");
197 fflush(stdout);
198 }
199 }
200
201 if (!gSkipCorrectnessTesting)
202 {
203 if (gWimpyMode)
204 vlog("Wimp pass");
205 else
206 vlog("passed");
207
208 vlog("\t%8.2f @ {%a, %a, %a}", maxError, maxErrorVal, maxErrorVal2,
209 maxErrorVal3);
210 }
211
212 vlog("\n");
213
214 return CL_SUCCESS;
215 }
216