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 #ifndef TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
17 #define TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
18
19 #include <type_traits>
20 #include <algorithm>
21 #include <string>
22 #include <cmath>
23
24 #include "../common.hpp"
25
26 #include "detail/base_func_type.hpp"
27 #include "generate_inputs.hpp"
28 #include "compare.hpp"
29
30 template<class IN1, class IN2, class IN3, class OUT1>
31 struct ternary_func : public detail::base_func_type<OUT1>
32 {
33 typedef IN1 in1_type;
34 typedef IN2 in2_type;
35 typedef IN3 in3_type;
36 typedef OUT1 out_type;
37
~ternary_functernary_func38 virtual ~ternary_func() {};
39 virtual std::string str() = 0;
40
decl_strternary_func41 std::string decl_str()
42 {
43 return type_name<OUT1>() + "(" + type_name<IN1>() + ", " + type_name<IN2>()+ ", " + type_name<IN3>() + ")";
44 }
45
is_in1_boolternary_func46 bool is_in1_bool()
47 {
48 return false;
49 }
50
is_in2_boolternary_func51 bool is_in2_bool()
52 {
53 return false;
54 }
55
is_in3_boolternary_func56 bool is_in3_bool()
57 {
58 return false;
59 }
60
min1ternary_func61 IN1 min1()
62 {
63 return detail::get_min<IN1>();
64 }
65
max1ternary_func66 IN1 max1()
67 {
68 return detail::get_max<IN1>();
69 }
70
min2ternary_func71 IN2 min2()
72 {
73 return detail::get_min<IN2>();
74 }
75
max2ternary_func76 IN2 max2()
77 {
78 return detail::get_max<IN2>();
79 }
80
min3ternary_func81 IN3 min3()
82 {
83 return detail::get_min<IN3>();
84 }
85
max3ternary_func86 IN3 max3()
87 {
88 return detail::get_max<IN3>();
89 }
90
in1_special_casesternary_func91 std::vector<IN1> in1_special_cases()
92 {
93 return { };
94 }
95
in2_special_casesternary_func96 std::vector<IN2> in2_special_cases()
97 {
98 return { };
99 }
100
in3_special_casesternary_func101 std::vector<IN3> in3_special_cases()
102 {
103 return { };
104 }
105
106 template<class T>
107 typename make_vector_type<cl_double, vector_size<T>::value>::type
deltaternary_func108 delta(const IN1& in1, const IN2& in2, const IN3& in3, const T& expected)
109 {
110 typedef
111 typename make_vector_type<cl_double, vector_size<T>::value>::type
112 delta_vector_type;
113 // Take care of unused variable warning
114 (void) in1;
115 (void) in2;
116 (void) in3;
117 auto e = detail::make_value<delta_vector_type>(1e-3);
118 return detail::multiply<delta_vector_type>(e, expected);
119 }
120 };
121
122 // -----------------------------------------------------------------------------------
123 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
124 // -----------------------------------------------------------------------------------
125 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
126 template <class func_type, class in1_type, class in2_type, class in3_type, class out_type>
generate_kernel_ternary(func_type func)127 std::string generate_kernel_ternary(func_type func)
128 {
129 std::string in1_value = "input1[gid]";
130 if(func.is_in1_bool())
131 {
132 std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
133 in1_value = "(input1[gid] != (int" + i + ")(0))";
134 }
135 std::string in2_value = "input2[gid]";
136 if(func.is_in2_bool())
137 {
138 std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
139 in2_value = "(input2[gid] != (int" + i + ")(0))";
140 }
141 std::string in3_value = "input3[gid]";
142 if(func.is_in3_bool())
143 {
144 std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
145 in3_value = "(input3[gid] != (int" + i + ")(0))";
146 }
147 std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_value + ")";
148 if(func.is_out_bool())
149 {
150 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
151 function_call = "convert_int" + i + "(" + func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_value + "))";
152 }
153 return
154 "__kernel void " + func.get_kernel_name() + "(global " + type_name<in1_type>() + " *input1,\n"
155 " global " + type_name<in2_type>() + " *input2,\n"
156 " global " + type_name<in3_type>() + " *input3,\n"
157 " global " + type_name<out_type>() + " *output)\n"
158 "{\n"
159 " size_t gid = get_global_id(0);\n"
160 " output[gid] = " + function_call + ";\n"
161 "}\n";
162 }
163 #else
164 template <class func_type, class in1_type, class in2_type, class in3_type, class out_type>
generate_kernel_ternary(func_type func)165 std::string generate_kernel_ternary(func_type func)
166 {
167 std::string headers = func.headers();
168 std::string in1_value = "input1[gid]";
169 if(func.is_in1_bool())
170 {
171 std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
172 in1_value = "(input1[gid] != (int" + i + ")(0))";
173 }
174 std::string in2_value = "input2[gid]";
175 if(func.is_in2_bool())
176 {
177 std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
178 in2_value = "(input2[gid] != (int" + i + ")(0))";
179 }
180 std::string in3_value = "input3[gid]";
181 if(func.is_in3_bool())
182 {
183 std::string i = vector_size<in3_type>::value == 1 ? "" : std::to_string(vector_size<in3_type>::value);
184 in3_value = "(input3[gid] != (int" + i + ")(0))";
185 }
186 std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_value + ")";
187 if(func.is_out_bool())
188 {
189 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
190 function_call = "convert_cast<int" + i + ">(" + func.str() + "(" + in1_value + ", " + in2_value + ", " + in3_value + "))";
191 }
192 if(func.is_out_bool() || func.is_in1_bool() || func.is_in2_bool() || func.is_in3_bool())
193 {
194 if(headers.find("#include <opencl_convert>") == std::string::npos)
195 {
196 headers += "#include <opencl_convert>\n";
197 }
198 }
199 return
200 "" + func.defs() +
201 "" + headers +
202 "#include <opencl_memory>\n"
203 "#include <opencl_work_item>\n"
204 "using namespace cl;\n"
205 "__kernel void " + func.get_kernel_name() + "(global_ptr<" + type_name<in1_type>() + "[]> input1,\n"
206 " global_ptr<" + type_name<in2_type>() + "[]> input2,\n"
207 " global_ptr<" + type_name<in3_type>() + "[]> input3,\n"
208 " global_ptr<" + type_name<out_type>() + "[]> output)\n"
209 "{\n"
210 " size_t gid = get_global_id(0);\n"
211 " output[gid] = " + function_call + ";\n"
212 "}\n";
213 }
214 #endif
215
216 template<class INPUT1, class INPUT2, class INPUT3, class OUTPUT, class ternary_op>
verify_ternary(const std::vector<INPUT1> & in1,const std::vector<INPUT2> & in2,const std::vector<INPUT3> & in3,const std::vector<OUTPUT> & out,ternary_op op)217 bool verify_ternary(const std::vector<INPUT1> &in1,
218 const std::vector<INPUT2> &in2,
219 const std::vector<INPUT3> &in3,
220 const std::vector<OUTPUT> &out,
221 ternary_op op)
222 {
223 for(size_t i = 0; i < in1.size(); i++)
224 {
225 auto expected = op(in1[i], in2[i], in3[i]);
226 if(!are_equal(expected, out[i], op.delta(in1[i], in2[i], in3[i], expected), op))
227 {
228 print_error_msg(expected, out[i], i, op);
229 return false;
230 }
231 }
232 return true;
233 }
234
235 template <class ternary_op>
test_ternary_func(cl_device_id device,cl_context context,cl_command_queue queue,size_t count,ternary_op op)236 int test_ternary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, ternary_op op)
237 {
238 cl_mem buffers[4];
239 cl_program program;
240 cl_kernel kernel;
241 size_t work_size[1];
242 int err;
243
244 typedef typename ternary_op::in1_type INPUT1;
245 typedef typename ternary_op::in2_type INPUT2;
246 typedef typename ternary_op::in3_type INPUT3;
247 typedef typename ternary_op::out_type OUTPUT;
248
249 // Don't run test for unsupported types
250 if(!(type_supported<INPUT1>(device)
251 && type_supported<INPUT2>(device)
252 && type_supported<INPUT3>(device)
253 && type_supported<OUTPUT>(device)))
254 {
255 return CL_SUCCESS;
256 }
257
258 std::string code_str = generate_kernel_ternary<ternary_op, INPUT1, INPUT2, INPUT3, OUTPUT>(op);
259 std::string kernel_name = op.get_kernel_name();
260
261 // -----------------------------------------------------------------------------------
262 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
263 // -----------------------------------------------------------------------------------
264 // Only OpenCL C++ to SPIR-V compilation
265 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
266 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
267 RETURN_ON_ERROR(err)
268 return err;
269 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
270 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
271 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
272 RETURN_ON_ERROR(err)
273 #else
274 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
275 RETURN_ON_ERROR(err)
276 #endif
277
278 std::vector<INPUT1> in1_spec_cases = op.in1_special_cases();
279 std::vector<INPUT2> in2_spec_cases = op.in2_special_cases();
280 std::vector<INPUT3> in3_spec_cases = op.in3_special_cases();
281 prepare_special_cases(in1_spec_cases, in2_spec_cases, in3_spec_cases);
282 std::vector<INPUT1> input1 = generate_input<INPUT1>(count, op.min1(), op.max1(), in1_spec_cases);
283 std::vector<INPUT2> input2 = generate_input<INPUT2>(count, op.min2(), op.max2(), in2_spec_cases);
284 std::vector<INPUT3> input3 = generate_input<INPUT3>(count, op.min3(), op.max3(), in3_spec_cases);
285 std::vector<OUTPUT> output = generate_output<OUTPUT>(count);
286
287 buffers[0] = clCreateBuffer(
288 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT1) * input1.size(), NULL, &err
289 );
290 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
291
292 buffers[1] = clCreateBuffer(
293 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT2) * input2.size(), NULL, &err
294 );
295 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
296
297 buffers[2] = clCreateBuffer(
298 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT3) * input3.size(), NULL, &err
299 );
300 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
301
302 buffers[3] = clCreateBuffer(
303 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(OUTPUT) * output.size(), NULL, &err
304 );
305 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
306
307 err = clEnqueueWriteBuffer(
308 queue, buffers[0], CL_TRUE, 0, sizeof(INPUT1) * input1.size(),
309 static_cast<void *>(input1.data()), 0, NULL, NULL
310 );
311 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer");
312
313 err = clEnqueueWriteBuffer(
314 queue, buffers[1], CL_TRUE, 0, sizeof(INPUT2) * input2.size(),
315 static_cast<void *>(input2.data()), 0, NULL, NULL
316 );
317 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer");
318
319 err = clEnqueueWriteBuffer(
320 queue, buffers[2], CL_TRUE, 0, sizeof(INPUT3) * input3.size(),
321 static_cast<void *>(input3.data()), 0, NULL, NULL
322 );
323 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer");
324
325 err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
326 err |= clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
327 err |= clSetKernelArg(kernel, 2, sizeof(buffers[2]), &buffers[2]);
328 err |= clSetKernelArg(kernel, 3, sizeof(buffers[3]), &buffers[3]);
329 RETURN_ON_CL_ERROR(err, "clSetKernelArg");
330
331 work_size[0] = count;
332 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
333 RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel");
334
335 err = clEnqueueReadBuffer(
336 queue, buffers[3], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
337 static_cast<void *>(output.data()), 0, NULL, NULL
338 );
339 RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");
340
341 if (!verify_ternary(input1, input2, input3, output, op))
342 {
343 RETURN_ON_ERROR_MSG(-1,
344 "test_%s %s(%s, %s, %s) failed", op.str().c_str(),
345 type_name<OUTPUT>().c_str(),
346 type_name<INPUT1>().c_str(),
347 type_name<INPUT2>().c_str(),
348 type_name<INPUT3>().c_str()
349 );
350 }
351 log_info(
352 "test_%s %s(%s, %s, %s) passed\n", op.str().c_str(),
353 type_name<OUTPUT>().c_str(),
354 type_name<INPUT1>().c_str(),
355 type_name<INPUT2>().c_str(),
356 type_name<INPUT3>().c_str()
357 );
358
359 clReleaseMemObject(buffers[0]);
360 clReleaseMemObject(buffers[1]);
361 clReleaseMemObject(buffers[2]);
362 clReleaseMemObject(buffers[3]);
363 clReleaseKernel(kernel);
364 clReleaseProgram(program);
365 return err;
366 }
367
368 #endif // TEST_CONFORMANCE_CLCPP_UTILS_TEST_TERNARY_HPP
369