22 #ifdef __SYCL_DEVICE_ONLY__
25 namespace __sycl_std = __host_std;
29 #define __FAST_MATH_GENFLOAT(T) \
30 (detail::is_genfloatd<T>::value || detail::is_genfloath<T>::value)
32 #define __FAST_MATH_GENFLOAT(T) (detail::is_genfloat<T>::value)
39 return __sycl_std::__invoke_acos<T>(x);
45 return __sycl_std::__invoke_acosh<T>(x);
51 return __sycl_std::__invoke_acospi<T>(x);
57 return __sycl_std::__invoke_asin<T>(x);
63 return __sycl_std::__invoke_asinh<T>(x);
69 return __sycl_std::__invoke_asinpi<T>(x);
75 return __sycl_std::__invoke_atan<T>(y_over_x);
80 detail::enable_if_t<detail::is_genfloat<T>::value, T>
atan2(T y, T x)
__NOEXC {
81 return __sycl_std::__invoke_atan2<T>(y, x);
87 return __sycl_std::__invoke_atanh<T>(x);
93 return __sycl_std::__invoke_atanpi<T>(x);
100 return __sycl_std::__invoke_atan2pi<T>(y, x);
104 template <
typename T>
106 return __sycl_std::__invoke_cbrt<T>(x);
110 template <
typename T>
112 return __sycl_std::__invoke_ceil<T>(x);
116 template <
typename T>
119 return __sycl_std::__invoke_copysign<T>(x, y);
123 template <
typename T>
125 return __sycl_std::__invoke_cos<T>(x);
129 template <
typename T>
131 return __sycl_std::__invoke_cosh<T>(x);
135 template <
typename T>
137 return __sycl_std::__invoke_cospi<T>(x);
141 template <
typename T>
143 return __sycl_std::__invoke_erfc<T>(x);
147 template <
typename T>
149 return __sycl_std::__invoke_erf<T>(x);
153 template <
typename T>
155 return __sycl_std::__invoke_exp<T>(x);
159 template <
typename T>
161 return __sycl_std::__invoke_exp2<T>(x);
165 template <
typename T>
167 return __sycl_std::__invoke_exp10<T>(x);
171 template <
typename T>
173 return __sycl_std::__invoke_expm1<T>(x);
177 template <
typename T>
179 return __sycl_std::__invoke_fabs<T>(x);
183 template <
typename T>
185 return __sycl_std::__invoke_fdim<T>(x, y);
189 template <
typename T>
191 return __sycl_std::__invoke_floor<T>(x);
195 template <
typename T>
198 return __sycl_std::__invoke_fma<T>(
a, b, c);
202 template <
typename T>
204 return __sycl_std::__invoke_fmax<T>(x, y);
208 template <
typename T>
209 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
211 return __sycl_std::__invoke_fmax<T>(x, T(y));
215 template <
typename T>
217 return __sycl_std::__invoke_fmin<T>(x, y);
221 template <
typename T>
222 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
224 return __sycl_std::__invoke_fmin<T>(x, T(y));
228 template <
typename T>
229 detail::enable_if_t<detail::is_genfloat<T>::value, T>
fmod(T x, T y)
__NOEXC {
230 return __sycl_std::__invoke_fmod<T>(x, y);
234 template <
typename T,
typename T2>
236 detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
238 detail::check_vector_size<T, T2>();
239 return __sycl_std::__invoke_fract<T>(x, iptr);
243 template <
typename T,
typename T2>
245 detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
247 detail::check_vector_size<T, T2>();
248 return __sycl_std::__invoke_frexp<T>(x,
exp);
252 template <
typename T>
254 return __sycl_std::__invoke_hypot<T>(x, y);
258 template <
typename T,
259 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
261 return __sycl_std::__invoke_ilogb<detail::change_base_type_t<T, int>>(x);
267 template <
typename T>
270 return __sycl_std::__invoke_ldexp<T>(x, k);
274 template <
typename T>
277 return __sycl_std::__invoke_ldexp<T>(x,
vec<
int, T::size()>(k));
281 template <
typename T,
typename T2>
283 detail::is_vgenfloat<T>::value && detail::is_intn<T2>::value, T>
285 detail::check_vector_size<T, T2>();
286 return __sycl_std::__invoke_ldexp<T>(x, k);
290 template <
typename T>
292 return __sycl_std::__invoke_lgamma<T>(x);
296 template <
typename T,
typename T2>
298 detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
300 detail::check_vector_size<T, T2>();
301 return __sycl_std::__invoke_lgamma_r<T>(x, signp);
305 template <
typename T>
307 return __sycl_std::__invoke_log<T>(x);
311 template <
typename T>
313 return __sycl_std::__invoke_log2<T>(x);
317 template <
typename T>
319 return __sycl_std::__invoke_log10<T>(x);
323 template <
typename T>
325 return __sycl_std::__invoke_log1p<T>(x);
329 template <
typename T>
331 return __sycl_std::__invoke_logb<T>(x);
335 template <
typename T>
338 return __sycl_std::__invoke_mad<T>(
a, b, c);
342 template <
typename T>
344 return __sycl_std::__invoke_maxmag<T>(x, y);
348 template <
typename T>
350 return __sycl_std::__invoke_minmag<T>(x, y);
354 template <
typename T,
typename T2>
356 detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
358 detail::check_vector_size<T, T2>();
359 return __sycl_std::__invoke_modf<T>(x, iptr);
362 template <
typename T,
363 typename = detail::enable_if_t<detail::is_nan_type<T>::value, T>>
365 return __sycl_std::__invoke_nan<detail::nan_return_t<T>>(
370 template <
typename T>
373 return __sycl_std::__invoke_nextafter<T>(x, y);
377 template <
typename T>
379 return __sycl_std::__invoke_pow<T>(x, y);
383 template <
typename T,
typename T2>
385 detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
387 detail::check_vector_size<T, T2>();
388 return __sycl_std::__invoke_pown<T>(x, y);
392 template <
typename T>
394 return __sycl_std::__invoke_powr<T>(x, y);
398 template <
typename T>
401 return __sycl_std::__invoke_remainder<T>(x, y);
405 template <
typename T,
typename T2>
407 detail::is_genfloat<T>::value && detail::is_genintptr<T2>::value, T>
409 detail::check_vector_size<T, T2>();
410 return __sycl_std::__invoke_remquo<T>(x, y, quo);
414 template <
typename T>
416 return __sycl_std::__invoke_rint<T>(x);
420 template <
typename T,
typename T2>
422 detail::is_genfloat<T>::value && detail::is_genint<T2>::value, T>
424 detail::check_vector_size<T, T2>();
425 return __sycl_std::__invoke_rootn<T>(x, y);
429 template <
typename T>
431 return __sycl_std::__invoke_round<T>(x);
435 template <
typename T>
437 return __sycl_std::__invoke_rsqrt<T>(x);
441 template <
typename T>
443 return __sycl_std::__invoke_sin<T>(x);
447 template <
typename T,
typename T2>
449 detail::is_genfloat<T>::value && detail::is_genfloatptr<T2>::value, T>
451 detail::check_vector_size<T, T2>();
452 return __sycl_std::__invoke_sincos<T>(x, cosval);
456 template <
typename T>
458 return __sycl_std::__invoke_sinh<T>(x);
462 template <
typename T>
464 return __sycl_std::__invoke_sinpi<T>(x);
468 template <
typename T>
470 return __sycl_std::__invoke_sqrt<T>(x);
474 template <
typename T>
476 return __sycl_std::__invoke_tan<T>(x);
480 template <
typename T>
482 return __sycl_std::__invoke_tanh<T>(x);
486 template <
typename T>
488 return __sycl_std::__invoke_tanpi<T>(x);
492 template <
typename T>
494 return __sycl_std::__invoke_tgamma<T>(x);
498 template <
typename T>
500 return __sycl_std::__invoke_trunc<T>(x);
505 template <
typename T>
508 return __sycl_std::__invoke_fclamp<T>(x, minval, maxval);
514 template <
typename T>
515 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
516 clamp(T x,
typename T::element_type minval,
517 typename T::element_type maxval)
__NOEXC {
518 return __sycl_std::__invoke_fclamp<T>(x, T(minval), T(maxval));
522 template <
typename T>
523 detail::enable_if_t<detail::is_genfloat<T>::value, T>
525 return __sycl_std::__invoke_degrees<T>(
radians);
529 template <
typename T>
531 return __sycl_std::__invoke_fabs<T>(x);
535 template <
typename T>
537 return __sycl_std::__invoke_fmax_common<T>(x, y);
543 template <
typename T>
544 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>(
max)(
545 T x,
typename T::element_type y)
__NOEXC {
546 return __sycl_std::__invoke_fmax_common<T>(x, T(y));
550 template <
typename T>
551 detail::enable_if_t<detail::is_genfloat<T>::value,
T>(
min)(
T x,
T y)
__NOEXC {
552 return __sycl_std::__invoke_fmin_common<T>(x, y);
558 template <
typename T>
559 detail::enable_if_t<detail::is_vgenfloat<T>::value,
T>(
min)(
560 T x,
typename T::element_type y)
__NOEXC {
561 return __sycl_std::__invoke_fmin_common<T>(x,
T(y));
565 template <
typename T>
568 return __sycl_std::__invoke_mix<T>(x, y,
a);
574 template <
typename T>
575 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
577 return __sycl_std::__invoke_mix<T>(x, y, T(
a));
581 template <
typename T>
582 detail::enable_if_t<detail::is_genfloat<T>::value, T>
584 return __sycl_std::__invoke_radians<T>(
degrees);
588 template <
typename T>
591 return __sycl_std::__invoke_step<T>(edge, x);
597 template <
typename T>
598 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
600 return __sycl_std::__invoke_step<T>(T(edge), x);
604 template <
typename T>
605 detail::enable_if_t<detail::is_genfloat<T>::value, T>
607 return __sycl_std::__invoke_smoothstep<T>(edge0, edge1, x);
613 template <
typename T>
614 detail::enable_if_t<detail::is_vgenfloat<T>::value, T>
615 smoothstep(
typename T::element_type edge0,
typename T::element_type edge1,
617 return __sycl_std::__invoke_smoothstep<T>(T(edge0), T(edge1), x);
621 template <
typename T>
623 return __sycl_std::__invoke_sign<T>(x);
628 template <
typename T>
630 return __sycl_std::__invoke_u_abs<T>(x);
634 template <
typename T>
635 detail::enable_if_t<detail::is_igeninteger<T>::value,
636 detail::make_unsigned_t<T>>
638 return __sycl_std::__invoke_s_abs<detail::make_unsigned_t<T>>(x);
642 template <
typename T>
645 return __sycl_std::__invoke_u_abs_diff<T>(x, y);
649 template <
typename T>
650 detail::enable_if_t<detail::is_igeninteger<T>::value,
651 detail::make_unsigned_t<T>>
653 return __sycl_std::__invoke_s_abs_diff<detail::make_unsigned_t<T>>(x, y);
657 template <
typename T>
660 return __sycl_std::__invoke_s_add_sat<T>(x, y);
664 template <
typename T>
667 return __sycl_std::__invoke_u_add_sat<T>(x, y);
671 template <
typename T>
674 return __sycl_std::__invoke_s_hadd<T>(x, y);
678 template <
typename T>
681 return __sycl_std::__invoke_u_hadd<T>(x, y);
685 template <
typename T>
688 return __sycl_std::__invoke_s_rhadd<T>(x, y);
692 template <
typename T>
695 return __sycl_std::__invoke_u_rhadd<T>(x, y);
699 template <
typename T>
700 detail::enable_if_t<detail::is_igeninteger<T>::value, T>
702 return __sycl_std::__invoke_s_clamp<T>(x, minval, maxval);
706 template <
typename T>
707 detail::enable_if_t<detail::is_ugeninteger<T>::value, T>
709 return __sycl_std::__invoke_u_clamp<T>(x, minval, maxval);
713 template <
typename T>
714 detail::enable_if_t<detail::is_vigeninteger<T>::value, T>
715 clamp(T x,
typename T::element_type minval,
716 typename T::element_type maxval)
__NOEXC {
717 return __sycl_std::__invoke_s_clamp<T>(x, T(minval), T(maxval));
721 template <
typename T>
722 detail::enable_if_t<detail::is_vugeninteger<T>::value, T>
723 clamp(T x,
typename T::element_type minval,
724 typename T::element_type maxval)
__NOEXC {
725 return __sycl_std::__invoke_u_clamp<T>(x, T(minval), T(maxval));
729 template <
typename T>
731 return __sycl_std::__invoke_clz<T>(x);
735 template <
typename T>
737 return __sycl_std::__invoke_ctz<T>(x);
743 template <
typename T>
745 "'sycl::ext::intel::ctz' is deprecated, use 'sycl::ctz' instead")
754 using namespace ext::intel;
758 template <
typename T>
761 return __sycl_std::__invoke_s_mad_hi<T>(x, y, z);
765 template <
typename T>
768 return __sycl_std::__invoke_u_mad_hi<T>(x, y, z);
772 template <
typename T>
775 return __sycl_std::__invoke_s_mad_sat<T>(
a, b, c);
779 template <
typename T>
782 return __sycl_std::__invoke_u_mad_sat<T>(
a, b, c);
786 template <
typename T>
787 detail::enable_if_t<detail::is_igeninteger<T>::value, T>(
max)(T x,
789 return __sycl_std::__invoke_s_max<T>(x, y);
793 template <
typename T>
794 detail::enable_if_t<detail::is_ugeninteger<T>::value,
T>(
max)(
T x,
796 return __sycl_std::__invoke_u_max<T>(x, y);
800 template <
typename T>
801 detail::enable_if_t<detail::is_vigeninteger<T>::value,
T>(
max)(
802 T x,
typename T::element_type y)
__NOEXC {
803 return __sycl_std::__invoke_s_max<T>(x,
T(y));
807 template <
typename T>
808 detail::enable_if_t<detail::is_vugeninteger<T>::value,
T>(
max)(
809 T x,
typename T::element_type y)
__NOEXC {
810 return __sycl_std::__invoke_u_max<T>(x,
T(y));
814 template <
typename T>
815 detail::enable_if_t<detail::is_igeninteger<T>::value,
T>(
min)(
T x,
817 return __sycl_std::__invoke_s_min<T>(x, y);
821 template <
typename T>
822 detail::enable_if_t<detail::is_ugeninteger<T>::value,
T>(
min)(
T x,
824 return __sycl_std::__invoke_u_min<T>(x, y);
828 template <
typename T>
829 detail::enable_if_t<detail::is_vigeninteger<T>::value,
T>(
min)(
830 T x,
typename T::element_type y)
__NOEXC {
831 return __sycl_std::__invoke_s_min<T>(x,
T(y));
835 template <
typename T>
836 detail::enable_if_t<detail::is_vugeninteger<T>::value,
T>(
min)(
837 T x,
typename T::element_type y)
__NOEXC {
838 return __sycl_std::__invoke_u_min<T>(x,
T(y));
842 template <
typename T>
845 return __sycl_std::__invoke_s_mul_hi<T>(x, y);
849 template <
typename T>
852 return __sycl_std::__invoke_u_mul_hi<T>(x, y);
856 template <
typename T>
859 return __sycl_std::__invoke_rotate<T>(v, i);
863 template <
typename T>
866 return __sycl_std::__invoke_s_sub_sat<T>(x, y);
870 template <
typename T>
873 return __sycl_std::__invoke_u_sub_sat<T>(x, y);
877 template <
typename T>
878 detail::enable_if_t<detail::is_ugeninteger8bit<T>::value,
879 detail::make_larger_t<T>>
881 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
885 template <
typename T,
typename T2>
886 detail::enable_if_t<detail::is_igeninteger8bit<T>::value &&
887 detail::is_ugeninteger8bit<T2>::value,
888 detail::make_larger_t<T>>
890 detail::check_vector_size<T, T2>();
891 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
895 template <
typename T>
896 detail::enable_if_t<detail::is_ugeninteger16bit<T>::value,
897 detail::make_larger_t<T>>
899 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
903 template <
typename T,
typename T2>
904 detail::enable_if_t<detail::is_igeninteger16bit<T>::value &&
905 detail::is_ugeninteger16bit<T2>::value,
906 detail::make_larger_t<T>>
908 detail::check_vector_size<T, T2>();
909 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
913 template <
typename T>
914 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value,
915 detail::make_larger_t<T>>
917 return __sycl_std::__invoke_u_upsample<detail::make_larger_t<T>>(hi, lo);
921 template <
typename T,
typename T2>
922 detail::enable_if_t<detail::is_igeninteger32bit<T>::value &&
923 detail::is_ugeninteger32bit<T2>::value,
924 detail::make_larger_t<T>>
926 detail::check_vector_size<T, T2>();
927 return __sycl_std::__invoke_s_upsample<detail::make_larger_t<T>>(hi, lo);
931 template <
typename T>
933 return __sycl_std::__invoke_popcount<T>(x);
938 template <
typename T>
939 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
941 return __sycl_std::__invoke_s_mad24<T>(x, y, z);
946 template <
typename T>
947 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
949 return __sycl_std::__invoke_u_mad24<T>(x, y, z);
953 template <
typename T>
954 detail::enable_if_t<detail::is_igeninteger32bit<T>::value, T>
956 return __sycl_std::__invoke_s_mul24<T>(x, y);
960 template <
typename T>
961 detail::enable_if_t<detail::is_ugeninteger32bit<T>::value, T>
963 return __sycl_std::__invoke_u_mul24<T>(x, y);
973 template <
typename T>
976 return __sycl_std::__invoke_cross<T>(p0, p1);
982 template <
typename T>
988 template <
typename T>
989 detail::enable_if_t<detail::is_vgengeofloat<T>::value,
float>
991 return __sycl_std::__invoke_Dot<float>(p0, p1);
995 template <
typename T>
996 detail::enable_if_t<detail::is_vgengeodouble<T>::value,
double>
998 return __sycl_std::__invoke_Dot<double>(p0, p1);
1002 template <
typename T>
1005 return __sycl_std::__invoke_Dot<half>(p0, p1);
1009 template <
typename T,
1010 typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1012 return __sycl_std::__invoke_distance<float>(p0, p1);
1016 template <
typename T,
1017 typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1019 return __sycl_std::__invoke_distance<double>(p0, p1);
1023 template <
typename T,
1024 typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1026 return __sycl_std::__invoke_distance<half>(p0, p1);
1030 template <
typename T,
1031 typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1033 return __sycl_std::__invoke_length<float>(p);
1037 template <
typename T,
1038 typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1040 return __sycl_std::__invoke_length<double>(p);
1044 template <
typename T,
1045 typename = detail::enable_if_t<detail::is_gengeohalf<T>::value, T>>
1047 return __sycl_std::__invoke_length<half>(p);
1051 template <
typename T>
1052 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1054 return __sycl_std::__invoke_normalize<T>(p);
1058 template <
typename T>
1059 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1061 return __sycl_std::__invoke_normalize<T>(p);
1065 template <
typename T>
1067 return __sycl_std::__invoke_normalize<T>(p);
1071 template <
typename T,
1072 typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1074 return __sycl_std::__invoke_fast_distance<float>(p0, p1);
1078 template <
typename T,
1079 typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1081 return __sycl_std::__invoke_fast_distance<double>(p0, p1);
1085 template <
typename T,
1086 typename = detail::enable_if_t<detail::is_gengeofloat<T>::value, T>>
1088 return __sycl_std::__invoke_fast_length<float>(p);
1092 template <
typename T,
1093 typename = detail::enable_if_t<detail::is_gengeodouble<T>::value, T>>
1095 return __sycl_std::__invoke_fast_length<double>(p);
1099 template <
typename T>
1100 detail::enable_if_t<detail::is_gengeofloat<T>::value, T>
1102 return __sycl_std::__invoke_fast_normalize<T>(p);
1106 template <
typename T>
1107 detail::enable_if_t<detail::is_gengeodouble<T>::value, T>
1109 return __sycl_std::__invoke_fast_normalize<T>(p);
1118 template <
typename T,
1119 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1130 template <
typename T,
1131 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1142 template <
typename T,
1143 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1154 template <
typename T,
1155 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1166 template <
typename T,
1167 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1178 template <
typename T,
1179 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1190 template <
typename T,
1191 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1202 template <
typename T,
1203 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1214 template <
typename T,
1215 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1226 template <
typename T,
1227 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1238 template <
typename T,
1239 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1250 template <
typename T,
1251 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1262 template <
typename T,
1263 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1274 template <
typename T,
1275 typename = detail::enable_if_t<detail::is_genfloat<T>::value, T>>
1282 template <
typename T>
1288 template <
typename T>
1296 template <
typename T>
1302 template <
typename T>
1310 template <
typename T>
1313 return __sycl_std::__invoke_bitselect<T>(
a, b, c);
1317 template <
typename T>
1320 return __sycl_std::__invoke_select<T>(
a, b,
static_cast<int>(c));
1324 template <
typename T,
typename T2>
1326 detail::is_geninteger<T>::value && detail::is_igeninteger<T2>::value, T>
1328 detail::check_vector_size<T, T2>();
1329 return __sycl_std::__invoke_select<T>(
a, b, c);
1333 template <
typename T,
typename T2>
1335 detail::is_geninteger<T>::value && detail::is_ugeninteger<T2>::value, T>
1337 detail::check_vector_size<T, T2>();
1338 return __sycl_std::__invoke_select<T>(
a, b, c);
1342 template <
typename T,
typename T2>
1344 detail::is_genfloatf<T>::value && detail::is_genint<T2>::value, T>
1346 detail::check_vector_size<T, T2>();
1347 return __sycl_std::__invoke_select<T>(
a, b, c);
1351 template <
typename T,
typename T2>
1353 detail::is_genfloatf<T>::value && detail::is_ugenint<T2>::value, T>
1355 detail::check_vector_size<T, T2>();
1356 return __sycl_std::__invoke_select<T>(
a, b, c);
1360 template <
typename T,
typename T2>
1362 detail::is_genfloatd<T>::value && detail::is_igeninteger64bit<T2>::value, T>
1364 detail::check_vector_size<T, T2>();
1365 return __sycl_std::__invoke_select<T>(
a, b, c);
1369 template <
typename T,
typename T2>
1371 detail::is_genfloatd<T>::value && detail::is_ugeninteger64bit<T2>::value, T>
1373 detail::check_vector_size<T, T2>();
1374 return __sycl_std::__invoke_select<T>(
a, b, c);
1378 template <
typename T,
typename T2>
1380 detail::is_genfloath<T>::value && detail::is_igeninteger16bit<T2>::value, T>
1382 detail::check_vector_size<T, T2>();
1383 return __sycl_std::__invoke_select<T>(
a, b, c);
1387 template <
typename T,
typename T2>
1389 detail::is_genfloath<T>::value && detail::is_ugeninteger16bit<T2>::value, T>
1391 detail::check_vector_size<T, T2>();
1392 return __sycl_std::__invoke_select<T>(
a, b, c);
1398 template <
typename T>
1400 return __sycl_std::__invoke_native_cos<T>(x);
1404 template <
typename T>
1407 return __sycl_std::__invoke_native_divide<T>(x, y);
1411 template <
typename T>
1413 return __sycl_std::__invoke_native_exp<T>(x);
1417 template <
typename T>
1419 return __sycl_std::__invoke_native_exp2<T>(x);
1423 template <
typename T>
1425 return __sycl_std::__invoke_native_exp10<T>(x);
1429 template <
typename T>
1431 return __sycl_std::__invoke_native_log<T>(x);
1435 template <
typename T>
1437 return __sycl_std::__invoke_native_log2<T>(x);
1441 template <
typename T>
1443 return __sycl_std::__invoke_native_log10<T>(x);
1447 template <
typename T>
1449 return __sycl_std::__invoke_native_powr<T>(x, y);
1453 template <
typename T>
1455 return __sycl_std::__invoke_native_recip<T>(x);
1459 template <
typename T>
1461 return __sycl_std::__invoke_native_rsqrt<T>(x);
1465 template <
typename T>
1467 return __sycl_std::__invoke_native_sin<T>(x);
1471 template <
typename T>
1473 return __sycl_std::__invoke_native_sqrt<T>(x);
1477 template <
typename T>
1479 return __sycl_std::__invoke_native_tan<T>(x);
1483 namespace half_precision {
1486 template <
typename T>
1488 return __sycl_std::__invoke_half_cos<T>(x);
1492 template <
typename T>
1495 return __sycl_std::__invoke_half_divide<T>(x, y);
1499 template <
typename T>
1501 return __sycl_std::__invoke_half_exp<T>(x);
1505 template <
typename T>
1507 return __sycl_std::__invoke_half_exp2<T>(x);
1511 template <
typename T>
1513 return __sycl_std::__invoke_half_exp10<T>(x);
1517 template <
typename T>
1519 return __sycl_std::__invoke_half_log<T>(x);
1523 template <
typename T>
1525 return __sycl_std::__invoke_half_log2<T>(x);
1529 template <
typename T>
1531 return __sycl_std::__invoke_half_log10<T>(x);
1535 template <
typename T>
1537 return __sycl_std::__invoke_half_powr<T>(x, y);
1541 template <
typename T>
1543 return __sycl_std::__invoke_half_recip<T>(x);
1547 template <
typename T>
1549 return __sycl_std::__invoke_half_rsqrt<T>(x);
1553 template <
typename T>
1555 return __sycl_std::__invoke_half_sin<T>(x);
1559 template <
typename T>
1561 return __sycl_std::__invoke_half_sqrt<T>(x);
1565 template <
typename T>
1567 return __sycl_std::__invoke_half_tan<T>(x);
1572 #ifdef __FAST_MATH__
1575 template <
typename T>
1576 detail::enable_if_t<detail::is_genfloatf<T>::value, T>
cos(T x)
__NOEXC {
1581 template <
typename T>
1582 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
exp(
T x)
__NOEXC {
1587 template <
typename T>
1588 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
exp2(
T x)
__NOEXC {
1593 template <
typename T>
1594 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
exp10(
T x)
__NOEXC {
1599 template <
typename T>
1600 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
log(
T x)
__NOEXC {
1605 template <
typename T>
1606 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
log2(
T x)
__NOEXC {
1611 template <
typename T>
1612 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
log10(
T x)
__NOEXC {
1617 template <
typename T>
1618 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
powr(
T x,
T y)
__NOEXC {
1623 template <
typename T>
1624 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
rsqrt(
T x)
__NOEXC {
1629 template <
typename T>
1630 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
sin(
T x)
__NOEXC {
1635 template <
typename T>
1636 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
sqrt(
T x)
__NOEXC {
1641 template <
typename T>
1642 detail::enable_if_t<detail::is_genfloatf<T>::value,
T>
tan(
T x)
__NOEXC {
1646 #endif // __FAST_MATH__
1650 #ifdef __SYCL_DEVICE_ONLY__
1658 extern SYCL_EXTERNAL lldiv_t lldiv(
long long int x,
long long int y);
1697 extern SYCL_EXTERNAL float remquof(
float x,
float y,
int *q);
1739 extern SYCL_EXTERNAL void *memset(
void *dest,
int c,
size_t n);
1740 extern SYCL_EXTERNAL int memcmp(
const void *s1,
const void *s2,
size_t n);
1744 extern SYCL_EXTERNAL void __assert_fail(
const char *expr,
const char *file,
1745 unsigned int line,
const char *func);
1760 extern SYCL_EXTERNAL float __complex__ cprojf(
float __complex__ z);
1761 extern SYCL_EXTERNAL double __complex__ cproj(
double __complex__ z);
1762 extern SYCL_EXTERNAL float __complex__ cexpf(
float __complex__ z);
1763 extern SYCL_EXTERNAL double __complex__ cexp(
double __complex__ z);
1764 extern SYCL_EXTERNAL float __complex__ clogf(
float __complex__ z);
1765 extern SYCL_EXTERNAL double __complex__ clog(
double __complex__ z);
1766 extern SYCL_EXTERNAL float __complex__ cpowf(
float __complex__ z);
1767 extern SYCL_EXTERNAL double __complex__ cpow(
double __complex__ z);
1768 extern SYCL_EXTERNAL float __complex__ csqrtf(
float __complex__ z);
1769 extern SYCL_EXTERNAL double __complex__ csqrt(
double __complex__ z);
1770 extern SYCL_EXTERNAL float __complex__ csinhf(
float __complex__ z);
1771 extern SYCL_EXTERNAL double __complex__ csinh(
double __complex__ z);
1772 extern SYCL_EXTERNAL float __complex__ ccoshf(
float __complex__ z);
1773 extern SYCL_EXTERNAL double __complex__ ccosh(
double __complex__ z);
1774 extern SYCL_EXTERNAL float __complex__ ctanhf(
float __complex__ z);
1775 extern SYCL_EXTERNAL double __complex__ ctanh(
double __complex__ z);
1776 extern SYCL_EXTERNAL float __complex__ csinf(
float __complex__ z);
1777 extern SYCL_EXTERNAL double __complex__ csin(
double __complex__ z);
1778 extern SYCL_EXTERNAL float __complex__ ccosf(
float __complex__ z);
1779 extern SYCL_EXTERNAL double __complex__ ccos(
double __complex__ z);
1780 extern SYCL_EXTERNAL float __complex__ ctanf(
float __complex__ z);
1781 extern SYCL_EXTERNAL double __complex__ ctan(
double __complex__ z);
1782 extern SYCL_EXTERNAL float __complex__ cacosf(
float __complex__ z);
1783 extern SYCL_EXTERNAL double __complex__ cacos(
double __complex__ z);
1784 extern SYCL_EXTERNAL float __complex__ cacoshf(
float __complex__ z);
1785 extern SYCL_EXTERNAL double __complex__ cacosh(
double __complex__ z);
1786 extern SYCL_EXTERNAL float __complex__ casinf(
float __complex__ z);
1787 extern SYCL_EXTERNAL double __complex__ casin(
double __complex__ z);
1788 extern SYCL_EXTERNAL float __complex__ casinhf(
float __complex__ z);
1789 extern SYCL_EXTERNAL double __complex__ casinh(
double __complex__ z);
1790 extern SYCL_EXTERNAL float __complex__ catanf(
float __complex__ z);
1791 extern SYCL_EXTERNAL double __complex__ catan(
double __complex__ z);
1792 extern SYCL_EXTERNAL float __complex__ catanhf(
float __complex__ z);
1793 extern SYCL_EXTERNAL double __complex__ catanh(
double __complex__ z);
1794 extern SYCL_EXTERNAL float __complex__ cpolarf(
float rho,
float theta);
1795 extern SYCL_EXTERNAL double __complex__ cpolar(
double rho,
double theta);
1796 extern SYCL_EXTERNAL float __complex__ __mulsc3(
float a,
float b,
float c,
1798 extern SYCL_EXTERNAL double __complex__ __muldc3(
double a,
double b,
double c,
1800 extern SYCL_EXTERNAL float __complex__ __divsc3(
float a,
float b,
float c,
1802 extern SYCL_EXTERNAL double __complex__ __divdc3(
float a,
float b,
float c,
1805 #elif defined(_WIN32)
1818 extern SYCL_EXTERNAL short _Exp(
double *px,
double y,
short eoff);
1824 extern SYCL_EXTERNAL short _FExp(
float *px,
float y,
short eoff);
1828 extern SYCL_EXTERNAL void _wassert(
const wchar_t *wexpr,
const wchar_t *wfile,
1832 #endif // __SYCL_DEVICE_ONLY__