DPC++ Runtime
Runtime libraries for oneAPI DPC++
types.hpp
Go to the documentation of this file.
1 //==-------------- types.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 // Types and type traits to implement Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
13 #include <sycl/detail/defines.hpp>
14 #include <sycl/detail/stl_type_traits.hpp> // to define C++14,17 extensions
19 #include <sycl/half_type.hpp>
20 
21 #if defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
22 #define __esimd_dbg_print(a) std::cout << ">>> " << #a << "\n"
23 #else
24 #define __esimd_dbg_print(a)
25 #endif // defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
26 
27 #include <cstdint>
28 
29 namespace sycl {
31 namespace ext::intel::esimd {
32 
33 // simd and simd_view_impl forward declarations
34 template <typename Ty, int N> class simd;
35 template <typename BaseTy, typename RegionTy> class simd_view;
36 
38 
39 namespace detail {
40 
41 // forward declarations of major internal simd classes
42 template <typename Ty, int N> class simd_mask_impl;
43 template <typename RawT, int N, class Derived, class SFINAE = void>
44 class simd_obj_impl;
45 
46 // @{
47 // Helpers for major simd classes, which don't require their definitions to
48 // compile. Error checking/SFINAE is not used as these are only used internally.
49 
50 using simd_mask_elem_type = unsigned short;
51 template <int N> using simd_mask_type = simd_mask_impl<simd_mask_elem_type, N>;
52 
53 template <typename T>
54 static inline constexpr bool is_esimd_scalar_v =
55  sycl::detail::is_arithmetic<T>::value;
56 
57 template <typename T>
58 using is_esimd_scalar = typename std::bool_constant<is_esimd_scalar_v<T>>;
59 
60 // @{
61 // Checks if given type T derives from simd_obj_impl or is equal to it.
62 template <typename T>
63 struct is_simd_obj_impl_derivative : public std::false_type {};
64 
65 // Specialization for the simd_obj_impl type itself.
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 {};
69 
70 // Specialization for all other types.
71 template <typename T, int N, template <typename, int> class Derived>
72 struct is_simd_obj_impl_derivative<Derived<T, N>>
73  : public std::conditional_t<
74  std::is_base_of_v<simd_obj_impl<__raw_t<T>, N, Derived<T, N>>,
75  Derived<T, N>>,
76  std::true_type, std::false_type> {};
77 
78 // Convenience shortcut.
79 template <typename T>
80 inline constexpr bool is_simd_obj_impl_derivative_v =
81  is_simd_obj_impl_derivative<T>::value;
82 // @}
83 
84 // @{
85 // "Resizes" given simd type \c T to given number of elements \c N.
86 template <class SimdT, int Ndst> struct resize_a_simd_type;
87 
88 // Specialization for the simd_obj_impl 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>>,
91  Ndst> {
92  using type = simd_obj_impl<__raw_t<T>, Ndst, SimdT<T, Ndst>>;
93 };
94 
95 // Specialization for the simd_obj_impl type derivatives.
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>;
99 };
100 
101 // Convenience shortcut.
102 template <class SimdT, int Ndst>
103 using resize_a_simd_type_t = typename resize_a_simd_type<SimdT, Ndst>::type;
104 // @}
105 
106 // @{
107 // Converts element type of given simd type \c SimdT to
108 // given scalar type \c NewElemT.
109 template <class SimdT, typename NewElemT> struct convert_simd_elem_type;
110 
111 // Specialization for the simd_obj_impl 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>>;
117 };
118 
119 // Specialization for the simd_obj_impl type derivatives.
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>;
124 };
125 
126 // Convenience shortcut.
127 template <class SimdT, typename NewElemT>
128 using convert_simd_elem_type_t =
129  typename convert_simd_elem_type<SimdT, NewElemT>::type;
130 
131 // @}
132 
133 // Constructs a simd type with the same template type as in \c SimdT, and
134 // given element type and number.
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>;
138 
139 // @}
140 
141 // must match simd_mask<N>::element_type
142 template <int N>
143 using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;
144 
145 // @{
146 // Checks if given type is a view of any simd type (simd or simd_mask).
147 template <typename Ty> struct is_any_simd_view_type : std::false_type {};
148 
149 template <typename BaseTy, typename RegionTy>
150 struct is_any_simd_view_type<simd_view<BaseTy, RegionTy>> : std::true_type {};
151 
152 template <typename Ty>
153 static inline constexpr bool is_any_simd_view_type_v =
154  is_any_simd_view_type<Ty>::value;
155 // @}
156 
157 // @{
158 // Check if a type is one of internal 'simd_xxx_impl' types exposing simd-like
159 // interfaces and behaving like a simd object type.
160 
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>;
164 // @}
165 
166 // @{
167 // Checks if given type is a any of the user-visible simd types (simd or
168 // simd_mask).
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;
174 
175 template <typename Ty> struct is_simd_mask_type : std::false_type {};
176 template <int N>
177 struct is_simd_mask_type<simd_mask_impl<simd_mask_elem_type, N>>
178  : std::true_type {};
179 template <typename Ty>
180 static inline constexpr bool is_simd_mask_type_v = is_simd_mask_type<Ty>::value;
181 // @}
182 
183 // @{
184 // Checks if given type is a view of the simd type.
185 template <typename Ty> struct is_simd_view_type_impl : std::false_type {};
186 
187 template <class BaseT, class RegionT>
188 struct is_simd_view_type_impl<simd_view<BaseT, RegionT>>
189  : std::conditional_t<is_simd_type_v<BaseT>, std::true_type,
190  std::false_type> {};
191 
192 template <class Ty>
193 struct is_simd_view_type : is_simd_view_type_impl<remove_cvref_t<Ty>> {};
194 
195 template <typename Ty>
196 static inline constexpr bool is_simd_view_type_v = is_simd_view_type<Ty>::value;
197 // @}
198 
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>;
202 
203 // @{
204 // Get the element type if it is a scalar, clang vector, simd or simd_view type.
205 
206 struct cant_deduce_element_type;
207 
208 template <class T, class SFINAE = void> struct element_type {
209  using type = cant_deduce_element_type;
210 };
211 
212 template <typename T>
213 struct element_type<T, std::enable_if_t<is_vectorizable_v<T>>> {
214  using type = remove_cvref_t<T>;
215 };
216 
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;
220 };
221 
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;
225 };
226 
227 template <typename T> using element_type_t = typename element_type<T>::type;
228 
229 // Determine element type of simd_obj_impl's Derived type w/o having to have
230 // complete instantiation of the Derived type (is required by element_type_t,
231 // hence can't be used here).
232 template <class T> struct simd_like_obj_info {
233  using element_type = T;
234  static inline constexpr int vector_length = 0;
235 };
236 
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;
240 };
241 
242 template <class T, int N> struct simd_like_obj_info<simd_mask_impl<T, N>> {
243  using element_type = simd_mask_elem_type; // equals T
244  static inline constexpr int vector_length = N;
245 };
246 
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;
251 };
252 
253 template <typename T>
254 using get_vector_element_type = typename simd_like_obj_info<T>::element_type;
255 
256 template <typename T>
257 static inline constexpr int get_vector_length =
258  simd_like_obj_info<T>::vector_length;
259 
260 // @}
261 
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>) {
266  return Val;
267  } else {
268  return __builtin_convertvector(Val, To);
269  }
270 }
271 
272 // calculates the number of elements in "To" type
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);
280  return ToN;
281  }
282 };
283 
284 // Change the element type of a simd vector.
285 template <typename ToEltTy, typename FromEltTy, int FromN,
286  typename = std::enable_if_t<is_vectorizable<ToEltTy>::value>>
287 ESIMD_INLINE typename std::conditional_t<
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) {
292  // Noop.
293  if constexpr (std::is_same<FromEltTy, ToEltTy>::value)
294  return Val;
295 
296  // Bitcast
297  constexpr int ToN = bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems();
298  using VTy = vector_type_t<ToEltTy, ToN>;
299  return reinterpret_cast<VTy>(Val);
300 }
301 
302 // Get computation type of a binary operator given its operand types:
303 // - if both types are arithmetic - return CPP's "common real type" of the
304 // computation (matches C++)
305 // - if both types are simd types, they must be of the same length N,
306 // and the returned type is simd<T, N>, where N is the "common real type" of
307 // the element type of the operands (diverges from clang)
308 // - otherwise, one type is simd and another is arithmetic - the simd type is
309 // returned (matches clang)
310 
311 struct invalid_computation_type;
312 
313 template <class T1, class T2, class SFINAE = void> struct computation_type {
314  using type = invalid_computation_type;
315 };
316 
317 template <class T1, class T2>
318 struct computation_type<T1, T2,
319  std::enable_if_t<is_valid_simd_elem_type_v<T1> &&
320  is_valid_simd_elem_type_v<T2>>> {
321 private:
322  template <class T> using tr = element_type_traits<T>;
323  template <class T>
324  using native_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>;
331 
332 public:
333  using type = std::conditional_t<
334  !is_wr1 && !is_wr2,
335  // T1 and T2 are both std C++ types - use std C++ type promotion
336  decltype(std::declval<T1>() + std::declval<T2>()),
338  std::is_same_v<T1, T2>,
339  // Types are the same wrapper type - return any
340  T1,
341  std::conditional_t<is_fp1 != is_fp2,
342  // One of the types is floating-point - return it
343  // (e.g. computation_type<int, sycl::half> will
344  // yield sycl::half)
345  std::conditional_t<is_fp1, T1, T2>,
346  // both are either floating point or integral -
347  // return result of C++ promotion of the native
348  // types
349  decltype(std::declval<native_t<T1>>() +
350  std::declval<native_t<T2>>())>>>;
351 };
352 
353 template <class T1, class T2>
354 struct computation_type<
355  T1, T2,
356  std::enable_if_t<is_simd_like_type_v<T1> || is_simd_like_type_v<T2>>> {
357 private:
358  using Ty1 = element_type_t<T1>;
359  using Ty2 = element_type_t<T2>;
360  using EltTy = typename computation_type<Ty1, Ty2>::type;
361 
362  static constexpr int N1 = [] {
363  if constexpr (is_simd_like_type_v<T1>) {
364  return T1::length;
365  } else {
366  return 0;
367  }
368  }();
369  static constexpr int N2 = [] {
370  if constexpr (is_simd_like_type_v<T2>) {
371  return T2::length;
372  } else {
373  return 0;
374  }
375  }();
376  static_assert((N1 == N2) || (N1 == 0) || (N2 == 0), "size mismatch");
377  static constexpr int N = N1 ? N1 : N2;
378 
379 public:
380  using type = simd<EltTy, N>;
381 };
382 
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;
386 
387 } // namespace detail
388 
390 
391 // Alias for backward compatibility.
392 template <int N> using mask_type_t = detail::simd_mask_storage_t<N>;
393 
394 } // namespace ext::intel::esimd
395 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
396 } // namespace sycl
T
common.hpp
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:34
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::ext::intel::esimd::simd_view
This class represents a reference to a sub-region of a base simd object.
Definition: types.hpp:35
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
stl_type_traits.hpp
defines.hpp
elem_type_traits.hpp
types_elementary.hpp
simd
Definition: simd.hpp:1027
std
Definition: accessor.hpp:3230
half_type.hpp
sycl::_V1::ext::intel::esimd::mask_type_t
detail::simd_mask_storage_t< N > mask_type_t
Definition: types.hpp:392
region.hpp
sycl::_V1::ext::intel::esimd::convert
ESIMD_INLINE simd< To, N > convert(const simd< From, N > &val)
Covert from a simd object with element type From to a simd object with element type To.
Definition: simd.hpp:184
sycl::_V1::detail::conditional_t
typename std::conditional< B, T, F >::type conditional_t
Definition: stl_type_traits.hpp:27