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__
25 #include <iostream>
26 #endif // __SYCL_DEVICE_ONLY__
27 
29 namespace __ESIMD_NS {
30 
33 
38 
40 
43 
53 template <typename Ty, int N>
54 class simd : public detail::simd_obj_impl<
55  detail::__raw_t<Ty>, N, simd<Ty, N>,
56  std::enable_if_t<detail::is_valid_simd_elem_type_v<Ty>>> {
57  using base_type = detail::simd_obj_impl<detail::__raw_t<Ty>, N, simd<Ty, N>>;
58 
59 public:
60  using base_type::base_type;
61  using element_type = Ty;
64  static constexpr int length = N;
65 
73  template <typename SimdT,
74  class = std::enable_if_t<__ESIMD_DNS::is_simd_type_v<SimdT> &&
75  (length == SimdT::length)>>
76  simd(const SimdT &RHS)
77  : base_type(detail::convert_vector<Ty, detail::element_type_t<SimdT>, N>(
78  RHS.data())) {
79  __esimd_dbg_print(simd(const SimdT &RHS));
80  }
81 
82  // Implicit conversion constructor from sycl::ext::oneapi::experimental::simd
83  template <
84  int N1 = N, class Ty1 = Ty,
85  class SFINAE = std::enable_if_t<
87  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
89  : simd(static_cast<raw_vector_type>(v)) {}
90 
96  template <typename T1,
97  class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
98  simd(T1 Val) : base_type(Val) {
99  __esimd_dbg_print(simd(T1 Val));
100  }
101 
107  template <class To, class T = simd,
109  (T::length == 1) && detail::is_valid_simd_elem_type_v<To>>>
110  operator To() const {
111  __esimd_dbg_print(operator To());
112  return detail::convert_scalar<To, element_type>(base_type::data()[0]);
113  }
114 
119  template <
120  int N1, class Ty1 = Ty,
121  class SFINAE = std::enable_if_t<
123  Ty>)&&!detail::is_wrapper_elem_type_v<Ty1>>>
125  return sycl::ext::oneapi::experimental::simd<Ty, N1>(base_type::data());
126  }
127 
131  *this += 1;
132  return *this;
133  }
134 
139  simd Ret(base_type::data());
140  operator++();
141  return Ret;
142  }
143 
147  *this -= 1;
148  return *this;
149  }
150 
155  simd Ret(base_type::data());
156  operator--();
157  return Ret;
158  }
159 
160 #define __ESIMD_DEF_SIMD_ARITH_UNARY_OP(ARITH_UNARY_OP, ID) \
161  template <class T1 = Ty> simd operator ARITH_UNARY_OP() const { \
162  static_assert(!std::is_unsigned_v<T1>, \
163  #ARITH_UNARY_OP "doesn't apply to unsigned types"); \
164  return simd{detail::vector_unary_op<detail::UnaryOp::ID, T1, N>( \
165  base_type::data())}; \
166  }
167 
172 #undef __ESIMD_DEF_SIMD_ARITH_UNARY_OP
173 };
174 
176 
179 
182 template <typename To, typename From, int N>
183 ESIMD_INLINE simd<To, N> convert(const simd<From, N> &val) {
184  if constexpr (std::is_same_v<To, From>)
185  return val;
186  else
187  return detail::convert_vector<To, From, N>(val.data());
188 }
190 
193 
196 template <int N> using simd_mask = detail::simd_mask_type<N>;
197 
199 
200 } // namespace __ESIMD_NS
201 } // __SYCL_INLINE_NAMESPACE(cl)
202 
206 template <typename Ty, int N>
207 std::ostream &operator<<(std::ostream &OS, const __ESIMD_NS::simd<Ty, N> &V)
208 #ifdef __SYCL_DEVICE_ONLY__
209  {}
210 #else
211 {
212  OS << "{";
213  for (int I = 0; I < N; I++) {
214  OS << V[I];
215  if (I < N - 1)
216  OS << ",";
217  }
218  OS << "}";
219  return OS;
220 }
221 #endif // __SYCL_DEVICE_ONLY__
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:160
T
simd_view.hpp
max_fixed_size
constexpr size_t max_fixed_size
Definition: simd.hpp:1008
cl::sycl::ext::intel::esimd::simd::simd
simd(const sycl::ext::oneapi::experimental::simd< Ty, N1 > &v)
Definition: simd.hpp:88
intrin.hpp
cl::sycl::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:174
cl::sycl::ext::intel::esimd::simd::operator++
simd & operator++()
Prefix increment, increments elements of this object.
Definition: simd.hpp:130
cl::sycl::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
cl::sycl::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:177
cl::sycl::ext::intel::esimd::simd::operator--
simd operator--(int)
Postfix decrement.
Definition: simd.hpp:154
cl::sycl::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:31
simd_obj_impl.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
types.hpp
simd_mask_impl.hpp
cl::sycl::ext::intel::esimd::simd::simd
simd(T1 Val)
Broadcast constructor with conversion.
Definition: simd.hpp:98
__esimd_dbg_print
#define __esimd_dbg_print(a)
Definition: types.hpp:22
simd
Definition: simd.hpp:1027
sycl::ext::oneapi::experimental::simd
std::experimental::simd< T, simd_abi::native_fixed_size< T, N > > simd
Definition: invoke_simd.hpp:85
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
sycl_util.hpp
cl::sycl::ext::intel::esimd::simd::simd
simd(const SimdT &RHS)
Implicit conversion constructor from another simd object of the same length.
Definition: simd.hpp:76
cl::sycl::ext::intel::esimd::simd::element_type
Ty element_type
Definition: simd.hpp:61
cl::sycl::plus
std::plus< T > plus
Definition: functional.hpp:18
invoke_simd.hpp
memory_intrin.hpp
cl::sycl::ext::intel::esimd::simd::raw_element_type
typename base_type::raw_element_type raw_element_type
Definition: simd.hpp:62
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::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:154
cl::sycl::ext::intel::esimd::simd::operator--
simd & operator--()
Prefix decrement, decrements elements of this object.
Definition: simd.hpp:146
operator<<
std::ostream & operator<<(std::ostream &OS, const sycl::ext::intel::esimd::simd< Ty, N > &V)
Definition: simd.hpp:207
cl::sycl::ext::intel::esimd::simd::operator++
simd operator++(int)
Postfix increment.
Definition: simd.hpp:138
cl::sycl::ext::intel::esimd::simd::raw_vector_type
typename base_type::raw_vector_type raw_vector_type
Definition: simd.hpp:63
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12