1 //
2 // Copyright (c) 2023 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 #include "CL/cl_half.h"
17 #include "harness/compat.h"
18 #include "harness/errorHelpers.h"
19 #include "harness/stringHelpers.h"
20
21 #include <stdio.h>
22 #include <stdlib.h>
23 #include <string.h>
24 #include <sys/types.h>
25 #include <sys/stat.h>
26
27 #include <algorithm>
28 #include <cstdint>
29 #include <map>
30 #include <vector>
31
32 #include "procs.h"
33
34 extern cl_half_rounding_mode halfRoundingMode;
35
36 #define HFF(num) cl_half_from_float(num, halfRoundingMode)
37 #define HTF(num) cl_half_to_float(num)
38
39 namespace {
40 const char *int2float_kernel_code = R"(
41 %s
42 __kernel void test_X2Y(__global TYPE_X *src, __global TYPE_Y *dst)
43 {
44 int tid = get_global_id(0);
45
46 dst[tid] = (TYPE_Y)src[tid];
47
48 })";
49
50 template <bool int2fp> struct TypesIterator
51 {
TypesIterator__anon70898bdb0111::TypesIterator52 TypesIterator(cl_device_id deviceID, cl_context context,
53 cl_command_queue queue, int num_elems, const char *test_name)
54 : context(context), queue(queue), test_name(test_name),
55 num_elements(num_elems)
56 {
57 fp16Support = is_extension_available(deviceID, "cl_khr_fp16");
58 fp64Support = is_extension_available(deviceID, "cl_khr_fp64");
59
60 type2name[sizeof(cl_half)] = std::make_pair("half", "short");
61 type2name[sizeof(cl_float)] = std::make_pair("float", "int");
62 type2name[sizeof(cl_double)] = std::make_pair("double", "long");
63
64 std::tuple<cl_float, cl_half, cl_double> it;
65 for_each_elem(it);
66 }
67
generate_random_inputs__anon70898bdb0111::TypesIterator68 template <typename T> void generate_random_inputs(std::vector<T> &v)
69 {
70 RandomSeed seed(gRandomSeed);
71
72 if (sizeof(T) == sizeof(cl_half))
73 {
74 // Bound generated half values to 0x1.ffcp+14(32752.0) which is the
75 // largest cl_half value smaller than the max value of cl_short,
76 // 32767.
77 if (int2fp)
78 {
79 auto random_generator = [&seed]() {
80 return (cl_short)get_random_float(
81 -MAKE_HEX_FLOAT(0x1.ffcp+14, 1.9990234375f, 14),
82 MAKE_HEX_FLOAT(0x1.ffcp+14, 1.9990234375f, 14), seed);
83 };
84 std::generate(v.begin(), v.end(), random_generator);
85 }
86 else
87 {
88 auto random_generator = [&seed]() {
89 return HFF(get_random_float(
90 -MAKE_HEX_FLOAT(0x1.ffcp+14, 1.9990234375f, 14),
91 MAKE_HEX_FLOAT(0x1.ffcp+14, 1.9990234375f, 14), seed));
92 };
93 std::generate(v.begin(), v.end(), random_generator);
94 }
95 }
96 else if (sizeof(T) == sizeof(cl_float))
97 {
98 auto random_generator = [&seed]() {
99 return get_random_float(-MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31),
100 MAKE_HEX_FLOAT(0x1.0p31f, 0x1, 31),
101 seed);
102 };
103 std::generate(v.begin(), v.end(), random_generator);
104 }
105 else if (sizeof(T) == sizeof(cl_double))
106 {
107 auto random_generator = [&seed]() {
108 return get_random_double(-MAKE_HEX_DOUBLE(0x1.0p63, 0x1, 63),
109 MAKE_HEX_DOUBLE(0x1.0p63, 0x1, 63),
110 seed);
111 };
112 std::generate(v.begin(), v.end(), random_generator);
113 }
114 }
115
equal_value__anon70898bdb0111::TypesIterator116 template <typename Tx, typename Ty> static bool equal_value(Tx a, Ty b)
117 {
118 return a == (Tx)b;
119 }
120
equal_value_from_half__anon70898bdb0111::TypesIterator121 static bool equal_value_from_half(cl_short a, cl_half b)
122 {
123 return a == (cl_short)HTF(b);
124 }
125
equal_value_to_half__anon70898bdb0111::TypesIterator126 static bool equal_value_to_half(cl_half a, cl_short b)
127 {
128 return a == HFF((float)b);
129 }
130
131
132 template <typename Tx, typename Ty>
verify_X2Y__anon70898bdb0111::TypesIterator133 int verify_X2Y(std::vector<Tx> input, std::vector<Ty> output)
134 {
135 if (std::is_same<Tx, cl_half>::value
136 || std::is_same<Ty, cl_half>::value)
137 {
138 bool res = true;
139 if (int2fp)
140 res = std::equal(output.begin(), output.end(), input.begin(),
141 equal_value_to_half);
142 else
143 res = std::equal(output.begin(), output.end(), input.begin(),
144 equal_value_from_half);
145
146 if (!res)
147 {
148 log_error("%s test failed\n", test_name.c_str());
149 return -1;
150 }
151 }
152 else
153 {
154 if (!std::equal(output.begin(), output.end(), input.begin(),
155 equal_value<Tx, Ty>))
156 {
157 log_error("%s test failed\n", test_name.c_str());
158 return -1;
159 }
160 }
161
162 log_info("%s test passed\n", test_name.c_str());
163 return 0;
164 }
165
test_X2Y__anon70898bdb0111::TypesIterator166 template <typename Tx, typename Ty> int test_X2Y()
167 {
168 clMemWrapper streams[2];
169 clProgramWrapper program;
170 clKernelWrapper kernel;
171 int err;
172
173 std::vector<Tx> input(num_elements);
174 std::vector<Ty> output(num_elements);
175
176 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
177 sizeof(Tx) * num_elements, nullptr, &err);
178 test_error(err, "clCreateBuffer failed.");
179 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
180 sizeof(Ty) * num_elements, nullptr, &err);
181 test_error(err, "clCreateBuffer failed.");
182
183 generate_random_inputs(input);
184
185 err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0,
186 sizeof(Tx) * num_elements, input.data(), 0,
187 nullptr, nullptr);
188 test_error(err, "clEnqueueWriteBuffer failed.");
189
190 std::string src_name = type2name[sizeof(Tx)].first;
191 std::string dst_name = type2name[sizeof(Tx)].second;
192 if (int2fp) std::swap(src_name, dst_name);
193
194 std::string build_options;
195 build_options.append("-DTYPE_X=").append(src_name.c_str());
196 build_options.append(" -DTYPE_Y=").append(dst_name.c_str());
197
198 std::string extension;
199 if (sizeof(Tx) == sizeof(cl_double))
200 extension = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
201
202 if (sizeof(Tx) == sizeof(cl_half))
203 extension = "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
204
205 std::string kernelSource =
206 str_sprintf(int2float_kernel_code, extension.c_str());
207 const char *ptr = kernelSource.c_str();
208
209 err = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
210 "test_X2Y", build_options.c_str());
211 test_error(err, "create_single_kernel_helper failed.");
212
213 err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]);
214 err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
215 test_error(err, "clSetKernelArg failed.");
216
217 size_t threads[] = { (size_t)num_elements };
218 err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, threads,
219 nullptr, 0, nullptr, nullptr);
220 test_error(err, "clEnqueueNDRangeKernel failed.");
221
222 err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
223 sizeof(Ty) * num_elements, output.data(), 0,
224 nullptr, nullptr);
225 test_error(err, "clEnqueueReadBuffer failed.");
226
227 err = verify_X2Y(input, output);
228
229 return err;
230 }
231
skip_type__anon70898bdb0111::TypesIterator232 template <typename T> bool skip_type()
233 {
234 if (std::is_same<double, T>::value && !fp64Support)
235 return true;
236 else if (std::is_same<cl_half, T>::value && !fp16Support)
237 return true;
238 return false;
239 }
240
iterate_type__anon70898bdb0111::TypesIterator241 template <std::size_t Cnt = 0, typename T> void iterate_type(const T &t)
242 {
243 bool doTest = !skip_type<T>();
244
245 if (doTest)
246 {
247 typedef typename std::conditional<
248 (sizeof(T) == sizeof(std::int16_t)), std::int16_t,
249 typename std::conditional<(sizeof(T) == sizeof(std::int32_t)),
250 std::int32_t,
251 std::int64_t>::type>::type U;
252 if (int2fp)
253 {
254 if (test_X2Y<U, T>())
255 throw std::runtime_error("test_X2Y failed\n");
256 }
257 else
258 {
259 if (test_X2Y<T, U>())
260 throw std::runtime_error("test_X2Y failed\n");
261 }
262 }
263 }
264
265 template <std::size_t Cnt = 0, typename... Tp>
266 inline typename std::enable_if<Cnt == sizeof...(Tp), void>::type
for_each_elem__anon70898bdb0111::TypesIterator267 for_each_elem(
268 const std::tuple<Tp...> &) // Unused arguments are given no names.
269 {}
270
271 template <std::size_t Cnt = 0, typename... Tp>
272 inline typename std::enable_if < Cnt<sizeof...(Tp), void>::type
273 for_each_elem(const std::tuple<Tp...> &t)
274 {
275 iterate_type<Cnt>(std::get<Cnt>(t));
276 for_each_elem<Cnt + 1, Tp...>(t);
277 }
278
279 protected:
280 cl_context context;
281 cl_command_queue queue;
282
283 cl_device_fp_config fpConfigHalf;
284 cl_device_fp_config fpConfigFloat;
285
286 bool fp16Support;
287 bool fp64Support;
288
289 std::map<size_t, std::pair<std::string, std::string>> type2name;
290
291 std::string test_name;
292 int num_elements;
293 };
294
295 }
296
test_int2fp(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)297 int test_int2fp(cl_device_id device, cl_context context, cl_command_queue queue,
298 int num_elements)
299 {
300 try
301 {
302 TypesIterator<true>(device, context, queue, num_elements, "INT2FP");
303 } catch (const std::runtime_error &e)
304 {
305 log_error("%s", e.what());
306 return TEST_FAIL;
307 }
308
309 return TEST_PASS;
310 }
311
test_fp2int(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)312 int test_fp2int(cl_device_id device, cl_context context, cl_command_queue queue,
313 int num_elements)
314 {
315 try
316 {
317 TypesIterator<false>(device, context, queue, num_elements, "FP2INT");
318 } catch (const std::runtime_error &e)
319 {
320 log_error("%s", e.what());
321 return TEST_FAIL;
322 }
323
324 return TEST_PASS;
325 }
326