DPC++ Runtime
Runtime libraries for oneAPI DPC++
simd_view.hpp
Go to the documentation of this file.
1 //==------------ - simd_view.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 view APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
15 
17 namespace __ESIMD_NS {
18 
21 
33 template <typename BaseTy,
34  typename RegionTy =
35  region1d_t<typename BaseTy::element_type, BaseTy::length, 1>>
36 class simd_view : public detail::simd_view_impl<BaseTy, RegionTy> {
38  template <typename, int, class, class> friend class detail::simd_obj_impl;
39  template <typename, int> friend class detail::simd_mask_impl;
40  template <typename, typename> friend class simd_view;
41  template <typename, int> friend class simd;
42  template <typename, typename> friend class detail::simd_view_impl;
43 
44 protected:
45  using BaseClass = detail::simd_view_impl<BaseTy, RegionTy>;
46  // Deduce the corresponding value type from its region type.
47  using ShapeTy = typename shape_type<RegionTy>::type;
48  using base_type = BaseTy;
49  template <typename ElT, int N>
50  using get_simd_t = typename BaseClass::template get_simd_t<ElT, N>;
52 
53 public:
54  static_assert(detail::is_simd_obj_impl_derivative_v<BaseTy>);
55 
56  static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
57 
59  using region_type = RegionTy;
60 
63  using element_type = typename ShapeTy::element_type;
64 
66  using value_type = get_simd_t<element_type, length>;
67 
69  using raw_vector_type =
70  detail::vector_type_t<detail::__raw_t<element_type>, length>;
71 
72 protected:
74  simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
75  simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}
77 
78 public:
80  simd_view(const simd_view &Other) = default;
81  simd_view(simd_view &&Other) = default;
82 
85  simd_view(BaseTy &Base) : BaseClass(Base) {}
86 
88  simd_view &operator=(const simd_view &Other) {
89  BaseClass::operator=(Other);
90  return *this;
91  }
92 
93  using BaseClass::operator--;
94  using BaseClass::operator++;
95  using BaseClass::operator=;
96 };
97 
98 #define __ESIMD_DEF_SCALAR_SIMD_VIEW_RELOP(RELOP) \
99  /* simd_view RELOP simd_view */ \
100  ESIMD_INLINE friend bool operator RELOP(const simd_view &X, \
101  const simd_view &Y) { \
102  return (element_type)X RELOP(element_type) Y; \
103  } \
104  \
105  /* simd_view RELOP SCALAR */ \
106  template <typename T1, \
107  std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>> \
108  ESIMD_INLINE friend bool operator RELOP(const simd_view &X, T1 Y) { \
109  return (element_type)X RELOP Y; \
110  } \
111  \
112  /* SCALAR RELOP simd_view */ \
113  template <typename T1, \
114  std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>> \
115  ESIMD_INLINE friend bool operator RELOP(T1 X, const simd_view &Y) { \
116  return X RELOP(element_type) Y; \
117  }
118 
130 template <typename BaseTy, class ViewedElemT>
131 class simd_view<BaseTy, region1d_scalar_t<ViewedElemT>>
132  : public detail::simd_view_impl<BaseTy, region1d_scalar_t<ViewedElemT>> {
133  template <typename, int, class, class> friend class detail::simd_obj_impl;
134  template <typename, typename> friend class detail::simd_view_impl;
135 
136 public:
137  using RegionTy = region1d_scalar_t<ViewedElemT>;
138  using BaseClass = detail::simd_view_impl<BaseTy, RegionTy>;
139  using ShapeTy = typename shape_type<RegionTy>::type;
140  static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
141  static_assert(1 == length, "length of this view is not equal to 1");
142  static_assert(std::is_same_v<typename ShapeTy::element_type, ViewedElemT>);
145  using element_type = ViewedElemT;
146  using base_type = BaseTy;
147  template <typename ElT, int N>
148  using get_simd_t = typename BaseClass::template get_simd_t<ElT, N>;
151 
152 private:
153  simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
154  simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}
155 
156 public:
158  simd_view(BaseTy &Base) : BaseClass(Base) {}
159 
160  operator element_type() const {
161  const auto v = BaseClass::read().data();
162  return detail::bitcast_to_wrapper_type<element_type>(std::move(v)[0]);
163  }
164 
165  using BaseClass::operator--;
166  using BaseClass::operator++;
167  using BaseClass::operator=;
168 
175 };
176 
177 // TODO: remove code duplication in two class specializations for a simd_view
178 // with a single element
179 
185 template <typename BaseTy, typename NestedRegion, class ViewedElemT>
186 class simd_view<BaseTy, std::pair<region1d_scalar_t<ViewedElemT>, NestedRegion>>
187  : public detail::simd_view_impl<
188  BaseTy, std::pair<region1d_scalar_t<ViewedElemT>, NestedRegion>> {
189  template <typename, int> friend class simd;
190  template <typename, typename> friend class detail::simd_view_impl;
191 
192 public:
193  using RegionTy = std::pair<region1d_scalar_t<ViewedElemT>, NestedRegion>;
194  using BaseClass = detail::simd_view_impl<BaseTy, RegionTy>;
195  using ShapeTy = typename shape_type<RegionTy>::type;
196  static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
197  static_assert(1 == length, "length of this view is not equal to 1");
198  static_assert(std::is_same_v<typename ShapeTy::element_type, ViewedElemT>);
201  using element_type = ViewedElemT;
202 
203 private:
204  simd_view(BaseTy &Base, RegionTy Region) : BaseClass(Base, Region) {}
205  simd_view(BaseTy &&Base, RegionTy Region) : BaseClass(Base, Region) {}
206 
207 public:
208  using BaseClass::operator=;
209 
210  operator element_type() const {
211  const auto v = BaseClass::read();
212  return detail::convert_scalar<element_type>(v[0]);
213  }
214 
221 };
222 
223 #undef __ESIMD_DEF_SCALAR_SIMD_VIEW_RELOP
224 
226 
227 } // namespace __ESIMD_NS
228 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::intel::esimd::simd_view::region_type
RegionTy region_type
The region type of this class.
Definition: simd_view.hpp:59
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::base_type
BaseTy base_type
Definition: simd_view.hpp:146
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::RegionTy
region1d_scalar_t< ViewedElemT > RegionTy
Definition: simd_view.hpp:137
cl::sycl::ext::intel::esimd::simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::ShapeTy
typename shape_type< RegionTy >::type ShapeTy
Definition: simd_view.hpp:195
cl::sycl::ext::intel::esimd::simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::BaseClass
detail::simd_view_impl< BaseTy, RegionTy > BaseClass
Definition: simd_view.hpp:194
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::ShapeTy
typename shape_type< RegionTy >::type ShapeTy
Definition: simd_view.hpp:139
cl::sycl::ext::intel::esimd::simd_view::raw_vector_type
detail::vector_type_t< detail::__raw_t< element_type >, length > raw_vector_type
The underlying builtin value type.
Definition: simd_view.hpp:70
cl::sycl::ext::intel::esimd::simd_view::value_type
get_simd_t< element_type, length > value_type
The simd type if reading the object.
Definition: simd_view.hpp:66
cl::sycl::ext::intel::esimd::simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::element_type
ViewedElemT 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:201
cl::sycl::ext::intel::esimd::simd_view::operator=
simd_view & operator=(const simd_view &Other)
Copy assignment operator.
Definition: simd_view.hpp:88
cl::sycl::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
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
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::simd_view
simd_view(BaseTy &Base)
Construct a complete view of a vector.
Definition: simd_view.hpp:158
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::element_type
ViewedElemT 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:145
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__ESIMD_DEF_SCALAR_SIMD_VIEW_RELOP
#define __ESIMD_DEF_SCALAR_SIMD_VIEW_RELOP(RELOP)
Definition: simd_view.hpp:98
types.hpp
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::value_type
get_simd_t< element_type, length > value_type
The simd type if reading the object.
Definition: simd_view.hpp:150
simd_view_impl.hpp
cl::sycl::ext::intel::esimd::simd_view::simd_view
simd_view(BaseTy &Base)
Construct a complete view of a vector.
Definition: simd_view.hpp:85
simd
Definition: simd.hpp:1027
cl::sycl::ext::intel::esimd::simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::RegionTy
std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > RegionTy
Definition: simd_view.hpp:193
std
Definition: accessor.hpp:2617
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::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::get_simd_t
typename BaseClass::template get_simd_t< ElT, N > get_simd_t
Definition: simd_view.hpp:148
cl::sycl::ext::intel::esimd::simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >::BaseClass
detail::simd_view_impl< BaseTy, RegionTy > BaseClass
Definition: simd_view.hpp:138
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
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12