This group combines types and functions specific to LSC, which is available in Intel GPUs starting from PVC and ACM. More...

Collaboration diagram for LSC-specific memory access APIs.:

Classes
class	sycl::_V1::ext::intel::experimental::esimd::config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks >
	Container class to hold parameters for `load2d/store2d` `functions` More...

Enumerations
enum class	sycl::_V1::ext::intel::esimd::native::lsc::atomic_op : uint8_t { sycl::_V1::ext::intel::esimd::native::lsc::inc = 0x08 , sycl::_V1::ext::intel::esimd::native::lsc::dec = 0x09 , sycl::_V1::ext::intel::esimd::native::lsc::load = 0x0a , sycl::_V1::ext::intel::esimd::native::lsc::store = 0x0b , sycl::_V1::ext::intel::esimd::native::lsc::add = 0x0c , sycl::_V1::ext::intel::esimd::native::lsc::sub = 0x0d , sycl::_V1::ext::intel::esimd::native::lsc::smin = 0x0e , sycl::_V1::ext::intel::esimd::native::lsc::smax = 0x0f , sycl::_V1::ext::intel::esimd::native::lsc::umin = 0x10 , sycl::_V1::ext::intel::esimd::native::lsc::umax = 0x11 , sycl::_V1::ext::intel::esimd::native::lsc::cmpxchg = 0x12 , sycl::_V1::ext::intel::esimd::native::lsc::fadd = 0x13 , sycl::_V1::ext::intel::esimd::native::lsc::fsub = 0x14 , sycl::_V1::ext::intel::esimd::native::lsc::fmin = 0x15 , sycl::_V1::ext::intel::esimd::native::lsc::fmax = 0x16 , sycl::_V1::ext::intel::esimd::native::lsc::fcmpxchg = 0x17 , sycl::_V1::ext::intel::esimd::native::lsc::bit_and = 0x18 , sycl::_V1::ext::intel::esimd::native::lsc::bit_or = 0x19 , sycl::_V1::ext::intel::esimd::native::lsc::bit_xor = 0x1a }
	LSC atomic operation codes. More...

Functions
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	SLM gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
	SLM gather. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})
	Transposed SLM gather with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru)
	Transposed SLM gather with 1 channel. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N NElts > pass_thru)
	USM pointer gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N NElts > pass_thru)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N NElts > pass_thru)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
	Accessor-based gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N *NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	USM pointer transposed gather with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, FlagsT)
	A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT={})
	USM pointer transposed gather with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})
	Accessor-based transposed gather with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, FlagsT flags)
	A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, FlagsT flags)

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT={})
	Accessor-based transposed gather with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT flags=FlagsT{})

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer prefetch gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > >	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, FlagsT={})
	USM pointer prefetch transposed gather with 1 channel. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > >	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based prefetch gather. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, FlagsT flags=FlagsT{})
	Accessor-based transposed prefetch gather with 1 channel. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_scatter (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	SLM scatter. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_store (uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags=FlagsT{})
	Transposed SLM scatter with 1 channel. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer scatter. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T p, Toffset offset, sycl::ext::intel::esimd::simd< T, N NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > >	sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based scatter. More...

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > >	sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	USM pointer transposed scatter with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags)
	A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	Accessor-based transposed scatter with 1 channel. More...

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags=FlagsT{})

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags)
	A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More...

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
	2D USM pointer block load. More...

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
	2D USM pointer block prefetch. More...

template<typename T , int BlockWidth, int BlockHeight = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false, false>()>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, sycl::ext::intel::esimd::simd< T, N > Vals)
	2D USM pointer block store. More...

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION sycl::ext::intel::esimd::simd< T, N >	sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload)
	A variation of `2D` stateless block load `with` parameters passed as `config_2d_mem_access` `object` Note: Compatibility with future hardware versions is not guaranteed. More...

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION void	sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload)
	A variation of `2D` stateless block prefetch `with` parameters passed as `config_2d_mem_access` `object` Note: Compatibility with future hardware versions is not guaranteed. More...

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION void	sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload, sycl::ext::intel::esimd::simd< T, N > Data)
	A variation of `2D` stateless block store `with` parameters passed as `config_2d_mem_access` `object` Note: Compatibility with future hardware versions is not guaranteed. More...

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...

template<sycl::ext::intel::esimd::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 OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&((Op !=sycl::ext::intel::esimd::atomic_op::store &&Op !=sycl::ext::intel::esimd::atomic_op::xchg)\|\|N==1), sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...

template<sycl::ext::intel::esimd::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 OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<sycl::ext::intel::esimd::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 Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1)

template<sycl::ext::intel::esimd::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< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&(Op==sycl::ext::intel::esimd::atomic_op::load\|\|__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write >), sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...

template<sycl::ext::intel::esimd::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 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of `lsc_atomic_update` that uses `local_accessor` as a parameter. More...

template<sycl::ext::intel::esimd::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< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...

template<sycl::ext::intel::esimd::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 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of `lsc_atomic_update` that uses `local_accessor` as a parameter. More...

template<sycl::ext::intel::esimd::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< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...

template<sycl::ext::intel::esimd::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 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of `lsc_atomic_update` that uses `local_accessor` as a parameter. More...

template<lsc_memory_kind Kind = lsc_memory_kind::untyped_global, lsc_fence_op FenceOp = lsc_fence_op::none, lsc_scope Scope = lsc_scope::group, int N = 16>
__ESIMD_API void	sycl::_V1::ext::intel::experimental::esimd::lsc_fence (sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Memory fence. More...

Detailed Description

This group combines types and functions specific to LSC, which is available in Intel GPUs starting from PVC and ACM.

Enumeration Type Documentation

◆ atomic_op

enum sycl::_V1::ext::intel::esimd::native::lsc::atomic_op : uint8_t

strong

LSC atomic operation codes.

atomic_update<native::lsc::atomic_op::inc>(...); is a short-cut to lsc_atomic_update<atomic_op::inc>(...); with default cache and data size controls.

Enumerator
inc
dec
load
store
add
sub
smin
smax
umin
umax
cmpxchg
fadd
fsub
fmin
fmax
fcmpxchg
bit_and
bit_or
bit_xor

Definition at line 39 of file common.hpp.

Function Documentation

◆ lsc_atomic_update() [1/14]

template<sycl::ext::intel::esimd::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<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
src0	is the first atomic operand (expected value).
src1	is the second atomic operand (new value).
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2650 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_atomic_update() [2/14]

template<sycl::ext::intel::esimd::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<__ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2590 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_atomic_update() [3/14]

template<sycl::ext::intel::esimd::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< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && (Op == sycl::ext::intel::esimd::atomic_op::load \|\| __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write>), sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2532 of file memory.hpp.

◆ lsc_atomic_update() [4/14]

template<sycl::ext::intel::esimd::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 >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Variant of lsc_atomic_update that uses local_accessor as a parameter.

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
src0	is the first atomic operand (expected value).
src1	is the second atomic operand (new value).
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2681 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_atomic_update() [5/14]

template<sycl::ext::intel::esimd::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 >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Variant of lsc_atomic_update that uses local_accessor as a parameter.

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2619 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_atomic_update() [6/14]

template<sycl::ext::intel::esimd::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 >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_rw_local_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

Variant of lsc_atomic_update that uses local_accessor as a parameter.

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2560 of file memory.hpp.

◆ lsc_atomic_update() [7/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 2, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.

Parameters

p	is the base pointer.
offsets	is the zero-based offsets.
src0	is the first atomic operand (expected value).
src1	is the second atomic operand (new value).
pred	predicates.

Definition at line 2469 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_atomic_update() [8/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 1, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.

Parameters

p	is the base pointer.
offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Definition at line 2411 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_atomic_update() [9/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 0, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.

Parameters

p	is the base pointer.
offsets	is the zero-based offsets.
pred	is predicates.

Definition at line 2371 of file memory.hpp.

◆ lsc_atomic_update() [10/14]

template<sycl::ext::intel::esimd::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 OffsetObjT , typename RegionTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 2, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 2483 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_atomic_update() [11/14]

template<sycl::ext::intel::esimd::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 OffsetObjT , typename RegionTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 1, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 2424 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_atomic_update() [12/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 2, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 2497 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_atomic_update() [13/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1 && ((Op != sycl::ext::intel::esimd::atomic_op::store && Op != sycl::ext::intel::esimd::atomic_op::xchg) \|\| N == 1), sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 2441 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_atomic_update() [14/14]

template<sycl::ext::intel::esimd::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 Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 0, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update	(	T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 2385 of file memory.hpp.

◆ lsc_block_load() [1/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		FlagsT	flags
	)

A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter.

Collects elements located at surface and returns them as a single simd object. When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.
flags	is the alignment specifier type tag.

Returns: is a vector of type T and size NElts. The elements of the returned vector for which the corresponding element in pred is 0 are undefined.

Definition at line 1134 of file memory.hpp.

◆ lsc_block_load() [2/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred,
		sycl::ext::intel::esimd::simd< T, NElts >	pass_thru,
		FlagsT	= `{}`
	)

Accessor-based transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at surface and returns them as a single simd object. When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.
FlagsT	is the alignment specifier type tag.

Parameters

acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.
pred	is operation predicate. Operation is skipped for index 'i' if pred[i] == 0 and the result element is taken from `pass_thru`[i]. Otherwise, the operation is performed.
pass_thru	contains the values copied to the result when the corresponding element from `pred` is zero.

Returns: is a vector of type T and size NElts

Definition at line 1195 of file memory.hpp.

◆ lsc_block_load() [3/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	flags = `FlagsT{}`
	)

Accessor-based transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at surface and returns them as a single simd object. When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

FlagsT	is the alignment specifier type tag.
acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.
pred	is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation.

Returns: is a vector of type T and size NElts. The elements of the returned vector for which the corresponding element in pred is 0 are undefined.

Definition at line 1075 of file memory.hpp.

◆ lsc_block_load() [4/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		uint32_t	offset,
		FlagsT	flags
	)

Definition at line 1149 of file memory.hpp.

◆ lsc_block_load() [5/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		uint32_t	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred,
		sycl::ext::intel::esimd::simd< T, NElts >	pass_thru,
		FlagsT	flags = `FlagsT{}`
	)

Definition at line 1213 of file memory.hpp.

◆ lsc_block_load() [6/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	AccessorTy	acc,
		uint32_t	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	flags = `FlagsT{}`
	)

Definition at line 1092 of file memory.hpp.

◆ lsc_block_load() [7/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	const T *	p,
		FlagsT
	)

A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter.

Accesses contiguous block of memory of NElts * S bytes starting from given address, where S is a byte size of an "element" defined by the DS template parameter. The maximum size of accessed block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
FlagsT	is the alignment specifier type tag.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.

Parameters

p	is the base pointer.

Returns: is a vector of type T and size NElts. The elements of the returned vector for which the corresponding element in pred is 0 are undefined.

Definition at line 983 of file memory.hpp.

◆ lsc_block_load() [8/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	const T *	p,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred,
		sycl::ext::intel::esimd::simd< T, NElts >	pass_thru,
		FlagsT	= `{}`
	)

USM pointer transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Accesses contiguous block of memory of NElts * S bytes starting from given address, where S is a byte size of an "element" defined by the DS template parameter. The maximum size of accessed block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
FlagsT	is the alignment specifier type tag.

Parameters

p	is the base pointer.
pred	is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed.
pass_thru	contains the vector which elements are copied to the returned result when the corresponding element of `pred` is 0.

Returns: is a vector of type T and size NElts.

Definition at line 1027 of file memory.hpp.

◆ lsc_block_load() [9/9]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT>, sycl::ext::intel::esimd::simd<T, NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_load	(	const T *	p,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	= `{}`
	)

USM pointer transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Accesses contiguous block of memory of NElts * S bytes starting from given address, where S is a byte size of an "element" defined by the DS template parameter. The maximum size of accessed block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8 bytes alignment is required for 64 bit data, 32 bit data and NElts equal to 128, 16 bit data and NElts equal to 256, 8 bit data and NElts equal to 512. Otherwise 4 bytes alignment is required.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
Flags	is the alignment specifier type tag.

Parameters

p	is the base pointer.
pred	is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation.

Returns: is a vector of type T and size NElts. The elements of the returned vector for which the corresponding element in pred is 0 are undefined.

Definition at line 942 of file memory.hpp.

◆ lsc_block_store() [1/5]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_store	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		FlagsT	flags
	)

A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter.

Scatters elements to surface. When DS equals lsc_data_size::u64 or sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-bytes aligned. Allowed values for the data size are lsc_data_size::u32, lsc_data_size::u64, lsc_data_size::u8, lsc_data_size::u16. When data size is either lsc_data_size::u8 or lsc_data_size::u16 the data is treated as 32 bit data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8 bytes alignment is required for 64 bit data, 32 bit data and NElts equal to 128, 16 bit data and NElts equal to 256, 8 bit data and NElts equal to 512. Otherwise 4 bytes alignment is required.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.
vals	is values to store.
flags	is the alignment specifier type tag.

Definition at line 1750 of file memory.hpp.

◆ lsc_block_store() [2/5]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_store	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	= `{}`
	)

Accessor-based transposed scatter with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_store.ugm

Scatters elements to surface. When DS equals lsc_data_size::u64 or sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-bytes aligned. Allowed values for the data size are lsc_data_size::u32, lsc_data_size::u64, lsc_data_size::u8, lsc_data_size::u16. When data size is either lsc_data_size::u8 or lsc_data_size::u16 the data is treated as 32 bit data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8 bytes alignment is required for 64 bit data, 32 bit data and NElts equal to 128, 16 bit data and NElts equal to 256, 8 bit data and NElts equal to 512. Otherwise 4 bytes alignment is required.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
Flags	is the alignment specifier type tag.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.
vals	is values to store.
pred	is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation.

Definition at line 1689 of file memory.hpp.

◆ lsc_block_store() [3/5]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_store	(	AccessorTy	acc,
		uint32_t	offset,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		FlagsT	flags = `FlagsT{}`
	)

Definition at line 1706 of file memory.hpp.

◆ lsc_block_store() [4/5]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_store	(	T *	p,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		FlagsT	flags
	)

A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter.

Scatters elements to specific address. When DS equals lsc_data_size::u64 or sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-bytes aligned. Allowed values for the data size are lsc_data_size::u32, lsc_data_size::u64, lsc_data_size::u8, lsc_data_size::u16. When data size is either lsc_data_size::u8 or lsc_data_size::u16 the data is treated as 32 bit data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8 bytes alignment is required for 64 bit data, 32 bit data and NElts equal to 128, 16 bit data and NElts equal to 256, 8 bit data and NElts equal to 512. Otherwise 4 bytes alignment is required.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.

Parameters

p	is the base pointer.
vals	is values to store.
flags	is the alignment specifier type tag.

Definition at line 1642 of file memory.hpp.

◆ lsc_block_store() [5/5]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_block_store	(	T *	p,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	= `{}`
	)

USM pointer transposed scatter with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_store.ugm

Scatters elements to specific address. When DS equals lsc_data_size::u64 or sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-bytes aligned. Allowed values for the data size are lsc_data_size::u32, lsc_data_size::u64, lsc_data_size::u8, lsc_data_size::u16. When data size is either lsc_data_size::u8 or lsc_data_size::u16 the data is treated as 32 bit data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8 bytes alignment is required for 64 bit data, 32 bit data and NElts equal to 128, 16 bit data and NElts equal to 256, 8 bit data and NElts equal to 512. Otherwise 4 bytes alignment is required.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size (unused/obsolete).
L1H	is L1 cache hint.
L2H	is L2 cache hint.
Flags	is the alignment specifier type tag.

Parameters

p	is the base pointer.
vals	is values to store.
pred	is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation.

Definition at line 1603 of file memory.hpp.

◆ lsc_fence()

template<lsc_memory_kind Kind = lsc_memory_kind::untyped_global, lsc_fence_op FenceOp = lsc_fence_op::none, lsc_scope Scope = lsc_scope::group, int N = 16>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_fence ( sycl::ext::intel::esimd::simd_mask< N > pred = 1 )

Memory fence.

Supported platforms: DG2, PVC

Template Parameters

Kind	is the Sfid shaded function.
FenceOp	is the fence operation.
Scope	is the operation scope.
N	is the number of channels (platform dependent).

Parameters

pred	is predicates.

Definition at line 2700 of file memory.hpp.

◆ lsc_gather() [1/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

Accessor-based gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at surface and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets in bytes.
pred	is predicates.
pass_thru	contains the vector which elements are copied to the returned result when the corresponding element of `pred` is 0.

Returns: is a vector of type T and size N * NElts

Definition at line 855 of file memory.hpp.

◆ lsc_gather() [2/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Accessor-based gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at surface and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets in bytes.
pred	is predicates.

Returns: is a vector of type T and size N * NElts

Definition at line 781 of file memory.hpp.

◆ lsc_gather() [3/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

Definition at line 895 of file memory.hpp.

◆ lsc_gather() [4/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 820 of file memory.hpp.

◆ lsc_gather() [5/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

USM pointer gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at specified address and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).

Parameters

p	is the base pointer.
offsets	is the zero-based offsets in bytes.
pred	is predicates.
pass_thru	contains the vector which elements are copied to the returned result when the corresponding element of `pred` is 0.

Returns: is a vector of type T and size N * NElts

Definition at line 701 of file memory.hpp.

◆ lsc_gather() [6/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

USM pointer gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at specified address and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).

Parameters

p	is the base pointer.
offsets	is the zero-based offsets in bytes.
pred	is predicates.

Returns: is a vector of type T and size N * NElts

Definition at line 668 of file memory.hpp.

◆ lsc_gather() [7/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

Definition at line 724 of file memory.hpp.

◆ lsc_gather() [8/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 714 of file memory.hpp.

◆ lsc_gather() [9/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

Definition at line 748 of file memory.hpp.

◆ lsc_gather() [10/10]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset>, sycl::ext::intel::esimd::simd<T, N * NElts> > sycl::_V1::ext::intel::experimental::esimd::lsc_gather	(	const T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 737 of file memory.hpp.

◆ lsc_load_2d() [1/2]

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>

ESIMD_INLINE SYCL_ESIMD_FUNCTION sycl::ext::intel::esimd::simd<T, N> sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d ( config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > & payload )

A variation of 2D stateless block load with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed.

Note: No software mitigation for hardware bugs is possible for this function.

Template Parameters

T	is the element data type
BlockWidth	the block width in number of elements
BlockHeight	block height in number of elements
NBlocks	Number of blocks
Transposed	is the transposed version or not.
Transformed	is apply VNNI transform or not.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

payload is config_2d_mem_access object holding all the data

Returns: is a vector of type T and size N, where N is getNextPowerOf2(Height) * Width * NBlocks, if transposed getNextPowerOf2(Width) * Height * NBlocks, otherwise

Definition at line 2100 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::check_cache_hints().

◆ lsc_load_2d() [2/2]

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N> sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d	(	const T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y
	)

2D USM pointer block load.

Supported platforms: PVC VISA instruction: lsc_load_block2d.ugm

Collects elements located at specified address and returns them as a single simd object.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
NBlocks	is the number of blocks.
Transposed	is the transposed version or not.
Transformed	is apply VNNI transform or not.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.

Returns: is a vector of type T and size N, where N is BlockWidth * BlockHeight * NBlocks, if transformed; otherwise, N = roundUpNextMultiple(BlockHeight, 4 / sizeof(T)) * getNextPowerOf2(BlockWidth) * NBlocks

Definition at line 1792 of file memory.hpp.

◆ lsc_prefetch() [1/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag, typename AccessorTy >

__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> && sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	AccessorTy	acc,
		__ESIMD_DNS::DeviceAccessorOffsetT	offset,
		FlagsT	flags = `FlagsT{}`
	)

Accessor-based transposed prefetch gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Prefetches elements located at surface of NElts * S bytes starting from given offset, where S is a byte size of an "element" defined by the DS template parameter. The maximum size of accessed block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
FlagsT	is the alignment specifier type tag.
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offset	is the zero-based offset in bytes.

Definition at line 1389 of file memory.hpp.

◆ lsc_prefetch() [2/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read> > sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Accessor-based prefetch gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Prefetches elements located at surface.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets in bytes.
pred	is predicates.

Definition at line 1325 of file memory.hpp.

◆ lsc_prefetch() [3/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API std::enable_if_t<sycl::ext::intel::esimd::is_simd_flag_type_v<FlagsT> > sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	const T *	p,
		FlagsT	= `{}`
	)

USM pointer prefetch transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Prefetches elements located at contiguous block of memory of NElts * S bytes starting from given address, where S is a byte size of an "element" defined by the DS template parameter. The maximum size of prefetched block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned. Also, 8-bytes alignment is required when the function has to load more than 256-bytes. In all other cases 4-byte alignment is required. When T is 1- or 2-byte type the data is treated as 4-byte data. Allowed NElts values for 64 bit data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 32 bit data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 16 bit data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 8 bit data are 4, 8, 12, 16, 32, 64, 128, 256, 512.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
FlagsT	is the alignment specifier type tag.

Parameters

p	is the base pointer.

Definition at line 1295 of file memory.hpp.

◆ lsc_prefetch() [4/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	const T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

USM pointer prefetch gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Prefetches elements located at specified address.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).

Parameters

p	is the base pointer.
offsets	is the zero-based offsets in bytes.
pred	is predicates.

Definition at line 1239 of file memory.hpp.

◆ lsc_prefetch() [5/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	const T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 1250 of file memory.hpp.

◆ lsc_prefetch() [6/6]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> > sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch	(	const T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 1260 of file memory.hpp.

◆ lsc_prefetch_2d() [1/2]

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>

ESIMD_INLINE SYCL_ESIMD_FUNCTION void sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d ( config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > & payload )

A variation of 2D stateless block prefetch with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed.

Note: No software mitigation for hardware bugs is possible for this function.

Template Parameters

T	is the element data type
BlockWidth	the block width in number of elements
BlockHeight	block height in number of elements
NBlocks	Number of blocks
Transposed	is the transposed version or not.
Transformed	is apply VNNI transform or not.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

payload is config_2d_mem_access object holding all the data

Definition at line 2196 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::check_cache_hints(), and sycl::_V1::ext::intel::esimd::prefetch().

◆ lsc_prefetch_2d() [2/2]

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d	(	const T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y
	)

2D USM pointer block prefetch.

Supported platforms: PVC VISA instruction: lsc_load_block2d.ugm

Prefetches elements located at specified address.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
NBlocks	is the number of blocks.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.

Definition at line 1826 of file memory.hpp.

◆ lsc_scatter() [1/5]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::experimental::esimd::lsc_scatter	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Accessor-based scatter.

Supported platforms: DG2, PVC VISA instruction: lsc_store.ugm

Scatters elements to surface.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).
AccessorTy	is the sycl::accessor type.

Parameters

acc	is the SYCL accessor.
offsets	is the zero-based offsets in bytes.
vals	is values to store.
pred	is predicates.

Definition at line 1523 of file memory.hpp.

◆ lsc_scatter() [2/5]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::experimental::esimd::lsc_scatter	(	AccessorTy	acc,
		sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 1560 of file memory.hpp.

◆ lsc_scatter() [3/5]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_scatter	(	T *	p,
		sycl::ext::intel::esimd::simd< Toffset, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

USM pointer scatter.

Supported platforms: DG2, PVC VISA instruction: lsc_store.ugm

Scatters elements to specific address.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the number of channels (platform dependent).

Parameters

p	is the base pointer.
offsets	is the zero-based offsets in bytes.
vals	is values to store.
pred	is predicates.

Definition at line 1469 of file memory.hpp.

◆ lsc_scatter() [4/5]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_scatter	(	T *	p,
		sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy >	offsets,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 1482 of file memory.hpp.

◆ lsc_scatter() [5/5]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && N == 1> sycl::_V1::ext::intel::experimental::esimd::lsc_scatter	(	T *	p,
		Toffset	offset,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

Definition at line 1493 of file memory.hpp.

◆ lsc_slm_atomic_update() [1/3]

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd< T, N >	src1,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

offsets	is the zero-based offsets.
src0	is the first atomic operand (expected value).
src1	is the second atomic operand (new value).
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2344 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ lsc_slm_atomic_update() [2/3]

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N >	src0,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2320 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ lsc_slm_atomic_update() [3/3]

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred
	)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

offsets	is the zero-based offsets.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 2298 of file memory.hpp.

◆ lsc_slm_block_load() [1/2]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API sycl::ext::intel::esimd::simd<T, NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load	(	uint32_t	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred,
		sycl::ext::intel::esimd::simd< T, NElts >	pass_thru
	)

Transposed SLM gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.slm

Collects elements located at slm and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).

Parameters

offset	is the zero-based offset for SLM buffer in bytes.
pred	is the predicate; if it contains 0, then the actual load is not performed and `pass_thru` is returned.
pass_thru	contains the vector that is returned if the parameter `pred` contains 0.

Returns: is a vector of type T and size NElts.

Definition at line 638 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::alignment.

◆ lsc_slm_block_load() [2/2]

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API sycl::ext::intel::esimd::simd<T, NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load	(	uint32_t	offset,
		sycl::ext::intel::esimd::simd_mask< 1 >	pred = `1`,
		FlagsT	flags = `FlagsT{}`
	)

Transposed SLM gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.slm

Collects elements located at slm and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size (unused/obsolete).

Parameters

offset	is the zero-based offset for SLM buffer in bytes.
pred	is the predicate; if it contains 0, then the actual load is not performed and the returned value is undefined.

Returns: is a vector of type T and size NElts

Definition at line 611 of file memory.hpp.

◆ lsc_slm_block_store()

template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_store	(	uint32_t	offset,
		sycl::ext::intel::esimd::simd< T, NElts >	vals,
		FlagsT	flags = `FlagsT{}`
	)

Transposed SLM scatter with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_store.slm

Scatters elements located to slm.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size (unused/obsolete).

Parameters

offset	is the zero-based offset for SLM buffer in bytes.
vals	is values to store.

Definition at line 1438 of file memory.hpp.

◆ lsc_slm_gather() [1/2]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred,
		sycl::ext::intel::esimd::simd< T, N *NElts >	pass_thru
	)

SLM gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.slm

Collects elements located at slm and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
N	is the number of channels (platform dependent).

Parameters

offsets	is the zero-based offsets for SLM buffer in bytes.
pred	is predicates.
pass_thru	values copied to the result when the corresponding element of `pred` is zero..

Returns: is a vector of type T and size N * NElts

Definition at line 587 of file memory.hpp.

◆ lsc_slm_gather() [2/2]

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>

__ESIMD_API sycl::ext::intel::esimd::simd<T, N * NElts> sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

SLM gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.slm

Collects elements located at slm and returns them as a single simd object.

Template Parameters

T	is element type.
NElts	is the number of elements to load per address.
DS	is the data size.
N	is the number of channels (platform dependent).

Parameters

offsets	is the zero-based offsets for SLM buffer in bytes.
pred	is predicates.

Returns: is a vector of type T and size N * NElts

Definition at line 561 of file memory.hpp.

◆ lsc_slm_scatter()

template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_slm_scatter	(	sycl::ext::intel::esimd::simd< uint32_t, N >	offsets,
		sycl::ext::intel::esimd::simd< T, N *NElts >	vals,
		sycl::ext::intel::esimd::simd_mask< N >	pred = `1`
	)

SLM scatter.

Supported platforms: DG2, PVC VISA instruction: lsc_store.slm

Scatters elements located to slm.

Template Parameters

T	is element type.
NElts	is the number of elements to store per address.
DS	is the data size.
N	is the number of channels (platform dependent).

Parameters

offsets	is the zero-based offsets for SLM buffer in bytes.
vals	is values to store.
pred	is predicates.

Definition at line 1418 of file memory.hpp.

◆ lsc_store_2d() [1/2]

template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>

ESIMD_INLINE SYCL_ESIMD_FUNCTION void sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d	(	config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &	payload,
		sycl::ext::intel::esimd::simd< T, N >	Data
	)

A variation of 2D stateless block store with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed.

Note: No software mitigation for hardware bugs is possible for this function.

Template Parameters

T	is the element data type
BlockWidth	the block width in number of elements
BlockHeight	block height in number of elements
NBlocks	Number of blocks
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

payload	is `config_2d_mem_access` `object` holding all the data
Data	is the data to be stored.

Definition at line 2242 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::check_cache_hints().

◆ lsc_store_2d() [2/2]

template<typename T , int BlockWidth, int BlockHeight = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false, false>()>

__ESIMD_API void sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d	(	T *	Ptr,
		unsigned	SurfaceWidth,
		unsigned	SurfaceHeight,
		unsigned	SurfacePitch,
		int	X,
		int	Y,
		sycl::ext::intel::esimd::simd< T, N >	Vals
	)

2D USM pointer block store.

Supported platforms: PVC VISA instruction: lsc_store_block2d.ugm

Stores elements at specified address.

Template Parameters

T	is element type.
BlockWidth	is the block width in number of elements.
BlockHeight	is the block height in number of elements.
L1H	is L1 cache hint.
L2H	is L2 cache hint.
N	is the data size

Parameters

Ptr	is the surface base address for this operation.
SurfaceWidth	is the surface width minus 1 in bytes
SurfaceHeight	is the surface height minus 1 in rows
SurfacePitch	is the surface pitch minus 1 in bytes
X	is zero based X-coordinate of the left upper rectangle corner in number of elements.
Y	is zero based Y-coordinate of the left upper rectangle corner in rows.
Vals	is a vector to store of type T and size N, where N = roundUpNextMultiple(BlockHeight, 4 / sizeof(T)) * getNextPowerOf2(BlockWidth) * NBlocks

Definition at line 1863 of file memory.hpp.

Classes

Enumerations

Functions

Detailed Description

Enumeration Type Documentation

◆ atomic_op

Function Documentation

◆ lsc_atomic_update() [1/14]

◆ lsc_atomic_update() [2/14]

◆ lsc_atomic_update() [3/14]

◆ lsc_atomic_update() [4/14]

◆ lsc_atomic_update() [5/14]

◆ lsc_atomic_update() [6/14]

◆ lsc_atomic_update() [7/14]

◆ lsc_atomic_update() [8/14]

◆ lsc_atomic_update() [9/14]

◆ lsc_atomic_update() [10/14]

◆ lsc_atomic_update() [11/14]

◆ lsc_atomic_update() [12/14]

◆ lsc_atomic_update() [13/14]

◆ lsc_atomic_update() [14/14]

◆ lsc_block_load() [1/9]

◆ lsc_block_load() [2/9]

◆ lsc_block_load() [3/9]

◆ lsc_block_load() [4/9]

◆ lsc_block_load() [5/9]

◆ lsc_block_load() [6/9]

◆ lsc_block_load() [7/9]

◆ lsc_block_load() [8/9]

◆ lsc_block_load() [9/9]

◆ lsc_block_store() [1/5]

◆ lsc_block_store() [2/5]

◆ lsc_block_store() [3/5]

◆ lsc_block_store() [4/5]

◆ lsc_block_store() [5/5]

◆ lsc_fence()

◆ lsc_gather() [1/10]

◆ lsc_gather() [2/10]

◆ lsc_gather() [3/10]

◆ lsc_gather() [4/10]

◆ lsc_gather() [5/10]

◆ lsc_gather() [6/10]

◆ lsc_gather() [7/10]

◆ lsc_gather() [8/10]

◆ lsc_gather() [9/10]

◆ lsc_gather() [10/10]

◆ lsc_load_2d() [1/2]

◆ lsc_load_2d() [2/2]

◆ lsc_prefetch() [1/6]

◆ lsc_prefetch() [2/6]

◆ lsc_prefetch() [3/6]

◆ lsc_prefetch() [4/6]

◆ lsc_prefetch() [5/6]

◆ lsc_prefetch() [6/6]

◆ lsc_prefetch_2d() [1/2]

◆ lsc_prefetch_2d() [2/2]

◆ lsc_scatter() [1/5]

◆ lsc_scatter() [2/5]

◆ lsc_scatter() [3/5]

◆ lsc_scatter() [4/5]

◆ lsc_scatter() [5/5]

◆ lsc_slm_atomic_update() [1/3]

◆ lsc_slm_atomic_update() [2/3]

◆ lsc_slm_atomic_update() [3/3]

◆ lsc_slm_block_load() [1/2]

◆ lsc_slm_block_load() [2/2]

◆ lsc_slm_block_store()

◆ lsc_slm_gather() [1/2]

◆ lsc_slm_gather() [2/2]

◆ lsc_slm_scatter()

◆ lsc_store_2d() [1/2]

◆ lsc_store_2d() [2/2]