DPC++ Runtime
Runtime libraries for oneAPI DPC++
simd_obj_impl.hpp
Go to the documentation of this file.
1 //==------------ - simd_obj_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 // Implement Explicit SIMD vector APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
20 
22 namespace __ESIMD_NS {
23 
26 
43 
45 
48 
53  template <typename VT, typename ET = detail::element_type_t<VT>>
54  static constexpr unsigned alignment = alignof(ET);
55 };
56 
61  template <typename VT> static constexpr unsigned alignment = alignof(VT);
62 };
63 
68 template <unsigned N> struct overaligned_tag {
69  static_assert(
70  detail::isPowerOf2(N),
71  "Alignment value N for overaligned_tag<N> must be a power of two");
72  template <typename> static constexpr unsigned alignment = N;
73 };
74 
75 inline constexpr element_aligned_tag element_aligned = {};
76 
77 inline constexpr vector_aligned_tag vector_aligned = {};
78 
79 template <unsigned N> inline constexpr overaligned_tag<N> overaligned = {};
80 
82 template <typename T> struct is_simd_flag_type : std::false_type {};
83 
84 template <> struct is_simd_flag_type<element_aligned_tag> : std::true_type {};
85 
86 template <> struct is_simd_flag_type<vector_aligned_tag> : std::true_type {};
87 
88 template <unsigned N>
89 struct is_simd_flag_type<overaligned_tag<N>> : std::true_type {};
90 
93 template <typename T>
94 static inline constexpr bool is_simd_flag_type_v = is_simd_flag_type<T>::value;
95 
97 
98 namespace detail {
99 
101 
102 // Functions to support efficient simd constructors - avoiding internal loop
103 // over elements.
104 template <class T, int N, size_t... Is>
105 constexpr vector_type_t<T, N> make_vector_impl(const T (&&Arr)[N],
106  std::index_sequence<Is...>) {
107  return vector_type_t<T, N>{Arr[Is]...};
108 }
109 
110 template <class T, int N>
111 constexpr vector_type_t<T, N> make_vector(const T (&&Arr)[N]) {
112  return make_vector_impl<T, N>(std::move(Arr), std::make_index_sequence<N>{});
113 }
114 
115 template <class T, int N, size_t... Is>
116 constexpr vector_type_t<T, N> make_vector_impl(T Base, T Stride,
117  std::index_sequence<Is...>) {
118  return vector_type_t<T, N>{(T)(Base + ((T)Is) * Stride)...};
119 }
120 
121 template <class T, int N>
122 constexpr vector_type_t<T, N> make_vector(T Base, T Stride) {
123  return make_vector_impl<T, N>(Base, Stride, std::make_index_sequence<N>{});
124 }
125 
127 
130 
153 template <typename RawTy, int N, class Derived, class SFINAE>
156 
157  // For the is_simd_obj_impl_derivative helper to work correctly, all derived
158  // classes must be templated by element type and number of elements. If fewer
159  // template arguments are needed, template aliases can be used
160  // (simd_mask_type).
161  //
162  template <typename, typename> friend class simd_view;
163  template <typename, typename> friend class simd_view_impl;
164  template <typename, int> friend class simd;
165  template <typename, int> friend class simd_mask_impl;
166 
168 
169 public:
171  using element_type = get_vector_element_type<Derived>;
172 
174  using raw_vector_type = vector_type_t<RawTy, N>;
175 
177  using raw_element_type = RawTy;
178 
180  static constexpr int length = N;
181 
182 protected:
184  using Ty = element_type;
185 
186  template <bool UseSet = true>
187  void init_from_array(const Ty (&&Arr)[N]) noexcept {
188  raw_vector_type tmp;
189 
190  if constexpr (is_wrapper_elem_type_v<Ty>) {
191  for (auto I = 0; I < N; ++I) {
192  tmp[I] = bitcast_to_raw_type(Arr[I]);
193  }
194  } else {
195  tmp = make_vector(std::move(Arr));
196  }
197  if constexpr (UseSet) {
198  set(std::move(tmp));
199  } else {
200  M_data = std::move(tmp);
201  }
202  }
203 
204  explicit operator raw_vector_type() const {
205  __esimd_dbg_print(explicit operator raw_vector_type());
206  return data();
207  }
208 
209 private:
210  Derived &cast_this_to_derived() { return reinterpret_cast<Derived &>(*this); }
211  const Derived &cast_this_to_derived() const {
212  return reinterpret_cast<const Derived &>(*this);
213  }
214 
216 
217 public:
220  simd_obj_impl() = default;
221 
224  simd_obj_impl(const simd_obj_impl &other) {
226  set(other.data());
227  }
228 
235  template <class Ty1, typename Derived1>
237  __esimd_dbg_print(simd_obj_impl(const simd_obj_impl... > &other));
238  set(convert_vector<Ty, element_type_t<Derived1>, N>(other.data()));
239  }
240 
245  set(Val);
246  }
247 
254  simd_obj_impl(Ty Val, Ty Step) noexcept {
255  __esimd_dbg_print(simd_obj_impl(Ty Val, Ty Step));
256  if constexpr (is_wrapper_elem_type_v<Ty> || !std::is_integral_v<Ty>) {
257  for (int i = 0; i < N; ++i) {
258  M_data[i] = bitcast_to_raw_type(Val);
259  Val = binary_op<BinOp::add, Ty>(Val, Step);
260  }
261  } else {
262  M_data = make_vector<Ty, N>(Val, Step);
263  }
264  }
265 
271  template <class T1,
272  class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
273  simd_obj_impl(T1 Val) noexcept {
275  M_data = bitcast_to_raw_type(detail::convert_scalar<Ty>(Val));
276  }
277 
282  template <int N1, class = std::enable_if_t<N1 == N>>
283  simd_obj_impl(const Ty (&&Arr)[N1]) noexcept {
284  __esimd_dbg_print(simd_obj_impl(const Ty(&&Arr)[N1]));
285  init_from_array<false /*init M_data w/o using set(...)*/>(std::move(Arr));
286  // It is OK not to mark a write to M_data with __esimd_vstore (via 'set')
287  // here because:
288  // - __esimd_vstore/vload are need only to mark ESIMD_PRIVATE variable
289  // access for the VC BE to generate proper code for them.
290  // - initializers are not allowed for ESIMD_PRIVATE vars, so only the
291  // default ctor can be used for them
292  }
293 
300  template <typename Flags = element_aligned_tag,
301  typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
302  simd_obj_impl(const Ty *ptr, Flags = {}) noexcept {
303  __esimd_dbg_print(simd_obj_impl(const Ty *ptr, Flags));
304  copy_from(ptr, Flags{});
305  }
306 
317  template <
318  typename AccessorT, typename Flags = element_aligned_tag,
319  typename = std::enable_if_t<
320  detail::is_sycl_accessor_with<AccessorT, accessor_mode_cap::can_read,
322  is_simd_flag_type_v<Flags>>>
323  simd_obj_impl(AccessorT acc, uint32_t offset, Flags = {}) noexcept {
324  __esimd_dbg_print(simd_obj_impl(AccessorT acc, uint32_t offset, Flags));
325  copy_from(acc, offset, Flags{});
326  }
327 
331  template <int N1> std::enable_if_t<N1 == N> copy_from(const Ty (&&Arr)[N1]) {
332  __esimd_dbg_print(copy_from(const Ty(&&Arr)[N1]));
333  init_from_array(std::move(Arr));
334  }
335 
338  template <typename T = simd_obj_impl,
340  operator Ty() const {
341  __esimd_dbg_print(operator Ty());
342  return bitcast_to_wrapper_type<Ty>(data()[0]);
343  }
344 
348 #ifndef __SYCL_DEVICE_ONLY__
349  return M_data;
350 #else
351  return __esimd_vload<RawTy, N>(&M_data);
352 #endif
353  }
354 
357  Derived read() const { return Derived{data()}; }
358 
361  Derived &write(const Derived &Val) {
362  set(Val.data());
363  return cast_this_to_derived();
364  }
365 
372  void merge(const Derived &Val, const simd_mask_type<N> &Mask) {
373  set(__esimd_wrregion<RawTy, N, N, 0 /*VS*/, N, 1, N>(data(), Val.data(), 0,
374  Mask.data()));
375  }
376 
384  void merge(const Derived &Val1, Derived Val2, const simd_mask_type<N> &Mask) {
385  Val2.merge(Val1, Mask);
386  set(Val2.data());
387  }
388 
395  template <typename EltTy> auto bit_cast_view() &[[clang::lifetimebound]] {
396  using TopRegionTy = compute_format_type_t<Derived, EltTy>;
397  using RetTy = simd_view<Derived, TopRegionTy>;
398  return RetTy{cast_this_to_derived(), TopRegionTy{0}};
399  }
400 
411  template <typename EltTy, int Height, int Width>
412  auto bit_cast_view() &[[clang::lifetimebound]] {
413  using TopRegionTy = compute_format_type_2d_t<Derived, EltTy, Height, Width>;
414  using RetTy = simd_view<Derived, TopRegionTy>;
415  return RetTy{cast_this_to_derived(), TopRegionTy{0, 0}};
416  }
417 
425  template <int Size, int Stride>
427  select(uint16_t Offset = 0) &[[clang::lifetimebound]] {
428  static_assert(Size > 1 || Stride == 1,
429  "Stride must be 1 in single-element region");
430  region1d_t<Ty, Size, Stride> Reg(Offset);
431  return {cast_this_to_derived(), std::move(Reg)};
432  }
433 
441  template <int Size, int Stride>
442  resize_a_simd_type_t<Derived, Size> select(uint16_t Offset = 0) && {
443  static_assert(Size > 1 || Stride == 1,
444  "Stride must be 1 in single-element region");
445  Derived &&Val = std::move(cast_this_to_derived());
446  return __esimd_rdregion<RawTy, N, Size, /*VS*/ 0, Size, Stride>(Val.data(),
447  Offset);
448  }
449 
453  Ty operator[](int i) const { return bitcast_to_wrapper_type<Ty>(data()[i]); }
454 
459  [[clang::lifetimebound]] {
460  return select<1, 1>(i);
461  }
462 
468  template <int Size>
469  resize_a_simd_type_t<Derived, Size>
470  iselect(const simd<uint16_t, Size> &Indices) {
471  vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(RawTy);
472  return __esimd_rdindirect<RawTy, N, Size>(data(), Offsets);
473  }
474 
478  void iupdate(ushort Index, Ty V) {
479  auto Val = data();
480  Val[Index] = bitcast_to_raw_type(V);
481  set(Val);
482  }
483 
489  template <int Size>
490  void iupdate(const simd<uint16_t, Size> &Indices,
491  const resize_a_simd_type_t<Derived, Size> &Val,
492  const simd_mask_type<Size> &Mask) {
493  vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(RawTy);
494  set(__esimd_wrindirect<RawTy, N, Size>(data(), Val.data(), Offsets,
495  Mask.data()));
496  }
497 
501  template <int Rep> resize_a_simd_type_t<Derived, Rep * N> replicate() const {
502  return replicate_w<Rep, N>(0);
503  }
504 
512  template <int Rep, int W>
513  resize_a_simd_type_t<Derived, Rep * W> replicate_w(uint16_t Offset) const {
514  return replicate_vs_w_hs<Rep, 0, W, 1>(Offset);
515  }
516 
526  template <int Rep, int VS, int W>
527  resize_a_simd_type_t<Derived, Rep * W> replicate_vs_w(uint16_t Offset) const {
528  return replicate_vs_w_hs<Rep, VS, W, 1>(Offset);
529  }
530 
579  template <int Rep, int VS, int W, int HS>
580  resize_a_simd_type_t<Derived, Rep * W>
581  replicate_vs_w_hs(uint16_t Offset) const {
582  return __esimd_rdregion<RawTy, N, Rep * W, VS, W, HS, N>(
583  data(), Offset * sizeof(RawTy));
584  }
585 
589  template <typename T1 = Ty,
590  typename = std::enable_if_t<std::is_integral<T1>::value>>
591  uint16_t any() const {
592  return __esimd_any<Ty, N>(data());
593  }
594 
598  template <typename T1 = Ty,
599  typename = std::enable_if_t<std::is_integral<T1>::value>>
600  uint16_t all() const {
601  return __esimd_all<Ty, N>(data());
602  }
603 
604 protected:
607  template <typename RTy, class ElemTy = __raw_t<typename RTy::element_type>>
608  ESIMD_INLINE void writeRegion(RTy Region,
609  const vector_type_t<ElemTy, RTy::length> &Val) {
610 
611  if constexpr (N * sizeof(RawTy) == RTy::length * sizeof(ElemTy))
612  // update the entire vector
613  set(bitcast<RawTy, ElemTy, RTy::length>(Val));
614  else {
615  static_assert(!RTy::Is_2D);
616  // If element type differs, do bitcast conversion first.
617  auto Base = bitcast<ElemTy, RawTy, N>(data());
618  constexpr int BN = (N * sizeof(RawTy)) / sizeof(ElemTy);
619  // Access the region information.
620  constexpr int M = RTy::Size_x;
621  constexpr int Stride = RTy::Stride_x;
622  uint16_t Offset = Region.M_offset_x * sizeof(ElemTy);
623 
624  // Merge and update.
625  auto Merged = __esimd_wrregion<ElemTy, BN, M,
626  /*VS*/ 0, M, Stride>(Base, Val, Offset);
627  // Convert back to the original element type, if needed.
628  set(bitcast<RawTy, ElemTy, BN>(Merged));
629  }
630  }
631 
634  template <typename TR, typename UR,
635  class ElemTy = __raw_t<typename TR::element_type>>
636  ESIMD_INLINE void writeRegion(std::pair<TR, UR> Region,
637  const vector_type_t<ElemTy, TR::length> &Val) {
638  // parent-region type
639  using PaTy = typename shape_type<UR>::type;
640  using BT = __raw_t<typename PaTy::element_type>;
641  constexpr int BN = PaTy::length;
642 
643  if constexpr (PaTy::Size_in_bytes == TR::Size_in_bytes) {
644  writeRegion(Region.second, bitcast<BT, ElemTy, TR::length>(Val));
645  } else {
646  // Recursively read the base
647  auto Base = readRegion<RawTy, N>(data(), Region.second);
648  // If element type differs, do bitcast conversion first.
649  auto Base1 = bitcast<ElemTy, BT, BN>(Base);
650  constexpr int BN1 = PaTy::Size_in_bytes / sizeof(ElemTy);
651 
652  if constexpr (!TR::Is_2D) {
653  // Access the region information.
654  constexpr int M = TR::Size_x;
655  constexpr int Stride = TR::Stride_x;
656  uint16_t Offset = Region.first.M_offset_x * sizeof(ElemTy);
657 
658  // Merge and update.
659  Base1 = __esimd_wrregion<ElemTy, BN1, M,
660  /*VS*/ 0, M, Stride>(Base1, Val, Offset);
661  } else {
662  static_assert(std::is_same<ElemTy, BT>::value);
663  // Read columns with non-trivial horizontal stride.
664  constexpr int M = TR::length;
665  constexpr int VS = PaTy::Size_x * TR::Stride_y;
666  constexpr int W = TR::Size_x;
667  constexpr int HS = TR::Stride_x;
668  constexpr int ParentWidth = PaTy::Size_x;
669 
670  // Compute the byte offset for the starting element.
671  uint16_t Offset = static_cast<uint16_t>(
672  (Region.first.M_offset_y * PaTy::Size_x + Region.first.M_offset_x) *
673  sizeof(ElemTy));
674 
675  // Merge and update.
676  Base1 = __esimd_wrregion<ElemTy, BN1, M, VS, W, HS, ParentWidth>(
677  Base1, Val, Offset);
678  }
679  // Convert back to the original element type, if needed.
680  auto Merged1 = bitcast<BT, ElemTy, BN1>(Base1);
681  // recursively write it back to the base
682  writeRegion(Region.second, Merged1);
683  }
684  }
685 
686 public:
697  template <typename Flags = element_aligned_tag, int ChunkSize = 32,
698  typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
699  ESIMD_INLINE void copy_from(const Ty *addr, Flags = {}) SYCL_ESIMD_FUNCTION;
700 
712  template <typename AccessorT, typename Flags = element_aligned_tag,
713  int ChunkSize = 32,
714  typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
715  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read,
717  copy_from(AccessorT acc, uint32_t offset, Flags = {}) SYCL_ESIMD_FUNCTION;
718 
725  template <typename Flags = element_aligned_tag, int ChunkSize = 32,
726  typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
727  ESIMD_INLINE void copy_to(Ty *addr, Flags = {}) const SYCL_ESIMD_FUNCTION;
728 
738  template <typename AccessorT, typename Flags = element_aligned_tag,
739  int ChunkSize = 32,
740  typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
741  ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write,
743  copy_to(AccessorT acc, uint32_t offset, Flags = {}) const SYCL_ESIMD_FUNCTION;
744 
745  // Unary operations.
746 
750  template <class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
751  Derived operator~() const {
752  return Derived{
753  detail::vector_unary_op<detail::UnaryOp::bit_not, T1, N>(data())};
754  }
755 
762  template <class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
763  simd_mask_type<N> operator!() const {
764  return *this == 0;
765  }
766 
767 #define __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(BINOP, OPASSIGN, COND) \
768  \
769  \
770  \
771  \
772  \
773  template <class T1, class SimdT, \
774  class = std::enable_if_t<(is_simd_type_v<Derived> == \
775  is_simd_type_v<SimdT>)&&COND>> \
776  Derived &operator OPASSIGN( \
777  const __ESIMD_DNS::simd_obj_impl<T1, N, SimdT> &RHS) { \
778  auto Res = *this BINOP RHS; \
779  using ResT = decltype(Res); \
780  set(__ESIMD_DNS::convert_vector<element_type, typename ResT::element_type, \
781  length>(Res.data())); \
782  return cast_this_to_derived(); \
783  } \
784  \
785  \
786  \
787  \
788  \
789  \
790  template <class SimdT1, class RegionT1, \
791  class T1 = typename RegionT1::element_type, \
792  class = std::enable_if_t< \
793  (is_simd_type_v<Derived> == \
794  is_simd_type_v<SimdT1>)&&(RegionT1::length == length) && \
795  COND>> \
796  Derived &operator OPASSIGN( \
797  const __ESIMD_NS::simd_view<SimdT1, RegionT1> &RHS) { \
798  auto Res = *this BINOP RHS.read(); \
799  using ResT = decltype(Res); \
800  set(__ESIMD_DNS::convert_vector<element_type, typename ResT::element_type, \
801  length>(Res.data())); \
802  return cast_this_to_derived(); \
803  } \
804  \
805  \
806  \
807  \
808  template <class T1, class = std::enable_if_t<COND>> \
809  Derived &operator OPASSIGN(T1 RHS) { \
810  if constexpr (is_simd_type_v<Derived>) { \
811  using RHSVecT = __ESIMD_DNS::construct_a_simd_type_t<Derived, T1, N>; \
812  return *this OPASSIGN RHSVecT(RHS); \
813  } else { \
814  return *this OPASSIGN Derived((RawTy)RHS); \
815  } \
816  }
817 
818 // Bitwise operations are defined for simd objects and masks, and both operands
819 // must be integral
820 #define __ESIMD_BITWISE_OP_FILTER \
821  std::is_integral_v<element_type> &&std::is_integral_v<T1>
834 #undef __ESIMD_BITWISE_OP_FILTER
835 
836 // Bit shift operations are defined only for simd objects (not for masks), and
837 // both operands must be integral
838 #define __ESIMD_SHIFT_OP_FILTER \
839  std::is_integral_v<element_type> &&std::is_integral_v<T1> \
840  &&__ESIMD_DNS::is_simd_type_v<Derived>
841 
849 #undef __ESIMD_SHIFT_OP_FILTER
850 
851 // Arithmetic operations are defined only for simd objects, and the second
852 // operand's element type must be vectorizable. This requirement for 'this'
853 // is fulfilled, because otherwise 'this' couldn't have been constructed.
854 #define __ESIMD_ARITH_OP_FILTER \
855  __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1>
856 
865 #undef __ESIMD_ARITH_OP_FILTER
866 #undef __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN
867 
868  // Getter for the test proxy member, if enabled
869  __ESIMD_DECLARE_TEST_PROXY_ACCESS
870 
871 private:
872  // The underlying data for this vector.
873  raw_vector_type M_data;
874 
875 protected:
876  // The test proxy if enabled
877  __ESIMD_DECLARE_TEST_PROXY
878 
879  void set(const raw_vector_type &Val) {
880 #ifndef __SYCL_DEVICE_ONLY__
881  M_data = Val;
882 #else
883  __esimd_vstore<RawTy, N>(&M_data, Val);
884 #endif
885  }
886 };
888 
889 } // namespace detail
890 
891 } // namespace __ESIMD_NS
892 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::merge
void merge(const Derived &Val, const simd_mask_type< N > &Mask)
"Merges" this object's value with another object: replaces part of the underlying data with the one t...
Definition: simd_obj_impl.hpp:372
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::replicate_vs_w_hs
resize_a_simd_type_t< Derived, Rep *W > replicate_vs_w_hs(uint16_t Offset) const
This function "replicates" a portion of this object's elements into a new object.
Definition: simd_obj_impl.hpp:581
__ESIMD_BITWISE_OP_FILTER
#define __ESIMD_BITWISE_OP_FILTER
Definition: simd_obj_impl.hpp:820
T
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(Ty Val, Ty Step) noexcept
Arithmetic progression constructor.
Definition: simd_obj_impl.hpp:254
simd_view.hpp
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::iupdate
void iupdate(const simd< uint16_t, Size > &Indices, const resize_a_simd_type_t< Derived, Size > &Val, const simd_mask_type< Size > &Mask)
Indirect update - update multiple elements with given variable indices.
Definition: simd_obj_impl.hpp:490
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::operator[]
Ty operator[](int i) const
Get value of this vector's element.
Definition: simd_obj_impl.hpp:453
simd_obj_impl
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::copy_from
std::enable_if_t< N1==N > copy_from(const Ty(&&Arr)[N1])
Initializes this object from an rvalue to an array with the same number of elements.
Definition: simd_obj_impl.hpp:331
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::bit_cast_view
auto bit_cast_view() &
Create a 2-dimensional view of this object.
Definition: simd_obj_impl.hpp:412
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::select
resize_a_simd_type_t< Derived, Size > select(uint16_t Offset=0) &&
Select and extract a subregion of this object's elements and return it as a new vector object.
Definition: simd_obj_impl.hpp:442
__ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN
#define __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(BINOP, OPASSIGN, COND)
Definition: simd_obj_impl.hpp:767
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(const simd_obj_impl< Ty1, N, Derived1, SFINAE > &other)
Implicit conversion constructor from another simd_obj_impl object.
Definition: simd_obj_impl.hpp:236
vector_aligned
constexpr vector_aligned_tag vector_aligned
Definition: simd.hpp:1036
element_aligned_tag
Definition: simd.hpp:1031
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::writeRegion
ESIMD_INLINE void writeRegion(RTy Region, const vector_type_t< ElemTy, RTy::length > &Val)
Write a simd_obj_impl-vector into a basic region of a simd_obj_impl object.
Definition: simd_obj_impl.hpp:608
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::bit_cast_view
auto bit_cast_view() &
Create a 1-dimensional view of this object.
Definition: simd_obj_impl.hpp:395
overaligned
constexpr overaligned_tag< _Np > overaligned
Definition: simd.hpp:1038
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::iupdate
void iupdate(ushort Index, Ty V)
Update single element with variable index.
Definition: simd_obj_impl.hpp:478
intrin.hpp
test_proxy.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::is_simd_flag_type
Checks if type is a simd load/store flag.
Definition: simd_obj_impl.hpp:82
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::replicate_w
resize_a_simd_type_t< Derived, Rep *W > replicate_w(uint16_t Offset) const
Shortcut to replicate_vs_w_hs with VS=0 and HS=1 to replicate a single "dense" (w/o gaps between elem...
Definition: simd_obj_impl.hpp:513
cl::sycl::ext::intel::esimd::overaligned_tag
overaligned_tag type.
Definition: simd_obj_impl.hpp:68
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::replicate
resize_a_simd_type_t< Derived, Rep *N > replicate() const
Replicates contents of this vector a number of times into a new vector.
Definition: simd_obj_impl.hpp:501
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::detail::simd_obj_impl::any
uint16_t any() const
See if any element is non-zero.
Definition: simd_obj_impl.hpp:591
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::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(AccessorT acc, uint32_t offset, Flags={}) noexcept
Accessor-based load constructor.
Definition: simd_obj_impl.hpp:323
cl::sycl::ext::intel::esimd::simd
The main simd vector class.
Definition: types.hpp:31
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::operator[]
simd_view< Derived, region1d_scalar_t< Ty > > operator[](int i)
Return writable view of a single element.
Definition: simd_obj_impl.hpp:458
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(const simd_obj_impl &other)
Copy constructor.
Definition: simd_obj_impl.hpp:224
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::iselect
resize_a_simd_type_t< Derived, Size > iselect(const simd< uint16_t, Size > &Indices)
Indirect select - select and extract multiple elements with given variable indices.
Definition: simd_obj_impl.hpp:470
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::all
uint16_t all() const
See if all elements are non-zero.
Definition: simd_obj_impl.hpp:600
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::writeRegion
ESIMD_INLINE void writeRegion(std::pair< TR, UR > Region, const vector_type_t< ElemTy, TR::length > &Val)
Write a simd_obj_impl-vector into a nested region of a simd_obj_impl object.
Definition: simd_obj_impl.hpp:636
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::operator~
Derived operator~() const
Per-element bitwise inversion, available in all subclasses, but only for integral element types (simd...
Definition: simd_obj_impl.hpp:751
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::data
raw_vector_type data() const
Definition: simd_obj_impl.hpp:346
cl::sycl::ext::intel::esimd::vector_aligned_tag
vector_aligned_tag type.
Definition: simd_obj_impl.hpp:60
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
elem_type_traits.hpp
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::operator!
simd_mask_type< N > operator!() const
Unary logical negation operator, available in all subclasses, but only for integral element types (si...
Definition: simd_obj_impl.hpp:763
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::element_type
get_vector_element_type< Derived > element_type
Element type of the derived (user) class.
Definition: simd_obj_impl.hpp:171
__ESIMD_SHIFT_OP_FILTER
#define __ESIMD_SHIFT_OP_FILTER
Definition: simd_obj_impl.hpp:838
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::set
__ESIMD_DECLARE_TEST_PROXY void set(const raw_vector_type &Val)
Definition: simd_obj_impl.hpp:879
__ESIMD_ARITH_OP_FILTER
#define __ESIMD_ARITH_OP_FILTER
Definition: simd_obj_impl.hpp:854
__esimd_dbg_print
#define __esimd_dbg_print(a)
Definition: types.hpp:22
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(const raw_vector_type &Val)
Implicit conversion constructor from a raw vector object.
Definition: simd_obj_impl.hpp:243
cl::sycl::ext::intel::esimd::element_aligned_tag
element_aligned_tag type.
Definition: simd_obj_impl.hpp:52
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::select
simd_view< Derived, region1d_t< Ty, Size, Stride > > select(uint16_t Offset=0) &
Select elements of this object into a subregion and create a 1D view for for it.
Definition: simd_obj_impl.hpp:427
sycl_util.hpp
element_aligned
constexpr element_aligned_tag element_aligned
Definition: simd.hpp:1035
cl::sycl::ext::intel::esimd::ushort
unsigned short ushort
Definition: common.hpp:83
cl::sycl::info::device
device
Definition: info_desc.hpp:53
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(T1 Val) noexcept
Broadcast constructor.
Definition: simd_obj_impl.hpp:273
memory_intrin.hpp
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::detail::simd_obj_impl::read
Derived read() const
Definition: simd_obj_impl.hpp:357
is_simd_flag_type_v
constexpr bool is_simd_flag_type_v
Definition: simd.hpp:1083
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(const Ty *ptr, Flags={}) noexcept
Pointer-based load constructor.
Definition: simd_obj_impl.hpp:302
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::write
Derived & write(const Derived &Val)
Replaces the underlying data with the one taken from another object.
Definition: simd_obj_impl.hpp:361
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::simd_obj_impl
simd_obj_impl(const Ty(&&Arr)[N1]) noexcept
Rvalue array-based constructor.
Definition: simd_obj_impl.hpp:283
type_format.hpp
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::merge
void merge(const Derived &Val1, Derived Val2, const simd_mask_type< N > &Mask)
Merges given two objects with a mask and writes resulting data into this object.
Definition: simd_obj_impl.hpp:384
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::ext::intel::esimd::detail::simd_obj_impl::replicate_vs_w
resize_a_simd_type_t< Derived, Rep *W > replicate_vs_w(uint16_t Offset) const
Shortcut to replicate_vs_w_hs with HS=1 to replicate dense blocks.
Definition: simd_obj_impl.hpp:527