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 
14 #include <CL/sycl/detail/stl_type_traits.hpp> // to define C++14,17 extensions
15 #include <CL/sycl/half_type.hpp>
18 
19 #if defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
20 #define __esimd_dbg_print(a) std::cout << ">>> " << #a << "\n"
21 #else
22 #define __esimd_dbg_print(a)
23 #endif // defined(__ESIMD_DBG_HOST) && !defined(__SYCL_DEVICE_ONLY__)
24 
25 #include <cstdint>
26 
28 namespace __ESIMD_NS {
29 
30 // simd and simd_view_impl forward declarations
31 template <typename Ty, int N> class simd;
32 template <typename BaseTy, typename RegionTy> class simd_view;
33 
35 
36 namespace detail {
37 
38 namespace csd = cl::sycl::detail;
39 
40 template <int N>
41 using uint_type_t = std::conditional_t<
42  N == 1, uint8_t,
44  N == 2, uint16_t,
45  std::conditional_t<N == 4, uint32_t,
46  std::conditional_t<N == 8, uint64_t, void>>>>;
47 
48 // forward declarations of major internal simd classes
49 template <typename Ty, int N> class simd_mask_impl;
50 template <typename RawT, int N, class Derived, class SFINAE = void>
51 class simd_obj_impl;
52 
53 // @{
54 // Helpers for major simd classes, which don't require their definitions to
55 // compile. Error checking/SFINAE is not used as these are only used internally.
56 
57 using simd_mask_elem_type = unsigned short;
58 template <int N> using simd_mask_type = simd_mask_impl<simd_mask_elem_type, N>;
59 
60 // @{
61 // Checks if given type T is a raw clang vector type, plus provides some info
62 // about it if it is.
63 
64 struct invalid_element_type;
65 
66 template <class T> struct is_clang_vector_type : std::false_type {
67  static inline constexpr int length = 0;
68  using element_type = invalid_element_type;
69 };
70 
71 template <class T, int N>
72 struct is_clang_vector_type<T __attribute__((ext_vector_type(N)))>
73  : std::true_type {
74  static inline constexpr int length = N;
75  using element_type = T;
76 };
77 template <class T>
78 static inline constexpr bool is_clang_vector_type_v =
79  is_clang_vector_type<T>::value;
80 
81 // @}
82 
83 template <typename T>
84 using remove_cvref_t = csd::remove_cv_t<csd::remove_reference_t<T>>;
85 
86 // is_esimd_arithmetic_type
87 template <class...> struct make_esimd_void {
88  using type = void;
89 };
90 template <class... Tys>
91 using __esimd_void_t = typename make_esimd_void<Tys...>::type;
92 
93 template <class Ty, class = void>
94 struct is_esimd_arithmetic_type : std::false_type {};
95 
96 template <class Ty>
97 struct is_esimd_arithmetic_type<
98  Ty, __esimd_void_t<std::enable_if_t<std::is_arithmetic_v<Ty>>,
99  decltype(std::declval<Ty>() + std::declval<Ty>()),
100  decltype(std::declval<Ty>() - std::declval<Ty>()),
101  decltype(std::declval<Ty>() * std::declval<Ty>()),
102  decltype(std::declval<Ty>() / std::declval<Ty>())>>
103  : std::true_type {};
104 
105 template <typename Ty>
106 static inline constexpr bool is_esimd_arithmetic_type_v =
107  is_esimd_arithmetic_type<Ty>::value;
108 
109 // is_vectorizable_type
110 template <typename Ty>
111 struct is_vectorizable : std::conditional_t<is_esimd_arithmetic_type_v<Ty>,
112  std::true_type, std::false_type> {};
113 
114 template <typename Ty>
115 static inline constexpr bool is_vectorizable_v = is_vectorizable<Ty>::value;
116 
117 template <typename T>
118 static inline constexpr bool is_esimd_scalar_v =
120 
121 template <typename T>
122 using is_esimd_scalar = typename std::bool_constant<is_esimd_scalar_v<T>>;
123 
124 // raw_vector_type, using clang vector type extension.
125 template <typename Ty, int N> struct raw_vector_type {
126  static_assert(!std::is_const<Ty>::value, "const element type not supported");
127  static_assert(is_vectorizable_v<Ty>, "element type not supported");
128  static_assert(N > 0, "zero-element vector not supported");
129 
130  static constexpr int length = N;
131  using type = Ty __attribute__((ext_vector_type(N)));
132 };
133 
134 template <typename Ty, int N>
135 using vector_type_t = typename raw_vector_type<Ty, N>::type;
136 
137 // @{
138 // Checks if given type T derives from simd_obj_impl or is equal to it.
139 template <typename T>
140 struct is_simd_obj_impl_derivative : public std::false_type {};
141 
142 // Specialization for the simd_obj_impl type itself.
143 template <typename RawT, int N, class Derived>
144 struct is_simd_obj_impl_derivative<simd_obj_impl<RawT, N, Derived>>
145  : public std::true_type {};
146 
147 template <class T, class SFINAE = void> struct element_type_traits;
148 template <class T>
149 using __raw_t = typename __ESIMD_DNS::element_type_traits<T>::RawT;
150 template <class T>
151 using __cpp_t = typename __ESIMD_DNS::element_type_traits<T>::EnclosingCppT;
152 
153 // Specialization for all other types.
154 template <typename T, int N, template <typename, int> class Derived>
155 struct is_simd_obj_impl_derivative<Derived<T, N>>
156  : public std::conditional_t<
157  std::is_base_of_v<simd_obj_impl<__raw_t<T>, N, Derived<T, N>>,
158  Derived<T, N>>,
159  std::true_type, std::false_type> {};
160 
161 // Convenience shortcut.
162 template <typename T>
163 inline constexpr bool is_simd_obj_impl_derivative_v =
164  is_simd_obj_impl_derivative<T>::value;
165 // @}
166 
167 // @{
168 // "Resizes" given simd type \c T to given number of elements \c N.
169 template <class SimdT, int Ndst> struct resize_a_simd_type;
170 
171 // Specialization for the simd_obj_impl type.
172 template <typename T, int Nsrc, int Ndst, template <typename, int> class SimdT>
173 struct resize_a_simd_type<simd_obj_impl<__raw_t<T>, Nsrc, SimdT<T, Nsrc>>,
174  Ndst> {
175  using type = simd_obj_impl<__raw_t<T>, Ndst, SimdT<T, Ndst>>;
176 };
177 
178 // Specialization for the simd_obj_impl type derivatives.
179 template <typename T, int Nsrc, int Ndst, template <typename, int> class SimdT>
180 struct resize_a_simd_type<SimdT<T, Nsrc>, Ndst> {
181  using type = SimdT<T, Ndst>;
182 };
183 
184 // Convenience shortcut.
185 template <class SimdT, int Ndst>
186 using resize_a_simd_type_t = typename resize_a_simd_type<SimdT, Ndst>::type;
187 // @}
188 
189 // @{
190 // Converts element type of given simd type \c SimdT to
191 // given scalar type \c NewElemT.
192 template <class SimdT, typename NewElemT> struct convert_simd_elem_type;
193 
194 // Specialization for the simd_obj_impl type.
195 template <typename OldElemT, int N, typename NewElemT,
196  template <typename, int> class SimdT>
197 struct convert_simd_elem_type<
198  simd_obj_impl<__raw_t<OldElemT>, N, SimdT<OldElemT, N>>, NewElemT> {
199  using type = simd_obj_impl<__raw_t<NewElemT>, N, SimdT<NewElemT, N>>;
200 };
201 
202 // Specialization for the simd_obj_impl type derivatives.
203 template <typename OldElemT, int N, typename NewElemT,
204  template <typename, int> class SimdT>
205 struct convert_simd_elem_type<SimdT<OldElemT, N>, NewElemT> {
206  using type = SimdT<NewElemT, N>;
207 };
208 
209 // Convenience shortcut.
210 template <class SimdT, typename NewElemT>
211 using convert_simd_elem_type_t =
212  typename convert_simd_elem_type<SimdT, NewElemT>::type;
213 
214 // @}
215 
216 // Constructs a simd type with the same template type as in \c SimdT, and
217 // given element type and number.
218 template <class SimdT, typename T, int N>
219 using construct_a_simd_type_t =
220  convert_simd_elem_type_t<resize_a_simd_type_t<SimdT, N>, T>;
221 
222 // @}
223 
224 // must match simd_mask<N>::element_type
225 template <int N>
226 using simd_mask_storage_t = vector_type_t<simd_mask_elem_type, N>;
227 
228 // @{
229 // Checks if given type is a view of any simd type (simd or simd_mask).
230 template <typename Ty> struct is_any_simd_view_type : std::false_type {};
231 
232 template <typename BaseTy, typename RegionTy>
233 struct is_any_simd_view_type<simd_view<BaseTy, RegionTy>> : std::true_type {};
234 
235 template <typename Ty>
236 static inline constexpr bool is_any_simd_view_type_v =
237  is_any_simd_view_type<Ty>::value;
238 // @}
239 
240 // @{
241 // Check if a type is one of internal 'simd_xxx_impl' types exposing simd-like
242 // interfaces and behaving like a simd object type.
243 
244 template <typename Ty>
245 static inline constexpr bool is_simd_like_type_v =
246  is_any_simd_view_type_v<Ty> || is_simd_obj_impl_derivative_v<Ty>;
247 // @}
248 
249 // @{
250 // Checks if given type is a any of the user-visible simd types (simd or
251 // simd_mask).
252 template <typename Ty> struct is_simd_type : std::false_type {};
253 template <typename ElTy, int N>
254 struct is_simd_type<simd<ElTy, N>> : std::true_type {};
255 template <typename Ty>
256 static inline constexpr bool is_simd_type_v = is_simd_type<Ty>::value;
257 
258 template <typename Ty> struct is_simd_mask_type : std::false_type {};
259 template <int N>
260 struct is_simd_mask_type<simd_mask_impl<simd_mask_elem_type, N>>
261  : std::true_type {};
262 template <typename Ty>
263 static inline constexpr bool is_simd_mask_type_v = is_simd_mask_type<Ty>::value;
264 // @}
265 
266 // @{
267 // Checks if given type is a view of the simd type.
268 template <typename Ty> struct is_simd_view_type_impl : std::false_type {};
269 
270 template <class BaseT, class RegionT>
271 struct is_simd_view_type_impl<simd_view<BaseT, RegionT>>
272  : std::conditional_t<is_simd_type_v<BaseT>, std::true_type,
273  std::false_type> {};
274 
275 template <class Ty>
276 struct is_simd_view_type : is_simd_view_type_impl<remove_cvref_t<Ty>> {};
277 
278 template <typename Ty>
279 static inline constexpr bool is_simd_view_type_v = is_simd_view_type<Ty>::value;
280 // @}
281 
282 template <typename T>
283 static inline constexpr bool is_simd_or_view_type_v =
284  is_simd_view_type_v<T> || is_simd_type_v<T>;
285 
286 // @{
287 // Get the element type if it is a scalar, clang vector, simd or simd_view type.
288 
289 struct cant_deduce_element_type;
290 
291 template <class T, class SFINAE = void> struct element_type {
292  using type = cant_deduce_element_type;
293 };
294 
295 template <typename T>
296 struct element_type<T, std::enable_if_t<is_vectorizable_v<T>>> {
297  using type = remove_cvref_t<T>;
298 };
299 
300 template <typename T>
301 struct element_type<T, std::enable_if_t<is_simd_like_type_v<T>>> {
302  using type = typename T::element_type;
303 };
304 
305 template <typename T>
306 struct element_type<T, std::enable_if_t<is_clang_vector_type_v<T>>> {
307  using type = typename is_clang_vector_type<T>::element_type;
308 };
309 
310 template <typename T> using element_type_t = typename element_type<T>::type;
311 
312 // Determine element type of simd_obj_impl's Derived type w/o having to have
313 // complete instantiation of the Derived type (is required by element_type_t,
314 // hence can't be used here).
315 template <class T> struct simd_like_obj_info {
316  using element_type = T;
317  static inline constexpr int vector_length = 0;
318 };
319 
320 template <class T, int N> struct simd_like_obj_info<simd<T, N>> {
321  using element_type = T;
322  static inline constexpr int vector_length = N;
323 };
324 
325 template <class T, int N> struct simd_like_obj_info<simd_mask_impl<T, N>> {
326  using element_type = simd_mask_elem_type; // equals T
327  static inline constexpr int vector_length = N;
328 };
329 
330 template <class BaseT, class RegionT>
331 struct simd_like_obj_info<simd_view<BaseT, RegionT>> {
332  using element_type = typename RegionT::element_type;
333  static inline constexpr int vector_length = RegionT::length;
334 };
335 
336 template <typename T>
337 using get_vector_element_type = typename simd_like_obj_info<T>::element_type;
338 
339 template <typename T>
340 static inline constexpr int get_vector_length =
341  simd_like_obj_info<T>::vector_length;
342 
343 // @}
344 
345 template <typename To, typename From>
346 std::enable_if_t<is_clang_vector_type_v<To> && is_clang_vector_type_v<From>, To>
347  ESIMD_INLINE convert(From Val) {
348  if constexpr (std::is_same_v<To, From>) {
349  return Val;
350  } else {
351  return __builtin_convertvector(Val, To);
352  }
353 }
354 
356 template <typename U> constexpr bool is_type() { return false; }
357 
358 template <typename U, typename T, typename... Ts> constexpr bool is_type() {
359  using UU = typename csd::remove_const_t<U>;
360  using TT = typename csd::remove_const_t<T>;
361  return std::is_same<UU, TT>::value || is_type<UU, Ts...>();
362 }
363 
364 // calculates the number of elements in "To" type
365 template <typename ToEltTy, typename FromEltTy, int FromN,
366  typename = csd::enable_if_t<is_vectorizable<ToEltTy>::value>>
367 struct bitcast_helper {
368  static inline constexpr int nToElems() {
369  constexpr int R1 = sizeof(ToEltTy) / sizeof(FromEltTy);
370  constexpr int R2 = sizeof(FromEltTy) / sizeof(ToEltTy);
371  constexpr int ToN = (R2 > 0) ? (FromN * R2) : (FromN / R1);
372  return ToN;
373  }
374 };
375 
376 // Change the element type of a simd vector.
377 template <typename ToEltTy, typename FromEltTy, int FromN,
378  typename = csd::enable_if_t<is_vectorizable<ToEltTy>::value>>
379 ESIMD_INLINE typename csd::conditional_t<
380  std::is_same<FromEltTy, ToEltTy>::value, vector_type_t<FromEltTy, FromN>,
381  vector_type_t<ToEltTy,
382  bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems()>>
383 bitcast(vector_type_t<FromEltTy, FromN> Val) {
384  // Noop.
385  if constexpr (std::is_same<FromEltTy, ToEltTy>::value)
386  return Val;
387 
388  // Bitcast
389  constexpr int ToN = bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems();
390  using VTy = vector_type_t<ToEltTy, ToN>;
391  return reinterpret_cast<VTy>(Val);
392 }
393 
394 } // namespace detail
395 
397 
398 // Alias for backward compatibility.
399 template <int N> using mask_type_t = detail::simd_mask_storage_t<N>;
400 
401 } // namespace __ESIMD_NS
402 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::is_arithmetic
Definition: type_traits.hpp:228
T
simd_obj_impl
stl_type_traits.hpp
cl::sycl::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
cl::sycl::detail::conditional_t
typename std::conditional< B, T, F >::type conditional_t
Definition: stl_type_traits.hpp:27
cl::sycl::ext::intel::esimd::simd_view
This class represents a reference to a sub-region of a base simd object.
Definition: types.hpp:32
cl::sycl::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:31
defines.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
simd
Definition: simd.hpp:1027
sycl::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:293
std
Definition: accessor.hpp:2617
cl::sycl::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:183
half_type.hpp
region.hpp
common.hpp
cl::sycl::detail
Definition: access.hpp:103
cl::sycl::ext::intel::esimd::simd_view::element_type
typename ShapeTy::element_type element_type
The element type of this class, which could be different from the element type of the base object typ...
Definition: simd_view.hpp:63
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::ext::intel::esimd::mask_type_t
detail::simd_mask_storage_t< N > mask_type_t
Definition: types.hpp:399
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12