DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins_relational.cpp
Go to the documentation of this file.
1 //==------ builtins_relational.cpp - SYCL built-in relational functions ----==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 // This file defines the host versions of functions defined
10 // in SYCL SPEC section - 4.13.7 Relational functions.
11 
12 #include "builtins_helper.hpp"
13 #include <CL/sycl/detail/stl_type_traits.hpp>
14 
15 #include <cmath>
16 
17 namespace s = cl::sycl;
18 namespace d = s::detail;
19 
20 __SYCL_INLINE_NAMESPACE(cl) {
21 namespace __host_std {
22 namespace {
23 
24 template <typename T> inline T __vFOrdEqual(T x, T y) {
25  return -static_cast<T>(x == y);
26 }
27 
28 template <typename T> inline T __sFOrdEqual(T x, T y) { return x == y; }
29 
30 template <typename T> inline T __vFUnordNotEqual(T x, T y) {
31  return -static_cast<T>(x != y);
32 }
33 
34 template <typename T> inline T __sFUnordNotEqual(T x, T y) { return x != y; }
35 
36 template <typename T> inline T __vFOrdGreaterThan(T x, T y) {
37  return -static_cast<T>(x > y);
38 }
39 
40 template <typename T> inline T __sFOrdGreaterThan(T x, T y) { return x > y; }
41 
42 template <typename T> inline T __vFOrdGreaterThanEqual(T x, T y) {
43  return -static_cast<T>(x >= y);
44 }
45 
46 template <typename T> inline T __sFOrdGreaterThanEqual(T x, T y) {
47  return x >= y;
48 }
49 
50 template <typename T> inline T __vFOrdLessThanEqual(T x, T y) {
51  return -static_cast<T>(x <= y);
52 }
53 
54 template <typename T> inline T __sFOrdLessThanEqual(T x, T y) { return x <= y; }
55 
56 template <typename T> inline T __vFOrdNotEqual(T x, T y) {
57  return -static_cast<T>((x < y) || (x > y));
58 }
59 
60 template <typename T> inline T __sFOrdNotEqual(T x, T y) {
61  return ((x < y) || (x > y));
62 }
63 
64 template <typename T> inline T __vLessOrGreater(T x, T y) {
65  return -static_cast<T>((x < y) || (x > y));
66 }
67 
68 template <typename T> inline T __sLessOrGreater(T x, T y) {
69  return ((x < y) || (x > y));
70 }
71 
72 template <typename T> s::cl_int inline __Any(T x) { return d::msbIsSet(x); }
73 template <typename T> s::cl_int inline __All(T x) { return d::msbIsSet(x); }
74 
75 template <typename T> inline T __vOrdered(T x, T y) {
76  return -static_cast<T>(
77  !(std::isunordered(d::cast_if_host_half(x), d::cast_if_host_half(y))));
78 }
79 
80 template <typename T> inline T __sOrdered(T x, T y) {
81  return !(std::isunordered(d::cast_if_host_half(x), d::cast_if_host_half(y)));
82 }
83 
84 template <typename T> inline T __vUnordered(T x, T y) {
85  return -(static_cast<T>(
86  std::isunordered(d::cast_if_host_half(x), d::cast_if_host_half(y))));
87 }
88 
89 template <typename T> inline T __sUnordered(T x, T y) {
90  return std::isunordered(d::cast_if_host_half(x), d::cast_if_host_half(y));
91 }
92 
93 template <typename T>
94 inline typename sycl::detail::enable_if_t<d::is_sgeninteger<T>::value, T>
95 __bitselect(T a, T b, T c) {
96  return (a & ~c) | (b & c);
97 }
98 
99 template <typename T> union databitset;
100 // cl_float
101 template <> union databitset<s::cl_float> {
102  static_assert(sizeof(s::cl_int) == sizeof(s::cl_float),
103  "size of cl_float is not equal to 32 bits(cl_int).");
104  s::cl_float f;
105  s::cl_int i;
106 };
107 
108 // cl_double
109 template <> union databitset<s::cl_double> {
110  static_assert(sizeof(s::cl_long) == sizeof(s::cl_double),
111  "size of cl_double is not equal to 64 bits(cl_long).");
112  s::cl_double f;
113  s::cl_long i;
114 };
115 
116 // cl_half
117 template <> union databitset<s::cl_half> {
118  static_assert(sizeof(s::cl_short) == sizeof(s::cl_half),
119  "size of cl_half is not equal to 16 bits(cl_short).");
120  s::cl_half f;
121  s::cl_short i;
122 };
123 
124 template <typename T>
125 typename sycl::detail::enable_if_t<d::is_sgenfloat<T>::value,
126  T> inline __bitselect(T a, T b, T c) {
127  databitset<T> ba;
128  ba.f = a;
129  databitset<T> bb;
130  bb.f = b;
131  databitset<T> bc;
132  bc.f = c;
133  databitset<T> br;
134  br.f = 0;
135  br.i = ((ba.i & ~bc.i) | (bb.i & bc.i));
136  return br.f;
137 }
138 
139 template <typename T, typename T2> inline T2 __select(T2 a, T2 b, T c) {
140  return (c ? b : a);
141 }
142 
143 template <typename T, typename T2> inline T2 __vselect(T2 a, T2 b, T c) {
144  return d::msbIsSet(c) ? b : a;
145 }
146 } // namespace
147 
148 // ---------- 4.13.7 Relational functions. Host implementations. ---------------
149 // FOrdEqual-isequal
150 __SYCL_EXPORT s::cl_int FOrdEqual(s::cl_float x, s::cl_float y) __NOEXC {
151  return __sFOrdEqual(x, y);
152 }
153 __SYCL_EXPORT s::cl_int FOrdEqual(s::cl_double x, s::cl_double y) __NOEXC {
154  return __sFOrdEqual(x, y);
155 }
156 __SYCL_EXPORT s::cl_int FOrdEqual(s::cl_half x, s::cl_half y) __NOEXC {
157  return __sFOrdEqual(x, y);
158 }
159 MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_int, s::cl_float, s::cl_float)
160 MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_long, s::cl_double, s::cl_double)
161 MAKE_1V_2V_FUNC(FOrdEqual, __vFOrdEqual, s::cl_short, s::cl_half, s::cl_half)
162 
163 // FUnordNotEqual-isnotequal
164 __SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_float x, s::cl_float y) __NOEXC {
165  return __sFUnordNotEqual(x, y);
166 }
167 __SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_double x, s::cl_double y) __NOEXC {
168  return __sFUnordNotEqual(x, y);
169 }
170 __SYCL_EXPORT s::cl_int FUnordNotEqual(s::cl_half x, s::cl_half y) __NOEXC {
171  return __sFUnordNotEqual(x, y);
172 }
173 MAKE_1V_2V_FUNC(FUnordNotEqual, __vFUnordNotEqual, s::cl_int, s::cl_float,
174  s::cl_float)
175 MAKE_1V_2V_FUNC(FUnordNotEqual, __vFUnordNotEqual, s::cl_long, s::cl_double,
176  s::cl_double)
177 MAKE_1V_2V_FUNC(FUnordNotEqual, __vFUnordNotEqual, s::cl_short, s::cl_half,
178  s::cl_half)
179 
180 // (FOrdGreaterThan) // isgreater
181 __SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_float x, s::cl_float y) __NOEXC {
182  return __sFOrdGreaterThan(x, y);
183 }
184 __SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_double x,
185  s::cl_double y) __NOEXC {
186  return __sFOrdGreaterThan(x, y);
187 }
188 __SYCL_EXPORT s::cl_int FOrdGreaterThan(s::cl_half x, s::cl_half y) __NOEXC {
189  return __sFOrdGreaterThan(x, y);
190 }
191 MAKE_1V_2V_FUNC(FOrdGreaterThan, __vFOrdGreaterThan, s::cl_int, s::cl_float,
192  s::cl_float)
193 MAKE_1V_2V_FUNC(FOrdGreaterThan, __vFOrdGreaterThan, s::cl_long, s::cl_double,
194  s::cl_double)
195 MAKE_1V_2V_FUNC(FOrdGreaterThan, __vFOrdGreaterThan, s::cl_short, s::cl_half,
196  s::cl_half)
197 
198 // (FOrdGreaterThanEqual) // isgreaterequal
199 __SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_float x,
200  s::cl_float y) __NOEXC {
201  return __sFOrdGreaterThanEqual(x, y);
202 }
203 __SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_double x,
204  s::cl_double y) __NOEXC {
205  return __sFOrdGreaterThanEqual(x, y);
206 }
207 __SYCL_EXPORT s::cl_int FOrdGreaterThanEqual(s::cl_half x,
208  s::cl_half y) __NOEXC {
209  return __sFOrdGreaterThanEqual(x, y);
210 }
211 MAKE_1V_2V_FUNC(FOrdGreaterThanEqual, __vFOrdGreaterThanEqual, s::cl_int,
212  s::cl_float, s::cl_float)
213 MAKE_1V_2V_FUNC(FOrdGreaterThanEqual, __vFOrdGreaterThanEqual, s::cl_long,
214  s::cl_double, s::cl_double)
215 MAKE_1V_2V_FUNC(FOrdGreaterThanEqual, __vFOrdGreaterThanEqual, s::cl_short,
216  s::cl_half, s::cl_half)
217 
218 // (FOrdLessThan) // isless
219 __SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC {
220  return (x < y);
221 }
222 __SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_double x, s::cl_double y) __NOEXC {
223  return (x < y);
224 }
225 __SYCL_EXPORT s::cl_int __vFOrdLessThan(s::cl_float x, s::cl_float y) __NOEXC {
226  return -(x < y);
227 }
228 __SYCL_EXPORT s::cl_long __vFOrdLessThan(s::cl_double x,
229  s::cl_double y) __NOEXC {
230  return -(x < y);
231 }
232 __SYCL_EXPORT s::cl_int FOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC {
233  return (x < y);
234 }
235 __SYCL_EXPORT s::cl_short __vFOrdLessThan(s::cl_half x, s::cl_half y) __NOEXC {
236  return -static_cast<s::cl_short>(x < y);
237 }
238 MAKE_1V_2V_FUNC(FOrdLessThan, __vFOrdLessThan, s::cl_int, s::cl_float,
239  s::cl_float)
240 MAKE_1V_2V_FUNC(FOrdLessThan, __vFOrdLessThan, s::cl_long, s::cl_double,
241  s::cl_double)
242 MAKE_1V_2V_FUNC(FOrdLessThan, __vFOrdLessThan, s::cl_short, s::cl_half,
243  s::cl_half)
244 
245 // (FOrdLessThanEqual) // islessequal
246 __SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_float x,
247  s::cl_float y) __NOEXC {
248  return __sFOrdLessThanEqual(x, y);
249 }
250 __SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_double x,
251  s::cl_double y) __NOEXC {
252  return __sFOrdLessThanEqual(x, y);
253 }
254 __SYCL_EXPORT s::cl_int FOrdLessThanEqual(s::cl_half x, s::cl_half y) __NOEXC {
255  return __sFOrdLessThanEqual(x, y);
256 }
257 MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_int, s::cl_float,
258  s::cl_float)
259 MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_long,
260  s::cl_double, s::cl_double)
261 MAKE_1V_2V_FUNC(FOrdLessThanEqual, __vFOrdLessThanEqual, s::cl_short,
262  s::cl_half, s::cl_half)
263 
264 // (FOrdNotEqual) // islessgreater
265 __SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_float x, s::cl_float y) __NOEXC {
266  return __sFOrdNotEqual(x, y);
267 }
268 __SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_double x, s::cl_double y) __NOEXC {
269  return __sFOrdNotEqual(x, y);
270 }
271 __SYCL_EXPORT s::cl_int FOrdNotEqual(s::cl_half x, s::cl_half y) __NOEXC {
272  return __sFOrdNotEqual(x, y);
273 }
274 MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_int, s::cl_float,
275  s::cl_float)
276 MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_long, s::cl_double,
277  s::cl_double)
278 MAKE_1V_2V_FUNC(FOrdNotEqual, __vFOrdNotEqual, s::cl_short, s::cl_half,
279  s::cl_half)
280 
281 // (LessOrGreater) // islessgreater
282 __SYCL_EXPORT s::cl_int LessOrGreater(s::cl_float x, s::cl_float y) __NOEXC {
283  return __sLessOrGreater(x, y);
284 }
285 __SYCL_EXPORT s::cl_int LessOrGreater(s::cl_double x, s::cl_double y) __NOEXC {
286  return __sLessOrGreater(x, y);
287 }
288 __SYCL_EXPORT s::cl_int LessOrGreater(s::cl_half x, s::cl_half y) __NOEXC {
289  return __sLessOrGreater(x, y);
290 }
291 MAKE_1V_2V_FUNC(LessOrGreater, __vLessOrGreater, s::cl_int, s::cl_float,
292  s::cl_float)
293 MAKE_1V_2V_FUNC(LessOrGreater, __vLessOrGreater, s::cl_long, s::cl_double,
294  s::cl_double)
295 MAKE_1V_2V_FUNC(LessOrGreater, __vLessOrGreater, s::cl_short, s::cl_half,
296  s::cl_half)
297 
298 // (IsFinite) // isfinite
299 __SYCL_EXPORT s::cl_int IsFinite(s::cl_float x) __NOEXC {
300  return std::isfinite(x);
301 }
302 __SYCL_EXPORT s::cl_int IsFinite(s::cl_double x) __NOEXC {
303  return std::isfinite(x);
304 }
305 __SYCL_EXPORT s::cl_int __vIsFinite(s::cl_float x) __NOEXC {
306  return -static_cast<s::cl_int>(std::isfinite(x));
307 }
308 __SYCL_EXPORT s::cl_long __vIsFinite(s::cl_double x) __NOEXC {
309  return -static_cast<s::cl_long>(std::isfinite(x));
310 }
311 __SYCL_EXPORT s::cl_int IsFinite(s::cl_half x) __NOEXC {
312  return std::isfinite(d::cast_if_host_half(x));
313 }
314 __SYCL_EXPORT s::cl_short __vIsFinite(s::cl_half x) __NOEXC {
315  return -static_cast<s::cl_int>(std::isfinite(d::cast_if_host_half(x)));
316 }
317 MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_int, s::cl_float)
318 MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_long, s::cl_double)
319 MAKE_1V_FUNC(IsFinite, __vIsFinite, s::cl_short, s::cl_half)
320 
321 // (IsInf) // isinf
322 __SYCL_EXPORT s::cl_int IsInf(s::cl_float x) __NOEXC { return std::isinf(x); }
323 __SYCL_EXPORT s::cl_int IsInf(s::cl_double x) __NOEXC { return std::isinf(x); }
324 __SYCL_EXPORT s::cl_int __vIsInf(s::cl_float x) __NOEXC {
325  return -static_cast<s::cl_int>(std::isinf(x));
326 }
327 __SYCL_EXPORT s::cl_long __vIsInf(s::cl_double x) __NOEXC {
328  return -static_cast<s::cl_long>(std::isinf(x));
329 }
330 __SYCL_EXPORT s::cl_int IsInf(s::cl_half x) __NOEXC {
331  return std::isinf(d::cast_if_host_half(x));
332 }
333 __SYCL_EXPORT s::cl_short __vIsInf(s::cl_half x) __NOEXC {
334  return -static_cast<s::cl_short>(std::isinf(d::cast_if_host_half(x)));
335 }
336 MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_int, s::cl_float)
337 MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_long, s::cl_double)
338 MAKE_1V_FUNC(IsInf, __vIsInf, s::cl_short, s::cl_half)
339 
340 // (IsNan) // isnan
341 __SYCL_EXPORT s::cl_int IsNan(s::cl_float x) __NOEXC { return std::isnan(x); }
342 __SYCL_EXPORT s::cl_int IsNan(s::cl_double x) __NOEXC { return std::isnan(x); }
343 __SYCL_EXPORT s::cl_int __vIsNan(s::cl_float x) __NOEXC {
344  return -static_cast<s::cl_int>(std::isnan(x));
345 }
346 __SYCL_EXPORT s::cl_long __vIsNan(s::cl_double x) __NOEXC {
347  return -static_cast<s::cl_long>(std::isnan(x));
348 }
349 
350 __SYCL_EXPORT s::cl_int IsNan(s::cl_half x) __NOEXC {
351  return std::isnan(d::cast_if_host_half(x));
352 }
353 __SYCL_EXPORT s::cl_short __vIsNan(s::cl_half x) __NOEXC {
354  return -static_cast<s::cl_short>(std::isnan(d::cast_if_host_half(x)));
355 }
356 MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_int, s::cl_float)
357 MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_long, s::cl_double)
358 MAKE_1V_FUNC(IsNan, __vIsNan, s::cl_short, s::cl_half)
359 
360 // (IsNormal) // isnormal
361 __SYCL_EXPORT s::cl_int IsNormal(s::cl_float x) __NOEXC {
362  return std::isnormal(x);
363 }
364 __SYCL_EXPORT s::cl_int IsNormal(s::cl_double x) __NOEXC {
365  return std::isnormal(x);
366 }
367 __SYCL_EXPORT s::cl_int __vIsNormal(s::cl_float x) __NOEXC {
368  return -static_cast<s::cl_int>(std::isnormal(x));
369 }
370 __SYCL_EXPORT s::cl_long __vIsNormal(s::cl_double x) __NOEXC {
371  return -static_cast<s::cl_long>(std::isnormal(x));
372 }
373 __SYCL_EXPORT s::cl_int IsNormal(s::cl_half x) __NOEXC {
374  return std::isnormal(d::cast_if_host_half(x));
375 }
376 __SYCL_EXPORT s::cl_short __vIsNormal(s::cl_half x) __NOEXC {
377  return -static_cast<s::cl_short>(std::isnormal(d::cast_if_host_half(x)));
378 }
379 MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_int, s::cl_float)
380 MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_long, s::cl_double)
381 MAKE_1V_FUNC(IsNormal, __vIsNormal, s::cl_short, s::cl_half)
382 
383 // (Ordered) // isordered
384 __SYCL_EXPORT s::cl_int Ordered(s::cl_float x, s::cl_float y) __NOEXC {
385  return __sOrdered(x, y);
386 }
387 __SYCL_EXPORT s::cl_int Ordered(s::cl_double x, s::cl_double y) __NOEXC {
388  return __sOrdered(x, y);
389 }
390 __SYCL_EXPORT s::cl_int Ordered(s::cl_half x, s::cl_half y) __NOEXC {
391  return __sOrdered(x, y);
392 }
393 MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_int, s::cl_float, s::cl_float)
394 MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_long, s::cl_double, s::cl_double)
395 MAKE_1V_2V_FUNC(Ordered, __vOrdered, s::cl_short, s::cl_half, s::cl_half)
396 
397 // (Unordered) // isunordered
398 __SYCL_EXPORT s::cl_int Unordered(s::cl_float x, s::cl_float y) __NOEXC {
399  return __sUnordered(x, y);
400 }
401 __SYCL_EXPORT s::cl_int Unordered(s::cl_double x, s::cl_double y) __NOEXC {
402  return __sUnordered(x, y);
403 }
404 __SYCL_EXPORT s::cl_int Unordered(s::cl_half x, s::cl_half y) __NOEXC {
405  return __sUnordered(x, y);
406 }
407 MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_int, s::cl_float, s::cl_float)
408 MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_long, s::cl_double, s::cl_double)
409 MAKE_1V_2V_FUNC(Unordered, __vUnordered, s::cl_short, s::cl_half, s::cl_half)
410 
411 // (SignBitSet) // signbit
412 __SYCL_EXPORT s::cl_int SignBitSet(s::cl_float x) __NOEXC {
413  return std::signbit(x);
414 }
415 __SYCL_EXPORT s::cl_int SignBitSet(s::cl_double x) __NOEXC {
416  return std::signbit(x);
417 }
418 __SYCL_EXPORT s::cl_int __vSignBitSet(s::cl_float x) __NOEXC {
419  return -static_cast<s::cl_int>(std::signbit(x));
420 }
421 __SYCL_EXPORT s::cl_long __vSignBitSet(s::cl_double x) __NOEXC {
422  return -static_cast<s::cl_long>(std::signbit(x));
423 }
424 __SYCL_EXPORT s::cl_int SignBitSet(s::cl_half x) __NOEXC {
425  return std::signbit(d::cast_if_host_half(x));
426 }
427 __SYCL_EXPORT s::cl_short __vSignBitSet(s::cl_half x) __NOEXC {
428  return -static_cast<s::cl_short>(std::signbit(d::cast_if_host_half(x)));
429 }
430 MAKE_1V_FUNC(SignBitSet, __vSignBitSet, s::cl_int, s::cl_float)
431 MAKE_1V_FUNC(SignBitSet, __vSignBitSet, s::cl_long, s::cl_double)
432 MAKE_1V_FUNC(SignBitSet, __vSignBitSet, s::cl_short, s::cl_half)
433 
434 // (Any) // any
435 MAKE_SR_1V_OR(Any, __Any, s::cl_int, s::cl_char)
436 MAKE_SR_1V_OR(Any, __Any, s::cl_int, s::cl_short)
437 MAKE_SR_1V_OR(Any, __Any, s::cl_int, s::cl_int)
438 MAKE_SR_1V_OR(Any, __Any, s::cl_int, s::cl_long)
439 
440 // (All) // all
441 MAKE_SR_1V_AND(All, __All, s::cl_int, s::cl_char)
442 MAKE_SR_1V_AND(All, __All, s::cl_int, s::cl_short)
443 MAKE_SR_1V_AND(All, __All, s::cl_int, s::cl_int)
444 MAKE_SR_1V_AND(All, __All, s::cl_int, s::cl_long)
445 
446 // (bitselect)
447 // Instantiate functions for the scalar types and vector types.
448 MAKE_SC_1V_2V_3V(bitselect, s::cl_float, s::cl_float, s::cl_float, s::cl_float)
449 MAKE_SC_1V_2V_3V(bitselect, s::cl_double, s::cl_double, s::cl_double,
450  s::cl_double)
451 MAKE_SC_1V_2V_3V(bitselect, s::cl_char, s::cl_char, s::cl_char, s::cl_char)
452 MAKE_SC_1V_2V_3V(bitselect, s::cl_uchar, s::cl_uchar, s::cl_uchar, s::cl_uchar)
453 MAKE_SC_1V_2V_3V(bitselect, s::cl_short, s::cl_short, s::cl_short, s::cl_short)
454 MAKE_SC_1V_2V_3V(bitselect, s::cl_ushort, s::cl_ushort, s::cl_ushort,
455  s::cl_ushort)
456 MAKE_SC_1V_2V_3V(bitselect, s::cl_int, s::cl_int, s::cl_int, s::cl_int)
457 MAKE_SC_1V_2V_3V(bitselect, s::cl_uint, s::cl_uint, s::cl_uint, s::cl_uint)
458 MAKE_SC_1V_2V_3V(bitselect, s::cl_long, s::cl_long, s::cl_long, s::cl_long)
459 MAKE_SC_1V_2V_3V(bitselect, s::cl_ulong, s::cl_ulong, s::cl_ulong, s::cl_ulong)
460 MAKE_SC_1V_2V_3V(bitselect, s::cl_half, s::cl_half, s::cl_half, s::cl_half)
461 
462 // (Select) // select
463 // for scalar: result = c ? b : a.
464 // for vector: result[i] = (MSB of c[i] is set)? b[i] : a[i]
465 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_float, s::cl_float,
466  s::cl_float, s::cl_int)
467 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_float, s::cl_float,
468  s::cl_float, s::cl_uint)
469 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_double, s::cl_double,
470  s::cl_double, s::cl_long)
471 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_double, s::cl_double,
472  s::cl_double, s::cl_ulong)
473 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_char, s::cl_char, s::cl_char,
474  s::cl_char)
475 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_char, s::cl_char, s::cl_char,
476  s::cl_uchar)
477 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_uchar, s::cl_uchar,
478  s::cl_uchar, s::cl_char)
479 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_uchar, s::cl_uchar,
480  s::cl_uchar, s::cl_uchar)
481 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_short, s::cl_short,
482  s::cl_short, s::cl_short)
483 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_short, s::cl_short,
484  s::cl_short, s::cl_ushort)
485 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_ushort, s::cl_ushort,
486  s::cl_ushort, s::cl_short)
487 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_ushort, s::cl_ushort,
488  s::cl_ushort, s::cl_ushort)
489 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_int, s::cl_int, s::cl_int,
490  s::cl_int)
491 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_int, s::cl_int, s::cl_int,
492  s::cl_uint)
493 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_uint, s::cl_uint, s::cl_uint,
494  s::cl_int)
495 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_uint, s::cl_uint, s::cl_uint,
496  s::cl_uint)
497 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_long, s::cl_long, s::cl_long,
498  s::cl_long)
499 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_long, s::cl_long, s::cl_long,
500  s::cl_ulong)
501 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_ulong, s::cl_ulong,
502  s::cl_ulong, s::cl_long)
503 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_ulong, s::cl_ulong,
504  s::cl_ulong, s::cl_ulong)
505 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_half, s::cl_half, s::cl_half,
506  s::cl_short)
507 MAKE_SC_FSC_1V_2V_3V_FV(select, __vselect, s::cl_half, s::cl_half, s::cl_half,
508  s::cl_ushort)
509 } // namespace __host_std
510 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::cast_if_host_half
T cast_if_host_half(T val)
Definition: half_type.hpp:617
cl::sycl::cl_double
double cl_double
Definition: aliases.hpp:89
cl::sycl::cl_half
half cl_half
Definition: aliases.hpp:87
cl::sycl::cl_uchar
std::uint8_t cl_uchar
Definition: aliases.hpp:80
MAKE_SR_1V_OR
#define MAKE_SR_1V_OR(Fun, Call, Ret, Arg1)
Definition: builtins_helper.hpp:188
cl::sycl::cl_long
std::int64_t cl_long
Definition: aliases.hpp:85
MAKE_SC_FSC_1V_2V_3V_FV
#define MAKE_SC_FSC_1V_2V_3V_FV(FunSc, FunV, Ret, Arg1, Arg2, Arg3)
Definition: builtins_helper.hpp:155
T
cl::sycl::select
detail::enable_if_t< detail::is_sgentype< T >::value, T > select(T a, T b, bool c) __NOEXC
Definition: builtins.hpp:1318
MAKE_SR_1V_AND
#define MAKE_SR_1V_AND(Fun, Call, Ret, Arg1)
Definition: builtins_helper.hpp:180
cl::sycl::ext::oneapi::experimental::matrix::matrix_use::b
@ b
cl::sycl::detail::msbIsSet
constexpr bool msbIsSet(const T x)
Definition: generic_type_traits.hpp:500
cl::sycl
Definition: access.hpp:14
cl::sycl::isnan
detail::common_rel_ret_t< T > isnan(T x) __NOEXC
Definition: builtins.hpp:1228
stl_type_traits.hpp
cl::sycl::cl_ushort
std::uint16_t cl_ushort
Definition: aliases.hpp:82
cl::sycl::detail::half_impl::half
Definition: half_type.hpp:329
cl::sycl::detail::Ordered
@ Ordered
Definition: queue_impl.hpp:52
cl::sycl::isunordered
detail::common_rel_ret_t< T > isunordered(T x, T y) __NOEXC
Definition: builtins.hpp:1264
__NOEXC
#define __NOEXC
Definition: builtins.hpp:18
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::signbit
detail::common_rel_ret_t< T > signbit(T x) __NOEXC
Definition: builtins.hpp:1276
cl::sycl::cl_char
std::int8_t cl_char
Definition: aliases.hpp:79
builtins_helper.hpp
cl::sycl::isinf
detail::common_rel_ret_t< T > isinf(T x) __NOEXC
Definition: builtins.hpp:1216
cl::sycl::cl_ulong
std::uint64_t cl_ulong
Definition: aliases.hpp:86
cl::sycl::cl_float
float cl_float
Definition: aliases.hpp:88
cl::sycl::image_channel_order::a
@ a
cl::__host_std::__vFUnordNotEqual
__vFUnordNotEqual
Definition: builtins_relational.cpp:175
MAKE_1V_2V_FUNC
#define MAKE_1V_2V_FUNC(Fun, Call, Ret, Arg1, Arg2)
Definition: builtins_helper.hpp:132
cl::sycl::cl_int
std::int32_t cl_int
Definition: aliases.hpp:83
cl::sycl::cl_uint
std::uint32_t cl_uint
Definition: aliases.hpp:84
cl::__host_std::FOrdEqual
s::cl_int FOrdEqual(s::cl_half x, s::cl_half y) __NOEXC
Definition: builtins_relational.cpp:156
cl::__host_std::FUnordNotEqual
s::cl_int FUnordNotEqual(s::cl_half x, s::cl_half y) __NOEXC
Definition: builtins_relational.cpp:170
cl::sycl::isnormal
detail::common_rel_ret_t< T > isnormal(T x) __NOEXC
Definition: builtins.hpp:1240
cl::sycl::cl_short
std::int16_t cl_short
Definition: aliases.hpp:81
MAKE_1V_FUNC
#define MAKE_1V_FUNC(Fun, Call, Ret, Arg1)
Definition: builtins_helper.hpp:121
MAKE_SC_1V_2V_3V
#define MAKE_SC_1V_2V_3V(Fun, Ret, Arg1, Arg2, Arg3)
Definition: builtins_helper.hpp:151
cl::sycl::isfinite
detail::common_rel_ret_t< T > isfinite(T x) __NOEXC
Definition: builtins.hpp:1204
All
@ All
Definition: usm_allocator.hpp:14
cl::sycl::bitselect
detail::enable_if_t< detail::is_gentype< T >::value, T > bitselect(T a, T b, T c) __NOEXC
Definition: builtins.hpp:1311
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12