29 return {x[start], x[start + 1]};
33 for (
size_t i = 0; i < N; i++)
39 for (
size_t i = 0; i < N; i++)
45 #ifdef __SYCL_DEVICE_ONLY__
52 #define __FAST_MATH_GENFLOAT(T) \
53 (detail::is_svgenfloatd<T>::value || detail::is_svgenfloath<T>::value)
54 #define __FAST_MATH_SGENFLOAT(T) \
55 (std::is_same_v<T, double> || std::is_same_v<T, half>)
57 #define __FAST_MATH_GENFLOAT(T) (detail::is_svgenfloat<T>::value)
58 #define __FAST_MATH_SGENFLOAT(T) (detail::is_sgenfloat<T>::value)
68 #define __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
70 for (size_t i = 0; i < N / 2; i++) { \
71 vec<T, 2> partial_res = \
72 __sycl_std::__invoke_##NAME<vec<T, 2>>(detail::to_vec2(x, i * 2)); \
73 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
76 res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1]); \
80 #define __SYCL_MATH_FUNCTION_OVERLOAD(NAME) \
81 template <typename T, size_t N> \
82 inline __SYCL_ALWAYS_INLINE \
83 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
84 NAME(marray<T, N> x) __NOEXC { \
85 __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
118 #undef __SYCL_MATH_FUNCTION_OVERLOAD
122 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
123 template <typename T, size_t N> \
124 inline __SYCL_ALWAYS_INLINE \
125 std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray<T, N>> \
126 NAME(marray<T, N> x) __NOEXC { \
127 __SYCL_MATH_FUNCTION_OVERLOAD_IMPL(NAME) \
142 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
143 #undef __SYCL_MATH_FUNCTION_OVERLOAD_IMPL
145 template <
typename T,
size_t N>
147 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<int, N>>
150 for (
size_t i = 0; i < N / 2; i++) {
156 res[N - 1] = __sycl_std::__invoke_ilogb<int>(x[N - 1]);
161 #define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
163 for (size_t i = 0; i < N / 2; i++) { \
164 auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
165 detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
166 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
169 res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1]); \
173 #define __SYCL_MATH_FUNCTION_2_OVERLOAD(NAME) \
174 template <typename T, size_t N> \
175 inline __SYCL_ALWAYS_INLINE \
176 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
177 NAME(marray<T, N> x, marray<T, N> y) __NOEXC { \
178 __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \
195 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD
197 template <
typename T,
size_t N>
203 #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL
205 #define __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(NAME) \
206 template <typename T, size_t N> \
207 inline __SYCL_ALWAYS_INLINE \
208 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
209 NAME(marray<T, N> x, T y) __NOEXC { \
211 sycl::vec<T, 2> y_vec{y, y}; \
212 for (size_t i = 0; i < N / 2; i++) { \
213 auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
214 detail::to_vec2(x, i * 2), y_vec); \
215 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
218 res[N - 1] = __sycl_std::__invoke_##NAME<T>(x[N - 1], y_vec[0]); \
227 #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD
229 template <
typename T,
size_t N>
231 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
235 for (
size_t i = 0; i < N; i++) {
236 res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k[i]);
241 template <
typename T,
size_t N>
243 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
246 for (
size_t i = 0; i < N; i++) {
247 res[i] = __sycl_std::__invoke_ldexp<T>(x[i], k);
252 #define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME) \
254 for (size_t i = 0; i < N; i++) { \
255 res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y[i]); \
259 template <
typename T,
size_t N>
261 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
266 template <
typename T,
size_t N>
268 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
273 #undef __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL
275 #define __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(NAME) \
277 for (size_t i = 0; i < N; i++) { \
278 res[i] = __sycl_std::__invoke_##NAME<T>(x[i], y); \
282 template <
typename T,
size_t N>
284 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
289 template <
typename T,
size_t N>
291 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>>
295 #undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL
297 #define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME) \
298 template <typename T, size_t N> \
299 inline __SYCL_ALWAYS_INLINE \
300 std::enable_if_t<detail::is_sgenfloat<T>::value, marray<T, N>> \
301 NAME(marray<T, N> x, marray<T, N> y, marray<T, N> z) __NOEXC { \
303 for (size_t i = 0; i < N / 2; i++) { \
304 auto partial_res = __sycl_std::__invoke_##NAME<vec<T, 2>>( \
305 detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2), \
306 detail::to_vec2(z, i * 2)); \
307 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<T, 2>)); \
311 __sycl_std::__invoke_##NAME<T>(x[N - 1], y[N - 1], z[N - 1]); \
319 #undef __SYCL_MATH_FUNCTION_3_OVERLOAD
322 template <
typename T>
323 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
acos(
T x)
__NOEXC {
324 return __sycl_std::__invoke_acos<T>(x);
328 template <
typename T>
329 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> acosh(
T x)
__NOEXC {
330 return __sycl_std::__invoke_acosh<T>(x);
334 template <
typename T>
335 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> acospi(
T x)
__NOEXC {
336 return __sycl_std::__invoke_acospi<T>(x);
340 template <
typename T>
341 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
asin(
T x)
__NOEXC {
342 return __sycl_std::__invoke_asin<T>(x);
346 template <
typename T>
347 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> asinh(
T x)
__NOEXC {
348 return __sycl_std::__invoke_asinh<T>(x);
352 template <
typename T>
353 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> asinpi(
T x)
__NOEXC {
354 return __sycl_std::__invoke_asinpi<T>(x);
358 template <
typename T>
359 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
atan(
T y_over_x)
__NOEXC {
360 return __sycl_std::__invoke_atan<T>(y_over_x);
364 template <
typename T>
365 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
atan2(
T y,
T x)
__NOEXC {
366 return __sycl_std::__invoke_atan2<T>(y, x);
370 template <
typename T>
371 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> atanh(
T x)
__NOEXC {
372 return __sycl_std::__invoke_atanh<T>(x);
376 template <
typename T>
377 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> atanpi(
T x)
__NOEXC {
378 return __sycl_std::__invoke_atanpi<T>(x);
382 template <
typename T>
383 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> atan2pi(
T y,
T x)
__NOEXC {
384 return __sycl_std::__invoke_atan2pi<T>(y, x);
388 template <
typename T>
389 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> cbrt(
T x)
__NOEXC {
390 return __sycl_std::__invoke_cbrt<T>(x);
394 template <
typename T>
395 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
ceil(
T x)
__NOEXC {
396 return __sycl_std::__invoke_ceil<T>(x);
400 template <
typename T>
401 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
copysign(
T x,
403 return __sycl_std::__invoke_copysign<T>(x, y);
407 template <
typename T>
409 return __sycl_std::__invoke_cos<T>(x);
413 template <
typename T>
414 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> cosh(
T x)
__NOEXC {
415 return __sycl_std::__invoke_cosh<T>(x);
419 template <
typename T>
420 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> cospi(
T x)
__NOEXC {
421 return __sycl_std::__invoke_cospi<T>(x);
425 template <
typename T>
426 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> erfc(
T x)
__NOEXC {
427 return __sycl_std::__invoke_erfc<T>(x);
431 template <
typename T>
432 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> erf(
T x)
__NOEXC {
433 return __sycl_std::__invoke_erf<T>(x);
437 template <
typename T>
439 return __sycl_std::__invoke_exp<T>(x);
443 template <
typename T>
445 return __sycl_std::__invoke_exp2<T>(x);
449 template <
typename T>
451 return __sycl_std::__invoke_exp10<T>(x);
455 template <
typename T>
456 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> expm1(
T x)
__NOEXC {
457 return __sycl_std::__invoke_expm1<T>(x);
461 template <
typename T>
462 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
fabs(
T x)
__NOEXC {
463 return __sycl_std::__invoke_fabs<T>(x);
467 template <
typename T>
468 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> fdim(
T x,
T y)
__NOEXC {
469 return __sycl_std::__invoke_fdim<T>(x, y);
473 template <
typename T>
474 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
floor(
T x)
__NOEXC {
475 return __sycl_std::__invoke_floor<T>(x);
479 template <
typename T>
480 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
fma(
T a,
T b,
482 return __sycl_std::__invoke_fma<T>(a, b, c);
486 template <
typename T>
487 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
fmax(
T x,
T y)
__NOEXC {
488 return __sycl_std::__invoke_fmax<T>(x, y);
492 template <
typename T>
493 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
495 return __sycl_std::__invoke_fmax<T>(x,
T(y));
499 template <
typename T>
500 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
fmin(
T x,
T y)
__NOEXC {
501 return __sycl_std::__invoke_fmin<T>(x, y);
505 template <
typename T>
506 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
508 return __sycl_std::__invoke_fmin<T>(x,
T(y));
512 template <
typename T>
513 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
fmod(
T x,
T y)
__NOEXC {
514 return __sycl_std::__invoke_fmod<T>(x, y);
518 template <
typename T,
typename T2>
520 detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value,
T>
522 detail::check_vector_size<T, T2>();
523 return __sycl_std::__invoke_fract<T>(x, iptr);
527 template <
typename T,
typename T2>
529 detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value,
T>
531 detail::check_vector_size<T, T2>();
532 return __sycl_std::__invoke_frexp<T>(x,
exp);
536 template <
typename T>
537 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> hypot(
T x,
T y)
__NOEXC {
538 return __sycl_std::__invoke_hypot<T>(x, y);
542 template <
typename T,
543 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
545 return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(
x);
551 template <
typename T>
552 std::enable_if_t<detail::is_sgenfloat<T>::value,
T>
ldexp(
T x,
int k)
__NOEXC {
553 return __sycl_std::__invoke_ldexp<T>(x, k);
557 template <
typename T>
558 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
ldexp(
T x,
int k)
__NOEXC {
559 return __sycl_std::__invoke_ldexp<T>(x, vec<
int, T::size()>(k));
563 template <
typename T,
typename T2>
564 std::enable_if_t<detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value,
567 detail::check_vector_size<T, T2>();
568 return __sycl_std::__invoke_ldexp<T>(x, k);
572 template <
typename T>
573 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> lgamma(
T x)
__NOEXC {
574 return __sycl_std::__invoke_lgamma<T>(x);
578 template <
typename T,
typename T2>
580 detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value,
T>
582 detail::check_vector_size<T, T2>();
583 return __sycl_std::__invoke_lgamma_r<T>(x, signp);
587 template <
typename T>
589 return __sycl_std::__invoke_log<T>(x);
593 template <
typename T>
595 return __sycl_std::__invoke_log2<T>(x);
599 template <
typename T>
601 return __sycl_std::__invoke_log10<T>(x);
605 template <
typename T>
606 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> log1p(
T x)
__NOEXC {
607 return __sycl_std::__invoke_log1p<T>(x);
611 template <
typename T>
612 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> logb(
T x)
__NOEXC {
613 return __sycl_std::__invoke_logb<T>(x);
617 template <
typename T>
618 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> mad(
T a,
T b,
620 return __sycl_std::__invoke_mad<T>(a, b, c);
624 template <
typename T>
625 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> maxmag(
T x,
T y)
__NOEXC {
626 return __sycl_std::__invoke_maxmag<T>(x, y);
630 template <
typename T>
631 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> minmag(
T x,
T y)
__NOEXC {
632 return __sycl_std::__invoke_minmag<T>(x, y);
636 template <
typename T,
typename T2>
638 detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value,
T>
640 detail::check_vector_size<T, T2>();
641 return __sycl_std::__invoke_modf<T>(x, iptr);
644 template <
typename T,
645 typename = std::enable_if_t<detail::is_nan_type<T>::value,
T>>
646 detail::nan_return_t<T> nan(
T nancode)
__NOEXC {
647 return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
648 detail::convert_data_type<T, detail::nan_argument_base_t<T>>()(nancode));
652 template <
typename T>
653 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> nextafter(
T x,
655 return __sycl_std::__invoke_nextafter<T>(x, y);
659 template <
typename T>
660 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
pow(
T x,
T y)
__NOEXC {
661 return __sycl_std::__invoke_pow<T>(x, y);
665 template <
typename T,
typename T2>
667 detail::is_svgenfloat<T>::value && detail::is_genint<T2>::value,
T>
669 detail::check_vector_size<T, T2>();
670 return __sycl_std::__invoke_pown<T>(x, y);
674 template <
typename T>
676 return __sycl_std::__invoke_powr<T>(x, y);
680 template <
typename T>
681 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> remainder(
T x,
683 return __sycl_std::__invoke_remainder<T>(x, y);
687 template <
typename T,
typename T2>
689 detail::is_svgenfloat<T>::value && detail::is_genintptr<T2>::value,
T>
691 detail::check_vector_size<T, T2>();
692 return __sycl_std::__invoke_remquo<T>(x, y, quo);
696 template <
typename T>
697 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
rint(
T x)
__NOEXC {
698 return __sycl_std::__invoke_rint<T>(x);
702 template <
typename T,
typename T2>
704 detail::is_svgenfloat<T>::value && detail::is_genint<T2>::value,
T>
706 detail::check_vector_size<T, T2>();
707 return __sycl_std::__invoke_rootn<T>(x, y);
711 template <
typename T>
712 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> round(
T x)
__NOEXC {
713 return __sycl_std::__invoke_round<T>(x);
717 template <
typename T>
719 return __sycl_std::__invoke_rsqrt<T>(x);
723 template <
typename T>
725 return __sycl_std::__invoke_sin<T>(x);
729 template <
typename T,
typename T2>
731 detail::is_svgenfloat<T>::value && detail::is_genfloatptr<T2>::value,
T>
733 detail::check_vector_size<T, T2>();
734 return __sycl_std::__invoke_sincos<T>(x, cosval);
738 template <
typename T>
739 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> sinh(
T x)
__NOEXC {
740 return __sycl_std::__invoke_sinh<T>(x);
744 template <
typename T>
745 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> sinpi(
T x)
__NOEXC {
746 return __sycl_std::__invoke_sinpi<T>(x);
750 template <
typename T>
752 return __sycl_std::__invoke_sqrt<T>(x);
756 template <
typename T>
758 return __sycl_std::__invoke_tan<T>(x);
762 template <
typename T>
763 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
tanh(
T x)
__NOEXC {
764 return __sycl_std::__invoke_tanh<T>(x);
768 template <
typename T>
769 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> tanpi(
T x)
__NOEXC {
770 return __sycl_std::__invoke_tanpi<T>(x);
774 template <
typename T>
775 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> tgamma(
T x)
__NOEXC {
776 return __sycl_std::__invoke_tgamma<T>(x);
780 template <
typename T>
781 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
trunc(
T x)
__NOEXC {
782 return __sycl_std::__invoke_trunc<T>(x);
789 #define __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARGPTR, \
792 for (int j = 0; j < N; j++) { \
795 address_space_cast<AddressSpace, IsDecorated, \
796 detail::marray_element_t<T2>>(&(*ARGPTR)[j])); \
800 #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \
801 NAME, ARG1, ARG2, ...) \
802 template <typename T, size_t N, typename T2, \
803 access::address_space AddressSpace, access::decorated IsDecorated> \
805 detail::is_svgenfloat<T>::value && \
806 detail::is_genfloatptr_marray<T2, AddressSpace, IsDecorated>::value, \
808 NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) \
810 __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
821 #undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENFLOATPTR_OVERLOAD
823 #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \
824 NAME, ARG1, ARG2, ...) \
825 template <typename T, size_t N, typename T2, \
826 access::address_space AddressSpace, access::decorated IsDecorated> \
828 detail::is_svgenfloat<T>::value && \
829 detail::is_genintptr_marray<T2, AddressSpace, IsDecorated>::value, \
831 NAME(marray<T, N> ARG1, multi_ptr<T2, AddressSpace, IsDecorated> ARG2) \
833 __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \
842 #undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD
844 #define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \
845 template <typename T, size_t N, typename T2, \
846 access::address_space AddressSpace, access::decorated IsDecorated> \
848 detail::is_svgenfloat<T>::value && \
849 detail::is_genintptr_marray<T2, AddressSpace, IsDecorated>::value, \
851 NAME(marray<T, N> x, marray<T, N> y, \
852 multi_ptr<T2, AddressSpace, IsDecorated> quo) __NOEXC { \
853 __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \
859 #undef __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD
861 #undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL
863 template <
typename T,
size_t N>
864 std::enable_if_t<detail::is_nan_type<T>::value,
865 marray<detail::nan_return_t<T>, N>>
866 nan(marray<T, N> nancode)
__NOEXC {
867 marray<detail::nan_return_t<T>, N> res;
868 for (
int j = 0; j < N; j++) {
869 res[j] = nan(nancode[j]);
876 template <
typename T>
877 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
clamp(
T x,
T minval,
879 return __sycl_std::__invoke_fclamp<T>(x, minval, maxval);
885 template <
typename T>
886 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
887 clamp(
T x,
typename T::element_type minval,
888 typename T::element_type maxval)
__NOEXC {
889 return __sycl_std::__invoke_fclamp<T>(x,
T(minval),
T(maxval));
893 template <
typename T>
894 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
896 return __sycl_std::__invoke_degrees<T>(radians);
900 template <
typename T>
901 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
abs(
T x)
__NOEXC {
902 return __sycl_std::__invoke_fabs<T>(x);
906 template <
typename T>
907 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>(
max)(
T x,
T y)
__NOEXC {
908 return __sycl_std::__invoke_fmax_common<T>(x, y);
914 template <
typename T>
915 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>(
max)(
916 T x,
typename T::element_type y)
__NOEXC {
917 return __sycl_std::__invoke_fmax_common<T>(x,
T(y));
921 template <
typename T>
922 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>(
min)(
T x,
T y)
__NOEXC {
923 return __sycl_std::__invoke_fmin_common<T>(x, y);
929 template <
typename T>
930 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>(
min)(
931 T x,
typename T::element_type y)
__NOEXC {
932 return __sycl_std::__invoke_fmin_common<T>(x,
T(y));
936 template <
typename T>
937 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> mix(
T x,
T y,
939 return __sycl_std::__invoke_mix<T>(x, y, a);
945 template <
typename T>
946 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
947 mix(
T x,
T y,
typename T::element_type a)
__NOEXC {
948 return __sycl_std::__invoke_mix<T>(x, y,
T(a));
952 template <
typename T>
953 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
955 return __sycl_std::__invoke_radians<T>(degrees);
959 template <
typename T>
960 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> step(
T edge,
T x)
__NOEXC {
961 return __sycl_std::__invoke_step<T>(edge, x);
967 template <
typename T>
968 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
969 step(
typename T::element_type edge,
T x)
__NOEXC {
970 return __sycl_std::__invoke_step<T>(
T(edge), x);
974 template <
typename T>
975 std::enable_if_t<detail::is_svgenfloat<T>::value,
T>
977 return __sycl_std::__invoke_smoothstep<T>(edge0, edge1, x);
983 template <
typename T>
984 std::enable_if_t<detail::is_vgenfloat<T>::value,
T>
985 smoothstep(
typename T::element_type edge0,
typename T::element_type edge1,
987 return __sycl_std::__invoke_smoothstep<T>(
T(edge0),
T(edge1), x);
991 template <
typename T>
992 std::enable_if_t<detail::is_svgenfloat<T>::value,
T> sign(
T x)
__NOEXC {
993 return __sycl_std::__invoke_sign<T>(x);
1000 #define __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
1002 for (int i = 0; i < T::size(); i++) { \
1003 res[i] = NAME(__VA_ARGS__); \
1007 #define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \
1008 template <typename T, \
1009 typename = std::enable_if_t<detail::is_mgenfloat<T>::value>> \
1010 T NAME(ARG) __NOEXC { \
1011 __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1018 #undef __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD
1020 #define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ...) \
1021 template <typename T, \
1022 typename = std::enable_if_t<detail::is_mgenfloat<T>::value>> \
1023 T NAME(ARG1, ARG2) __NOEXC { \
1024 __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1031 detail::marray_element_t<T> y,
1035 detail::marray_element_t<T> y,
1039 detail::marray_element_t<T> edge,
1042 #undef __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD
1044 #define __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(NAME, ARG1, ARG2, ARG3, \
1046 template <typename T, \
1047 typename = std::enable_if_t<detail::is_mgenfloat<T>::value>> \
1048 T NAME(ARG1, ARG2, ARG3) __NOEXC { \
1049 __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1053 x[i], minval[i], maxval[i])
1055 clamp,
T x, detail::marray_element_t<T> minval,
1056 detail::marray_element_t<T> maxval, x[i], minval, maxval)
1060 detail::marray_element_t<T> a,
1063 edge0[i], edge1[i], x[i])
1065 smoothstep, detail::marray_element_t<T> edge0,
1066 detail::marray_element_t<T> edge1,
T x, edge0, edge1, x[i])
1068 #undef __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD
1069 #undef __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL
1073 template <
typename T>
1074 std::enable_if_t<detail::is_ugeninteger<T>::value,
T>
abs(
T x)
__NOEXC {
1075 return __sycl_std::__invoke_u_abs<T>(x);
1079 template <
typename T>
1080 std::enable_if_t<detail::is_igeninteger<T>::value,
T>
abs(
T x)
__NOEXC {
1081 auto res = __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(
x);
1082 if constexpr (detail::is_vigeninteger<T>::value) {
1083 return res.template convert<detail::vector_element_t<T>>();
1085 return detail::make_signed_t<decltype(res)>(res);
1089 template <
typename T>
1090 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> abs_diff(
T x,
1092 return __sycl_std::__invoke_u_abs_diff<T>(x, y);
1096 template <
typename T>
1097 std::enable_if_t<detail::is_igeninteger<T>::value, detail::make_unsigned_t<T>>
1099 return __sycl_std::__invoke_s_abs_diff<detail::make_unsigned_t<T>>(
x,
y);
1103 template <
typename T>
1104 std::enable_if_t<detail::is_igeninteger<T>::value,
T> add_sat(
T x,
1106 return __sycl_std::__invoke_s_add_sat<T>(x, y);
1110 template <
typename T>
1111 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> add_sat(
T x,
1113 return __sycl_std::__invoke_u_add_sat<T>(x, y);
1117 template <
typename T>
1118 std::enable_if_t<detail::is_igeninteger<T>::value,
T>
hadd(
T x,
T y)
__NOEXC {
1119 return __sycl_std::__invoke_s_hadd<T>(x, y);
1123 template <
typename T>
1124 std::enable_if_t<detail::is_ugeninteger<T>::value,
T>
hadd(
T x,
T y)
__NOEXC {
1125 return __sycl_std::__invoke_u_hadd<T>(x, y);
1129 template <
typename T>
1130 std::enable_if_t<detail::is_igeninteger<T>::value,
T> rhadd(
T x,
T y)
__NOEXC {
1131 return __sycl_std::__invoke_s_rhadd<T>(x, y);
1135 template <
typename T>
1136 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> rhadd(
T x,
T y)
__NOEXC {
1137 return __sycl_std::__invoke_u_rhadd<T>(x, y);
1141 template <
typename T>
1142 std::enable_if_t<detail::is_igeninteger<T>::value,
T>
clamp(
T x,
T minval,
1144 return __sycl_std::__invoke_s_clamp<T>(x, minval, maxval);
1148 template <
typename T>
1149 std::enable_if_t<detail::is_ugeninteger<T>::value,
T>
clamp(
T x,
T minval,
1151 return __sycl_std::__invoke_u_clamp<T>(x, minval, maxval);
1155 template <
typename T>
1156 std::enable_if_t<detail::is_vigeninteger<T>::value,
T>
1157 clamp(
T x,
typename T::element_type minval,
1158 typename T::element_type maxval)
__NOEXC {
1159 return __sycl_std::__invoke_s_clamp<T>(x,
T(minval),
T(maxval));
1163 template <
typename T>
1164 std::enable_if_t<detail::is_vugeninteger<T>::value,
T>
1165 clamp(
T x,
typename T::element_type minval,
1166 typename T::element_type maxval)
__NOEXC {
1167 return __sycl_std::__invoke_u_clamp<T>(x,
T(minval),
T(maxval));
1171 template <
typename T>
1172 std::enable_if_t<detail::is_geninteger<T>::value,
T> clz(
T x)
__NOEXC {
1173 return __sycl_std::__invoke_clz<T>(x);
1177 template <
typename T>
1178 std::enable_if_t<detail::is_geninteger<T>::value,
T> ctz(
T x)
__NOEXC {
1179 return __sycl_std::__invoke_ctz<T>(x);
1183 namespace ext::intel {
1184 template <
typename T>
1186 "'sycl::ext::intel::ctz' is deprecated, use 'sycl::ctz' instead")
1188 return sycl::ctz(x);
1193 using namespace ext::intel;
1197 template <
typename T>
1198 std::enable_if_t<detail::is_igeninteger<T>::value,
T> mad_hi(
T x,
T y,
1200 return __sycl_std::__invoke_s_mad_hi<T>(x, y, z);
1204 template <
typename T>
1205 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> mad_hi(
T x,
T y,
1207 return __sycl_std::__invoke_u_mad_hi<T>(x, y, z);
1211 template <
typename T>
1212 std::enable_if_t<detail::is_igeninteger<T>::value,
T> mad_sat(
T a,
T b,
1214 return __sycl_std::__invoke_s_mad_sat<T>(a, b, c);
1218 template <
typename T>
1219 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> mad_sat(
T a,
T b,
1221 return __sycl_std::__invoke_u_mad_sat<T>(a, b, c);
1225 template <
typename T>
1226 std::enable_if_t<detail::is_igeninteger<T>::value,
T>(
max)(
T x,
T y)
__NOEXC {
1227 return __sycl_std::__invoke_s_max<T>(x, y);
1231 template <
typename T>
1232 std::enable_if_t<detail::is_ugeninteger<T>::value,
T>(
max)(
T x,
T y)
__NOEXC {
1233 return __sycl_std::__invoke_u_max<T>(x, y);
1237 template <
typename T>
1238 std::enable_if_t<detail::is_vigeninteger<T>::value,
T>(
max)(
1239 T x,
typename T::element_type y)
__NOEXC {
1240 return __sycl_std::__invoke_s_max<T>(x,
T(y));
1244 template <
typename T>
1245 std::enable_if_t<detail::is_vugeninteger<T>::value,
T>(
max)(
1246 T x,
typename T::element_type y)
__NOEXC {
1247 return __sycl_std::__invoke_u_max<T>(x,
T(y));
1251 template <
typename T>
1252 std::enable_if_t<detail::is_igeninteger<T>::value,
T>(
min)(
T x,
T y)
__NOEXC {
1253 return __sycl_std::__invoke_s_min<T>(x, y);
1257 template <
typename T>
1258 std::enable_if_t<detail::is_ugeninteger<T>::value,
T>(
min)(
T x,
T y)
__NOEXC {
1259 return __sycl_std::__invoke_u_min<T>(x, y);
1263 template <
typename T>
1264 std::enable_if_t<detail::is_vigeninteger<T>::value,
T>(
min)(
1265 T x,
typename T::element_type y)
__NOEXC {
1266 return __sycl_std::__invoke_s_min<T>(x,
T(y));
1270 template <
typename T>
1271 std::enable_if_t<detail::is_vugeninteger<T>::value,
T>(
min)(
1272 T x,
typename T::element_type y)
__NOEXC {
1273 return __sycl_std::__invoke_u_min<T>(x,
T(y));
1277 template <
typename T>
1278 std::enable_if_t<detail::is_igeninteger<T>::value,
T> mul_hi(
T x,
T y)
__NOEXC {
1279 return __sycl_std::__invoke_s_mul_hi<T>(x, y);
1283 template <
typename T>
1284 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> mul_hi(
T x,
T y)
__NOEXC {
1285 return __sycl_std::__invoke_u_mul_hi<T>(x, y);
1289 template <
typename T>
1290 std::enable_if_t<detail::is_geninteger<T>::value,
T> rotate(
T v,
T i)
__NOEXC {
1291 return __sycl_std::__invoke_rotate<T>(v, i);
1295 template <
typename T>
1296 std::enable_if_t<detail::is_igeninteger<T>::value,
T> sub_sat(
T x,
1298 return __sycl_std::__invoke_s_sub_sat<T>(x, y);
1302 template <
typename T>
1303 std::enable_if_t<detail::is_ugeninteger<T>::value,
T> sub_sat(
T x,
1305 return __sycl_std::__invoke_u_sub_sat<T>(x, y);
1309 template <
typename T>
1310 std::enable_if_t<detail::is_ugeninteger8bit<T>::value, detail::make_larger_t<T>>
1312 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1316 template <
typename T,
typename T2>
1317 std::enable_if_t<detail::is_igeninteger8bit<T>::value &&
1318 detail::is_ugeninteger8bit<T2>::value,
1319 detail::make_larger_t<T>>
1321 detail::check_vector_size<T, T2>();
1322 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1326 template <
typename T>
1327 std::enable_if_t<detail::is_ugeninteger16bit<T>::value,
1328 detail::make_larger_t<T>>
1330 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1334 template <
typename T,
typename T2>
1335 std::enable_if_t<detail::is_igeninteger16bit<T>::value &&
1336 detail::is_ugeninteger16bit<T2>::value,
1337 detail::make_larger_t<T>>
1339 detail::check_vector_size<T, T2>();
1340 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1344 template <
typename T>
1345 std::enable_if_t<detail::is_ugeninteger32bit<T>::value,
1346 detail::make_larger_t<T>>
1348 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
1352 template <
typename T,
typename T2>
1353 std::enable_if_t<detail::is_igeninteger32bit<T>::value &&
1354 detail::is_ugeninteger32bit<T2>::value,
1355 detail::make_larger_t<T>>
1357 detail::check_vector_size<T, T2>();
1358 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
1362 template <
typename T>
1364 return __sycl_std::__invoke_popcount<T>(x);
1369 template <
typename T>
1370 std::enable_if_t<detail::is_igeninteger32bit<T>::value,
T> mad24(
T x,
T y,
1372 return __sycl_std::__invoke_s_mad24<T>(x, y, z);
1377 template <
typename T>
1378 std::enable_if_t<detail::is_ugeninteger32bit<T>::value,
T> mad24(
T x,
T y,
1380 return __sycl_std::__invoke_u_mad24<T>(x, y, z);
1384 template <
typename T>
1385 std::enable_if_t<detail::is_igeninteger32bit<T>::value,
T> mul24(
T x,
1387 return __sycl_std::__invoke_s_mul24<T>(x, y);
1391 template <
typename T>
1392 std::enable_if_t<detail::is_ugeninteger32bit<T>::value,
T> mul24(
T x,
1394 return __sycl_std::__invoke_u_mul24<T>(x, y);
1401 #define __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
1403 for (int j = 0; j < N; j++) { \
1404 res[j] = NAME(__VA_ARGS__); \
1409 #define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD(NAME, ARG, ...) \
1410 template <typename T, size_t N> \
1411 std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1412 marray<T, N> ARG) __NOEXC { \
1413 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1416 #define __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD(NAME, ARG, ...) \
1417 template <typename T, size_t N> \
1418 std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1419 marray<T, N> ARG) __NOEXC { \
1420 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1426 #undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_I_OVERLOAD
1427 #undef __SYCL_MARRAY_INTEGER_FUNCTION_ABS_U_OVERLOAD
1429 #define __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \
1430 template <typename T, size_t N> \
1431 std::enable_if_t<detail::is_geninteger<T>::value, marray<T, N>> NAME( \
1432 marray<T, N> ARG) __NOEXC { \
1433 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1440 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UNOP_OVERLOAD
1442 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
1443 template <typename T, size_t N> \
1444 std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1445 marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1446 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1449 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD(NAME, ARG1, \
1451 template <typename T, size_t N> \
1452 std::enable_if_t<detail::is_igeninteger<T>::value, \
1453 marray<detail::make_unsigned_t<T>, N>> \
1454 NAME(marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1455 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1458 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
1459 template <typename T, size_t N> \
1460 std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1461 marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1462 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1465 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD( \
1466 NAME, ARG1, ARG2, ...) \
1467 template <typename T, size_t N> \
1468 std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1469 marray<T, N> ARG1, T ARG2) __NOEXC { \
1470 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1473 #define __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD( \
1474 NAME, ARG1, ARG2, ...) \
1475 template <typename T, size_t N> \
1476 std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1477 marray<T, N> ARG1, T ARG2) __NOEXC { \
1478 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1509 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_2ND_ARG_SCALAR_OVERLOAD
1510 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_2ND_ARG_SCALAR_OVERLOAD
1511 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_OVERLOAD
1512 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_I_RET_U_OVERLOAD
1513 #undef __SYCL_MARRAY_INTEGER_FUNCTION_BINOP_U_OVERLOAD
1515 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD(NAME, ARG1, ARG2, \
1517 template <typename T, size_t N> \
1518 std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1519 marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1520 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1523 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD(NAME, ARG1, ARG2, \
1525 template <typename T, size_t N> \
1526 std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1527 marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1528 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1531 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
1532 NAME, ARG1, ARG2, ARG3, ...) \
1533 template <typename T, size_t N> \
1534 std::enable_if_t<detail::is_ugeninteger<T>::value, marray<T, N>> NAME( \
1535 marray<T, N> ARG1, T ARG2, T ARG3) __NOEXC { \
1536 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1539 #define __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD( \
1540 NAME, ARG1, ARG2, ARG3, ...) \
1541 template <typename T, size_t N> \
1542 std::enable_if_t<detail::is_igeninteger<T>::value, marray<T, N>> NAME( \
1543 marray<T, N> ARG1, T ARG2, T ARG3) __NOEXC { \
1544 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1548 minval[j], maxval[j])
1550 minval[j], maxval[j])
1552 clamp, x, minval, maxval, x[j], minval, maxval)
1554 clamp, x, minval, maxval, x[j], minval, maxval)
1564 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_2ND_3RD_ARGS_SCALAR_OVERLOAD
1565 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_2ND_3RD_ARGS_SCALAR_OVERLOAD
1566 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_I_OVERLOAD
1567 #undef __SYCL_MARRAY_INTEGER_FUNCTION_TEROP_U_OVERLOAD
1570 #define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD(NAME, ARG1, ARG2, \
1572 template <typename T, size_t N> \
1573 std::enable_if_t<detail::is_ugeninteger32bit<T>::value, marray<T, N>> NAME( \
1574 marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1575 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1578 #define __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD(NAME, ARG1, ARG2, \
1580 template <typename T, size_t N> \
1581 std::enable_if_t<detail::is_igeninteger32bit<T>::value, marray<T, N>> NAME( \
1582 marray<T, N> ARG1, marray<T, N> ARG2, marray<T, N> ARG3) __NOEXC { \
1583 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1591 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_I_OVERLOAD
1592 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MAD24_U_OVERLOAD
1595 #define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD(NAME, ARG1, ARG2, ...) \
1596 template <typename T, size_t N> \
1597 std::enable_if_t<detail::is_ugeninteger32bit<T>::value, marray<T, N>> NAME( \
1598 marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1599 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1602 #define __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD(NAME, ARG1, ARG2, ...) \
1603 template <typename T, size_t N> \
1604 std::enable_if_t<detail::is_igeninteger32bit<T>::value, marray<T, N>> NAME( \
1605 marray<T, N> ARG1, marray<T, N> ARG2) __NOEXC { \
1606 __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \
1612 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_I_OVERLOAD
1613 #undef __SYCL_MARRAY_INTEGER_FUNCTION_MUL24_U_OVERLOAD
1614 #undef __SYCL_MARRAY_INTEGER_FUNCTION_OVERLOAD_IMPL
1618 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1619 detail::make_larger_t<marray<T, N>> res; \
1620 for (int j = 0; j < N; j++) { \
1621 res[j] = NAME(hi[j], lo[j]); \
1626 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD(NAME, KBIT) \
1627 template <typename T, size_t N> \
1628 std::enable_if_t<detail::is_ugeninteger##KBIT<T>::value, \
1629 detail::make_larger_t<marray<T, N>>> \
1630 NAME(marray<T, N> hi, marray<T, N> lo) __NOEXC { \
1631 __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1634 #define __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD(NAME, KBIT) \
1635 template <typename T, typename T2, size_t N> \
1636 std::enable_if_t<detail::is_igeninteger##KBIT<T>::value && \
1637 detail::is_ugeninteger##KBIT<T2>::value, \
1638 detail::make_larger_t<marray<T, N>>> \
1639 NAME(marray<T, N> hi, marray<T2, N> lo) __NOEXC { \
1640 __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL(NAME) \
1650 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_IU_OVERLOAD
1651 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_UU_OVERLOAD
1652 #undef __SYCL_MARRAY_INTEGER_FUNCTION_UPSAMPLE_OVERLOAD_IMPL
1661 template <
typename T>
1662 std::enable_if_t<detail::is_gencross<T>::value,
T> cross(
T p0,
T p1)
__NOEXC {
1663 return __sycl_std::__invoke_cross<T>(p0, p1);
1669 template <
typename T>
1670 std::enable_if_t<detail::is_sgenfloat<T>::value,
T> dot(
T p0,
T p1)
__NOEXC {
1675 template <
typename T>
1676 std::enable_if_t<detail::is_vgengeofloat<T>::value,
float> dot(
T p0,
1678 return __sycl_std::__invoke_Dot<float>(p0, p1);
1682 template <
typename T>
1683 std::enable_if_t<detail::is_vgengeodouble<T>::value,
double> dot(
T p0,
1685 return __sycl_std::__invoke_Dot<double>(p0, p1);
1689 template <
typename T>
1690 std::enable_if_t<detail::is_vgengeohalf<T>::value,
half> dot(
T p0,
1692 return __sycl_std::__invoke_Dot<half>(p0, p1);
1696 template <
typename T,
1697 typename = std::enable_if_t<detail::is_gengeofloat<T>::value,
T>>
1699 return __sycl_std::__invoke_distance<float>(p0, p1);
1703 template <
typename T,
1704 typename = std::enable_if_t<detail::is_gengeodouble<T>::value,
T>>
1706 return __sycl_std::__invoke_distance<double>(p0, p1);
1710 template <
typename T,
1711 typename = std::enable_if_t<detail::is_gengeohalf<T>::value,
T>>
1713 return __sycl_std::__invoke_distance<half>(p0, p1);
1717 template <
typename T,
1718 typename = std::enable_if_t<detail::is_gengeofloat<T>::value,
T>>
1720 return __sycl_std::__invoke_length<float>(p);
1724 template <
typename T,
1725 typename = std::enable_if_t<detail::is_gengeodouble<T>::value,
T>>
1727 return __sycl_std::__invoke_length<double>(p);
1731 template <
typename T,
1732 typename = std::enable_if_t<detail::is_gengeohalf<T>::value,
T>>
1734 return __sycl_std::__invoke_length<half>(p);
1738 template <
typename T>
1739 std::enable_if_t<detail::is_gengeofloat<T>::value,
T> normalize(
T p)
__NOEXC {
1740 return __sycl_std::__invoke_normalize<T>(p);
1744 template <
typename T>
1745 std::enable_if_t<detail::is_gengeodouble<T>::value,
T> normalize(
T p)
__NOEXC {
1746 return __sycl_std::__invoke_normalize<T>(p);
1750 template <
typename T>
1751 std::enable_if_t<detail::is_gengeohalf<T>::value,
T> normalize(
T p)
__NOEXC {
1752 return __sycl_std::__invoke_normalize<T>(p);
1756 template <
typename T,
1757 typename = std::enable_if_t<detail::is_gengeofloat<T>::value,
T>>
1758 float fast_distance(
T p0,
T p1)
__NOEXC {
1759 return __sycl_std::__invoke_fast_distance<float>(p0, p1);
1763 template <
typename T,
1764 typename = std::enable_if_t<detail::is_gengeodouble<T>::value,
T>>
1765 double fast_distance(
T p0,
T p1)
__NOEXC {
1766 return __sycl_std::__invoke_fast_distance<double>(p0, p1);
1770 template <
typename T,
1771 typename = std::enable_if_t<detail::is_gengeofloat<T>::value,
T>>
1773 return __sycl_std::__invoke_fast_length<float>(p);
1777 template <
typename T,
1778 typename = std::enable_if_t<detail::is_gengeodouble<T>::value,
T>>
1780 return __sycl_std::__invoke_fast_length<double>(p);
1784 template <
typename T>
1785 std::enable_if_t<detail::is_gengeofloat<T>::value,
T>
1787 return __sycl_std::__invoke_fast_normalize<T>(p);
1791 template <
typename T>
1792 std::enable_if_t<detail::is_gengeodouble<T>::value,
T>
1794 return __sycl_std::__invoke_fast_normalize<T>(p);
1799 #define __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL(NAME, ...) \
1800 vec<detail::marray_element_t<T>, T::size()> result_v; \
1801 result_v = NAME(__VA_ARGS__); \
1802 return detail::to_marray(result_v);
1804 template <
typename T>
1805 std::enable_if_t<detail::is_gencrossmarray<T>::value,
T> cross(
T p0,
1811 template <
typename T>
1812 std::enable_if_t<detail::is_gengeomarray<T>::value,
T> normalize(
T p)
__NOEXC {
1816 template <
typename T>
1817 std::enable_if_t<detail::is_gengeomarrayfloat<T>::value,
T>
1823 #undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_OVERLOAD_IMPL
1825 #define __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD(NAME) \
1826 template <typename T> \
1827 std::enable_if_t<detail::is_gengeomarray<T>::value, \
1828 detail::marray_element_t<T>> \
1829 NAME(T p0, T p1) __NOEXC { \
1830 return NAME(detail::to_vec(p0), detail::to_vec(p1)); \
1838 #undef __SYCL_MARRAY_GEOMETRIC_FUNCTION_IS_GENGEOMARRAY_BINOP_OVERLOAD
1840 template <
typename T>
1841 std::enable_if_t<detail::is_gengeomarray<T>::value, detail::marray_element_t<T>>
1843 return __sycl_std::__invoke_length<detail::marray_element_t<T>>(
1847 template <
typename T>
1848 std::enable_if_t<detail::is_gengeomarrayfloat<T>::value,
1849 detail::marray_element_t<T>>
1854 template <
typename T>
1855 std::enable_if_t<detail::is_gengeomarrayfloat<T>::value,
1856 detail::marray_element_t<T>>
1864 template <
typename T,
1865 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1866 detail::common_rel_ret_t<T> isequal(
T x,
T y)
__NOEXC {
1867 return detail::RelConverter<T>::apply(
1868 __sycl_std::__invoke_FOrdEqual<detail::internal_rel_ret_t<T>>(x, y));
1871 template <
typename T,
1872 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1873 detail::common_rel_ret_t<T> isnotequal(
T x,
T y)
__NOEXC {
1874 return detail::RelConverter<T>::apply(
1875 __sycl_std::__invoke_FUnordNotEqual<detail::internal_rel_ret_t<T>>(x, y));
1878 template <
typename T,
1879 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1880 detail::common_rel_ret_t<T> isgreater(
T x,
T y)
__NOEXC {
1881 return detail::RelConverter<T>::apply(
1882 __sycl_std::__invoke_FOrdGreaterThan<detail::internal_rel_ret_t<T>>(x,
1886 template <
typename T,
1887 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1888 detail::common_rel_ret_t<T> isgreaterequal(
T x,
T y)
__NOEXC {
1889 return detail::RelConverter<T>::apply(
1890 __sycl_std::__invoke_FOrdGreaterThanEqual<detail::internal_rel_ret_t<T>>(
1894 template <
typename T,
1895 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1896 detail::common_rel_ret_t<T> isless(
T x,
T y)
__NOEXC {
1897 return detail::RelConverter<T>::apply(
1898 __sycl_std::__invoke_FOrdLessThan<detail::internal_rel_ret_t<T>>(x, y));
1901 template <
typename T,
1902 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1903 detail::common_rel_ret_t<T> islessequal(
T x,
T y)
__NOEXC {
1904 return detail::RelConverter<T>::apply(
1905 __sycl_std::__invoke_FOrdLessThanEqual<detail::internal_rel_ret_t<T>>(x,
1909 template <
typename T,
1910 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1911 detail::common_rel_ret_t<T> islessgreater(
T x,
T y)
__NOEXC {
1912 return detail::RelConverter<T>::apply(
1913 __sycl_std::__invoke_FOrdNotEqual<detail::internal_rel_ret_t<T>>(x, y));
1916 template <
typename T,
1917 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1918 detail::common_rel_ret_t<T> isfinite(
T x)
__NOEXC {
1919 return detail::RelConverter<T>::apply(
1920 __sycl_std::__invoke_IsFinite<detail::internal_rel_ret_t<T>>(x));
1923 template <
typename T,
1924 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1925 detail::common_rel_ret_t<T> isinf(
T x)
__NOEXC {
1926 return detail::RelConverter<T>::apply(
1927 __sycl_std::__invoke_IsInf<detail::internal_rel_ret_t<T>>(x));
1930 template <
typename T,
1931 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1933 return detail::RelConverter<T>::apply(
1934 __sycl_std::__invoke_IsNan<detail::internal_rel_ret_t<T>>(x));
1937 template <
typename T,
1938 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1939 detail::common_rel_ret_t<T> isnormal(
T x)
__NOEXC {
1940 return detail::RelConverter<T>::apply(
1941 __sycl_std::__invoke_IsNormal<detail::internal_rel_ret_t<T>>(x));
1944 template <
typename T,
1945 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1946 detail::common_rel_ret_t<T> isordered(
T x,
T y)
__NOEXC {
1947 return detail::RelConverter<T>::apply(
1948 __sycl_std::__invoke_Ordered<detail::internal_rel_ret_t<T>>(x, y));
1951 template <
typename T,
1952 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1953 detail::common_rel_ret_t<T> isunordered(
T x,
T y)
__NOEXC {
1954 return detail::RelConverter<T>::apply(
1955 __sycl_std::__invoke_Unordered<detail::internal_rel_ret_t<T>>(x, y));
1958 template <
typename T,
1959 typename = std::enable_if_t<detail::is_svgenfloat<T>::value,
T>>
1960 detail::common_rel_ret_t<T> signbit(
T x)
__NOEXC {
1961 return detail::RelConverter<T>::apply(
1962 __sycl_std::__invoke_SignBitSet<detail::internal_rel_ret_t<T>>(x));
1967 #define __SYCL_MARRAY_RELATIONAL_FUNCTION_BINOP_OVERLOAD(NAME) \
1968 template <typename T, \
1969 typename = std::enable_if_t<detail::is_mgenfloat<T>::value>> \
1970 sycl::marray<bool, T::size()> NAME(T x, T y) __NOEXC { \
1971 sycl::marray<bool, T::size()> res; \
1972 for (int i = 0; i < x.size(); i++) { \
1973 res[i] = NAME(x[i], y[i]); \
1978 #define __SYCL_MARRAY_RELATIONAL_FUNCTION_UNOP_OVERLOAD(NAME) \
1979 template <typename T, \
1980 typename = std::enable_if_t<detail::is_mgenfloat<T>::value>> \
1981 sycl::marray<bool, T::size()> NAME(T x) __NOEXC { \
1982 sycl::marray<bool, T::size()> res; \
1983 for (int i = 0; i < x.size(); i++) { \
1984 res[i] = NAME(x[i]); \
2005 template <
typename T>
2006 std::enable_if_t<detail::is_sigeninteger<T>::value,
bool> any(
T x)
__NOEXC {
2011 template <
typename T>
2012 std::enable_if_t<detail::is_vigeninteger<T>::value,
int> any(
T x)
__NOEXC {
2013 return detail::rel_sign_bit_test_ret_t<T>(
2014 __sycl_std::__invoke_Any<detail::rel_sign_bit_test_ret_t<T>>(
2015 detail::rel_sign_bit_test_arg_t<T>(x)));
2019 template <
typename T>
2020 std::enable_if_t<detail::is_sigeninteger<T>::value,
bool>
all(
T x)
__NOEXC {
2025 template <
typename T>
2026 std::enable_if_t<detail::is_vigeninteger<T>::value,
int>
all(
T x)
__NOEXC {
2027 return detail::rel_sign_bit_test_ret_t<T>(
2028 __sycl_std::__invoke_All<detail::rel_sign_bit_test_ret_t<T>>(
2029 detail::rel_sign_bit_test_arg_t<T>(x)));
2033 template <
typename T>
2034 std::enable_if_t<detail::is_gentype<T>::value,
T> bitselect(
T a,
T b,
2036 return __sycl_std::__invoke_bitselect<T>(a, b, c);
2040 template <
typename T>
2041 std::enable_if_t<detail::is_sgentype<T>::value,
T> select(
T a,
T b,
2043 constexpr
size_t SizeT =
sizeof(
T);
2049 using get_select_opencl_builtin_c_arg_type =
typename std::conditional_t<
2054 (detail::is_contained<
2055 T, detail::type_list<long, unsigned long>>::value &&
2056 (SizeT == 4 || SizeT == 8)),
2061 std::conditional_t<SizeT == 8, long long, void>>>>>;
2063 return __sycl_std::__invoke_select<T>(
2064 a, b,
static_cast<get_select_opencl_builtin_c_arg_type
>(c));
2068 template <
typename T,
typename T2>
2070 detail::is_geninteger<T>::value && detail::is_igeninteger<T2>::value,
T>
2072 detail::check_vector_size<T, T2>();
2073 return __sycl_std::__invoke_select<T>(a, b, c);
2077 template <
typename T,
typename T2>
2079 detail::is_geninteger<T>::value && detail::is_ugeninteger<T2>::value,
T>
2081 detail::check_vector_size<T, T2>();
2082 return __sycl_std::__invoke_select<T>(a, b, c);
2086 template <
typename T,
typename T2>
2088 detail::is_svgenfloatf<T>::value && detail::is_genint<T2>::value,
T>
2090 detail::check_vector_size<T, T2>();
2091 return __sycl_std::__invoke_select<T>(a, b, c);
2095 template <
typename T,
typename T2>
2097 detail::is_svgenfloatf<T>::value && detail::is_ugenint<T2>::value,
T>
2099 detail::check_vector_size<T, T2>();
2100 return __sycl_std::__invoke_select<T>(a, b, c);
2104 template <
typename T,
typename T2>
2105 std::enable_if_t<detail::is_svgenfloatd<T>::value &&
2106 detail::is_igeninteger64bit<T2>::value,
2109 detail::check_vector_size<T, T2>();
2110 return __sycl_std::__invoke_select<T>(a, b, c);
2114 template <
typename T,
typename T2>
2115 std::enable_if_t<detail::is_svgenfloatd<T>::value &&
2116 detail::is_ugeninteger64bit<T2>::value,
2119 detail::check_vector_size<T, T2>();
2120 return __sycl_std::__invoke_select<T>(a, b, c);
2124 template <
typename T,
typename T2>
2125 std::enable_if_t<detail::is_svgenfloath<T>::value &&
2126 detail::is_igeninteger16bit<T2>::value,
2129 detail::check_vector_size<T, T2>();
2130 return __sycl_std::__invoke_select<T>(a, b, c);
2134 template <
typename T,
typename T2>
2135 std::enable_if_t<detail::is_svgenfloath<T>::value &&
2136 detail::is_ugeninteger16bit<T2>::value,
2139 detail::check_vector_size<T, T2>();
2140 return __sycl_std::__invoke_select<T>(a, b, c);
2145 template <
typename T,
size_t N>
2146 std::enable_if_t<detail::is_sigeninteger<T>::value,
bool>
2148 return std::any_of(
x.begin(),
x.end(), [](
T i) { return any(i); });
2151 template <
typename T,
size_t N>
2152 std::enable_if_t<detail::is_sigeninteger<T>::value,
bool>
2154 return std::all_of(
x.begin(),
x.end(), [](
T i) { return all(i); });
2157 template <
typename T,
size_t N>
2158 std::enable_if_t<detail::is_gentype<T>::value, marray<T, N>>
2159 bitselect(marray<T, N> a, marray<T, N> b, marray<T, N> c)
__NOEXC {
2161 for (
int i = 0; i < N; i++) {
2162 res[i] = bitselect(a[i], b[i], c[i]);
2167 template <
typename T,
size_t N>
2168 std::enable_if_t<detail::is_gentype<T>::value, marray<T, N>>
2169 select(marray<T, N> a, marray<T, N> b, marray<bool, N> c)
__NOEXC {
2171 for (
int i = 0; i < N; i++) {
2172 res[i] = select(a[i], b[i], c[i]);
2180 #define __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD(NAME) \
2181 template <size_t N> \
2182 inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) \
2184 marray<float, N> res; \
2185 for (size_t i = 0; i < N / 2; i++) { \
2186 auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
2187 detail::to_vec2(x, i * 2)); \
2188 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
2191 res[N - 1] = __sycl_std::__invoke_native_##NAME<float>(x[N - 1]); \
2209 #undef __SYCL_NATIVE_MATH_FUNCTION_OVERLOAD
2211 #define __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD(NAME) \
2212 template <size_t N> \
2213 inline __SYCL_ALWAYS_INLINE marray<float, N> NAME( \
2214 marray<float, N> x, marray<float, N> y) __NOEXC { \
2215 marray<float, N> res; \
2216 for (size_t i = 0; i < N / 2; i++) { \
2217 auto partial_res = __sycl_std::__invoke_native_##NAME<vec<float, 2>>( \
2218 detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
2219 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
2223 __sycl_std::__invoke_native_##NAME<float>(x[N - 1], y[N - 1]); \
2231 #undef __SYCL_NATIVE_MATH_FUNCTION_2_OVERLOAD
2234 template <
typename T>
2235 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
cos(
T x)
__NOEXC {
2236 return __sycl_std::__invoke_native_cos<T>(x);
2240 template <
typename T>
2241 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> divide(
T x,
T y)
__NOEXC {
2242 return __sycl_std::__invoke_native_divide<T>(x, y);
2246 template <
typename T>
2247 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp(
T x)
__NOEXC {
2248 return __sycl_std::__invoke_native_exp<T>(x);
2252 template <
typename T>
2253 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp2(
T x)
__NOEXC {
2254 return __sycl_std::__invoke_native_exp2<T>(x);
2258 template <
typename T>
2259 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> exp10(
T x)
__NOEXC {
2260 return __sycl_std::__invoke_native_exp10<T>(x);
2264 template <
typename T>
2265 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log(
T x)
__NOEXC {
2266 return __sycl_std::__invoke_native_log<T>(x);
2270 template <
typename T>
2271 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log2(
T x)
__NOEXC {
2272 return __sycl_std::__invoke_native_log2<T>(x);
2276 template <
typename T>
2277 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> log10(
T x)
__NOEXC {
2278 return __sycl_std::__invoke_native_log10<T>(x);
2282 template <
typename T>
2283 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
powr(
T x,
T y)
__NOEXC {
2284 return __sycl_std::__invoke_native_powr<T>(x, y);
2288 template <
typename T>
2289 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> recip(
T x)
__NOEXC {
2290 return __sycl_std::__invoke_native_recip<T>(x);
2294 template <
typename T>
2295 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
rsqrt(
T x)
__NOEXC {
2296 return __sycl_std::__invoke_native_rsqrt<T>(x);
2300 template <
typename T>
2301 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sin(
T x)
__NOEXC {
2302 return __sycl_std::__invoke_native_sin<T>(x);
2306 template <
typename T>
2307 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sqrt(
T x)
__NOEXC {
2308 return __sycl_std::__invoke_native_sqrt<T>(x);
2312 template <
typename T>
2313 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> tan(
T x)
__NOEXC {
2314 return __sycl_std::__invoke_native_tan<T>(x);
2318 namespace half_precision {
2320 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(NAME) \
2321 template <size_t N> \
2322 inline __SYCL_ALWAYS_INLINE marray<float, N> NAME(marray<float, N> x) \
2324 marray<float, N> res; \
2325 for (size_t i = 0; i < N / 2; i++) { \
2326 auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
2327 detail::to_vec2(x, i * 2)); \
2328 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
2331 res[N - 1] = __sycl_std::__invoke_half_##NAME<float>(x[N - 1]); \
2349 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD
2351 #define __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD(NAME) \
2352 template <size_t N> \
2353 inline __SYCL_ALWAYS_INLINE marray<float, N> NAME( \
2354 marray<float, N> x, marray<float, N> y) __NOEXC { \
2355 marray<float, N> res; \
2356 for (size_t i = 0; i < N / 2; i++) { \
2357 auto partial_res = __sycl_std::__invoke_half_##NAME<vec<float, 2>>( \
2358 detail::to_vec2(x, i * 2), detail::to_vec2(y, i * 2)); \
2359 std::memcpy(&res[i * 2], &partial_res, sizeof(vec<float, 2>)); \
2363 __sycl_std::__invoke_half_##NAME<float>(x[N - 1], y[N - 1]); \
2371 #undef __SYCL_HALF_PRECISION_MATH_FUNCTION_2_OVERLOAD
2374 template <
typename T>
2375 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
cos(
T x)
__NOEXC {
2376 return __sycl_std::__invoke_half_cos<T>(x);
2380 template <
typename T>
2381 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> divide(
T x,
T y)
__NOEXC {
2382 return __sycl_std::__invoke_half_divide<T>(x, y);
2386 template <
typename T>
2387 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp(
T x)
__NOEXC {
2388 return __sycl_std::__invoke_half_exp<T>(x);
2392 template <
typename T>
2393 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp2(
T x)
__NOEXC {
2394 return __sycl_std::__invoke_half_exp2<T>(x);
2398 template <
typename T>
2399 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> exp10(
T x)
__NOEXC {
2400 return __sycl_std::__invoke_half_exp10<T>(x);
2404 template <
typename T>
2405 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log(
T x)
__NOEXC {
2406 return __sycl_std::__invoke_half_log<T>(x);
2410 template <
typename T>
2411 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log2(
T x)
__NOEXC {
2412 return __sycl_std::__invoke_half_log2<T>(x);
2416 template <
typename T>
2417 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> log10(
T x)
__NOEXC {
2418 return __sycl_std::__invoke_half_log10<T>(x);
2422 template <
typename T>
2423 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
powr(
T x,
T y)
__NOEXC {
2424 return __sycl_std::__invoke_half_powr<T>(x, y);
2428 template <
typename T>
2429 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> recip(
T x)
__NOEXC {
2430 return __sycl_std::__invoke_half_recip<T>(x);
2434 template <
typename T>
2435 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
rsqrt(
T x)
__NOEXC {
2436 return __sycl_std::__invoke_half_rsqrt<T>(x);
2440 template <
typename T>
2441 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sin(
T x)
__NOEXC {
2442 return __sycl_std::__invoke_half_sin<T>(x);
2446 template <
typename T>
2447 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sqrt(
T x)
__NOEXC {
2448 return __sycl_std::__invoke_half_sqrt<T>(x);
2452 template <
typename T>
2453 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> tan(
T x)
__NOEXC {
2454 return __sycl_std::__invoke_half_tan<T>(x);
2459 #ifdef __FAST_MATH__
2462 #define __SYCL_MATH_FUNCTION_OVERLOAD_FM(NAME) \
2463 template <typename T, size_t N> \
2464 inline __SYCL_ALWAYS_INLINE \
2465 std::enable_if_t<std::is_same_v<T, float>, marray<T, N>> \
2466 NAME(marray<T, N> x) __NOEXC { \
2467 return native::NAME(x); \
2481 #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM
2483 template <
typename T,
size_t N>
2485 std::enable_if_t<std::is_same_v<T, float>, marray<T, N>>
2491 template <
typename T>
2492 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
cos(
T x)
__NOEXC {
2497 template <
typename T>
2498 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp(
T x)
__NOEXC {
2503 template <
typename T>
2504 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
exp2(
T x)
__NOEXC {
2509 template <
typename T>
2510 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> exp10(
T x)
__NOEXC {
2511 return native::exp10(x);
2515 template <
typename T>
2516 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log(
T x)
__NOEXC {
2521 template <
typename T>
2522 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
log2(
T x)
__NOEXC {
2527 template <
typename T>
2528 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> log10(
T x)
__NOEXC {
2529 return native::log10(x);
2533 template <
typename T>
2534 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
powr(
T x,
T y)
__NOEXC {
2539 template <
typename T>
2540 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
rsqrt(
T x)
__NOEXC {
2545 template <
typename T>
2546 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sin(
T x)
__NOEXC {
2551 template <
typename T>
2552 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T>
sqrt(
T x)
__NOEXC {
2557 template <
typename T>
2558 std::enable_if_t<detail::is_svgenfloatf<T>::value,
T> tan(
T x)
__NOEXC {
2559 return native::tan(x);
2562 #endif // __FAST_MATH__
2566 #ifdef __SYCL_DEVICE_ONLY__
2664 __imf_ullmax(
unsigned long long int x,
unsigned long long int y);
2666 __imf_ullmin(
unsigned long long int x,
unsigned long long int y);
2673 __imf_brevll(
unsigned long long int x);
2675 __imf_byte_perm(
unsigned int x,
unsigned int y,
unsigned int s);
2685 __imf_usad(
unsigned int x,
unsigned int y,
unsigned int z);
2700 __imf_umul64hi(
unsigned long long int x,
unsigned long long int y);
2832 __imf_bfloat162ushort_rd(uint16_t x);
2834 __imf_bfloat162ushort_rn(uint16_t x);
2836 __imf_bfloat162ushort_ru(uint16_t x);
2838 __imf_bfloat162ushort_rz(uint16_t x);
2840 __imf_bfloat162ull_rd(uint16_t x);
2842 __imf_bfloat162ull_rn(uint16_t x);
2844 __imf_bfloat162ull_ru(uint16_t x);
2846 __imf_bfloat162ull_rz(uint16_t x);
2865 __imf_ushort2bfloat16_rd(
unsigned short x);
2867 __imf_ushort2bfloat16_rn(
unsigned short x);
2869 __imf_ushort2bfloat16_ru(
unsigned short x);
2871 __imf_ushort2bfloat16_rz(
unsigned short x);
2877 __imf_ull2bfloat16_rd(
unsigned long long x);
2879 __imf_ull2bfloat16_rn(
unsigned long long x);
2881 __imf_ull2bfloat16_ru(
unsigned long long x);
2883 __imf_ull2bfloat16_rz(
unsigned long long x);
2899 __imf_bfloat16_as_ushort(uint16_t x);
2902 __imf_ushort_as_bfloat16(
unsigned short x);
2953 __imf_ull2double_rd(
unsigned long long int x);
2955 __imf_ull2double_rn(
unsigned long long int x);
2957 __imf_ull2double_ru(
unsigned long long int x);
2959 __imf_ull2double_rz(
unsigned long long int x);
2961 __imf_double2ull_rd(
double x);
2963 __imf_double2ull_rn(
double x);
2965 __imf_double2ull_ru(
double x);
2967 __imf_double2ull_rz(
double x);
3193 double c,
double d);
3199 #elif defined(_WIN32)
3223 const wchar_t *wfile,
unsigned line);
3226 #endif // __SYCL_DEVICE_ONLY__