DPC++ Runtime
Runtime libraries for oneAPI DPC++
memory.hpp File Reference
Include dependency graph for memory.hpp:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Namespaces

 sycl
 
 sycl::_V1
 
 sycl::_V1::ext
 
 sycl::_V1::ext::intel
 
 sycl::_V1::ext::intel::esimd
 
 sycl::_V1::ext::intel::esimd::detail
 

Macros

#define __ESIMD_FP_ATOMIC_OP_TYPE_CHECK(T)
 

Typedefs

using sycl::_V1::ext::intel::esimd::detail::DeviceAccessorOffsetT = uint32_t
 

Enumerations

enum  sycl::_V1::ext::intel::esimd::fence_mask : uint8_t {
  sycl::_V1::ext::intel::esimd::global_coherent_fence = 0x1 , sycl::_V1::ext::intel::esimd::l3_flush_instructions = 0x2 , sycl::_V1::ext::intel::esimd::l3_flush_texture_data = 0x4 , sycl::_V1::ext::intel::esimd::l3_flush_constant_data = 0x8 ,
  sycl::_V1::ext::intel::esimd::l3_flush_rw_data = 0x10 , sycl::_V1::ext::intel::esimd::local_barrier = 0x20 , sycl::_V1::ext::intel::esimd::l1_flush_ro_data = 0x40 , sycl::_V1::ext::intel::esimd::sw_barrier = 0x80
}
 Represetns a bit mask to control behavior of esimd::fence. More...
 

Functions

template<typename AccessorTy >
__ESIMD_API SurfaceIndex sycl::_V1::ext::intel::esimd::get_surface_index (AccessorTy acc)
 Get surface index corresponding to a SYCL accessor. More...
 
template<typename RT , typename T , int N>
ESIMD_INLINE simd< RT, N > sycl::_V1::ext::intel::esimd::detail::lsc_format_input (simd< T, N > Vals)
 
template<typename T , typename T1 , int N>
ESIMD_INLINE simd< T, N > sycl::_V1::ext::intel::esimd::detail::lsc_format_ret (simd< T1, N > Vals)
 
template<typename T , int NElts, lsc_data_size DS, cache_hint L1H, cache_hint L2H, int N, typename OffsetT >
__ESIMD_API simd< T, N *NElts > sycl::_V1::ext::intel::esimd::detail::gather_impl (const T *p, simd< OffsetT, N > offsets, simd_mask< N > pred)
 USM pointer gather. More...
 
template<typename T , int NElts, lsc_data_size DS, cache_hint L1H, cache_hint L2H, int N, typename OffsetT >
__ESIMD_API simd< T, N *NElts > sycl::_V1::ext::intel::esimd::detail::gather_impl (const T *p, simd< OffsetT, N > offsets, simd_mask< N > pred, simd< T, N *NElts > pass_thru)
 USM pointer gather. More...
 
template<typename T , int NElts, lsc_data_size DS, cache_hint L1H, cache_hint L2H, int N, typename Toffset >
__ESIMD_API void sycl::_V1::ext::intel::esimd::detail::scatter_impl (T *p, simd< Toffset, N > offsets, simd< T, N *NElts > vals, simd_mask< N > pred)
 USM pointer scatter. More...
 
constexpr bool sycl::_V1::ext::intel::esimd::detail::isMaskedGatherScatterLLVMAvailable ()
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> 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 = {}); // (usm-ga-1) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-4) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
 template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (usm-ga-5) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
 template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-ga-6) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
template<typename T , int N, int VS = 1, typename OffsetObjT , typename OffsetRegionT , 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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd_view< OffsetObjT, OffsetRegionT > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetObjT, typename OffsetRegionT, typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, simd_view<OffsetObjT, OffsetRegionT> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-7) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
template<typename T , int N, int VS = 1, typename OffsetObjT , typename OffsetRegionT , 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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd_view< OffsetObjT, OffsetRegionT > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 simd <T, N> gather(const T *p, simd_view<OffsetObjT, OffsetRegionT> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-8) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
template<typename T , int N, int VS = 1, typename OffsetObjT , typename OffsetRegionT , 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 > > sycl::_V1::ext::intel::esimd::gather (const T *p, simd_view< OffsetObjT, OffsetRegionT > byte_offsets, PropertyListT props={})
 simd <T, N> gather(const T *p, simd_view<OffsetObjT, OffsetRegionT> byte_offsets, PropertyListT props = {}); // (usm-ga-9) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
 
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< Tx, N > > sycl::_V1::ext::intel::esimd::gather (const Tx *p, Toffset offset, simd_mask< N > mask=1)
 A variation of gather API with offsets represented as scalar. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 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, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-1) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
 Writes ("scatters") elements of the input vector to different memory locations. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 Writes ("scatters") elements of the input vector to different memory locations. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT 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, PropertyListT props = {}); // (usm-sc-4) More...
 
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 > sycl::_V1::ext::intel::esimd::scatter (Tx *p, Toffset offset, simd< Tx, N > vals, simd_mask< N > mask=1)
 A variation of scatter API with offsets represented as scalar. More...
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename FlagsT >
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< FlagsT >, simd< T, NElts > > sycl::_V1::ext::intel::esimd::detail::block_load_impl (const T *p, simd_mask< 1 > pred, FlagsT flags)
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename FlagsT >
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< FlagsT >, simd< T, NElts > > sycl::_V1::ext::intel::esimd::detail::block_load_impl (const T *p, simd_mask< 1 > pred, simd< T, NElts > pass_thru, FlagsT flags)
 USM pointer transposed gather with 1 channel. More...
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename AccessorT , typename FlagsT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&is_simd_flag_type_v< FlagsT >, simd< T, NElts > > sycl::_V1::ext::intel::esimd::detail::block_load_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd_mask< 1 > pred, FlagsT flags)
 Accessor-based transposed gather with 1 channel. More...
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename AccessorT , typename FlagsT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&is_simd_flag_type_v< FlagsT >, simd< T, NElts > > sycl::_V1::ext::intel::esimd::detail::block_load_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd_mask< 1 > pred, simd< T, NElts > pass_thru, FlagsT flags)
 Accessor-based transposed gather with 1 channel. More...
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename FlagsT >
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< FlagsT > > sycl::_V1::ext::intel::esimd::detail::block_store_impl (T *p, simd< T, NElts > vals, simd_mask< 1 > pred, FlagsT flags)
 
template<typename T , int NElts, cache_hint L1H, cache_hint L2H, typename AccessorT , typename FlagsT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&is_simd_flag_type_v< FlagsT > > sycl::_V1::ext::intel::esimd::detail::block_store_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd< T, NElts > vals, simd_mask< 1 > pred, FlagsT flags)
 
template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags > > sycl::_V1::ext::intel::esimd::block_store (Tx *addr, simd< Tx, N > vals, Flags)
 Stores elements of the vector vals to a contiguous block of memory at the given address addr. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, PropertyListT props={})
 Each of the following block load functions loads a contiguous memory block from the address referenced by the USM pointer 'ptr', or from 'ptr + offset', where 'offset' is the offset in bytes (not in elements!). More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, size_t byte_offset, PropertyListT props={})
 simd<T, N> block_load(const T* ptr, size_t byte_offset, props={}); // (usm-bl-2) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> block_load(const T* ptr, simd_mask<1> pred, props={}); // (usm-bl-3) This function loads a contiguous memory block from USM pointer ptr. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, size_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> block_load(const T* ptr, size_t byte_offset, simd_mask<1> pred, props={}); // (usm-bl-4) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
 simd<T, N> block_load(const T* ptr, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (usm-bl-5) This function loads a contiguous memory block from USM pointer ptr. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const T *ptr, size_t byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
 simd<T, N> block_load(const T* ptr, size_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (usm-bl-6) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (const Tx *addr, Flags)
 Loads a contiguous block of memory from the given memory address addr and returns the loaded data as a vector. More...
 
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 > sycl::_V1::ext::intel::esimd::block_load (AccessorTy acc, detail::DeviceAccessorOffsetT byte_offset, Flags flags)
 Loads a contiguous block of memory from the given accessor acc and byte_offset and returns the loaded data as a vector. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, PropertyListT props={})
 Each of the following block load functions loads a contiguous memory block from the address referenced by accessor 'acc', or from 'acc + byte_offset', The parameter 'pred' is the one element predicate. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, PropertyListT={})
 simd<T, N> block_load(AccessorT acc, props = {}); // (acc-bl-2) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT={})
 simd<T, N> block_load(AccessorT acc, OffsetT byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props = {}); // (acc-bl-3) This function loads a contiguous memory block referenced by accessor acc and the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> block_load(AccessorT acc, OffsetT byte_offset, simd_mask<1> pred, props = {}); // (acc-bl-4) This function loads a contiguous memory block referenced by accessor acc and the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT={})
 simd<T, N> block_load(AccessorT acc, simd_mask<1> pred, simd<T, N> pass_thru, props = {}); // (acc-bl-5) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT acc, simd_mask< 1 > pred, PropertyListT={})
 simd<T, N> block_load(AccessorT acc, simd_mask<1> pred, props = {}); // (acc-bl-6) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (T *ptr, simd< T, N > vals, PropertyListT props={})
 Each of the following block store functions stores a contiguous memory block to the address referenced by the USM pointer 'ptr', or from 'ptr + offset', where 'offset' is the offset in bytes (not in elements!) with data specified by 'vals'. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (T *ptr, size_t byte_offset, simd< T, N > vals, PropertyListT props={})
 void block_store(T* ptr, size_t byte_offset, // (usm-bs-2) simd<T, N> vals, props={}); This function stores a contiguous memory block to USM pointer ptr and byte-offset byte_offset with data specified by vals. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (T *ptr, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(T* ptr, simd<T, N> vals, // (usm-bs-3) simd_mask<1> pred, props={}); This function stores a contiguous memory block to USM pointer ptr with data specified by vals. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (T *ptr, size_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(T* ptr, size_t byte_offset, // (usm-bs-4) simd<T, N> vals, simd_mask<1> pred, props={}); This function stores a contiguous memory block to USM pointer ptr and byte-offset byte_offset with data specified by vals. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd< T, N > vals, PropertyListT props={})
 Each of the following block_store functions stores the vector 'vals' to a contiguous memory block at the address referenced by accessor 'acc', or from 'acc + byte_offset', The parameter 'pred' is the one element predicate. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, simd< T, N > vals, PropertyListT props={})
 void block_store(AccessorT acc, simd<T, N> vals, props = {}); // (acc-bs-2) This function stores a contiguous memory block to accessor acc with data specified by vals and implied offset=0. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(AccessorT acc, OffsetT byte_offset, // (acc-bs-3) simd<T, N> vals, simd_mask<1> pred, props = {}); This function stores a contiguous memory block to accessor acc and byte_offset with data specified by vals. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(AccessorT acc, simd<T, N> vals, // (acc-bs-4) simd_mask<1> pred, props = {}); This function stores a contiguous memory block to accessor acc with data specified by vals and implied offset=0. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > byte_offsets, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
 Accessor-based gather. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, detail::DeviceAccessorOffsetT glob_offset)
 Loads and broadcasts the element located at acc and byte offset glob_offset to a vector and returns it as a simd object. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-1) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) Supported platforms: DG2, PVC in most cases. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) Supported platforms: DG2, PVC in most cases. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-4) This function is identical to (acc-ga-1) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (acc-ga-5) This function is identical to (acc-ga-2) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< OffsetT, N > byte_offsets, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-ga-6) This function is identical to (acc-ga-3) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 Accessor-based scatter. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-2) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-4) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< detail::DeviceAccessorOffsetT, N > offsets, simd< T, N > vals, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
 Writes elements of a simd object into an accessor at given offsets. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, detail::DeviceAccessorOffsetT glob_offset, simd< T, N > vals, simd_mask< N > mask=1)
 
template<typename T , typename AccessorTy >
__ESIMD_API T sycl::_V1::ext::intel::esimd::scalar_load (AccessorTy acc, detail::DeviceAccessorOffsetT offset)
 Load a scalar value from an accessor. More...
 
template<typename T , typename AccessorTy >
__ESIMD_API void sycl::_V1::ext::intel::esimd::scalar_store (AccessorTy acc, detail::DeviceAccessorOffsetT offset, T val)
 Store a scalar value into an accessor. More...
 
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)> sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, simd< Toffset, N > offsets, simd_mask< N > mask=1)
 Gather and transpose pixels from given memory locations defined by the base pointer p and offsets. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API simd< T, N *get_num_channels_enabled(RGBAMask)> sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, simd_view< OffsetObjT, RegionTy > offsets, simd_mask< N > mask=1)
 A variation of gather_rgba API with offsets represented as simd_view object. More...
 
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)> > sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, Toffset offset, simd_mask< N > mask=1)
 A variation of gather_rgba API with offsets represented as scalar. More...
 
template<rgba_channel_mask M>
static void sycl::_V1::ext::intel::esimd::detail::validate_rgba_write_channel_mask ()
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API void sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd< Toffset, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
 Transpose and scatter pixels to given memory locations defined by the base pointer p and offsets. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API void sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
 A variation of scatter_rgba API with offsets represented as simd_view object. More...
 
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 > sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, Toffset offset, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
 A variation of scatter_rgba API with offsets represented as scalar. More...
 
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 > sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
 
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)> > sycl::_V1::ext::intel::esimd::gather_rgba (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > offsets, detail::DeviceAccessorOffsetT global_offset=0, simd_mask< N > mask=1)
 Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. More...
 
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 > > sycl::_V1::ext::intel::esimd::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)
 Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
 
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, unsigned NumSrc, bool IsLSC = false>
constexpr void sycl::_V1::ext::intel::esimd::detail::check_atomic ()
 Check the legality of an atomic call in terms of size and type. More...
 
template<uint32_t SLMSize>
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_init ()
 Declare per-work-group slm size. More...
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_init (uint32_t size)
 Declare per-work-group slm size. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-1) simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-2) simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (slm-ga-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-2) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N/VS > byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (slm-ga-3) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N > byte_offsets, simd_mask< N > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-4) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
 template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (slm-ga-5) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N > byte_offsets, PropertyListT props={})
 template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (slm-ga-6) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-7) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-8) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_gather (OffsetSimdViewT byte_offsets, PropertyListT props={})
 simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (slm-ga-9) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
 
template<typename T >
__ESIMD_API T sycl::_V1::ext::intel::esimd::slm_scalar_load (uint32_t offset)
 Load a scalar value from the Shared Local Memory. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_scatter (simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename PropertyListT = empty_properties_t> void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-sc-1) void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-2) More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_scatter (simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename PropertyListT = empty_properties_t> void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-2) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_scatter (OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> void slm_scatter( OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-sc-3) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_scatter (OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
 void slm_scatter( OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-4) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
 
template<typename T >
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_scalar_store (uint32_t offset, T val)
 Store a scalar value into the Shared Local Memory. More...
 
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)> > sycl::_V1::ext::intel::esimd::slm_gather_rgba (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
template<typename T , int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4)> sycl::_V1::ext::intel::esimd::slm_scatter_rgba (simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(Mask)> vals, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_load (uint32_t byte_offset, Flags)
 Loads a contiguous block of SLM memory referenced by the given byte-offset offset, then returns the loaded data as a simd object. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_load (uint32_t byte_offset, PropertyListT props={})
 Each of the following slm_block_load functions loads a contiguous memory block from SLM (Shared Local Memory) and the byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_load (uint32_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> slm_block_load(uint32_t byte_offset, simd_mask<N> pred, props = {}); // (slm-bl-2) Loads a contiguous memory block from SLM (Shared Local Memory) at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_load (uint32_t offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
 simd<T, N> slm_block_load(uint32_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (slm-bl-3) Loads a contiguous memory block from SLM (Shared Local Memory) at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, uint32_t byte_offset, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, props={}); // (lacc-bl-1) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, props={}); // (lacc-bl-2) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, uint32_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, simd_mask<1> pred, props={}); // (lacc-bl-3) Loads a contiguous memory block from SLM (Shared Local Memory) associated the local accessor lacc at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, simd_mask< 1 > pred, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, simd_mask<1> pred, props={}); // (lacc-bl-4) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, uint32_t byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (lacc-bl-5) Loads a contiguous memory block from SLM (Shared Local Memory) associated the local accessor lacc at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorT lacc, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
 simd<T, N> block_load(local_accessor lacc, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (lacc-bl-6) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
 
template<typename T , int N, typename Flags >
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags > > sycl::_V1::ext::intel::esimd::slm_block_store (uint32_t offset, simd< T, N > vals, Flags)
 Stores elements of the vector vals to a contiguous block of SLM memory at the given byte-offset offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_store (uint32_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 Each of the following slm_block_store functions stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) at the byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::slm_block_store (uint32_t byte_offset, simd< T, N > vals, PropertyListT props={})
 void slm_block_store(uint32_t byte_offset, simd<T, N> vals, // (slm-bs-2) props = {}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT lacc, uint32_t byte_offset, simd< T, N > vals, PropertyListT props={})
 void block_store(local_accessor lacc, uint32_t byte_offset, // (lacc-bs-1) simd<T, N> vals, props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc at the given byte_offset. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT lacc, simd< T, N > vals, PropertyListT props={})
 void block_store(local_accessor lacc, simd<T, N> vals, // (lacc-bs-2) props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT lacc, uint32_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(local_accessor lacc, uint32_t byte_offset, // (lacc-bs-3) simd<T, N> vals, simd_mask<1> pred, props={}); More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT lacc, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
 void block_store(local_accessor lacc, simd<T, N> vals, // (lacc-bs-4) simd_mask<1> pred, props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc. More...
 
template<typename T , sycl::ext::intel::esimd::atomic_op Op>
constexpr int sycl::_V1::ext::intel::esimd::detail::lsc_to_internal_atomic_op ()
 
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 > > sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl (simd< uint32_t, N > offsets, simd_mask< N > pred)
 SLM atomic. More...
 
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 > > sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl (simd< uint32_t, N > offsets, simd< T, N > src0, simd_mask< N > pred)
 SLM atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API simd< T, N > sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl (simd< uint32_t, N > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 SLM atomic. More...
 
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > byte_offset, simd_mask< N > mask=1)
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd_mask< N > mask=1)
 simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd_mask<N> pred = 1); // (lacc-au0-1) Atomically updates N memory locations in SLM ssociated with the local accessor lacc at the given byte_offset, and returns a vector of old values found at the memory locations before update. More...
 
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > byte_offset, simd< T, N > src0, simd_mask< N > mask=1)
 One argument variant of the atomic update operation. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd< T, N > src0, simd_mask< N > mask=1)
 simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0, simd_mask<1> pred = 1); // (lacc-au1-1) More...
 
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
 Two argument variant of the atomic update operation. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
 simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<1> pred = 1); // (lacc-au2-1) More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, cache_hint L1H, cache_hint L2H, typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==0, simd< T, N > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (T *p, simd< Toffset, N > offsets, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, cache_hint L1H, cache_hint L2H, typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==1, simd< T, N > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (T *p, simd< Toffset, N > offsets, simd< T, N > src0, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, cache_hint L1H, cache_hint L2H, typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==2, simd< T, N > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (T *p, simd< Toffset, N > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, 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 > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offsets, simd_mask< N > pred)
 Accessor-based atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, cache_hint L1H, cache_hint L2H, 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 > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > pred)
 Accessor-based atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, cache_hint L1H, cache_hint L2H, 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 > > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 Accessor-based atomic. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, simd_mask< N > mask, PropertyListT props={})
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd<Toffset, N> byte_offset, props = {}); /// (usm-au0-2) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, RegionTy> byte_offset, simd_mask<N> mask, props = {}); /// (usm-au0-3) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, RegionTy > byte_offset, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, RegionTy> byte_offset, props = {}); /// (usm-au0-4) More...
 
template<atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< T, N > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, Toffset byte_offset, simd_mask< N > mask=1)
 A variation of atomic_update API with offset represented as scalar. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, PropertyListT props={})
 simd<T, N> atomic_update(T *ptr, simd<Toffset, N> byte_offset, simd<T, N> src0, props = {}); // (usm-au1-2) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, simd_mask<N> mask, props = {}); // (usm-au1-3) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, props = {}); // (usm-au1-4) More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (Tx *p, Toffset byte_offset, simd< Tx, N > src0, simd_mask< N > mask)
 A variation of atomic_update API with offset represented as scalar object. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
 Atomically updates N memory locations represented by a USM pointer and a vector of offsets relative to the pointer, and returns a vector of old values found at the memory locations before update. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (usm-au2-2) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename OffsetRegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, OffsetRegionTy > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<N> mask, props = {}) // (usm-au2-3) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename OffsetRegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< OffsetObjT, OffsetRegionTy > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
 simd<T, N> atomic_update(T *p, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}) // (usm-au2-4) More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< Tx, N > > sycl::_V1::ext::intel::esimd::atomic_update (Tx *p, Toffset byte_offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 A variation of atomic_update API with byte_offset represented as scalar. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd_mask< N > mask, PropertyListT props={})
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd<OffsetObjT, N> byte_offset, props = {}); /// (acc-au0-2) A variation of atomic_update API without mask operand More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > byte_offset, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd_view<OffsetObjT, RegionTy> byte_offset, simd_mask<N> mask, props = {}); /// (acc-au0-3) A variation of atomic_update API with offsets represented as simd_view object. More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > byte_offset, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd_view<OffsetObjT, RegionTy> byte_offset, props = {}); /// (acc-au0-4) A variation of atomic_update API with offsets represented as simd_view object and no mask operand. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, Toffset byte_offset, simd_mask< N > mask)
 A variation of atomic_update API with offset represented as scalar. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, uint32_t byte_offset, simd_mask< N > mask)
 A variation of atomic_update API with byte_offset represented as scalar using local_accessor. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd<Toffset, N> byte_offset, simd<T, N> src0, props = {}); // (acc-au1-2) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, simd_mask<N> mask, props = {}); // (acc-au1-3) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename RegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > byte_offset, simd< T, N > src0, PropertyListT props={})
 simd<T, N> atomic_update(AccessorT acc, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, props = {}); // (acc-au1-4) More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, Toffset offset, simd< T, N > src0, simd_mask< N > mask)
 A variation of atomic_update API with offset represented as scalar object. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, uint32_t offset, simd< Tx, N > src0, simd_mask< N > mask)
 A variation of atomic_update API with offset represented as scalar object and uses local_accessor. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
  More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
 simd<T, N> atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (acc-au2-2) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename OffsetRegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, OffsetRegionTy > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
 simd<T, N> atomic_update(AccessorTy acc, simd_view<OffsetObjT, OffsetRegionTy> byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<N> mask, props = {}); // (acc-au2-3) More...
 
template<atomic_op Op, typename T , int N, typename OffsetObjT , typename AccessorTy , typename OffsetRegionTy , 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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, simd_view< OffsetObjT, OffsetRegionTy > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
 simd<T, N> atomic_update(AccessorTy acc, simd_view<OffsetObjT, OffsetRegionTy>, byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (acc-au2-4) More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, Toffset offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as scalar. More...
 
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 > > sycl::_V1::ext::intel::esimd::atomic_update (AccessorTy acc, uint32_t offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as scalar and local_accessor is used. More...
 
template<uint8_t cntl>
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence ()
 esimd::fence sets the memory read/write order. More...
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence (fence_mask 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 sycl::_V1::ext::intel::esimd::fence ()
 Memory fence. More...
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::barrier ()
 Generic work-group barrier. More...
 
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API simd< T, m *N > sycl::_V1::ext::intel::esimd::media_block_load (AccessorTy acc, unsigned x, unsigned y)
 Media block load. More...
 
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API void sycl::_V1::ext::intel::esimd::media_block_store (AccessorTy acc, unsigned x, unsigned y, simd< T, m *N > vals)
 Media block store. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_load (AccessorTy acc, uint32_t byte_offset, Flags flags)
 Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte_offset, then returns the loaded data as a simd object. More...
 
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 > > sycl::_V1::ext::intel::esimd::block_store (AccessorT acc, uint32_t offset, simd< T, N > vals, Flags flags)
 Variant of block_store that uses local accessor as a parameter. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 Variant of gather that uses local accessor as a parameter template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> 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 = {}); // (lacc-ga-1) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) Supported platforms: DG2, PVC in most cases. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) Supported platforms: DG2, PVC in most cases. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-4) This function is identical to (lacc-ga-1) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (lacc-ga-5) This function is identical to (lacc-ga-2) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, simd< uint32_t, N > byte_offsets, PropertyListT props={})
 template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (lacc-ga-6) This function is identical to (lacc-ga-3) except that vector size is fixed to 1. More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
 
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 > > sycl::_V1::ext::intel::esimd::gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset, simd_mask< N > mask=1)
 Variant of gather that uses local accessor as a parameter. More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 Variant of scatter that uses local accessor as a parameter template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-1) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-2) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-3) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-4) More...
 
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 > > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset, simd_mask< N > mask=1)
 Variant of scatter that uses local accessor as a parameter. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-1) void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
 template <typename T, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
 template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (usm-pf-3) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
 template <typename T, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-pf-4) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-5) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetSimdViewT byte_offsets, PropertyListT props={})
 template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-pf-6) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetT byte_offset, simd_mask< 1 > mask, PropertyListT props={})
 template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-7) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, OffsetT byte_offset, PropertyListT props={})
 template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}); // (usm-pf-8) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, simd_mask< 1 > mask, PropertyListT props={})
 template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}); //(usm-pf-9) Supported platforms: DG2, PVC only. More...
 
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 > > sycl::_V1::ext::intel::esimd::prefetch (const T *p, PropertyListT props={})
 template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, PropertyListT props = {}); // (usm-pf-10) Supported platforms: DG2, PVC only. More...
 
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)> > sycl::_V1::ext::intel::esimd::gather_rgba (AccessorT acc, simd< uint32_t, N > offsets, uint32_t global_offset=0, simd_mask< N > mask=1)
 Variant of gather_rgba that uses local accessor as a parameter. More...
 
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 > > sycl::_V1::ext::intel::esimd::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)
 Variant of scatter_rgba that uses local accessor as a parameter Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
 
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 sycl::ext::intel::esimd::simd< T1, n1 > sycl::_V1::ext::intel::esimd::raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msg_dst, sycl::ext::intel::esimd::simd< T2, n2 > msg_src0, sycl::ext::intel::esimd::simd< T3, n3 > msg_src1, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
 Raw sends. More...
 
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 sycl::ext::intel::esimd::simd< T1, n1 > sycl::_V1::ext::intel::esimd::raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msg_dst, sycl::ext::intel::esimd::simd< T2, n2 > msg_src0, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
 Raw send. More...
 
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 sycl::_V1::ext::intel::esimd::raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msg_src0, sycl::ext::intel::esimd::simd< T2, n2 > msg_src1, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
 Raw sends. More...
 
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 sycl::_V1::ext::intel::esimd::raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msg_src0, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
 Raw send. More...
 

Macro Definition Documentation

◆ __ESIMD_FP_ATOMIC_OP_TYPE_CHECK

#define __ESIMD_FP_ATOMIC_OP_TYPE_CHECK (   T)
Value:
static_assert(is_type<T, float, sycl::half, double>(), \
"float, double or sycl::half type is expected");

Definition at line 4048 of file memory.hpp.