DPC++ Runtime
Runtime libraries for oneAPI DPC++
simd.hpp
Go to the documentation of this file.
1 //==------------ - simd.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 // Implement Explicit SIMD vector APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
15 
21 
23 
24 #ifndef __SYCL_DEVICE_ONLY__
26 #endif // __SYCL_DEVICE_ONLY__
27 
28 namespace sycl {
30 namespace ext::intel::esimd {
31 
34 
39 
41 
44 
54 template <typename Ty, int N>
55 class simd : public detail::simd_obj_impl<
56  detail::__raw_t<Ty>, N, simd<Ty, N>,
57  std::enable_if_t<detail::is_valid_simd_elem_type_v<Ty>>> {
58  using base_type = detail::simd_obj_impl<detail::__raw_t<Ty>, N, simd<Ty, N>>;
59 
60 public:
61  using base_type::base_type;
62  using element_type = Ty;
65  static constexpr int length = N;
66 
74  template <typename SimdT,
75  class = std::enable_if_t<__ESIMD_DNS::is_simd_type_v<SimdT> &&
76  (length == SimdT::length)>>
77  simd(const SimdT &RHS)
78  : base_type(detail::convert_vector<Ty, detail::element_type_t<SimdT>, N>(
79  RHS.data())) {
80  __esimd_dbg_print(simd(const SimdT &RHS));
81  }
82 
83  // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
84  template <
85  int N1 = N, class Ty1 = Ty,
86  class SFINAE = std::enable_if_t<
88  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
89  simd(const sycl::ext::oneapi::experimental::simd<Ty, N1> &v)
90  : simd(static_cast<raw_vector_type>(v)) {}
91 
97  template <typename T1,
98  class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
99  simd(T1 Val) : base_type(Val) {
100  __esimd_dbg_print(simd(T1 Val));
101  }
102 
108  template <class To, class T = simd,
110  (T::length == 1) && detail::is_valid_simd_elem_type_v<To>>>
111  operator To() const {
112  __esimd_dbg_print(operator To());
113  return detail::convert_scalar<To, element_type>(base_type::data()[0]);
114  }
115 
120  template <
121  int N1, class Ty1 = Ty,
122  class SFINAE = std::enable_if_t<
124  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
125  operator sycl::ext::oneapi::experimental::simd<Ty, N1>() {
126  return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
127  }
128 
132  *this += 1;
133  return *this;
134  }
135 
140  simd Ret(base_type::data());
141  operator++();
142  return Ret;
143  }
144 
148  *this -= 1;
149  return *this;
150  }
151 
156  simd Ret(base_type::data());
157  operator--();
158  return Ret;
159  }
160 
161 #define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID) \
162  template <class T1 = Ty> simd operator ARITH_UNARY_OP() const { \
163  static_assert(!std::is_unsigned_v<T1>, \
164  #ARITH_UNARY_OP "doesn't apply to unsigned types"); \
165  return simd{detail::vector_unary_op<detail::UnaryOp::ID, T1, N>( \
166  base_type::data())}; \
167  }
168 
173 #undef __ESIMD_DEF_SIMD_ARITH_UNARY_OP
174 };
175 
177 
180 
183 template <typename To, typename From, int N>
184 ESIMD_INLINE simd<To, N> convert(const simd<From, N> &val) {
185  if constexpr (std::is_same_v<std::remove_const_t<To>,
186  std::remove_const_t<From>>)
187  return val;
188  else
189  return detail::convert_vector<To, From, N>(val.data());
190 }
192 
195 
198 template <int N> using simd_mask = detail::simd_mask_type<N>;
199 
201 
202 } // namespace ext::intel::esimd
203 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
204 } // namespace sycl
205 
209 template <typename Ty, int N>
210 std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
211 #ifdef __SYCL_DEVICE_ONLY__
212  {}
213 #else
214 {
215  OS << "{";
216  for (int I = 0; I < N; I++) {
217  OS << V[I];
218  if (I < N - 1)
219  OS << ",";
220  }
221  OS << "}";
222  return OS;
223 }
224 #endif // __SYCL_DEVICE_ONLY__
sycl::_V1::ext::intel::esimd::simd::raw_vector_type
typename base_type::raw_vector_type raw_vector_type
Definition: simd.hpp:64
simd_mask
Definition: simd.hpp:1029
__ESIMD_DEF_SIMD_ARITH_UNARY_OP
#define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID)
Definition: simd.hpp:161
simd_view.hpp
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
max_fixed_size
constexpr size_t max_fixed_size
Definition: simd.hpp:1008
sycl::_V1::ext::intel::esimd::simd::simd
simd(T1 Val)
Broadcast constructor with conversion.
Definition: simd.hpp:99
sycl::_V1::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:34
sycl::_V1::ext::intel::esimd::detail::simd_obj_impl::raw_vector_type
vector_type_t< RawTy, N > raw_vector_type
The underlying raw storage vector data type.
Definition: simd_obj_impl.hpp:175
sycl::_V1::ext::oneapi::plus
std::plus< T > plus
Definition: functional.hpp:19
sycl::_V1::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:13
sycl::_V1::ext::intel::esimd::simd::element_type
Ty element_type
Definition: simd.hpp:62
sycl::_V1::ext::intel::esimd::detail::simd_obj_impl
This is a base class for all ESIMD simd classes with real storage (simd, simd_mask_impl).
Definition: simd_obj_impl.hpp:155
intrin.hpp
sycl::_V1::ext::intel::esimd::detail::simd_obj_impl::raw_element_type
RawTy raw_element_type
The element type of the raw storage vector.
Definition: simd_obj_impl.hpp:178
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
__esimd_dbg_print
#define __esimd_dbg_print(a)
Definition: types.hpp:24
sycl::_V1::ext::intel::esimd::simd::operator--
simd & operator--()
Prefix decrement, decrements elements of this object.
Definition: simd.hpp:147
sycl::_V1::ext::intel::esimd::simd::operator++
simd & operator++()
Prefix increment, increments elements of this object.
Definition: simd.hpp:131
types.hpp
simd_obj_impl.hpp
sycl::_V1::ext::intel::esimd::simd::simd
simd(const SimdT &RHS)
Implicit conversion constructor from another simd object of the same length.
Definition: simd.hpp:77
simd_mask_impl.hpp
sycl::_V1::ext::oneapi::experimental::simd
std::experimental::simd< T, simd_abi::native_fixed_size< T, N > > simd
Definition: invoke_simd.hpp:90
sycl::_V1::ext::intel::esimd::simd::raw_element_type
typename base_type::raw_element_type raw_element_type
Definition: simd.hpp:63
sycl::_V1::ext::intel::esimd::simd::operator--
simd operator--(int)
Postfix decrement.
Definition: simd.hpp:155
sycl::_V1::ext::intel::esimd::simd::operator++
simd operator++(int)
Postfix increment.
Definition: simd.hpp:139
iostream_proxy.hpp
simd
Definition: simd.hpp:1027
sycl_util.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
invoke_simd.hpp
memory_intrin.hpp
sycl::_V1::ext::intel::esimd::simd::simd
simd(const sycl::ext::oneapi::experimental::simd< Ty, N1 > &v)
Definition: simd.hpp:89
operator<<
std::ostream & operator<<(std::ostream &OS, const sycl::ext::intel::esimd::simd< Ty, N > &V)
Definition: simd.hpp:210