21 #if defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
22 #define __esimd_dbg_print(a) std::cout << ">>> " << #a << "\n"
24 #define __esimd_dbg_print(a)
25 #endif // defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
31 namespace ext::intel::esimd {
34 template <
typename Ty,
int N>
class simd;
35 template <
typename BaseTy,
typename RegionTy>
class simd_view;
42 template <
typename Ty,
int N>
class simd_mask_impl;
43 template <
typename RawT,
int N,
class Derived,
class SFINAE =
void>
50 using simd_mask_elem_type =
unsigned short;
51 template <
int N>
using simd_mask_type = simd_mask_impl<simd_mask_elem_type, N>;
54 static inline constexpr
bool is_esimd_scalar_v =
55 sycl::detail::is_arithmetic<T>::value;
58 using is_esimd_scalar =
typename std::bool_constant<is_esimd_scalar_v<T>>;
63 struct is_simd_obj_impl_derivative :
public std::false_type {};
66 template <
typename RawT,
int N,
class Derived>
67 struct is_simd_obj_impl_derivative<simd_obj_impl<RawT, N, Derived>>
68 :
public std::true_type {};
71 template <
typename T,
int N,
template <
typename,
int>
class Derived>
72 struct is_simd_obj_impl_derivative<Derived<T, N>>
74 std::is_base_of_v<simd_obj_impl<__raw_t<T>, N, Derived<T, N>>,
76 std::true_type, std::false_type> {};
80 inline constexpr
bool is_simd_obj_impl_derivative_v =
81 is_simd_obj_impl_derivative<T>::value;
86 template <
class SimdT,
int Ndst>
struct resize_a_simd_type;
89 template <
typename T,
int Nsrc,
int Ndst,
template <
typename,
int>
class SimdT>
90 struct resize_a_simd_type<simd_obj_impl<__raw_t<T>, Nsrc, SimdT<T, Nsrc>>,
92 using type = simd_obj_impl<__raw_t<T>, Ndst, SimdT<T, Ndst>>;
96 template <
typename T,
int Nsrc,
int Ndst,
template <
typename,
int>
class SimdT>
97 struct resize_a_simd_type<SimdT<
T, Nsrc>, Ndst> {
98 using type = SimdT<T, Ndst>;
102 template <
class SimdT,
int Ndst>
103 using resize_a_simd_type_t =
typename resize_a_simd_type<SimdT, Ndst>::type;
109 template <
class SimdT,
typename NewElemT>
struct convert_simd_elem_type;
112 template <
typename OldElemT,
int N,
typename NewElemT,
113 template <
typename,
int>
class SimdT>
114 struct convert_simd_elem_type<
115 simd_obj_impl<__raw_t<OldElemT>, N, SimdT<OldElemT, N>>, NewElemT> {
116 using type = simd_obj_impl<__raw_t<NewElemT>, N, SimdT<NewElemT, N>>;
120 template <
typename OldElemT,
int N,
typename NewElemT,
121 template <
typename,
int>
class SimdT>
122 struct convert_simd_elem_type<SimdT<OldElemT, N>, NewElemT> {
123 using type = SimdT<NewElemT, N>;
127 template <
class SimdT,
typename NewElemT>
128 using convert_simd_elem_type_t =
129 typename convert_simd_elem_type<SimdT, NewElemT>::type;
135 template <
class SimdT,
typename T,
int N>
136 using construct_a_simd_type_t =
137 convert_simd_elem_type_t<resize_a_simd_type_t<SimdT, N>, T>;
143 using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;
147 template <
typename Ty>
struct is_any_simd_view_type : std::false_type {};
149 template <
typename BaseTy,
typename RegionTy>
150 struct is_any_simd_view_type<simd_view<BaseTy, RegionTy>> : std::true_type {};
152 template <
typename Ty>
153 static inline constexpr
bool is_any_simd_view_type_v =
154 is_any_simd_view_type<Ty>::value;
161 template <
typename Ty>
162 static inline constexpr
bool is_simd_like_type_v =
163 is_any_simd_view_type_v<Ty> || is_simd_obj_impl_derivative_v<Ty>;
169 template <
typename Ty>
struct is_simd_type : std::false_type {};
170 template <
typename ElTy,
int N>
171 struct is_simd_type<
simd<ElTy, N>> : std::true_type {};
172 template <
typename Ty>
173 static inline constexpr
bool is_simd_type_v = is_simd_type<Ty>::value;
175 template <
typename Ty>
struct is_simd_mask_type : std::false_type {};
177 struct is_simd_mask_type<simd_mask_impl<simd_mask_elem_type, N>>
179 template <
typename Ty>
180 static inline constexpr
bool is_simd_mask_type_v = is_simd_mask_type<Ty>::value;
185 template <
typename Ty>
struct is_simd_view_type_impl : std::false_type {};
187 template <
class BaseT,
class RegionT>
188 struct is_simd_view_type_impl<simd_view<BaseT, RegionT>>
193 struct is_simd_view_type : is_simd_view_type_impl<remove_cvref_t<Ty>> {};
195 template <
typename Ty>
196 static inline constexpr
bool is_simd_view_type_v = is_simd_view_type<Ty>::value;
199 template <
typename T>
200 static inline constexpr
bool is_simd_or_view_type_v =
201 is_simd_view_type_v<T> || is_simd_type_v<T>;
206 struct cant_deduce_element_type;
208 template <
class T,
class SFINAE =
void>
struct element_type {
209 using type = cant_deduce_element_type;
212 template <
typename T>
214 using type = remove_cvref_t<T>;
217 template <
typename T>
218 struct element_type<T,
std::
enable_if_t<is_simd_like_type_v<T>>> {
219 using type =
typename T::element_type;
222 template <
typename T>
223 struct element_type<T,
std::
enable_if_t<is_clang_vector_type_v<T>>> {
224 using type =
typename is_clang_vector_type<T>::element_type;
227 template <
typename T>
using element_type_t =
typename element_type<T>::type;
232 template <
class T>
struct simd_like_obj_info {
233 using element_type = T;
234 static inline constexpr
int vector_length = 0;
237 template <
class T,
int N>
struct simd_like_obj_info<
simd<T, N>> {
238 using element_type = T;
239 static inline constexpr
int vector_length = N;
242 template <
class T,
int N>
struct simd_like_obj_info<simd_mask_impl<T, N>> {
243 using element_type = simd_mask_elem_type;
244 static inline constexpr
int vector_length = N;
247 template <
class BaseT,
class RegionT>
248 struct simd_like_obj_info<simd_view<BaseT, RegionT>> {
249 using element_type =
typename RegionT::element_type;
250 static inline constexpr
int vector_length = RegionT::length;
253 template <
typename T>
254 using get_vector_element_type =
typename simd_like_obj_info<T>::element_type;
256 template <
typename T>
257 static inline constexpr
int get_vector_length =
258 simd_like_obj_info<T>::vector_length;
262 template <
typename To,
typename From>
263 std::enable_if_t<is_clang_vector_type_v<To> && is_clang_vector_type_v<From>, To>
264 ESIMD_INLINE
convert(From Val) {
265 if constexpr (std::is_same_v<To, From>) {
268 return __builtin_convertvector(Val, To);
273 template <
typename ToEltTy,
typename FromEltTy,
int FromN,
274 typename = std::enable_if_t<is_vectorizable<ToEltTy>::value>>
275 struct bitcast_helper {
276 static inline constexpr
int nToElems() {
277 constexpr
int R1 =
sizeof(ToEltTy) /
sizeof(FromEltTy);
278 constexpr
int R2 =
sizeof(FromEltTy) /
sizeof(ToEltTy);
279 constexpr
int ToN = (R2 > 0) ? (FromN * R2) : (FromN / R1);
285 template <
typename ToEltTy,
typename FromEltTy,
int FromN,
286 typename = std::enable_if_t<is_vectorizable<ToEltTy>::value>>
288 std::is_same<FromEltTy, ToEltTy>::value, vector_type_t<FromEltTy, FromN>,
289 vector_type_t<ToEltTy,
290 bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems()>>
291 bitcast(vector_type_t<FromEltTy, FromN> Val) {
293 if constexpr (std::is_same<FromEltTy, ToEltTy>::value)
297 constexpr
int ToN = bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems();
298 using VTy = vector_type_t<ToEltTy, ToN>;
299 return reinterpret_cast<VTy
>(Val);
311 struct invalid_computation_type;
313 template <
class T1,
class T2,
class SFINAE =
void>
struct computation_type {
314 using type = invalid_computation_type;
317 template <
class T1,
class T2>
318 struct computation_type<T1, T2,
320 is_valid_simd_elem_type_v<T2>>> {
322 template <
class T>
using tr = element_type_traits<T>;
325 std::conditional_t<tr<T>::use_native_cpp_ops,
typename tr<T>::RawT,
326 typename tr<T>::EnclosingCppT>;
327 static inline constexpr
bool is_wr1 = is_wrapper_elem_type_v<T1>;
328 static inline constexpr
bool is_wr2 = is_wrapper_elem_type_v<T2>;
329 static inline constexpr
bool is_fp1 = is_generic_floating_point_v<T1>;
330 static inline constexpr
bool is_fp2 = is_generic_floating_point_v<T2>;
336 decltype(std::declval<T1>() + std::declval<T2>()),
338 std::is_same_v<T1, T2>,
345 std::conditional_t<is_fp1, T1, T2>,
349 decltype(std::declval<native_t<T1>>() +
350 std::declval<native_t<T2>>())>>>;
353 template <
class T1,
class T2>
354 struct computation_type<
356 std::
enable_if_t<is_simd_like_type_v<T1> || is_simd_like_type_v<T2>>> {
358 using Ty1 = element_type_t<T1>;
359 using Ty2 = element_type_t<T2>;
360 using EltTy =
typename computation_type<Ty1, Ty2>::type;
362 static constexpr
int N1 = [] {
363 if constexpr (is_simd_like_type_v<T1>) {
369 static constexpr
int N2 = [] {
370 if constexpr (is_simd_like_type_v<T2>) {
376 static_assert((N1 == N2) || (N1 == 0) || (N2 == 0),
"size mismatch");
377 static constexpr
int N = N1 ? N1 : N2;
383 template <
class T1,
class T2 = T1>
384 using computation_type_t =
385 typename computation_type<remove_cvref_t<T1>, remove_cvref_t<T2>>::type;
392 template <
int N>
using mask_type_t = detail::simd_mask_storage_t<N>;