sycl::_V1::ext::intel::esimd Namespace Reference

Namespaces
	detail
	native
	xmx

Classes
struct	saturation_on_tag
	Gen hardware supports applying saturation to results of certain operations. More...
struct	saturation_off_tag
	This type tag represents "saturation off" behavior. More...
struct	element_aligned_tag
	element_aligned_tag type. More...
struct	vector_aligned_tag
	vector_aligned_tag type. More...
struct	overaligned_tag
	overaligned_tag type. More...
struct	is_simd_flag_type
	Checks if type is a simd load/store flag. More...
struct	is_simd_flag_type< element_aligned_tag >
struct	is_simd_flag_type< vector_aligned_tag >
struct	is_simd_flag_type< overaligned_tag< N > >
struct	is_simd_flag_type< detail::dqword_element_aligned_tag >
class	simd
	The main simd vector class. More...
class	simd_view
	This class represents a reference to a sub-region of a base simd object. More...
class	properties
struct	cache_hint_L1_key
	The 'cache_hint_L1', 'cache_hint_L2' and 'cache_hint_L3' properties are used to specify L1, L2, L3 cache hints available in target device. More...
struct	cache_hint_L2_key
struct	cache_hint_L3_key
class	simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >
	This is a specialization of simd_view class with a single element. More...
class	simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >
	This is a specialization of nested simd_view class with a single element. More...
class	slm_allocator
	RAII-style class used to implement "semi-dynamic" SLM allocation. More...

Typedefs
using	uchar = unsigned char
using	ushort = unsigned short
using	uint = unsigned int
using	SurfaceIndex = unsigned int
	Surface index type. More...
template<int N>
using	mask_type_t = detail::simd_mask_storage_t< N >
using	alignment_key = sycl::ext::oneapi::experimental::alignment_key
	The 'alignment' property is used to specify the alignment of memory accessed in ESIMD memory operation such as block_load(). More...
using	default_cache_hint_L1 = cache_hint_L1_key::value_t< cache_hint::none >
using	default_cache_hint_L2 = cache_hint_L2_key::value_t< cache_hint::none >
using	default_cache_hint_L3 = cache_hint_L3_key::value_t< cache_hint::none >
template<int N>
using	simd_mask = detail::simd_mask_type< N >
	Represents a simd mask os size N. More...

Enumerations
enum class	rgba_channel : uint8_t { R , G , B , A }
	Represents a pixel's channel. More...
enum class	raw_send_eot : uint8_t { not_eot = 0 , eot = 1 }
	Specify if end of thread should be set. More...
enum class	raw_send_sendc : uint8_t { not_sendc = 0 , sendc = 1 }
	Specify if sendc should be used. More...
enum class	rgba_channel_mask : uint8_t {   R = detail::chR , G = detail::chG , GR = detail::chG | detail::chR , B = detail::chB ,   BR = detail::chB | detail::chR , BG = detail::chB | detail::chG , BGR = detail::chB | detail::chG | detail::chR , A = detail::chA ,   AR = detail::chA | detail::chR , AG = detail::chA | detail::chG , AGR = detail::chA | detail::chG | detail::chR , AB = detail::chA | detail::chB ,   ABR = detail::chA | detail::chB | detail::chR , ABG = detail::chA | detail::chB | detail::chG , ABGR = detail::chA | detail::chB | detail::chG | detail::chR }
	Represents a pixel's channel mask - all possible combinations of enabled channels. More...
enum class	atomic_op : uint8_t {   add = 0x0 , sub = 0x1 , inc = 0x2 , dec = 0x3 ,   umin = 0x4 , umax = 0x5 , xchg = 0x6 , cmpxchg = 0x7 ,   bit_and = 0x8 , bit_or = 0x9 , bit_xor = 0xa , smin = 0xb ,   smax = 0xc , fmax = 0x10 , fmin = 0x11 , fcmpxchg = 0x12 ,   fcmpwr = fcmpxchg , fadd = 0x13 , fsub = 0x14 , load = 0x15 ,   store = 0x16 , predec = 0xff }
	Represents an atomic operation. More...
enum class	fence_scope : uint8_t {   group = 0 , local = 1 , tile = 2 , gpu = 3 ,   gpus = 4 , system = 5 , system_acquire = 6 }
	The scope that fence() operation should apply to. More...
enum class	fence_flush_op : uint8_t { none = 0 , evict = 1 , invalidate = 2 , clean = 4 }
	The cache flush operation to apply to caches after fence() is complete. More...
enum class	memory_kind : uint8_t { global = 0 , image = 2 , local = 3 }
	The target memory kind for fence() operation. More...
enum class	bfn_t : uint8_t { x = 0xAA , y = 0xCC , z = 0xF0 }
	This enum is used to encode all possible logical operations performed on the 3 input operands. More...
enum	fence_mask : uint8_t {   global_coherent_fence = 0x1 , l2_flush_instructions = 0x2 , __SYCL_DEPRECATED , l2_flush_texture_data = 0x4 ,   __SYCL_DEPRECATED , l2_flush_constant_data = 0x8 , __SYCL_DEPRECATED , l2_flush_rw_data = 0x10 ,   __SYCL_DEPRECATED , local_barrier = 0x20 , l1_flush_ro_data = 0x40 , __SYCL_DEPRECATED }
	Represetns a bit mask to control behavior of esimd::fence. More...
enum class	cache_level : uint8_t { L1 = 1 , L2 = 2 , L3 = 3 }
	L1, L2 or L3 cache hint levels. L3 is reserved for future use. More...
enum class	cache_hint : uint8_t {   none = 0 , uncached = 1 , cached = 2 , write_back = 3 ,   write_through = 4 , streaming = 5 , read_invalidate = 6 , const_cached = 7 }
	L1, L2 or L3 cache hints. More...

Functions
template<class T , int N>
__ESIMD_API simd< T, N >	merge (simd< T, N > a, simd< T, N > b, simd_mask< N > m)
	"Merges" elements of the input simd object according to the merge mask. More...
template<int N>
__ESIMD_API simd_mask< N >	merge (simd_mask< N > a, simd_mask< N > b, simd_mask< N > m)
	"Merges" elements of the input masks according to the merge mask. More...
template<class BaseT1 , class BaseT2 , class RegionT1 , class RegionT2 , class = std::enable_if_t< (shape_type<RegionT1>::length == shape_type<RegionT2>::length) && std::is_same_v<detail::element_type_t<BaseT1>, detail::element_type_t<BaseT2>>>>
__ESIMD_API auto	merge (simd_view< BaseT1, RegionT1 > v1, simd_view< BaseT2, RegionT2 > v2, simd_mask< shape_type< RegionT1 >::length > m)
	"Merges" elements of vectors referenced by the input views. More...
constexpr int	is_channel_enabled (rgba_channel_mask M, rgba_channel Ch)
constexpr int	get_num_channels_enabled (rgba_channel_mask M)
template<typename T0 , typename T1 , int SZ>
__ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v< T0 >||std::is_same_v< T1, T0 >, simd< T0, SZ > >	saturate (simd< T1, SZ > src)
	Conversion of input vector elements of type T1 into vector of elements of type T0 with saturation. More...
template<typename TRes , typename TArg , int SZ>
ESIMD_DETAIL __ESIMD_API std::enable_if_t< !std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value, simd< TRes, SZ > >	abs (simd< TArg, SZ > src0)
	Get absolute value (vector version) More...
template<typename TRes , typename TArg >
__ESIMD_API std::enable_if_t<!std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value &&detail::is_esimd_scalar< TRes >::value &&detail::is_esimd_scalar< TArg >::value, std::remove_const_t< TRes > >	abs (TArg src0)
	Get absolute value (scalar version) More...
template<typename T1 , int SZ>
__ESIMD_API simd< T1, SZ >	abs (simd< T1, SZ > src0)
	Get absolute value (vector version). More...
template<typename T1 >
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T1 >::value, std::remove_const_t< T1 > >	abs (T1 src0)
	Get absolute value (scalar version). More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API	simd< T, SZ > (max)(simd< T
	Selects component-wise the maximum of the two vectors. More...
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >()	max (simd< T, SZ > src0, T src1, Sat sat={})
__ESIMD_API simd< T, SZ > Sat class T()	max (T src0, T src1, Sat sat={})
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API	simd< T, SZ > (min)(simd< T
	Selects component-wise the minimum of the two vectors. More...
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >()	min (simd< T, SZ > src0, T src1, Sat sat={})
__ESIMD_API simd< T, SZ > Sat class T()	min (T src0, T src1, Sat sat={})
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	inv (simd< T, N > src, Sat sat={})
	Inversion - calculates (1/x). More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	inv (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	log2 (simd< T, N > src, Sat sat={})
	Logarithm base 2. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	log2 (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	exp2 (simd< T, N > src, Sat sat={})
	Exponent base 2. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	exp2 (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sqrt (simd< T, N > src, Sat sat={})
	Square root. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sqrt (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	sqrt_ieee (simd< T, N > src, Sat sat={})
	IEEE754-compliant square root. Supports float and double. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API T	sqrt_ieee (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	rsqrt (simd< T, N > src, Sat sat={})
	Square root reciprocal - calculates 1/sqrt(x). More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	rsqrt (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	sin (simd< T, N > src, Sat sat={})
	Sine. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	sin (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	cos (simd< T, N > src, Sat sat={})
	Cosine. More...
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	cos (T src, Sat sat={})
	Scalar version. More...
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	pow (simd< T, N > src0, simd< U, N > src1, Sat sat={})
	Power - calculates src0 in power of src1. More...
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N >	pow (simd< T, N > src0, U src1, Sat sat={})
	(vector, scalar) version. More...
template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T	pow (T src0, U src1, Sat sat={})
	(scalar, scalar) version. More...
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	div_ieee (simd< T, N > src0, simd< U, N > src1, Sat sat={})
	IEEE754-compliant floating-point division. Supports float and double. More...
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N >	div_ieee (simd< T, N > src0, U src1, Sat sat={})
	(vector, scalar) version. More...
template<class T , class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API T	div_ieee (T src0, U src1, Sat sat={})
	(scalar, scalar) version. More...
template<class T , int SZ, class Sat = saturation_off_tag>
ESIMD_DETAIL ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ >	log (simd< T, SZ > src0, Sat sat={})
	Computes the natural logarithm of the given argument. More...
template<class T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T	log (T src0, Sat sat={})
template<class T , int SZ, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ >	exp (simd< T, SZ > src0, Sat sat={})
	Computes e raised to the power of the given argument. More...
template<class T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T	exp (T src0, Sat sat={})
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ >	rndd (sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
	Round-down (also known as floor). More...
template<typename T , class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API T	rndd (float src0, Sat sat={})
	Scalar version. More...
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ >	rndu (sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
	Round-up (also known as ceil). More...
template<typename T , class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API T	rndu (float src0, Sat sat={})
	Scalar version. More...
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ >	rnde (sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
	Round-to-even (also known as round). More...
template<typename T , class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API T	rnde (float src0, Sat sat={})
	Scalar version. More...
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ >	rndz (sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
	Round-to-zero (also known as trunc). More...
template<typename T , class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API T	rndz (float src0, Sat sat={})
	Scalar version. More...
template<typename RT , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ >	floor (const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
	"Floor" operation, vector version - alias of rndd. More...
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE RT	floor (float src0, Sat sat={})
	"Floor" operation, scalar version - alias of rndd. More...
template<typename RT , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ >	ceil (const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
	"Ceiling" operation, vector version - alias of rndu. More...
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE RT	ceil (float src0, Sat sat={})
	"Ceiling" operation, scalar version - alias of rndu. More...
template<typename RT , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ >	trunc (const sycl::ext::intel::esimd::simd< float, SZ > &src0, Sat sat={})
	Round to integral value using the round to zero rounding mode (vector version). More...
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API RT	trunc (float src0, Sat sat={})
	Round to integral value using the round to zero rounding mode (scalar version). More...
template<int N>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<(N==8||N==16||N==32), uint >	pack_mask (simd_mask< N > src0)
	Pack a simd_mask into a single unsigned 32-bit integer value. More...
template<int N>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<(N==8||N==16||N==32), simd_mask< N > >	unpack_mask (uint src0)
	Unpack an unsigned 32-bit integer value into a simd_mask. More...
template<int N>
__ESIMD_API std::enable_if_t<(N !=8 &&N !=16 &&N< 32), uint >	pack_mask (simd_mask< N > src0)
	pack_mask specialization when the number of elements N is not 8, 16 or 32. More...
template<typename T , int N>
__ESIMD_API std::enable_if_t<(std::is_same_v< T, ushort >||std::is_same_v< T, uint >)&&(N > 0 &&N<=32), uint >	ballot (simd< T, N > mask)
	Compare source vector elements against zero and return a bitfield combining the comparison result. More...
template<typename T , int N>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)<=4), simd< uint32_t, N > >	cbit (simd< T, N > src)
	Count number of bits set in the source operand per element. More...
template<typename T >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)<=4), uint32_t >	cbit (T src)
	Scalar version of cbit - both input and output are scalars rather than vectors. More...
template<typename BaseTy , typename RegionTy >
__ESIMD_API std::enable_if_t< std::is_integral< typename simd_view< BaseTy, RegionTy >::element_type >::value &&(sizeof(typename simd_view< BaseTy, RegionTy >::element_type)<=4) &&(simd_view< BaseTy, RegionTy >::length==1), uint32_t >	cbit (simd_view< BaseTy, RegionTy > src)
	Scalar version of cbit, that takes simd_view object as an argument, e.g. More...
template<typename T , int N>
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)==4), simd< T, N > >	fbl (simd< T, N > src)
	Find the per element number of the first bit set in the source operand starting from the least significant bit. More...
template<typename T >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)==4), T >	fbl (T src)
	Scalar version of fbl - both input and output are scalars rather than vectors. More...
template<typename BaseTy , typename RegionTy >
__ESIMD_API std::enable_if_t< std::is_integral< typename simd_view< BaseTy, RegionTy >::element_type >::value &&(sizeof(typename simd_view< BaseTy, RegionTy >::element_type)==4) &&(simd_view< BaseTy, RegionTy >::length==1), typename simd_view< BaseTy, RegionTy >::element_type >	fbl (simd_view< BaseTy, RegionTy > src)
	Scalar version of fbl, that takes simd_view object as an argument, e.g. More...
template<typename T , int N>
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_signed< T >::value &&(sizeof(T)==4), simd< T, N > >	fbh (simd< T, N > src)
	Find the per element number of the first bit set in the source operand starting from the most significant bit (sign bit is skipped). More...
template<typename T , int N>
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&!std::is_signed< T >::value &&(sizeof(T)==4), simd< T, N > >	fbh (simd< T, N > src)
	Find the per element number of the first bit set in the source operand starting from the most significant bit (sign bit is counted). More...
template<typename T >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)==4), T >	fbh (T src)
	Scalar version of fbh - both input and output are scalars rather than vectors. More...
template<typename BaseTy , typename RegionTy >
__ESIMD_API std::enable_if_t< std::is_integral< typename simd_view< BaseTy, RegionTy >::element_type >::value &&(sizeof(typename simd_view< BaseTy, RegionTy >::element_type)==4) &&(simd_view< BaseTy, RegionTy >::length==1), typename simd_view< BaseTy, RegionTy >::element_type >	fbh (simd_view< BaseTy, RegionTy > src)
	Scalar version of fbh, that takes simd_view object as an argument, e.g. More...
template<typename T1 , typename T2 , typename T3 , typename T4 , int N, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_dword_type< T1 >::value &&detail::is_dword_type< T2 >::value &&detail::is_dword_type< T3 >::value &&detail::is_dword_type< T4 >::value, simd< T1, N > >	dp4a (simd< T2, N > src0, simd< T3, N > src1, simd< T4, N > src2, Sat sat={})
	DP4A. More...
template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0	hmax (simd< T1, SZ > v)
	ESIMD_DETAIL. More...
template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0	hmin (simd< T1, SZ > v)
	Performs 'minimum' operation reduction over elements of the input vector, that is, returns the minimal vector element. More...
template<typename T0 , typename T1 , int SZ, typename BinaryOperation >
ESIMD_INLINE ESIMD_NODEBUG T0	reduce (simd< T1, SZ > v, BinaryOperation op)
	Performs reduction over elements of the input vector. More...
static constexpr bfn_t	operator~ (bfn_t x)
static constexpr bfn_t	operator| (bfn_t x, bfn_t y)
static constexpr bfn_t	operator& (bfn_t x, bfn_t y)
static constexpr bfn_t	operator^ (bfn_t x, bfn_t y)
template<bfn_t FuncControl, typename T , int N>
__ESIMD_API std::enable_if_t< std::is_integral_v< T >, sycl::ext::intel::esimd::simd< T, N > >	bfn (sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd< T, N > src2)
	Performs binary function computation with three vector operands. More...
template<bfn_t FuncControl, typename T >
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T >::value &&std::is_integral_v< T >, T >	bfn (T src0, T src1, T src2)
	Performs binary function computation with three scalar operands. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	addc (sycl::ext::intel::esimd::simd< uint32_t, N > &carry, sycl::ext::intel::esimd::simd< uint32_t, N > src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
	Performs add with carry of 2 unsigned 32-bit vectors. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	addc (sycl::ext::intel::esimd::simd< uint32_t, N > &carry, sycl::ext::intel::esimd::simd< uint32_t, N > src0, uint32_t src1)
	Performs add with carry of a unsigned 32-bit vector and scalar. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	addc (sycl::ext::intel::esimd::simd< uint32_t, N > &carry, uint32_t src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
	Performs add with carry of a unsigned 32-bit scalar and vector. More...
__ESIMD_API uint32_t	addc (uint32_t &carry, uint32_t src0, uint32_t src1)
	Performs add with carry of a unsigned 32-bit scalars. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	subb (sycl::ext::intel::esimd::simd< uint32_t, N > &borrow, sycl::ext::intel::esimd::simd< uint32_t, N > src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
	Performs substraction with borrow of 2 unsigned 32-bit vectors. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	subb (sycl::ext::intel::esimd::simd< uint32_t, N > &borrow, sycl::ext::intel::esimd::simd< uint32_t, N > src0, uint32_t src1)
	Performs substraction with borrow of unsigned 32-bit vector and scalar. More...
template<int N>
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N >	subb (sycl::ext::intel::esimd::simd< uint32_t, N > &borrow, uint32_t src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
	Performs substraction with borrow of unsigned 32-bit scalar and vector. More...
__ESIMD_API uint32_t	subb (uint32_t &borrow, uint32_t src0, uint32_t src1)
	Performs substraction with borrow of 2 unsigned 32-bit scalars. More...
__ESIMD_API uint64_t	rdtsc ()
	rdtsc - get the value of timestamp counter. More...
template<typename AccessorTy >
__ESIMD_API SurfaceIndex	get_surface_index (AccessorTy acc)
	Get surface index corresponding to a SYCL accessor. More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-1) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-2) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-ga-3) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-4) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (usm-ga-5) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	gather (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-ga-6) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	gather (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (usm-ga-7) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	gather (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-ga-8) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	gather (const T *p, OffsetSimdViewT byte_offsets, PropertyListT props={})
	simd <T, N> gather(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-ga-9) Loads ("gathers") elements of the type 'T' from memory locations addressed by the base pointer p and byte offsets byte_offsets, and returns the loaded elements. More...
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< Tx, N > >	gather (const Tx *p, Toffset offset, simd_mask< N > mask=1)
	A variation of gather API with offsets represented as scalar. More...
template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(T *p, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-sc-1) More...
template<typename T , int N, int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (T *p, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	Writes ("scatters") elements of the input vector to different memory locations. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Writes ("scatters") elements of the input vector to different memory locations. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (T *p, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(T *p, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (usm-sc-4) More...
template<typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	scatter (Tx *p, Toffset offset, simd< Tx, N > vals, simd_mask< N > mask=1)
	A variation of scatter API with offsets represented as scalar. More...
template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags > >	block_store (Tx *addr, simd< Tx, N > vals, Flags)
	Stores elements of the vector vals to a contiguous block of memory at the given address addr. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, PropertyListT props={})
	Each of the following block load functions loads a contiguous memory block from the address referenced by the USM pointer 'ptr', or from 'ptr + offset', where 'offset' is the offset in bytes (not in elements!). More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, size_t byte_offset, PropertyListT props={})
	simd<T, N> block_load(const T* ptr, size_t byte_offset, props={}); // (usm-bl-2) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> block_load(const T* ptr, simd_mask<1> pred, props={}); // (usm-bl-3) This function loads a contiguous memory block from USM pointer ptr. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, size_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> block_load(const T* ptr, size_t byte_offset, simd_mask<1> pred, props={}); // (usm-bl-4) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
	simd<T, N> block_load(const T* ptr, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (usm-bl-5) This function loads a contiguous memory block from USM pointer ptr. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (const T *ptr, size_t byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
	simd<T, N> block_load(const T* ptr, size_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (usm-bl-6) This function loads a contiguous memory block from address referenced by USM pointer ptr and the given byte_offset. More...
template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags >, simd< Tx, N > >	block_load (const Tx *addr, Flags)
	Loads a contiguous block of memory from the given memory address addr and returns the loaded data as a vector. More...
template<typename Tx , int N, typename AccessorTy , typename Flags = vector_aligned_tag, typename = std::enable_if_t< is_simd_flag_type_v<Flags> && detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read>>, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N >	block_load (AccessorTy acc, detail::DeviceAccessorOffsetT byte_offset, Flags flags)
	Loads a contiguous block of memory from the given accessor acc and byte_offset and returns the loaded data as a vector. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, PropertyListT props={})
	Each of the following block load functions loads a contiguous memory block from the address referenced by accessor 'acc', or from 'acc + byte_offset', The parameter 'pred' is the one element predicate. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, PropertyListT={})
	simd<T, N> block_load(AccessorT acc, props = {}); // (acc-bl-2) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT={})
	simd<T, N> block_load(AccessorT acc, OffsetT byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props = {}); // (acc-bl-3) This function loads a contiguous memory block referenced by accessor acc and the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> block_load(AccessorT acc, OffsetT byte_offset, simd_mask<1> pred, props = {}); // (acc-bl-4) This function loads a contiguous memory block referenced by accessor acc and the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT={})
	simd<T, N> block_load(AccessorT acc, simd_mask<1> pred, simd<T, N> pass_thru, props = {}); // (acc-bl-5) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	block_load (AccessorT acc, simd_mask< 1 > pred, PropertyListT={})
	simd<T, N> block_load(AccessorT acc, simd_mask<1> pred, props = {}); // (acc-bl-6) This function loads a contiguous memory block referenced by accessor acc and implied offset=0. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_property_list_v< PropertyListT > >	block_store (T *ptr, simd< T, N > vals, PropertyListT={})
	Each of the following block store functions stores a contiguous memory block to the address referenced by the USM pointer 'ptr', or from 'ptr + offset', where 'offset' is the offset in bytes (not in elements!) with data specified by 'vals'. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (T *ptr, size_t byte_offset, simd< T, N > vals, PropertyListT props={})
	void block_store(T* ptr, size_t byte_offset, // (usm-bs-2) simd<T, N> vals, props={}); This function stores a contiguous memory block to USM pointer ptr and byte-offset byte_offset with data specified by vals. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_property_list_v< PropertyListT > >	block_store (T *ptr, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT={})
	void block_store(T* ptr, simd<T, N> vals, // (usm-bs-3) simd_mask<1> pred, props={}); This function stores a contiguous memory block to USM pointer ptr with data specified by vals. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (T *ptr, size_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	void block_store(T* ptr, size_t byte_offset, // (usm-bs-4) simd<T, N> vals, simd_mask<1> pred, props={}); This function stores a contiguous memory block to USM pointer ptr and byte-offset byte_offset with data specified by vals. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	block_store (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd< T, N > vals, PropertyListT props={})
	Each of the following block_store functions stores the vector 'vals' to a contiguous memory block at the address referenced by accessor 'acc', or from 'acc + byte_offset', The parameter 'pred' is the one element predicate. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	block_store (AccessorT acc, simd< T, N > vals, PropertyListT props={})
	void block_store(AccessorT acc, simd<T, N> vals, props = {}); // (acc-bs-2) This function stores a contiguous memory block to accessor acc with data specified by vals and implied offset=0. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	block_store (AccessorT acc, detail::DeviceAccessorOffsetT byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	void block_store(AccessorT acc, OffsetT byte_offset, // (acc-bs-3) simd<T, N> vals, simd_mask<1> pred, props = {}); This function stores a contiguous memory block to accessor acc and byte_offset with data specified by vals. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	block_store (AccessorT acc, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	void block_store(AccessorT acc, simd<T, N> vals, // (acc-bs-4) simd_mask<1> pred, props = {}); This function stores a contiguous memory block to accessor acc with data specified by vals and implied offset=0. More...
template<typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	gather (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > byte_offsets, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
	Accessor-based gather. More...
template<typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N > >	gather (AccessorT acc, detail::DeviceAccessorOffsetT glob_offset)
	Loads and broadcasts the element located at acc and byte offset glob_offset to a vector and returns it as a simd object. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-1) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-ga-2) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-ga-3) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (acc-ga-4) This function is identical to (acc-ga-1) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N > byte_offsets, MaskT mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (acc-ga-5) This function is identical to (acc-ga-2) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (acc-ga-6) This function is identical to (acc-ga-3) except that vector size is fixed to 1. More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Accessor-based scatter. More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorTy acc, simd< OffsetT, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-2) More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-sc-3) More...
template<typename T , int N, int VS = 1, typename AccessorTy , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorTy acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorTy, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (acc-sc-4) More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) &&detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > >	scatter (AccessorTy acc, simd< detail::DeviceAccessorOffsetT, N > offsets, simd< T, N > vals, detail::DeviceAccessorOffsetT glob_offset, simd_mask< N > mask=1)
	Writes elements of a simd object into an accessor at given offsets. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(detail::isPowerOf2(N, 32)) &&detail::is_device_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > >	scatter (AccessorTy acc, detail::DeviceAccessorOffsetT glob_offset, simd< T, N > vals, simd_mask< N > mask=1)
template<typename T , typename AccessorTy >
__ESIMD_API T	scalar_load (AccessorTy acc, detail::DeviceAccessorOffsetT offset)
	Load a scalar value from an accessor. More...
template<typename T , typename AccessorTy >
__ESIMD_API void	scalar_store (AccessorTy acc, detail::DeviceAccessorOffsetT offset, T val)
	Store a scalar value into an accessor. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API simd< T, N *get_num_channels_enabled(RGBAMask)>	gather_rgba (const T *p, simd< Toffset, N > offsets, simd_mask< N > mask=1)
	Gather and transpose pixels from given memory locations defined by the base pointer p and offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	gather_rgba (const T *p, OffsetSimdViewT offsets, simd_mask< N > mask=1)
	A variation of gather_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	gather_rgba (const T *p, Toffset offset, simd_mask< N > mask=1)
	A variation of gather_rgba API with offsets represented as scalar. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API void	scatter_rgba (T *p, simd< Toffset, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	Transpose and scatter pixels to given memory locations defined by the base pointer p and offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename OffsetSimdViewT , typename RegionTy >
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > >	scatter_rgba (T *p, OffsetSimdViewT offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	A variation of scatter_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	scatter_rgba (T *p, Toffset offset, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	A variation of scatter_rgba API with offsets represented as scalar. More...
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 >	scatter_rgba (T *p, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<((N==8||N==16||N==32) &&sizeof(T)==4 &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >), simd< T, N *get_num_channels_enabled(RGBAMask)> >	gather_rgba (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > offsets, detail::DeviceAccessorOffsetT global_offset=0, simd_mask< N > mask=1)
	Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	scatter_rgba (AccessorT acc, simd< detail::DeviceAccessorOffsetT, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, detail::DeviceAccessorOffsetT global_offset=0, simd_mask< N > mask=1)
	Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
template<uint32_t SLMSize>
__ESIMD_API void	slm_init ()
	Declare per-work-group slm size. More...
__ESIMD_API void	slm_init (uint32_t size)
	Declare per-work-group slm size. More...
template<typename T , int N, int VS, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-1) simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-2) simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (slm-ga-3) More...
template<typename T , int N, int VS, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-2) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (slm-ga-3) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N > byte_offsets, simd_mask< N > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-4) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {});// (slm-ga-5) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (simd< uint32_t, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename PropertyListT = empty_properties_t> simd<T, N> slm_gather(simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (slm-ga-6) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (slm-ga-7) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-ga-8) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_gather (OffsetSimdViewT byte_offsets, PropertyListT props={})
	simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (slm-ga-9) Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. More...
template<typename T >
__ESIMD_API T	slm_scalar_load (uint32_t offset)
	Load a scalar value from the Shared Local Memory. More...
template<typename T , int N, int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_scatter (simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename PropertyListT = empty_properties_t> void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-sc-1) void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-2) More...
template<typename T , int N, int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_scatter (simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename PropertyListT = empty_properties_t> void slm_scatter(simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-2) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_scatter (OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> void slm_scatter( OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (slm-sc-3) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_scatter (OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	void slm_scatter( OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (slm-sc-4) Stores ("scatters") elements of the type 'T' to Shared Local Memory locations addressed by byte offsets byte_offsets. More...
template<typename T >
__ESIMD_API void	slm_scalar_store (uint32_t offset, T val)
	Store a scalar value into the Shared Local Memory. More...
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4), simd< T, N *get_num_channels_enabled(RGBAMask)> >	slm_gather_rgba (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
	Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
template<typename T , int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4)>	slm_scatter_rgba (simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(Mask)> vals, simd_mask< N > mask=1)
	Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
template<typename T , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags >, simd< T, N > >	slm_block_load (uint32_t byte_offset, Flags)
	Loads a contiguous block of SLM memory referenced by the given byte-offset offset, then returns the loaded data as a simd object. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_block_load (uint32_t byte_offset, PropertyListT props={})
	Each of the following slm_block_load functions loads a contiguous memory block from SLM (Shared Local Memory) and the byte_offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_block_load (uint32_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> slm_block_load(uint32_t byte_offset, simd_mask<N> pred, props = {}); // (slm-bl-2) Loads a contiguous memory block from SLM (Shared Local Memory) at the given byte_offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	slm_block_load (uint32_t offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
	simd<T, N> slm_block_load(uint32_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (slm-bl-3) Loads a contiguous memory block from SLM (Shared Local Memory) at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, uint32_t byte_offset, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, props={}); // (lacc-bl-1) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, props={}); // (lacc-bl-2) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, uint32_t byte_offset, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, simd_mask<1> pred, props={}); // (lacc-bl-3) Loads a contiguous memory block from SLM (Shared Local Memory) associated the local accessor lacc at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, simd_mask< 1 > pred, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, simd_mask<1> pred, props={}); // (lacc-bl-4) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, uint32_t byte_offset, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, uint32_t byte_offset, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (lacc-bl-5) Loads a contiguous memory block from SLM (Shared Local Memory) associated the local accessor lacc at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	block_load (AccessorT lacc, simd_mask< 1 > pred, simd< T, N > pass_thru, PropertyListT props={})
	simd<T, N> block_load(local_accessor lacc, simd_mask<1> pred, simd<T, N> pass_thru, props={}); // (lacc-bl-6) Loads a contiguous memory block from SLM (Shared Local Memory) associated with the local accessor lacc at zero offset. More...
template<typename T , int N, typename Flags >
__ESIMD_API std::enable_if_t< is_simd_flag_type_v< Flags > >	slm_block_store (uint32_t offset, simd< T, N > vals, Flags)
	Stores elements of the vector vals to a contiguous block of SLM memory at the given byte-offset offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_block_store (uint32_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	Each of the following slm_block_store functions stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) at the byte_offset. More...
template<typename T , int N, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	slm_block_store (uint32_t byte_offset, simd< T, N > vals, PropertyListT props={})
	void slm_block_store(uint32_t byte_offset, simd<T, N> vals, // (slm-bs-2) props = {}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (AccessorT lacc, uint32_t byte_offset, simd< T, N > vals, PropertyListT props={})
	void block_store(local_accessor lacc, uint32_t byte_offset, // (lacc-bs-1) simd<T, N> vals, props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc at the given byte_offset. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (AccessorT lacc, simd< T, N > vals, PropertyListT props={})
	void block_store(local_accessor lacc, simd<T, N> vals, // (lacc-bs-2) props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (AccessorT lacc, uint32_t byte_offset, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	void block_store(local_accessor lacc, uint32_t byte_offset, // (lacc-bs-3) simd<T, N> vals, simd_mask<1> pred, props={}); More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	block_store (AccessorT lacc, simd< T, N > vals, simd_mask< 1 > pred, PropertyListT props={})
	void block_store(local_accessor lacc, simd<T, N> vals, // (lacc-bs-4) simd_mask<1> pred, props={}); Stores the vector vals to a contiguous memory block in SLM (Shared Local Memory) associated with the local accessor lacc. More...
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > >	slm_atomic_update (simd< uint32_t, N > byte_offset, simd_mask< N > mask=1)
	More...
template<atomic_op Op, typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_local_accessor_v< AccessorT >, simd< T, N > >	atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd_mask< N > mask=1)
	simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd_mask<N> pred = 1); // (lacc-au0-1) Atomically updates N memory locations in SLM ssociated with the local accessor lacc at the given byte_offset, and returns a vector of old values found at the memory locations before update. More...
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > >	slm_atomic_update (simd< uint32_t, N > byte_offset, simd< T, N > src0, simd_mask< N > mask=1)
	One argument variant of the atomic update operation. More...
template<atomic_op Op, typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_local_accessor_v< AccessorT >, simd< T, N > >	atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd< T, N > src0, simd_mask< N > mask=1)
	simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0, simd_mask<1> pred = 1); // (lacc-au1-1) More...
template<atomic_op Op, typename T , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > >	slm_atomic_update (simd< uint32_t, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
	Two argument variant of the atomic update operation. More...
template<atomic_op Op, typename T , int N, typename AccessorT >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_local_accessor_v< AccessorT >, simd< T, N > >	atomic_update (AccessorT lacc, simd< uint32_t, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
	simd<T, N> atomic_update(local_accessor lacc, simd<uint32_t, N> byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<1> pred = 1); // (lacc-au2-1) More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, simd_mask< N > mask, PropertyListT props={})
	More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, PropertyListT props={})
	simd<T, N> atomic_update(T *p, simd<Toffset, N> byte_offset, props = {}); /// (usm-au0-2) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT offsets, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, simd_mask<N> mask, props = {}); /// (usm-au0-3) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT byte_offset, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, props = {}); /// (usm-au0-4) More...
template<atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< T, N > >	atomic_update (T *p, Toffset byte_offset, simd_mask< N > mask=1)
	A variation of atomic_update API with offset represented as scalar. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
	More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, PropertyListT props={})
	simd<T, N> atomic_update(T *ptr, simd<Toffset, N> byte_offset, simd<T, N> src0, props = {}); // (usm-au1-2) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT offsets, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0, simd_mask<N> mask, props = {}); // (usm-au1-3) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT offsets, simd< T, N > src0, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0, props = {}); // (usm-au1-4) More...
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&((Op !=atomic_op::store &&Op !=atomic_op::xchg)||N==1), simd< Tx, N > >	atomic_update (Tx *p, Toffset byte_offset, simd< Tx, N > src0, simd_mask< N > mask)
	A variation of atomic_update API with offset represented as scalar object. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
	Atomically updates N memory locations represented by a USM pointer and a vector of offsets relative to the pointer, and returns a vector of old values found at the memory locations before update. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
	simd<T, N> atomic_update(T *p, simd<Toffset, N> byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (usm-au2-2) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<N> mask, props = {}) // (usm-au2-3) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (T *p, OffsetSimdViewT byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
	simd<T, N> atomic_update(T *p, OffsetSimdViewT byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}) // (usm-au2-4) More...
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< Tx, N > >	atomic_update (Tx *p, Toffset byte_offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	A variation of atomic_update API with byte_offset represented as scalar. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd_mask< N > mask, PropertyListT props={})
	More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, simd<Toffset, N> byte_offset, props = {}); /// (acc-au0-2) A variation of atomic_update API without mask operand More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, OffsetSimdViewT byte_offset, simd_mask<N> mask, props = {}); /// (acc-au0-3) A variation of atomic_update API with offsets represented as simd_view object. More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, OffsetSimdViewT byte_offset, props = {}); /// (acc-au0-4) A variation of atomic_update API with offsets represented as simd_view object and no mask operand. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, Toffset byte_offset, simd_mask< N > mask)
	A variation of atomic_update API with offset represented as scalar. More...
template<atomic_op Op, typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, uint32_t byte_offset, simd_mask< N > mask)
	A variation of atomic_update API with byte_offset represented as scalar using local_accessor. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
	More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, simd<Toffset, N> byte_offset, simd<T, N> src0, props = {}); // (acc-au1-2) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, simd< T, N > src0, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, OffsetSimdViewT byte_offset, simd<T, N> src0, simd_mask<N> mask, props = {}); // (acc-au1-3) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, simd< T, N > src0, PropertyListT props={})
	simd<T, N> atomic_update(AccessorT acc, OffsetSimdViewT byte_offset, simd<T, N> src0, props = {}); // (acc-au1-4) More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&((Op !=atomic_op::store &&Op !=atomic_op::xchg)||N==1), simd< T, N > >	atomic_update (AccessorTy acc, Toffset offset, simd< T, N > src0, simd_mask< N > mask)
	A variation of atomic_update API with offset represented as scalar object. More...
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy > &&((Op !=atomic_op::store &&Op !=atomic_op::xchg)||N==1), simd< Tx, N > >	atomic_update (AccessorTy acc, uint32_t offset, simd< Tx, N > src0, simd_mask< N > mask)
	A variation of atomic_update API with offset represented as scalar object and uses local_accessor. More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&std::is_integral_v< Toffset > &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
	More...
template<atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
	simd<T, N> atomic_update(AccessorTy acc, simd<Toffset, N> byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (acc-au2-2) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask, PropertyListT props={})
	simd<T, N> atomic_update(AccessorTy acc, OffsetSimdViewT byte_offset, simd<T, N> src0, simd<T, N> src1, simd_mask<N> mask, props = {}); // (acc-au2-3) More...
template<atomic_op Op, typename T , int N, typename OffsetSimdViewT , typename AccessorTy , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > &&detail::is_simd_view_type_v< OffsetSimdViewT >, simd< T, N > >	atomic_update (AccessorTy acc, OffsetSimdViewT byte_offset, simd< T, N > src0, simd< T, N > src1, PropertyListT props={})
	simd<T, N> atomic_update(AccessorTy acc, OffsetSimdViewT, byte_offset, simd<T, N> src0, simd<T, N> src1, props = {}); // (acc-au2-4) More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, Toffset offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as scalar. More...
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, uint32_t offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as scalar and local_accessor is used. More...
template<uint8_t cntl>
__ESIMD_API void	fence ()
	esimd::fence sets the memory read/write order. More...
__ESIMD_API void	fence (fence_mask cntl)
template<memory_kind Kind = memory_kind::global, fence_flush_op FenceOp = fence_flush_op::none, fence_scope Scope = fence_scope::group>
__ESIMD_API void	fence ()
	Memory fence. More...
__ESIMD_API void	barrier ()
	Generic work-group barrier. More...
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API simd< T, m *N >	media_block_load (AccessorTy acc, unsigned x, unsigned y)
	Media block load. More...
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API void	media_block_store (AccessorTy acc, unsigned x, unsigned y, simd< T, m *N > vals)
	Media block store. More...
template<typename T , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read > &&is_simd_flag_type_v< Flags >, simd< T, N > >	block_load (AccessorTy acc, uint32_t byte_offset, Flags flags)
	Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte_offset, then returns the loaded data as a simd object. More...
template<typename T , int N, typename AccessorT , typename Flags >
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&is_simd_flag_type_v< Flags > >	block_store (AccessorT acc, uint32_t offset, simd< T, N > vals, Flags flags)
	Variant of block_store that uses local accessor as a parameter. More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	Variant of gather that uses local accessor as a parameter template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-1) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-2) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, int VS, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (lacc-ga-3) Supported platforms: DG2, PVC in most cases. More...
template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-4) This function is identical to (lacc-ga-1) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename MaskT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&std::is_same_v< MaskT, simd_mask< N >> &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N > byte_offsets, MaskT mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props // (lacc-ga-5) This function is identical to (lacc-ga-2) except that vector size is fixed to 1. More...
template<typename T , int N, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, simd< uint32_t, N > byte_offsets, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename PropertyListT = empty_properties_t> simd<T, N> gather(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (lacc-ga-6) This function is identical to (lacc-ga-3) except that vector size is fixed to 1. More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, simd< T, N > pass_thru, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t<(detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT >), simd< T, N > >	gather (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_read >, simd< T, N > >	gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset, simd_mask< N > mask=1)
	Variant of gather that uses local accessor as a parameter. More...
template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	Variant of scatter that uses local accessor as a parameter template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-1) More...
template<typename T , int N, int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorT acc, simd< uint32_t, N/VS > byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-2) More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-sc-3) More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	scatter (AccessorT acc, OffsetSimdViewT byte_offsets, simd< T, N > vals, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void scatter(AccessorT acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {}); // (lacc-sc-4) More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write > >	scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset, simd_mask< N > mask=1)
	Variant of scatter that uses local accessor as a parameter. More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-1) void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) More...
template<typename T , int N, int VS, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int VS, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (usm-pf-2) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (usm-pf-3) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd<OffsetT, N> byte_offsets, PropertyListT props = {}); // (usm-pf-4) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (usm-pf-5) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (usm-pf-6) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, OffsetT byte_offset, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (usm-pf-7) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, OffsetT byte_offset, PropertyListT props={})
	template <typename T, int VS = 1, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(const T *p, OffsetT byte_offset, PropertyListT props = {}); // (usm-pf-8) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, simd_mask<1> mask, PropertyListT props = {}); //(usm-pf-9) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (const T *p, PropertyListT props={})
	template <typename T, int VS = 1, typename PropertyListT = empty_properties_t> void prefetch(const T *p, PropertyListT props = {}); // (usm-pf-10) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-1) void prefetch(AccessorT acc, simd<OffsetT, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2) More...
template<typename T , int N, int VS, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, simd< OffsetT, N/VS > byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N / VS> byte_offsets, PropertyListT props = {}); // (acc-pf-2) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, simd< OffsetT, N > byte_offsets, simd_mask< N > mask, PropertyListT props={})
	template <typename T, int N, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, simd_mask<N> mask, PropertyListT props = {}); // (acc-pf-3) Supported platforms: DG2, PVC only. More...
template<typename T , int N, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, simd< OffsetT, N > byte_offsets, PropertyListT props={})
	template <typename T, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd<uint32_t, N> byte_offsets, PropertyListT props = {}); // (acc-pf-4) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask< N/VS > mask, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (acc-pf-5) Supported platforms: DG2, PVC only. More...
template<typename T , int N, int VS = 1, typename OffsetSimdViewT , typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&detail::is_simd_view_type_v< OffsetSimdViewT > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props={})
	template <typename T, int N, int VS = 1, typename OffsetSimdViewT, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (acc-pf-6) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, OffsetT byte_offset, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, simd_mask<1> mask, PropertyListT props = {}); // (acc-pf-7) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename OffsetT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< std::is_integral_v< OffsetT > &&detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, OffsetT byte_offset, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename OffsetT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, uint32_t byte_offset, PropertyListT props = {}); // (acc-pf-8) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, simd_mask< 1 > mask, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, simd_mask<1> mask, PropertyListT props = {}); //(acc-pf-9) Supported platforms: DG2, PVC only. More...
template<typename T , int VS = 1, typename AccessorT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch (AccessorT acc, PropertyListT props={})
	template <typename T, int VS = 1, typename AccessorT, typename PropertyListT = empty_properties_t> void prefetch(AccessorT acc, PropertyListT props = {}); // (acc-pf-10) Supported platforms: DG2, PVC only. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT >, simd< T, N > >	load_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props={})
	template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>(), typename PropertyListT = empty_properties_t> simd<T, N> load_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block load. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	prefetch_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props={})
	template <typename T, int BlockWidth, int BlockHeight = 1, int NBlocks = 1, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>(), typename PropertyListT = empty_properties_t> void prefetch_2d(const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, PropertyListT props = {}); 2D USM pointer block prefetch. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false , false >(), typename PropertyListT = oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< ext::oneapi::experimental::is_property_list_v< PropertyListT > >	store_2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, simd< T, N > Vals, PropertyListT props={})
	2D USM pointer block store. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	gather_rgba (AccessorT acc, simd< uint32_t, N > offsets, uint32_t global_offset=0, simd_mask< N > mask=1)
	Variant of gather_rgba that uses local accessor as a parameter. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > >	scatter_rgba (AccessorT acc, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, uint32_t global_offset=0, simd_mask< N > mask=1)
	Variant of scatter_rgba that uses local accessor as a parameter Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
template<uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_src1, uint8_t num_dst, raw_send_eot eot = raw_send_eot::not_eot, raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1 , int n1, typename T2 , int n2, typename T3 , int n3>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 >	raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msg_dst, sycl::ext::intel::esimd::simd< T2, n2 > msg_src0, sycl::ext::intel::esimd::simd< T3, n3 > msg_src1, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
	Raw sends. More...
template<uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_dst, raw_send_eot eot = raw_send_eot::not_eot, raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1 , int n1, typename T2 , int n2>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 >	raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msg_dst, sycl::ext::intel::esimd::simd< T2, n2 > msg_src0, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
	Raw send. More...
template<uint8_t exec_size, uint8_t sfid, uint8_t num_src0, uint8_t num_src1, raw_send_eot eot = raw_send_eot::not_eot, raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1 , int n1, typename T2 , int n2>
__ESIMD_API void	raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msg_src0, sycl::ext::intel::esimd::simd< T2, n2 > msg_src1, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
	Raw sends. More...
template<uint8_t exec_size, uint8_t sfid, uint8_t num_src0, raw_send_eot eot = raw_send_eot::not_eot, raw_send_sendc sendc = raw_send_sendc::not_sendc, typename T1 , int n1>
__ESIMD_API void	raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msg_src0, uint32_t ex_desc, uint32_t msg_desc, sycl::ext::intel::esimd::simd_mask< exec_size > mask=1)
	Raw send. More...
template<typename To , typename From , int N>
ESIMD_INLINE simd< To, N >	convert (const simd< From, N > &val)
	Covert from a simd object with element type From to a simd object with element type To. More...
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > offset, simd_mask< N > mask)
	LSC version of no argument variant of the atomic_update - accepts native::lsc::atomic_op instead of atomic_op as atomic operation template argument. More...
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > >	atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > >	atomic_update (T *p, Toffset offset, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > offset, simd< T, N > src0, simd_mask< N > mask)
	LSC version of the single-argument atomic update. More...
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > >	atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > >	atomic_update (T *p, Toffset offset, simd< T, N > src0, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > >	atomic_update (T *p, simd< Toffset, N > offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask)
	LSC version of the two-argument atomic update. More...
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > >	atomic_update (T *p, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, 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 > >	atomic_update (T *p, Toffset offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > offsets, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, Toffset offset, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< T, N > src0, simd_mask< N > mask)
	LSC version of the single-argument atomic update. More...
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, Toffset offset, simd< T, N > src0, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask)
	LSC version of the two-argument atomic update. More...
template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< OffsetObjT, RegionTy > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask)
template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	atomic_update (AccessorTy acc, Toffset offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask)

Variables
static constexpr saturation_off_tag	saturation_off {}
	Type tag object representing "saturation off" behavior. More...
static constexpr saturation_on_tag	saturation_on {}
	Type tag object representing "saturation on" behavior. More...
constexpr element_aligned_tag	element_aligned = {}
constexpr vector_aligned_tag	vector_aligned = {}
template<unsigned N>
constexpr overaligned_tag< N >	overaligned = {}
template<typename T >
static constexpr bool	is_simd_flag_type_v = is_simd_flag_type<T>::value
	Checks if given type is a simd load/store flag. More...
__ESIMD_API SZ	src0
__ESIMD_API SZ simd< T, SZ >	src1
__ESIMD_API SZ simd< T, SZ > Sat	sat
__ESIMD_API SZ simd< T, SZ > Sat int	SZ
__ESIMD_API SZ simd< T, SZ > Sat int class	Sat
template<int K>
constexpr alignment_key::value_t< K >	alignment
template<cache_hint Hint>
constexpr cache_hint_L1_key::value_t< Hint >	cache_hint_L1
template<cache_hint Hint>
constexpr cache_hint_L2_key::value_t< Hint >	cache_hint_L2
template<cache_hint Hint>
constexpr cache_hint_L3_key::value_t< Hint >	cache_hint_L3

Typedef Documentation

alignment_key

using sycl::_V1::ext::intel::esimd::alignment_key = typedef sycl::ext::oneapi::experimental::alignment_key

The 'alignment' property is used to specify the alignment of memory accessed in ESIMD memory operation such as block_load().

Case1: ESIMD memory operation accepts only USM pointer parameter. The 'alignment' property specifies the alignment of the memory accessed by the USM pointer. Case2: ESIMD memory operation accepts two parameters: either (accessor + offset) or (USM pointer + offset). The 'alignment' property specifies the alignment of the memory accesed by those two parameters.

Definition at line 95 of file memory_properties.hpp.

default_cache_hint_L1

using sycl::_V1::ext::intel::esimd::default_cache_hint_L1 = typedef cache_hint_L1_key::value_t<cache_hint::none>

Definition at line 152 of file memory_properties.hpp.

default_cache_hint_L2

using sycl::_V1::ext::intel::esimd::default_cache_hint_L2 = typedef cache_hint_L2_key::value_t<cache_hint::none>

Definition at line 153 of file memory_properties.hpp.

default_cache_hint_L3

using sycl::_V1::ext::intel::esimd::default_cache_hint_L3 = typedef cache_hint_L3_key::value_t<cache_hint::none>

Definition at line 154 of file memory_properties.hpp.

mask_type_t

template<int N>

using sycl::_V1::ext::intel::esimd::mask_type_t = typedef detail::simd_mask_storage_t<N>

Definition at line 391 of file types.hpp.

Enumeration Type Documentation

cache_hint

enum sycl::_V1::ext::intel::esimd::cache_hint : uint8_t

strong

L1, L2 or L3 cache hints.

Enumerator
none
uncached	load/store/atomic: do not cache data to cache;
cached
write_back	store: write data into cache level and mark the cache line as "dirty". Upon eviction, the "dirty" data will be written into the furthest subsequent cache;
write_through	store: immediately write data to the subsequent furthest cache, marking the cache line in the current cache as "not dirty";
streaming	load: cache data to cache using the evict-first policy to minimize cache pollution caused by temporary streaming data that may only be accessed once or twice; store/atomic: same as write-through, but use the evict-first policy to limit cache pollution by streaming;
read_invalidate	load: asserts that the cache line containing the data will not be read again until it’s overwritten, therefore the load operation can invalidate the cache line and discard "dirty" data. If the assertion is violated (the cache line is read again) then behavior is undefined.
const_cached	load, L2 cache only, next gen GPU after Xe required: asserts that the L2 cache line containing the data will not be written until all invocations of the shader or kernel execution are finished. If the assertion is violated (the cache line is written), the behavior is undefined.

Definition at line 27 of file memory_properties.hpp.

cache_level

enum sycl::_V1::ext::intel::esimd::cache_level : uint8_t

strong

L1, L2 or L3 cache hint levels. L3 is reserved for future use.

Enumerator
L1
L2
L3

Definition at line 24 of file memory_properties.hpp.

fence_flush_op

enum sycl::_V1::ext::intel::esimd::fence_flush_op : uint8_t

strong

The cache flush operation to apply to caches after fence() is complete.

Supported platforms: DG2, PVC

Enumerator
none
evict	no operation;
invalidate	R/W: evict dirty lines; R/W and RO: invalidate clean lines.
clean	R/W and RO: invalidate all clean lines;.

Definition at line 384 of file common.hpp.

fence_scope

enum sycl::_V1::ext::intel::esimd::fence_scope : uint8_t

strong

The scope that fence() operation should apply to.

Supported platforms: DG2, PVC

Enumerator
group	Wait until all previous memory transactions from this thread are observed within the local thread-group.
local	Wait until all previous memory transactions from this thread are observed within the local sub-slice.
tile	Wait until all previous memory transactions from this thread are observed in the local tile.
gpu	Wait until all previous memory transactions from this thread are observed in the local GPU.
gpus	Wait until all previous memory transactions from this thread are observed across all GPUs in the system.
system	Global memory data-port only: wait until all previous memory transactions from this thread are observed at the "system" level.
system_acquire	Global memory data-port only: for GPUs that do not follow PCIe Write ordering for downstream writes targeting device memory, this op will commit to device memory all downstream and peer writes that have reached the device.

Definition at line 350 of file common.hpp.

memory_kind

enum sycl::_V1::ext::intel::esimd::memory_kind : uint8_t

strong

The target memory kind for fence() operation.

Supported platforms: DG2, PVC

Enumerator
global
image	untyped global memory
local	image (also known as typed global memory)

Definition at line 397 of file common.hpp.

Function Documentation

atomic_update() [1/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 2 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd< Toffset, N >	offset,
		simd< T, N >	src0,
		simd< T, N >	src1,
		simd_mask< N >	mask
	)

LSC version of the two-argument atomic update.

Definition at line 2880 of file memory.hpp.

References src0, and src1.

atomic_update() [2/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd< Toffset, N >	offset,
		simd< T, N >	src0,
		simd_mask< N >	mask
	)

LSC version of the single-argument atomic update.

Definition at line 2844 of file memory.hpp.

References src0.

atomic_update() [3/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 0 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd< Toffset, N >	offset,
		simd_mask< N >	mask
	)

Definition at line 2810 of file memory.hpp.

atomic_update() [4/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 2 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd_view< OffsetObjT, RegionTy >	offsets,
		simd< T, N >	src0,
		simd< T, N >	src1,
		simd_mask< N >	mask
	)

Definition at line 2894 of file memory.hpp.

References src0, and src1.

atomic_update() [5/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 1 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd_view< OffsetObjT, RegionTy >	offsets,
		simd< T, N >	src0,
		simd_mask< N >	mask
	)

Definition at line 2855 of file memory.hpp.

References src0.

atomic_update() [6/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy , typename AccessorTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 0 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd_view< OffsetObjT, RegionTy >	offsets,
		simd_mask< N >	mask
	)

Definition at line 2820 of file memory.hpp.

atomic_update() [7/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 2 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		Toffset	offset,
		simd< T, N >	src0,
		simd< T, N >	src1,
		simd_mask< N >	mask
	)

Definition at line 2906 of file memory.hpp.

References src0, and src1.

atomic_update() [8/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		Toffset	offset,
		simd< T, N >	src0,
		simd_mask< N >	mask
	)

Definition at line 2867 of file memory.hpp.

References src0.

atomic_update() [9/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 0 && __ESIMD_DNS::is_rw_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		Toffset	offset,
		simd_mask< N >	mask
	)

Definition at line 2832 of file memory.hpp.

atomic_update() [10/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >

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

LSC version of the two-argument atomic update.

Definition at line 2776 of file memory.hpp.

References src0, and src1.

atomic_update() [11/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	T *	p,
		simd< Toffset, N >	offset,
		simd< T, N >	src0,
		simd_mask< N >	mask
	)

LSC version of the single-argument atomic update.

Definition at line 2746 of file memory.hpp.

References src0.

atomic_update() [12/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >

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

LSC version of no argument variant of the atomic_update - accepts native::lsc::atomic_op instead of atomic_op as atomic operation template argument.

Definition at line 2718 of file memory.hpp.

atomic_update() [13/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >

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

Definition at line 2788 of file memory.hpp.

References src0, and src1.

atomic_update() [14/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >

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

Definition at line 2756 of file memory.hpp.

References src0.

atomic_update() [15/18]

template<native::lsc::atomic_op Op, typename T , int N, typename OffsetObjT , typename RegionTy >

__ESIMD_API std::enable_if_t<__ESIMD_DNS::get_num_args<Op>) == 0, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	T *	p,
		simd_view< OffsetObjT, RegionTy >	offsets,
		simd_mask< N >	mask = 1
	)

Definition at line 2726 of file memory.hpp.

atomic_update() [16/18]

template<native::lsc::atomic_op Op, typename T , int N, 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::esimd::atomic_update	(	T *	p,
		Toffset	offset,
		simd< T, N >	src0,
		simd< T, N >	src1,
		simd_mask< N >	mask = 1
	)

Definition at line 2798 of file memory.hpp.

References src0, and src1.

atomic_update() [17/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >

__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	T *	p,
		Toffset	offset,
		simd< T, N >	src0,
		simd_mask< N >	mask = 1
	)

Definition at line 2766 of file memory.hpp.

References src0.

atomic_update() [18/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >

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

Definition at line 2736 of file memory.hpp.

fbh()

template<typename T , int N>

__ESIMD_API std::enable_if_t<std::is_integral<T>::value && !std::is_signed<T>::value && (sizeof(T) == 4), simd<T, N> > sycl::_V1::ext::intel::esimd::fbh

(

simd< T, N >

src

)

Find the per element number of the first bit set in the source operand starting from the most significant bit (sign bit is counted).

Parameters

src0	the source operand to count bits in.

Returns

a vector of the same type as the source operand, where each element is set to the number first bit set in corresponding element of the source operand. 0xFFFFffff is returned for an element equal to 0.

Definition at line 804 of file math.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data().

gather() [1/3]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, PropertyListT props = {}); // (lacc-ga-9) This function is identical to (lacc-ga-3) except that the byte_offsets is represented as simd_view.

Definition at line 8407 of file memory.hpp.

gather() [2/3]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename AccessorT, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PropertyListT props = {}); // (lacc-ga-8) This function is identical to (lacc-ga-2) except that the byte_offsets is represented as simd_view.

Definition at line 8385 of file memory.hpp.

gather() [3/3]

template<typename T , int N, int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::detail::empty_properties_t>

__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorT	acc,
		OffsetSimdViewT	byte_offsets,
		simd_mask< N/VS >	mask,
		simd< T, N >	pass_thru,
		PropertyListT	props = {}
	)

template <typename T, int N, int VS = 1, typename OffsetSimdViewT,

simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); // (lacc-ga-7) This function is identical to (lacc-ga-1) except that the byte_offsets is represented as simd_view.

Definition at line 8362 of file memory.hpp.

gather_rgba()

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read>, simd<T, N * get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba	(	AccessorT	acc,
		simd< uint32_t, N >	offsets,
		uint32_t	global_offset = 0,
		simd_mask< N >	mask = 1
	)

Variant of gather_rgba that uses local accessor as a parameter.

Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. Up to 4 32-bit data elements may be accessed at each address depending on the channel mask RGBAMask. Each pixel's address must be 4-byte aligned. For usage examples, see usm_gather_rgba above, the only difference would be the usage of an accessor instead of a usm pointer.

Template Parameters

RGBAMask	A pixel's channel mask.
AccessorT	The accessor type for the memory to be loaded/gathered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
N	Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.

Parameters

acc	The accessor representing memory address of the access.
offsets	Byte offsets of the pixels relative to the base pointer.
global_offset	Byte offset of the pixels relative to the base pointer.
mask	Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.

Returns

Read data - up to N*4 values of type Tx.

Definition at line 9492 of file memory.hpp.

pack_mask()

template<int N>

__ESIMD_API std::enable_if_t<(N != 8 && N != 16 && N < 32), uint> sycl::_V1::ext::intel::esimd::pack_mask

(

simd_mask< N >

src0

)

pack_mask specialization when the number of elements N is not 8, 16 or 32.

Definition at line 667 of file math.hpp.

References pack_mask(), and src0.

scatter()

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorTy, detail::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		simd< uint32_t, N >	offsets,
		simd< T, N >	vals,
		uint32_t	glob_offset,
		simd_mask< N >	mask = 1
	)

Variant of scatter that uses local accessor as a parameter.

Writes elements of a simd object into an accessor at given offsets. An element can be a 1, 2 or 4-byte value.

Template Parameters

T	Element type; can only be a 1,2,4-byte integer, sycl::half or float.
N	The number of vector elements. Can be 1, 8, 16 or 32.
AccessorTy	The accessor type.

Parameters

acc	The accessor to scatter to.
offsets	Per-element offsets in bytes.
vals	Values to write.
glob_offset	Offset in bytes added to each individual element's offset to compute actual memory access offset for that element.
mask	Memory access mask. Elements with zero corresponding mask's predicate are not accessed.

Definition at line 8650 of file memory.hpp.

scatter_rgba()

template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>

__ESIMD_API std::enable_if_t<detail::is_local_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write> > sycl::_V1::ext::intel::esimd::scatter_rgba	(	AccessorT	acc,
		simd< uint32_t, N >	offsets,
		simd< T, N *get_num_channels_enabled(RGBAMask)>	vals,
		uint32_t	global_offset = 0,
		simd_mask< N >	mask = 1
	)

Variant of scatter_rgba that uses local accessor as a parameter Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector.

See usm_gather_rgba for information about the operation semantics and parameter restrictions/interdependencies.

Template Parameters

RGBAMask	Pixel's channel mask.
AccessorT	The accessor type for the memory to be stored/scattered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
N	The number of elements to access.

Parameters

offsets	Byte offsets of each element.
vals	values to be written.
global_offset	Byte offset of the pixels relative to the base pointer.
mask	Operation mask. All-1 by default.

Definition at line 9519 of file memory.hpp.

Variable Documentation

alignment

template<int K>

constexpr alignment_key::value_t<K> sycl::_V1::ext::intel::esimd::alignment

inlineconstexpr

Definition at line 97 of file memory_properties.hpp.

Referenced by block_load(), and sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load().

cache_hint_L1

template<cache_hint Hint>

constexpr cache_hint_L1_key::value_t<Hint> sycl::_V1::ext::intel::esimd::cache_hint_L1

inlineconstexpr

Definition at line 129 of file memory_properties.hpp.

cache_hint_L2

template<cache_hint Hint>

constexpr cache_hint_L2_key::value_t<Hint> sycl::_V1::ext::intel::esimd::cache_hint_L2

inlineconstexpr

Definition at line 131 of file memory_properties.hpp.

cache_hint_L3

template<cache_hint Hint>

constexpr cache_hint_L3_key::value_t<Hint> sycl::_V1::ext::intel::esimd::cache_hint_L3

inlineconstexpr

Definition at line 133 of file memory_properties.hpp.