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_MATH_FUNCS_LOG_FUNCS_HPP
17 #define TEST_CONFORMANCE_CLCPP_MATH_FUNCS_LOG_FUNCS_HPP
18
19 #include <type_traits>
20 #include <cmath>
21
22 #include "common.hpp"
23
24 namespace detail
25 {
26
27 // This function reads values of FP_ILOGB0 and FP_ILOGBNAN macros defined on the device.
28 // OpenCL C++ Spec:
29 // The value of FP_ILOGB0 shall be either {INT_MIN} or {INT_MAX}. The value of FP_ILOGBNAN
30 // shall be either {INT_MAX} or {INT_MIN}.
get_ilogb_nan_zero(cl_device_id device,cl_context context,cl_command_queue queue,cl_int & ilogb_nan,cl_int & ilogb_zero)31 int get_ilogb_nan_zero(cl_device_id device, cl_context context, cl_command_queue queue, cl_int& ilogb_nan, cl_int& ilogb_zero)
32 {
33 cl_mem buffers[1];
34 cl_program program;
35 cl_kernel kernel;
36 size_t work_size[1];
37 int err;
38
39 std::string code_str =
40 "__kernel void get_ilogb_nan_zero(__global int *out)\n"
41 "{\n"
42 " out[0] = FP_ILOGB0;\n"
43 " out[1] = FP_ILOGBNAN;\n"
44 "}\n";
45 std::string kernel_name("get_ilogb_nan_zero");
46
47 err = create_opencl_kernel(context, &program, &kernel, code_str, kernel_name, "-cl-std=CL2.0", false);
48 RETURN_ON_ERROR(err)
49
50 std::vector<cl_int> output = generate_output<cl_int>(2);
51
52 buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * output.size(), NULL, &err);
53 RETURN_ON_CL_ERROR(err, "clCreateBuffer")
54
55 err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);
56 RETURN_ON_CL_ERROR(err, "clSetKernelArg");
57
58 work_size[0] = 1;
59 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, NULL);
60 RETURN_ON_CL_ERROR(err, "clEnqueueNDRangeKernel");
61
62 err = clEnqueueReadBuffer(
63 queue, buffers[0], CL_TRUE, 0, sizeof(cl_int) * output.size(),
64 static_cast<void *>(output.data()), 0, NULL, NULL
65 );
66 RETURN_ON_CL_ERROR(err, "clEnqueueReadBuffer");
67
68 // Save
69 ilogb_zero = output[0];
70 ilogb_nan = output[1];
71
72 clReleaseMemObject(buffers[0]);
73 clReleaseKernel(kernel);
74 clReleaseProgram(program);
75 return err;
76 }
77
78 } // namespace detail
79
80 struct logarithmic_func_ilogb : public unary_func<cl_float, cl_int>
81 {
logarithmic_func_ilogblogarithmic_func_ilogb82 logarithmic_func_ilogb(cl_int ilogb_nan, cl_int ilogb_zero)
83 : m_ilogb_nan(ilogb_nan), m_ilogb_zero(ilogb_zero)
84 {
85
86 }
87
strlogarithmic_func_ilogb88 std::string str()
89 {
90 return "ilogb";
91 }
92
headerslogarithmic_func_ilogb93 std::string headers()
94 {
95 return "#include <opencl_math>\n";
96 }
97
operator ()logarithmic_func_ilogb98 cl_int operator()(const cl_float& x)
99 {
100 if((std::isnan)(x))
101 {
102 return m_ilogb_nan;
103 }
104 else if(x == 0.0 || x == -0.0)
105 {
106 return m_ilogb_zero;
107 }
108 static_assert(
109 sizeof(cl_int) == sizeof(int),
110 "Tests assumes that sizeof(cl_int) == sizeof(int)"
111 );
112 return (std::ilogb)(x);
113 }
114
min1logarithmic_func_ilogb115 cl_float min1()
116 {
117 return -100.0f;
118 }
119
max1logarithmic_func_ilogb120 cl_float max1()
121 {
122 return 1000.0f;
123 }
124
in1_special_caseslogarithmic_func_ilogb125 std::vector<cl_float> in1_special_cases()
126 {
127 return {
128 cl_float(0.0f),
129 cl_float(-0.0f),
130 cl_float(1.0f),
131 cl_float(-1.0f),
132 cl_float(2.0f),
133 cl_float(-2.0f),
134 std::numeric_limits<cl_float>::infinity(),
135 -std::numeric_limits<cl_float>::infinity(),
136 std::numeric_limits<cl_float>::quiet_NaN()
137 };
138 }
139 private:
140 cl_int m_ilogb_nan;
141 cl_int m_ilogb_zero;
142 };
143
144 // gentype log(gentype x);
145 // gentype logb(gentype x);
146 // gentype log2(gentype x);
147 // gentype log10(gentype x);
148 // gentype log1p(gentype x);
149 // group_name, func_name, reference_func, use_ulp, ulp, ulp_for_embedded, max_delta, min1, max1
150 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, log, std::log, true, 3.0f, 4.0f, 0.001f, -10.0f, 1000.0f)
151 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, logb, std::logb, true, 0.0f, 0.0f, 0.001f, -10.0f, 1000.0f)
152 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, log2, std::log2, true, 3.0f, 4.0f, 0.001f, -10.0f, 1000.0f)
153 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, log10, std::log10, true, 3.0f, 4.0f, 0.001f, -10.0f, 1000.0f)
154 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, log1p, std::log1p, true, 2.0f, 4.0f, 0.001f, -10.0f, 1000.0f)
155
156 // gentype lgamma(gentype x);
157 // OpenCL C++ Spec.:
158 // The ULP values for built-in math functions lgamma and lgamma_r is currently undefined.
159 // Because of that we don't check ULP and set acceptable delta to 0.2f (20%).
160 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, lgamma, std::lgamma, false, 0.0f, 0.0f, 0.2f, -10.0f, 1000.0f)
161
162 // gentype lgamma_r(gentype x, intn* signp);
163 // OpenCL C++ Spec.:
164 // The ULP values for built-in math functions lgamma and lgamma_r is currently undefined.
165 // Because of that we don't check ULP and set acceptable delta to 0.2f (20%).
166 //
167 // Note:
168 // We DO NOT test if sign of the gamma function return by lgamma_r is correct.
169 MATH_FUNCS_DEFINE_UNARY_FUNC(logarithmic, lgamma_r, std::lgamma, false, 0.0f, 0.0f, 0.2f, -10.0f, 1000.0f)
170
171 // We need to specialize generate_kernel_unary<>() function template for logarithmic_func_lgamma_r
172 // because it takes two arguments, but only one of it is input, the 2nd one is used to return
173 // the sign of the gamma function.
174 // -----------------------------------------------------------------------------------
175 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
176 // -----------------------------------------------------------------------------------
177 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
178 template <>
generate_kernel_unary(logarithmic_func_lgamma_r func)179 std::string generate_kernel_unary<logarithmic_func_lgamma_r, cl_float, cl_float>(logarithmic_func_lgamma_r func)
180 {
181 return
182 "__kernel void test_lgamma_r(global float *input, global float *output)\n"
183 "{\n"
184 " size_t gid = get_global_id(0);\n"
185 " int sign;\n"
186 " output[gid] = lgamma_r(input[gid], &sign);\n"
187 "}\n";
188 }
189 #else
190 template <>
191 std::string generate_kernel_unary<logarithmic_func_lgamma_r, cl_float, cl_float>(logarithmic_func_lgamma_r func)
192 {
193 return
194 "" + func.defs() +
195 "" + func.headers() +
196 "#include <opencl_memory>\n"
197 "#include <opencl_work_item>\n"
198 "using namespace cl;\n"
199 "__kernel void test_lgamma_r(global_ptr<float[]> input, global_ptr<float[]> output)\n"
200 "{\n"
201 " size_t gid = get_global_id(0);\n"
202 " int sign;\n"
203 " output[gid] = lgamma_r(input[gid], &sign);\n"
204 "}\n";
205 }
206 #endif
207
208 // logarithmic functions
AUTO_TEST_CASE(test_logarithmic_funcs)209 AUTO_TEST_CASE(test_logarithmic_funcs)
210 (cl_device_id device, cl_context context, cl_command_queue queue, int n_elems)
211 {
212 int error = CL_SUCCESS;
213 int last_error = CL_SUCCESS;
214
215 // Check for EMBEDDED_PROFILE
216 bool is_embedded_profile = false;
217 char profile[128];
218 error = clGetDeviceInfo(device, CL_DEVICE_PROFILE, sizeof(profile), (void *)&profile, NULL);
219 RETURN_ON_CL_ERROR(error, "clGetDeviceInfo")
220 if (std::strcmp(profile, "EMBEDDED_PROFILE") == 0)
221 is_embedded_profile = true;
222
223 // Write values of FP_ILOGB0 and FP_ILOGBNAN, which are macros defined on the device, to
224 // ilogb_zero and ilogb_nan.
225 cl_int ilogb_nan = 0;
226 cl_int ilogb_zero = 0;
227 error = detail::get_ilogb_nan_zero(device, context, queue, ilogb_nan, ilogb_zero);
228 RETURN_ON_ERROR_MSG(error, "detail::get_ilogb_nan_zero function failed");
229
230 // intn ilogb(gentype x);
231 TEST_UNARY_FUNC_MACRO((logarithmic_func_ilogb(ilogb_nan, ilogb_zero)))
232
233 // gentype log(gentype x);
234 // gentype logb(gentype x);
235 // gentype log2(gentype x);
236 // gentype log10(gentype x);
237 // gentype log1p(gentype x);
238 TEST_UNARY_FUNC_MACRO((logarithmic_func_log(is_embedded_profile)))
239 TEST_UNARY_FUNC_MACRO((logarithmic_func_logb(is_embedded_profile)))
240 TEST_UNARY_FUNC_MACRO((logarithmic_func_log2(is_embedded_profile)))
241 TEST_UNARY_FUNC_MACRO((logarithmic_func_log10(is_embedded_profile)))
242 TEST_UNARY_FUNC_MACRO((logarithmic_func_log1p(is_embedded_profile)))
243
244 // gentype lgamma(gentype x);
245 TEST_UNARY_FUNC_MACRO((logarithmic_func_lgamma(is_embedded_profile)))
246
247 // gentype lgamma(gentype x);
248 //
249 // Note:
250 // We DO NOT test if sign of the gamma function return by lgamma_r is correct
251 TEST_UNARY_FUNC_MACRO((logarithmic_func_lgamma_r(is_embedded_profile)))
252
253 if(error != CL_SUCCESS)
254 {
255 return -1;
256 }
257 return error;
258 }
259
260 #endif // TEST_CONFORMANCE_CLCPP_MATH_FUNCS_LOG_FUNCS_HPP
261