14 inline namespace _V1 {
16 template <
typename... Ts>
struct last_int_rest_same {
17 static constexpr
bool value = []() constexpr {
18 constexpr
auto N =
sizeof...(Ts);
19 using first_type =
typename first_type<Ts...>::type;
20 if (!builtin_same_shape_v<first_type>)
24 std::conditional_t<is_vec_or_swizzle_v<first_type>, int32_t,
int>;
28 ? builtin_same_shape_v<Ts>
30 && std::is_same_v<get_elem_type_t<Ts>, int_type>
31 : builtin_same_or_swizzle_v<first_type, Ts>)));
34 template <
typename... Ts>
struct last_intptr_rest_same {
35 static constexpr
bool value = []() constexpr {
36 constexpr
auto N =
sizeof...(Ts);
37 using first_type =
typename first_type<Ts...>::type;
38 if (!builtin_same_shape_v<first_type>)
42 std::conditional_t<is_vec_or_swizzle_v<first_type>, int32_t,
int>;
47 std::is_pointer_v<Ts> ||
48 (is_multi_ptr_v<Ts> && has_writeable_addr_space_v<Ts>)) &&
53 : builtin_same_or_swizzle_v<first_type, Ts>)));
58 non_scalar_only, same_elem_type)
60 fp_elem_type, any_shape, same_elem_type)
62 non_scalar_only, last_int_rest_same)
64 fp_elem_type, scalar_only, last_intptr_rest_same)
66 fp_elem_type, non_scalar_only, last_intptr_rest_same)
70 template <
typename... Ts>
71 inline constexpr
bool builtin_enable_math_allow_scalar_v =
77 #ifdef __SYCL_DEVICE_ONLY__
82 #define BUILTIN_GENF_DEVICE_COMMON(NUM_ARGS, NAME, VEC_IMPL) \
83 inline double NAME(NUM_ARGS##_TYPE_ARG(double)) { \
84 return __spirv_ocl_##NAME(NUM_ARGS##_ARG); \
86 inline half NAME(NUM_ARGS##_TYPE_ARG(half)) { \
87 return __spirv_ocl_##NAME(NUM_ARGS##_CONVERTED_ARG); \
89 DEVICE_IMPL_TEMPLATE(NUM_ARGS, NAME, builtin_enable_math_t, VEC_IMPL)
91 #define BUILTIN_GENF(NUM_ARGS, NAME) \
92 inline float NAME(NUM_ARGS##_TYPE_ARG(float)) { \
93 return __spirv_ocl_##NAME(NUM_ARGS##_ARG); \
95 BUILTIN_GENF_DEVICE_COMMON(NUM_ARGS, NAME, __spirv_ocl_##NAME)
97 #define BUILTIN_GENF_NATIVE_OPT(NUM_ARGS, NAME) \
99 template <typename T> \
100 decltype(auto) maybe_fmf_##NAME(NUM_ARGS##_TYPE_ARG(T)) { \
101 if constexpr (use_fast_math_v<T>) { \
102 return __spirv_ocl_native_##NAME(NUM_ARGS##_ARG); \
104 return __spirv_ocl_##NAME(NUM_ARGS##_ARG); \
108 inline float NAME(NUM_ARGS##_TYPE_ARG(float)) { \
109 return detail::maybe_fmf_##NAME(NUM_ARGS##_CONVERTED_ARG); \
111 BUILTIN_GENF_DEVICE_COMMON(NUM_ARGS, NAME, detail::maybe_fmf_##NAME)
114 #define BUILTIN_GENF(NUM_ARGS, NAME) \
115 FOR_EACH2(HOST_IMPL_SCALAR, NUM_ARGS, NAME, FP_TYPES) \
116 HOST_IMPL_TEMPLATE(NUM_ARGS, NAME, builtin_enable_math_t, math, \
120 #define BUILTIN_GENF_NATIVE_OPT(NUM_ARGS, NAME) BUILTIN_GENF(NUM_ARGS, NAME)
166 #define BUILTIN_GENF_SCALAR_2ND(NAME) \
167 BUILTIN_GENF(TWO_ARGS, NAME) \
168 template <typename T> \
169 detail::builtin_enable_math_t<T> NAME(T x, detail::get_elem_type_t<T> y) { \
170 return NAME(detail::simplify_if_swizzle_t<T>{x}, \
171 detail::simplify_if_swizzle_t<T>{y}); \
174 BUILTIN_GENF_SCALAR_2ND(fmax)
175 BUILTIN_GENF_SCALAR_2ND(fmin)
177 #undef BUILTIN_GENF_SCALAR_2ND
179 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
cos)
180 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
exp)
181 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
exp10)
182 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
exp2)
183 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
log)
184 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
log10)
185 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
log2)
186 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
rsqrt)
187 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
sin)
188 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
sqrt)
189 BUILTIN_GENF_NATIVE_OPT(ONE_ARG,
tan)
190 BUILTIN_GENF_NATIVE_OPT(TWO_ARGS,
powr)
192 #undef BUILTIN_GENF_NATIVE_OPT
194 #undef BUILTIN_GENF_DEVICE_COMMON
197 template <
typename T0,
typename T1>
198 inline constexpr
bool builtin_ptr_check_v =
199 (std::is_pointer_v<T1> ||
200 (is_multi_ptr_v<T1> && has_writeable_addr_space_v<T1>)) &&
201 is_valid_elem_type_v<T1, simplify_if_swizzle_t<T0>>;
203 template <
typename T0,
typename T1>
204 inline constexpr
bool builtin_enable_ptr_v =
205 builtin_enable_math_allow_scalar_v<T0> && builtin_ptr_check_v<T0, T1>;
207 template <
typename T0,
typename T1>
208 using builtin_enable_ptr_scalar_t =
209 std::enable_if_t<builtin_enable_ptr_v<T0, T1> && is_scalar_arithmetic_v<T0>,
212 template <
typename T0,
typename T1>
213 using builtin_enable_ptr_non_scalar_t =
214 std::enable_if_t<builtin_enable_ptr_v<T0, T1> &&
215 !is_scalar_arithmetic_v<T0>,
216 simplify_if_swizzle_t<T0>>;
218 template <
typename FuncTy,
typename PtrTy,
typename... Ts>
219 auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) {
220 using T0 =
typename first_type<Ts...>::type;
222 simplify_if_swizzle_t<T0>
r{};
227 if constexpr (is_multi_ptr_v<PtrTy>)
229 get_multi_ptr_decoration_v<PtrTy>>(&(*p)[0]);
234 constexpr
auto N = T0::size();
235 if constexpr (N <= 16)
236 loop<N>([&](auto i) {
r[i] = F(xs[i]..., p0 + i); });
238 for (
size_t i = 0; i < N; ++i)
239 r[i] = F(xs[i]..., p0 + i);
244 #define LAST_PTR_SCALAR(NUM_ARGS, NAME, SCALAR_ENABLER, TYPE) \
245 template <typename PtrTy> \
246 detail::SCALAR_ENABLER<TYPE, PtrTy> NAME( \
247 SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPE_ARG)(TYPE), PtrTy p) { \
248 return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
250 #define BUILTIN_LAST_PTR_COMMON(NUM_ARGS, NAME, SCALAR_ENABLER, \
251 NON_SCALAR_ENABLER) \
252 FOR_EACH3(LAST_PTR_SCALAR, NUM_ARGS, NAME, SCALAR_ENABLER, FP_TYPES) \
253 template <SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPENAME_TYPE), typename PtrTy> \
254 detail::NON_SCALAR_ENABLER<SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), \
256 NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) { \
257 return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
260 #if __SYCL_DEVICE_ONLY__
261 #define BUILTIN_LAST_PTR(NUM_ARGS, NAME, SCALAR_ENABLER, NON_SCALAR_ENABLER) \
263 template <NUM_ARGS##_TYPENAME_TYPE> \
264 auto NAME##_impl(NUM_ARGS##_TEMPLATE_TYPE_ARG_REF) { \
265 if constexpr (is_marray_v<T0>) { \
266 return builtin_delegate_ptr_impl( \
267 [](auto... xs) { return NAME##_impl(xs...); }, \
268 NUM_ARGS##_ARG_ROTATED); \
270 return __spirv_ocl_##NAME(NUM_ARGS##_CONVERTED_ARG); \
274 BUILTIN_LAST_PTR_COMMON(NUM_ARGS, NAME, SCALAR_ENABLER, NON_SCALAR_ENABLER)
276 #define BUILTIN_LAST_PTR(NUM_ARGS, NAME, SCALAR_ENABLER, NON_SCALAR_ENABLER) \
277 BUILTIN_LAST_PTR_COMMON(NUM_ARGS, NAME, SCALAR_ENABLER, NON_SCALAR_ENABLER)
281 template <
typename T>
282 using builtin_last_raw_intptr_t =
285 change_elements_t<std::conditional_t<is_marray_v<T>, int, int32_t>,
286 simplify_if_swizzle_t<T>> *;
289 #ifdef __SYCL_DEVICE_ONLY__
290 #define BUILTIN_LAST_INTPTR(NUM_ARGS, NAME) \
291 BUILTIN_LAST_PTR(NUM_ARGS, NAME, builtin_enable_last_intptr_scalar_t, \
292 builtin_enable_last_intptr_non_scalar_t)
294 #define LAST_INT_PTR_DECLARE_SCALAR(NUM_ARGS, NAME, TYPE) \
295 __SYCL_EXPORT TYPE NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPE)(TYPE), \
297 #define BUILTIN_LAST_INTPTR(NUM_ARGS, NAME) \
299 FOR_EACH2(LAST_INT_PTR_DECLARE_SCALAR, NUM_ARGS, NAME, FP_TYPES) \
300 template <SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPENAME_TYPE), typename PtrTy> \
301 auto NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG_REF), \
303 if constexpr (is_multi_ptr_v<PtrTy>) { \
305 return NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p.get_raw()); \
307 return builtin_delegate_ptr_impl( \
308 [](auto... xs) { return NAME##_impl(xs...); }, p, \
309 SYCL_CONCAT(LESS_ONE(NUM_ARGS), _SIMPLIFIED_ARG)); \
313 BUILTIN_LAST_PTR(NUM_ARGS, NAME, builtin_enable_last_intptr_scalar_t, \
314 builtin_enable_last_intptr_non_scalar_t)
317 BUILTIN_LAST_INTPTR(TWO_ARGS, frexp)
318 BUILTIN_LAST_INTPTR(TWO_ARGS, lgamma_r)
319 BUILTIN_LAST_INTPTR(THREE_ARGS, remquo)
321 #undef BUILTIN_LAST_INTPTR
322 #undef LAST_INT_PTR_DECLARE_SCALAR
324 #ifndef __SYCL_DEVICE_ONLY__
326 template <
typename T0,
typename T1>
auto fract_impl(T0 &x, T1 &y) {
327 auto flr =
floor(simplify_if_swizzle_t<T0>{
x});
329 return fmin(x - flr,
nextafter(simplify_if_swizzle_t<T0>{1.0},
330 simplify_if_swizzle_t<T0>{0.0}));
334 BUILTIN_LAST_PTR(TWO_ARGS, fract, builtin_enable_ptr_scalar_t,
335 builtin_enable_ptr_non_scalar_t)
337 #ifndef __SYCL_DEVICE_ONLY__
339 __SYCL_EXPORT
float modf_impl(
float,
float *);
340 __SYCL_EXPORT
double modf_impl(
double,
double *);
342 template <
typename T0,
typename T1>
auto modf_impl(T0 &x, T1 &&y) {
346 return modf_impl(x,
y.get_raw());
348 return builtin_delegate_ptr_impl(
349 [](
auto x,
auto y) {
return modf_impl(x, y); },
y,
350 simplify_if_swizzle_t<T0>{
x});
355 BUILTIN_LAST_PTR(TWO_ARGS, modf, builtin_enable_ptr_scalar_t,
356 builtin_enable_ptr_non_scalar_t)
358 #undef BUILTIN_LAST_PTR
360 #ifdef __SYCL_DEVICE_ONLY__
361 #define BUILTIN_MATH_LAST_INT(NAME) \
362 float NAME(float x, int y) { return __spirv_ocl_##NAME(x, y); } \
363 double NAME(double x, int y) { return __spirv_ocl_##NAME(x, y); } \
364 half NAME(half x, int y) { \
365 return __spirv_ocl_##NAME(detail::builtins::convert_arg(x), y); \
367 DEVICE_IMPL_TEMPLATE(TWO_ARGS, NAME, builtin_enable_last_int_t, \
370 #define SCALAR_EXTERN_LAST_INT(NAME, TYPE) \
371 inline TYPE NAME(TYPE x, int y) { \
372 extern SYCL_BUILTIN_EXPORT TYPE __##NAME##_impl(TYPE, int); \
373 return __##NAME##_impl(x, y); \
375 #define BUILTIN_MATH_LAST_INT(NAME) \
376 FOR_EACH1(SCALAR_EXTERN_LAST_INT, NAME, FP_TYPES) \
377 HOST_IMPL_TEMPLATE(TWO_ARGS, NAME, builtin_enable_last_int_t, math, \
384 template <
typename T> detail::builtin_enable_math_t<T> ldexp(T
x,
int y) {
386 detail::simplify_if_swizzle_t<T>{
x},
387 detail::change_elements_t<int, detail::simplify_if_swizzle_t<T>>{
y});
390 #undef BUILTIN_MATH_LAST_INT
391 #undef SCALAR_EXTERN_LAST_INT
394 #ifdef __SYCL_DEVICE_ONLY__
395 template <
typename T0,
typename T1>
auto sincos_impl(T0 &x, T1 &&y) {
396 if constexpr (is_marray_v<T0>) {
397 return builtin_delegate_ptr_impl(
398 [](
auto... xs) {
return sincos_impl(xs...); },
y,
x);
400 using detail::builtins::convert_arg;
401 if constexpr (use_fast_math_v<T0>) {
409 *
y = __spirv_ocl_native_cos(convert_arg(x));
410 return __spirv_ocl_native_sin(convert_arg(x));
412 return __spirv_ocl_sincos(convert_arg(x), convert_arg(y));
417 __SYCL_EXPORT
float sincos_impl(
float,
float *);
418 __SYCL_EXPORT
double sincos_impl(
double,
double *);
420 template <
typename T0,
typename T1>
auto sincos_impl(T0 &x, T1 &&y) {
424 return sincos_impl(x,
y.get_raw());
426 return builtin_delegate_ptr_impl(
427 [](
auto... xs) {
return sincos_impl(xs...); },
y,
428 simplify_if_swizzle_t<T0>{
x});
433 BUILTIN_LAST_PTR_COMMON(TWO_ARGS, sincos, builtin_enable_ptr_scalar_t,
434 builtin_enable_ptr_non_scalar_t)
436 #undef BUILTIN_LAST_PTR_COMMON
437 #undef LAST_PTR_SCALAR
440 template <
typename T>
441 struct ilogb_ret_traits
442 : change_elements<std::conditional_t<is_vec_or_swizzle_v<T>, int32_t, int>,
444 template <
typename T>
445 using builtin_enable_ilogb_t =
446 std::enable_if_t<builtin_enable_math_allow_scalar_v<T>,
447 typename ilogb_ret_traits<T>::type>;
450 #ifdef __SYCL_DEVICE_ONLY__
451 inline int ilogb(
float x) {
return __spirv_ocl_ilogb(
x); }
452 inline int ilogb(
double x) {
return __spirv_ocl_ilogb(
x); }
453 inline int ilogb(
half x) {
454 return __spirv_ocl_ilogb(detail::builtins::convert_arg(
x));
458 inline int ilogb(
float x) {
460 return __ilogb_impl(
x);
462 inline int ilogb(
double x) {
464 return __ilogb_impl(
x);
466 inline int ilogb(
half x) {
468 return __ilogb_impl(
x);
477 template <
typename T>
480 typename map_type<get_elem_type_t<T>,
483 uint16_t,
half>::type,
487 template <
typename T>
488 using builtin_enable_nan_t = std::enable_if_t<
489 (((is_vec_or_swizzle_v<T> || is_marray_v<T>)) &&
490 check_type_in_v<get_elem_type_t<T>, uint32_t, uint64_t, uint16_t>),
491 nan_elem_result_type<T>>;
494 #ifdef __SYCL_DEVICE_ONLY__
495 inline float nan(uint32_t
x) {
496 return __spirv_ocl_nan(detail::builtins::convert_arg(
x));
498 inline double nan(uint64_t
x) {
499 return __spirv_ocl_nan(detail::builtins::convert_arg(
x));
501 inline half nan(uint16_t
x) {
502 return __spirv_ocl_nan(detail::builtins::convert_arg(
x));
506 inline float nan(uint32_t) {
return std::numeric_limits<float>::quiet_NaN(); }
507 inline double nan(uint64_t) {
return std::numeric_limits<float>::quiet_NaN(); }
509 inline half nan(uint16_t) {
return std::numeric_limits<half>::quiet_NaN(); }
510 template <
typename T> detail::builtin_enable_nan_t<T> nan(T
x) {
#define BUILTIN_CREATE_ENABLER(NAME, RET_TYPE_TRAIT, ELEM_TYPE_CHECKER, SHAPE_CHECKER, EXTRA_CONDITIONS)
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
__ESIMD_API simd< T, N > pow(simd< T, N > src0, simd< U, N > src1, Sat sat={})
Power - calculates src0 in power of src1.
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
__ESIMD_API simd< T, N > exp2(simd< T, N > src, Sat sat={})
Exponent base 2.
#define HOST_IMPL_TEMPLATE(NUM_ARGS, NAME, ENABLER, FUNC_CLASS, RET_TYPE_TRAITS)
#define SYCL_BUILTIN_EXPORT
#define DEVICE_IMPL_TEMPLATE(NUM_ARGS, NAME, ENABLER,...)
#define BUILTIN_GENF(NUM_ARGS, NAME)
#define BUILTIN_MATH_LAST_INT(NAME, IMPL)
constexpr bool check_type_in_v
constexpr bool builtin_same_shape_v
auto builtin_delegate_to_scalar(FuncTy F, const Ts &...x)
typename get_elem_type< T >::type get_elem_type_t
typename change_elements< NewElemT, T >::type change_elements_t
constexpr bool is_multi_ptr_v
constexpr bool builtin_same_or_swizzle_v
constexpr bool is_swizzle_v
std::enable_if_t< std::is_same_v< Tp, float >, float > copysign(Tp x, Tp y)
__SYCL_ALWAYS_INLINE std::enable_if_t< sycl::detail::is_svgenfloatf_v< T >||sycl::detail::is_svgenfloath_v< T >, T > tanh(T x) __NOEXC
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > atanh(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > tan(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > asin(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > cosh(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > asinh(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > acos(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > acosh(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > log10(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > atan(const complex< _Tp > &__x)
__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY std::enable_if_t< is_genfloat< _Tp >::value, complex< _Tp > > sinh(const complex< _Tp > &__x)
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fmin(T x, T y)
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fabs(T x)
return(x >=T(0)) ? T(std sinpi
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > log(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > sin(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > cos(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > exp(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
sycl::detail::half_impl::half half
std::enable_if_t< std::is_floating_point_v< ValueT >||std::is_same_v< sycl::half, ValueT >, ValueT > cbrt(ValueT val)
cbrt function wrapper.