memory.hpp

Go to the documentation of this file.

//==-------------- memory.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 memory-access APIs.
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/ext/intel/esimd/common.hpp>
#include <sycl/ext/intel/esimd/detail/memory_intrin.hpp>
#include <sycl/ext/intel/esimd/detail/types.hpp>
#include <sycl/ext/intel/esimd/detail/util.hpp>
#include <sycl/ext/intel/esimd/memory_properties.hpp>
#include <sycl/ext/intel/esimd/simd.hpp>
#include <sycl/ext/intel/esimd/simd_view.hpp>
#include <sycl/half_type.hpp>
 
#include <algorithm>
#include <cstdint>
 
namespace sycl {
inline namespace _V1 {
namespace ext::intel::esimd {
 
 
 
 
 
 
 
template <typename AccessorTy>
__ESIMD_API SurfaceIndex get_surface_index(AccessorTy acc) {
  if constexpr (std::is_same_v<detail::LocalAccessorMarker, AccessorTy> ||
                sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>) {
    return detail::SLM_BTI;
  } else {
#ifdef __ESIMD_FORCE_STATELESS_MEM
    static_assert(sycl::detail::acc_properties::is_image_accessor_v<AccessorTy>,
                  "The function get_surface_index() is available only for "
                  "image- and local-accessors in stateless-only memory mode. "
                  "Consider using "
                  "-fno-sycl-esimd-force-stateless-mem compilation switch.");
#endif // __ESIMD_FORCE_STATELESS_MEM
    return __esimd_get_surface_index(
        detail::AccessorPrivateProxy::getQualifiedPtrOrImageObj(acc));
  }
}
 
namespace detail {
 
// Format u8 and u16 to u8u32 and u16u32 by doing garbage-extension.
template <typename RT, typename T, int N>
ESIMD_INLINE simd<RT, N> lsc_format_input(simd<T, N> Vals) {
  if constexpr (sizeof(T) == 1) {
    // Extend bytes to RT.
    return Vals.template bit_cast_view<uint8_t>();
  } else if constexpr (sizeof(T) == 2) {
    // Extend words to RT.
    return Vals.template bit_cast_view<uint16_t>();
  } else {
    return Vals.template bit_cast_view<RT>();
  }
}
 
// Format u8u32 and u16u32 back to u8 and u16.
template <typename T, typename T1, int N>
ESIMD_INLINE simd<T, N> lsc_format_ret(simd<T1, N> Vals) {
  auto Formatted = Vals.template bit_cast_view<T>();
  if constexpr (sizeof(T) == sizeof(T1)) {
    return Formatted;
  } else {
    constexpr int Stride = Formatted.length / N;
    return Formatted.template select<N, Stride>(0);
  }
}
 
template <typename PropertyListT, cache_level Level>
constexpr cache_hint getCacheHintForIntrin() {
  static_assert(Level == cache_level::L1 || Level == cache_level::L2,
                "ESIMD/GENX intrinsics accept only L1/L2 cache hints");
  if constexpr (Level == cache_level::L1) {
    return getPropertyValue<PropertyListT, cache_hint_L1_key>(cache_hint::none);
  } else {
    return getPropertyValue<PropertyListT, cache_hint_L2_key>(cache_hint::none);
  }
}
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          int N, typename OffsetT>
__ESIMD_API simd<T, N * NElts> gather_impl(const T *p, simd<OffsetT, N> offsets,
                                           simd_mask<N> pred,
                                           simd<T, N * NElts> pass_thru) {
  static_assert(std::is_integral_v<OffsetT>, "Unsupported offset type");
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<NElts>();
  constexpr auto Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<uintptr_t, N> Addrs = reinterpret_cast<uintptr_t>(p);
  Addrs += convert<uintptr_t>(offsets);
  simd<MsgT, N * NElts> PassThruExpanded = lsc_format_input<MsgT>(pass_thru);
  simd<MsgT, N * NElts> Result =
      __esimd_lsc_load_merge_stateless<MsgT, L1H, L2H, AddressScale, ImmOffset,
                                       EDS, VS, Transposed, N>(
          pred.data(), Addrs.data(), PassThruExpanded.data());
  return lsc_format_ret<T>(Result);
}
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          int N, typename Toffset>
__ESIMD_API void scatter_impl(T *p, simd<Toffset, N> offsets,
                              simd<T, N * NElts> vals, simd_mask<N> pred) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::store, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p);
  addrs += convert<uintptr_t>(offsets);
  simd<MsgT, N * NElts> Tmp = lsc_format_input<MsgT, T>(vals);
  __esimd_lsc_store_stateless<MsgT, L1H, L2H, AddressScale, ImmOffset, EDS, VS,
                              Transposed, N>(pred.data(), addrs.data(),
                                             Tmp.data());
}
 
// Returns true iff it is Ok to use llvm.masked.gather and llvm.masked.scatter.
// By default (without use specifying __ESIMD_GATHER_SCATTER_LLVM_IR) it is
// not used because of an issue in GPU driver, which does not recognize
// those operations in SPIR-V when they are used in mixed (scalar and vector)
// kernels using invoke_simd() API.
constexpr bool isMaskedGatherScatterLLVMAvailable() {
#ifdef __ESIMD_GATHER_SCATTER_LLVM_IR
  return true;
#else
  return false;
#endif
}
 
} // namespace detail
 
 
#ifndef __ESIMD_GATHER_SCATTER_LLVM_IR
#endif // __ESIMD_GATHER_SCATTER_LLVM_IR
template <typename T, int N, int VS, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  static_assert(std::is_integral_v<OffsetT>, "Unsupported offset type");
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "gather() requires at least element-size alignment");
 
  // Use LSC lowering if cache-hints are used or VS > 1. Also, if
  // llvm.masked.gather is not available, then LSC is the only lowering option.
  if constexpr (detail::has_cache_hints<PropertyListT>() || VS > 1 ||
                !detail::isMaskedGatherScatterLLVMAvailable()) {
    static_assert(VS == 1 || sizeof(T) >= 4,
                  "VS > 1 is supprted only for 4- and 8-byte elements");
    return detail::gather_impl<T, VS, detail::lsc_data_size::default_size,
                               PropertyListT>(p, byte_offsets, mask, pass_thru);
  } else {
    simd<uint64_t, N> Addrs(reinterpret_cast<uint64_t>(p));
    Addrs = Addrs + convert<uint64_t>(byte_offsets);
 
    using MsgT = detail::__raw_t<T>;
    return __esimd_gather_ld<MsgT, N, Alignment>(
        Addrs.data(), mask.data(),
        sycl::bit_cast<__ESIMD_DNS::vector_type_t<MsgT, N>>(pass_thru.data()));
  }
}
 
template <typename T, int N, int VS, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask,
       PropertyListT props = {}) {
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "gather() requires at least element-size alignment");
 
  if constexpr (detail::has_cache_hints<PropertyListT>() || VS > 1 ||
                detail::isMaskedGatherScatterLLVMAvailable() ||
                !detail::isPowerOf2(N, 32)) {
    simd<T, N> PassThru; // it is intentionally undefined
    return gather<T, N, VS>(p, byte_offsets, mask, PassThru, props);
  } else {
    simd<uintptr_t, N> Addrs = reinterpret_cast<uintptr_t>(p);
    Addrs += convert<uintptr_t>(byte_offsets);
    using MsgT = detail::__raw_t<T>;
    if constexpr (sizeof(T) == 1) {
      auto Ret = __esimd_svm_gather<MsgT, N, detail::ElemsPerAddrEncoding<4>(),
                                    detail::ElemsPerAddrEncoding<1>()>(
          Addrs.data(), mask.data());
      detail::check_rdregion_params<N * 4, N, /*VS*/ 0, N, 4>();
      return __esimd_rdregion<MsgT, N * 4, N, /*VS*/ 0, N, 4>(Ret, 0);
    } else if constexpr (sizeof(T) == 2) {
      auto Ret = __esimd_svm_gather<MsgT, N, detail::ElemsPerAddrEncoding<2>(),
                                    detail::ElemsPerAddrEncoding<2>()>(
          Addrs.data(), mask.data());
      detail::check_rdregion_params<N * 2, N, /*VS*/ 0, N, 2>();
      return __esimd_rdregion<MsgT, N * 2, N, /*VS*/ 0, N, 2>(Ret, 0);
    } else {
      return __esimd_svm_gather<MsgT, N, detail::ElemsPerAddrEncoding<1>(),
                                detail::ElemsPerAddrEncoding<1>()>(Addrs.data(),
                                                                   mask.data());
    }
  }
}
 
template <typename T, int N, int VS, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N / VS> byte_offsets,
       PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  return gather<T, N, VS>(p, byte_offsets, Mask, props);
}
 
template <typename T, int N, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  constexpr int VS = 1;
  return gather<T, N, VS>(p, byte_offsets, mask, pass_thru, props);
}
 
template <typename T, int N, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask,
       PropertyListT props = {}) {
  constexpr int VS = 1;
  return gather<T, N, VS>(p, byte_offsets, mask, props);
}
 
template <typename T, int N, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}) {
  constexpr int VS = 1;
  return gather<T, N, VS>(p, byte_offsets, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  return gather<T, N, VS>(p, byte_offsets.read(), mask, pass_thru, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       PropertyListT props = {}) {
  return gather<T, N, VS>(p, byte_offsets.read(), mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}) {
  return gather<T, N, VS>(p, byte_offsets.read(), props);
}
 
template <typename Tx, int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, simd<Tx, N>>
gather(const Tx *p, Toffset offset, simd_mask<N> mask = 1) {
  return gather<Tx, N>(p, simd<Toffset, N>(offset), mask);
}
 
 
 
 
 
 
template <typename T, int N, int VS = 1, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
  static_assert(std::is_integral_v<OffsetT>, "Unsupported offset type");
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "scatter() requires at least element-size alignment");
 
  // Use LSC lowering if cache-hints are used or VS > 1.
  if constexpr (detail::has_cache_hints<PropertyListT>() || VS > 1 ||
                (!__ESIMD_DNS::isPowerOf2(N, 32) &&
                 !detail::isMaskedGatherScatterLLVMAvailable())) {
    static_assert(VS == 1 || sizeof(T) >= 4,
                  "VS > 1 is supprted only for 4- and 8-byte elements");
    return detail::scatter_impl<T, VS, detail::lsc_data_size::default_size,
                                PropertyListT>(p, byte_offsets, vals, mask);
  } else if constexpr (detail::isMaskedGatherScatterLLVMAvailable()) {
    simd<uint64_t, N> Addrs(reinterpret_cast<uint64_t>(p));
    Addrs = Addrs + convert<uint64_t>(byte_offsets);
    using MsgT = detail::__raw_t<T>;
    __esimd_scatter_st<MsgT, N, Alignment>(
        sycl::bit_cast<__ESIMD_DNS::vector_type_t<MsgT, N>>(vals.data()),
        Addrs.data(), mask.data());
  } else {
    using Tx = detail::__raw_t<T>;
    simd<uint64_t, N> byte_offsets_i = convert<uint64_t>(byte_offsets);
    simd<uint64_t, N> addrs(reinterpret_cast<uint64_t>(p));
    addrs = addrs + byte_offsets_i;
    if constexpr (sizeof(T) == 1) {
      detail::check_wrregion_params<N * 4, N, /*VS*/ 0, N, 4>();
      simd<T, N * 4> D = __esimd_wrregion<Tx, N * 4, N, /*VS*/ 0, N, 4>(
          D.data(), vals.data(), 0);
      __esimd_svm_scatter<Tx, N, detail::ElemsPerAddrEncoding<4>(),
                          detail::ElemsPerAddrEncoding<1>()>(
          addrs.data(), D.data(), mask.data());
    } else if constexpr (sizeof(T) == 2) {
      detail::check_wrregion_params<N * 2, N, /*VS*/ 0, N, 2>();
      simd<Tx, N * 2> D = __esimd_wrregion<Tx, N * 2, N, /*VS*/ 0, N, 2>(
          D.data(), vals.data(), 0);
      __esimd_svm_scatter<Tx, N, detail::ElemsPerAddrEncoding<2>(),
                          detail::ElemsPerAddrEncoding<2>()>(
          addrs.data(), D.data(), mask.data());
    } else
      __esimd_svm_scatter<Tx, N, detail::ElemsPerAddrEncoding<1>(),
                          detail::ElemsPerAddrEncoding<1>()>(
          addrs.data(), vals.data(), mask.data());
  }
}
 
// template <typename T, int N, int VS = 1, typename OffsetT,
//    typename PropertyListT = empty_properties_t>
// void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals,
//  PropertyListT props = {});                               // (usm-sc-2)
template <typename T, int N, int VS = 1, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(p, byte_offsets, vals, Mask, props);
}
 
// template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
//    typename PropertyListT = empty_properties_t>
// void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals,
//  simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-3)
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
  scatter<T, N, VS>(p, byte_offsets.read(), vals, mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(p, byte_offsets.read(), vals, Mask, props);
}
 
template <typename Tx, int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && N == 1>
scatter(Tx *p, Toffset offset, simd<Tx, N> vals, simd_mask<N> mask = 1) {
  scatter<Tx, N, 1>(p, simd<Toffset, N>(offset), vals, mask);
}
 
namespace detail {
// Accessors may get either 32-bit offset or 64-bit depending on
// the -fsycl-esimd-force-stateles-mem mode setting.
#ifdef __ESIMD_FORCE_STATELESS_MEM
using DeviceAccessorOffsetT = uint64_t;
#else
using DeviceAccessorOffsetT = uint32_t;
#endif
 
template <typename T, int NElts, typename PropertyListT>
__ESIMD_API std::enable_if_t<is_property_list_v<PropertyListT>, simd<T, NElts>>
block_load_impl(const T *p, simd_mask<1> pred, simd<T, NElts> pass_thru) {
  // Verify input template arguments.
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr size_t Alignment =
      PropertyListT::template get_property<alignment_key>().value;
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
 
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size ActualDS =
      Use64BitData ? lsc_data_size::u64 : lsc_data_size::u32;
  constexpr lsc_vector_size VS = to_lsc_vector_size<FactoredNElts>();
  constexpr auto Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uintptr_t, N> Addrs = reinterpret_cast<uintptr_t>(p);
  simd<LoadElemT, FactoredNElts> PassThru =
      pass_thru.template bit_cast_view<LoadElemT>();
  simd<LoadElemT, FactoredNElts> Result =
      __esimd_lsc_load_merge_stateless<LoadElemT, L1H, L2H, AddressScale,
                                       ImmOffset, ActualDS, VS, Transposed, N>(
          pred.data(), Addrs.data(), PassThru.data());
  return Result.template bit_cast_view<T>();
}
 
template <typename T, int NElts, typename PropertyListT, typename AccessorT>
__ESIMD_API
    std::enable_if_t<detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_read> &&
                         is_property_list_v<PropertyListT>,
                     simd<T, NElts>>
    block_load_impl(AccessorT acc, DeviceAccessorOffsetT offset,
                    simd_mask<1> pred) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  simd<T, NElts> PassThru; // Intentionally undefined.
  return block_load_impl<T, NElts, PropertyListT>(
      accessorToPointer<T>(acc, offset), pred, PassThru);
#else  // !__ESIMD_FORCE_STATELESS_MEM
  // Verify input template arguments.
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size ActualDS =
      Use64BitData ? lsc_data_size::u64 : lsc_data_size::u32;
  constexpr auto VS = to_lsc_vector_size<FactoredNElts>();
  constexpr auto Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, N> Offsets = offset;
  auto SI = get_surface_index(acc);
  simd<LoadElemT, FactoredNElts> Result =
      __esimd_lsc_load_bti<LoadElemT, L1H, L2H, AddressScale, ImmOffset,
                           ActualDS, VS, Transposed, N>(pred.data(),
                                                        Offsets.data(), SI);
  return Result.template bit_cast_view<T>();
#endif // !__ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int NElts, typename PropertyListT, typename AccessorT>
__ESIMD_API
    std::enable_if_t<detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_read> &&
                         is_property_list_v<PropertyListT>,
                     simd<T, NElts>>
    block_load_impl(AccessorT acc, DeviceAccessorOffsetT offset,
                    simd_mask<1> pred, simd<T, NElts> pass_thru) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return block_load_impl<T, NElts, PropertyListT>(
      accessorToPointer<T>(acc, offset), pred, pass_thru);
#else  // !__ESIMD_FORCE_STATELESS_MEM
  // Verify input template arguments.
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr size_t Alignment =
      PropertyListT::template get_property<alignment_key>().value;
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size ActualDS =
      Use64BitData ? lsc_data_size::u64 : lsc_data_size::u32;
  constexpr auto VS = to_lsc_vector_size<FactoredNElts>();
  constexpr auto Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, N> Offsets = offset;
  auto SI = get_surface_index(acc);
  simd<LoadElemT, FactoredNElts> PassThru =
      pass_thru.template bit_cast_view<LoadElemT>();
  simd<LoadElemT, FactoredNElts> Result =
      __esimd_lsc_load_merge_bti<LoadElemT, L1H, L2H, AddressScale, ImmOffset,
                                 ActualDS, VS, Transposed, N>(
          pred.data(), Offsets.data(), SI, PassThru.data());
  return Result.template bit_cast_view<T>();
#endif // !__ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int NElts, typename PropertyListT>
__ESIMD_API std::enable_if_t<detail::is_property_list_v<PropertyListT>>
block_store_impl(T *p, simd<T, NElts> vals, simd_mask<1> pred) {
  detail::check_cache_hints<cache_action::store, PropertyListT>();
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed store");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can store QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
 
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
 
  check_lsc_vector_size<FactoredNElts>();
 
  using StoreType = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size ActualDS =
      Use64BitData ? lsc_data_size::u64 : lsc_data_size::u32;
  constexpr lsc_vector_size VS = to_lsc_vector_size<FactoredNElts>();
  constexpr auto Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
  simd<uintptr_t, N> Addrs = reinterpret_cast<uintptr_t>(p);
 
  __esimd_lsc_store_stateless<StoreType, L1H, L2H, AddressScale, ImmOffset,
                              ActualDS, VS, Transposed, N>(
      pred.data(), Addrs.data(),
      sycl::bit_cast<__ESIMD_DNS::vector_type_t<StoreType, FactoredNElts>>(
          vals.data()));
}
 
template <typename T, int NElts, typename PropertyListT, typename AccessorT>
__ESIMD_API
    std::enable_if_t<detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_write> &&
                     detail::is_property_list_v<PropertyListT>>
    block_store_impl(AccessorT acc, DeviceAccessorOffsetT offset,
                     simd<T, NElts> vals, simd_mask<1> pred) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  block_store_impl<T, NElts, PropertyListT>(accessorToPointer<T>(acc, offset),
                                            vals, pred);
#else
  // Verify input template arguments.
  check_cache_hints<cache_action::store, PropertyListT>();
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > static_cast<size_t>(1)
          ? sizeof(uint32_t) / sizeof(T)
          : static_cast<size_t>(1);
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed store");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can store QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using StoreElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size ActualDS =
      Use64BitData ? lsc_data_size::u64 : lsc_data_size::u32;
  constexpr auto VS = to_lsc_vector_size<FactoredNElts>();
  constexpr auto Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, N> Offsets = offset;
  auto SI = get_surface_index(acc);
 
  __esimd_lsc_store_bti<StoreElemT, L1H, L2H, AddressScale, ImmOffset, ActualDS,
                        VS, Transposed, N>(
      pred.data(), Offsets.data(),
      sycl::bit_cast<__ESIMD_DNS::vector_type_t<StoreElemT, FactoredNElts>>(
          vals.data()),
      SI);
#endif
}
 
} // namespace detail
 
template <typename Tx, int N,
          typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags>>
block_store(Tx *addr, simd<Tx, N> vals, Flags) {
  using T = typename detail::__raw_t<Tx>;
  using VecT = typename simd<T, N>::raw_vector_type;
  constexpr size_t Align = Flags::template alignment<simd<T, N>>;
  __esimd_svm_block_st<T, N, Align>(reinterpret_cast<VecT *>(addr),
                                    vals.data());
}
 
 
 
 
 
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
block_load(const T *ptr, PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  if constexpr (detail::has_cache_hints<PropertyListT>()) {
    simd<T, N> PassThru; // Intentionally undefined.
    simd_mask<1> Mask = 1;
    return detail::block_load_impl<T, N, NewPropertyListT>(ptr, Mask, PassThru);
  } else {
    constexpr size_t Alignment =
        NewPropertyListT::template get_property<alignment_key>().value;
    return block_load<T, N>(ptr, overaligned_tag<Alignment>{});
  }
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
block_load(const T *ptr, size_t byte_offset, PropertyListT props = {}) {
  const T *AdjustedPtr = reinterpret_cast<const T *>(
      reinterpret_cast<const int8_t *>(ptr) + byte_offset);
  return block_load<T, N>(AdjustedPtr, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API
    std::enable_if_t<detail::is_property_list_v<PropertyListT>, simd<T, N>>
    block_load(const T *ptr, simd_mask<1> pred, PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  simd<T, N> PassThru; // Intentionally uninitialized.
  return detail::block_load_impl<T, N, NewPropertyListT>(ptr, pred, PassThru);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
block_load(const T *ptr, size_t byte_offset, simd_mask<1> pred,
           PropertyListT props = {}) {
  const T *AdjustedPtr = reinterpret_cast<const T *>(
      reinterpret_cast<const int8_t *>(ptr) + byte_offset);
  return block_load<T, N>(AdjustedPtr, pred, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
block_load(const T *ptr, simd_mask<1> pred, simd<T, N> pass_thru,
           PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  return detail::block_load_impl<T, N, NewPropertyListT>(ptr, pred, pass_thru);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
block_load(const T *ptr, size_t byte_offset, simd_mask<1> pred,
           simd<T, N> pass_thru, PropertyListT props = {}) {
  const T *AdjustedPtr = reinterpret_cast<const T *>(
      reinterpret_cast<const int8_t *>(ptr) + byte_offset);
  return block_load<T, N>(AdjustedPtr, pred, pass_thru, props);
}
 
template <typename Tx, int N,
          typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags>, simd<Tx, N>>
block_load(const Tx *addr, Flags) {
  using T = typename detail::__raw_t<Tx>;
  using VecT = typename simd<T, N>::raw_vector_type;
  constexpr size_t Align = Flags::template alignment<simd<T, N>>;
  return __esimd_svm_block_ld<T, N, Align>(
      reinterpret_cast<const VecT *>(addr));
}
 
template <typename Tx, int N, typename AccessorTy,
          typename Flags = vector_aligned_tag,
          typename = std::enable_if_t<
              is_simd_flag_type_v<Flags> &&
              detail::is_device_accessor_with_v<
                  AccessorTy, detail::accessor_mode_cap::can_read>>,
          class T = detail::__raw_t<Tx>>
__ESIMD_API simd<Tx, N> block_load(AccessorTy acc,
                                   detail::DeviceAccessorOffsetT byte_offset,
                                   Flags flags) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return block_load<Tx, N>(__ESIMD_DNS::accessorToPointer<Tx>(acc, byte_offset),
                           flags);
#else
  std::ignore = flags;
  constexpr unsigned Sz = sizeof(T) * N;
  static_assert(Sz >= detail::OperandSize::OWORD,
                "block size must be at least 1 oword");
  static_assert(Sz % detail::OperandSize::OWORD == 0,
                "block size must be whole number of owords");
  static_assert(detail::isPowerOf2(Sz / detail::OperandSize::OWORD),
                "block must be 1, 2, 4 or 8 owords long");
  static_assert(Sz <= 8 * detail::OperandSize::OWORD,
                "block size must be at most 8 owords");
 
  auto surf_ind = __esimd_get_surface_index(
      detail::AccessorPrivateProxy::getQualifiedPtrOrImageObj(acc));
 
  if constexpr (Flags::template alignment<simd<T, N>> >=
                detail::OperandSize::OWORD) {
    return __esimd_oword_ld<T, N>(surf_ind, byte_offset >> 4);
  } else {
    return __esimd_oword_ld_unaligned<T, N>(surf_ind, byte_offset);
  }
#endif
}
 
 
 
 
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, detail::DeviceAccessorOffsetT byte_offset,
           PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return block_load<T, N>(detail::accessorToPointer<T>(acc, byte_offset),
                          props);
#else  // !__ESIMD_FORCE_STATELESS_MEM
  // If the alignment property is not passed, then assume the pointer
  // is element-aligned.
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
 
  // Legacy surface index loads must be 1, 2, 4 or 8 owords long.
  constexpr size_t Size = sizeof(T) * N;
  constexpr size_t OWord = detail::OperandSize::OWORD;
  constexpr bool IsLegacySize = Size == OWord || Size == 2 * OWord ||
                                Size == 4 * OWord || Size == 8 * OWord;
 
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  if constexpr (detail::has_cache_hints<PropertyListT>() || !IsLegacySize) {
    return detail::block_load_impl<T, N, NewPropertyListT>(acc, byte_offset,
                                                           simd_mask<1>(1));
  } else {
    constexpr size_t Alignment =
        NewPropertyListT::template get_property<alignment_key>().value;
    return block_load<T, N>(acc, byte_offset, overaligned_tag<Alignment>{});
  }
#endif // !__ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, PropertyListT /* props */ = {}) {
  // Create new properties without the alignment property passed in 'props',
  // and add alignment<16> as it is usable and most favourable in this case.
  using NewPropertyListT =
      detail::add_or_replace_alignment_property_t<PropertyListT, 16>;
  return block_load<T, N>(acc, 0, NewPropertyListT{});
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, detail::DeviceAccessorOffsetT byte_offset,
           simd_mask<1> pred, simd<T, N> pass_thru,
           PropertyListT /* props */ = {}) {
  // If the alignment property is not passed, then assume the byte_offset
  // is element-aligned and is at least 4-bytes.
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  return detail::block_load_impl<T, N, NewPropertyListT>(acc, byte_offset, pred,
                                                         pass_thru);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, detail::DeviceAccessorOffsetT byte_offset,
           simd_mask<1> pred, PropertyListT props = {}) {
  simd<T, N> PassThru; // Intentionally uninitialized.
  return block_load<T, N>(acc, byte_offset, pred, PassThru, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, simd_mask<1> pred, simd<T, N> pass_thru,
           PropertyListT /* props */ = {}) {
  // Create new properties without the alignment property passed in 'props',
  // and add alignment<16> as it is usable and most favourable in this case.
  using NewPropertyListT =
      detail::add_or_replace_alignment_property_t<PropertyListT, 16>;
  return block_load<T, N>(acc, 0, pred, pass_thru, NewPropertyListT{});
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_device_accessor_with_v<AccessorT,
                                          detail::accessor_mode_cap::can_read>,
    simd<T, N>>
block_load(AccessorT acc, simd_mask<1> pred, PropertyListT /* props */ = {}) {
  // Create new properties without the alignment property passed in 'props',
  // and add alignment<16> as it is usable and most favourable in this case.
  using NewPropertyListT =
      detail::add_or_replace_alignment_property_t<PropertyListT, 16>;
  simd<T, N> PassThru; // Intentionally uninitialized.
  return block_load<T, N>(acc, 0, pred, PassThru, NewPropertyListT{});
}
 
 
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<detail::is_property_list_v<PropertyListT>>
block_store(T *ptr, simd<T, N> vals, PropertyListT /* props */ = {}) {
  if constexpr (detail::has_cache_hints<PropertyListT>()) {
    constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
    using NewPropertyListT =
        detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
    simd_mask<1> Mask = 1;
    detail::block_store_impl<T, N, NewPropertyListT>(ptr, vals, Mask);
  } else {
    // If the alignment property is not passed, then assume the pointer
    // is OWORD-aligned.
    constexpr size_t Alignment =
        detail::getPropertyValue<PropertyListT, alignment_key>(
            detail::OperandSize::OWORD);
    block_store<T, N>(ptr, vals, overaligned_tag<Alignment>{});
  }
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(T *ptr, size_t byte_offset, simd<T, N> vals,
            PropertyListT props = {}) {
  T *AdjustedPtr =
      reinterpret_cast<T *>(reinterpret_cast<int8_t *>(ptr) + byte_offset);
  block_store<T, N>(AdjustedPtr, vals, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<detail::is_property_list_v<PropertyListT>>
block_store(T *ptr, simd<T, N> vals, simd_mask<1> pred,
            PropertyListT /* props */ = {}) {
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  detail::block_store_impl<T, N, NewPropertyListT>(ptr, vals, pred);
}
 
//  the minimally required element-size alignment otherwise.
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(T *ptr, size_t byte_offset, simd<T, N> vals, simd_mask<1> pred,
            PropertyListT props = {}) {
  T *AdjustedPtr =
      reinterpret_cast<T *>(reinterpret_cast<int8_t *>(ptr) + byte_offset);
  block_store<T, N>(AdjustedPtr, vals, pred, props);
}
 
 
 
 
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_write>>
block_store(AccessorT acc, detail::DeviceAccessorOffsetT byte_offset,
            simd<T, N> vals, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  block_store<T, N>(detail::accessorToPointer<T>(acc, byte_offset), vals,
                    props);
#else
  constexpr int DefaultLSCAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(
          DefaultLSCAlignment);
  constexpr bool AlignmentRequiresLSC =
      PropertyListT::template has_property<alignment_key>() && Alignment < 16;
  using Tx = detail::__raw_t<T>;
  constexpr unsigned Sz = sizeof(Tx) * N;
  constexpr bool SzRequiresLSC =
      Sz < detail::OperandSize::OWORD || Sz % detail::OperandSize::OWORD != 0 ||
      !detail::isPowerOf2(Sz / detail::OperandSize::OWORD) ||
      Sz > 8 * detail::OperandSize::OWORD;
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                AlignmentRequiresLSC || SzRequiresLSC) {
    using NewPropertyListT =
        detail::add_alignment_property_t<PropertyListT, DefaultLSCAlignment>;
    simd_mask<1> Mask = 1;
    detail::block_store_impl<T, N, NewPropertyListT>(acc, byte_offset, vals,
                                                     Mask);
  } else {
    auto surf_ind = __esimd_get_surface_index(
        detail::AccessorPrivateProxy::getQualifiedPtrOrImageObj(acc));
    __esimd_oword_st<Tx, N>(surf_ind, byte_offset >> 4, vals.data());
  }
#endif
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_write>>
block_store(AccessorT acc, simd<T, N> vals, PropertyListT props = {}) {
  // Create new properties without the alignment property passed in 'props',
  // and add alignment<16> as it is usable and most favourable in this case.
  using NewPropertyListT =
      detail::add_or_replace_alignment_property_t<PropertyListT, 16>;
  block_store<T, N>(acc, 0, vals, NewPropertyListT{});
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_write>>
block_store(AccessorT acc, detail::DeviceAccessorOffsetT byte_offset,
            simd<T, N> vals, simd_mask<1> pred, PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = (sizeof(T) <= 4) ? 4 : sizeof(T);
  using NewPropertyListT =
      detail::add_alignment_property_t<PropertyListT, DefaultAlignment>;
  detail::block_store_impl<T, N, NewPropertyListT>(acc, byte_offset, vals,
                                                   pred);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_write>>
block_store(AccessorT acc, simd<T, N> vals, simd_mask<1> pred,
            PropertyListT props = {}) {
  // Create new properties without the alignment property passed in 'props',
  // and add alignment<16> as it is usable and most favourable in this case.
  using NewPropertyListT =
      detail::add_or_replace_alignment_property_t<PropertyListT, 16>;
  block_store<T, N>(acc, 0, vals, pred, NewPropertyListT{});
}
 
 
 
 
// Implementations of accessor-based gather and scatter functions
namespace detail {
template <typename T, int N, typename AccessorTy>
ESIMD_INLINE ESIMD_NODEBUG std::enable_if_t<
    std::is_same_v<detail::LocalAccessorMarker, AccessorTy> ||
    is_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write>>
scatter_impl(AccessorTy acc, simd<T, N> vals, simd<uint32_t, N> offsets,
             uint32_t glob_offset, simd_mask<N> mask) {
 
  static_assert(detail::isPowerOf2(N, 32), "Unexpected vector length");
  if constexpr (sizeof(T) == 8) {
    scatter_impl<uint32_t, N>(
        acc, vals.template bit_cast_view<uint32_t>().template select<N, 2>(0),
        offsets, glob_offset, mask);
    scatter_impl<uint32_t, N>(
        acc, vals.template bit_cast_view<uint32_t>().template select<N, 2>(1),
        offsets, glob_offset + sizeof(uint32_t), mask);
  } else {
    constexpr int TypeSizeLog2 = detail::ElemsPerAddrEncoding<sizeof(T)>();
    // TODO (performance) use hardware-supported scale once BE supports it
    constexpr int16_t scale = 0;
    const auto si = __ESIMD_NS::get_surface_index(acc);
 
    if constexpr (sizeof(T) < 4) {
      using Tint = std::conditional_t<std::is_integral_v<T>, T,
                                      detail::uint_type_t<sizeof(T)>>;
      using Treal = __raw_t<T>;
      simd<Tint, N> vals_int = bitcast<Tint, Treal, N>(std::move(vals).data());
      using PromoT = typename std::conditional_t<std::is_signed<Tint>::value,
                                                 int32_t, uint32_t>;
      const simd<PromoT, N> promo_vals = convert<PromoT>(std::move(vals_int));
      __esimd_scatter_scaled<PromoT, N, decltype(si), TypeSizeLog2, scale>(
          mask.data(), si, glob_offset, offsets.data(), promo_vals.data());
    } else {
      using Treal = __raw_t<T>;
      if constexpr (!std::is_same_v<Treal, T>) {
        simd<Treal, N> Values = vals.template bit_cast_view<Treal>();
        __esimd_scatter_scaled<Treal, N, decltype(si), TypeSizeLog2, scale>(
            mask.data(), si, glob_offset, offsets.data(), Values.data());
      } else {
        __esimd_scatter_scaled<T, N, decltype(si), TypeSizeLog2, scale>(
            mask.data(), si, glob_offset, offsets.data(), vals.data());
      }
    }
  }
}
 
#ifndef __ESIMD_FORCE_STATELESS_MEM
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          int N, typename AccessorTy, typename OffsetT>
__ESIMD_API std::enable_if_t<
    is_device_accessor_with_v<AccessorTy, accessor_mode_cap::can_write>>
scatter_impl(AccessorTy acc, simd<OffsetT, N> offsets, simd<T, N * NElts> vals,
             simd_mask<N> pred) {
  static_assert(std::is_integral_v<OffsetT>,
                "Scatter must have integral byte_offset type");
  static_assert(sizeof(OffsetT) <= 4,
                "Implicit truncation of 64-bit byte_offset to 32-bit is "
                "disabled. Use -fsycl-esimd-force-stateless-mem or explicitly "
                "convert offsets to a 32-bit vector");
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::store, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCNElts = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<MsgT, N * NElts> Tmp = lsc_format_input<MsgT, T>(vals);
  simd<uint32_t, N> ByteOffsets32 = convert<uint32_t>(offsets);
  auto si = get_surface_index(acc);
  __esimd_lsc_store_bti<MsgT, L1H, L2H, AddressScale, ImmOffset, EDS, LSCNElts,
                        Transposed, N>(pred.data(), ByteOffsets32.data(),
                                       Tmp.data(), si);
}
#endif // __ESIMD_FORCE_STATELESS_MEM
 
template <typename T, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
    (std::is_same_v<detail::LocalAccessorMarker, AccessorTy> ||
     is_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_read>),
    simd<T, N>>
gather_impl(AccessorTy acc, simd<uint32_t, N> offsets, uint32_t glob_offset,
            simd_mask<N> mask) {
  static_assert(detail::isPowerOf2(N, 32), "Unexpected vector length");
 
  if constexpr (sizeof(T) == 8) {
    simd<T, N> Res;
    Res.template bit_cast_view<uint32_t>().template select<N, 2>(0) =
        gather_impl<uint32_t, N>(acc, offsets, glob_offset, mask);
    Res.template bit_cast_view<uint32_t>().template select<N, 2>(1) =
        gather_impl<uint32_t, N>(acc, offsets, glob_offset + sizeof(uint32_t),
                                 mask);
    return Res;
  } else {
    using Treal = __raw_t<T>;
    constexpr int TypeSizeLog2 = detail::ElemsPerAddrEncoding<sizeof(T)>();
    // TODO (performance) use hardware-supported scale once BE supports it
    constexpr uint32_t scale = 0;
    const auto si = get_surface_index(acc);
    if constexpr (sizeof(T) < 4) {
      using Tint = std::conditional_t<std::is_integral_v<T>, T,
                                      detail::uint_type_t<sizeof(T)>>;
 
      static_assert(std::is_integral<Tint>::value,
                    "only integral 1- & 2-byte types are supported");
      using PromoT = typename std::conditional_t<std::is_signed<Tint>::value,
                                                 int32_t, uint32_t>;
      simd<PromoT, N> promo_vals =
          __esimd_gather_masked_scaled2<PromoT, N, decltype(si), TypeSizeLog2,
                                        scale>(si, glob_offset, offsets.data(),
                                               mask.data());
      auto Res = convert<Tint>(promo_vals);
 
      if constexpr (!std::is_same_v<Tint, T>) {
        return detail::bitcast<Treal, Tint, N>(Res.data());
      } else {
        return Res;
      }
    } else {
      simd<Treal, N> Res = __esimd_gather_masked_scaled2<Treal, N, decltype(si),
                                                         TypeSizeLog2, scale>(
          si, glob_offset, offsets.data(), mask.data());
      if constexpr (!std::is_same_v<Treal, T>) {
        return Res.template bit_cast_view<T>();
      } else {
        return Res;
      }
    }
  }
}
 
#ifndef __ESIMD_FORCE_STATELESS_MEM
template <typename T, int N, int VS, typename PropertyListT, lsc_data_size DS,
          typename OffsetT, typename AccessorT>
__ESIMD_API std::enable_if_t<
    is_device_accessor_with_v<AccessorT, accessor_mode_cap::can_read>,
    simd<T, N>>
gather_impl(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
            simd_mask<N / VS> pred, simd<T, N> pass_thru) {
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
  static_assert(std::is_integral_v<OffsetT>,
                "Gather must have integral byte_offset type");
  static_assert(sizeof(OffsetT) <= 4,
                "Implicit truncation of 64-bit byte_offset to 32-bit is "
                "disabled. Use -fsycl-esimd-force-stateless-mem or explicitly "
                "convert offsets to a 32-bit vector");
  static_assert(VS == 1 || sizeof(T) >= 4,
                "VS > 1 is supprted only for 4- and 8-byte elements");
  check_lsc_vector_size<VS>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<VS>();
  constexpr auto Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  auto SI = get_surface_index(acc);
  simd<uint32_t, N / VS> ByteOffsets32 = convert<uint32_t>(byte_offsets);
  simd<MsgT, N> PassThruExpanded = lsc_format_input<MsgT>(pass_thru);
  simd<MsgT, N> Result =
      __esimd_lsc_load_merge_bti<MsgT, L1H, L2H, AddressScale, ImmOffset, EDS,
                                 LSCVS, Transposed, N / VS>(
          pred.data(), ByteOffsets32.data(), SI, PassThruExpanded.data());
  return lsc_format_ret<T>(Result);
}
#endif // __ESIMD_FORCE_STATELESS_MEM
 
template <typename T, int NElts, lsc_data_size DS, int N>
__ESIMD_API simd<T, N * NElts> slm_gather_impl(simd<uint32_t, N> offsets,
                                               simd_mask<N> pred,
                                               simd<T, N * NElts> pass_thru) {
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<MsgT, N * NElts> PassThruExpanded = lsc_format_input<MsgT>(pass_thru);
  simd<MsgT, N * NElts> Result =
      __esimd_lsc_load_merge_slm<MsgT, cache_hint::none, cache_hint::none,
                                 AddressScale, ImmOffset, EDS, LSCVS,
                                 Transposed, N>(pred.data(), offsets.data(),
                                                PassThruExpanded.data());
  return lsc_format_ret<T>(Result);
}
 
template <typename T, int NElts, lsc_data_size DS, int N>
__ESIMD_API void slm_scatter_impl(simd<uint32_t, N> offsets,
                                  simd<T, N * NElts> vals, simd_mask<N> pred) {
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<MsgT, N * NElts> Tmp = lsc_format_input<MsgT, T>(vals);
  __esimd_lsc_store_slm<MsgT, cache_hint::none, cache_hint::none, AddressScale,
                        ImmOffset, EDS, LSCVS, Transposed, N>(
      pred.data(), offsets.data(), Tmp.data());
}
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          int N, typename Toffset>
__ESIMD_API void prefetch_impl(const T *p, simd<Toffset, N> byte_offsets,
                               simd_mask<N> pred) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::prefetch, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p);
  addrs += convert<uintptr_t>(byte_offsets);
  __esimd_lsc_prefetch_stateless<MsgT, L1H, L2H, AddressScale, ImmOffset, EDS,
                                 LSCVS, Transposed, N>(pred.data(),
                                                       addrs.data());
}
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>>
prefetch_impl(const T *p, Toffset offset, simd_mask<1> pred) {
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::prefetch, PropertyListT>();
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can prefetch QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
 
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = finalize_data_size<LoadElemT, DS>();
 
  static_assert(
      EDS == lsc_data_size::u32 || EDS == lsc_data_size::u64,
      "Transposed prefetch is supported only for data size u32 or u64");
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<FactoredNElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p) + offset;
  __esimd_lsc_prefetch_stateless<LoadElemT, L1H, L2H, AddressScale, ImmOffset,
                                 EDS, LSCVS, Transposed, N>(pred.data(),
                                                            addrs.data());
}
 
#ifndef __ESIMD_FORCE_STATELESS_MEM
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          int N, typename AccessorTy, typename OffsetT>
__ESIMD_API std::enable_if_t<
    is_device_accessor_with_v<AccessorTy, accessor_mode_cap::can_read>>
prefetch_impl(AccessorTy acc, simd<OffsetT, N> byte_offsets,
              simd_mask<N> pred) {
  static_assert(std::is_integral_v<OffsetT>,
                "Prefetch must have integral byte_offset type");
  static_assert(sizeof(OffsetT) <= 4,
                "Implicit truncation of 64-bit byte_offset to 32-bit is "
                "disabled. Use -fsycl-esimd-force-stateless-mem or explicitly "
                "convert offsets to a 32-bit vector");
  check_lsc_vector_size<NElts>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::prefetch, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<NElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  simd<uint32_t, N> ByteOffsets32 = convert<uint32_t>(byte_offsets);
  auto SI = get_surface_index(acc);
  __esimd_lsc_prefetch_bti<MsgT, L1H, L2H, AddressScale, ImmOffset, EDS, LSCVS,
                           Transposed, N>(pred.data(), ByteOffsets32.data(),
                                          SI);
}
 
template <typename T, int NElts, lsc_data_size DS, typename PropertyListT,
          typename AccessorTy, typename OffsetT>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<OffsetT> &&
    is_device_accessor_with_v<AccessorTy, accessor_mode_cap::can_read>>
prefetch_impl(AccessorTy acc, OffsetT byte_offset, simd_mask<1> pred) {
  static_assert(sizeof(OffsetT) <= 4,
                "Implicit truncation of 64-bit byte_offset to 32-bit is "
                "disabled. Use -fsycl-esimd-force-stateless-mem or explicitly "
                "convert offsets to a 32-bit vector");
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::prefetch, PropertyListT>();
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(NElts > 0 && NElts % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (NElts * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (NElts * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || NElts * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredNElts = NElts / SmallIntFactor;
  check_lsc_vector_size<FactoredNElts>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
 
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = finalize_data_size<LoadElemT, DS>();
 
  static_assert(
      EDS == lsc_data_size::u32 || EDS == lsc_data_size::u64,
      "Transposed prefetch is supported only for data size u32 or u64");
  constexpr lsc_vector_size LSCVS = to_lsc_vector_size<FactoredNElts>();
  constexpr lsc_data_order Transposed = lsc_data_order::transpose;
  constexpr int N = 1;
 
  simd<uint32_t, N> offsets = byte_offset;
  auto SI = get_surface_index(acc);
  __esimd_lsc_prefetch_bti<LoadElemT, L1H, L2H, AddressScale, ImmOffset, EDS,
                           LSCVS, Transposed, N>(pred.data(), offsets.data(),
                                                 SI);
}
#endif // __ESIMD_FORCE_STATELESS_MEM
 
// Compute the data size for 2d block load or store.
template <typename T, int NBlocks, int Height, int Width, bool Transposed,
          bool Transformed>
constexpr int get_lsc_block_2d_data_size() {
  if constexpr (Transformed)
    return roundUpNextMultiple<Height, 4 / sizeof(T)>() *
           getNextPowerOf2<Width>() * NBlocks;
  return Width * Height * NBlocks;
}
 
#ifndef __ESIMD_DWORD_BLOCK_2D_WIDTH_SCALE
#define __ESIMD_DWORD_BLOCK_2D_WIDTH_SCALE (1)
#endif
 
#ifndef __ESIMD_BLOCK_2D_WIDTH_CHECK
#define __ESIMD_BLOCK_2D_WIDTH_CHECK(OP, BLOCK_WIDTH, NBLOCKS, SIZE)           \
  static_assert((BLOCK_WIDTH) * (NBLOCKS) * (SIZE) <= 64,                      \
                "Unsupported block width");
#endif
 
enum class block_2d_op { prefetch, load, store };
 
// Compile-time checks for lsc_load_2d/prefetch_2d/store_2d restrictions.
template <typename T, int BlockWidth, int BlockHeight, int NBlocks,
          bool Transposed, bool Transformed, block_2d_op Op>
constexpr void check_lsc_block_2d_restrictions() {
  constexpr int GRFByteSize = BlockWidth * BlockHeight * NBlocks * sizeof(T);
  static_assert(BlockWidth > 0, "Block width must be positive");
  static_assert(BlockHeight > 0, "Block height must be positive");
  // Restrictions based on documentation.
  if constexpr (Op == block_2d_op::store)
    static_assert(GRFByteSize <= 512, "2D store supports 512 bytes max");
  else
    static_assert(GRFByteSize <= 2048,
                  "2D load/prefetch supports 2048 bytes max");
  static_assert(!Transposed || !Transformed,
                "Transposed and transformed is not supported");
  static_assert((sizeof(T) * BlockWidth) % 4 == 0,
                "Block width must be aligned by DW");
  if constexpr (Transposed) {
    static_assert(NBlocks == 1, "Transposed expected to be 1 block only");
    static_assert(sizeof(T) == 4 || sizeof(T) == 8,
                  "Transposed load is supported only for data size u32 or u64");
    static_assert(sizeof(T) == 8 ? BlockHeight == 8
                                 : BlockHeight >= 1 && BlockHeight <= 32,
                  "Unsupported block height");
    static_assert(sizeof(T) == 8
                      ? __ESIMD_DNS::isPowerOf2(BlockWidth, 4)
                      : BlockWidth >= 1 &&
                            BlockWidth <=
                                8 * __ESIMD_DWORD_BLOCK_2D_WIDTH_SCALE,
                  "Unsupported block width");
  } else if constexpr (Transformed) {
    static_assert(sizeof(T) == 1 || sizeof(T) == 2,
                  "VNNI transform is supported only for data size u8 or u16");
    static_assert(__ESIMD_DNS::isPowerOf2(NBlocks, 4),
                  "Unsupported number of blocks");
    static_assert(BlockHeight * sizeof(T) >= 4 && BlockHeight <= 32,
                  "Unsupported block height");
    static_assert(BlockWidth * sizeof(T) >= 4 && BlockWidth <= 16 &&
                      BlockWidth * NBlocks * sizeof(T) <= 64,
                  "Unsupported block width");
  } else {
    if constexpr (Op == block_2d_op::store) {
      static_assert(NBlocks == 1, "Unsupported number of blocks for 2D store");
      static_assert(BlockHeight <= 8, "Unsupported block height for store");
    } else {
      static_assert(
          __ESIMD_DNS::isPowerOf2(NBlocks, sizeof(T) == 1 ? 4 : 8 / sizeof(T)),
          "Unsupported number of blocks for 2D load/prefetch");
      static_assert(BlockHeight <= 32, "Unsupported block height for load");
    }
    static_assert(BlockWidth * sizeof(T) >= 4, "Unsupported block width");
    __ESIMD_BLOCK_2D_WIDTH_CHECK(Op, BlockWidth, NBlocks, sizeof(T));
  }
}
#undef __ESIMD_DWORD_BLOCK_2D_WIDTH_SCALE
#undef __ESIMD_BLOCK_2D_WIDTH_CHECK
 
template <
    typename T, int BlockWidth, int BlockHeight, int NBlocks, bool Transposed,
    bool Transformed, typename PropertyListT,
    int N = get_lsc_block_2d_data_size<__raw_t<T>, NBlocks, BlockHeight,
                                       BlockWidth, Transposed, Transformed>()>
__ESIMD_API simd<T, N> load_2d_impl(const T *Ptr, unsigned SurfaceWidth,
                                    unsigned SurfaceHeight,
                                    unsigned SurfacePitch, int X, int Y) {
 
  check_cache_hints<cache_action::load, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  using RawT = __raw_t<T>;
  check_lsc_block_2d_restrictions<RawT, BlockWidth, BlockHeight, NBlocks,
                                  Transposed, Transformed, block_2d_op::load>();
  // For Load BlockWidth is padded up to the next power-of-two value.
  // For Load with Transpose the pre-operation BlockHeight is padded up
  // to the next power-of-two value.
  // For Load with Transform pre-operation BlockHeight is padded up to
  // multiple of K, where K = 4B / sizeof(T).
  constexpr int ElemsPerDword = 4 / sizeof(RawT);
  constexpr int GRFRowSize = Transposed    ? BlockHeight
                             : Transformed ? BlockWidth * ElemsPerDword
                                           : BlockWidth;
  constexpr int GRFRowPitch = getNextPowerOf2<GRFRowSize>();
  constexpr int GRFColSize =
      Transposed
          ? BlockWidth
          : (Transformed ? (BlockHeight + ElemsPerDword - 1) / ElemsPerDword
                         : BlockHeight);
  constexpr int GRFBlockSize = GRFRowPitch * GRFColSize;
  constexpr int GRFBlockPitch =
      roundUpNextMultiple<64 / sizeof(RawT), GRFBlockSize>();
  constexpr int ActualN = NBlocks * GRFBlockPitch;
 
  constexpr int DstBlockElements = GRFColSize * GRFRowSize;
  constexpr int DstElements = DstBlockElements * NBlocks;
 
  static_assert(N == ActualN || N == DstElements, "Incorrect element count");
  simd_mask<ActualN> Mask = 1;
  constexpr lsc_data_size DS =
      finalize_data_size<RawT, lsc_data_size::default_size>();
  uintptr_t Addr = reinterpret_cast<uintptr_t>(Ptr);
  constexpr lsc_data_order Transpose =
      Transposed ? lsc_data_order::transpose : lsc_data_order::nontranspose;
  simd<RawT, ActualN> Raw =
      __esimd_lsc_load2d_stateless<RawT, L1H, L2H, DS, Transpose, NBlocks,
                                   BlockWidth, BlockHeight, Transformed,
                                   ActualN>(Mask.data(), Addr, SurfaceWidth,
                                            SurfaceHeight, SurfacePitch, X, Y);
 
  if constexpr (ActualN == N) {
    return Raw;
  } else {
    // HW restrictions force data which is read to contain padding filled with
    // zeros for 2d lsc loads. This code eliminates such padding.
 
    // For example, 2D block load of 5 elements of 1 byte data type will
    // take 8 bytes per row for each block.
    //
    // +----+----+----+----+----+----+-----+-----+
    // | 00 | 01 | 02 | 03 | 04 | 05 | 06* | 07* |
    // +----+----+----+----+----+----+-----+-----+
    // | 10 | 11 | 12 | 13 | 14 | 15 | 16* | 17* |
    // +----+----+----+----+----+----+-----+-----+
    // | 20 | 21 | 22 | 23 | 24 | 25 | 26* | 27* |
    // +----+----+----+----+----+----+-----+-----+
    // | 30 | 31 | 32 | 33 | 34 | 35 | 36* | 37* |
    // +----+----+----+----+----+----+-----+-----+
    // * signifies the padded element.
 
    simd<RawT, DstElements> Dst;
 
    for (auto i = 0; i < NBlocks; i++) {
      auto DstBlock =
          Dst.template select<DstBlockElements, 1>(i * DstBlockElements);
 
      auto RawBlock = Raw.template select<GRFBlockSize, 1>(i * GRFBlockPitch);
      DstBlock =
          RawBlock.template bit_cast_view<RawT, GRFColSize, GRFRowPitch>()
              .template select<GRFColSize, 1, GRFRowSize, 1>(0, 0)
              .template bit_cast_view<RawT>();
    }
 
    return Dst;
  }
}
 
template <typename T, int BlockWidth, int BlockHeight, int NBlocks,
          typename PropertyListT,
          int N = get_lsc_block_2d_data_size<__raw_t<T>, NBlocks, BlockHeight,
                                             BlockWidth, false /*Transposed*/,
                                             false /*Transformed*/>()>
__ESIMD_API void prefetch_2d_impl(const T *Ptr, unsigned SurfaceWidth,
                                  unsigned SurfaceHeight, unsigned SurfacePitch,
                                  int X, int Y) {
  using RawT = __raw_t<T>;
  check_cache_hints<cache_action::prefetch, PropertyListT>();
  check_lsc_block_2d_restrictions<RawT, BlockWidth, BlockHeight, NBlocks, false,
                                  false, block_2d_op::prefetch>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr lsc_data_size DS =
      finalize_data_size<RawT, lsc_data_size::default_size>();
  uintptr_t Addr = reinterpret_cast<uintptr_t>(Ptr);
  constexpr lsc_data_order Transpose = lsc_data_order::nontranspose;
  simd_mask<N> Mask = 1;
  __esimd_lsc_prefetch2d_stateless<RawT, L1H, L2H, DS, Transpose, NBlocks,
                                   BlockWidth, BlockHeight, false, N>(
      Mask.data(), Addr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y);
}
 
template <typename T, int BlockWidth, int BlockHeight, typename PropertyListT,
          int N = detail::get_lsc_block_2d_data_size<
              __raw_t<T>, 1u, BlockHeight, BlockWidth, false /*Transposed*/,
              false /*Transformed*/>()>
__ESIMD_API void store_2d_impl(T *Ptr, unsigned SurfaceWidth,
                               unsigned SurfaceHeight, unsigned SurfacePitch,
                               int X, int Y, simd<T, N> Vals) {
  using RawT = __raw_t<T>;
  __ESIMD_DNS::check_cache_hints<__ESIMD_DNS::cache_action::store,
                                 PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  check_lsc_block_2d_restrictions<RawT, BlockWidth, BlockHeight, 1, false,
                                  false, block_2d_op::store>();
  constexpr lsc_data_size DS =
      finalize_data_size<RawT, lsc_data_size::default_size>();
  uintptr_t Addr = reinterpret_cast<uintptr_t>(Ptr);
  constexpr lsc_data_order Transpose = lsc_data_order::nontranspose;
 
  constexpr int Pitch = getNextPowerOf2<BlockWidth>();
  constexpr int NElts = BlockHeight * Pitch;
  simd<RawT, NElts> Raw;
  simd_mask<NElts> Mask = 1;
 
  if constexpr (NElts == N) {
    Raw = Vals;
  } else {
    // For store with padding, allocate the block with padding, and place
    // original data there.
    auto Data2D = Vals.template bit_cast_view<RawT, BlockHeight, BlockWidth>();
    auto Raw2D = Raw.template bit_cast_view<RawT, BlockHeight, Pitch>();
    Raw2D.template select<BlockHeight, 1, BlockWidth, 1>(0, 0) = Data2D;
  }
 
  __esimd_lsc_store2d_stateless<RawT, L1H, L2H, DS, Transpose, 1u, BlockWidth,
                                BlockHeight, false, NElts>(
      Mask.data(), Addr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y,
      Raw.data());
}
 
} // namespace detail
 
 
 
// Dev note: the argument \p glob_offset of this function does not have
// a default value to not conflict with more generic variant (acc-ga-3)
// defined below. This restriction though requires adding an additional
// variant: simd<T, N> gather(acc, glob_offset) to support calls that require
// implicit conversion of a scalar offset to a vector of offsets, e.g.
// 'res = gather<T, N>(acc, 0);'
template <typename T, int N, typename AccessorT>
__ESIMD_API
    std::enable_if_t<detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_read>,
                     simd<T, N>>
    gather(AccessorT acc, simd<detail::DeviceAccessorOffsetT, N> byte_offsets,
           detail::DeviceAccessorOffsetT glob_offset, simd_mask<N> mask = 1) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return gather<T, N>(__ESIMD_DNS::accessorToPointer<T>(acc, glob_offset),
                      byte_offsets, mask);
#else
  if constexpr (!detail::isPowerOf2(N, 32)) {
    // Requires DG2 or PVC.
    simd<T, N> PassThru; // Intentionally undefined
    byte_offsets += glob_offset;
    return detail::gather_impl<T, N, 1,
                               oneapi::experimental::empty_properties_t,
                               detail::lsc_data_size::default_size>(
        acc, byte_offsets, mask, PassThru);
  } else {
    return detail::gather_impl<T, N>(acc, byte_offsets, glob_offset, mask);
  }
#endif // __ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int N, typename AccessorT>
__ESIMD_API
    std::enable_if_t<detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_read>,
                     simd<T, N>>
    gather(AccessorT acc, detail::DeviceAccessorOffsetT glob_offset) {
  simd<detail::DeviceAccessorOffsetT, N> ByteOffsets = 0;
  return gather<T, N>(acc, ByteOffsets, glob_offset);
}
 
#ifdef __ESIMD_FORCE_STATELESS_MEM
template <typename T, int N, typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_read> &&
        std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>,
    simd<T, N>>
gather(AccessorTy acc, simd<Toffset, N> offsets, uint64_t glob_offset,
       simd_mask<N> mask = 1) {
  return gather<T, N>(acc, convert<uint64_t>(offsets), glob_offset, mask);
}
#endif
 
//            typename PropertyListT = empty_properties_t>
 
template <typename T, int N, int VS, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
       simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return gather<T, N, VS>(detail::accessorToPointer<T>(acc), byte_offsets, mask,
                          pass_thru, props);
#else
  return detail::gather_impl<T, N, VS, PropertyListT,
                             detail::lsc_data_size::default_size>(
      acc, byte_offsets, mask, pass_thru);
#endif // __ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int N, int VS, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
       simd_mask<N / VS> mask, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return gather<T, N, VS>(detail::accessorToPointer<T>(acc), byte_offsets, mask,
                          props);
#else
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "gather() requires at least element-size alignment");
 
  if constexpr (detail::has_cache_hints<PropertyListT>() || VS > 1 ||
                !(detail::isPowerOf2(N, 32))) {
    simd<T, N> PassThru; // Intentionally undefined
    return detail::gather_impl<T, N, VS, PropertyListT,
                               detail::lsc_data_size::default_size>(
        acc, byte_offsets, mask, PassThru);
  } else {
    return detail::gather_impl<T, N>(acc, byte_offsets, 0, mask);
  }
#endif // __ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int N, int VS, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
       PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  return gather<T, N, VS>(acc, byte_offsets, Mask, props);
}
 
// Dev note: the mask type was turned into template parameter `MaskT` to
// avoid the conflicts of this prototype with the old gather() function
// accepting a 'global_offset' parameter and avoid 'ambiguous call' errors
// for calls like this: gather(acc, byte_offsets_simd, 0, mask);
template <typename T, int N, typename AccessorT, typename OffsetT,
          typename MaskT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     std::is_same_v<MaskT, simd_mask<N>> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N> byte_offsets, MaskT mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  return gather<T, N, 1>(acc, byte_offsets, mask, pass_thru, props);
}
 
// Dev note: the mask type was turned into template parameter `MaskT` to
// avoid the conflicts of this prototype with the old gather() function
// accepting a 'global_offset' parameter and avoid 'ambiguous call' errors
// for calls like this: gather(acc, byte_offsets_simd, 0);
template <typename T, int N, typename AccessorT, typename OffsetT,
          typename MaskT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     std::is_same_v<MaskT, simd_mask<N>> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N> byte_offsets, MaskT mask,
       PropertyListT props = {}) {
  return gather<T, N, 1>(acc, byte_offsets, mask, props);
}
 
template <typename T, int N, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}) {
  return gather<T, N, 1>(acc, byte_offsets, props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), mask, pass_thru, props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), mask, props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), props);
}
 
 
 
template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  scatter<T, N, VS>(__ESIMD_DNS::accessorToPointer<T>(acc), byte_offsets, vals,
                    mask, props);
#else
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "gather() requires at least element-size alignment");
 
  if constexpr (detail::has_cache_hints<PropertyListT>() || VS > 1 ||
                !detail::isPowerOf2(N, 32)) {
    detail::scatter_impl<T, VS, detail::lsc_data_size::default_size,
                         PropertyListT>(acc, byte_offsets, vals, mask);
  } else {
    detail::scatter_impl<T, N, AccessorTy>(acc, vals, byte_offsets, 0, mask);
  }
 
#endif // __ESIMD_FORCE_STATELESS_MEM
}
template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(acc, byte_offsets, vals, Mask, props);
}
 
template <typename T, int N, int VS = 1, typename AccessorTy,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_write> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
  scatter<T, N, VS>(acc, byte_offsets.read(), vals, mask, props);
}
 
template <typename T, int N, int VS = 1, typename AccessorTy,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_write> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(acc, byte_offsets.read(), vals, Mask, props);
}
 
template <typename T, int N, typename AccessorTy>
__ESIMD_API
    std::enable_if_t<(detail::isPowerOf2(N, 32)) &&
                     detail::is_device_accessor_with_v<
                         AccessorTy, detail::accessor_mode_cap::can_write>>
    scatter(AccessorTy acc, simd<detail::DeviceAccessorOffsetT, N> offsets,
            simd<T, N> vals, detail::DeviceAccessorOffsetT glob_offset,
            simd_mask<N> mask = 1) {
  offsets += glob_offset;
  scatter<T, N>(acc, offsets, vals, mask);
}
 
template <typename T, int N, typename AccessorTy>
__ESIMD_API
    std::enable_if_t<(detail::isPowerOf2(N, 32)) &&
                     detail::is_device_accessor_with_v<
                         AccessorTy, detail::accessor_mode_cap::can_write>>
    scatter(AccessorTy acc, detail::DeviceAccessorOffsetT glob_offset,
            simd<T, N> vals, simd_mask<N> mask = 1) {
  simd<detail::DeviceAccessorOffsetT, N> ByteOffsets = 0;
  scatter<T, N>(acc, ByteOffsets, vals, glob_offset, mask);
}
 
#ifdef __ESIMD_FORCE_STATELESS_MEM
template <typename T, int N, typename AccessorTy, typename Toffset>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorTy,
                                      detail::accessor_mode_cap::can_write> &&
    std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>>
scatter(AccessorTy acc, simd<Toffset, N> offsets, simd<T, N> vals,
        uint64_t glob_offset, simd_mask<N> mask = 1) {
  scatter<T, N, AccessorTy>(acc, convert<uint64_t>(offsets), vals, glob_offset,
                            mask);
}
#endif
 
template <typename T, typename AccessorTy>
__ESIMD_API T scalar_load(AccessorTy acc,
                          detail::DeviceAccessorOffsetT offset) {
  const simd<T, 1> Res =
      gather<T, 1, AccessorTy>(acc, simd<decltype(offset), 1>(offset));
  return Res[0];
}
 
template <typename T, typename AccessorTy>
__ESIMD_API void scalar_store(AccessorTy acc,
                              detail::DeviceAccessorOffsetT offset, T val) {
  scatter<T, 1, AccessorTy>(acc, simd<decltype(offset), 1>(offset),
                            simd<T, 1>(val));
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename Toffset>
__ESIMD_API simd<T, N * get_num_channels_enabled(RGBAMask)>
gather_rgba(const T *p, simd<Toffset, N> offsets, simd_mask<N> mask = 1) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  static_assert((N == 8 || N == 16 || N == 32), "Unsupported value of N");
  static_assert(sizeof(T) == 4, "Unsupported size of type T");
  simd<uint64_t, N> offsets_i = convert<uint64_t>(offsets);
  simd<uint64_t, N> addrs(reinterpret_cast<uint64_t>(p));
  addrs = addrs + offsets_i;
  return __esimd_svm_gather4_scaled<detail::__raw_t<T>, N, RGBAMask>(
      addrs.data(), mask.data());
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename OffsetSimdViewT, typename RegionTy>
__ESIMD_API std::enable_if_t<detail::is_simd_view_type_v<OffsetSimdViewT>,
                             simd<T, N * get_num_channels_enabled(RGBAMask)>>
gather_rgba(const T *p, OffsetSimdViewT offsets, simd_mask<N> mask = 1) {
  return gather_rgba<RGBAMask, T, N>(p, offsets.read(), mask);
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>,
                             simd<T, N * get_num_channels_enabled(RGBAMask)>>
gather_rgba(const T *p, Toffset offset, simd_mask<N> mask = 1) {
  return gather_rgba<RGBAMask, T, N>(p, simd<Toffset, N>(offset), mask);
}
 
namespace detail {
template <rgba_channel_mask M> static void validate_rgba_write_channel_mask() {
  using CM = rgba_channel_mask;
  static_assert(
      (M == CM::ABGR || M == CM::BGR || M == CM::GR || M == CM::R) &&
      "Only ABGR, BGR, GR, R channel masks are valid in write operations");
}
} // namespace detail
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename Toffset>
__ESIMD_API void
scatter_rgba(T *p, simd<Toffset, N> offsets,
             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
             simd_mask<N> mask = 1) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  static_assert((N == 8 || N == 16 || N == 32), "Unsupported value of N");
  static_assert(sizeof(T) == 4, "Unsupported size of type T");
  detail::validate_rgba_write_channel_mask<RGBAMask>();
  simd<uint64_t, N> offsets_i = convert<uint64_t>(offsets);
  simd<uint64_t, N> addrs(reinterpret_cast<uint64_t>(p));
  addrs = addrs + offsets_i;
  __esimd_svm_scatter4_scaled<detail::__raw_t<T>, N, RGBAMask>(
      addrs.data(), vals.data(), mask.data());
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename OffsetSimdViewT, typename RegionTy>
__ESIMD_API std::enable_if_t<detail::is_simd_view_type_v<OffsetSimdViewT>>
scatter_rgba(T *p, OffsetSimdViewT offsets,
             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
             simd_mask<N> mask = 1) {
  scatter_rgba<RGBAMask, T, N>(p, offsets.read(), vals, mask);
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T,
          int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && N == 1>
scatter_rgba(T *p, Toffset offset,
             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
             simd_mask<N> mask = 1) {
  scatter_rgba<RGBAMask, T, N>(p, simd<Toffset, N>(offset), vals, mask);
}
 
template <typename T, int N, rgba_channel_mask RGBAMask>
__SYCL_DEPRECATED("use scatter_rgba<rgba_channel_mask>()")
__ESIMD_API std::
    enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4> scatter_rgba(
        T *p, simd<uint32_t, N> offsets,
        simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
        simd_mask<N> mask = 1) {
  scatter_rgba<RGBAMask>(p, offsets, vals, mask);
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type>
__ESIMD_API
    std::enable_if_t<((N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
                      detail::is_device_accessor_with_v<
                          AccessorT, detail::accessor_mode_cap::can_read>),
                     simd<T, N * get_num_channels_enabled(RGBAMask)>>
    gather_rgba(AccessorT acc, simd<detail::DeviceAccessorOffsetT, N> offsets,
                detail::DeviceAccessorOffsetT global_offset = 0,
                simd_mask<N> mask = 1) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return gather_rgba<RGBAMask>(
      __ESIMD_DNS::accessorToPointer<T>(acc, global_offset), offsets, mask);
#else
  // TODO (performance) use hardware-supported scale once BE supports it
  constexpr uint32_t Scale = 0;
  const auto SI = get_surface_index(acc);
  return __esimd_gather4_masked_scaled2<detail::__raw_t<T>, N, RGBAMask,
                                        decltype(SI), Scale>(
      SI, global_offset, offsets.data(), mask.data());
#endif
}
 
#ifdef __ESIMD_FORCE_STATELESS_MEM
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type, typename Toffset>
__ESIMD_API std::enable_if_t<
    ((N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
     detail::is_device_accessor_with_v<AccessorT,
                                       detail::accessor_mode_cap::can_read> &&
     std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>),
    simd<T, N * get_num_channels_enabled(RGBAMask)>>
gather_rgba(AccessorT acc, simd<Toffset, N> offsets, uint64_t global_offset = 0,
            simd_mask<N> mask = 1) {
  return gather_rgba<RGBAMask, AccessorT, N, T>(acc, convert<uint64_t>(offsets),
                                                global_offset, mask);
}
#endif
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type>
__ESIMD_API
    std::enable_if_t<(N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
                     detail::is_device_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_write>>
    scatter_rgba(AccessorT acc, simd<detail::DeviceAccessorOffsetT, N> offsets,
                 simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
                 detail::DeviceAccessorOffsetT global_offset = 0,
                 simd_mask<N> mask = 1) {
  detail::validate_rgba_write_channel_mask<RGBAMask>();
#ifdef __ESIMD_FORCE_STATELESS_MEM
  scatter_rgba<RGBAMask>(__ESIMD_DNS::accessorToPointer<T>(acc, global_offset),
                         offsets, vals, mask);
#else
  // TODO (performance) use hardware-supported scale once BE supports it
  constexpr uint32_t Scale = 0;
  const auto SI = get_surface_index(acc);
  __esimd_scatter4_scaled<T, N, decltype(SI), RGBAMask, Scale>(
      mask.data(), SI, global_offset, offsets.data(), vals.data());
#endif
}
 
#ifdef __ESIMD_FORCE_STATELESS_MEM
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type, typename Toffset>
__ESIMD_API std::enable_if_t<
    (N == 8 || N == 16 || N == 32) && sizeof(T) == 4 &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_write> &&
    std::is_integral_v<Toffset> && !std::is_same_v<Toffset, uint64_t>>
scatter_rgba(AccessorT acc, simd<Toffset, N> offsets,
             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
             uint64_t global_offset = 0, simd_mask<N> mask = 1) {
  scatter_rgba<RGBAMask, AccessorT, N, T>(acc, convert<uint64_t>(offsets), vals,
                                          global_offset, mask);
}
#endif
 
namespace detail {
 
#ifndef __ESIMD_FP_ATOMIC_OP_TYPE_CHECK
#define __ESIMD_FP_ATOMIC_OP_TYPE_CHECK(T)                                     \
  static_assert(is_type<T, float, sycl::half, double>(),                       \
                "float, double or sycl::half type is expected");
#endif // __ESIMD_FP_ATOMIC_OP_TYPE_CHECK
 
template <__ESIMD_NS::atomic_op Op, typename T, int N, unsigned NumSrc,
          bool IsLSC = false>
constexpr void check_atomic() {
 
  static_assert(sizeof(T) > 1, "Unsupported data type");
 
  // LSC atomic operation is supported for any width.
  if constexpr (!IsLSC)
    static_assert((detail::isPowerOf2(N, 32)),
                  "Execution size 1, 2, 4, 8, 16, 32 are supported");
 
  static_assert(NumSrc == __ESIMD_DNS::get_num_args<Op>(),
                "Wrong number of operands");
  constexpr bool IsInt2BytePlus =
      std::is_integral_v<T> && (sizeof(T) >= sizeof(uint16_t));
 
  if constexpr (Op == __ESIMD_NS::atomic_op::xchg ||
                Op == __ESIMD_NS::atomic_op::cmpxchg ||
                Op == __ESIMD_NS::atomic_op::predec ||
                Op == __ESIMD_NS::atomic_op::inc ||
                Op == __ESIMD_NS::atomic_op::dec) {
 
    static_assert(IsInt2BytePlus, "Integral 16-bit or wider type is expected");
  }
  // FP ops (are always delegated to native::lsc::<Op>)
  if constexpr (Op == __ESIMD_NS::atomic_op::fmax ||
                Op == __ESIMD_NS::atomic_op::fmin ||
                Op == __ESIMD_NS::atomic_op::fadd ||
                Op == __ESIMD_NS::atomic_op::fsub ||
                Op == __ESIMD_NS::atomic_op::fcmpxchg) {
    __ESIMD_FP_ATOMIC_OP_TYPE_CHECK(T);
  }
  if constexpr (Op == __ESIMD_NS::atomic_op::add ||
                Op == __ESIMD_NS::atomic_op::sub ||
                Op == __ESIMD_NS::atomic_op::umin ||
                Op == __ESIMD_NS::atomic_op::umax ||
                Op == __ESIMD_NS::atomic_op::bit_and ||
                Op == __ESIMD_NS::atomic_op::bit_or ||
                Op == __ESIMD_NS::atomic_op::bit_xor ||
                Op == __ESIMD_NS::atomic_op::smin ||
                Op == __ESIMD_NS::atomic_op::smax) {
    static_assert(IsInt2BytePlus, "Integral 16-bit or wider type is expected");
    constexpr bool IsSignedMinmax = (Op == __ESIMD_NS::atomic_op::smin) ||
                                    (Op == __ESIMD_NS::atomic_op::smax);
    constexpr bool IsUnsignedMinmax = (Op == __ESIMD_NS::atomic_op::umin) ||
                                      (Op == __ESIMD_NS::atomic_op::umax);
 
    if constexpr (IsSignedMinmax || IsUnsignedMinmax) {
      constexpr bool SignOK = std::is_signed_v<T> == IsSignedMinmax;
      static_assert(SignOK, "Signed/unsigned integer type expected for "
                            "signed/unsigned min/max operation");
    }
  }
}
#undef __ESIMD_FP_ATOMIC_OP_TYPE_CHECK
} // namespace detail
 
 
template <uint32_t SLMSize> __ESIMD_API void slm_init() {
  __esimd_slm_init(SLMSize);
}
 
//  with esimd::slm_allocator() class.
__ESIMD_API void slm_init(uint32_t size) { __esimd_slm_init(size); }
 
 
#ifndef __ESIMD_GATHER_SCATTER_LLVM_IR
#endif // __ESIMD_GATHER_SCATTER_LLVM_IR
template <typename T, int N, int VS,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask,
           simd<T, N> pass_thru, PropertyListT props = {}) {
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "slm_gather() requires at least element-size alignment");
 
  // Use LSC lowering if VS > 1. Also, if masked gather is
  // not available, then LSC is the only lowering option.
  if constexpr (VS > 1 || !detail::isMaskedGatherScatterLLVMAvailable()) {
    return __ESIMD_DNS::slm_gather_impl<T, VS,
                                        detail::lsc_data_size::default_size>(
        byte_offsets, mask, pass_thru);
  } else {
    if constexpr (sizeof(T) == 8) {
      simd<T, N> Res;
      Res.template bit_cast_view<uint32_t>().template select<N, 2>(0) =
          __esimd_slm_gather_ld<uint32_t, N, Alignment>(
              byte_offsets.data(), mask.data(),
              (pass_thru.template bit_cast_view<uint32_t>()
                   .template select<N, 2>(0))
                  .data());
      simd<uint32_t, N / VS> Offset = byte_offsets + sizeof(uint32_t);
      Res.template bit_cast_view<uint32_t>().template select<N, 2>(1) =
          __esimd_slm_gather_ld<uint32_t, N, sizeof(uint32_t)>(
              Offset.data(), mask.data(),
              (pass_thru.template bit_cast_view<uint32_t>()
                   .template select<N, 2>(1))
                  .data());
      return Res;
    } else {
      using MsgT = detail::__raw_t<T>;
      return __esimd_slm_gather_ld<MsgT, N, Alignment>(
          byte_offsets.data(), mask.data(), pass_thru.data());
    }
  }
}
 
template <typename T, int N, int VS,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask,
           PropertyListT props = {}) {
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "slm_gather() requires at least element-size alignment");
 
  if constexpr (VS > 1 || (!detail::isPowerOf2(N, 32) &&
                           !detail::isMaskedGatherScatterLLVMAvailable())) {
    simd<T, N> PassThru; // Intentionally undefined
    return detail::slm_gather_impl<T, VS, detail::lsc_data_size::default_size>(
        byte_offsets, mask, PassThru);
  } else if constexpr (detail::isMaskedGatherScatterLLVMAvailable()) {
    if constexpr (sizeof(T) == 8) {
      simd<T, N> Res;
      simd<uint32_t, N> PassThru; // it is intentionally undefined
 
      Res.template bit_cast_view<uint32_t>().template select<N, 2>(0) =
          __esimd_slm_gather_ld<uint32_t, N, Alignment>(
              byte_offsets.data(), mask.data(), PassThru.data());
      simd<uint32_t, N / VS> Offset = byte_offsets + sizeof(uint32_t);
      Res.template bit_cast_view<uint32_t>().template select<N, 2>(1) =
          __esimd_slm_gather_ld<uint32_t, N, sizeof(uint32_t)>(
              Offset.data(), mask.data(), PassThru.data());
      return Res;
    } else {
      using MsgT = detail::__raw_t<T>;
      simd<MsgT, N> PassThru; // it is intentionally undefined
      return __esimd_slm_gather_ld<MsgT, N, Alignment>(
          byte_offsets.data(), mask.data(), PassThru.data());
    }
  } else {
    detail::LocalAccessorMarker acc;
    return detail::gather_impl<T, N>(acc, byte_offsets, 0, mask);
  }
}
 
template <typename T, int N, int VS,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  return slm_gather<T, N, VS>(byte_offsets, Mask, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask,
           simd<T, N> pass_thru, PropertyListT props = {}) {
  constexpr int VS = 1;
  return slm_gather<T, N, VS>(byte_offsets, mask, pass_thru, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask,
           PropertyListT props = {}) {
  constexpr int VS = 1;
  return slm_gather<T, N, VS>(byte_offsets, mask, props);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_gather(simd<uint32_t, N> byte_offsets, PropertyListT props = {}) {
  constexpr int VS = 1;
  return slm_gather<T, N, VS>(byte_offsets, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
slm_gather(OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
           simd<T, N> pass_thru, PropertyListT props = {}) {
  return slm_gather<T, N, VS>(byte_offsets.read(), mask, pass_thru, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
slm_gather(OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
           PropertyListT props = {}) {
  return slm_gather<T, N, VS>(byte_offsets.read(), mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
slm_gather(OffsetSimdViewT byte_offsets, PropertyListT props = {}) {
  return slm_gather<T, N, VS>(byte_offsets.read(), props);
}
 
template <typename T> __ESIMD_API T slm_scalar_load(uint32_t offset) {
  const simd<T, 1> Res = slm_gather<T, 1>(simd<uint32_t, 1>(offset));
  return Res[0];
}
 
 
template <typename T, int N, int VS = 1,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals,
            simd_mask<N / VS> mask, PropertyListT props = {}) {
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
 
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(sizeof(T));
  static_assert(Alignment >= sizeof(T),
                "slm_scatter() requires at least element-size alignment");
 
  // Use LSC lowering if VS > 1.
  if constexpr (VS > 1 || (!detail::isPowerOf2(N, 32) &&
                           !detail::isMaskedGatherScatterLLVMAvailable())) {
    __ESIMD_DNS::slm_scatter_impl<T, VS, detail::lsc_data_size::default_size>(
        byte_offsets, vals, mask);
  } else if constexpr (detail::isMaskedGatherScatterLLVMAvailable()) {
    if constexpr (sizeof(T) == 8) {
      __esimd_slm_scatter_st<uint32_t, N, Alignment>(
          vals.template bit_cast_view<uint32_t>()
              .template select<N, 2>(0)
              .data(),
          byte_offsets.data(), mask.data());
      simd<uint32_t, N / VS> Offset = byte_offsets + sizeof(uint32_t);
      __esimd_slm_scatter_st<uint32_t, N, sizeof(uint32_t)>(
          vals.template bit_cast_view<uint32_t>()
              .template select<N, 2>(1)
              .data(),
          Offset.data(), mask.data());
 
    } else {
      using MsgT = detail::__raw_t<T>;
      __esimd_slm_scatter_st<MsgT, N, Alignment>(
          sycl::bit_cast<__ESIMD_DNS::vector_type_t<MsgT, N>>(vals.data()),
          byte_offsets.data(), mask.data());
    }
  } else {
    detail::LocalAccessorMarker acc;
    detail::scatter_impl<T, N>(acc, vals, byte_offsets, 0, mask);
  }
}
 
template <typename T, int N, int VS = 1,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals,
            PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  slm_scatter<T, N, VS>(byte_offsets, vals, Mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_scatter(OffsetSimdViewT byte_offsets, simd<T, N> vals,
            simd_mask<N / VS> mask, PropertyListT props = {}) {
  slm_scatter<T, N, VS>(byte_offsets.read(), vals, mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_scatter(OffsetSimdViewT byte_offsets, simd<T, N> vals,
            PropertyListT props = {}) {
  return slm_scatter<T, N, VS>(byte_offsets.read(), vals, props);
}
 
template <typename T>
__ESIMD_API void slm_scalar_store(uint32_t offset, T val) {
  slm_scatter<T, 1>(simd<uint32_t, 1>(offset), simd<T, 1>(val), 1);
}
 
template <typename T, int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4),
                             simd<T, N * get_num_channels_enabled(RGBAMask)>>
slm_gather_rgba(simd<uint32_t, N> offsets, simd_mask<N> mask = 1) {
  const auto SI = __ESIMD_NS::get_surface_index(detail::LocalAccessorMarker());
  return __esimd_gather4_masked_scaled2<T, N, RGBAMask>(
      SI, 0 /*global_offset*/, offsets.data(), mask.data());
}
 
template <typename T, int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4)>
slm_scatter_rgba(simd<uint32_t, N> offsets,
                 simd<T, N * get_num_channels_enabled(Mask)> vals,
                 simd_mask<N> mask = 1) {
  detail::validate_rgba_write_channel_mask<Mask>();
  const auto si = __ESIMD_NS::get_surface_index(detail::LocalAccessorMarker());
  constexpr int16_t Scale = 0;
  constexpr int global_offset = 0;
  __esimd_scatter4_scaled<T, N, decltype(si), Mask, Scale>(
      mask.data(), si, global_offset, offsets.data(), vals.data());
}
 
template <typename T, int N,
          typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags>, simd<T, N>>
slm_block_load(uint32_t byte_offset, Flags) {
  constexpr size_t Align = Flags::template alignment<simd<T, N>>;
  return __esimd_slm_block_ld<detail::__raw_t<T>, N, Align>(byte_offset);
}
 
 
 
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_block_load(uint32_t byte_offset, PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = detail::OperandSize::OWORD;
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
  return __esimd_slm_block_ld<detail::__raw_t<T>, N, Alignment>(byte_offset);
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_block_load(uint32_t byte_offset, simd_mask<1> pred,
               PropertyListT props = {}) {
  // Verify input template arguments.
  constexpr size_t DefaultAlignment = sizeof(T) <= 4 ? 4 : sizeof(T);
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(N > 0 && N % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (N * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (N * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || N * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredN = N / SmallIntFactor;
  detail::check_lsc_vector_size<FactoredN>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
 
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr detail::lsc_data_size DS =
      Use64BitData ? detail::lsc_data_size::u64 : detail::lsc_data_size::u32;
  constexpr auto VS = detail::to_lsc_vector_size<FactoredN>();
  constexpr auto Transposed = detail::lsc_data_order::transpose;
  constexpr int NLanes = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, NLanes> Offsets = byte_offset;
  simd<LoadElemT, FactoredN> Result =
      __esimd_lsc_load_slm<LoadElemT, cache_hint::none, cache_hint::none,
                           AddressScale, ImmOffset, DS, VS, Transposed, NLanes>(
          pred.data(), Offsets.data());
  return Result.template bit_cast_view<T>();
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
slm_block_load(uint32_t offset, simd_mask<1> pred, simd<T, N> pass_thru,
               PropertyListT props = {}) {
  // Verify input template arguments.
  constexpr size_t DefaultAlignment = sizeof(T) <= 4 ? 4 : sizeof(T);
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
  static_assert(N > 0 && N % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed load");
 
  // If alignment >= 8 and (N * sizeof(T)) % 8 == 0) we can load QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (N * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || N * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredN = N / SmallIntFactor;
  detail::check_lsc_vector_size<FactoredN>();
 
  // Prepare template arguments for the call of intrinsic.
  using LoadElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
 
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr detail::lsc_data_size DS =
      Use64BitData ? detail::lsc_data_size::u64 : detail::lsc_data_size::u32;
  constexpr auto VS = detail::to_lsc_vector_size<FactoredN>();
  constexpr auto Transposed = detail::lsc_data_order::transpose;
  constexpr int NLanes = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, NLanes> Offsets = offset;
  simd<LoadElemT, FactoredN> PassThru =
      pass_thru.template bit_cast_view<LoadElemT>();
  simd<LoadElemT, FactoredN> Result =
      __esimd_lsc_load_merge_slm<LoadElemT, cache_hint::none, cache_hint::none,
                                 AddressScale, ImmOffset, DS, VS, Transposed,
                                 NLanes>(pred.data(), Offsets.data(),
                                         PassThru.data());
  return Result.template bit_cast_view<T>();
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, uint32_t byte_offset, PropertyListT props = {}) {
  byte_offset += detail::localAccessorToOffset(lacc);
  return slm_block_load<T, N>(byte_offset, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, PropertyListT props = {}) {
  return slm_block_load<T, N>(detail::localAccessorToOffset(lacc), props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, uint32_t byte_offset, simd_mask<1> pred,
           PropertyListT props = {}) {
  byte_offset += detail::localAccessorToOffset(lacc);
  return slm_block_load<T, N>(byte_offset, pred, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, simd_mask<1> pred, PropertyListT props = {}) {
  return slm_block_load<T, N>(detail::localAccessorToOffset(lacc), pred, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, uint32_t byte_offset, simd_mask<1> pred,
           simd<T, N> pass_thru, PropertyListT props = {}) {
  byte_offset += __ESIMD_DNS::localAccessorToOffset(lacc);
  return slm_block_load<T, N>(byte_offset, pred, pass_thru, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_read> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
block_load(AccessorT lacc, simd_mask<1> pred, simd<T, N> pass_thru,
           PropertyListT props = {}) {
  return slm_block_load<T, N>(__ESIMD_DNS::localAccessorToOffset(lacc), pred,
                              pass_thru, props);
}
 
template <typename T, int N, typename Flags>
__ESIMD_API std::enable_if_t<is_simd_flag_type_v<Flags>>
slm_block_store(uint32_t offset, simd<T, N> vals, Flags) {
  constexpr size_t Align = Flags::template alignment<simd<T, N>>;
  __esimd_slm_block_st<detail::__raw_t<T>, N, Align>(offset, vals.data());
}
 
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_block_store(uint32_t byte_offset, simd<T, N> vals, simd_mask<1> pred,
                PropertyListT props = {}) {
  // Verify input template arguments.
  constexpr size_t DefaultAlignment = sizeof(T) <= 4 ? 4 : sizeof(T);
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
  static_assert(
      (Alignment >= __ESIMD_DNS::OperandSize::DWORD && sizeof(T) <= 4) ||
          (Alignment >= __ESIMD_DNS::OperandSize::QWORD && sizeof(T) > 4),
      "Incorrect alignment for the data type");
 
  constexpr int SmallIntFactor64Bit = sizeof(uint64_t) / sizeof(T);
  constexpr int SmallIntFactor32Bit =
      sizeof(uint32_t) / sizeof(T) > 1 ? sizeof(uint32_t) / sizeof(T) : 1;
 
  static_assert(N > 0 && N % SmallIntFactor32Bit == 0,
                "Number of elements is not supported by Transposed store");
 
  // If alignment >= 8 and (N * sizeof(T)) % 8 == 0) we can store QWORDs.
  // Don't do it for 4-byte vectors (unless it is greater than 256-bytes),
  // because it would require a bit-cast, which is supposed to be NO-OP, but
  // might confuse GPU BE sometimes. 1- and 2-byte vectors are casted anyways.
  constexpr bool Use64BitData =
      Alignment >= __ESIMD_DNS::OperandSize::QWORD &&
      (N * sizeof(T)) % sizeof(uint64_t) == 0 &&
      (sizeof(T) != sizeof(uint32_t) || N * sizeof(T) > 256);
  constexpr int SmallIntFactor =
      Use64BitData ? SmallIntFactor64Bit : SmallIntFactor32Bit;
  constexpr int FactoredN = N / SmallIntFactor;
  detail::check_lsc_vector_size<FactoredN>();
 
  // Prepare template arguments for the call of intrinsic.
  using StoreElemT = __ESIMD_DNS::__raw_t<
      std::conditional_t<SmallIntFactor == 1, T,
                         std::conditional_t<Use64BitData, uint64_t, uint32_t>>>;
 
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr detail::lsc_data_size DS =
      Use64BitData ? detail::lsc_data_size::u64 : detail::lsc_data_size::u32;
  constexpr auto VS = detail::to_lsc_vector_size<FactoredN>();
  constexpr auto Transposed = detail::lsc_data_order::transpose;
  constexpr int NLanes = 1;
 
  // Prepare non-template arguments and call the intrinsic.
  simd<uint32_t, NLanes> Offsets = byte_offset;
  __esimd_lsc_store_slm<StoreElemT, cache_hint::none, cache_hint::none,
                        AddressScale, ImmOffset, DS, VS, Transposed, NLanes>(
      pred.data(), Offsets.data(),
      sycl::bit_cast<__ESIMD_DNS::vector_type_t<StoreElemT, FactoredN>>(
          vals.data()));
}
 
template <typename T, int N,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
slm_block_store(uint32_t byte_offset, simd<T, N> vals,
                PropertyListT props = {}) {
  constexpr size_t DefaultAlignment = detail::OperandSize::OWORD;
  constexpr size_t Alignment =
      detail::getPropertyValue<PropertyListT, alignment_key>(DefaultAlignment);
  using StoreElemT = detail::__raw_t<T>;
  __esimd_slm_block_st<StoreElemT, N, Alignment>(
      byte_offset,
      sycl::bit_cast<__ESIMD_DNS::vector_type_t<StoreElemT, N>>(vals.data()));
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(AccessorT lacc, uint32_t byte_offset, simd<T, N> vals,
            PropertyListT props = {}) {
  byte_offset += detail::localAccessorToOffset(lacc);
  slm_block_store<T, N>(byte_offset, vals, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(AccessorT lacc, simd<T, N> vals, PropertyListT props = {}) {
  slm_block_store<T, N>(detail::localAccessorToOffset(lacc), vals, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(AccessorT lacc, uint32_t byte_offset, simd<T, N> vals,
            simd_mask<1> pred, PropertyListT props = {}) {
  byte_offset += detail::localAccessorToOffset(lacc);
  slm_block_store<T, N>(byte_offset, vals, pred, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
block_store(AccessorT lacc, simd<T, N> vals, simd_mask<1> pred,
            PropertyListT props = {}) {
  slm_block_store<T, N>(detail::localAccessorToOffset(lacc), vals, pred, props);
}
namespace detail {
 
// lsc_atomic_update() operations may share atomic_op values for data types
// of the same (fp vs integral) class for convenience (e.g. re-use 'fmax' for
// all FP types). In fact those data types may require using different internal
// opcodes. This function returns the corresponding internal opcode for
// the input type 'T' and operation 'Op'.
template <typename T, __ESIMD_NS::atomic_op Op>
constexpr int lsc_to_internal_atomic_op() {
  constexpr __ESIMD_NS::native::lsc::atomic_op LSCOp =
      __ESIMD_DNS::to_lsc_atomic_op<Op>();
  return static_cast<int>(LSCOp);
}
 
 
template <atomic_op Op, typename T, int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 0, simd<T, N>>
slm_atomic_update_impl(simd<uint32_t, N> offsets, simd_mask<N> pred) {
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_slm_0<MsgT, IOp, cache_hint::none, cache_hint::none,
                                AddressScale, ImmOffset, EDS, VS, Transposed,
                                N>(pred.data(), offsets.data());
  return lsc_format_ret<T>(Tmp);
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 1, simd<T, N>>
slm_atomic_update_impl(simd<uint32_t, N> offsets, simd<T, N> src0,
                       simd_mask<N> pred) {
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  if constexpr (std::is_same_v<T, double> || std::is_same_v<T, float>) {
    return __esimd_lsc_xatomic_slm_1<T, IOp, cache_hint::none, cache_hint::none,
                                     AddressScale, ImmOffset, EDS, VS,
                                     Transposed, N>(pred.data(), offsets.data(),
                                                    src0.data());
  } else {
    using MsgT = typename lsc_expand_type<T>::type;
    simd<MsgT, N> Msg_data = lsc_format_input<MsgT>(src0);
    simd<MsgT, N> Tmp =
        __esimd_lsc_xatomic_slm_1<MsgT, IOp, cache_hint::none, cache_hint::none,
                                  AddressScale, ImmOffset, EDS, VS, Transposed,
                                  N>(pred.data(), offsets.data(),
                                     Msg_data.data());
    return lsc_format_ret<T>(Tmp);
  }
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS>
__ESIMD_API simd<T, N> slm_atomic_update_impl(simd<uint32_t, N> offsets,
                                              simd<T, N> src0, simd<T, N> src1,
                                              simd_mask<N> pred) {
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  if constexpr (std::is_same_v<T, double> || std::is_same_v<T, float>) {
    return __esimd_lsc_xatomic_slm_2<T, IOp, cache_hint::none, cache_hint::none,
                                     AddressScale, ImmOffset, EDS, VS,
                                     Transposed, N>(pred.data(), offsets.data(),
                                                    src0.data(), src1.data());
  } else {
    using MsgT = typename lsc_expand_type<T>::type;
    simd<MsgT, N> Msg_data0 = lsc_format_input<MsgT>(src0);
    simd<MsgT, N> Msg_data1 = lsc_format_input<MsgT>(src1);
    simd<MsgT, N> Tmp =
        __esimd_lsc_xatomic_slm_2<MsgT, IOp, cache_hint::none, cache_hint::none,
                                  AddressScale, ImmOffset, EDS, VS, Transposed,
                                  N>(pred.data(), offsets.data(),
                                     Msg_data0.data(), Msg_data1.data());
    return lsc_format_ret<T>(Tmp);
  }
}
 
} // namespace detail
 
 
 
 
 
 
template <atomic_op Op, typename T, int N>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0, simd<T, N>>
slm_atomic_update(simd<uint32_t, N> byte_offset, simd_mask<N> mask = 1) {
  // 2 byte, 8 byte types, non-power of two, and operations wider than
  // 32 are supported only by LSC.
  if constexpr (sizeof(T) == 2 || sizeof(T) == 8 ||
                !__ESIMD_DNS::isPowerOf2(N, 32)) {
    return slm_atomic_update_impl<Op, T, N,
                                  detail::lsc_data_size::default_size>(
        byte_offset, mask);
  } else if constexpr (Op == atomic_op::load) {
    if constexpr (std::is_integral_v<T>) {
      return slm_atomic_update<atomic_op::bit_or, T, N>(byte_offset,
                                                        simd<T, N>(0), mask);
    } else {
      using Tint = detail::uint_type_t<sizeof(T)>;
      simd<Tint, N> Res = slm_atomic_update<atomic_op::bit_or, Tint, N>(
          byte_offset, simd<Tint, N>(0), mask);
      return Res.template bit_cast_view<T>();
    }
  } else {
    detail::check_atomic<Op, T, N, 0>();
    const auto si = get_surface_index(detail::LocalAccessorMarker());
    return __esimd_dword_atomic0<Op, T, N>(mask.data(), si, byte_offset.data());
  }
}
 
template <atomic_op Op, typename T, int N, typename AccessorT>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
                                 __ESIMD_DNS::is_rw_local_accessor_v<AccessorT>,
                             simd<T, N>>
atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset,
              simd_mask<N> mask = 1) {
  byte_offset += detail::localAccessorToOffset(lacc);
  return slm_atomic_update<Op, T, N>(byte_offset, mask);
}
 
 
 
 
 
template <atomic_op Op, typename T, int N>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1, simd<T, N>>
slm_atomic_update(simd<uint32_t, N> byte_offset, simd<T, N> src0,
                  simd_mask<N> mask = 1) {
  // Non-LSC atomic_update supports only 4-byte int vector operations with
  // 1,2,4,8,16,32 vector length. Non-LSC supports only 'store' for FP types.
  if constexpr (Op == atomic_op::fmin || Op == atomic_op::fmax ||
                Op == atomic_op::fadd || Op == atomic_op::fsub ||
                sizeof(T) != 4 || !__ESIMD_DNS::isPowerOf2(N, 32)) {
    return slm_atomic_update_impl<Op, T, N,
                                  detail::lsc_data_size::default_size>(
        byte_offset, src0, mask);
  } else if constexpr (Op == atomic_op::store) {
    if constexpr (std::is_integral_v<T>) {
      return slm_atomic_update<atomic_op::xchg, T, N>(byte_offset, src0, mask);
    } else {
      using Tint = detail::uint_type_t<sizeof(T)>;
      simd<Tint, N> Res = slm_atomic_update<atomic_op::xchg, Tint, N>(
          byte_offset, src0.template bit_cast_view<Tint>(), mask);
      return Res.template bit_cast_view<T>();
    }
  } else {
    detail::check_atomic<Op, T, N, 1>();
    const auto si = get_surface_index(detail::LocalAccessorMarker());
    return __esimd_dword_atomic1<Op, T, N>(mask.data(), si, byte_offset.data(),
                                           src0.data());
  }
}
 
template <atomic_op Op, typename T, int N, typename AccessorT>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 1 &&
                                 __ESIMD_DNS::is_rw_local_accessor_v<AccessorT>,
                             simd<T, N>>
atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0,
              simd_mask<N> mask = 1) {
  byte_offset += detail::localAccessorToOffset(lacc);
  return slm_atomic_update<Op, T, N>(byte_offset, src0, mask);
}
 
 
 
 
template <atomic_op Op, typename T, int N>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2, simd<T, N>>
slm_atomic_update(simd<uint32_t, N> byte_offset, simd<T, N> src0,
                  simd<T, N> src1, simd_mask<N> mask = 1) {
  // Non-LSC atomic_update supports only 4-byte int vector operations with
  // 1,2,4,8,16,32 vector length.
  if constexpr (sizeof(T) != 4 || Op == atomic_op::fcmpxchg ||
                !__ESIMD_DNS::isPowerOf2(N, 32)) {
    // 2-argument lsc_atomic_update arguments order matches the standard one -
    // expected value first, then new value. But atomic_update uses reverse
    // order, hence the src1/src0 swap.
    return detail::slm_atomic_update_impl<Op, T, N,
                                          detail::lsc_data_size::default_size>(
        byte_offset, src1, src0, mask);
  } else {
    detail::check_atomic<Op, T, N, 2>();
    const auto si = get_surface_index(detail::LocalAccessorMarker());
    return __esimd_dword_atomic2<Op, T, N>(mask.data(), si, byte_offset.data(),
                                           src0.data(), src1.data());
  }
}
 
template <atomic_op Op, typename T, int N, typename AccessorT>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 2 &&
                                 __ESIMD_DNS::is_rw_local_accessor_v<AccessorT>,
                             simd<T, N>>
atomic_update(AccessorT lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0,
              simd<T, N> src1, simd_mask<N> mask = 1) {
  byte_offset += detail::localAccessorToOffset(lacc);
  return slm_atomic_update<Op, T, N>(byte_offset, src0, src1, mask);
}
 
 
namespace detail {
 
template <atomic_op Op, typename T, int N, lsc_data_size DS,
          typename PropertyListT, typename Toffset>
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 0, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd_mask<N> pred) {
  static_assert(sizeof(T) > 1, "Unsupported data type");
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
  check_lsc_data_size<T, DS>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p);
  addrs += convert<uintptr_t>(offsets);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_stateless_0<MsgT, IOp, L1H, L2H, AddressScale,
                                      ImmOffset, EDS, VS, Transposed, N>(
          pred.data(), addrs.data());
  return lsc_format_ret<T>(Tmp);
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS,
          typename PropertyListT, typename Toffset>
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 1, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
                   simd_mask<N> pred) {
  static_assert(sizeof(T) > 1, "Unsupported data type");
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<MsgT, N> Msg_data = lsc_format_input<MsgT>(src0);
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p);
  addrs += convert<uintptr_t>(offsets);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_stateless_1<MsgT, IOp, L1H, L2H, AddressScale,
                                      ImmOffset, EDS, VS, Transposed, N>(
          pred.data(), addrs.data(), Msg_data.data());
  return lsc_format_ret<T>(Tmp);
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS,
          typename PropertyListT, typename Toffset>
__ESIMD_API std::enable_if_t<get_num_args<Op>() == 2, simd<T, N>>
atomic_update_impl(T *p, simd<Toffset, N> offsets, simd<T, N> src0,
                   simd<T, N> src1, simd_mask<N> pred) {
  static_assert(sizeof(T) > 1, "Unsupported data type");
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<MsgT, N> Msg_data0 = lsc_format_input<MsgT>(src0);
  simd<MsgT, N> Msg_data1 = lsc_format_input<MsgT>(src1);
  simd<uintptr_t, N> addrs = reinterpret_cast<uintptr_t>(p);
  addrs += convert<uintptr_t>(offsets);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_stateless_2<MsgT, IOp, L1H, L2H, AddressScale,
                                      ImmOffset, EDS, VS, Transposed, N>(
          pred.data(), addrs.data(), Msg_data0.data(), Msg_data1.data());
  return lsc_format_ret<T>(Tmp);
}
 
template <atomic_op Op, typename T, int N,
          lsc_data_size DS = lsc_data_size::default_size,
          typename PropertyListT, typename AccessorTy, typename Toffset>
__ESIMD_API
    std::enable_if_t<get_num_args<Op>() == 0 &&
                         __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
                     simd<T, N>>
    atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offsets,
                       simd_mask<N> pred) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update_impl<Op, T, N, DS, PropertyListT>(
      accessorToPointer<T>(acc), byte_offsets, pred);
#else
  static_assert(sizeof(T) > 1, "Unsupported data type");
  static_assert(std::is_integral_v<Toffset> && sizeof(Toffset) == 4,
                "Unsupported offset type");
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 0, /*IsLSC*/ true>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  auto si = get_surface_index(acc);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_bti_0<MsgT, IOp, L1H, L2H, AddressScale, ImmOffset,
                                EDS, VS, Transposed, N>(
          pred.data(), byte_offsets.data(), si);
  return lsc_format_ret<T>(Tmp);
#endif
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS,
          typename PropertyListT, typename AccessorTy, typename Toffset>
__ESIMD_API
    std::enable_if_t<get_num_args<Op>() == 1 &&
                         __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
                     simd<T, N>>
    atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offset,
                       simd<T, N> src0, simd_mask<N> pred) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update_impl<Op, T, N, DS, PropertyListT>(
      accessorToPointer<T>(acc), byte_offset, src0, pred);
#else
  static_assert(sizeof(T) > 1, "Unsupported data type");
  static_assert(std::is_integral_v<Toffset> && sizeof(Toffset) == 4,
                "Unsupported offset type");
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 1, /*IsLSC*/ true>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<MsgT, N> Src0Msg = lsc_format_input<MsgT>(src0);
  auto si = get_surface_index(acc);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_bti_1<MsgT, IOp, L1H, L2H, AddressScale, ImmOffset,
                                EDS, VS, Transposed, N>(
          pred.data(), byte_offset.data(), Src0Msg.data(), si);
  return lsc_format_ret<T>(Tmp);
#endif
}
 
template <atomic_op Op, typename T, int N, lsc_data_size DS,
          typename PropertyListT, typename AccessorTy, typename Toffset>
__ESIMD_API
    std::enable_if_t<get_num_args<Op>() == 2 &&
                         __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
                     simd<T, N>>
    atomic_update_impl(AccessorTy acc, simd<Toffset, N> byte_offset,
                       simd<T, N> src0, simd<T, N> src1, simd_mask<N> pred) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update_impl<Op, T, N, DS, PropertyListT>(
      __ESIMD_DNS::accessorToPointer<T>(acc), byte_offset, src0, src1, pred);
#else
  static_assert(std::is_integral_v<Toffset> && sizeof(Toffset) == 4,
                "Unsupported offset type");
  check_lsc_vector_size<1>();
  check_lsc_data_size<T, DS>();
  check_atomic<Op, T, N, 2, /*IsLSC*/ true>();
  check_cache_hints<cache_action::atomic, PropertyListT>();
  constexpr auto L1H = getCacheHintForIntrin<PropertyListT, cache_level::L1>();
  constexpr auto L2H = getCacheHintForIntrin<PropertyListT, cache_level::L2>();
  constexpr uint16_t AddressScale = 1;
  constexpr int ImmOffset = 0;
  constexpr lsc_data_size EDS = expand_data_size(finalize_data_size<T, DS>());
  constexpr lsc_vector_size VS = to_lsc_vector_size<1>();
  constexpr lsc_data_order Transposed = lsc_data_order::nontranspose;
  using MsgT = typename lsc_expand_type<T>::type;
  constexpr int IOp = lsc_to_internal_atomic_op<T, Op>();
  simd<MsgT, N> Msg_data0 = lsc_format_input<MsgT>(src0);
  simd<MsgT, N> Msg_data1 = lsc_format_input<MsgT>(src1);
  auto si = get_surface_index(acc);
  simd<MsgT, N> Tmp =
      __esimd_lsc_xatomic_bti_2<MsgT, IOp, L1H, L2H, AddressScale, ImmOffset,
                                EDS, VS, Transposed, N>(
          pred.data(), byte_offset.data(), Msg_data0.data(), Msg_data1.data(),
          si);
  return lsc_format_ret<T>(Tmp);
#endif
}
} // namespace detail
 
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, simd_mask<N> mask,
              PropertyListT props = {}) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
 
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                !__ESIMD_DNS::isPowerOf2(N, 32) || sizeof(T) < 4) {
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT, Toffset>(
        p, byte_offset, mask);
  } else if constexpr (N == 16 || N == 32) {
    // TODO: In fact GPU BE supports legalization for any N, even for
    // non-power-of-2, but it is implemented with an error now. For example,
    // N=17 is emulated as 2 calls (N=16 and N=1), while it must be 3 calls:
    // (N=8, N=8, N=1). I.e. Gen12 atomic instruction supports only N up to 8
    // and GPU thinks now it is up to 16.
    // Thus we emulate N=16 with 2 calls with N=8 each.
    // N=32 is emulated with 4 calls with N=8 each.
    // Task1: Remove the special-case emulation for N=16 and N=32 below when
    // GPU driver fixes the error.
    // Task2: remove the condition "!__ESIMD_DNS::isPowerOf2(N, 32)" above
    // and let svm.atomic for any N.
 
    simd<T, N> Res;
    for (int I = 0; I < N; I += 8) {
      simd_mask<8> Mask8 = mask.template select<8, 1>(I);
      simd<Toffset, 8> ByteOffset8 = byte_offset.template select<8, 1>(I);
      Res.template select<8, 1>(I) =
          atomic_update<Op, T, 8>(p, ByteOffset8, Mask8, props);
    }
    return Res;
  } else if constexpr (Op == atomic_op::load) {
    if constexpr (std::is_integral_v<T>) {
      return atomic_update<atomic_op::bit_or, T, N>(p, byte_offset,
                                                    simd<T, N>(0), mask, props);
    } else {
      using Tint = detail::uint_type_t<sizeof(T)>;
      simd<Tint, N> Res = atomic_update<atomic_op::bit_or, Tint, N>(
          reinterpret_cast<Tint *>(p), byte_offset, simd<Tint, N>(0), mask,
          props);
      return Res.template bit_cast_view<T>();
    }
  } else {
    detail::check_atomic<Op, T, N, 0>();
    simd<uintptr_t, N> vAddr(reinterpret_cast<uintptr_t>(p));
    simd<uintptr_t, N> offset_i1 = convert<uintptr_t>(byte_offset);
    vAddr += offset_i1;
    using Tx = typename detail::__raw_t<T>;
    return __esimd_svm_atomic0<Op, Tx, N>(vAddr.data(), mask.data());
  }
}
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, byte_offset, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT offsets, simd_mask<N> mask,
              PropertyListT props = {}) {
  return atomic_update<Op, T, N>(p, offsets.read(), mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT byte_offset, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, byte_offset.read(), mask, props);
}
 
template <atomic_op Op, typename T, int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, simd<T, N>>
atomic_update(T *p, Toffset byte_offset, simd_mask<N> mask = 1) {
  return atomic_update<Op, T, N>(p, simd<Toffset, N>(byte_offset), mask);
}
 
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd_mask<N> mask, PropertyListT props = {}) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
 
  // Auto-convert FP atomics to LSC version.
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                (Op == atomic_op::fmin) || (Op == atomic_op::fmax) ||
                (Op == atomic_op::fadd) || (Op == atomic_op::fsub) ||
                !__ESIMD_DNS::isPowerOf2(N, 32) || sizeof(T) < 4) {
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT, Toffset>(
        p, byte_offset, src0, mask);
  } else if constexpr (N == 16 || N == 32) {
    // TODO: In fact GPU BE supports legalization for any N, even for
    // non-power-of-2, but it is implemented with an error now. For example,
    // N=17 is emulated as 2 calls (N=16 and N=1), while it must be 3 calls:
    // (N=8, N=8, N=1). I.e. Gen12 atomic instruction supports only N up to 8
    // and GPU thinks now it is up to 16.
    // Thus we emulate N=16 with 2 calls with N=8 each.
    // N=32 is emulated with 4 calls with N=8 each.
    // Task1: Remove the special-case emulation for N=16 and N=32 below when
    // GPU driver fixes the error.
    // Task2: remove the condition "!__ESIMD_DNS::isPowerOf2(N, 32)" above
    // and let svm.atomic for any N.
    simd<T, N> Res;
    for (int I = 0; I < N; I += 8) {
      simd_mask<8> Mask8 = mask.template select<8, 1>(I);
      simd<Toffset, 8> ByteOffset8 = byte_offset.template select<8, 1>(I);
      simd<T, 8> Src08 = src0.template select<8, 1>(I);
      Res.template select<8, 1>(I) =
          atomic_update<Op, T, 8>(p, ByteOffset8, Src08, Mask8, props);
    }
    return Res;
  } else if constexpr (Op == atomic_op::store) {
    if constexpr (std::is_integral_v<T>) {
      return atomic_update<atomic_op::xchg, T, N>(p, byte_offset, src0, mask,
                                                  props);
    } else {
      using Tint = detail::uint_type_t<sizeof(T)>;
      simd<Tint, N> Res = atomic_update<atomic_op::xchg, Tint, N>(
          reinterpret_cast<Tint *>(p), byte_offset,
          src0.template bit_cast_view<Tint>(), mask, props);
      return Res.template bit_cast_view<T>();
    }
  } else {
    detail::check_atomic<Op, T, N, 1>();
    simd<uintptr_t, N> vAddr(reinterpret_cast<uintptr_t>(p));
    simd<uintptr_t, N> offset_i1 = convert<uintptr_t>(byte_offset);
    vAddr += offset_i1;
 
    using Tx = typename detail::__raw_t<T>;
    return __esimd_svm_atomic1<Op, Tx, N>(vAddr.data(), src0.data(),
                                          mask.data());
  }
}
 
 
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, byte_offset, src0, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT offsets, simd<T, N> src0, simd_mask<N> mask,
              PropertyListT props = {}) {
  return atomic_update<Op, T, N>(p, offsets.read(), src0, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT offsets, simd<T, N> src0,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, offsets.read(), src0, mask, props);
}
 
template <atomic_op Op, typename Tx, int N, typename Toffset>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<Toffset> &&
        ((Op != atomic_op::store && Op != atomic_op::xchg) || N == 1),
    simd<Tx, N>>
atomic_update(Tx *p, Toffset byte_offset, simd<Tx, N> src0, simd_mask<N> mask) {
  return atomic_update<Op, Tx, N>(p, simd<Toffset, N>(byte_offset), src0, mask);
}
 
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd<T, N> src1, simd_mask<N> mask, PropertyListT props = {}) {
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
 
  // Use LSC atomic when cache hints are present, FP atomics is used,
  // non-power of two length is used, or operation width greater than 32, or the
  // data size is less than 4 bytes.
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                Op == atomic_op::fcmpxchg || !__ESIMD_DNS::isPowerOf2(N, 32) ||
                sizeof(T) < 4) {
    // 2-argument lsc_atomic_update arguments order matches the standard one -
    // expected value first, then new value. But atomic_update uses reverse
    // order, hence the src1/src0 swap.
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT, Toffset>(
        p, byte_offset, src1, src0, mask);
  } else if constexpr (N == 16 || N == 32) {
    // TODO: In fact GPU BE supports legalization for any N, even for
    // non-power-of-2, but it is implemented with an error now. For example,
    // N=17 is emulated as 2 calls (N=16 and N=1), while it must be 3 calls:
    // (N=8, N=8, N=1). I.e. Gen12 atomic instruction supports only N up to 8
    // and GPU thinks now it is up to 16.
    // Thus we emulate N=16 with 2 calls with N=8 each.
    // N=32 is emulated with 4 calls with N=8 each.
    // Task1: Remove the special-case emulation for N=16 and N=32 below when
    // GPU driver fixes the error.
    // Task2: remove the condition "!__ESIMD_DNS::isPowerOf2(N, 32)" above
    // and let svm.atomic for any N.
    simd<T, N> Res;
    for (int I = 0; I < N; I += 8) {
      simd_mask<8> Mask8 = mask.template select<8, 1>(I);
      simd<Toffset, 8> ByteOffset8 = byte_offset.template select<8, 1>(I);
      simd<T, 8> Src08 = src0.template select<8, 1>(I);
      simd<T, 8> Src18 = src1.template select<8, 1>(I);
      Res.template select<8, 1>(I) =
          atomic_update<Op, T, 8>(p, ByteOffset8, Src08, Src18, Mask8, props);
    }
    return Res;
  } else {
    detail::check_atomic<Op, T, N, 2>();
    simd<uintptr_t, N> vAddr(reinterpret_cast<uintptr_t>(p));
    simd<uintptr_t, N> offset_i1 = convert<uintptr_t>(byte_offset);
    vAddr += offset_i1;
    using Tx = typename detail::__raw_t<T>;
    return __esimd_svm_atomic2<Op, Tx, N>(vAddr.data(), src0.data(),
                                          src1.data(), mask.data());
  }
}
 
//
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd<T, N> src1, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, byte_offset, src0, src1, mask, props);
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0,
              simd<T, N> src1, simd_mask<N> mask, PropertyListT props = {}) {
  return atomic_update<Op, T, N>(p, byte_offset.read(), src0, src1, mask,
                                 props);
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0,
              simd<T, N> src1, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(p, byte_offset.read(), src0, src1, mask,
                                 props);
}
 
template <atomic_op Op, typename Tx, int N, typename Toffset>
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, simd<Tx, N>>
atomic_update(Tx *p, Toffset byte_offset, simd<Tx, N> src0, simd<Tx, N> src1,
              simd_mask<N> mask) {
  return atomic_update<Op, Tx, N>(p, simd<Toffset, N>(byte_offset), src0, src1,
                                  mask);
}
 
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd_mask<N> mask,
              PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update<Op, T, N>(__ESIMD_DNS::accessorToPointer<T>(acc),
                                 byte_offset, mask, props);
#else
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
 
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                !detail::isPowerOf2(N, 32) || sizeof(T) < 4) {
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT>(
        acc, byte_offset, mask);
  } else {
    if constexpr (Op == atomic_op::load) {
      if constexpr (std::is_integral_v<T>) {
        return atomic_update<atomic_op::bit_or, T, N>(
            acc, byte_offset, simd<T, N>(0), mask, props);
      } else {
        using Tint = detail::uint_type_t<sizeof(T)>;
        simd<Tint, N> Res = atomic_update<atomic_op::bit_or, Tint, N>(
            acc, byte_offset, simd<Tint, N>(0), mask, props);
        return Res.template bit_cast_view<T>();
      }
    } else {
      detail::check_atomic<Op, T, N, 0>();
      static_assert(sizeof(Toffset) == 4, "Only 32 bit offset is supported");
 
      static_assert(sizeof(T) == 4, "Only 32 bit data is supported");
      const auto si = get_surface_index(acc);
      using Tx = typename detail::__raw_t<T>;
      return __esimd_dword_atomic0<Op, Tx, N>(mask.data(), si,
                                              byte_offset.data());
    }
  }
#endif
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset, mask, props);
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd_mask<N> mask,
              PropertyListT props = {}) {
  return atomic_update<Op, T, N>(acc, byte_offset.read(), mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 0 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset.read(), mask, props);
}
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
                         __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
                     simd<T, N>>
    atomic_update(AccessorTy acc, Toffset byte_offset, simd_mask<N> mask) {
  return atomic_update<Op, T, N>(acc, simd<Toffset, N>(byte_offset), mask);
}
 
template <atomic_op Op, typename T, int N, typename AccessorTy>
__ESIMD_API
    std::enable_if_t<__ESIMD_DNS::get_num_args<Op>() == 0 &&
                         __ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
                     simd<T, N>>
    atomic_update(AccessorTy acc, uint32_t byte_offset, simd_mask<N> mask) {
  return atomic_update<Op, T, N>(acc, simd<uint32_t, N>(byte_offset), mask);
}
 
 
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd_mask<N> mask, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update<Op, T, N>(__ESIMD_DNS::accessorToPointer<T>(acc),
                                 byte_offset, src0, mask, props);
#else
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  static_assert(sizeof(Toffset) == 4, "Only 32 bit offset is supported");
  // Auto-convert FP atomics to LSC version.
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                Op == atomic_op::fmin || Op == atomic_op::fmax ||
                Op == atomic_op::fadd || Op == atomic_op::fsub ||
                !__ESIMD_DNS::isPowerOf2(N, 32) || sizeof(T) < 4) {
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT>(
        acc, byte_offset, src0, mask);
  } else if constexpr (Op == atomic_op::store) {
    if constexpr (std::is_integral_v<T>) {
      return atomic_update<atomic_op::xchg, T, N>(acc, byte_offset, src0, mask,
                                                  props);
    } else {
      using Tint = detail::uint_type_t<sizeof(T)>;
      simd<Tint, N> Res = atomic_update<atomic_op::xchg, Tint, N>(
          acc, byte_offset, src0.template bit_cast_view<Tint>(), mask, props);
      return Res.template bit_cast_view<T>();
    }
  } else {
    detail::check_atomic<Op, T, N, 1>();
    static_assert(sizeof(T) == 4, "Only 32 bit data is supported");
    const auto si = __ESIMD_NS::get_surface_index(acc);
    using Tx = typename detail::__raw_t<T>;
    return __esimd_dword_atomic1<Op, Tx, N>(
        mask.data(), si, byte_offset.data(),
        sycl::bit_cast<__ESIMD_DNS::vector_type_t<Tx, N>>(src0.data()));
  }
#endif
}
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset, src0, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd<T, N> src0,
              simd_mask<N> mask, PropertyListT props = {}) {
  return atomic_update<Op, T, N>(acc, byte_offset.read(), src0, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 1 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd<T, N> src0,
              PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset.read(), src0, mask, props);
}
 
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ((Op != atomic_op::store && Op != atomic_op::xchg) || N == 1),
    simd<T, N>>
atomic_update(AccessorTy acc, Toffset offset, simd<T, N> src0,
              simd_mask<N> mask) {
  return atomic_update<Op, T, N>(acc, simd<Toffset, N>(offset), src0, mask);
}
 
template <atomic_op Op, typename Tx, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::is_rw_local_accessor_v<AccessorTy> &&
        ((Op != atomic_op::store && Op != atomic_op::xchg) || N == 1),
    simd<Tx, N>>
atomic_update(AccessorTy acc, uint32_t offset, simd<Tx, N> src0,
              simd_mask<N> mask) {
  return atomic_update<Op, Tx, N>(acc, simd<uint32_t, N>(offset), src0, mask);
}
 
//                simd_mask<N> mask,props = {});                 // (acc-au2-1)
 
//                simd_mask<N> mask,props = {});                 // (acc-au2-1)
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 && std::is_integral_v<Toffset> &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd<T, N> src1, simd_mask<N> mask, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  return atomic_update<Op, T, N>(__ESIMD_DNS::accessorToPointer<T>(acc),
                                 byte_offset, src0, src1, mask, props);
#else
  static_assert(std::is_integral_v<Toffset>, "Unsupported offset type");
  static_assert(sizeof(Toffset) == 4, "Only 32 bit offset is supported");
  // Use LSC atomic when cache hints are present, FP atomics is used,
  // non-power of two length is used, operation width greater than 32, or the
  // data size is less than 4 bytes,
  if constexpr (detail::has_cache_hints<PropertyListT>() ||
                Op == atomic_op::fcmpxchg || !__ESIMD_DNS::isPowerOf2(N, 32) ||
                sizeof(T) < 4) {
    // 2-argument lsc_atomic_update arguments order matches the standard one -
    // expected value first, then new value. But atomic_update uses reverse
    // order, hence the src1/src0 swap.
    return detail::atomic_update_impl<
        Op, T, N, detail::lsc_data_size::default_size, PropertyListT>(
        acc, byte_offset, src1, src0, mask);
  } else {
    detail::check_atomic<Op, T, N, 2>();
    static_assert(sizeof(T) == 4, "Only 32 bit data is supported");
    const auto si = __ESIMD_NS::get_surface_index(acc);
    using Tx = typename detail::__raw_t<T>;
    return __esimd_dword_atomic2<Op, Tx, N>(
        mask.data(), si, byte_offset.data(),
        sycl::bit_cast<__ESIMD_DNS::vector_type_t<Tx, N>>(src0.data()),
        sycl::bit_cast<__ESIMD_DNS::vector_type_t<Tx, N>>(src1.data()));
  }
#endif
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename Toffset,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT>,
    simd<T, N>>
atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0,
              simd<T, N> src1, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset, src0, src1, mask, props);
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd<T, N> src0,
              simd<T, N> src1, simd_mask<N> mask, PropertyListT props = {}) {
  return atomic_update<Op, T, N>(acc, byte_offset.read(), src0, src1, mask,
                                 props);
}
 
//    Other properties are ignored.
template <atomic_op Op, typename T, int N, typename OffsetSimdViewT,
          typename AccessorTy,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    __ESIMD_DNS::get_num_args<Op>() == 2 &&
        __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy> &&
        ext::oneapi::experimental::is_property_list_v<PropertyListT> &&
        detail::is_simd_view_type_v<OffsetSimdViewT>,
    simd<T, N>>
atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd<T, N> src0,
              simd<T, N> src1, PropertyListT props = {}) {
  simd_mask<N> mask = 1;
  return atomic_update<Op, T, N>(acc, byte_offset.read(), src0, src1, mask,
                                 props);
}
 
template <atomic_op Op, typename Tx, int N, typename Toffset,
          typename AccessorTy>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>,
                             simd<Tx, N>>
atomic_update(AccessorTy acc, Toffset offset, simd<Tx, N> src0,
              simd<Tx, N> src1, simd_mask<N> mask) {
  return atomic_update<Op, Tx, N>(acc, simd<Toffset, N>(offset), src0, src1,
                                  mask);
}
 
template <atomic_op Op, typename Tx, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>,
                             simd<Tx, N>>
atomic_update(AccessorTy acc, uint32_t offset, simd<Tx, N> src0,
              simd<Tx, N> src1, simd_mask<N> mask) {
  return atomic_update<Op, Tx, N>(acc, simd<uint32_t, N>(offset), src0, src1,
                                  mask);
}
 
 
 
enum fence_mask : uint8_t {
  global_coherent_fence = 0x1,
  l2_flush_instructions = 0x2,
  l3_flush_instructions __SYCL_DEPRECATED(
      "it means L2 here, use l2_flush_instructions") = l2_flush_instructions,
  l2_flush_texture_data = 0x4,
  l3_flush_texture_data __SYCL_DEPRECATED(
      "it means L2 here, use l2_flush_texture_data") = l2_flush_texture_data,
  l2_flush_constant_data = 0x8,
  l3_flush_constant_data __SYCL_DEPRECATED(
      "it means L2 here, use l2_flush_constant_data") = l2_flush_constant_data,
  l2_flush_rw_data = 0x10,
  l3_flush_rw_data __SYCL_DEPRECATED("it means L2 here, use l2_flush_rw_data") =
      l2_flush_rw_data,
  local_barrier = 0x20,
  l1_flush_ro_data = 0x40,
  sw_barrier __SYCL_DEPRECATED("reserved - this enum is ignored") = 0x80
};
 
template <uint8_t cntl> __ESIMD_API void fence() { __esimd_fence(cntl); }
 
__SYCL_DEPRECATED("use fence<fence_mask>()")
__ESIMD_API void fence(fence_mask cntl) { __esimd_fence(cntl); }
 
template <memory_kind Kind = memory_kind::global,
          fence_flush_op FenceOp = fence_flush_op::none,
          fence_scope Scope = fence_scope::group>
__ESIMD_API void fence() {
  static_assert(
      Kind != memory_kind::local ||
          (FenceOp == fence_flush_op::none && Scope == fence_scope::group),
      "SLM fence must have 'none' lsc_fence_op and 'group' scope");
  constexpr int N = 16;
  simd_mask<N> Mask = 1;
  __esimd_lsc_fence<static_cast<uint8_t>(Kind), static_cast<uint8_t>(FenceOp),
                    static_cast<uint8_t>(Scope), N>(Mask.data());
}
 
__ESIMD_API void barrier() {
  __esimd_fence(fence_mask::global_coherent_fence | fence_mask::local_barrier);
  __esimd_barrier();
}
 
 
template <typename T, int m, int N, typename AccessorTy, unsigned plane = 0>
__ESIMD_API simd<T, m * N> media_block_load(AccessorTy acc, unsigned x,
                                            unsigned y) {
  constexpr unsigned Width = N * sizeof(T);
  static_assert(Width * m <= 256u,
                "data does not fit into a single dataport transaction");
  static_assert(Width <= 64u, "valid block width is in range [1, 64]");
  static_assert(m <= 64u, "valid block height is in range [1, 64]");
  static_assert(plane <= 3u, "valid plane index is in range [0, 3]");
 
  const auto si = __ESIMD_NS::get_surface_index(acc);
  using SurfIndTy = decltype(si);
  constexpr unsigned int RoundedWidth =
      Width < 4 ? 4 : detail::getNextPowerOf2<Width>();
  constexpr int BlockWidth = sizeof(T) * N;
  constexpr int Mod = 0;
 
  if constexpr (Width < RoundedWidth) {
    constexpr unsigned int n1 = RoundedWidth / sizeof(T);
    simd<T, m * n1> temp =
        __esimd_media_ld<T, m, n1, Mod, SurfIndTy, (int)plane, BlockWidth>(
            si, x, y);
    return temp.template select<m, 1, N, 1>(0, 0);
  } else {
    return __esimd_media_ld<T, m, N, Mod, SurfIndTy, (int)plane, BlockWidth>(
        si, x, y);
  }
}
 
template <typename T, int m, int N, typename AccessorTy, unsigned plane = 0>
__ESIMD_API void media_block_store(AccessorTy acc, unsigned x, unsigned y,
                                   simd<T, m * N> vals) {
  constexpr unsigned Width = N * sizeof(T);
  static_assert(Width * m <= 256u,
                "data does not fit into a single dataport transaction");
  static_assert(Width <= 64u, "valid block width is in range [1, 64]");
  static_assert(m <= 64u, "valid block height is in range [1, 64]");
  static_assert(plane <= 3u, "valid plane index is in range [0, 3]");
  const auto si = __ESIMD_NS::get_surface_index(acc);
  using SurfIndTy = decltype(si);
  constexpr unsigned int RoundedWidth =
      Width < 4 ? 4 : detail::getNextPowerOf2<Width>();
  constexpr unsigned int n1 = RoundedWidth / sizeof(T);
  constexpr int BlockWidth = sizeof(T) * N;
  constexpr int Mod = 0;
 
  if constexpr (Width < RoundedWidth) {
    simd<T, m * n1> temp;
    auto temp_ref = temp.template bit_cast_view<T, m, n1>();
    auto vals_ref = vals.template bit_cast_view<T, m, N>();
    temp_ref.template select<m, 1, N, 1>() = vals_ref;
    __esimd_media_st<T, m, n1, Mod, SurfIndTy, plane, BlockWidth>(si, x, y,
                                                                  temp.data());
  } else {
    __esimd_media_st<T, m, N, Mod, SurfIndTy, plane, BlockWidth>(si, x, y,
                                                                 vals.data());
  }
}
 
template <typename T, int N, typename AccessorTy,
          typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API
    std::enable_if_t<detail::is_local_accessor_with_v<
                         AccessorTy, detail::accessor_mode_cap::can_read> &&
                         is_simd_flag_type_v<Flags>,
                     simd<T, N>>
    block_load(AccessorTy acc, uint32_t byte_offset, Flags flags) {
  return slm_block_load<T, N>(byte_offset + detail::localAccessorToOffset(acc),
                              flags);
}
 
template <typename T, int N, typename AccessorT, typename Flags>
__ESIMD_API
    std::enable_if_t<detail::is_local_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_write> &&
                     is_simd_flag_type_v<Flags>>
    block_store(AccessorT acc, uint32_t offset, simd<T, N> vals, Flags flags) {
  slm_block_store<T, N>(offset + __ESIMD_DNS::localAccessorToOffset(acc), vals,
                        flags);
}
 
//            typename PropertyListT = empty_properties_t>
 
template <typename T, int N, int VS, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets,
       simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}) {
  return slm_gather<T, N, VS>(byte_offsets +
                                  __ESIMD_DNS::localAccessorToOffset(acc),
                              mask, pass_thru, props);
}
 
template <typename T, int N, int VS, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets,
       simd_mask<N / VS> mask, PropertyListT props = {}) {
  return slm_gather<T, N, VS>(
      byte_offsets + __ESIMD_DNS::localAccessorToOffset(acc), mask, props);
}
 
template <typename T, int N, int VS, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets,
       PropertyListT props = {}) {
  return slm_gather<T, N, VS>(
      byte_offsets + __ESIMD_DNS::localAccessorToOffset(acc), props);
}
 
// Dev note: the mask type was turned into template parameter `MaskT` to
// avoid the conflicts of this prototype with the old gather() function
// accepting a 'global_offset' parameter and avoid 'ambiguous call' errors
// for calls like this: gather(acc, byte_offsets_simd, 0, mask);
template <typename T, int N, typename AccessorT, typename MaskT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     std::is_same_v<MaskT, simd_mask<N>> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N> byte_offsets, MaskT mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  return slm_gather<T, N>(byte_offsets +
                              __ESIMD_DNS::localAccessorToOffset(acc),
                          mask, pass_thru, props);
}
 
// Dev note: the mask type was turned into template parameter `MaskT` to
// avoid the conflicts of this prototype with the old gather() function
// accepting a 'global_offset' parameter and avoid 'ambiguous call' errors
// for calls like this: gather(acc, byte_offsets_simd, 0);
template <typename T, int N, typename AccessorT, typename MaskT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     std::is_same_v<MaskT, simd_mask<N>> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N> byte_offsets, MaskT mask,
       PropertyListT props = {}) {
  return slm_gather<T, N>(
      byte_offsets + __ESIMD_DNS::localAccessorToOffset(acc), mask, props);
}
 
template <typename T, int N, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, simd<uint32_t, N> byte_offsets,
       PropertyListT props = {}) {
  return slm_gather<T, N>(
      byte_offsets + __ESIMD_DNS::localAccessorToOffset(acc), props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       simd<T, N> pass_thru, PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), mask, pass_thru, props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
       PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), mask, props);
}
 
//            typename PropertyListT = empty_properties_t>
template <typename T, int N, int VS = 1, typename AccessorT,
          typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    (detail::is_local_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
     detail::is_simd_view_type_v<OffsetSimdViewT> &&
     ext::oneapi::experimental::is_property_list_v<PropertyListT>),
    simd<T, N>>
gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}) {
  return gather<T, N, VS>(acc, byte_offsets.read(), props);
}
 
template <typename T, int N, typename AccessorTy>
__ESIMD_API
    std::enable_if_t<detail::is_local_accessor_with_v<
                         AccessorTy, detail::accessor_mode_cap::can_read>,
                     simd<T, N>>
    gather(AccessorTy acc, simd<uint32_t, N> offsets, uint32_t glob_offset,
           simd_mask<N> mask = 1) {
  return slm_gather<T, N>(
      offsets + glob_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask);
}
 
 
 
 
 
 
template <typename T, int N, int VS = 1, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
  slm_scatter<T, N, VS>(byte_offsets + __ESIMD_DNS::localAccessorToOffset(acc),
                        vals, mask, props);
}
 
template <typename T, int N, int VS = 1, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(acc, byte_offsets, vals, Mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        simd_mask<N / VS> mask, PropertyListT props = {}) {
  scatter<T, N, VS>(acc, byte_offsets.read(), vals, mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT,
                                     detail::accessor_mode_cap::can_write> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals,
        PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  scatter<T, N, VS>(acc, byte_offsets.read(), vals, Mask, props);
}
 
template <typename T, int N, typename AccessorTy>
__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v<
    AccessorTy, detail::accessor_mode_cap::can_write>>
scatter(AccessorTy acc, simd<uint32_t, N> offsets, simd<T, N> vals,
        uint32_t glob_offset, simd_mask<N> mask = 1) {
  slm_scatter<T, N>(offsets + glob_offset +
                        __ESIMD_DNS::localAccessorToOffset(acc),
                    vals, mask);
}
 
 
template <typename T, int N, int VS, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask,
         PropertyListT props = {}) {
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
  detail::prefetch_impl<T, VS, detail::lsc_data_size::default_size,
                        PropertyListT>(p, byte_offsets, mask);
}
 
template <typename T, int N, int VS, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets,
         PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  prefetch<T, N, VS>(p, byte_offsets, Mask, props);
}
 
template <typename T, int N, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask,
         PropertyListT props = {}) {
  constexpr int VS = 1;
  prefetch<T, N, VS>(p, byte_offsets, mask, props);
}
 
template <typename T, int N, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}) {
  constexpr int VS = 1;
  prefetch<T, N, VS>(p, byte_offsets, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
         PropertyListT props = {}) {
  prefetch<T, N, VS>(p, byte_offsets.read(), mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}) {
  prefetch<T, N, VS>(p, byte_offsets.read(), props);
}
 
 
template <typename T, int VS = 1, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<OffsetT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask,
         PropertyListT props = {}) {
  detail::prefetch_impl<T, VS, detail::lsc_data_size::default_size,
                        PropertyListT>(p, byte_offset, mask);
}
 
template <typename T, int VS = 1, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<OffsetT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}) {
  simd_mask<1> Mask = 1;
  prefetch<T, VS>(p, byte_offset, Mask, props);
}
 
template <typename T, int VS = 1,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}) {
  prefetch<T, VS>(p, 0, mask, props);
}
 
template <typename T, int VS = 1,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(const T *p, PropertyListT props = {}) {
  simd_mask<1> Mask = 1;
  prefetch<T, VS>(p, 0, Mask, props);
}
 
 
template <typename T, int N, int VS, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
         simd_mask<N / VS> mask, PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  prefetch<T, N, VS>(detail::accessorToPointer<T>(acc), byte_offsets, mask,
                     props);
#else
  static_assert(N / VS >= 1 && N % VS == 0, "N must be divisible by VS");
  detail::prefetch_impl<T, VS, detail::lsc_data_size::default_size,
                        PropertyListT>(acc, byte_offsets, mask);
#endif // __ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int N, int VS, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets,
         PropertyListT props = {}) {
  simd_mask<N / VS> Mask = 1;
  prefetch<T, N, VS>(acc, byte_offsets, Mask, props);
}
 
template <typename T, int N, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask,
         PropertyListT props = {}) {
  constexpr int VS = 1;
  prefetch<T, N, VS>(acc, byte_offsets, mask, props);
}
 
template <typename T, int N, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, simd<OffsetT, N> byte_offsets,
         PropertyListT props = {}) {
  constexpr int VS = 1;
  prefetch<T, N, VS>(acc, byte_offsets, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask,
         PropertyListT props = {}) {
  prefetch<T, N, VS>(acc, byte_offsets.read(), mask, props);
}
 
template <typename T, int N, int VS = 1, typename OffsetSimdViewT,
          typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    detail::is_simd_view_type_v<OffsetSimdViewT> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, OffsetSimdViewT byte_offsets,
         PropertyListT props = {}) {
  prefetch<T, N, VS>(acc, byte_offsets.read(), props);
}
 
template <typename T, int VS = 1, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<OffsetT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, OffsetT byte_offset, simd_mask<1> mask,
         PropertyListT props = {}) {
#ifdef __ESIMD_FORCE_STATELESS_MEM
  prefetch<T, VS>(detail::accessorToPointer<T>(acc), byte_offset, mask, props);
#else
  detail::prefetch_impl<T, VS, detail::lsc_data_size::default_size,
                        PropertyListT>(acc, byte_offset, mask);
#endif // __ESIMD_FORCE_STATELESS_MEM
}
 
template <typename T, int VS = 1, typename AccessorT, typename OffsetT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    std::is_integral_v<OffsetT> &&
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, OffsetT byte_offset, PropertyListT props = {}) {
  simd_mask<1> Mask = 1;
  prefetch<T, VS>(acc, byte_offset, Mask, props);
}
 
template <typename T, int VS = 1, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, simd_mask<1> mask, PropertyListT props = {}) {
  prefetch<T, VS>(acc, 0, mask, props);
}
 
template <typename T, int VS = 1, typename AccessorT,
          typename PropertyListT =
              ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<
    detail::is_device_accessor_with_v<AccessorT,
                                      detail::accessor_mode_cap::can_read> &&
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch(AccessorT acc, PropertyListT props = {}) {
  simd_mask<1> Mask = 1;
  prefetch<T, VS>(acc, 0, Mask, props);
}
 
template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
          bool Transposed = false, bool Transformed = false,
          int N = detail::get_lsc_block_2d_data_size<
              T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(),
          typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>, simd<T, N>>
load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
        unsigned SurfacePitch, int X, int Y, PropertyListT props = {}) {
  return detail::load_2d_impl<T, BlockWidth, BlockHeight, NBlocks, Transposed,
                              Transformed, PropertyListT>(
      Ptr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y);
}
 
template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1,
          int N = detail::get_lsc_block_2d_data_size<
              T, NBlocks, BlockHeight, BlockWidth, false /*Transposed*/,
              false /*Transformed*/>(),
          typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
prefetch_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
            unsigned SurfacePitch, int X, int Y, PropertyListT props = {}) {
  detail::prefetch_2d_impl<T, BlockWidth, BlockHeight, NBlocks, PropertyListT>(
      Ptr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y);
}
 
template <typename T, int BlockWidth, int BlockHeight = 1,
          int N = detail::get_lsc_block_2d_data_size<
              T, 1u, BlockHeight, BlockWidth, false /*Transposed*/,
              false /*Transformed*/>(),
          typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t<
    ext::oneapi::experimental::is_property_list_v<PropertyListT>>
store_2d(T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight,
         unsigned SurfacePitch, int X, int Y, simd<T, N> Vals,
         PropertyListT props = {}) {
  detail::store_2d_impl<T, BlockWidth, BlockHeight, PropertyListT>(
      Ptr, SurfaceWidth, SurfaceHeight, SurfacePitch, X, Y, Vals);
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type>
__ESIMD_API
    std::enable_if_t<detail::is_local_accessor_with_v<
                         AccessorT, detail::accessor_mode_cap::can_read>,
                     simd<T, N * get_num_channels_enabled(RGBAMask)>>
    gather_rgba(AccessorT acc, simd<uint32_t, N> offsets,
                uint32_t global_offset = 0, simd_mask<N> mask = 1) {
  return slm_gather_rgba<T, N, RGBAMask>(
      offsets + global_offset + __ESIMD_DNS::localAccessorToOffset(acc), mask);
}
 
template <rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR,
          typename AccessorT, int N,
          typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v<
    AccessorT, detail::accessor_mode_cap::can_write>>
scatter_rgba(AccessorT acc, simd<uint32_t, N> offsets,
             simd<T, N * get_num_channels_enabled(RGBAMask)> vals,
             uint32_t global_offset = 0, simd_mask<N> mask = 1) {
  detail::validate_rgba_write_channel_mask<RGBAMask>();
  slm_scatter_rgba<T, N, RGBAMask>(offsets + global_offset +
                                       __ESIMD_DNS::localAccessorToOffset(acc),
                                   vals, mask);
}
 
 
template <uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_src1,
          uint8_t num_dst, raw_send_eot eot = raw_send_eot::not_eot,
          raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1, int n1,
          typename T2, int n2, typename T3, int n3>
__ESIMD_API __ESIMD_NS::simd<T1, n1>
raw_sends(__ESIMD_NS::simd<T1, n1> msg_dst, __ESIMD_NS::simd<T2, n2> msg_src0,
          __ESIMD_NS::simd<T3, n3> msg_src1, uint32_t ex_desc,
          uint32_t msg_desc, __ESIMD_NS::simd_mask<exec_size> mask = 1) {
  constexpr unsigned _Width1 = n1 * sizeof(T1);
  static_assert(_Width1 % 32 == 0, "Invalid size for raw send rspVar");
  constexpr unsigned _Width2 = n2 * sizeof(T2);
  static_assert(_Width2 % 32 == 0, "Invalid size for raw send msg_src0");
  constexpr unsigned _Width3 = n3 * sizeof(T3);
  static_assert(_Width3 % 32 == 0, "Invalid size for raw send msg_src1");
 
  using ElemT1 = __ESIMD_DNS::__raw_t<T1>;
  using ElemT2 = __ESIMD_DNS::__raw_t<T2>;
  using ElemT3 = __ESIMD_DNS::__raw_t<T3>;
 
  constexpr uint8_t modifier =
      ((eot == raw_send_eot::eot) << 1) | (sendc == raw_send_sendc::sendc);
 
  return __esimd_raw_sends2<ElemT1, n1, ElemT2, n2, ElemT3, n3, exec_size>(
      modifier, exec_size, mask.data(), num_src0, num_src1, num_dst, sfid,
      ex_desc, msg_desc, msg_src0.data(), msg_src1.data(), msg_dst.data());
}
 
template <uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_dst,
          raw_send_eot eot = raw_send_eot::not_eot,
          raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1, int n1,
          typename T2, int n2>
__ESIMD_API __ESIMD_NS::simd<T1, n1>
raw_send(__ESIMD_NS::simd<T1, n1> msg_dst, __ESIMD_NS::simd<T2, n2> msg_src0,
         uint32_t ex_desc, uint32_t msg_desc,
         __ESIMD_NS::simd_mask<exec_size> mask = 1) {
  constexpr unsigned _Width1 = n1 * sizeof(T1);
  static_assert(_Width1 % 32 == 0, "Invalid size for raw send rspVar");
  constexpr unsigned _Width2 = n2 * sizeof(T2);
  static_assert(_Width2 % 32 == 0, "Invalid size for raw send msg_src0");
 
  using ElemT1 = __ESIMD_DNS::__raw_t<T1>;
  using ElemT2 = __ESIMD_DNS::__raw_t<T2>;
 
  constexpr uint8_t modifier =
      ((eot == raw_send_eot::eot) << 1) | (sendc == raw_send_sendc::sendc);
  return __esimd_raw_send2<ElemT1, n1, ElemT2, n2, exec_size>(
      modifier, exec_size, mask.data(), num_src0, num_dst, sfid, ex_desc,
      msg_desc, msg_src0.data(), msg_dst.data());
}
 
template <uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_src1,
          raw_send_eot eot = raw_send_eot::not_eot,
          raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1, int n1,
          typename T2, int n2>
__ESIMD_API void raw_sends(__ESIMD_NS::simd<T1, n1> msg_src0,
                           __ESIMD_NS::simd<T2, n2> msg_src1, uint32_t ex_desc,
                           uint32_t msg_desc,
                           __ESIMD_NS::simd_mask<exec_size> mask = 1) {
  constexpr unsigned _Width1 = n1 * sizeof(T1);
  static_assert(_Width1 % 32 == 0, "Invalid size for raw send msg_src0");
  constexpr unsigned _Width2 = n2 * sizeof(T2);
  static_assert(_Width2 % 32 == 0, "Invalid size for raw send msg_src1");
 
  using ElemT1 = __ESIMD_DNS::__raw_t<T1>;
  using ElemT2 = __ESIMD_DNS::__raw_t<T2>;
 
  constexpr uint8_t modifier =
      ((eot == raw_send_eot::eot) << 1) | (sendc == raw_send_sendc::sendc);
  __esimd_raw_sends2_noresult<ElemT1, n1, ElemT2, n2, exec_size>(
      modifier, exec_size, mask.data(), num_src0, num_src1, sfid, ex_desc,
      msg_desc, msg_src0.data(), msg_src1.data());
}
 
template <uint8_t exec_size, uint8_t sfid, uint8_t num_src0,
          raw_send_eot eot = raw_send_eot::not_eot,
          raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1, int n1>
__ESIMD_API void raw_send(__ESIMD_NS::simd<T1, n1> msg_src0, uint32_t ex_desc,
                          uint32_t msg_desc,
                          __ESIMD_NS::simd_mask<exec_size> mask = 1) {
  constexpr unsigned _Width1 = n1 * sizeof(T1);
  static_assert(_Width1 % 32 == 0, "Invalid size for raw send msg_src0");
  using ElemT1 = __ESIMD_DNS::__raw_t<T1>;
  constexpr uint8_t modifier =
      ((eot == raw_send_eot::eot) << 1) | (sendc == raw_send_sendc::sendc);
  __esimd_raw_send2_noresult<ElemT1, n1, exec_size>(
      modifier, exec_size, mask.data(), num_src0, sfid, ex_desc, msg_desc,
      msg_src0.data());
}
 
 
 
 
namespace detail {
// -- Outlined implementations of simd_obj_impl class memory access APIs.
 
template <typename T, int N, class T1, class SFINAE>
template <typename Flags, int ChunkSize, typename>
void simd_obj_impl<T, N, T1, SFINAE>::copy_from(
    const simd_obj_impl<T, N, T1, SFINAE>::element_type *Addr,
    Flags) SYCL_ESIMD_FUNCTION {
  using UT = simd_obj_impl<T, N, T1, SFINAE>::element_type;
  constexpr unsigned Size = sizeof(T) * N;
  constexpr unsigned Align = Flags::template alignment<T1>;
 
  constexpr unsigned BlockSize = OperandSize::OWORD * 8;
  constexpr unsigned NumBlocks = Size / BlockSize;
  constexpr unsigned RemSize = Size % BlockSize;
 
  if constexpr (Align >= OperandSize::DWORD && Size % OperandSize::OWORD == 0 &&
                detail::isPowerOf2(RemSize / OperandSize::OWORD)) {
    if constexpr (NumBlocks > 0) {
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      ForHelper<NumBlocks>::unroll([BlockN, Addr, this](unsigned Block) {
        select<BlockN, 1>(Block * BlockN) =
            block_load<UT, BlockN, Flags>(Addr + (Block * BlockN), Flags{});
      });
    }
    if constexpr (RemSize > 0) {
      constexpr unsigned RemN = RemSize / sizeof(T);
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      select<RemN, 1>(NumBlocks * BlockN) =
          block_load<UT, RemN, Flags>(Addr + (NumBlocks * BlockN), Flags{});
    }
  } else if constexpr (sizeof(T) == 8) {
    simd<int32_t, N * 2> BC(reinterpret_cast<const int32_t *>(Addr), Flags{});
    bit_cast_view<int32_t>() = BC;
  } else {
    constexpr unsigned NumChunks = N / ChunkSize;
    if constexpr (NumChunks > 0) {
      simd<uint32_t, ChunkSize> Offsets(0u, sizeof(T));
      ForHelper<NumChunks>::unroll([Addr, &Offsets, this](unsigned Block) {
        select<ChunkSize, 1>(Block * ChunkSize) =
            gather<UT, ChunkSize>(Addr + (Block * ChunkSize), Offsets);
      });
    }
    constexpr unsigned RemN = N % ChunkSize;
    if constexpr (RemN > 0) {
      if constexpr (RemN == 1) {
        select<1, 1>(NumChunks * ChunkSize) = Addr[NumChunks * ChunkSize];
      } else if constexpr (RemN == 8 || RemN == 16) {
        simd<uint32_t, RemN> Offsets(0u, sizeof(T));
        select<RemN, 1>(NumChunks * ChunkSize) =
            gather<UT, RemN>(Addr + (NumChunks * ChunkSize), Offsets);
      } else {
        constexpr int N1 = RemN < 8 ? 8 : RemN < 16 ? 16 : 32;
        simd_mask_type<N1> Pred(0);
        Pred.template select<RemN, 1>() = 1;
        simd<uint32_t, N1> Offsets(0u, sizeof(T));
        simd<UT, N1> Vals =
            gather<UT, N1>(Addr + (NumChunks * ChunkSize), Offsets, Pred);
        select<RemN, 1>(NumChunks * ChunkSize) =
            Vals.template select<RemN, 1>();
      }
    }
  }
}
 
template <typename T, int N, class T1, class SFINAE>
template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
ESIMD_INLINE void simd_obj_impl<T, N, T1, SFINAE>::copy_to_impl(
    AccessorT acc, TOffset offset) const SYCL_ESIMD_FUNCTION {
  using UT = simd_obj_impl<T, N, T1, SFINAE>::element_type;
  constexpr unsigned Size = sizeof(T) * N;
  constexpr unsigned Align = Flags::template alignment<T1>;
 
  constexpr unsigned BlockSize = OperandSize::OWORD * 8;
  constexpr unsigned NumBlocks = Size / BlockSize;
  constexpr unsigned RemSize = Size % BlockSize;
 
  simd<UT, N> Tmp{data()};
  if constexpr (Align >= OperandSize::OWORD && Size % OperandSize::OWORD == 0 &&
                detail::isPowerOf2(RemSize / OperandSize::OWORD)) {
    if constexpr (NumBlocks > 0) {
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      ForHelper<NumBlocks>::unroll([BlockN, acc, offset, &Tmp](unsigned Block) {
        block_store<UT, BlockN, AccessorT>(
            acc, offset + (Block * BlockSize),
            Tmp.template select<BlockN, 1>(Block * BlockN));
      });
    }
    if constexpr (RemSize > 0) {
      constexpr unsigned RemN = RemSize / sizeof(T);
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      block_store<UT, RemN, AccessorT>(
          acc, offset + (NumBlocks * BlockSize),
          Tmp.template select<RemN, 1>(NumBlocks * BlockN));
    }
  } else if constexpr (sizeof(T) == 8) {
    simd<int32_t, N * 2> BC = Tmp.template bit_cast_view<int32_t>();
    BC.copy_to(acc, offset, Flags{});
  } else {
    constexpr unsigned NumChunks = N / ChunkSize;
    if constexpr (NumChunks > 0) {
      simd<TOffset, ChunkSize> Offsets(0u, sizeof(T));
      ForHelper<NumChunks>::unroll([acc, offset, &Offsets,
                                    &Tmp](unsigned Block) {
        scatter<UT, ChunkSize, AccessorT>(
            acc, Offsets, Tmp.template select<ChunkSize, 1>(Block * ChunkSize),
            offset + (Block * ChunkSize * sizeof(T)));
      });
    }
    constexpr unsigned RemN = N % ChunkSize;
    if constexpr (RemN > 0) {
      if constexpr (RemN == 1 || RemN == 8 || RemN == 16) {
        simd<TOffset, RemN> Offsets(0u, sizeof(T));
        scatter<UT, RemN, AccessorT>(
            acc, Offsets, Tmp.template select<RemN, 1>(NumChunks * ChunkSize),
            offset + (NumChunks * ChunkSize * sizeof(T)));
      } else {
        constexpr int N1 = RemN < 8 ? 8 : RemN < 16 ? 16 : 32;
        simd_mask_type<N1> Pred(0);
        Pred.template select<RemN, 1>() = 1;
        simd<UT, N1> Vals;
        Vals.template select<RemN, 1>() =
            Tmp.template select<RemN, 1>(NumChunks * ChunkSize);
        simd<TOffset, N1> Offsets(0u, sizeof(T));
        scatter<UT, N1, AccessorT>(acc, Offsets, Vals,
                                   offset + (NumChunks * ChunkSize * sizeof(T)),
                                   Pred);
      }
    }
  }
}
 
template <typename T, int N, class T1, class SFINAE>
template <int ChunkSize, typename Flags, typename AccessorT, typename TOffset>
ESIMD_INLINE void simd_obj_impl<T, N, T1, SFINAE>::copy_from_impl(
    AccessorT acc, TOffset offset) SYCL_ESIMD_FUNCTION {
  using UT = simd_obj_impl<T, N, T1, SFINAE>::element_type;
  static_assert(sizeof(UT) == sizeof(T));
  constexpr unsigned Size = sizeof(T) * N;
  constexpr unsigned Align = Flags::template alignment<T1>;
 
  constexpr unsigned BlockSize = OperandSize::OWORD * 8;
  constexpr unsigned NumBlocks = Size / BlockSize;
  constexpr unsigned RemSize = Size % BlockSize;
 
  if constexpr (Align >= OperandSize::DWORD && Size % OperandSize::OWORD == 0 &&
                detail::isPowerOf2(RemSize / OperandSize::OWORD)) {
    if constexpr (NumBlocks > 0) {
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      ForHelper<NumBlocks>::unroll([BlockN, acc, offset, this](unsigned Block) {
        select<BlockN, 1>(Block * BlockN) =
            block_load<UT, BlockN, AccessorT, Flags>(
                acc, offset + (Block * BlockSize), Flags{});
      });
    }
    if constexpr (RemSize > 0) {
      constexpr unsigned RemN = RemSize / sizeof(T);
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      select<RemN, 1>(NumBlocks * BlockN) =
          block_load<UT, RemN, AccessorT, Flags>(
              acc, offset + (NumBlocks * BlockSize), Flags{});
    }
  } else if constexpr (sizeof(T) == 8) {
    simd<int32_t, N * 2> BC(acc, offset, Flags{});
    bit_cast_view<int32_t>() = BC;
  } else {
    constexpr unsigned NumChunks = N / ChunkSize;
    if constexpr (NumChunks > 0) {
      simd<TOffset, ChunkSize> Offsets(0u, sizeof(T));
      ForHelper<NumChunks>::unroll(
          [acc, offset, &Offsets, this](unsigned Block) {
            select<ChunkSize, 1>(Block * ChunkSize) =
                gather<UT, ChunkSize, AccessorT>(
                    acc, Offsets, offset + (Block * ChunkSize * sizeof(T)));
          });
    }
    constexpr unsigned RemN = N % ChunkSize;
    if constexpr (RemN > 0) {
      if constexpr (RemN == 1 || RemN == 8 || RemN == 16) {
        simd<TOffset, RemN> Offsets(0u, sizeof(T));
        select<RemN, 1>(NumChunks * ChunkSize) = gather<UT, RemN, AccessorT>(
            acc, Offsets, offset + (NumChunks * ChunkSize * sizeof(T)));
      } else {
        constexpr int N1 = RemN < 8 ? 8 : RemN < 16 ? 16 : 32;
        simd_mask_type<N1> Pred(0);
        Pred.template select<RemN, 1>() = 1;
        simd<TOffset, N1> Offsets(0u, sizeof(T));
        simd<UT, N1> Vals = gather<UT, N1>(
            acc, Offsets, offset + (NumChunks * ChunkSize * sizeof(T)), Pred);
        select<RemN, 1>(NumChunks * ChunkSize) =
            Vals.template select<RemN, 1>();
      }
    }
  }
}
 
template <typename T, int N, class T1, class SFINAE>
template <typename AccessorT, typename Flags, int ChunkSize, typename>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_read, void>
simd_obj_impl<T, N, T1, SFINAE>::copy_from(AccessorT acc,
                                           detail::DeviceAccessorOffsetT offset,
                                           Flags) SYCL_ESIMD_FUNCTION {
 
  copy_from_impl<ChunkSize, Flags>(acc, offset);
}
 
template <typename T, int N, class T1, class SFINAE>
template <typename AccessorT, typename Flags, int ChunkSize, typename>
ESIMD_INLINE std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_read>,
    void>
simd_obj_impl<T, N, T1, SFINAE>::copy_from(AccessorT acc, uint32_t offset,
                                           Flags) SYCL_ESIMD_FUNCTION {
 
  copy_from_impl<ChunkSize, Flags>(acc, offset);
}
 
template <typename T, int N, class T1, class SFINAE>
template <typename Flags, int ChunkSize, typename>
void simd_obj_impl<T, N, T1, SFINAE>::copy_to(
    simd_obj_impl<T, N, T1, SFINAE>::element_type *Addr,
    Flags) const SYCL_ESIMD_FUNCTION {
  using UT = simd_obj_impl<T, N, T1, SFINAE>::element_type;
  constexpr unsigned Size = sizeof(T) * N;
  constexpr unsigned Align = Flags::template alignment<T1>;
 
  constexpr unsigned BlockSize = OperandSize::OWORD * 8;
  constexpr unsigned NumBlocks = Size / BlockSize;
  constexpr unsigned RemSize = Size % BlockSize;
 
  simd<UT, N> Tmp{data()};
  if constexpr (Align >= OperandSize::OWORD && Size % OperandSize::OWORD == 0 &&
                detail::isPowerOf2(RemSize / OperandSize::OWORD)) {
    if constexpr (NumBlocks > 0) {
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      ForHelper<NumBlocks>::unroll([BlockN, Addr, &Tmp](unsigned Block) {
        block_store<UT, BlockN>(Addr + (Block * BlockN),
                                Tmp.template select<BlockN, 1>(Block * BlockN));
      });
    }
    if constexpr (RemSize > 0) {
      constexpr unsigned RemN = RemSize / sizeof(T);
      constexpr unsigned BlockN = BlockSize / sizeof(T);
      block_store<UT, RemN>(Addr + (NumBlocks * BlockN),
                            Tmp.template select<RemN, 1>(NumBlocks * BlockN));
    }
  } else if constexpr (sizeof(T) == 8) {
    simd<int32_t, N * 2> BC = Tmp.template bit_cast_view<int32_t>();
    BC.copy_to(reinterpret_cast<int32_t *>(Addr), Flags{});
  } else {
    constexpr unsigned NumChunks = N / ChunkSize;
    if constexpr (NumChunks > 0) {
      simd<uint32_t, ChunkSize> Offsets(0u, sizeof(T));
      ForHelper<NumChunks>::unroll([Addr, &Offsets, &Tmp](unsigned Block) {
        scatter<UT, ChunkSize>(
            Addr + (Block * ChunkSize), Offsets,
            Tmp.template select<ChunkSize, 1>(Block * ChunkSize));
      });
    }
    constexpr unsigned RemN = N % ChunkSize;
    if constexpr (RemN > 0) {
      if constexpr (RemN == 1) {
        Addr[NumChunks * ChunkSize] = Tmp[NumChunks * ChunkSize];
      } else if constexpr (RemN == 8 || RemN == 16) {
        // TODO: GPU runtime may handle scatter of 16 byte elements
        // incorrectly. The code below is a workaround which must be deleted
        // once GPU runtime is fixed.
        if constexpr (sizeof(T) == 1 && RemN == 16) {
          if constexpr (Align % OperandSize::DWORD > 0) {
            ForHelper<RemN>::unroll([Addr, &Tmp](unsigned Index) {
              Addr[Index + NumChunks * ChunkSize] =
                  Tmp[Index + NumChunks * ChunkSize];
            });
          } else {
            simd_mask_type<8> Pred(0);
            simd<int32_t, 8> Vals;
            Pred.template select<4, 1>() = 1;
            Vals.template select<4, 1>() =
                Tmp.template bit_cast_view<int32_t>().template select<4, 1>(
                    NumChunks * ChunkSize);
 
            simd<uint32_t, 8> Offsets(0u, sizeof(int32_t));
            scatter<int32_t, 8>(
                reinterpret_cast<int32_t *>(Addr + (NumChunks * ChunkSize)),
                Offsets, Vals, Pred);
          }
        } else {
          simd<uint32_t, RemN> Offsets(0u, sizeof(T));
          scatter<UT, RemN>(
              Addr + (NumChunks * ChunkSize), Offsets,
              Tmp.template select<RemN, 1>(NumChunks * ChunkSize));
        }
      } else {
        constexpr int N1 = RemN < 8 ? 8 : RemN < 16 ? 16 : 32;
        simd_mask_type<N1> Pred(0);
        Pred.template select<RemN, 1>() = 1;
        simd<UT, N1> Vals;
        Vals.template select<RemN, 1>() =
            Tmp.template select<RemN, 1>(NumChunks * ChunkSize);
        simd<uint32_t, N1> Offsets(0u, sizeof(T));
        scatter<UT, N1>(Addr + (NumChunks * ChunkSize), Offsets, Vals, Pred);
      }
    }
  }
}
 
template <typename T, int N, class T1, class SFINAE>
template <typename AccessorT, typename Flags, int ChunkSize, typename>
ESIMD_INLINE EnableIfAccessor<AccessorT, accessor_mode_cap::can_write, void>
simd_obj_impl<T, N, T1, SFINAE>::copy_to(AccessorT acc,
                                         detail::DeviceAccessorOffsetT offset,
                                         Flags) const SYCL_ESIMD_FUNCTION {
  copy_to_impl<ChunkSize, Flags>(acc, offset);
}
 
template <typename T, int N, class T1, class SFINAE>
template <typename AccessorT, typename Flags, int ChunkSize, typename>
ESIMD_INLINE std::enable_if_t<
    detail::is_local_accessor_with_v<AccessorT, accessor_mode_cap::can_write>,
    void>
simd_obj_impl<T, N, T1, SFINAE>::copy_to(AccessorT acc, uint32_t offset,
                                         Flags) const SYCL_ESIMD_FUNCTION {
  copy_to_impl<ChunkSize, Flags>(acc, offset);
}
 
} // namespace detail
 
} // namespace ext::intel::esimd
} // namespace _V1
} // namespace sycl