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 experimental Explicit SIMD math intrinsics.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
21 
22 #include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
23 #include <sycl/ext/intel/esimd/detail/math_intrin.hpp>
24 #include <sycl/ext/intel/esimd/detail/types.hpp>
25 #include <sycl/ext/intel/esimd/xmx/common.hpp>
26 
27 #define __ESIMD_raw_vec_t(T, SZ) \
28  sycl::ext::intel::esimd::detail::vector_type_t< \
29  sycl::ext::intel::esimd::detail::__raw_t<T>, SZ>
30 #define __ESIMD_cpp_vec_t(T, SZ) \
31  sycl::ext::intel::esimd::detail::vector_type_t< \
32  sycl::ext::intel::esimd::detail::__cpp_t<T>, SZ>
33 
34 template <typename T, int SZ>
35 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
36  __esimd_umulh(__ESIMD_raw_vec_t(T, SZ) src0,
37  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
38 template <typename T, int SZ>
39 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
40  __esimd_smulh(__ESIMD_raw_vec_t(T, SZ) src0,
41  __ESIMD_raw_vec_t(T, SZ) src1) __ESIMD_INTRIN_END;
42 
43 template <int SZ>
44 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
45 __esimd_frc(__ESIMD_DNS::vector_type_t<float, SZ> src0) __ESIMD_INTRIN_END;
46 
47 template <typename T, int SZ>
48 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
49  __esimd_lzd(__ESIMD_raw_vec_t(T, SZ) src0) __ESIMD_INTRIN_END;
50 
51 template <typename T0, typename T1, int SZ>
52 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
53  __esimd_bfrev(__ESIMD_raw_vec_t(T1, SZ) src0) __ESIMD_INTRIN_END;
54 
55 template <typename T0, int SZ>
56 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
57  __esimd_bfi(__ESIMD_raw_vec_t(T0, SZ) src0, __ESIMD_raw_vec_t(T0, SZ) src1,
58  __ESIMD_raw_vec_t(T0, SZ) src2,
59  __ESIMD_raw_vec_t(T0, SZ) src3) __ESIMD_INTRIN_END;
60 
61 template <typename T0, int SZ>
62 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
63  __esimd_sbfe(__ESIMD_raw_vec_t(T0, SZ) src0, __ESIMD_raw_vec_t(T0, SZ) src1,
64  __ESIMD_raw_vec_t(T0, SZ) src2) __ESIMD_INTRIN_END;
65 
66 template <typename T, int N>
67 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, N)
68  __esimd_dp4(__ESIMD_raw_vec_t(T, N) v1,
69  __ESIMD_raw_vec_t(T, N) v2) __ESIMD_INTRIN_END;
70 
71 template <__ESIMD_XMX_NS::dpas_argument_type src1_precision,
72  __ESIMD_XMX_NS::dpas_argument_type src2_precision, int systolic_depth,
73  int repeat_count, typename T, typename T0, typename T1, typename T2,
74  int N, int N1, int N2, int res_sign = std::is_signed_v<T>,
75  int acc_sign = std::is_signed_v<T0>>
76 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
77 __esimd_dpas2(__ESIMD_DNS::vector_type_t<T0, N> src0,
78  __ESIMD_DNS::vector_type_t<T1, N1> src1,
79  __ESIMD_DNS::vector_type_t<T2, N2> src2) __ESIMD_INTRIN_END;
80 
81 template <int Info, typename T, typename T1, typename T2, int N, int N1, int N2>
82 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
83 __esimd_dpas_nosrc0(__ESIMD_DNS::vector_type_t<T1, N1> src1,
84  __ESIMD_DNS::vector_type_t<T2, N2> src2) __ESIMD_INTRIN_END;
85 
86 template <int Info, typename T, typename T1, typename T2, int N, int N1, int N2>
87 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N>
88 __esimd_dpasw(__ESIMD_DNS::vector_type_t<T, N> src0,
89  __ESIMD_DNS::vector_type_t<T1, N1> src1,
90  __ESIMD_DNS::vector_type_t<T2, N2> src2) __ESIMD_INTRIN_END;
91 
92 template <int Info, typename T, typename T1, typename T2, int N, int N1, int N2>
93 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<T, N> __esimd_dpasw_nosrc0(
94  __ESIMD_DNS::vector_type_t<T1, N1> src1,
95  __ESIMD_DNS::vector_type_t<T2, N2> src2) __ESIMD_INTRIN_END;
96 
97 template <typename T, int N>
98 __ESIMD_INTRIN std::pair<__ESIMD_DNS::vector_type_t<T, N>,
99  __ESIMD_DNS::vector_type_t<T, N>>
100 __esimd_addc(__ESIMD_DNS::vector_type_t<T, N> src0,
101  __ESIMD_DNS::vector_type_t<T, N> src1) __ESIMD_INTRIN_END;
102 
103 template <typename T, int N>
104 __ESIMD_INTRIN std::pair<__ESIMD_DNS::vector_type_t<T, N>,
105  __ESIMD_DNS::vector_type_t<T, N>>
106 __esimd_subb(__ESIMD_DNS::vector_type_t<T, N> src0,
107  __ESIMD_DNS::vector_type_t<T, N> src1) __ESIMD_INTRIN_END;
108 
109 template <uint8_t FuncControl, typename T, int N>
110 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, N)
111  __esimd_bfn(__ESIMD_raw_vec_t(T, N) src0, __ESIMD_raw_vec_t(T, N) src1,
112  __ESIMD_raw_vec_t(T, N) src2) __ESIMD_INTRIN_END;
113 
114 template <int N>
115 __ESIMD_INTRIN __ESIMD_raw_vec_t(sycl::half, N)
116  __esimd_srnd(__ESIMD_DNS::vector_type_t<float, N> src1,
117  __ESIMD_DNS::vector_type_t<uint16_t, N> src2)
118  __ESIMD_INTRIN_END;
119 
120 template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_fma(T, T, T);
121 template <typename T, int N>
122 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
123  __spirv_ocl_fma(__ESIMD_raw_vec_t(T, N) a, __ESIMD_raw_vec_t(T, N) b,
124  __ESIMD_raw_vec_t(T, N) c) __ESIMD_INTRIN_END;
125 
126 template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_popcount(T);
127 template <typename T, int N>
128 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
129  __spirv_ocl_popcount(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
130 
131 template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_ctz(T);
132 template <typename T, int N>
133 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
134  __spirv_ocl_ctz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
135 
136 template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_ocl_clz(T);
137 template <typename T, int N>
138 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
139  __spirv_ocl_clz(__ESIMD_raw_vec_t(T, N) src0) __ESIMD_INTRIN_END;
140 
141 template <typename T> extern __DPCPP_SYCL_EXTERNAL T __spirv_FRem(T);
142 template <typename T, int N>
143 extern __DPCPP_SYCL_EXTERNAL __ESIMD_raw_vec_t(T, N)
144  __spirv_FRem(__ESIMD_raw_vec_t(T, N) src0,
145  __ESIMD_raw_vec_t(T, N) src1) __ESIMD_INTRIN_END;
146 
147 #undef __ESIMD_raw_vec_t
148 #undef __ESIMD_cpp_vec_t
149 
sycl::_V1::detail::half_impl::half
Definition: half_type.hpp:186
__DPCPP_SYCL_EXTERNAL
#define __DPCPP_SYCL_EXTERNAL
Definition: defines_elementary.hpp:33
math_intrin.hpp
defines_elementary.hpp
sycl::_V1::ext::intel::esimd::src1
__ESIMD_API SZ simd< T, SZ > src1
Definition: math.hpp:187
sycl::_V1::ext::intel::esimd::SZ
__ESIMD_API SZ simd< T, SZ > Sat int SZ
Definition: math.hpp:222
sycl::_V1::ext::intel::esimd::src0
__ESIMD_API SZ src0
Definition: math.hpp:187
sycl::_V1::ext::intel::esimd::xmx::dpas_argument_type
dpas_argument_type
Describes the element types in the input matrices.
Definition: common.hpp:23
sycl::_V1::a
auto a
Definition: integer_functions.cpp:179
sycl::_V1::b
auto autodecltype(a) b
Definition: integer_functions.cpp:179