simd_obj_impl.hpp

Go to the documentation of this file.

//==------------ - simd_obj_impl.hpp - DPC++ Explicit SIMD API -------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// Implement Explicit SIMD vector APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/aspects.hpp>
#include <sycl/ext/intel/esimd/detail/elem_type_traits.hpp>
#include <sycl/ext/intel/esimd/detail/intrin.hpp>
#include <sycl/ext/intel/esimd/detail/memory_intrin.hpp>
#include <sycl/ext/intel/esimd/detail/sycl_util.hpp>
#include <sycl/ext/intel/esimd/detail/test_proxy.hpp>
#include <sycl/ext/intel/esimd/detail/type_format.hpp>
#include <sycl/ext/intel/esimd/simd_view.hpp>
 
namespace sycl {
inline namespace _V1 {
namespace ext::intel::esimd {
 
 
 
 
 
struct element_aligned_tag {
  template <typename VT, typename ET = detail::element_type_t<VT>>
  static constexpr unsigned alignment = alignof(ET);
};
 
struct vector_aligned_tag {
  template <typename VT> static constexpr unsigned alignment = alignof(VT);
};
 
template <unsigned N> struct overaligned_tag {
  static_assert(
      detail::isPowerOf2(N),
      "Alignment value N for overaligned_tag<N> must be a power of two");
  template <typename> static constexpr unsigned alignment = N;
};
 
inline constexpr element_aligned_tag element_aligned = {};
 
inline constexpr vector_aligned_tag vector_aligned = {};
 
template <unsigned N> inline constexpr overaligned_tag<N> overaligned = {};
 
template <typename T> struct is_simd_flag_type : std::false_type {};
 
template <> struct is_simd_flag_type<element_aligned_tag> : std::true_type {};
 
template <> struct is_simd_flag_type<vector_aligned_tag> : std::true_type {};
 
template <unsigned N>
struct is_simd_flag_type<overaligned_tag<N>> : std::true_type {};
 
template <typename T>
static inline constexpr bool is_simd_flag_type_v = is_simd_flag_type<T>::value;
 
 
namespace detail {
 
struct dqword_element_aligned_tag {
  template <typename VT, typename ET = detail::element_type_t<VT>>
  static constexpr unsigned alignment = alignof(ET) > 4 ? alignof(ET) : 4;
};
 
inline constexpr dqword_element_aligned_tag dqword_element_aligned = {};
 
// Functions to support efficient simd constructors - avoiding internal loop
// over elements.
template <class T, int N, size_t... Is>
constexpr vector_type_t<T, N> make_vector_impl(const T (&&Arr)[N],
                                               std::index_sequence<Is...>) {
  return vector_type_t<T, N>{Arr[Is]...};
}
 
template <class T, int N>
constexpr vector_type_t<T, N> make_vector(const T (&&Arr)[N]) {
  return make_vector_impl<T, N>(std::move(Arr), std::make_index_sequence<N>{});
}
 
template <class T, int N, size_t... Is>
constexpr auto make_vector_impl(T Base, T Stride, std::index_sequence<Is...>) {
  if constexpr (std::is_integral_v<T> && N <= 3) {
    // This sequence is a bit more efficient for integral types and N <= 3.
    return vector_type_t<T, N>{(T)(Base + ((T)Is) * Stride)...};
  } else {
    using CppT = typename element_type_traits<T>::EnclosingCppT;
    CppT BaseCpp = Base;
    CppT StrideCpp = Stride;
    vector_type_t<CppT, N> VBase = BaseCpp;
    vector_type_t<CppT, N> VStride = StrideCpp;
    vector_type_t<CppT, N> VStrideCoef{(CppT)(Is)...};
    vector_type_t<CppT, N> Result{VBase + VStride * VStrideCoef};
    return wrapper_type_converter<T>::template to_vector<N>(Result);
  }
}
 
template <class T, int N> constexpr auto make_vector(T Base, T Stride) {
  return make_vector_impl<T, N>(Base, Stride, std::make_index_sequence<N>{});
}
 
 
 
template <typename RawTy, int N, class Derived, class SFINAE>
#ifndef __SYCL_DEVICE_ONLY__
class simd_obj_impl {
#else
class [[__sycl_detail__::__uses_aspects__(
    sycl::aspect::ext_intel_esimd)]] simd_obj_impl {
#endif
 
  // For the is_simd_obj_impl_derivative helper to work correctly, all derived
  // classes must be templated by element type and number of elements. If fewer
  // template arguments are needed, template aliases can be used
  // (simd_mask_type).
  //
  template <typename, typename> friend class simd_view;
  template <typename, typename> friend class simd_view_impl;
  template <typename, int> friend class simd;
  template <typename, int> friend class simd_mask_impl;
 
 
public:
  using element_type = get_vector_element_type<Derived>;
 
  using raw_vector_type = vector_type_t<RawTy, N>;
 
  using raw_element_type = RawTy;
 
  static constexpr int length = N;
 
protected:
  using Ty = element_type;
 
  template <bool UseSet = true>
  void init_from_array(const Ty (&&Arr)[N]) noexcept {
    raw_vector_type tmp;
 
    if constexpr (is_wrapper_elem_type_v<Ty>) {
      for (auto I = 0; I < N; ++I) {
        tmp[I] = bitcast_to_raw_type(Arr[I]);
      }
    } else {
      tmp = make_vector(std::move(Arr));
    }
    if constexpr (UseSet) {
      set(std::move(tmp));
    } else {
      M_data = std::move(tmp);
    }
  }
 
  explicit operator raw_vector_type() const {
    __esimd_dbg_print(explicit operator raw_vector_type());
    return data();
  }
 
private:
  Derived &cast_this_to_derived() { return reinterpret_cast<Derived &>(*this); }
  const Derived &cast_this_to_derived() const {
    return reinterpret_cast<const Derived &>(*this);
  }
 
 
public:
  simd_obj_impl() = default;
 
  simd_obj_impl(const simd_obj_impl &other) {
    __esimd_dbg_print(simd_obj_impl(const simd_obj_impl &other));
    set(other.data());
  }
 
  template <class Ty1, typename Derived1>
  simd_obj_impl(const simd_obj_impl<Ty1, N, Derived1, SFINAE> &other) {
    __esimd_dbg_print(simd_obj_impl(const simd_obj_impl... > &other));
    set(convert_vector<Ty, element_type_t<Derived1>, N>(other.data()));
  }
 
  simd_obj_impl(const raw_vector_type &Val) {
    __esimd_dbg_print(simd_obj_impl(const raw_vector_type &Val));
    set(Val);
  }
 
  simd_obj_impl(Ty Base, Ty Step) noexcept {
    __esimd_dbg_print(simd_obj_impl(Ty Base, Ty Step));
    M_data = make_vector<Ty, N>(Base, Step);
  }
 
  template <class T1,
            class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
  simd_obj_impl(T1 Val) noexcept {
    __esimd_dbg_print(simd_obj_impl(T1 Val));
    M_data = bitcast_to_raw_type(detail::convert_scalar<Ty>(Val));
  }
 
  template <int N1, class = std::enable_if_t<N1 == N>>
  simd_obj_impl(const Ty (&&Arr)[N1]) noexcept {
    __esimd_dbg_print(simd_obj_impl(const Ty(&&Arr)[N1]));
    init_from_array<false /*init M_data w/o using set(...)*/>(std::move(Arr));
    // It is OK not to mark a write to M_data with __esimd_vstore (via 'set')
    // here because:
    // - __esimd_vstore/vload are need only to mark ESIMD_PRIVATE variable
    //   access for the VC BE to generate proper code for them.
    // - initializers are not allowed for ESIMD_PRIVATE vars, so only the
    //   default ctor can be used for them
  }
 
  template <typename Flags,
            typename std::enable_if_t<is_simd_flag_type_v<Flags>, bool> = true>
  simd_obj_impl(const Ty *ptr, Flags) noexcept {
    __esimd_dbg_print(simd_obj_impl(const Ty *ptr, Flags));
    copy_from(ptr, Flags{});
  }
 
  template <typename PropertyListT = oneapi::experimental::empty_properties_t,
            typename std::enable_if_t<
                ext::oneapi::experimental::is_property_list_v<PropertyListT>,
                bool> = true>
  simd_obj_impl(const Ty *ptr, PropertyListT = {}) noexcept {
    __esimd_dbg_print(simd_obj_impl(const Ty *ptr, PropertyListT));
    using NewPropertyListT =
        detail::add_alignment_property_t<PropertyListT, sizeof(Ty)>;
    copy_from(ptr, NewPropertyListT{});
  }
 
  template <
      typename AccessorT, typename Flags,
      typename std::enable_if_t<
          detail::is_accessor_with_v<AccessorT, accessor_mode_cap::can_read> &&
              is_simd_flag_type_v<Flags>,
          bool> = true>
  simd_obj_impl(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
                uint64_t offset,
#else
                uint32_t offset,
#endif
                Flags) noexcept {
    __esimd_dbg_print(simd_obj_impl(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
                                    uint64_t offset,
#else
                                    uint32_t offset,
#endif
                                    Flags));
    copy_from(acc, offset, Flags{});
  }
 
  template <
      typename AccessorT,
      typename PropertyListT = oneapi::experimental::empty_properties_t,
      typename std::enable_if_t<
          detail::is_accessor_with_v<AccessorT, accessor_mode_cap::can_read> &&
              ext::oneapi::experimental::is_property_list_v<PropertyListT>,
          bool> = true>
  simd_obj_impl(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
                uint64_t offset,
#else
                uint32_t offset,
#endif
                PropertyListT = {}) noexcept {
    __esimd_dbg_print(simd_obj_impl(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
                                    uint64_t offset,
#else
                                    uint32_t offset,
#endif
                                    PropertyListT));
    using NewPropertyListT =
        detail::add_alignment_property_t<PropertyListT, sizeof(Ty)>;
    copy_from(acc, offset, NewPropertyListT{});
  }
 
  Derived &operator=(const simd_obj_impl &other) noexcept {
    set(other.data());
    return cast_this_to_derived();
  }
 
  template <typename T = simd_obj_impl,
            typename = std::enable_if_t<T::length == 1>>
  operator Ty() const {
    __esimd_dbg_print(operator Ty());
    return bitcast_to_wrapper_type<Ty>(data()[0]);
  }
 
  raw_vector_type data() const {
    __esimd_dbg_print(raw_vector_type data());
#ifndef __SYCL_DEVICE_ONLY__
    return M_data;
#else
    return __esimd_vload<RawTy, N>(&M_data);
#endif
  }
 
  raw_vector_type &data_ref() { return M_data; }
 
  Derived read() const { return Derived{data()}; }
 
  Derived &write(const Derived &Val) {
    set(Val.data());
    return cast_this_to_derived();
  }
 
  void merge(const Derived &Val, const simd_mask_type<N> &Mask) {
    check_wrregion_params<N, N, 0 /*VS*/, N, 1>();
    set(__esimd_wrregion<RawTy, N, N, 0 /*VS*/, N, 1, N>(data(), Val.data(), 0,
                                                         Mask.data()));
  }
 
  void merge(const Derived &Val1, Derived Val2, const simd_mask_type<N> &Mask) {
    Val2.merge(Val1, Mask);
    set(Val2.data());
  }
 
  template <typename EltTy> auto bit_cast_view() &[[clang::lifetimebound]] {
    using TopRegionTy = compute_format_type_t<Derived, EltTy>;
    using RetTy = simd_view<Derived, TopRegionTy>;
    return RetTy{cast_this_to_derived(), TopRegionTy{0}};
  }
 
  template <typename EltTy, int Height, int Width>
  auto bit_cast_view() &[[clang::lifetimebound]] {
    using TopRegionTy = compute_format_type_2d_t<Derived, EltTy, Height, Width>;
    using RetTy = simd_view<Derived, TopRegionTy>;
    return RetTy{cast_this_to_derived(), TopRegionTy{0, 0}};
  }
 
  template <int Size, int Stride>
  simd_view<Derived, region1d_t<Ty, Size, Stride>>
  select(uint16_t Offset = 0) &[[clang::lifetimebound]] {
    static_assert(Size > 1 || Stride == 1,
                  "Stride must be 1 in single-element region");
    region1d_t<Ty, Size, Stride> Reg(Offset);
    return {cast_this_to_derived(), std::move(Reg)};
  }
 
  template <int Size, int Stride>
  resize_a_simd_type_t<Derived, Size> select(uint16_t Offset = 0) && {
    static_assert(Size > 1 || Stride == 1,
                  "Stride must be 1 in single-element region");
    Derived &&Val = std::move(cast_this_to_derived());
    check_rdregion_params<N, Size, /*VS*/ 0, Size, Stride>();
    return __esimd_rdregion<RawTy, N, Size, /*VS*/ 0, Size, Stride>(Val.data(),
                                                                    Offset);
  }
 
  Ty operator[](int i) const { return bitcast_to_wrapper_type<Ty>(data()[i]); }
 
  simd_view<Derived, region1d_scalar_t<Ty>> operator[](int i)
      [[clang::lifetimebound]] {
    return select<1, 1>(i);
  }
 
  template <int Size>
  resize_a_simd_type_t<Derived, Size>
  iselect(const simd<uint16_t, Size> &Indices) {
    vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(RawTy);
    return __esimd_rdindirect<RawTy, N, Size>(data(), Offsets);
  }
 
  void iupdate(ushort Index, Ty V) {
    auto Val = data();
    Val[Index] = bitcast_to_raw_type(V);
    set(Val);
  }
 
  template <int Size>
  void iupdate(const simd<uint16_t, Size> &Indices,
               const resize_a_simd_type_t<Derived, Size> &Val,
               const simd_mask_type<Size> &Mask) {
    vector_type_t<uint16_t, Size> Offsets = Indices.data() * sizeof(RawTy);
    set(__esimd_wrindirect<RawTy, N, Size>(data(), Val.data(), Offsets,
                                           Mask.data()));
  }
 
  template <int Rep> resize_a_simd_type_t<Derived, Rep * N> replicate() const {
    return replicate_w<Rep, N>(0);
  }
 
  template <int Rep, int W>
  resize_a_simd_type_t<Derived, Rep * W> replicate_w(uint16_t Offset) const {
    return replicate_vs_w_hs<Rep, 0, W, 1>(Offset);
  }
 
  template <int Rep, int VS, int W>
  resize_a_simd_type_t<Derived, Rep * W> replicate_vs_w(uint16_t Offset) const {
    return replicate_vs_w_hs<Rep, VS, W, 1>(Offset);
  }
 
  template <int Rep, int VS, int W, int HS>
  resize_a_simd_type_t<Derived, Rep * W>
  replicate_vs_w_hs(uint16_t Offset) const {
    check_rdregion_params<N, Rep * W, VS, W, HS>();
    return __esimd_rdregion<RawTy, N, Rep * W, VS, W, HS, N>(
        data(), Offset * sizeof(RawTy));
  }
 
  template <typename T1 = Ty,
            typename = std::enable_if_t<std::is_integral_v<T1>>>
  uint16_t any() const {
    return __esimd_any<Ty, N>(data());
  }
 
  template <typename T1 = Ty,
            typename = std::enable_if_t<std::is_integral_v<T1>>>
  uint16_t all() const {
    return __esimd_all<Ty, N>(data());
  }
 
protected:
  template <typename RTy, class ElemTy = __raw_t<typename RTy::element_type>>
  ESIMD_INLINE void writeRegion(RTy Region,
                                const vector_type_t<ElemTy, RTy::length> &Val) {
 
    if constexpr (N * sizeof(RawTy) == RTy::length * sizeof(ElemTy))
      // update the entire vector
      set(bitcast<RawTy, ElemTy, RTy::length>(Val));
    else {
      static_assert(!RTy::Is_2D);
      // If element type differs, do bitcast conversion first.
      auto Base = bitcast<ElemTy, RawTy, N>(data());
      constexpr int BN = (N * sizeof(RawTy)) / sizeof(ElemTy);
      // Access the region information.
      constexpr int M = RTy::Size_x;
      constexpr int Stride = RTy::Stride_x;
      uint16_t Offset = Region.M_offset_x * sizeof(ElemTy);
      check_wrregion_params<BN, M, /*VS*/ 0, M, Stride>();
      // Merge and update.
      auto Merged = __esimd_wrregion<ElemTy, BN, M,
                                     /*VS*/ 0, M, Stride>(Base, Val, Offset);
      // Convert back to the original element type, if needed.
      set(bitcast<RawTy, ElemTy, BN>(Merged));
    }
  }
 
  template <typename TR, typename UR,
            class ElemTy = __raw_t<typename TR::element_type>>
  ESIMD_INLINE void writeRegion(std::pair<TR, UR> Region,
                                const vector_type_t<ElemTy, TR::length> &Val) {
    // parent-region type
    using PaTy = typename shape_type<UR>::type;
    using BT = __raw_t<typename PaTy::element_type>;
    constexpr int BN = PaTy::length;
 
    if constexpr (PaTy::Size_in_bytes == TR::Size_in_bytes) {
      writeRegion(Region.second, bitcast<BT, ElemTy, TR::length>(Val));
    } else {
      // Recursively read the base
      auto Base = readRegion<RawTy, N>(data(), Region.second);
      // If element type differs, do bitcast conversion first.
      auto Base1 = bitcast<ElemTy, BT, BN>(Base);
      constexpr int BN1 = PaTy::Size_in_bytes / sizeof(ElemTy);
 
      if constexpr (!TR::Is_2D) {
        // Access the region information.
        constexpr int M = TR::Size_x;
        constexpr int Stride = TR::Stride_x;
        uint16_t Offset = Region.first.M_offset_x * sizeof(ElemTy);
 
        check_wrregion_params<BN1, M, /*VS*/ 0, M, Stride>();
        // Merge and update.
        Base1 = __esimd_wrregion<ElemTy, BN1, M,
                                 /*VS*/ 0, M, Stride>(Base1, Val, Offset);
      } else {
        static_assert(std::is_same<ElemTy, BT>::value);
        // Read columns with non-trivial horizontal stride.
        constexpr int M = TR::length;
        constexpr int VS = PaTy::Size_x * TR::Stride_y;
        constexpr int W = TR::Size_x;
        constexpr int HS = TR::Stride_x;
        constexpr int ParentWidth = PaTy::Size_x;
 
        // Compute the byte offset for the starting element.
        uint16_t Offset = static_cast<uint16_t>(
            (Region.first.M_offset_y * PaTy::Size_x + Region.first.M_offset_x) *
            sizeof(ElemTy));
 
        check_wrregion_params<BN1, M, VS, W, HS>();
        // Merge and update.
        Base1 = __esimd_wrregion<ElemTy, BN1, M, VS, W, HS, ParentWidth>(
            Base1, Val, Offset);
      }
      // Convert back to the original element type, if needed.
      auto Merged1 = bitcast<BT, ElemTy, BN1>(Base1);
      // recursively write it back to the base
      writeRegion(Region.second, Merged1);
    }
  }
 
public:
  template <typename Flags = element_aligned_tag, int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<is_simd_flag_type_v<Flags>>
  copy_from(const Ty *addr, Flags) SYCL_ESIMD_FUNCTION;
 
  template <int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_from(const Ty *addr, PropertyListT = {}) SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, typename Flags = element_aligned_tag,
            int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<detail::is_device_accessor_with_v<
                                    AccessorT, accessor_mode_cap::can_read> &&
                                is_simd_flag_type_v<Flags>>
  copy_from(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
            uint64_t offset,
#else
            uint32_t offset,
#endif
            Flags) SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      detail::is_device_accessor_with_v<AccessorT,
                                        accessor_mode_cap::can_read> &&
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_from(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
            uint64_t offset,
#else
            uint32_t offset,
#endif
            PropertyListT = {}) SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, typename Flags = element_aligned_tag,
            int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<detail::is_local_accessor_with_v<
                                    AccessorT, accessor_mode_cap::can_read> &&
                                    is_simd_flag_type_v<Flags>,
                                void>
  copy_from(AccessorT acc, uint32_t offset, Flags) SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      detail::is_local_accessor_with_v<AccessorT,
                                       accessor_mode_cap::can_read> &&
          ext::oneapi::experimental::is_property_list_v<PropertyListT>,
      void>
  copy_from(AccessorT acc, uint32_t offset,
            PropertyListT = {}) SYCL_ESIMD_FUNCTION;
 
  template <typename Flags, int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<is_simd_flag_type_v<Flags>>
  copy_to(Ty *addr, Flags) const SYCL_ESIMD_FUNCTION;
 
  template <int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_to(Ty *addr, PropertyListT = {}) const SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, typename Flags, int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<detail::is_device_accessor_with_v<
                                    AccessorT, accessor_mode_cap::can_write> &&
                                is_simd_flag_type_v<Flags>>
  copy_to(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
          uint64_t offset,
#else
          uint32_t offset,
#endif
          Flags) const SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      detail::is_device_accessor_with_v<AccessorT,
                                        accessor_mode_cap::can_write> &&
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_to(AccessorT acc,
#ifdef __ESIMD_FORCE_STATELESS_MEM
          uint64_t offset,
#else
          uint32_t offset,
#endif
          PropertyListT = {}) const SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, typename Flags, int ChunkSize = 32>
  ESIMD_INLINE std::enable_if_t<detail::is_local_accessor_with_v<
                                    AccessorT, accessor_mode_cap::can_write> &&
                                    is_simd_flag_type_v<Flags>,
                                void>
  copy_to(AccessorT acc, uint32_t offset, Flags) const SYCL_ESIMD_FUNCTION;
 
  template <typename AccessorT, int ChunkSize = 32,
            typename PropertyListT = oneapi::experimental::empty_properties_t>
  ESIMD_INLINE std::enable_if_t<
      detail::is_local_accessor_with_v<AccessorT,
                                       accessor_mode_cap::can_write> &&
          ext::oneapi::experimental::is_property_list_v<PropertyListT>,
      void>
  copy_to(AccessorT acc, uint32_t offset,
          PropertyListT = {}) const SYCL_ESIMD_FUNCTION;
 
  // Unary operations.
 
  template <class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
  Derived operator~() const {
    return Derived{
        detail::vector_unary_op<detail::UnaryOp::bit_not, T1, N>(data())};
  }
 
  template <class T1 = Ty, class = std::enable_if_t<std::is_integral_v<T1>>>
  simd_mask_type<N> operator!() const {
    return *this == 0;
  }
 
#define __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(BINOP, OPASSIGN, COND)              \
                                                                               \
 \
 \
 \
 \
  template <class T1, class SimdT,                                             \
            class = std::enable_if_t<(is_simd_type_v<Derived> ==               \
                                      is_simd_type_v<SimdT>)&&COND>>           \
  Derived &operator OPASSIGN(                                                  \
      const __ESIMD_DNS::simd_obj_impl<T1, N, SimdT> &RHS) {                   \
    auto Res = *this BINOP RHS;                                                \
    using ResT = decltype(Res);                                                \
    set(__ESIMD_DNS::convert_vector<element_type, typename ResT::element_type, \
                                    length>(Res.data()));                      \
    return cast_this_to_derived();                                             \
  }                                                                            \
                                                                               \
 \
 \
 \
 \
 \
  template <class SimdT1, class RegionT1,                                      \
            class T1 = typename RegionT1::element_type,                        \
            class = std::enable_if_t<                                          \
                (is_simd_type_v<Derived> ==                                    \
                 is_simd_type_v<SimdT1>)&&(RegionT1::length == length) &&      \
                COND>>                                                         \
  Derived &operator OPASSIGN(                                                  \
      const __ESIMD_NS::simd_view<SimdT1, RegionT1> &RHS) {                    \
    auto Res = *this BINOP RHS.read();                                         \
    using ResT = decltype(Res);                                                \
    set(__ESIMD_DNS::convert_vector<element_type, typename ResT::element_type, \
                                    length>(Res.data()));                      \
    return cast_this_to_derived();                                             \
  }                                                                            \
                                                                               \
 \
 \
 \
  template <class T1, class = std::enable_if_t<COND>>                          \
  Derived &operator OPASSIGN(T1 RHS) {                                         \
    if constexpr (is_simd_type_v<Derived>) {                                   \
      using RHSVecT = __ESIMD_DNS::construct_a_simd_type_t<Derived, T1, N>;    \
      return *this OPASSIGN RHSVecT(RHS);                                      \
    } else {                                                                   \
      return *this OPASSIGN Derived((RawTy)RHS);                               \
    }                                                                          \
  }
 
// Bitwise operations are defined for simd objects and masks, and both operands
// must be integral
#define __ESIMD_BITWISE_OP_FILTER                                              \
  std::is_integral_v<element_type> &&std::is_integral_v<T1>
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(^, ^=, __ESIMD_BITWISE_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(|, |=, __ESIMD_BITWISE_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(&, &=, __ESIMD_BITWISE_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(%, %=, __ESIMD_BITWISE_OP_FILTER)
#undef __ESIMD_BITWISE_OP_FILTER
 
// Bit shift operations are defined only for simd objects (not for masks), and
// both operands must be integral
#define __ESIMD_SHIFT_OP_FILTER                                                \
  std::is_integral_v<element_type> &&std::is_integral_v<T1>                    \
      &&__ESIMD_DNS::is_simd_type_v<Derived>
 
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(<<, <<=, __ESIMD_SHIFT_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(>>, >>=, __ESIMD_SHIFT_OP_FILTER)
#undef __ESIMD_SHIFT_OP_FILTER
 
// Arithmetic operations are defined only for simd objects, and the second
// operand's element type must be vectorizable. This requirement for 'this'
// is fulfilled, because otherwise 'this' couldn't have been constructed.
#define __ESIMD_ARITH_OP_FILTER                                                \
  __ESIMD_DNS::is_simd_type_v<Derived> &&__ESIMD_DNS::is_vectorizable_v<T1>
 
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(+, +=, __ESIMD_ARITH_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(-, -=, __ESIMD_ARITH_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(*, *=, __ESIMD_ARITH_OP_FILTER)
  __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN(/, /=, __ESIMD_ARITH_OP_FILTER)
#undef __ESIMD_ARITH_OP_FILTER
#undef __ESIMD_DEF_SIMD_OBJ_IMPL_OPASSIGN
 
  // Getter for the test proxy member, if enabled
  __ESIMD_DECLARE_TEST_PROXY_ACCESS
 
private:
  // The underlying data for this vector.
  raw_vector_type M_data;
 
  template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
  ESIMD_INLINE std::enable_if_t<is_simd_flag_type_v<Flags>>
  copy_to_impl(AccessorT acc, TOffset offset) const SYCL_ESIMD_FUNCTION;
  template <int ChunkSize, typename PropertyListT, typename AccessorT,
            typename TOffset>
  ESIMD_INLINE std::enable_if_t<
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_to_impl(AccessorT acc, TOffset offset,
               PropertyListT = {}) const SYCL_ESIMD_FUNCTION;
  template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
  ESIMD_INLINE std::enable_if_t<is_simd_flag_type_v<Flags>>
  copy_from_impl(AccessorT acc, TOffset offset) SYCL_ESIMD_FUNCTION;
  template <int ChunkSize, typename PropertyListT, typename AccessorT,
            typename TOffset>
  ESIMD_INLINE std::enable_if_t<
      ext::oneapi::experimental::is_property_list_v<PropertyListT>>
  copy_from_impl(AccessorT acc, TOffset offset,
                 PropertyListT = {}) SYCL_ESIMD_FUNCTION;
 
protected:
  // The test proxy if enabled
  __ESIMD_DECLARE_TEST_PROXY
 
  void set(const raw_vector_type &Val) {
#ifndef __SYCL_DEVICE_ONLY__
    M_data = Val;
#else
    __esimd_vstore<RawTy, N>(&M_data, Val);
#endif
  }
};
 
} // namespace detail
 
template <>
struct is_simd_flag_type<detail::dqword_element_aligned_tag> : std::true_type {
};
} // namespace ext::intel::esimd
} // namespace _V1
} // namespace sycl