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 
18 namespace __ESIMD_DNS {
19 
22 
37 template <typename BaseTy,
38  typename RegionTy =
39  region1d_t<typename BaseTy::element_type, BaseTy::length, 1>>
41 public:
44  using Derived = simd_view<BaseTy, RegionTy>;
45 
46 protected:
48 
49  template <typename, int, class, class> friend class simd_obj_impl;
50  template <typename, int> friend class simd;
51  template <typename, typename> friend class simd_view_impl;
52  template <typename, int> friend class simd_mask_impl;
53 
54  static_assert(is_simd_obj_impl_derivative_v<BaseTy>);
55 
56 protected:
57  // Deduce the corresponding value type from its region type.
58  using ShapeTy = typename shape_type<RegionTy>::type;
59  static constexpr int length = ShapeTy::Size_x * ShapeTy::Size_y;
60 
61  using base_type = BaseTy;
62  template <typename ElT, int N>
63  using get_simd_t = construct_a_simd_type_t<base_type, ElT, N>;
64 
66  using region_type = RegionTy;
67 
69 
70 public:
73  using element_type = typename ShapeTy::element_type;
74 
76  using raw_element_type = __raw_t<element_type>;
77 
79  using value_type = get_simd_t<element_type, length>;
80 
82  using raw_vector_type = vector_type_t<__raw_t<element_type>, length>;
83 
84 private:
86 
87  Derived &cast_this_to_derived() { return reinterpret_cast<Derived &>(*this); }
88 
89 protected:
90  simd_view_impl(BaseTy &Base, RegionTy Region)
91  : M_base(Base), M_region(Region) {}
92 
93  simd_view_impl(BaseTy &Base) : M_base(Base), M_region(RegionTy(0)) {}
94 
96 
97 public:
99  simd_view_impl(const simd_view_impl &Other) = default;
100 
102  simd_view_impl(simd_view_impl &&Other) = default;
103 
109  template <typename ToTy, class T = BaseTy,
110  class = std::enable_if_t<is_simd_type_v<T>>>
111  inline operator simd<ToTy, length>() const {
112  if constexpr (std::is_same_v<element_type, ToTy>)
113  return read();
114  else
115  return convert_vector<ToTy, element_type, length>(read().data());
116  }
117 
121  template <class T = BaseTy, class = std::enable_if_t<is_simd_mask_type_v<T>>>
122  inline operator simd_mask_type<length>() const {
123  return read();
124  }
125 
127  static constexpr bool is1D() { return !ShapeTy::Is_2D; }
129  static constexpr bool is2D() { return ShapeTy::Is_2D; }
131  static constexpr int getSizeX() { return ShapeTy::Size_x; }
133  static constexpr int getStrideX() { return ShapeTy::Stride_x; }
135  static constexpr int getSizeY() { return ShapeTy::Size_y; }
137  static constexpr int getStrideY() { return ShapeTy::Stride_y; }
138 
141  constexpr uint16_t getOffsetX() const {
142  return getTopRegion(M_region).M_offset_x;
143  }
144 
147  constexpr uint16_t getOffsetY() const {
148  return getTopRegion(M_region).M_offset_y;
149  }
150 
154  value_type read() const {
155  using BT = typename BaseTy::element_type;
156  constexpr int BN = BaseTy::length;
157  return value_type{readRegion<BT, BN>(M_base.data(), M_region)};
158  }
159 
161  typename value_type::raw_vector_type data() const { return read().data(); }
162 
167  Derived &write(const value_type &Val) {
168  M_base.writeRegion(M_region, Val.data());
169  return cast_this_to_derived();
170  }
171 
177  void merge(const value_type &Val, const simd_mask_type<length> &Mask) {
178  merge(Val, read(), Mask);
179  }
180 
187  void merge(const value_type &Val1, value_type Val2,
188  const simd_mask_type<length> &Mask) {
189  Val2.merge(Val1, Mask);
190  write(Val2.read());
191  }
192 
198  template <typename EltTy> auto bit_cast_view() {
199  using TopRegionTy = detail::compute_format_type_t<Derived, EltTy>;
200  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
201  using RetTy = simd_view<BaseTy, NewRegionTy>;
202  TopRegionTy TopReg(0);
203  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
204  }
205 
216  template <typename EltTy, int Height, int Width> auto bit_cast_view() {
217  using TopRegionTy =
218  detail::compute_format_type_2d_t<Derived, EltTy, Height, Width>;
219  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
220  using RetTy = simd_view<BaseTy, NewRegionTy>;
221  TopRegionTy TopReg(0, 0);
222  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
223  }
224 
231  template <int Size, int Stride, typename T = Derived,
232  typename = sycl::detail::enable_if_t<T::is1D()>>
233  auto select(uint16_t Offset = 0) {
234  using TopRegionTy = region1d_t<element_type, Size, Stride>;
235  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
236  using RetTy = simd_view<BaseTy, NewRegionTy>;
237  TopRegionTy TopReg(Offset);
238  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
239  }
240 
241  // clang-format off
280  // clang-format on
281  template <int SizeY, int StrideY, int SizeX, int StrideX,
282  typename T = Derived,
283  typename = sycl::detail::enable_if_t<T::is2D()>>
284  auto select(uint16_t OffsetY = 0, uint16_t OffsetX = 0) {
285  using TopRegionTy =
286  region2d_t<element_type, SizeY, StrideY, SizeX, StrideX>;
287  using NewRegionTy = std::pair<TopRegionTy, RegionTy>;
288  using RetTy = simd_view<BaseTy, NewRegionTy>;
289  TopRegionTy TopReg(OffsetY, OffsetX);
290  return RetTy{this->M_base, std::make_pair(TopReg, M_region)};
291  }
292 #define __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN(BINOP, OPASSIGN, COND) \
293  \
294  /* OPASSIGN simd_obj_impl */ \
295  template <class T1, int N1, class SimdT1, class T = element_type, \
296  class SimdT = BaseTy, \
297  class = \
298  std::enable_if_t<(is_simd_type_v<SimdT> == \
299  is_simd_type_v<SimdT1>)&&(N1 == length) && \
300  COND>> \
301  Derived &operator OPASSIGN(const simd_obj_impl<T1, N1, SimdT1> &RHS) { \
302  auto Res = read() BINOP RHS; \
303  write(Res); \
304  return cast_this_to_derived(); \
305  } \
306  \
307  /* OPASSIGN simd_view_impl */ \
308  template <class SimdT1, class RegionT1, \
309  class T1 = \
310  typename __ESIMD_NS::shape_type<RegionT1>::element_type, \
311  class T = element_type, class SimdT = BaseTy, \
312  class = std::enable_if_t< \
313  (is_simd_type_v<SimdT> == is_simd_type_v<SimdT1>)&&( \
314  length == __ESIMD_NS::shape_type<RegionT1>::length) && \
315  COND>> \
316  Derived &operator OPASSIGN(const simd_view_impl<SimdT1, RegionT1> &RHS) { \
317  *this OPASSIGN RHS.read(); \
318  return cast_this_to_derived(); \
319  } \
320  \
321  /* OPASSIGN scalar */ \
322  template <class T1, class T = element_type, class SimdT = BaseTy, \
323  class = std::enable_if_t<COND>> \
324  Derived &operator OPASSIGN(T1 RHS) { \
325  auto Res = read() BINOP RHS; \
326  write(Res); \
327  return cast_this_to_derived(); \
328  }
329 
330 #define __ESIMD_BITWISE_OP_FILTER std::is_integral_v<T> &&std::is_integral_v<T1>
335 #undef __ESIMD_BITWISE_OP_FILTER
336 
337 #define __ESIMD_SHIFT_OP_FILTER \
338  std::is_integral_v<T> &&std::is_integral_v<T1> &&is_simd_type_v<SimdT>
339 
342 #undef __ESIMD_SHIFT_OP_FILTER
343 
344 #define __ESIMD_ARITH_OP_FILTER \
345  is_valid_simd_elem_type_v<T> &&is_valid_simd_elem_type_v<T1> \
346  &&is_simd_type_v<SimdT>
347 
352 
353 #undef __ESIMD_ARITH_OP_FILTER
354 #undef __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN
355 
356 #define __ESIMD_DEF_UNARY_OP(UNARY_OP, COND) \
357  template <class T = element_type, class SimdT = BaseTy, \
358  class = std::enable_if_t<COND>> \
359  auto operator UNARY_OP() { \
360  auto V = UNARY_OP(read().data()); \
361  return get_simd_t<element_type, length>(V); \
362  }
363  __ESIMD_DEF_UNARY_OP(~, std::is_integral_v<T> &&is_simd_type_v<SimdT>)
364  __ESIMD_DEF_UNARY_OP(+, is_simd_type_v<SimdT>)
365  __ESIMD_DEF_UNARY_OP(-, is_simd_type_v<SimdT>)
366 
367 #undef __ESIMD_DEF_UNARY_OP
368 
370  template <class T = element_type,
371  class = std::enable_if_t<std::is_integral_v<T>>>
372  auto operator!() {
373  using MaskVecT = typename simd_mask_type<length>::raw_vector_type;
374  auto V = read().data() == 0;
375  return simd_mask_type<length>{__builtin_convertvector(V, MaskVecT) &
376  MaskVecT(1)};
377  }
378 
382  return write(Other.read());
383  }
384 
389  Derived &operator=(const Derived &Other) { return write(Other.read()); }
390 
395  Derived &operator=(const value_type &Val) { return write(Val); }
396 
402  __esimd_move_test_proxy(Other);
403  return write(Other.read());
404  }
405 
408  __esimd_move_test_proxy(Other);
409  return write(Other.read());
410  }
411 
420  template <class T, int N, class SimdT,
421  class = std::enable_if_t<(is_simd_type_v<SimdT> ==
422  is_simd_type_v<BaseTy>)&&(length ==
423  SimdT::length)>>
425  return write(convert_vector<element_type, typename SimdT::element_type, N>(
426  Other.data()));
427  }
428 
434  template <class T1, class = std::enable_if_t<is_valid_simd_elem_type_v<T1>>>
435  Derived &operator=(T1 RHS) {
436  return write(value_type(convert_scalar<element_type>(RHS)));
437  }
438 
442  *this += 1;
443  return cast_this_to_derived();
444  }
445 
450  value_type Ret(read());
451  operator++();
452  return Ret;
453  }
454 
458  *this -= 1;
459  return cast_this_to_derived();
460  }
461 
466  value_type Ret(read());
467  operator--();
468  return Ret;
469  }
470 
474  template <typename T = Derived,
475  typename = sycl::detail::enable_if_t<T::is2D()>>
476  auto row(int i) {
477  return select<1, 1, getSizeX(), 1>(i, 0)
478  .template bit_cast_view<element_type>();
479  }
480 
485  template <typename T = Derived,
486  typename = sycl::detail::enable_if_t<T::is2D()>>
487  auto column(int i) {
488  return select<getSizeY(), 1, 1, 1>(0, i);
489  }
490 
494  template <typename T = Derived,
495  typename = sycl::detail::enable_if_t<T::is1D()>>
496  element_type operator[](int i) const {
497  const auto v = read();
498  return v[i];
499  }
500 
504  template <typename T = Derived,
505  typename = sycl::detail::enable_if_t<T::is1D()>>
506  auto operator[](int i) {
507  return select<1, 1>(i);
508  }
509 
511  template <int Rep> get_simd_t<element_type, Rep> replicate() {
512  return read().template replicate<Rep>();
513  }
514 
518  template <int Rep, int W>
519  get_simd_t<element_type, Rep * W> replicate_w(uint16_t OffsetX) {
520  return replicate_vs_w<Rep, 0, W>(0, OffsetX);
521  }
522 
526  template <int Rep, int W>
527  get_simd_t<element_type, Rep * W> replicate_w(uint16_t OffsetY,
528  uint16_t OffsetX) {
529  return replicate_vs_w<Rep, 0, W>(OffsetY, OffsetX);
530  }
531 
532  // clang-format off
536  // clang-format on
537  template <int Rep, int VS, int W>
538  get_simd_t<element_type, Rep * W> replicate_vs_w(uint16_t OffsetX) {
539  return replicate_vs_w_hs<Rep, VS, W, 1>(0, OffsetX);
540  }
541 
542  // clang-format off
546  // clang-format on
547  template <int Rep, int VS, int W>
548  get_simd_t<element_type, Rep * W> replicate_vs_w(uint16_t OffsetY,
549  uint16_t OffsetX) {
550  return replicate_vs_w_hs<Rep, VS, W, 1>(OffsetY, OffsetX);
551  }
552 
555  template <int Rep, int VS, int W, int HS>
556  get_simd_t<element_type, Rep * W> replicate_vs_w_hs(uint16_t OffsetX) {
557  return read().template replicate_vs_w_hs<Rep, VS, W, HS>(OffsetX);
558  }
559 
564  template <int Rep, int VS, int W, int HS>
565  get_simd_t<element_type, Rep * W> replicate_vs_w_hs(uint16_t OffsetY,
566  uint16_t OffsetX) {
567  constexpr int RowSize = is2D() ? getSizeX() : 0;
568  return read().template replicate_vs_w_hs<Rep, VS, W, HS>(OffsetY * RowSize +
569  OffsetX);
570  }
571 
573  template <typename T1 = element_type, typename T2 = BaseTy,
574  typename = std::enable_if_t<std::is_integral<T1>::value, T2>>
575  uint16_t any() {
576  return read().any();
577  }
578 
580  template <typename T1 = element_type, typename T2 = BaseTy,
581  typename = std::enable_if_t<std::is_integral<T1>::value, T2>>
582  uint16_t all() {
583  return read().all();
584  }
585 
587 public:
588  // Getter for the test proxy member, if enabled
589  __ESIMD_DECLARE_TEST_PROXY_ACCESS
590 
591 protected:
592  // The reference to the base object, which must be a simd object
593  BaseTy &M_base;
594 
595  // The test proxy if enabled
596  __ESIMD_DECLARE_TEST_PROXY
597 
598  // The region applied on the base object. Its type could be
599  // - region1d_t
600  // - region2d_t
601  // - std::pair<top_region_type, base_region_type>
602  //
603  RegionTy M_region;
605 };
606 
608 
609 } // namespace __ESIMD_DNS
610 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::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
cl::__ESIMD_DNS::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:395
cl::__ESIMD_DNS::simd_view_impl::getStrideY
static constexpr int getStrideY()
Get element stride of the view along Y dimension.
Definition: simd_view_impl.hpp:137
cl::__ESIMD_DNS::simd_view_impl::operator!
auto operator!()
Unary logical negeation operator. Applies only to integer element types.
Definition: simd_view_impl.hpp:372
cl::__ESIMD_DNS::simd_view_impl::getSizeY
static constexpr int getSizeY()
Get number of elements in the view along Y dimension.
Definition: simd_view_impl.hpp:135
__ESIMD_BITWISE_OP_FILTER
#define __ESIMD_BITWISE_OP_FILTER
Definition: simd_view_impl.hpp:330
cl::__ESIMD_DNS::simd_view_impl::operator++
value_type operator++(int)
Postfix increment.
Definition: simd_view_impl.hpp:449
cl::__ESIMD_DNS::simd_view_impl::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:79
cl::__ESIMD_DNS::simd_view_impl::data
value_type::raw_vector_type data() const
Definition: simd_view_impl.hpp:161
T
cl::__ESIMD_DNS::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:506
cl::__ESIMD_DNS::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:187
cl::__ESIMD_DNS::simd_view_impl::row
auto row(int i)
Reference a row from a 2D region.
Definition: simd_view_impl.hpp:476
simd_obj_impl
cl::__ESIMD_DNS::simd_view_impl::select
auto select(uint16_t OffsetY=0, uint16_t OffsetX=0)
2D region select.
Definition: simd_view_impl.hpp:284
cl::__ESIMD_DNS::simd_view_impl::is2D
static constexpr bool is2D()
Tells whether this view is 2-dimensional.
Definition: simd_view_impl.hpp:129
__ESIMD_ARITH_OP_FILTER
#define __ESIMD_ARITH_OP_FILTER
Definition: simd_view_impl.hpp:344
cl::__ESIMD_DNS::simd_view_impl::getStrideX
static constexpr int getStrideX()
Get element stride of the view along X dimension.
Definition: simd_view_impl.hpp:133
__ESIMD_DEF_UNARY_OP
#define __ESIMD_DEF_UNARY_OP(UNARY_OP, COND)
Definition: simd_view_impl.hpp:356
intrin.hpp
cl::sycl::detail::write
void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize, unsigned WIOffset, const char *Str, unsigned Len, unsigned Padding=0)
Definition: stream.hpp:110
test_proxy.hpp
cl::__ESIMD_DNS::simd_view_impl::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:82
cl::__ESIMD_DNS::simd_view_impl
Base class for "simd view" types.
Definition: simd_view_impl.hpp:40
cl::__ESIMD_DNS::simd_view_impl::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:73
cl::__ESIMD_DNS::simd_view_impl::bit_cast_view
auto bit_cast_view()
Create a 1-dimensional view of the target region.
Definition: simd_view_impl.hpp:198
cl::sycl::length
float length(T p) __NOEXC
Definition: builtins.hpp:1032
cl::__ESIMD_DNS::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:548
cl::__ESIMD_DNS::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:511
cl::__ESIMD_DNS::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:407
cl::__ESIMD_DNS::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:424
cl::__ESIMD_DNS::simd_view_impl::operator=
simd_view_impl & operator=(const simd_view_impl &Other)
Copy assignment.
Definition: simd_view_impl.hpp:381
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::__ESIMD_DNS::simd_view_impl::any
uint16_t any()
Applies simd_obj_impl::any operation to the target region.
Definition: simd_view_impl.hpp:575
cl::__ESIMD_DNS::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:55
cl::__ESIMD_DNS::simd_view_impl::operator++
Derived & operator++()
Prefix increment.
Definition: simd_view_impl.hpp:441
cl::__ESIMD_DNS::simd_view_impl::operator=
Derived & operator=(const Derived &Other)
Assignment from an object of the derived class.
Definition: simd_view_impl.hpp:389
cl::__ESIMD_DNS::simd_view_impl::all
uint16_t all()
Applies simd_obj_impl::all operation to the target region.
Definition: simd_view_impl.hpp:582
cl::__ESIMD_DNS::simd_view_impl::is1D
static constexpr bool is1D()
Tells whether this view is 1-dimensional.
Definition: simd_view_impl.hpp:127
cl::__ESIMD_DNS::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:141
cl::__ESIMD_DNS::simd_view_impl::select
auto select(uint16_t Offset=0)
1D region select.
Definition: simd_view_impl.hpp:233
cl::__ESIMD_DNS::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:565
cl::__ESIMD_DNS::simd_view_impl::operator[]
element_type operator[](int i) const
Read a single element from the target 1D region.
Definition: simd_view_impl.hpp:496
cl::__ESIMD_DNS::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:556
simd
Definition: simd.hpp:1027
cl::__ESIMD_DNS::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:147
cl::__ESIMD_DNS::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:538
cl::__ESIMD_DNS::simd_view_impl::operator--
Derived & operator--()
Prefix decrement.
Definition: simd_view_impl.hpp:457
__ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN
#define __ESIMD_DEF_SIMD_VIEW_IMPL_OPASSIGN(BINOP, OPASSIGN, COND)
Definition: simd_view_impl.hpp:292
cl::__ESIMD_DNS::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:519
cl::__ESIMD_DNS::simd_view_impl::Derived
simd_view< BaseTy, RegionTy > Derived
The only type which is supposed to extend this one and be used in user code.
Definition: simd_view_impl.hpp:44
cl::__ESIMD_DNS::simd_view_impl::column
auto column(int i)
Reference a column from a 2D region.
Definition: simd_view_impl.hpp:487
cl::__ESIMD_DNS::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:177
cl::__ESIMD_DNS::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:527
cl::__ESIMD_DNS::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:167
cl::__ESIMD_DNS::simd_view_impl::operator=
Derived & operator=(T1 RHS)
Broadcast assignment of a scalar with conversion.
Definition: simd_view_impl.hpp:435
cl::__ESIMD_DNS::simd_view_impl::operator=
Derived & operator=(Derived &&Other)
Assignment from an rvalue object of the derived class.
Definition: simd_view_impl.hpp:401
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::__ESIMD_DNS::simd_view_impl::operator--
value_type operator--(int)
Postfix decrement.
Definition: simd_view_impl.hpp:465
cl::__ESIMD_DNS::simd_view_impl::bit_cast_view
auto bit_cast_view()
Create a 2-dimensional view of the target region.
Definition: simd_view_impl.hpp:216
__ESIMD_SHIFT_OP_FILTER
#define __ESIMD_SHIFT_OP_FILTER
Definition: simd_view_impl.hpp:337
cl::__ESIMD_DNS::simd_view_impl::raw_element_type
__raw_t< element_type > raw_element_type
Corresponding "raw" (storage) type for the element type.
Definition: simd_view_impl.hpp:76
cl::__ESIMD_DNS::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:154
type_format.hpp
cl::__ESIMD_DNS::simd_view_impl::getSizeX
static constexpr int getSizeX()
Get number of elements in the view along X dimension.
Definition: simd_view_impl.hpp:131
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12