DPC++ Runtime
Runtime libraries for oneAPI DPC++
math_intrin.hpp
Go to the documentation of this file.
1 //==------------ math_intrin.hpp - DPC++ Explicit SIMD API -----------------==//
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 // Declares Explicit SIMD math intrinsics used to implement working with
9 // the SIMD classes objects.
10 //===----------------------------------------------------------------------===//
11 
12 #pragma once
13 
15 
16 #include <sycl/builtins.hpp>
17 #include <sycl/ext/intel/esimd/common.hpp>
18 #include <sycl/ext/intel/esimd/detail/elem_type_traits.hpp>
19 #include <sycl/ext/intel/esimd/detail/types.hpp>
20 #include <sycl/ext/intel/esimd/detail/util.hpp>
21 
22 #include <cstdint>
23 
24 #define __ESIMD_raw_vec_t(T, SZ) \
25  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, SZ>
26 #define __ESIMD_cpp_vec_t(T, SZ) \
27  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__cpp_t<T>, SZ>
28 
29 // saturation intrinsics
30 template <typename T0, typename T1, int SZ>
31 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
32  __esimd_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
33 
34 template <typename T0, typename T1, int SZ>
35 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
36  __esimd_fptoui_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
37 
38 template <typename T0, typename T1, int SZ>
39 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
40  __esimd_fptosi_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
41 
42 template <typename T0, typename T1, int SZ>
43 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
44  __esimd_uutrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
45 
46 template <typename T0, typename T1, int SZ>
47 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
48  __esimd_ustrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
49 
50 template <typename T0, typename T1, int SZ>
51 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
52  __esimd_sutrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
53 
54 template <typename T0, typename T1, int SZ>
55 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
56  __esimd_sstrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src) __ESIMD_INTRIN_END;
57 
58 template <typename T, int SZ>
59 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
60  __esimd_abs(__ESIMD_raw_vec_t(T, SZ) src0) __ESIMD_INTRIN_END;
61 
63 template <typename T, int SZ>
64 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
65  __esimd_fmax(__ESIMD_raw_vec_t(T, SZ) src0,
66  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
67 template <typename T, int SZ>
68 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
69  __esimd_umax(__ESIMD_raw_vec_t(T, SZ) src0,
70  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
71 template <typename T, int SZ>
72 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
73  __esimd_smax(__ESIMD_raw_vec_t(T, SZ) src0,
74  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
75 
77 template <typename T, int SZ>
78 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
79  __esimd_fmin(__ESIMD_raw_vec_t(T, SZ) src0,
80  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
81 template <typename T, int SZ>
82 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
83  __esimd_umin(__ESIMD_raw_vec_t(T, SZ) src0,
84  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
85 template <typename T, int SZ>
86 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
87  __esimd_smin(__ESIMD_raw_vec_t(T, SZ) src0,
88  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
89 
90 template <typename T, int SZ>
91 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<unsigned int, SZ>
92  __esimd_cbit(__ESIMD_raw_vec_t(T, SZ) src0) __ESIMD_INTRIN_END;
93 
94 template <typename T0, int SZ>
95 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
96  __esimd_fbl(__ESIMD_raw_vec_t(T0, SZ) src0) __ESIMD_INTRIN_END;
97 
98 template <typename T0, int SZ>
99 __ESIMD_INTRIN __ESIMD_raw_vec_t(int, SZ)
100  __esimd_sfbh(__ESIMD_raw_vec_t(T0, SZ) src0) __ESIMD_INTRIN_END;
101 
102 template <typename T0, int SZ>
103 __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, SZ)
104  __esimd_ufbh(__ESIMD_raw_vec_t(T0, SZ) src0) __ESIMD_INTRIN_END;
105 
106 #define __ESIMD_UNARY_EXT_MATH_INTRIN(name) \
107  template <class T, int SZ> \
108  __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ) \
109  __esimd_##name(__ESIMD_raw_vec_t(T, SZ) src) __ESIMD_INTRIN_END
110 
111 __ESIMD_UNARY_EXT_MATH_INTRIN(inv);
112 __ESIMD_UNARY_EXT_MATH_INTRIN(log);
113 __ESIMD_UNARY_EXT_MATH_INTRIN(exp);
114 __ESIMD_UNARY_EXT_MATH_INTRIN(sqrt);
115 __ESIMD_UNARY_EXT_MATH_INTRIN(ieee_sqrt);
116 __ESIMD_UNARY_EXT_MATH_INTRIN(rsqrt);
117 __ESIMD_UNARY_EXT_MATH_INTRIN(sin);
118 __ESIMD_UNARY_EXT_MATH_INTRIN(cos);
119 
120 #undef __ESIMD_UNARY_EXT_MATH_INTRIN
121 
122 template <class T, int SZ>
123 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
124  __esimd_pow(__ESIMD_raw_vec_t(T, SZ) src0,
125  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
126 
127 template <class T, int SZ>
128 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
129  __esimd_ieee_div(__ESIMD_raw_vec_t(T, SZ) src0,
130  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
131 
132 template <int SZ>
133 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
134 __esimd_rndd(__ESIMD_DNS::vector_type_t<float, SZ> src0) __ESIMD_INTRIN_END;
135 template <int SZ>
136 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
137 __esimd_rndu(__ESIMD_DNS::vector_type_t<float, SZ> src0) __ESIMD_INTRIN_END;
138 template <int SZ>
139 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
140 __esimd_rnde(__ESIMD_DNS::vector_type_t<float, SZ> src0) __ESIMD_INTRIN_END;
141 template <int SZ>
142 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
143 __esimd_rndz(__ESIMD_DNS::vector_type_t<float, SZ> src0) __ESIMD_INTRIN_END;
144 
145 template <int N>
146 __ESIMD_INTRIN uint32_t __esimd_pack_mask(
147  __ESIMD_DNS::vector_type_t<uint16_t, N> src0) __ESIMD_INTRIN_END;
148 
149 template <int N>
150 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<uint16_t, N>
151 __esimd_unpack_mask(uint32_t src0) __ESIMD_INTRIN_END;
152 
153 template <typename T1, typename T2, typename T3, typename T4, int N>
154 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
155  __esimd_uudp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
156  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
157 
158 template <typename T1, typename T2, typename T3, typename T4, int N>
159 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
160  __esimd_usdp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
161  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
162 
163 template <typename T1, typename T2, typename T3, typename T4, int N>
164 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
165  __esimd_sudp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
166  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
167 
168 template <typename T1, typename T2, typename T3, typename T4, int N>
169 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
170  __esimd_ssdp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
171  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
172 
173 template <typename T1, typename T2, typename T3, typename T4, int N>
174 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
175  __esimd_uudp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
176  __ESIMD_raw_vec_t(T3, N) src1,
177  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
178 
179 template <typename T1, typename T2, typename T3, typename T4, int N>
180 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
181  __esimd_usdp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
182  __ESIMD_raw_vec_t(T3, N) src1,
183  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
184 
185 template <typename T1, typename T2, typename T3, typename T4, int N>
186 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
187  __esimd_sudp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
188  __ESIMD_raw_vec_t(T3, N) src1,
189  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
190 
191 template <typename T1, typename T2, typename T3, typename T4, int N>
192 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
193  __esimd_ssdp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
194  __ESIMD_raw_vec_t(T3, N) src1,
195  __ESIMD_raw_vec_t(T4, N) src2) __ESIMD_INTRIN_END;
196 __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, 4)
197  __esimd_timestamp() __ESIMD_INTRIN_END;
198 
199 #ifdef __SYCL_DEVICE_ONLY__
200 
201 // lane-id for reusing scalar math functions.
202 // Depending upon the SIMT mode(8/16/32), the return value is
203 // in the range of 0-7, 0-15, or 0-31.
204 __ESIMD_INTRIN int __esimd_lane_id();
205 
206 // Wrapper for designating a scalar region of code that will be
207 // vectorized by the backend compiler.
208 #define __ESIMD_SIMT_BEGIN(N, lane) \
209  [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE [[intel::sycl_esimd_vectorize(N)]] { \
210  int lane = __esimd_lane_id();
211 #define __ESIMD_SIMT_END \
212  } \
213  ();
214 
215 #define ESIMD_MATH_INTRINSIC_IMPL(type, func) \
216  template <int SZ> \
217  __ESIMD_INTRIN __ESIMD_raw_vec_t(type, SZ) \
218  ocl_##func(__ESIMD_raw_vec_t(type, SZ) src0) { \
219  __ESIMD_raw_vec_t(type, SZ) retv; \
220  __ESIMD_SIMT_BEGIN(SZ, lane) \
221  retv[lane] = sycl::func(src0[lane]); \
222  __ESIMD_SIMT_END \
223  return retv; \
224  }
225 
226 namespace sycl {
227 inline namespace _V1 {
228 namespace ext::intel::esimd::detail {
229 // TODO support half vectors in std sycl math functions.
230 ESIMD_MATH_INTRINSIC_IMPL(float, sin)
231 ESIMD_MATH_INTRINSIC_IMPL(float, cos)
232 ESIMD_MATH_INTRINSIC_IMPL(float, exp)
233 ESIMD_MATH_INTRINSIC_IMPL(float, log)
234 } // namespace ext::intel::esimd::detail
235 } // namespace _V1
236 } // namespace sycl
237 
238 #undef __ESIMD_SIMT_BEGIN
239 #undef __ESIMD_SIMT_END
240 #undef ESIMD_MATH_INTRINSIC_IMPL
241 
242 #endif // #ifdef __SYCL_DEVICE_ONLY__
243 
244 #undef __ESIMD_raw_vec_t
245 #undef __ESIMD_cpp_vec_t
246 
elem_type_traits.hpp
sycl::_V1::ext::intel::esimd::cos
__ESIMD_API simd< T, N > cos(simd< T, N > src, Sat sat={})
Cosine.
Definition: math.hpp:404
sycl::_V1::ext::intel::esimd::sqrt
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
Definition: math.hpp:388
sycl::_V1::ext::intel::esimd::rsqrt
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
Definition: math.hpp:396
sycl::_V1::ext::intel::esimd::inv
__ESIMD_API simd< T, N > inv(simd< T, N > src, Sat sat={})
Inversion - calculates (1/x).
Definition: math.hpp:374
sycl::_V1::ext::intel::esimd::sin
__ESIMD_API simd< T, N > sin(simd< T, N > src, Sat sat={})
Sine.
Definition: math.hpp:400
sycl::_V1::ext::intel::esimd::src1
__ESIMD_API SZ simd< T, SZ > src1
Definition: math.hpp:179
sycl::_V1::ext::intel::esimd::log
ESIMD_DETAIL ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > log(simd< T, SZ > src0, Sat sat={})
Computes the natural logarithm of the given argument.
Definition: math.hpp:469
sycl::_V1::ext::intel::esimd::SZ
__ESIMD_API SZ simd< T, SZ > Sat int SZ
Definition: math.hpp:210
sycl::_V1::ext::intel::esimd::exp
ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > exp(simd< T, SZ > src0, Sat sat={})
Computes e raised to the power of the given argument.
Definition: math.hpp:490
sycl::_V1::ext::intel::esimd::src0
__ESIMD_API SZ src0
Definition: math.hpp:179
sycl::_V1::log
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > log(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:59
sycl::_V1::sin
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > sin(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:37
sycl::_V1::cos
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > cos(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:26
sycl::_V1::exp
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > exp(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
Definition: builtins_esimd.hpp:48
sycl
Definition: access.hpp:18