DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins.hpp
Go to the documentation of this file.
1 //==----------- builtins.hpp - SYCL built-in functions ---------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <CL/__spirv/spirv_ops.hpp>
12 #include <sycl/detail/common.hpp>
14 
15 #include <type_traits>
16 
17 // TODO Decide whether to mark functions with this attribute.
18 #define __NOEXC /*noexcept*/
19 
20 #ifdef __SYCL_DEVICE_ONLY__
21 #define __FUNC_PREFIX_OCL __spirv_ocl_
22 #define __FUNC_PREFIX_CORE __spirv_
23 #define __SYCL_EXTERN_IT1(Ret, prefix, call, Arg1)
24 #define __SYCL_EXTERN_IT2(Ret, prefix, call, Arg1, Arg2)
25 #define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, Arg)
26 #define __SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3)
27 #else
28 #define __FUNC_PREFIX_OCL sycl_host_
29 #define __FUNC_PREFIX_CORE sycl_host_
30 #define __SYCL_EXTERN_IT1(Ret, prefix, call, Arg) \
31  extern Ret __SYCL_PPCAT(prefix, call)(Arg)
32 #define __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, Arg) \
33  extern Ret __SYCL_PPCAT(prefix, call)(Arg, Arg)
34 #define __SYCL_EXTERN_IT2(Ret, prefix, call, Arg1, Arg2) \
35  extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2)
36 #define __SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3) \
37  extern Ret __SYCL_PPCAT(prefix, call)(Arg1, Arg2, Arg3)
38 #endif
39 
40 #define __SYCL_PPCAT_NX(A, B) A##B
41 #define __SYCL_PPCAT(A, B) __SYCL_PPCAT_NX(A, B)
42 
43 #define __SYCL_MAKE_CALL_ARG1(call, prefix) \
44  template <typename R, typename T1> \
45  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1) __NOEXC { \
46  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
47  using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
48  __SYCL_EXTERN_IT1(Ret, prefix, call, Arg1); \
49  Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(t1); \
50  Ret ret = __SYCL_PPCAT(prefix, call)(arg1); \
51  return sycl::detail::convertDataToType<Ret, R>(ret); \
52  }
53 
54 #define __SYCL_MAKE_CALL_ARG2(call, prefix) \
55  template <typename R, typename T1, typename T2> \
56  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1, T2 t2) __NOEXC { \
57  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
58  using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
59  using Arg2 = sycl::detail::ConvertToOpenCLType_t<T2>; \
60  __SYCL_EXTERN_IT2(Ret, prefix, call, Arg1, Arg2); \
61  Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(t1); \
62  Arg2 arg2 = sycl::detail::convertDataToType<T2, Arg2>(t2); \
63  Ret ret = __SYCL_PPCAT(prefix, call)(arg1, arg2); \
64  return sycl::detail::convertDataToType<Ret, R>(ret); \
65  }
66 
67 #define __SYCL_MAKE_CALL_ARG2_SAME(call, prefix) \
68  template <typename R, typename T> \
69  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T t1, T t2) __NOEXC { \
70  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
71  using Arg = sycl::detail::ConvertToOpenCLType_t<T>; \
72  __SYCL_EXTERN_IT2_SAME(Ret, prefix, call, Arg); \
73  Arg arg1 = sycl::detail::convertDataToType<T, Arg>(t1); \
74  Arg arg2 = sycl::detail::convertDataToType<T, Arg>(t2); \
75  Ret ret = __SYCL_PPCAT(prefix, call)(arg1, arg2); \
76  return sycl::detail::convertDataToType<Ret, R>(ret); \
77  }
78 
79 #define __SYCL_MAKE_CALL_ARG2_SAME_RESULT(call, prefix) \
80  template <typename T> \
81  inline __SYCL_ALWAYS_INLINE T __invoke_##call(T v1, T v2) __NOEXC { \
82  using Type = sycl::detail::ConvertToOpenCLType_t<T>; \
83  __SYCL_EXTERN_IT2_SAME(Type, prefix, call, Type); \
84  Type arg1 = sycl::detail::convertDataToType<T, Type>(v1); \
85  Type arg2 = sycl::detail::convertDataToType<T, Type>(v2); \
86  Type ret = __SYCL_PPCAT(prefix, call)(arg1, arg2); \
87  return sycl::detail::convertDataToType<Type, T>(ret); \
88  }
89 
90 #define __SYCL_MAKE_CALL_ARG3(call, prefix) \
91  template <typename R, typename T1, typename T2, typename T3> \
92  inline __SYCL_ALWAYS_INLINE R __invoke_##call(T1 t1, T2 t2, T3 t3) __NOEXC { \
93  using Ret = sycl::detail::ConvertToOpenCLType_t<R>; \
94  using Arg1 = sycl::detail::ConvertToOpenCLType_t<T1>; \
95  using Arg2 = sycl::detail::ConvertToOpenCLType_t<T2>; \
96  using Arg3 = sycl::detail::ConvertToOpenCLType_t<T3>; \
97  __SYCL_EXTERN_IT3(Ret, prefix, call, Arg1, Arg2, Arg3); \
98  Arg1 arg1 = sycl::detail::convertDataToType<T1, Arg1>(t1); \
99  Arg2 arg2 = sycl::detail::convertDataToType<T2, Arg2>(t2); \
100  Arg3 arg3 = sycl::detail::convertDataToType<T3, Arg3>(t3); \
101  Ret ret = __SYCL_PPCAT(prefix, call)(arg1, arg2, arg3); \
102  return sycl::detail::convertDataToType<Ret, R>(ret); \
103  }
104 
105 #ifndef __SYCL_DEVICE_ONLY__
106 namespace __host_std {
107 #endif // __SYCL_DEVICE_ONLY__
108 /* ----------------- 4.13.3 Math functions. ---------------------------------*/
204 /* --------------- 4.13.4 Integer functions. --------------------------------*/
239 /* --------------- 4.13.5 Common functions. ---------------------------------*/
249 /* --------------- 4.13.6 Geometric Functions. ------------------------------*/
259 /* --------------- 4.13.7 Relational functions. -----------------------------*/
261 __SYCL_MAKE_CALL_ARG2_SAME(FUnordNotEqual, __FUNC_PREFIX_CORE) // isnotequal
262 __SYCL_MAKE_CALL_ARG2_SAME(FOrdGreaterThan, __FUNC_PREFIX_CORE) // isgreater
263 __SYCL_MAKE_CALL_ARG2_SAME(FOrdGreaterThanEqual,
264  __FUNC_PREFIX_CORE) // isgreaterequal
266 __SYCL_MAKE_CALL_ARG2_SAME(FOrdLessThanEqual, __FUNC_PREFIX_CORE) // islessequal
267 __SYCL_MAKE_CALL_ARG2_SAME(FOrdNotEqual, __FUNC_PREFIX_CORE) // islessgreater
268 __SYCL_MAKE_CALL_ARG2_SAME(LessOrGreater, __FUNC_PREFIX_CORE) // islessgreater
272 __SYCL_MAKE_CALL_ARG1(IsNormal, __FUNC_PREFIX_CORE) // isnormal
274 __SYCL_MAKE_CALL_ARG2_SAME(Unordered, __FUNC_PREFIX_CORE) // isunordered
279 __SYCL_MAKE_CALL_ARG3(select, __FUNC_PREFIX_OCL) // select
280 #ifndef __SYCL_DEVICE_ONLY__
281 } // namespace __host_std
282 #endif
283 
284 #undef __NOEXC
285 #undef __SYCL_MAKE_CALL_ARG1
286 #undef __SYCL_MAKE_CALL_ARG2
287 #undef __SYCL_MAKE_CALL_ARG2_SAME
288 #undef __SYCL_MAKE_CALL_ARG2_SAME_RESULT
289 #undef __SYCL_MAKE_CALL_ARG3
290 #undef __SYCL_PPCAT_NX
291 #undef __SYCL_PPCAT
292 #undef __FUNC_PREFIX_OCL
293 #undef __FUNC_PREFIX_CORE
294 #undef __SYCL_EXTERN_IT1
295 #undef __SYCL_EXTERN_IT2
296 #undef __SYCL_EXTERN_IT2_SAME
297 #undef __SYCL_EXTERN_IT3
#define __FUNC_PREFIX_OCL
Definition: builtins.hpp:28
#define __SYCL_MAKE_CALL_ARG2(call, prefix)
Definition: builtins.hpp:54
#define __FUNC_PREFIX_CORE
Definition: builtins.hpp:29
#define __SYCL_MAKE_CALL_ARG3(call, prefix)
Definition: builtins.hpp:90
#define __SYCL_MAKE_CALL_ARG2_SAME_RESULT(call, prefix)
Definition: builtins.hpp:79
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ > trunc(const sycl::ext::intel::esimd::simd< float, SZ > &src0, Sat sat={})
Round to integral value using the round to zero rounding mode (vector version).
Definition: math.hpp:614
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > floor(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Floor" operation, vector version - alias of rndd.
Definition: math.hpp:581
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > ceil(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Ceiling" operation, vector version - alias of rndu.
Definition: math.hpp:594
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
Definition: math.hpp:389
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
Definition: math.hpp:397
__ESIMD_API simd< T, N > pow(simd< T, N > src0, simd< U, N > src1, Sat sat={})
Power - calculates src0 in power of src1.
Definition: math.hpp:442
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
Definition: math.hpp:381
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > asin(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1243
__ESIMD_API sycl::ext::intel::esimd::simd< float, SZ > sincos(sycl::ext::intel::esimd::simd< float, SZ > &dstcos, U src0, Sat sat={})
Definition: math.hpp:1146
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > acos(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1205
ESIMD_DETAIL ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< T, SZ > atan(sycl::ext::intel::esimd::simd< T, SZ > src0)
Definition: math.hpp:1161
sycl::ext::intel::esimd::simd< float, N > fmod(sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1386
sycl::ext::intel::esimd::simd< float, N > atan2(sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1353
__FUNC_PREFIX_CORE __FUNC_PREFIX_CORE __FUNC_PREFIX_CORE __FUNC_PREFIX_CORE __FUNC_PREFIX_CORE __SYCL_MAKE_CALL_ARG1(IsFinite, __FUNC_PREFIX_CORE) __SYCL_MAKE_CALL_ARG1(IsInf
__SYCL_MAKE_CALL_ARG2_SAME(Dot, __FUNC_PREFIX_CORE) __SYCL_MAKE_CALL_ARG2_SAME_RESULT(FMul
__SYCL_ALWAYS_INLINE std::enable_if_t< std::is_same_v< T, half >||std::is_same_v< T, float >, sycl::marray< T, N > > tanh(sycl::marray< T, N > x) __NOEXC
Definition: builtins.hpp:116
__SYCL_ALWAYS_INLINE sycl::marray< half, N > exp2(sycl::marray< half, N > x) __NOEXC
Definition: builtins.hpp:158
std::enable_if_t< detail::is_bf16_storage_type< T >::value, T > fabs(T x)
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_ALWAYS_INLINE std::enable_if_t< __FAST_MATH_SGENFLOAT(T), marray< T, N > > powr(marray< T, N > x, marray< T, N > y) __NOEXC
Definition: builtins.hpp:168
@ All
int popcount(const simd_mask< _Tp, _Abi > &) noexcept