DPC++ Runtime
Runtime libraries for oneAPI DPC++
simd_view_impl.hpp
Go to the documentation of this file.
1 //==------------ - simd_view_impl.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 // Implementation detail of Explicit SIMD vector view class.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
16 
17 namespace sycl {
19 namespace ext::intel::esimd::detail {
20 
23 
38 template <typename BaseTy,
39  typename RegionTy =
40  region1d_t<typename BaseTy::element_type, BaseTy::length, 1>>
42 public:
46 
47 protected:
49 
50  template <typename, int, class, class> friend class simd_obj_impl;
51  template <typename, int> friend class simd;
52  template <typename, typename> friend class simd_view_impl;
53  template <typename, int> friend class simd_mask_impl;
54 
55  static_assert(is_simd_obj_impl_derivative_v<BaseTy>);
56 
57 protected:
58  // Deduce the corresponding value type from its region type.
59  using ShapeTy = typename shape_type<RegionTy>::type;
60  static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
61 
62  using base_type = BaseTy;
63  template <typename ElT, int N>
64  using get_simd_t = construct_a_simd_type_t<base_type, ElT, N>;
65 
67  using region_type = RegionTy;
68 
70 
71 public:
74  using element_type = typename ShapeTy::element_type;
75 
77  using raw_element_type = __raw_t<element_type>;
78 
80  using value_type = get_simd_t<element_type, length>;
81 
83  using raw_vector_type = vector_type_t<__raw_t<element_type>, length>;
84 
85 private:
87 
88  Derived &cast_this_to_derived() { return reinterpret_cast<Derived &>(*this); }
89 
90 protected:
91  simd_view_impl(BaseTy &Base, RegionTy Region)
92  : M_base(Base), M_region(Region) {}
93 
94  simd_view_impl(BaseTy &Base) : M_base(Base), M_region(RegionTy(0)) {}
95 
97 
98 public:
100  simd_view_impl(const simd_view_impl &Other) = default;
101 
103  simd_view_impl(simd_view_impl &&Other) = default;
104 
110  template <typename ToTy, class T = BaseTy,
111  class = std::enable_if_t<is_simd_type_v<T>>>
112  inline operator simd<ToTy, length>() const {
113  if constexpr (std::is_same_v<element_type, ToTy>)
114  return read();
115  else
116  return convert_vector<ToTy, element_type, length>(read().data());
117  }
118 
122  template <class T = BaseTy, class = std::enable_if_t<is_simd_mask_type_v<T>>>
123  inline operator simd_mask_type<length>() const {
124  return read();
125  }
126 
128  static constexpr bool is1D() { return !ShapeTy::Is_2D; }
130  static constexpr bool is2D() { return ShapeTy::Is_2D; }
132  static constexpr int getSizeX() { return ShapeTy::Size_x; }
134  static constexpr int getStrideX() { return ShapeTy::Stride_x; }
136  static constexpr int getSizeY() { return ShapeTy::Size_y; }
138  static constexpr int getStrideY() { return ShapeTy::Stride_y; }
139 
142  constexpr uint16_t getOffsetX() const {
143  return getTopRegion(M_region).M_offset_x;
144  }
145 
148  constexpr uint16_t getOffsetY() const {
149  return getTopRegion(M_region).M_offset_y;
150  }
151 
155  value_type read() const {
156  using BT = typename BaseTy::element_type;
157  constexpr int BN = BaseTy::length;
158  return value_type{readRegion<BT, BN>(M_base.data(), M_region)};
159  }
160 
162  typename value_type::raw_vector_type data() const { return read().data(); }
163 
168  Derived &write(const value_type &Val) {
169  M_base.writeRegion(M_region, Val.data());
170  return cast_this_to_derived();
171  }
172 
178  void merge(const value_type &Val, const simd_mask_type<length> &Mask) {
179  merge(Val, read(), Mask);
180  }
181 
188  void merge(const value_type &Val1, value_type Val2,
189  const simd_mask_type<length> &Mask) {
190  Val2.merge(Val1, Mask);
191  write(Val2.read());
192  }
193 
199  template <typename EltTy> auto bit_cast_view() {
200  using TopRegionTy = detail::compute_format_type_t<Derived, EltTy>;
201  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
202  using RetTy = simd_view<BaseTy, NewRegionTy>;
203  TopRegionTy TopReg(0);
204  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
205  }
206 
217  template <typename EltTy, int Height, int Width> auto bit_cast_view() {
218  using TopRegionTy =
219  detail::compute_format_type_2d_t<Derived, EltTy, Height, Width>;
220  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
221  using RetTy = simd_view<BaseTy, NewRegionTy>;
222  TopRegionTy TopReg(0, 0);
223  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
224  }
225 
232  template <int Size, int Stride, typename T = Derived,
233  typename = sycl::detail::enable_if_t<T::is1D()>>
234  auto select(uint16_t Offset = 0) {
235  using TopRegionTy = region1d_t<element_type, Size, Stride>;
236  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
237  using RetTy = simd_view<BaseTy, NewRegionTy>;
238  TopRegionTy TopReg(Offset);
239  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
240  }
241 
242  // clang-format off
281  // clang-format on
282  template <int SizeY, int StrideY, int SizeX, int StrideX,
283  typename T = Derived,
284  typename = sycl::detail::enable_if_t<T::is2D()>>
285  auto select(uint16_t OffsetY = 0, uint16_t OffsetX = 0) {
286  using TopRegionTy =
287  region2d_t<element_type, SizeY, StrideY, SizeX, StrideX>;
288  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
289  using RetTy = simd_view<BaseTy, NewRegionTy>;
290  TopRegionTy TopReg(OffsetY, OffsetX);
291  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
292  }
293 #define __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN(BINOP, OPASSIGN, COND) \
294  \
295  /* OPASSIGN simd_obj_impl */ \
296  template <class T1, int N1, class SimdT1, class T = element_type, \
297  class SimdT = BaseTy, \
298  class = \
299  std::enable_if_t<(is_simd_type_v<SimdT> == \
300  is_simd_type_v<SimdT1>)&&(N1 == length) && \
301  COND>> \
302  Derived &operator OPASSIGN(const simd_obj_impl<T1, N1, SimdT1> &RHS) { \
303  auto Res = read() BINOP RHS; \
304  write(Res); \
305  return cast_this_to_derived(); \
306  } \
307  \
308  /* OPASSIGN simd_view_impl */ \
309  template <class SimdT1, class RegionT1, \
310  class T1 = \
311  typename __ESIMD_NS::shape_type<RegionT1>::element_type, \
312  class T = element_type, class SimdT = BaseTy, \
313  class = std::enable_if_t< \
314  (is_simd_type_v<SimdT> == is_simd_type_v<SimdT1>)&&( \
315  length == __ESIMD_NS::shape_type<RegionT1>::length) && \
316  COND>> \
317  Derived &operator OPASSIGN(const simd_view_impl<SimdT1, RegionT1> &RHS) { \
318  *this OPASSIGN RHS.read(); \
319  return cast_this_to_derived(); \
320  } \
321  \
322  /* OPASSIGN scalar */ \
323  template <class T1, class T = element_type, class SimdT = BaseTy, \
324  class = std::enable_if_t<COND>> \
325  Derived &operator OPASSIGN(T1 RHS) { \
326  auto Res = read() BINOP RHS; \
327  write(Res); \
328  return cast_this_to_derived(); \
329  }
330 
331 #define __ESIMD_BITWISE_OP_FILTER std::is_integral_v<T> &&std::is_integral_v<T1>
336 #undef __ESIMD_BITWISE_OP_FILTER
337 
338 #define __ESIMD_SHIFT_OP_FILTER \
339  std::is_integral_v<T> &&std::is_integral_v<T1> &&is_simd_type_v<SimdT>
340 
343 #undef __ESIMD_SHIFT_OP_FILTER
344 
345 #define __ESIMD_ARITH_OP_FILTER \
346  is_valid_simd_elem_type_v<T> &&is_valid_simd_elem_type_v<T1> \
347  &&is_simd_type_v<SimdT>
348 
353 
354 #undef __ESIMD_ARITH_OP_FILTER
355 #undef __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN
356 
357 #define __ESIMD_DEF_UNARY_OP(UNARY_OP, COND) \
358  template <class T = element_type, class SimdT = BaseTy, \
359  class = std::enable_if_t<COND>> \
360  auto operator UNARY_OP() { \
361  auto V = UNARY_OP(read().data()); \
362  return get_simd_t<element_type, length>(V); \
363  }
364  __ESIMD_DEF_UNARY_OP(~, std::is_integral_v<T> &&is_simd_type_v<SimdT>)
365  __ESIMD_DEF_UNARY_OP(+, is_simd_type_v<SimdT>)
366  __ESIMD_DEF_UNARY_OP(-, is_simd_type_v<SimdT>)
367 
368 #undef __ESIMD_DEF_UNARY_OP
369 
371  template <class T = element_type,
372  class = std::enable_if_t<std::is_integral_v<T>>>
373  auto operator!() {
374  using MaskVecT = typename simd_mask_type<length>::raw_vector_type;
375  auto V = read().data() == 0;
376  return simd_mask_type<length>{__builtin_convertvector(V, MaskVecT) &
377  MaskVecT(1)};
378  }
379 
383  return write(Other.read());
384  }
385 
390  Derived &operator=(const Derived &Other) { return write(Other.read()); }
391 
396  Derived &operator=(const value_type &Val) { return write(Val); }
397 
403  __esimd_move_test_proxy(Other);
404  return write(Other.read());
405  }
406 
409  __esimd_move_test_proxy(Other);
410  return write(Other.read());
411  }
412 
421  template <class T, int N, class SimdT,
422  class = std::enable_if_t<(is_simd_type_v<SimdT> ==
423  is_simd_type_v<BaseTy>)&&(length ==
424  SimdT::length)>>
426  return write(convert_vector<element_type, typename SimdT::element_type, N>(
427  Other.data()));
428  }
429 
435  template <class T1, class = std::enable_if_t<is_valid_simd_elem_type_v<T1>>>
436  Derived &operator=(T1 RHS) {
437  return write(value_type(convert_scalar<element_type>(RHS)));
438  }
439 
443  *this += 1;
444  return cast_this_to_derived();
445  }
446 
451  value_type Ret(read());
452  operator++();
453  return Ret;
454  }
455 
459  *this -= 1;
460  return cast_this_to_derived();
461  }
462 
467  value_type Ret(read());
468  operator--();
469  return Ret;
470  }
471 
475  template <typename T = Derived,
476  typename = sycl::detail::enable_if_t<T::is2D()>>
477  auto row(int i) {
478  return select<1, 1, getSizeX(), 1>(i, 0)
479  .template bit_cast_view<element_type>();
480  }
481 
486  template <typename T = Derived,
487  typename = sycl::detail::enable_if_t<T::is2D()>>
488  auto column(int i) {
489  return select<getSizeY(), 1, 1, 1>(0, i);
490  }
491 
495  template <typename T = Derived,
496  typename = sycl::detail::enable_if_t<T::is1D()>>
497  element_type operator[](int i) const {
498  const auto v = read();
499  return v[i];
500  }
501 
505  template <typename T = Derived,
506  typename = sycl::detail::enable_if_t<T::is1D()>>
507  auto operator[](int i) {
508  return select<1, 1>(i);
509  }
510 
512  template <int Rep> get_simd_t<element_type, Rep> replicate() {
513  return read().template replicate<Rep>();
514  }
515 
519  template <int Rep, int W>
520  get_simd_t<element_type, Rep * W> replicate_w(uint16_t OffsetX) {
521  return replicate_vs_w<Rep, 0, W>(0, OffsetX);
522  }
523 
527  template <int Rep, int W>
528  get_simd_t<element_type, Rep * W> replicate_w(uint16_t OffsetY,
529  uint16_t OffsetX) {
530  return replicate_vs_w<Rep, 0, W>(OffsetY, OffsetX);
531  }
532 
533  // clang-format off
537  // clang-format on
538  template <int Rep, int VS, int W>
539  get_simd_t<element_type, Rep * W> replicate_vs_w(uint16_t OffsetX) {
540  return replicate_vs_w_hs<Rep, VS, W, 1>(0, OffsetX);
541  }
542 
543  // clang-format off
547  // clang-format on
548  template <int Rep, int VS, int W>
549  get_simd_t<element_type, Rep * W> replicate_vs_w(uint16_t OffsetY,
550  uint16_t OffsetX) {
551  return replicate_vs_w_hs<Rep, VS, W, 1>(OffsetY, OffsetX);
552  }
553 
556  template <int Rep, int VS, int W, int HS>
557  get_simd_t<element_type, Rep * W> replicate_vs_w_hs(uint16_t OffsetX) {
558  return read().template replicate_vs_w_hs<Rep, VS, W, HS>(OffsetX);
559  }
560 
565  template <int Rep, int VS, int W, int HS>
566  get_simd_t<element_type, Rep * W> replicate_vs_w_hs(uint16_t OffsetY,
567  uint16_t OffsetX) {
568  constexpr int RowSize = is2D() ? getSizeX() : 0;
569  return read().template replicate_vs_w_hs<Rep, VS, W, HS>(OffsetY * RowSize +
570  OffsetX);
571  }
572 
574  template <typename T1 = element_type, typename T2 = BaseTy,
575  typename = std::enable_if_t<std::is_integral<T1>::value, T2>>
576  uint16_t any() {
577  return read().any();
578  }
579 
581  template <typename T1 = element_type, typename T2 = BaseTy,
582  typename = std::enable_if_t<std::is_integral<T1>::value, T2>>
583  uint16_t all() {
584  return read().all();
585  }
586 
588 public:
589  // Getter for the test proxy member, if enabled
590  __ESIMD_DECLARE_TEST_PROXY_ACCESS
591 
592 protected:
593  // The reference to the base object, which must be a simd object
594  BaseTy &M_base;
595 
596  // The test proxy if enabled
597  __ESIMD_DECLARE_TEST_PROXY
598 
599  // The region applied on the base object. Its type could be
600  // - region1d_t
601  // - region2d_t
602  // - std::pair<top_region_type, base_region_type>
603  //
604  RegionTy M_region;
606 };
607 
609 
610 } // namespace ext::intel::esimd::detail
611 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
612 } // namespace sycl
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator--
value_type operator--(int)
Postfix decrement.
Definition: simd_view_impl.hpp:466
sycl::_V1::ext::intel::esimd::detail::simd_view_impl< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::element_type
typename ShapeTy::element_type element_type
Element type of this view, may differ from the element type of the target object.
Definition: simd_view_impl.hpp:74
__ESIMD_BITWISE_OP_FILTER
#define __ESIMD_BITWISE_OP_FILTER
Definition: simd_view_impl.hpp:331
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
Derived & operator=(Derived &&Other)
Assignment from an rvalue object of the derived class.
Definition: simd_view_impl.hpp:402
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::select
auto select(uint16_t OffsetY=0, uint16_t OffsetX=0)
2D region select.
Definition: simd_view_impl.hpp:285
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::is1D
static constexpr bool is1D()
Tells whether this view is 1-dimensional.
Definition: simd_view_impl.hpp:128
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getStrideX
static constexpr int getStrideX()
Get element stride of the view along X dimension.
Definition: simd_view_impl.hpp:134
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
simd_view_impl & operator=(const simd_view_impl &Other)
Copy assignment.
Definition: simd_view_impl.hpp:382
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::any
uint16_t any()
Applies simd_obj_impl::any operation to the target region.
Definition: simd_view_impl.hpp:576
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::merge
void merge(const value_type &Val, const simd_mask_type< length > &Mask)
"Merges" the viewed target's region with given value according to a per-element mask.
Definition: simd_view_impl.hpp:178
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::write
Derived & write(const value_type &Val)
Assigns a new value to the viewed target's region.
Definition: simd_view_impl.hpp:168
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_vs_w
get_simd_t< element_type, Rep *W > replicate_vs_w(uint16_t OffsetY, uint16_t OffsetX)
Shortcut to replicate_vs_w_hs<int Rep, int VS, int W, int Hs>(uint16_t OffsetY, uint16_t OffsetX) wit...
Definition: simd_view_impl.hpp:549
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::merge
void merge(const value_type &Val1, value_type Val2, const simd_mask_type< length > &Mask)
"Merges" given values according to a per-element mask and writes the result to the viewed target's re...
Definition: simd_view_impl.hpp:188
sycl::_V1::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:34
__ESIMD_ARITH_OP_FILTER
#define __ESIMD_ARITH_OP_FILTER
Definition: simd_view_impl.hpp:345
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::all
uint16_t all()
Applies simd_obj_impl::all operation to the target region.
Definition: simd_view_impl.hpp:583
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_view
This class represents a reference to a sub-region of a base simd object.
Definition: types.hpp:35
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
sycl::_V1::ext::intel::esimd::detail::simd_view_impl< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::raw_vector_type
vector_type_t< __raw_t< element_type >, length > raw_vector_type
The underlying builtin vector type of the the viewed region.
Definition: simd_view_impl.hpp:83
sycl::_V1::ext::intel::esimd::detail::simd_obj_impl::data
raw_vector_type data() const
Definition: simd_obj_impl.hpp:339
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator[]
element_type operator[](int i) const
Read a single element from the target 1D region.
Definition: simd_view_impl.hpp:497
__ESIMD_DEF_UNARY_OP
#define __ESIMD_DEF_UNARY_OP(UNARY_OP, COND)
Definition: simd_view_impl.hpp:357
intrin.hpp
test_proxy.hpp
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator++
Derived & operator++()
Prefix increment.
Definition: simd_view_impl.hpp:442
sycl::_V1::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:109
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getSizeX
static constexpr int getSizeX()
Get number of elements in the view along X dimension.
Definition: simd_view_impl.hpp:132
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate
get_simd_t< element_type, Rep > replicate()
Applies simd_obj_impl::replicate to the target region.
Definition: simd_view_impl.hpp:512
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
Derived & operator=(T1 RHS)
Broadcast assignment of a scalar with conversion.
Definition: simd_view_impl.hpp:436
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::column
auto column(int i)
Reference a column from a 2D region.
Definition: simd_view_impl.hpp:488
sycl::_V1::ext::intel::esimd::detail::simd_view_impl< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::raw_element_type
__raw_t< element_type > raw_element_type
Corresponding "raw" (storage) type for the element type.
Definition: simd_view_impl.hpp:77
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getSizeY
static constexpr int getSizeY()
Get number of elements in the view along Y dimension.
Definition: simd_view_impl.hpp:136
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::bit_cast_view
auto bit_cast_view()
Create a 2-dimensional view of the target region.
Definition: simd_view_impl.hpp:217
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_vs_w_hs
get_simd_t< element_type, Rep *W > replicate_vs_w_hs(uint16_t OffsetX)
Applies simd_obj_impl::replicate_vs_w_hs<int Rep, int VS, int W, int HS> to the target region.
Definition: simd_view_impl.hpp:557
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::row
auto row(int i)
Reference a row from a 2D region.
Definition: simd_view_impl.hpp:477
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getOffsetY
constexpr uint16_t getOffsetY() const
Get the offset of the first element of the view within the parent object along Y dimension.
Definition: simd_view_impl.hpp:148
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_vs_w_hs
get_simd_t< element_type, Rep *W > replicate_vs_w_hs(uint16_t OffsetY, uint16_t OffsetX)
Applies simd_obj_impl::replicate_vs_w_hs<int Rep, int VS, int W, int HS> to the target region.
Definition: simd_view_impl.hpp:566
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
Derived & operator=(const Derived &Other)
Assignment from an object of the derived class.
Definition: simd_view_impl.hpp:390
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_vs_w
get_simd_t< element_type, Rep *W > replicate_vs_w(uint16_t OffsetX)
Shortcut to replicate_vs_w_hs<int Rep, int VS, int W, int Hs>(uint16_t OffsetY, uint16_t OffsetX) wit...
Definition: simd_view_impl.hpp:539
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_w
get_simd_t< element_type, Rep *W > replicate_w(uint16_t OffsetX)
Shortcut to replicate_vs_w<int Rep, int Vs, int W>(uint16_t) with with Vs = 0.
Definition: simd_view_impl.hpp:520
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator[]
auto operator[](int i)
Return a writeable view of a single element in the target 1D region.
Definition: simd_view_impl.hpp:507
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator!
auto operator!()
Unary logical negeation operator. Applies only to integer element types.
Definition: simd_view_impl.hpp:373
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::data
value_type::raw_vector_type data() const
Definition: simd_view_impl.hpp:162
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
Derived & operator=(const simd_obj_impl< T, N, SimdT > &Other)
Assignment with element type conversion from a vector object.
Definition: simd_view_impl.hpp:425
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::bit_cast_view
auto bit_cast_view()
Create a 1-dimensional view of the target region.
Definition: simd_view_impl.hpp:199
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getOffsetX
constexpr uint16_t getOffsetX() const
Get the offset of the first element of the view within the parent object along X dimension.
Definition: simd_view_impl.hpp:142
sycl::_V1::ext::intel::esimd::detail::simd_view_impl< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >::value_type
get_simd_t< element_type, length > value_type
The simd type of the viewed region of the target object.
Definition: simd_view_impl.hpp:80
__ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN
#define __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN(BINOP, OPASSIGN, COND)
Definition: simd_view_impl.hpp:293
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::replicate_w
get_simd_t< element_type, Rep *W > replicate_w(uint16_t OffsetY, uint16_t OffsetX)
Shortcut to replicate_vs_w<int Rep, int Vs, int W>(uint16_t, uint16_t) with Vs = 0.
Definition: simd_view_impl.hpp:528
sycl::_V1::ext::intel::esimd::simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >
This is a specialization of nested simd_view class with a single element.
Definition: simd_view.hpp:187
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
simd_view_impl & operator=(simd_view_impl &&Other)
Move assignment operator. Updates the target region viewed by this object.
Definition: simd_view_impl.hpp:408
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::is2D
static constexpr bool is2D()
Tells whether this view is 2-dimensional.
Definition: simd_view_impl.hpp:130
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator--
Derived & operator--()
Prefix decrement.
Definition: simd_view_impl.hpp:458
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator=
Derived & operator=(const value_type &Val)
Assignment from a value_type object - simd or simd_mask.
Definition: simd_view_impl.hpp:396
__ESIMD_SHIFT_OP_FILTER
#define __ESIMD_SHIFT_OP_FILTER
Definition: simd_view_impl.hpp:338
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::getStrideY
static constexpr int getStrideY()
Get element stride of the view along Y dimension.
Definition: simd_view_impl.hpp:138
sycl::_V1::ext::intel::esimd::detail::simd_mask_impl
This class is a simd_obj_impl specialization representing a simd mask, which is basically a simd_obj_...
Definition: simd_mask_impl.hpp:56
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::read
value_type read() const
Reads the viewed region from the target w/o any conversion and returns as an object of the value_type...
Definition: simd_view_impl.hpp:155
sycl::_V1::ext::intel::esimd::merge
__ESIMD_API simd< T, N > merge(simd< T, N > a, simd< T, N > b, simd_mask< N > m)
"Merges" elements of the input simd object according to the merge mask.
Definition: alt_ui.hpp:28
type_format.hpp
sycl::_V1::ext::intel::esimd::detail::simd_view_impl
Base class for "simd view" types.
Definition: simd_view_impl.hpp:41
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::operator++
value_type operator++(int)
Postfix increment.
Definition: simd_view_impl.hpp:450
sycl::_V1::ext::intel::esimd::detail::simd_view_impl::select
auto select(uint16_t Offset=0)
1D region select.
Definition: simd_view_impl.hpp:234