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 namespace sycl {
25 inline namespace _V1 {
26 namespace ext::intel::esimd {
27 
30 
35 
37 
40 
50 template <typename Ty, int N>
51 class simd : public detail::simd_obj_impl<
52  detail::__raw_t<Ty>, N, simd<Ty, N>,
53  std::enable_if_t<detail::is_valid_simd_elem_type_v<Ty>>> {
55 
56 public:
57  using base_type::base_type;
58  using element_type = Ty;
61  static constexpr int length = N;
62 
70  template <typename SimdT,
71  class = std::enable_if_t<__ESIMD_DNS::is_simd_type_v<SimdT> &&
72  (length == SimdT::length)>>
73  simd(const SimdT &RHS)
74  : base_type(detail::convert_vector<Ty, detail::element_type_t<SimdT>, N>(
75  RHS.data())) {
76  __esimd_dbg_print(simd(const SimdT &RHS));
77  }
78 
79  // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
80  template <
81  int N1 = N, class Ty1 = Ty,
82  class SFINAE = std::enable_if_t<
84  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
86  : simd(static_cast<raw_vector_type>(v)) {}
87 
93  template <typename T1,
94  class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
95  simd(T1 Val) : base_type(Val) {
96  __esimd_dbg_print(simd(T1 Val));
97  }
98 
104  template <class To, class T = simd,
105  class = std::enable_if_t<(T::length == 1) &&
106  detail::is_valid_simd_elem_type_v<To>>>
107  operator To() const {
108  __esimd_dbg_print(operator To());
109  return detail::convert_scalar<To, element_type>(base_type::data()[0]);
110  }
111 
116  template <
117  int N1, class Ty1 = Ty,
118  class SFINAE = std::enable_if_t<
120  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
123  }
124 
126  simd &operator=(const simd &other) noexcept {
127  return base_type::operator=(other);
128  }
129 
133  *this += 1;
134  return *this;
135  }
136 
141  simd Ret(base_type::data());
142  operator++();
143  return Ret;
144  }
145 
149  *this -= 1;
150  return *this;
151  }
152 
157  simd Ret(base_type::data());
158  operator--();
159  return Ret;
160  }
161 
162 #define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID) \
163  template <class T1 = Ty> simd operator ARITH_UNARY_OP() const { \
164  static_assert(!std::is_unsigned_v<T1>, \
165  #ARITH_UNARY_OP "doesn't apply to unsigned types"); \
166  return simd{detail::vector_unary_op<detail::UnaryOp::ID, T1, N>( \
167  base_type::data())}; \
168  }
169 
174 #undef __ESIMD_DEF_SIMD_ARITH_UNARY_OP
175 };
176 
178 
181 
184 template <typename To, typename From, int N>
185 ESIMD_INLINE simd<To, N> convert(const simd<From, N> &val) {
186  if constexpr (std::is_same_v<std::remove_const_t<To>,
187  std::remove_const_t<From>>)
188  return val;
189  else
190  return detail::convert_vector<To, From, N>(val.data());
191 }
193 
196 
199 template <int N> using simd_mask = detail::simd_mask_type<N>;
200 
202 
203 } // namespace ext::intel::esimd
204 } // namespace _V1
205 } // namespace sycl
206 
210 template <typename Ty, int N>
211 std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
212 #ifdef __SYCL_DEVICE_ONLY__
213  {}
214 #else
215 {
216  __ESIMD_UNSUPPORTED_ON_HOST;
217 }
218 #endif // __SYCL_DEVICE_ONLY__
This is a base class for all ESIMD simd classes with real storage (simd, simd_mask_impl).
Derived & operator=(const simd_obj_impl &other) noexcept
Copy assignment operator.
RawTy raw_element_type
The element type of the raw storage vector.
vector_type_t< RawTy, N > raw_vector_type
The underlying raw storage vector data type.
The main simd vector class.
Definition: simd.hpp:53
typename base_type::raw_vector_type raw_vector_type
Definition: simd.hpp:60
simd & operator++()
Prefix increment, increments elements of this object.
Definition: simd.hpp:132
simd(const SimdT &RHS)
Implicit conversion constructor from another simd object of the same length.
Definition: simd.hpp:73
static constexpr int length
Definition: simd.hpp:61
typename base_type::raw_element_type raw_element_type
Definition: simd.hpp:59
simd & operator=(const simd &other) noexcept
Copy assignment operator.
Definition: simd.hpp:126
simd(const sycl::ext::oneapi::experimental::simd< Ty, N1 > &v)
Definition: simd.hpp:85
simd(T1 Val)
Broadcast constructor with conversion.
Definition: simd.hpp:95
simd operator++(int)
Postfix increment.
Definition: simd.hpp:140
simd operator--(int)
Postfix decrement.
Definition: simd.hpp:156
simd & operator--()
Prefix decrement, decrements elements of this object.
Definition: simd.hpp:148
#define __esimd_dbg_print(a)
Definition: types.hpp:24
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
std::ostream & operator<<(std::ostream &OS, const sycl::ext::intel::esimd::simd< Ty, N > &V)
Prints a simd object to an output stream.
Definition: simd.hpp:211
std::experimental::simd< T, simd_abi::native_fixed_size< T, N > > simd
std::plus< T > plus
Definition: functional.hpp:18
Definition: access.hpp:18
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
constexpr size_t max_fixed_size
Definition: simd.hpp:1011
#define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID)
Definition: simd.hpp:162