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