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_UNARY_HPP
17 #define TEST_CONFORMANCE_CLCPP_UTILS_TEST_UNARY_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 OUT1>
31 struct unary_func : public detail::base_func_type<OUT1>
32 {
33 typedef IN1 in_type;
34 typedef OUT1 out_type;
35
~unary_funcunary_func36 virtual ~unary_func() {};
37 virtual std::string str() = 0;
38
39 // Return string with function type, for example: int(float).
decl_strunary_func40 std::string decl_str()
41 {
42 return type_name<OUT1>() + "(" + type_name<IN1>() + ")";
43 }
44
45 // Return true if IN1 type in OpenCL kernel should be treated
46 // as bool type; false otherwise.
is_in1_boolunary_func47 bool is_in1_bool()
48 {
49 return false;
50 }
51
52 // Return min value that can be used as a first argument.
min1unary_func53 IN1 min1()
54 {
55 return detail::get_min<IN1>();
56 }
57
58 // Return max value that can be used as a first argument.
max1unary_func59 IN1 max1()
60 {
61 return detail::get_max<IN1>();
62 }
63
64 // This returns a list of special cases input values we want to
65 // test.
in_special_casesunary_func66 std::vector<IN1> in_special_cases()
67 {
68 return { };
69 }
70
71 // Max error. Error should be raised if
72 // abs(result - expected) > delta(.., expected)
73 //
74 // Default value: 0.001 * expected
75 //
76 // (This effects how are_equal() function works,
77 // it may not have effect if verify() method in derived
78 // class does not use are_equal() function.)
79 //
80 // Only for FP numbers/vectors
81 template<class T>
82 typename make_vector_type<cl_double, vector_size<T>::value>::type
deltaunary_func83 delta(const IN1& in1, const T& expected)
84 {
85 typedef
86 typename make_vector_type<cl_double, vector_size<T>::value>::type
87 delta_vector_type;
88 // Take care of unused variable warning
89 (void) in1;
90 auto e = detail::make_value<delta_vector_type>(1e-3);
91 return detail::multiply<delta_vector_type>(e, expected);
92 }
93 };
94
95 // -----------------------------------------------------------------------------------
96 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
97 // -----------------------------------------------------------------------------------
98 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
99 template <class func_type, class in_type, class out_type>
generate_kernel_unary(func_type func)100 std::string generate_kernel_unary(func_type func)
101 {
102 std::string in1_value = "input[gid]";
103 // Convert uintN to boolN values
104 if(func.is_in1_bool())
105 {
106 std::string i = vector_size<in_type>::value == 1 ? "" : std::to_string(vector_size<in_type>::value);
107 in1_value = "(input[gid] != (int" + i + ")(0))";
108 }
109 std::string function_call = func.str() + "(" + in1_value + ");";
110 // Convert boolN result of funtion func_type to uintN
111 if(func.is_out_bool())
112 {
113 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
114 function_call = "convert_int" + i + "(" + func.str() + "(" + in1_value + "))";
115 }
116 return
117 "__kernel void " + func.get_kernel_name() + "(global " + type_name<in_type>() + " *input, global " + type_name<out_type>() + " *output)\n"
118 "{\n"
119 " size_t gid = get_global_id(0);\n"
120 " output[gid] = " + function_call + ";\n"
121 "}\n";
122 }
123 #else
124 template <class func_type, class in_type, class out_type>
generate_kernel_unary(func_type func)125 std::string generate_kernel_unary(func_type func)
126 {
127 std::string headers = func.headers();
128 std::string in1_value = "input[gid]";
129 if(func.is_in1_bool())
130 {
131 std::string i = vector_size<in_type>::value == 1 ? "" : std::to_string(vector_size<in_type>::value);
132 in1_value = "(input[gid] != (int" + i + ")(0))";
133 }
134 std::string function_call = func.str() + "(" + in1_value + ")";
135 if(func.is_out_bool())
136 {
137 std::string i = vector_size<out_type>::value == 1 ? "" : std::to_string(vector_size<out_type>::value);
138 function_call = "convert_cast<int" + i + ">(" + func.str() + "(" + in1_value + "))";
139 }
140 if(func.is_out_bool() || func.is_in1_bool())
141 {
142 if(headers.find("#include <opencl_convert>") == std::string::npos)
143 {
144 headers += "#include <opencl_convert>\n";
145 }
146 }
147 return
148 "" + func.defs() +
149 "" + headers +
150 "#include <opencl_memory>\n"
151 "#include <opencl_work_item>\n"
152 "using namespace cl;\n"
153 "__kernel void " + func.get_kernel_name() + "(global_ptr<" + type_name<in_type>() + "[]> input,"
154 "global_ptr<" + type_name<out_type>() + "[]> output)\n"
155 "{\n"
156 " size_t gid = get_global_id(0);\n"
157 " output[gid] = " + function_call + ";\n"
158 "}\n";
159 }
160 #endif
161
162 template<class INPUT, class OUTPUT, class unary_op>
verify_unary(const std::vector<INPUT> & in,const std::vector<OUTPUT> & out,unary_op op)163 bool verify_unary(const std::vector<INPUT> &in, const std::vector<OUTPUT> &out, unary_op op)
164 {
165 for(size_t i = 0; i < in.size(); i++)
166 {
167 auto expected = op(in[i]);
168 if(!are_equal(expected, out[i], op.delta(in[i], expected), op))
169 {
170 print_error_msg(expected, out[i], i, op);
171 return false;
172 }
173 }
174 return true;
175 }
176
177 template <class unary_op>
test_unary_func(cl_device_id device,cl_context context,cl_command_queue queue,size_t count,unary_op op)178 int test_unary_func(cl_device_id device, cl_context context, cl_command_queue queue, size_t count, unary_op op)
179 {
180 cl_mem buffers[2];
181 cl_program program;
182 cl_kernel kernel;
183 size_t work_size[1];
184 int err;
185
186 typedef typename unary_op::in_type INPUT;
187 typedef typename unary_op::out_type OUTPUT;
188
189 // Don't run test for unsupported types
190 if(!(type_supported<INPUT>(device) && type_supported<OUTPUT>(device)))
191 {
192 return CL_SUCCESS;
193 }
194
195 std::string code_str = generate_kernel_unary<unary_op, INPUT, OUTPUT>(op);
196 std::string kernel_name = op.get_kernel_name();
197
198 // -----------------------------------------------------------------------------------
199 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
200 // -----------------------------------------------------------------------------------
201 // Only OpenCL C++ to SPIR-V compilation
202 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
203 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
204 RETURN_ON_ERROR(err)
205 return err;
206 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
207 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
208 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
209 RETURN_ON_ERROR(err)
210 #else
211 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name);
212 RETURN_ON_ERROR(err)
213 #endif
214
215 std::vector<INPUT> input = generate_input<INPUT>(count, op.min1(), op.max1(), op.in_special_cases());
216 std::vector<OUTPUT> output = generate_output<OUTPUT>(count);
217
218 buffers[0] = clCreateBuffer(
219 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT) * input.size(), NULL, &err
220 );
221 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
222
223 buffers[1] = clCreateBuffer(
224 context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(OUTPUT) * output.size(), NULL, &err
225 );
226 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
227
228 err = clEnqueueWriteBuffer(
229 queue, buffers[0], CL_TRUE, 0, sizeof(INPUT) * input.size(),
230 static_cast<void *>(input.data()), 0, NULL, NULL
231 );
232 RETURN_ON_CL_ERROR(err, "clEnqueueWriteBuffer");
233
234 err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
235 err |= clSetKernelArg(kernel, 1, sizeof(buffers[1]), &buffers[1]);
236 RETURN_ON_CL_ERROR(err, "clSetKernelArg");
237
238 work_size[0] = count;
239 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
240 RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel");
241
242 err = clEnqueueReadBuffer(
243 queue, buffers[1], CL_TRUE, 0, sizeof(OUTPUT) * output.size(),
244 static_cast<void *>(output.data()), 0, NULL, NULL
245 );
246 RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");
247
248 if (!verify_unary(input, output, op))
249 {
250 RETURN_ON_ERROR_MSG(-1, "test_%s %s(%s) failed", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());
251 }
252 log_info("test_%s %s(%s) passed\n", op.str().c_str(), type_name<OUTPUT>().c_str(), type_name<INPUT>().c_str());
253
254 clReleaseMemObject(buffers[0]);
255 clReleaseMemObject(buffers[1]);
256 clReleaseKernel(kernel);
257 clReleaseProgram(program);
258 return err;
259 }
260
261 #endif // TEST_CONFORMANCE_CLCPP_UTILS_TEST_UNARY_HPP
262