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::puts(">>> " #a);
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 {
30 inline namespace _V1 {
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 = std::is_arithmetic_v<T>;
55 
56 template <typename T>
57 using is_esimd_scalar = typename std::bool_constant<is_esimd_scalar_v<T>>;
58 
59 // @{
60 // Checks if given type T derives from simd_obj_impl or is equal to it.
61 template <typename T>
62 struct is_simd_obj_impl_derivative : public std::false_type {};
63 
64 // Specialization for the simd_obj_impl type itself.
65 template <typename RawT, int N, class Derived>
66 struct is_simd_obj_impl_derivative<simd_obj_impl<RawT, N, Derived>>
67  : public std::true_type {};
68 
69 // Specialization for all other types.
70 template <typename T, int N, template <typename, int> class Derived>
71 struct is_simd_obj_impl_derivative<Derived<T, N>>
72  : public std::conditional_t<
73  std::is_base_of_v<simd_obj_impl<__raw_t<T>, N, Derived<T, N>>,
74  Derived<T, N>>,
75  std::true_type, std::false_type> {};
76 
77 // Convenience shortcut.
78 template <typename T>
79 inline constexpr bool is_simd_obj_impl_derivative_v =
80  is_simd_obj_impl_derivative<T>::value;
81 // @}
82 
83 // @{
84 // "Resizes" given simd type \c T to given number of elements \c N.
85 template <class SimdT, int Ndst> struct resize_a_simd_type;
86 
87 // Specialization for the simd_obj_impl type.
88 template <typename T, int Nsrc, int Ndst, template <typename, int> class SimdT>
89 struct resize_a_simd_type<simd_obj_impl<__raw_t<T>, Nsrc, SimdT<T, Nsrc>>,
90  Ndst> {
91  using type = simd_obj_impl<__raw_t<T>, Ndst, SimdT<T, Ndst>>;
92 };
93 
94 // Specialization for the simd_obj_impl type derivatives.
95 template <typename T, int Nsrc, int Ndst, template <typename, int> class SimdT>
96 struct resize_a_simd_type<SimdT<T, Nsrc>, Ndst> {
97  using type = SimdT<T, Ndst>;
98 };
99 
100 // Convenience shortcut.
101 template <class SimdT, int Ndst>
102 using resize_a_simd_type_t = typename resize_a_simd_type<SimdT, Ndst>::type;
103 // @}
104 
105 // @{
106 // Converts element type of given simd type \c SimdT to
107 // given scalar type \c NewElemT.
108 template <class SimdT, typename NewElemT> struct convert_simd_elem_type;
109 
110 // Specialization for the simd_obj_impl type.
111 template <typename OldElemT, int N, typename NewElemT,
112  template <typename, int> class SimdT>
113 struct convert_simd_elem_type<
114  simd_obj_impl<__raw_t<OldElemT>, N, SimdT<OldElemT, N>>, NewElemT> {
115  using type = simd_obj_impl<__raw_t<NewElemT>, N, SimdT<NewElemT, N>>;
116 };
117 
118 // Specialization for the simd_obj_impl type derivatives.
119 template <typename OldElemT, int N, typename NewElemT,
120  template <typename, int> class SimdT>
121 struct convert_simd_elem_type<SimdT<OldElemT, N>, NewElemT> {
122  using type = SimdT<NewElemT, N>;
123 };
124 
125 // Convenience shortcut.
126 template <class SimdT, typename NewElemT>
127 using convert_simd_elem_type_t =
128  typename convert_simd_elem_type<SimdT, NewElemT>::type;
129 
130 // @}
131 
132 // Constructs a simd type with the same template type as in \c SimdT, and
133 // given element type and number.
134 template <class SimdT, typename T, int N>
135 using construct_a_simd_type_t =
136  convert_simd_elem_type_t<resize_a_simd_type_t<SimdT, N>, T>;
137 
138 // @}
139 
140 // must match simd_mask<N>::element_type
141 template <int N>
142 using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;
143 
144 // @{
145 // Checks if given type is a view of any simd type (simd or simd_mask).
146 template <typename Ty> struct is_any_simd_view_type : std::false_type {};
147 
148 template <typename BaseTy, typename RegionTy>
149 struct is_any_simd_view_type<simd_view<BaseTy, RegionTy>> : std::true_type {};
150 
151 template <typename Ty>
152 static inline constexpr bool is_any_simd_view_type_v =
153  is_any_simd_view_type<Ty>::value;
154 // @}
155 
156 // @{
157 // Check if a type is one of internal 'simd_xxx_impl' types exposing simd-like
158 // interfaces and behaving like a simd object type.
159 
160 template <typename Ty>
161 static inline constexpr bool is_simd_like_type_v =
162  is_any_simd_view_type_v<Ty> || is_simd_obj_impl_derivative_v<Ty>;
163 // @}
164 
165 // @{
166 // Checks if given type is a any of the user-visible simd types (simd or
167 // simd_mask).
168 template <typename Ty> struct is_simd_type : std::false_type {};
169 template <typename ElTy, int N>
170 struct is_simd_type<simd<ElTy, N>> : std::true_type {};
171 template <typename Ty>
172 static inline constexpr bool is_simd_type_v = is_simd_type<Ty>::value;
173 
174 template <typename Ty> struct is_simd_mask_type : std::false_type {};
175 template <int N>
176 struct is_simd_mask_type<simd_mask_impl<simd_mask_elem_type, N>>
177  : std::true_type {};
178 template <typename Ty>
179 static inline constexpr bool is_simd_mask_type_v = is_simd_mask_type<Ty>::value;
180 // @}
181 
182 // @{
183 // Checks if given type is a view of the simd type.
184 template <typename Ty> struct is_simd_view_type_impl : std::false_type {};
185 
186 template <class BaseT, class RegionT>
187 struct is_simd_view_type_impl<simd_view<BaseT, RegionT>>
188  : std::conditional_t<is_simd_type_v<BaseT>, std::true_type,
189  std::false_type> {};
190 
191 template <class Ty>
192 struct is_simd_view_type : is_simd_view_type_impl<remove_cvref_t<Ty>> {};
193 
194 template <typename Ty>
195 static inline constexpr bool is_simd_view_type_v = is_simd_view_type<Ty>::value;
196 // @}
197 
198 template <typename T>
199 static inline constexpr bool is_simd_or_view_type_v =
200  is_simd_view_type_v<T> || is_simd_type_v<T>;
201 
202 // @{
203 // Get the element type if it is a scalar, clang vector, simd or simd_view type.
204 
205 struct cant_deduce_element_type;
206 
207 template <class T, class SFINAE = void> struct element_type {
208  using type = cant_deduce_element_type;
209 };
210 
211 template <typename T>
212 struct element_type<T, std::enable_if_t<is_vectorizable_v<T>>> {
213  using type = remove_cvref_t<T>;
214 };
215 
216 template <typename T>
217 struct element_type<T, std::enable_if_t<is_simd_like_type_v<T>>> {
218  using type = typename T::element_type;
219 };
220 
221 template <typename T>
222 struct element_type<T, std::enable_if_t<is_clang_vector_type_v<T>>> {
223  using type = typename is_clang_vector_type<T>::element_type;
224 };
225 
226 template <typename T> using element_type_t = typename element_type<T>::type;
227 
228 // Determine element type of simd_obj_impl's Derived type w/o having to have
229 // complete instantiation of the Derived type (is required by element_type_t,
230 // hence can't be used here).
231 template <class T> struct simd_like_obj_info {
232  using element_type = T;
233  static constexpr int vector_length = 0;
234 };
235 
236 template <class T, int N> struct simd_like_obj_info<simd<T, N>> {
237  using element_type = T;
238  static constexpr int vector_length = N;
239 };
240 
241 template <class T, int N> struct simd_like_obj_info<simd_mask_impl<T, N>> {
242  using element_type = simd_mask_elem_type; // equals T
243  static constexpr int vector_length = N;
244 };
245 
246 template <class BaseT, class RegionT>
247 struct simd_like_obj_info<simd_view<BaseT, RegionT>> {
248  using element_type = typename RegionT::element_type;
249  static constexpr int vector_length = RegionT::length;
250 };
251 
252 template <typename T>
253 using get_vector_element_type = typename simd_like_obj_info<T>::element_type;
254 
255 template <typename T>
256 static inline constexpr int get_vector_length =
257  simd_like_obj_info<T>::vector_length;
258 
259 // @}
260 
261 template <typename To, typename From>
262 std::enable_if_t<is_clang_vector_type_v<To> && is_clang_vector_type_v<From>, To>
263  ESIMD_INLINE convert(From Val) {
264  if constexpr (std::is_same_v<To, From>) {
265  return Val;
266  } else {
267  return __builtin_convertvector(Val, To);
268  }
269 }
270 
271 // calculates the number of elements in "To" type
272 template <typename ToEltTy, typename FromEltTy, int FromN,
273  typename = std::enable_if_t<is_vectorizable<ToEltTy>::value>>
274 struct bitcast_helper {
275  static constexpr int nToElems() {
276  constexpr int R1 = sizeof(ToEltTy) / sizeof(FromEltTy);
277  constexpr int R2 = sizeof(FromEltTy) / sizeof(ToEltTy);
278  constexpr int ToN = (R2 > 0) ? (FromN * R2) : (FromN / R1);
279  return ToN;
280  }
281 };
282 
283 // Change the element type of a simd vector.
284 template <typename ToEltTy, typename FromEltTy, int FromN,
285  typename = std::enable_if_t<is_vectorizable<ToEltTy>::value>>
286 ESIMD_INLINE typename std::conditional_t<
287  std::is_same_v<FromEltTy, ToEltTy>, vector_type_t<FromEltTy, FromN>,
288  vector_type_t<ToEltTy,
289  bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems()>>
290 bitcast(vector_type_t<FromEltTy, FromN> Val) {
291  // Noop.
292  if constexpr (std::is_same_v<FromEltTy, ToEltTy>)
293  return Val;
294 
295  // Bitcast
296  constexpr int ToN = bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems();
297  using VTy = vector_type_t<ToEltTy, ToN>;
298  return reinterpret_cast<VTy>(Val);
299 }
300 
301 // Get computation type of a binary operator given its operand types:
302 // - if both types are arithmetic - return CPP's "common real type" of the
303 // computation (matches C++)
304 // - if both types are simd types, they must be of the same length N,
305 // and the returned type is simd<T, N>, where N is the "common real type" of
306 // the element type of the operands (diverges from clang)
307 // - otherwise, one type is simd and another is arithmetic - the simd type is
308 // returned (matches clang)
309 
310 struct invalid_computation_type;
311 
312 template <class T1, class T2, class SFINAE = void> struct computation_type {
313  using type = invalid_computation_type;
314 };
315 
316 template <class T1, class T2>
317 struct computation_type<T1, T2,
318  std::enable_if_t<is_valid_simd_elem_type_v<T1> &&
319  is_valid_simd_elem_type_v<T2>>> {
320 private:
321  template <class T> using tr = element_type_traits<T>;
322  template <class T>
323  using native_t =
324  std::conditional_t<tr<T>::use_native_cpp_ops, typename tr<T>::RawT,
325  typename tr<T>::EnclosingCppT>;
326  static constexpr bool is_wr1 = is_wrapper_elem_type_v<T1>;
327  static constexpr bool is_wr2 = is_wrapper_elem_type_v<T2>;
328  static constexpr bool is_fp1 = is_generic_floating_point_v<T1>;
329  static constexpr bool is_fp2 = is_generic_floating_point_v<T2>;
330 
331 public:
332  using type = std::conditional_t<
333  !is_wr1 && !is_wr2,
334  // T1 and T2 are both std C++ types - use std C++ type promotion
335  decltype(std::declval<T1>() + std::declval<T2>()),
336  std::conditional_t<
337  std::is_same_v<T1, T2>,
338  // Types are the same wrapper type - return any
339  T1,
340  std::conditional_t<is_fp1 != is_fp2,
341  // One of the types is floating-point - return it
342  // (e.g. computation_type<int, sycl::half> will
343  // yield sycl::half)
344  std::conditional_t<is_fp1, T1, T2>,
345  // both are either floating point or integral -
346  // return result of C++ promotion of the native
347  // types
348  decltype(std::declval<native_t<T1>>() +
349  std::declval<native_t<T2>>())>>>;
350 };
351 
352 template <class T1, class T2>
353 struct computation_type<
354  T1, T2,
355  std::enable_if_t<is_simd_like_type_v<T1> || is_simd_like_type_v<T2>>> {
356 private:
357  using Ty1 = element_type_t<T1>;
358  using Ty2 = element_type_t<T2>;
359  using EltTy = typename computation_type<Ty1, Ty2>::type;
360 
361  static constexpr int N1 = [] {
362  if constexpr (is_simd_like_type_v<T1>) {
363  return T1::length;
364  } else {
365  return 0;
366  }
367  }();
368  static constexpr int N2 = [] {
369  if constexpr (is_simd_like_type_v<T2>) {
370  return T2::length;
371  } else {
372  return 0;
373  }
374  }();
375  static_assert((N1 == N2) || (N1 == 0) || (N2 == 0), "size mismatch");
376  static constexpr int N = N1 ? N1 : N2;
377 
378 public:
379  using type = simd<EltTy, N>;
380 };
381 
382 template <class T1, class T2 = T1>
383 using computation_type_t =
384  typename computation_type<remove_cvref_t<T1>, remove_cvref_t<T2>>::type;
385 
386 } // namespace detail
387 
389 
390 // Alias for backward compatibility.
391 template <int N> using mask_type_t = detail::simd_mask_storage_t<N>;
392 
393 } // namespace ext::intel::esimd
394 } // namespace _V1
395 } // namespace sycl
Definition: simd.hpp:1387
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:185
detail::simd_mask_storage_t< N > mask_type_t
Definition: types.hpp:391
std::conditional_t< std::is_same_v< ElementType, half >, sycl::detail::half_impl::BIsRepresentationT, ElementType > element_type
Definition: multi_ptr.hpp:752
Definition: access.hpp:18