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_BINARY_HPP
17 #define TEST_CONFORMANCE_CLCPP_UTILS_TEST_BINARY_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 OUT1>
31 struct binary_func : public detail::base_func_type<OUT1>
32 {
33 typedef IN1 in1_type;
34 typedef IN2 in2_type;
35 typedef OUT1 out_type;
36
~binary_funcbinary_func37 virtual ~binary_func() {};
38 virtual std::string str() = 0;
39
decl_strbinary_func40 std::string decl_str()
41 {
42 return type_name<OUT1>() + "(" + type_name<IN1>() + ", " + type_name<IN2>() + ")";
43 }
44
is_in1_boolbinary_func45 bool is_in1_bool()
46 {
47 return false;
48 }
49
is_in2_boolbinary_func50 bool is_in2_bool()
51 {
52 return false;
53 }
54
min1binary_func55 IN1 min1()
56 {
57 return detail::get_min<IN1>();
58 }
59
max1binary_func60 IN1 max1()
61 {
62 return detail::get_max<IN1>();
63 }
64
min2binary_func65 IN2 min2()
66 {
67 return detail::get_min<IN2>();
68 }
69
max2binary_func70 IN2 max2()
71 {
72 return detail::get_max<IN2>();
73 }
74
in1_special_casesbinary_func75 std::vector<IN1> in1_special_cases()
76 {
77 return { };
78 }
79
in2_special_casesbinary_func80 std::vector<IN2> in2_special_cases()
81 {
82 return { };
83 }
84
85 template<class T>
86 typename make_vector_type<cl_double, vector_size<T>::value>::type
deltabinary_func87 delta(const IN1& in1, const IN2& in2, const T& expected)
88 {
89 typedef
90 typename make_vector_type<cl_double, vector_size<T>::value>::type
91 delta_vector_type;
92 // Take care of unused variable warning
93 (void) in1;
94 (void) in2;
95 auto e = detail::make_value<delta_vector_type>(1e-3);
96 return detail::multiply<delta_vector_type>(e, expected);
97 }
98 };
99
100 // -----------------------------------------------------------------------------------
101 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
102 // -----------------------------------------------------------------------------------
103 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
104 template <class func_type, class in1_type, class in2_type, class out_type>
generate_kernel_binary(func_type func)105 std::string generate_kernel_binary(func_type func)
106 {
107 std::string in1_value = "input1[gid]";
108 if(func.is_in1_bool())
109 {
110 std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
111 in1_value = "(input1[gid] != (int" + i + ")(0))";
112 }
113 std::string in2_value = "input2[gid]";
114 if(func.is_in2_bool())
115 {
116 std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
117 in2_value = "(input2[gid] != (int" + i + ")(0))";
118 }
119 std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ")";
120 if(func.is_out_bool())
121 {
122 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
123 function_call = "convert_int" + i + "(" + func.str() + "(" + in1_value + ", " + in2_value + "))";
124 }
125 return
126 "__kernel void " + func.get_kernel_name() + "(global " + type_name<in1_type>() + " *input1,\n"
127 " global " + type_name<in2_type>() + " *input2,\n"
128 " global " + type_name<out_type>() + " *output)\n"
129 "{\n"
130 " size_t gid = get_global_id(0);\n"
131 " output[gid] = " + function_call + ";\n"
132 "}\n";
133 }
134 #else
135 template <class func_type, class in1_type, class in2_type, class out_type>
generate_kernel_binary(func_type func)136 std::string generate_kernel_binary(func_type func)
137 {
138 std::string headers = func.headers();
139 std::string in1_value = "input1[gid]";
140 if(func.is_in1_bool())
141 {
142 std::string i = vector_size<in1_type>::value == 1 ? "" : std::to_string(vector_size<in1_type>::value);
143 in1_value = "(input1[gid] != (int" + i + ")(0))";
144 }
145 std::string in2_value = "input2[gid]";
146 if(func.is_in2_bool())
147 {
148 std::string i = vector_size<in2_type>::value == 1 ? "" : std::to_string(vector_size<in2_type>::value);
149 in2_value = "(input2[gid] != (int" + i + ")(0))";
150 }
151 std::string function_call = func.str() + "(" + in1_value + ", " + in2_value + ")";
152 if(func.is_out_bool())
153 {
154 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
155 function_call = "convert_cast<int" + i + ">(" + func.str() + "(" + in1_value + ", " + in2_value + "))";
156 }
157 if(func.is_out_bool() || func.is_in1_bool() || func.is_in2_bool())
158 {
159 if(headers.find("#include <opencl_convert>") == std::string::npos)
160 {
161 headers += "#include <opencl_convert>\n";
162 }
163 }
164 return
165 "" + func.defs() +
166 "" + headers +
167 "#include <opencl_memory>\n"
168 "#include <opencl_work_item>\n"
169 "using namespace cl;\n"
170 "__kernel void " + func.get_kernel_name() + "(global_ptr<" + type_name<in1_type>() + "[]> input1,\n"
171 " global_ptr<" + type_name<in2_type>() + "[]> input2,\n"
172 " global_ptr<" + type_name<out_type>() + "[]> output)\n"
173 "{\n"
174 " size_t gid = get_global_id(0);\n"
175 " output[gid] = " + function_call + ";\n"
176 "}\n";
177 }
178 #endif
179
180 template<class INPUT1, class INPUT2, class OUTPUT, class binary_op>
verify_binary(const std::vector<INPUT1> & in1,const std::vector<INPUT2> & in2,const std::vector<OUTPUT> & out,binary_op op)181 bool verify_binary(const std::vector<INPUT1> &in1,
182 const std::vector<INPUT2> &in2,
183 const std::vector<OUTPUT> &out,
184 binary_op op)
185 {
186 for(size_t i = 0; i < in1.size(); i++)
187 {
188 auto expected = op(in1[i], in2[i]);
189 if(!are_equal(expected, out[i], op.delta(in1[i], in2[i], expected), op))
190 {
191 print_error_msg(expected, out[i], i, op);
192 return false;
193 }
194 }
195 return true;
196 }
197
198 template <class binary_op>
test_binary_func(cl_device_id device,cl_context context,cl_command_queue queue,size_t count,binary_op op)199 int test_binary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, binary_op op)
200 {
201 cl_mem buffers[3];
202 cl_program program;
203 cl_kernel kernel;
204 size_t work_size[1];
205 int err;
206
207 typedef typename binary_op::in1_type INPUT1;
208 typedef typename binary_op::in2_type INPUT2;
209 typedef typename binary_op::out_type OUTPUT;
210
211 // Don't run test for unsupported types
212 if(!(type_supported<INPUT1>(device)
213 && type_supported<INPUT2>(device)
214 && type_supported<OUTPUT>(device)))
215 {
216 return CL_SUCCESS;
217 }
218
219 std::string code_str = generate_kernel_binary<binary_op, INPUT1, INPUT2, OUTPUT>(op);
220 std::string kernel_name = op.get_kernel_name();
221
222 // -----------------------------------------------------------------------------------
223 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
224 // -----------------------------------------------------------------------------------
225 // Only OpenCL C++ to SPIR-V compilation
226 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
227 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
228 RETURN_ON_ERROR(err)
229 return err;
230 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
231 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
232 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
233 RETURN_ON_ERROR(err)
234 #else
235 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
236 RETURN_ON_ERROR(err)
237 #endif
238
239 std::vector<INPUT1> in1_spec_cases = op.in1_special_cases();
240 std::vector<INPUT2> in2_spec_cases = op.in2_special_cases();
241 prepare_special_cases(in1_spec_cases, in2_spec_cases);
242 std::vector<INPUT1> input1 = generate_input<INPUT1>(count, op.min1(), op.max1(), in1_spec_cases);
243 std::vector<INPUT2> input2 = generate_input<INPUT2>(count, op.min2(), op.max2(), in2_spec_cases);
244 std::vector<OUTPUT> output = generate_output<OUTPUT>(count);
245
246 buffers[0] = clCreateBuffer(
247 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT1) * input1.size(), NULL, &err
248 );
249 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
250
251 buffers[1] = clCreateBuffer(
252 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT2) * input2.size(), NULL, &err
253 );
254 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
255
256 buffers[2] = clCreateBuffer(
257 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(OUTPUT) * output.size(), NULL, &err
258 );
259 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
260
261 err = clEnqueueWriteBuffer(
262 queue, buffers[0], CL_TRUE, 0, sizeof(INPUT1) * input1.size(),
263 static_cast<void *>(input1.data()), 0, NULL, NULL
264 );
265 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer")
266
267 err = clEnqueueWriteBuffer(
268 queue, buffers[1], CL_TRUE, 0, sizeof(INPUT2) * input2.size(),
269 static_cast<void *>(input2.data()), 0, NULL, NULL
270 );
271 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer")
272
273 err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
274 err |= clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
275 err |= clSetKernelArg(kernel, 2, sizeof(buffers[2]), &buffers[2]);
276 RETURN_ON_CL_ERROR(err, "clSetKernelArg");
277
278 work_size[0] = count;
279 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
280 RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel");
281
282 err = clEnqueueReadBuffer(
283 queue, buffers[2], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
284 static_cast<void *>(output.data()), 0, NULL, NULL
285 );
286 RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");
287
288 if (!verify_binary(input1, input2, output, op))
289 {
290 RETURN_ON_ERROR_MSG(-1,
291 "test_%s %s(%s, %s) failed", op.str().c_str(),
292 type_name<OUTPUT>().c_str(), type_name<INPUT1>().c_str(), type_name<INPUT2>().c_str()
293 );
294 }
295 log_info(
296 "test_%s %s(%s, %s) passed\n", op.str().c_str(),
297 type_name<OUTPUT>().c_str(), type_name<INPUT1>().c_str(), type_name<INPUT2>().c_str()
298 );
299
300 clReleaseMemObject(buffers[0]);
301 clReleaseMemObject(buffers[1]);
302 clReleaseMemObject(buffers[2]);
303 clReleaseKernel(kernel);
304 clReleaseProgram(program);
305 return err;
306 }
307
308 #endif // TEST_CONFORMANCE_CLCPP_UTILS_TEST_BINARY_HPP
309