DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins_legacy_marray_vec.hpp
Go to the documentation of this file.
1 //==--- builtins_legacy_marray_vec.hpp - Old SYCL built-in nd definitions --==//
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 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
10 #error "Legacy builtins must not be used in preview."
11 #endif
12 
13 #pragma once
14 
15 #include <sycl/access/access.hpp> // for address_space, decorated
16 #include <sycl/aliases.hpp> // for half
17 #include <sycl/builtins_legacy_scalar.hpp> // for scalar builtin variants
18 #include <sycl/builtins_utils_vec.hpp> // for to_vec, to_marray...
19 #include <sycl/detail/boolean.hpp> // for Boolean
20 #include <sycl/detail/builtins.hpp> // for __invoke_select, __in...
21 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
22 #include <sycl/detail/generic_type_traits.hpp> // for is_svgenfloat, is_sge...
23 #include <sycl/detail/type_list.hpp> // for is_contained, type_list
24 #include <sycl/detail/type_traits.hpp> // for make_larger_t, marray...
25 #include <sycl/half_type.hpp> // for half, intel
26 #include <sycl/marray.hpp> // for marray
27 #include <sycl/multi_ptr.hpp> // for address_space_cast
28 #include <sycl/types.hpp> // for vec
29 
30 namespace sycl {
31 inline namespace _V1 {
32 
33 #ifdef __SYCL_DEVICE_ONLY__
34 #define __sycl_std
35 #else
36 namespace __sycl_std = __host_std;
37 #endif
38 
39 #ifdef __FAST_MATH__
40 #define __FAST_MATH_GENFLOAT(T) \
41  (detail::is_svgenfloatd_v<T> || detail::is_svgenfloath_v<T>)
42 #define __FAST_MATH_SGENFLOAT(T) \
43  (std::is_same_v<T, double> || std::is_same_v<T, half>)
44 #else
45 #define __FAST_MATH_GENFLOAT(T) (detail::is_svgenfloat_v<T>)
46 #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat_v<T>)
47 #endif
48 
49 /* ------------------ 4.13.3 Math functions. ---------------------------------*/
50 // These macros for marray math function implementations use vectorizations of
51 // size two as a simple general optimization. A more complex implementation
52 // using larger vectorizations for large marray sizes is possible; however more
53 // testing is required in order to ascertain the performance implications for
54 // all backends.
55 #define __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
56  marray<T, N> res; \
57  for (size_t i = 0; i < N / 2; i++) { \
58  vec<T, 2> partial_res = \
59  __sycl_std::__invoke_##NAME<vec<T, 2>>(detail::to_vec2(x, i * 2)); \
60  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
61  } \
62  if (N % 2) { \
63  res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1]); \
64  } \
65  return res;
66 
67 #define __SYCL_MATH_FUNCTION_OVERLOAD(NAME) \
68  template <typename T, size_t N> \
69  inline __SYCL_ALWAYS_INLINE \
70  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>> \
71  NAME(marray<T, N> x) { \
72  __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
73  }
74 
104 
105 #undef __SYCL_MATH_FUNCTION_OVERLOAD
106 
107 // __SYCL_MATH_FUNCTION_OVERLOAD_FM cases are replaced by corresponding native
108 // implementations when the -ffast-math flag is used with float.
109 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
110  template <typename T, size_t N> \
111  inline __SYCL_ALWAYS_INLINE \
112  std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray<T, N>> \
113  NAME(marray<T, N> x) { \
114  __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
115  }
116 
128 
129 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
130 #undef __SYCL_MATH_FUNCTION_OVERLOAD_IMPL
131 
132 template <typename T, size_t N>
134  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<int, N>>
136  marray<int, N> res;
137  for (size_t i = 0; i < N / 2; i++) {
138  vec<int, 2> partial_res =
139  __sycl_std::__invoke_ilogb<vec<int, 2>>(detail::to_vec2(x, i * 2));
140  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<int, 2>));
141  }
142  if (N % 2) {
143  res[N - 1] = __sycl_std::__invoke_ilogb<int>(x[N - 1]);
144  }
145  return res;
146 }
147 
148 #define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
149  marray<T, N> res; \
150  for (size_t i = 0; i < N / 2; i++) { \
151  auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
152  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
153  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
154  } \
155  if (N % 2) { \
156  res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1]); \
157  } \
158  return res;
159 
160 #define __SYCL_MATH_FUNCTION_2_OVERLOAD(NAME) \
161  template <typename T, size_t N> \
162  inline __SYCL_ALWAYS_INLINE \
163  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>> \
164  NAME(marray<T, N> x, marray<T, N> y) { \
165  __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
166  }
167 
181 
182 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD
183 
184 template <typename T, size_t N>
186  std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray<T, N>>
189 }
190 
191 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL
192 
193 #define __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(NAME) \
194  template <typename T, size_t N> \
195  inline __SYCL_ALWAYS_INLINE \
196  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>> \
197  NAME(marray<T, N> x, T y) { \
198  marray<T, N> res; \
199  sycl::vec<T, 2> y_vec{y, y}; \
200  for (size_t i = 0; i < N / 2; i++) { \
201  auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
202  detail::to_vec2(x, i * 2), y_vec); \
203  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
204  } \
205  if (N % 2) { \
206  res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y_vec[0]); \
207  } \
208  return res; \
209  }
210 
212 // clang-format off
214 
215 #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD
216 
217 template <typename T, size_t N>
219  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>>
221  // clang-format on
222  marray<T, N> res;
223  for (size_t i = 0; i < N; i++) {
224  res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k[i]);
225  }
226  return res;
227 }
228 
229 template <typename T, size_t N>
231  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>>
232  ldexp(marray<T, N> x, int k) {
233  marray<T, N> res;
234  for (size_t i = 0; i < N; i++) {
235  res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k);
236  }
237  return res;
238 }
239 
240 #define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME) \
241  marray<T, N> res; \
242  for (size_t i = 0; i < N; i++) { \
243  res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y[i]); \
244  } \
245  return res;
246 
247 template <typename T, size_t N>
249  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>>
252 }
253 
254 template <typename T, size_t N>
256  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>>
259 
260 #undef __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL
261 
262 #define __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(NAME) \
263  marray<T, N> res; \
264  for (size_t i = 0; i < N; i++) { \
265  res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y); \
266  } \
267  return res;
268 
269 template <typename T, size_t N>
270 inline __SYCL_ALWAYS_INLINE std::enable_if_t<detail::is_sgenfloat_v<T>,
272  int y) {
274 }
275 
276 template <typename T, size_t N>
278  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>>
281 }
282 
283 #undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL
284 
285 #define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME) \
286  template <typename T, size_t N> \
287  inline __SYCL_ALWAYS_INLINE \
288  std::enable_if_t<detail::is_sgenfloat_v<T>, marray<T, N>> \
289  NAME(marray<T, N> x, marray<T, N> y, marray<T, N> z) { \
290  marray<T, N> res; \
291  for (size_t i = 0; i < N / 2; i++) { \
292  auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
293  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2), \
294  detail::to_vec2(z, i * 2)); \
295  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
296  } \
297  if (N % 2) { \
298  res[N - 1] = \
299  __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1], z[N - 1]); \
300  } \
301  return res; \
302  }
303 
307 
308 #undef __SYCL_MATH_FUNCTION_3_OVERLOAD
309 
310 // svgenfloat fmax (svgenfloat x, sgenfloat y)
311 template <typename T>
312 std::enable_if_t<detail::is_vgenfloat_v<T>, T>
313 fmax(T x, typename T::element_type y) {
314  return __sycl_std::__invoke_fmax<T>(x, T(y));
315 }
316 
317 // svgenfloat fmin (svgenfloat x, sgenfloat y)
318 template <typename T>
319 std::enable_if_t<detail::is_vgenfloat_v<T>, T>
320 fmin(T x, typename T::element_type y) {
321  return __sycl_std::__invoke_fmin<T>(x, T(y));
322 }
323 
324 // vgenfloat ldexp (vgenfloat x, int k)
325 template <typename T>
326 std::enable_if_t<detail::is_vgenfloat_v<T>, T> ldexp(T x, int k) {
327  return __sycl_std::__invoke_ldexp<T>(x, vec<int, T::size()>(k));
328 }
329 
330 // vgenfloat ldexp (vgenfloat x, genint k)
331 template <typename T, typename T2>
332 std::enable_if_t<detail::is_vgenfloat_v<T> && detail::is_intn_v<T2>, T>
333 ldexp(T x, T2 k) {
334  detail::check_vector_size<T, T2>();
335  return __sycl_std::__invoke_ldexp<T>(x, k);
336 }
337 
338 // other marray math functions
339 
340 // TODO: can be optimized in the way marray math functions above are optimized
341 // (usage of vec<T, 2>)
342 #define __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARGPTR, \
343  ...) \
344  marray<T, N> res; \
345  for (int j = 0; j < N; j++) { \
346  res[j] = \
347  NAME(__VA_ARGS__, \
348  address_space_cast<AddressSpace, IsDecorated, \
349  detail::marray_element_t<T2>>(&(*ARGPTR)[j])); \
350  } \
351  return res;
352 
353 #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \
354  NAME, ARG1, ARG2, ...) \
355  template <typename T, size_t N, typename T2, \
356  access::address_space AddressSpace, access::decorated IsDecorated> \
357  std::enable_if_t< \
358  detail::is_svgenfloat_v<T> && \
359  detail::is_genfloatptr_marray_v<T2, AddressSpace, IsDecorated>, \
360  marray<T, N>> \
361  NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) { \
362  __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
363  __VA_ARGS__) \
364  }
365 
367  x[j])
369  x[j])
371  cosval, x[j])
372 
373 #undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENFLOATPTR_OVERLOAD
374 
375 #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \
376  NAME, ARG1, ARG2, ...) \
377  template <typename T, size_t N, typename T2, \
378  access::address_space AddressSpace, access::decorated IsDecorated> \
379  std::enable_if_t< \
380  detail::is_svgenfloat_v<T> && \
381  detail::is_genintptr_marray_v<T2, AddressSpace, IsDecorated>, \
382  marray<T, N>> \
383  NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) { \
384  __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
385  __VA_ARGS__) \
386  }
387 
389  x[j])
391  x[j])
392 
393 #undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD
394 
395 #define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \
396  template <typename T, size_t N, typename T2, \
397  access::address_space AddressSpace, access::decorated IsDecorated> \
398  std::enable_if_t< \
399  detail::is_svgenfloat_v<T> && \
400  detail::is_genintptr_marray_v<T2, AddressSpace, IsDecorated>, \
401  marray<T, N>> \
402  NAME(marray<T, N> x, marray<T, N> y, \
403  multi_ptr<T2, AddressSpace, IsDecorated> quo) { \
404  __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \
405  __VA_ARGS__) \
406  }
407 
409 
410 #undef __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD
411 
412 #undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL
413 
414 template <typename T, size_t N>
415 std::enable_if_t<detail::is_nan_type_v<T> &&
416  detail::is_non_deprecated_nan_type_v<T>,
418 nan(marray<T, N> nancode) {
420  for (int j = 0; j < N; j++) {
421  res[j] = nan(nancode[j]);
422  }
423  return res;
424 }
425 template <typename T, size_t N>
427  "This is a deprecated argument type for SYCL nan built-in function.")
428 std::enable_if_t<detail::is_nan_type_v<T> &&
429  !detail::is_non_deprecated_nan_type_v<T>,
430  marray<detail::nan_return_t<T>, N>> nan(marray<T, N> nancode) {
432  for (int j = 0; j < N; j++) {
433  res[j] = nan(nancode[j]);
434  }
435  return res;
436 }
437 
438 /* --------------- 4.13.5 Common functions. ---------------------------------*/
439 // vgenfloath clamp (vgenfloath x, half minval, half maxval)
440 // vgenfloatf clamp (vgenfloatf x, float minval, float maxval)
441 // vgenfloatd clamp (vgenfloatd x, double minval, double maxval)
442 template <typename T>
443 std::enable_if_t<detail::is_vgenfloat_v<T>, T>
445  return __sycl_std::__invoke_fclamp<T>(x, T(minval), T(maxval));
446 }
447 
448 // vgenfloatf max (vgenfloatf x, float y)
449 // vgenfloatd max (vgenfloatd x, double y)
450 // vgenfloath max (vgenfloath x, half y)
451 template <typename T>
452 std::enable_if_t<detail::is_vgenfloat_v<T>, T>(max)(
453  T x, typename T::element_type y) {
454  return __sycl_std::__invoke_fmax_common<T>(x, T(y));
455 }
456 
457 // vgenfloatf min (vgenfloatf x, float y)
458 // vgenfloatd min (vgenfloatd x, double y)
459 // vgenfloath min (vgenfloath x, half y)
460 template <typename T>
461 std::enable_if_t<detail::is_vgenfloat_v<T>, T>(min)(
462  T x, typename T::element_type y) {
463  return __sycl_std::__invoke_fmin_common<T>(x, T(y));
464 }
465 
466 // vgenfloatf mix (vgenfloatf x, vgenfloatf y, float a)
467 // vgenfloatd mix (vgenfloatd x, vgenfloatd y, double a)
468 // vgenfloatd mix (vgenfloath x, vgenfloath y, half a)
469 template <typename T>
470 std::enable_if_t<detail::is_vgenfloat_v<T>, T> mix(T x, T y,
471  typename T::element_type a) {
472  return __sycl_std::__invoke_mix<T>(x, y, T(a));
473 }
474 
475 // vgenfloatf step (float edge, vgenfloatf x)
476 // vgenfloatd step (double edge, vgenfloatd x)
477 // vgenfloatd step (half edge, vgenfloath x)
478 template <typename T>
479 std::enable_if_t<detail::is_vgenfloat_v<T>, T>
480 step(typename T::element_type edge, T x) {
481  return __sycl_std::__invoke_step<T>(T(edge), x);
482 }
483 
484 // vgenfloatf smoothstep (float edge0, float edge1, vgenfloatf x)
485 // vgenfloatd smoothstep (double edge0, double edge1, vgenfloatd x)
486 // vgenfloath smoothstep (half edge0, half edge1, vgenfloath x)
487 template <typename T>
488 std::enable_if_t<detail::is_vgenfloat_v<T>, T>
490  T x) {
491  return __sycl_std::__invoke_smoothstep<T>(T(edge0), T(edge1), x);
492 }
493 
494 // marray common functions
495 
496 // TODO: can be optimized in the way math functions are optimized (usage of
497 // vec<T, 2>)
498 #define __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
499  T res; \
500  for (int i = 0; i < T::size(); i++) { \
501  res[i] = NAME(__VA_ARGS__); \
502  } \
503  return res;
504 
505 #define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \
506  template <typename T, \
507  typename = std::enable_if_t<detail::is_mgenfloat_v<T>>> \
508  T NAME(ARG) { \
509  __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
510  }
511 
515 
516 #undef __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD
517 
518 #define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ...) \
519  template <typename T, \
520  typename = std::enable_if_t<detail::is_mgenfloat_v<T>>> \
521  T NAME(ARG1, ARG2) { \
522  __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
523  }
524 
525 // min and max may be defined as macros, so we wrap them in parentheses to avoid
526 // errors.
529  detail::marray_element_t<T> y,
530  x[i], y)
533  detail::marray_element_t<T> y,
534  x[i], y)
535 __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i])
537  detail::marray_element_t<T> edge,
538  T x, edge, x[i])
539 
540 #undef __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD
541 
542 #define __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(NAME, ARG1, ARG2, ARG3, \
543  ...) \
544  template <typename T, \
545  typename = std::enable_if_t<detail::is_mgenfloat_v<T>>> \
546  T NAME(ARG1, ARG2, ARG3) { \
547  __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
548  }
549 
551  x[i], minval[i], maxval[i])
553  detail::marray_element_t<T> minval,
554  detail::marray_element_t<T> maxval,
555  x[i], minval, maxval)
557  a[i])
559  detail::marray_element_t<T> a,
560  x[i], y[i], a)
562  edge0[i], edge1[i], x[i])
564  detail::marray_element_t<T> edge0,
565  detail::marray_element_t<T> edge1,
566  T x, edge0, edge1, x[i])
567 
568 #undef __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD
569 #undef __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL
570 
571 /* --------------- 4.13.4 Integer functions. --------------------------------*/
572 // igeninteger abs (geninteger x)
573 template <typename T>
574 std::enable_if_t<detail::is_igeninteger_v<T>, T> abs(T x) {
575  auto res = __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
576  if constexpr (detail::is_vigeninteger_v<T>) {
577  return res.template convert<detail::vector_element_t<T>>();
578  } else
579  return detail::make_signed_t<decltype(res)>(res);
580 }
581 
582 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
583 template <typename T>
584 std::enable_if_t<detail::is_vigeninteger_v<T>, T>
586  return __sycl_std::__invoke_s_clamp<T>(x, T(minval), T(maxval));
587 }
588 
589 // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval)
590 template <typename T>
591 std::enable_if_t<detail::is_vugeninteger_v<T>, T>
593  return __sycl_std::__invoke_u_clamp<T>(x, T(minval), T(maxval));
594 }
595 
596 // igeninteger max (vigeninteger x, sigeninteger y)
597 template <typename T>
598 std::enable_if_t<detail::is_vigeninteger_v<T>, T>(max)(
599  T x, typename T::element_type y) {
600  return __sycl_std::__invoke_s_max<T>(x, T(y));
601 }
602 
603 // vugeninteger max (vugeninteger x, sugeninteger y)
604 template <typename T>
605 std::enable_if_t<detail::is_vugeninteger_v<T>, T>(max)(
606  T x, typename T::element_type y) {
607  return __sycl_std::__invoke_u_max<T>(x, T(y));
608 }
609 
610 // vigeninteger min (vigeninteger x, sigeninteger y)
611 template <typename T>
612 std::enable_if_t<detail::is_vigeninteger_v<T>, T>(min)(
613  T x, typename T::element_type y) {
614  return __sycl_std::__invoke_s_min<T>(x, T(y));
615 }
616 
617 // vugeninteger min (vugeninteger x, sugeninteger y)
618 template <typename T>
619 std::enable_if_t<detail::is_vugeninteger_v<T>, T>(min)(
620  T x, typename T::element_type y) {
621  return __sycl_std::__invoke_u_min<T>(x, T(y));
622 }
623 
624 // marray integer functions
625 
626 // TODO: can be optimized in the way math functions are optimized (usage of
627 // vec<T, 2>)
628 #define __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
629  marray<T, N> res; \
630  for (int j = 0; j < N; j++) { \
631  res[j] = NAME(__VA_ARGS__); \
632  } \
633  return res;
634 
635 // Keep NAME for readability
636 #define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD(NAME, ARG, ...) \
637  template <typename T, size_t N> \
638  std::enable_if_t<detail::is_ugeninteger_v<T>, marray<T, N>> NAME( \
639  marray<T, N> ARG) { \
640  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
641  }
642 
643 #define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(NAME, ARG, ...) \
644  template <typename T, size_t N> \
645  std::enable_if_t<detail::is_igeninteger_v<T>, marray<T, N>> NAME( \
646  marray<T, N> ARG) { \
647  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
648  }
649 
652 
653 #undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD
654 #undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD
655 
656 #define __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \
657  template <typename T, size_t N> \
658  std::enable_if_t<detail::is_geninteger_v<T>, marray<T, N>> NAME( \
659  marray<T, N> ARG) { \
660  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
661  }
662 
666 
667 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD
668 
669 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
670  template <typename T, size_t N> \
671  std::enable_if_t<detail::is_ugeninteger_v<T>, marray<T, N>> NAME( \
672  marray<T, N> ARG1, marray<T, N> ARG2) { \
673  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
674  }
675 
676 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD(NAME, ARG1, \
677  ARG2, ...) \
678  template <typename T, size_t N> \
679  std::enable_if_t<detail::is_igeninteger_v<T>, \
680  marray<detail::make_unsigned_t<T>, N>> \
681  NAME(marray<T, N> ARG1, marray<T, N> ARG2) { \
682  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
683  }
684 
685 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
686  template <typename T, size_t N> \
687  std::enable_if_t<detail::is_igeninteger_v<T>, marray<T, N>> NAME( \
688  marray<T, N> ARG1, marray<T, N> ARG2) { \
689  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
690  }
691 
692 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD( \
693  NAME, ARG1, ARG2, ...) \
694  template <typename T, size_t N> \
695  std::enable_if_t<detail::is_ugeninteger_v<T>, marray<T, N>> NAME( \
696  marray<T, N> ARG1, T ARG2) { \
697  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
698  }
699 
700 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD( \
701  NAME, ARG1, ARG2, ...) \
702  template <typename T, size_t N> \
703  std::enable_if_t<detail::is_igeninteger_v<T>, marray<T, N>> NAME( \
704  marray<T, N> ARG1, T ARG2) { \
705  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
706  }
707 
710  y[j])
720  x[j], y)
722  x[j], y)
726  x[j], y)
728  x[j], y)
735 
736 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD
737 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD
738 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD
739 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD
740 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD
741 
742 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(NAME, ARG1, ARG2, \
743  ARG3, ...) \
744  template <typename T, size_t N> \
745  std::enable_if_t<detail::is_ugeninteger_v<T>, marray<T, N>> NAME( \
746  marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) { \
747  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
748  }
749 
750 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(NAME, ARG1, ARG2, \
751  ARG3, ...) \
752  template <typename T, size_t N> \
753  std::enable_if_t<detail::is_igeninteger_v<T>, marray<T, N>> NAME( \
754  marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) { \
755  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
756  }
757 
758 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
759  NAME, ARG1, ARG2, ARG3, ...) \
760  template <typename T, size_t N> \
761  std::enable_if_t<detail::is_ugeninteger_v<T>, marray<T, N>> NAME( \
762  marray<T, N> ARG1, T ARG2, T ARG3) { \
763  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
764  }
765 
766 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
767  NAME, ARG1, ARG2, ARG3, ...) \
768  template <typename T, size_t N> \
769  std::enable_if_t<detail::is_igeninteger_v<T>, marray<T, N>> NAME( \
770  marray<T, N> ARG1, T ARG2, T ARG3) { \
771  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
772  }
773 
775  minval[j], maxval[j])
777  minval[j], maxval[j])
779  clamp, x, minval, maxval, x[j], minval, maxval)
781  clamp, x, minval, maxval, x[j], minval, maxval)
783  c[j])
785  c[j])
787  c[j])
789  c[j])
790 
791 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD
792 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD
793 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD
794 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD
795 
796 // Keep NAME for readability
797 #define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD(NAME, ARG1, ARG2, \
798  ARG3, ...) \
799  template <typename T, size_t N> \
800  std::enable_if_t<detail::is_ugeninteger32bit_v<T>, marray<T, N>> NAME( \
801  marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) { \
802  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
803  }
804 
805 #define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD(NAME, ARG1, ARG2, \
806  ARG3, ...) \
807  template <typename T, size_t N> \
808  std::enable_if_t<detail::is_igeninteger32bit_v<T>, marray<T, N>> NAME( \
809  marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) { \
810  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
811  }
812 
814  z[j])
816  z[j])
817 
818 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD
819 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD
820 
821 // Keep NAME for readability
822 #define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
823  template <typename T, size_t N> \
824  std::enable_if_t<detail::is_ugeninteger32bit_v<T>, marray<T, N>> NAME( \
825  marray<T, N> ARG1, marray<T, N> ARG2) { \
826  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
827  }
828 
829 #define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
830  template <typename T, size_t N> \
831  std::enable_if_t<detail::is_igeninteger32bit_v<T>, marray<T, N>> NAME( \
832  marray<T, N> ARG1, marray<T, N> ARG2) { \
833  __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
834  }
835 
838 
839 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD
840 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD
841 #undef __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL
842 
843 // TODO: can be optimized in the way math functions are optimized (usage of
844 // vec<T, 2>)
845 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
846  detail::make_larger_t<marray<T, N>> res; \
847  for (int j = 0; j < N; j++) { \
848  res[j] = NAME(hi[j], lo[j]); \
849  } \
850  return res;
851 
852 // Keep NAME for readability
853 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(NAME, KBIT) \
854  template <typename T, size_t N> \
855  std::enable_if_t<detail::is_ugeninteger##KBIT##_v<T>, \
856  detail::make_larger_t<marray<T, N>>> \
857  NAME(marray<T, N> hi, marray<T, N> lo) { \
858  __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
859  }
860 
861 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(NAME, KBIT) \
862  template <typename T, typename T2, size_t N> \
863  std::enable_if_t<detail::is_igeninteger##KBIT##_v<T> && \
864  detail::is_ugeninteger##KBIT##_v<T2>, \
865  detail::make_larger_t<marray<T, N>>> \
866  NAME(marray<T, N> hi, marray<T2, N> lo) { \
867  __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
868  }
869 
876 
877 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD
878 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD
879 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL
880 
881 /* --------------- 4.13.6 Geometric Functions. ------------------------------*/
882 // float3 cross (float3 p0, float3 p1)
883 // float4 cross (float4 p0, float4 p1)
884 // double3 cross (double3 p0, double3 p1)
885 // double4 cross (double4 p0, double4 p1)
886 // half3 cross (half3 p0, half3 p1)
887 // half4 cross (half4 p0, half4 p1)
888 template <typename T>
889 std::enable_if_t<detail::is_gencross_v<T>, T> cross(T p0, T p1) {
890  return __sycl_std::__invoke_cross<T>(p0, p1);
891 }
892 
893 // float dot (vgengeofloat p0, vgengeofloat p1)
894 template <typename T>
895 std::enable_if_t<detail::is_vgengeofloat_v<T>, float> dot(T p0, T p1) {
896  return __sycl_std::__invoke_Dot<float>(p0, p1);
897 }
898 
899 // double dot (vgengeodouble p0, vgengeodouble p1)
900 template <typename T>
901 std::enable_if_t<detail::is_vgengeodouble_v<T>, double> dot(T p0, T p1) {
902  return __sycl_std::__invoke_Dot<double>(p0, p1);
903 }
904 
905 // half dot (vgengeohalf p0, vgengeohalf p1)
906 template <typename T>
907 std::enable_if_t<detail::is_vgengeohalf_v<T>, half> dot(T p0, T p1) {
908  return __sycl_std::__invoke_Dot<half>(p0, p1);
909 }
910 
911 // float distance (gengeofloat p0, gengeofloat p1)
912 template <typename T,
913  typename = std::enable_if_t<detail::is_gengeofloat_v<T>, T>>
914 float distance(T p0, T p1) {
915  return __sycl_std::__invoke_distance<float>(p0, p1);
916 }
917 
918 // double distance (gengeodouble p0, gengeodouble p1)
919 template <typename T,
920  typename = std::enable_if_t<detail::is_gengeodouble_v<T>, T>>
921 double distance(T p0, T p1) {
922  return __sycl_std::__invoke_distance<double>(p0, p1);
923 }
924 
925 // half distance (gengeohalf p0, gengeohalf p1)
926 template <typename T,
927  typename = std::enable_if_t<detail::is_gengeohalf_v<T>, T>>
928 half distance(T p0, T p1) {
929  return __sycl_std::__invoke_distance<half>(p0, p1);
930 }
931 
932 // float length (gengeofloat p)
933 template <typename T,
934  typename = std::enable_if_t<detail::is_gengeofloat_v<T>, T>>
935 float length(T p) {
936  return __sycl_std::__invoke_length<float>(p);
937 }
938 
939 // double length (gengeodouble p)
940 template <typename T,
941  typename = std::enable_if_t<detail::is_gengeodouble_v<T>, T>>
942 double length(T p) {
943  return __sycl_std::__invoke_length<double>(p);
944 }
945 
946 // half length (gengeohalf p)
947 template <typename T,
948  typename = std::enable_if_t<detail::is_gengeohalf_v<T>, T>>
949 half length(T p) {
950  return __sycl_std::__invoke_length<half>(p);
951 }
952 
953 // gengeofloat normalize (gengeofloat p)
954 template <typename T>
955 std::enable_if_t<detail::is_gengeofloat_v<T>, T> normalize(T p) {
956  return __sycl_std::__invoke_normalize<T>(p);
957 }
958 
959 // gengeodouble normalize (gengeodouble p)
960 template <typename T>
961 std::enable_if_t<detail::is_gengeodouble_v<T>, T> normalize(T p) {
962  return __sycl_std::__invoke_normalize<T>(p);
963 }
964 
965 // gengeohalf normalize (gengeohalf p)
966 template <typename T>
967 std::enable_if_t<detail::is_gengeohalf_v<T>, T> normalize(T p) {
968  return __sycl_std::__invoke_normalize<T>(p);
969 }
970 
971 // float fast_distance (gengeofloat p0, gengeofloat p1)
972 template <typename T,
973  typename = std::enable_if_t<detail::is_gengeofloat_v<T>, T>>
974 float fast_distance(T p0, T p1) {
975  return __sycl_std::__invoke_fast_distance<float>(p0, p1);
976 }
977 
978 // double fast_distance (gengeodouble p0, gengeodouble p1)
979 template <typename T,
980  typename = std::enable_if_t<detail::is_gengeodouble_v<T>, T>>
981 __SYCL_DEPRECATED("fast_distance for double precision types is non-standard "
982  "and has been deprecated")
983 double fast_distance(T p0, T p1) {
984  return __sycl_std::__invoke_fast_distance<double>(p0, p1);
985 }
986 
987 // float fast_length (gengeofloat p)
988 template <typename T,
989  typename = std::enable_if_t<detail::is_gengeofloat_v<T>, T>>
990 float fast_length(T p) {
991  return __sycl_std::__invoke_fast_length<float>(p);
992 }
993 
994 // double fast_length (gengeodouble p)
995 template <typename T,
996  typename = std::enable_if_t<detail::is_gengeodouble_v<T>, T>>
997 __SYCL_DEPRECATED("fast_length for double precision types is non-standard "
998  "and has been deprecated")
999 double fast_length(T p) {
1000  return __sycl_std::__invoke_fast_length<double>(p);
1001 }
1002 
1003 // gengeofloat fast_normalize (gengeofloat p)
1004 template <typename T>
1005 std::enable_if_t<detail::is_gengeofloat_v<T>, T> fast_normalize(T p) {
1006  return __sycl_std::__invoke_fast_normalize<T>(p);
1007 }
1008 
1009 // gengeodouble fast_normalize (gengeodouble p)
1010 template <typename T>
1011 __SYCL_DEPRECATED("fast_normalize for double precision types is non-standard "
1012  "and has been deprecated")
1013 std::enable_if_t<detail::is_gengeodouble_v<T>, T> fast_normalize(T p) {
1014  return __sycl_std::__invoke_fast_normalize<T>(p);
1015 }
1016 
1017 // marray geometric functions
1018 
1019 #define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
1020  vec<detail::marray_element_t<T>, T::size()> result_v; \
1021  result_v = NAME(__VA_ARGS__); \
1022  return detail::to_marray(result_v);
1023 
1024 template <typename T>
1025 std::enable_if_t<detail::is_gencrossmarray_v<T>, T> cross(T p0, T p1) {
1027  detail::to_vec(p1))
1028 }
1029 
1030 template <typename T>
1031 std::enable_if_t<detail::is_gengeomarray_v<T>, T> normalize(T p) {
1033 }
1034 
1035 template <typename T>
1036 std::enable_if_t<detail::is_gengeomarrayfloat_v<T>, T> fast_normalize(T p) {
1038  detail::to_vec(p))
1039 }
1040 
1041 #undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL
1042 
1043 #define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME) \
1044  template <typename T> \
1045  std::enable_if_t<detail::is_gengeomarray_v<T>, detail::marray_element_t<T>> \
1046  NAME(T p0, T p1) { \
1047  return NAME(detail::to_vec(p0), detail::to_vec(p1)); \
1048  }
1049 
1050 // clang-format off
1053 // clang-format on
1054 
1055 #undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD
1056 
1057 template <typename T>
1058 std::enable_if_t<detail::is_gengeomarray_v<T>, detail::marray_element_t<T>>
1059 length(T p) {
1060  return __sycl_std::__invoke_length<detail::marray_element_t<T>>(
1061  detail::to_vec(p));
1062 }
1063 
1064 template <typename T>
1065 std::enable_if_t<detail::is_gengeomarrayfloat_v<T>, detail::marray_element_t<T>>
1066 fast_distance(T p0, T p1) {
1068 }
1069 
1070 template <typename T>
1071 std::enable_if_t<detail::is_gengeomarrayfloat_v<T>, detail::marray_element_t<T>>
1073  return fast_length(detail::to_vec(p));
1074 }
1075 
1076 /* SYCL 1.2.1 ---- 4.13.7 Relational functions. -----------------------------*/
1077 /* SYCL 2020 ---- 4.17.9 Relational functions. -----------------------------*/
1078 
1079 // marray relational functions
1080 
1081 #define __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(NAME) \
1082  template <typename T, \
1083  typename = std::enable_if_t<detail::is_mgenfloat_v<T>>> \
1084  sycl::marray<bool, T::size()> NAME(T x, T y) { \
1085  sycl::marray<bool, T::size()> res; \
1086  for (int i = 0; i < x.size(); i++) { \
1087  res[i] = NAME(x[i], y[i]); \
1088  } \
1089  return res; \
1090  }
1091 
1092 #define __SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(NAME) \
1093  template <typename T, \
1094  typename = std::enable_if_t<detail::is_mgenfloat_v<T>>> \
1095  sycl::marray<bool, T::size()> NAME(T x) { \
1096  sycl::marray<bool, T::size()> res; \
1097  for (int i = 0; i < x.size(); i++) { \
1098  res[i] = NAME(x[i]); \
1099  } \
1100  return res; \
1101  }
1102 
1117 
1118 // int any (vigeninteger x)
1119 template <typename T>
1120 std::enable_if_t<detail::is_vigeninteger_v<T>, int> any(T x) {
1121  return detail::rel_sign_bit_test_ret_t<T>(
1122  __sycl_std::__invoke_Any<detail::rel_sign_bit_test_ret_t<T>>(
1123  detail::rel_sign_bit_test_arg_t<T>(x)));
1124 }
1125 
1126 // int all (vigeninteger x)
1127 template <typename T>
1128 std::enable_if_t<detail::is_vigeninteger_v<T>, int> all(T x) {
1129  return detail::rel_sign_bit_test_ret_t<T>(
1130  __sycl_std::__invoke_All<detail::rel_sign_bit_test_ret_t<T>>(
1131  detail::rel_sign_bit_test_arg_t<T>(x)));
1132 }
1133 
1134 // other marray relational functions
1135 
1136 template <typename T, size_t N>
1137 std::enable_if_t<detail::is_sigeninteger_v<T>, bool> any(marray<T, N> x) {
1138  return std::any_of(x.begin(), x.end(), [](T i) { return any(i); });
1139 }
1140 
1141 template <typename T, size_t N>
1142 std::enable_if_t<detail::is_sigeninteger_v<T>, bool> all(marray<T, N> x) {
1143  return std::all_of(x.begin(), x.end(), [](T i) { return all(i); });
1144 }
1145 
1146 template <typename T, size_t N>
1147 std::enable_if_t<detail::is_gentype_v<T>, marray<T, N>>
1149  marray<T, N> res;
1150  for (int i = 0; i < N; i++) {
1151  res[i] = bitselect(a[i], b[i], c[i]);
1152  }
1153  return res;
1154 }
1155 
1156 template <typename T, size_t N>
1157 std::enable_if_t<detail::is_gentype_v<T>, marray<T, N>>
1159  marray<T, N> res;
1160  for (int i = 0; i < N; i++) {
1161  res[i] = select(a[i], b[i], c[i]);
1162  }
1163  return res;
1164 }
1165 
1166 namespace native {
1167 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1168 
1169 #define __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(NAME) \
1170  template <size_t N> \
1171  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) { \
1172  marray<float, N> res; \
1173  for (size_t i = 0; i < N / 2; i++) { \
1174  auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
1175  detail::to_vec2(x, i * 2)); \
1176  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1177  } \
1178  if (N % 2) { \
1179  res[N - 1] = __sycl_std::__invoke_native_##NAME<float>(x[N - 1]); \
1180  } \
1181  return res; \
1182  }
1183 
1196 
1197 #undef __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD
1198 
1199 #define __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(NAME) \
1200  template <size_t N> \
1201  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x, \
1202  marray<float, N> y) { \
1203  marray<float, N> res; \
1204  for (size_t i = 0; i < N / 2; i++) { \
1205  auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
1206  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
1207  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1208  } \
1209  if (N % 2) { \
1210  res[N - 1] = \
1211  __sycl_std::__invoke_native_##NAME<float>(x[N - 1], y[N - 1]); \
1212  } \
1213  return res; \
1214  }
1215 
1218 
1219 #undef __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD
1220 
1221 } // namespace native
1222 namespace half_precision {
1223 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
1224 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(NAME) \
1225  template <size_t N> \
1226  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) { \
1227  marray<float, N> res; \
1228  for (size_t i = 0; i < N / 2; i++) { \
1229  auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
1230  detail::to_vec2(x, i * 2)); \
1231  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1232  } \
1233  if (N % 2) { \
1234  res[N - 1] = __sycl_std::__invoke_half_##NAME<float>(x[N - 1]); \
1235  } \
1236  return res; \
1237  }
1238 
1251 
1252 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD
1253 
1254 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(NAME) \
1255  template <size_t N> \
1256  inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x, \
1257  marray<float, N> y) { \
1258  marray<float, N> res; \
1259  for (size_t i = 0; i < N / 2; i++) { \
1260  auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
1261  detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
1262  std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
1263  } \
1264  if (N % 2) { \
1265  res[N - 1] = \
1266  __sycl_std::__invoke_half_##NAME<float>(x[N - 1], y[N - 1]); \
1267  } \
1268  return res; \
1269  }
1270 
1273 
1274 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD
1275 
1276 } // namespace half_precision
1277 
1278 #ifdef __FAST_MATH__
1279 /* ----------------- -ffast-math functions. ---------------------------------*/
1280 
1281 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
1282  template <typename T, size_t N> \
1283  inline __SYCL_ALWAYS_INLINE \
1284  std::enable_if_t<std::is_same_v<T, float>, marray<T, N>> \
1285  NAME(marray<T, N> x) { \
1286  return native::NAME(x); \
1287  }
1288 
1300 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
1301 
1302 template <typename T, size_t N>
1303 inline __SYCL_ALWAYS_INLINE
1304  std::enable_if_t<std::is_same_v<T, float>, marray<T, N>>
1305  powr(marray<T, N> x, marray<T, N> y) {
1306  return native::powr(x, y);
1307 }
1308 
1309 #endif // __FAST_MATH__
1310 } // namespace _V1
1311 } // namespace sycl
#define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(NAME, KBIT)
#define __SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(NAME)
#define __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(NAME)
#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD(NAME, ARG,...)
#define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME)
#define __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(NAME)
#define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME)
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD( NAME, ARG1, ARG2, ARG3,...)
#define __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(NAME)
#define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(NAME, KBIT)
#define __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(NAME)
#define __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(NAME)
#define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD(NAME, ARG1, ARG2,...)
#define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME)
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(NAME, ARG1, ARG2,...)
#define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME,...)
#define __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(NAME)
#define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME)
#define __FAST_MATH_SGENFLOAT(T)
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(NAME, ARG1, ARG2,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD( NAME, ARG1, ARG2,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD(NAME, ARG1, ARG2, ARG3,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(NAME, ARG,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(NAME, ARG,...)
#define __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(NAME)
#define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(NAME, ARG1, ARG2, ARG3,...)
#define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME)
#define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARG,...)
#define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD(NAME, ARG1, ARG2,...)
#define __SYCL_MATH_FUNCTION_2_OVERLOAD(NAME)
#define __SYCL_MATH_FUNCTION_OVERLOAD(NAME)
Provides a cross-platform math array class template that works on SYCL devices as well as in host C++...
Definition: marray.hpp:48
defined(__INTEL_PREVIEW_BREAKING_CHANGES)
Definition: types.hpp:346
#define __SYCL_ALWAYS_INLINE
typename T::value_type marray_element_t
constexpr bool is_nan_type_v
vec< T, N > to_vec(marray< T, N > X)
constexpr bool is_gengeodouble_v
vec< T, 2 > to_vec2(marray< T, N > X, size_t Start)
typename nan_types< T, T >::ret_type nan_return_t
constexpr bool is_non_deprecated_nan_type_v
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > cos(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp10(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > sqrt(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp2(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log10(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > rsqrt(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > sin(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > divide(T x, T y)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > tan(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log2(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > recip(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > cos(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > recip(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log2(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > divide(T x, T y)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > sin(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp10(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > tan(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > log10(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > sqrt(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > rsqrt(T x)
std::enable_if_t< detail::is_svgenfloatf_v< T >, T > exp2(T x)
std::enable_if_t< detail::is_gentype_v< T >, marray< T, N > > select(marray< T, N > a, marray< T, N > b, marray< bool, N > c)
detail::common_rel_ret_t< T > isless(T x, T y)
std::enable_if_t< detail::is_vgenfloat_v< T >, T > fmin(T x, typename T::element_type y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > tgamma(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > cbrt(T x)
std::enable_if_t< detail::is_igeninteger_v< T >, T > sub_sat(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > logb(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > acos(T x)
__SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD(abs_diff, x, y, x[j], y[j]) __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD((max)
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > log10(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > rint(T x)
detail::common_rel_ret_t< T > isfinite(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > ceil(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > acosh(T x)
float distance(T p0, T p1)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > degrees(T radians)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > fdim(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > acospi(T x)
T detail::marray_element_t< T > y T T minval
y y maxval[j] maxval c[j] c[j] z[j] __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD(mad24, x, y, z, x[j], y[j], z[j]) template< typename T > std
return(x >=T(0)) ? T(std sinpi
detail::common_rel_ret_t< T > isunordered(T x, T y)
std::enable_if_t< detail::is_igeninteger_v< T >, T > mad_sat(T a, T b, T c)
x[j] x[j] size_t N std::enable_if_t< detail::is_nan_type_v< T > &&detail::is_non_deprecated_nan_type_v< T >, marray< detail::nan_return_t< T >, N > > nan(marray< T, N > nancode)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > asin(T x)
std::enable_if_t< __FAST_MATH_GENFLOAT(T) &&detail::is_genfloatptr_v< T2 >, T > sincos(T x, T2 cosval)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > lgamma(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > minmag(T x, T y)
T detail::marray_element_t< T > y T T T maxval
std::enable_if_t< detail::is_svgenfloat_v< T >, T > cospi(T x)
detail::common_rel_ret_t< T > islessgreater(T x, T y)
std::enable_if_t< detail::is_geninteger_v< T >, T > popcount(T x)
T detail::marray_element_t< T > y T T T maxval[i] T T T a[i] T T edge1
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > log2(T x)
std::enable_if_t< detail::is_svgenfloat_v< T > &&detail::is_genintptr_v< T2 >, T > remquo(T x, T y, T2 quo)
detail::common_rel_ret_t< T > signbit(T x)
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > sqrt(T x)
std::enable_if_t< detail::is_gencrossmarray_v< T >, T > cross(T p0, T p1)
std::enable_if_t< detail::is_geninteger_v< T >, T > rotate(T v, T i)
T detail::marray_element_t< T > y T T T maxval[i] T T T a
std::enable_if_t< detail::is_geninteger_v< T >, T > clz(T x)
y y maxval[j] maxval b
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > tan(T x)
std::enable_if_t< detail::is_vigeninteger_v< T >, int > all(T x)
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > log(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
__SYCL_ALWAYS_INLINE std::enable_if_t< detail::is_sgenfloat_v< T >, marray< T, N > > ldexp(marray< T, N > x, marray< int, N > k)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > expm1(T x)
std::enable_if_t< detail::is_svgenfloat_v< T > &&detail::is_genintptr_v< T2 >, T > lgamma_r(T x, T2 signp)
detail::common_rel_ret_t< T > isinf(T x)
detail::common_rel_ret_t< T > isnan(T x)
std::enable_if_t< detail::is_igeninteger32bit_v< T >, T > mad24(T x, T y, T z)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > mad(T a, T b, T c)
std::enable_if_t< detail::is_vgenfloat_v< T >, T > step(typename T::element_type edge, T x)
y y maxval[j] maxval c
float fast_distance(T p0, T p1)
std::enable_if_t< detail::is_gengeofloat_v< T >, T > fast_normalize(T p)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > pow(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > tanpi(T x)
std::enable_if_t< detail::is_igeninteger32bit_v< T >, T > mul24(T x, T y)
std::enable_if_t< detail::is_vgenfloat_v< T >, T > fmax(T x, typename T::element_type y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > sign(T x)
std::enable_if_t< detail::is_ugeninteger8bit_v< T >, detail::make_larger_t< T > > upsample(T hi, T lo)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > cosh(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > fabs(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > erf(T x)
std::enable_if_t< detail::is_gentype_v< T >, marray< T, N > > bitselect(marray< T, N > a, marray< T, N > b, marray< T, N > c)
T detail::marray_element_t< T > y T T T maxval[i] T T T a[i] T edge0
y y maxval[j] maxval c[j] c[j] z
std::enable_if_t< detail::is_svgenfloat_v< T >, T > erfc(T x)
__SYCL_ALWAYS_INLINE std::enable_if_t< detail::is_sgenfloat_v< T >, marray< T, N > > rootn(marray< T, N > x, marray< int, N > y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > fma(T a, T b, T c)
y __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD((max), x, y, x[j], y) __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD((min)
std::enable_if_t< detail::is_gengeofloat_v< T >, T > normalize(T p)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > tanh(T x)
__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((min), T x, detail::marray_element_t< T > y, x[i], y) __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD((max)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > floor(T x)
std::enable_if_t< detail::is_svgenfloat_v< T > &&detail::is_genfloatptr_v< T2 >, T > fract(T x, T2 iptr)
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > sin(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
std::enable_if_t< detail::is_svgenfloat_v< T >, T > atanpi(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > asinh(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > remainder(T x, T y)
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > cos(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
std::enable_if_t< detail::is_svgenfloat_v< T >, T > atanh(T x)
std::enable_if_t< detail::is_igeninteger_v< T >, T > rhadd(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > hypot(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T > &&detail::is_genintptr_v< T2 >, T > frexp(T x, T2 exp)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > log1p(T x)
detail::common_rel_ret_t< T > isequal(T x, T y)
std::enable_if_t< detail::is_vigeninteger_v< T >, int > any(T x)
__SYCL_DEPRECATED("This is a deprecated argument type for SYCL nan built-in function.") std
__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(fract, x, iptr, x[j]) __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(modf
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > rsqrt(T x)
y y maxval[j] maxval __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD(clamp, x, minval, maxval, x[j], minval, maxval) __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(mad_hi
T detail::marray_element_t< T > y T T T maxval[i] __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(clamp, T x, detail::marray_element_t< T > minval, detail::marray_element_t< T > maxval, x[i], minval, maxval) __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix
std::enable_if_t< detail::is_svgenfloat_v< T >, T > sinh(T x)
std::enable_if_t< detail::is_svgenfloat_v< T > &&detail::is_genfloatptr_v< T2 >, T > modf(T x, T2 iptr)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > fmod(T x, T y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > copysign(T x, T y)
__SYCL_ALWAYS_INLINE std::enable_if_t< detail::is_sgenfloat_v< T >, marray< int, N > > ilogb(marray< T, N > x)
std::enable_if_t< detail::is_vgengeofloat_v< T >, float > dot(T p0, T p1)
std::enable_if_t< __FAST_MATH_GENFLOAT(T), T > exp2(T x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > atan(T y_over_x)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > atan2(T y, T x)
std::conditional_t< std::is_same_v< ElementType, half >, sycl::detail::half_impl::BIsRepresentationT, ElementType > element_type
Definition: multi_ptr.hpp:752
std::enable_if_t< detail::is_ugeninteger_v< T >, T > abs(T x)
x[j] x[j] __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(lgamma_r, x, signp, x[j]) template< typename T
__SYCL_ALWAYS_INLINE std::enable_if_t< detail::is_sgenfloat_v< T >, marray< T, N > > pown(marray< T, N > x, marray< int, N > y)
std::enable_if_t< detail::is_svgenfloat_v< T >, T > trunc(T x)
y y maxval[j] __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(clamp, x, minval, maxval, x[j], minval[j], maxval[j]) __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD(clamp
std::enable_if_t< detail::is_svgenfloat_v< T >, T > round(T x)
std::enable_if_t< detail::is_geninteger_v< T >, T > ctz(T x)
std::enable_if_t< detail::is_igeninteger_v< T >, T > add_sat(T x, T y)
detail::common_rel_ret_t< T > isnormal(T x)
detail::common_rel_ret_t< T > isgreater(T x, T y)
Definition: access.hpp:18
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept