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

Namespaces
	detail
	native
	xmx

Classes
struct	element_aligned_tag
	element_aligned_tag type. More...
struct	is_simd_flag_type
	Checks if type is a simd load/store flag. More...
struct	is_simd_flag_type< detail::dqword_element_aligned_tag >
struct	is_simd_flag_type< element_aligned_tag >
struct	is_simd_flag_type< overaligned_tag< N > >
struct	is_simd_flag_type< vector_aligned_tag >
struct	overaligned_tag
	overaligned_tag type. More...
struct	saturation_off_tag
	This type tag represents "saturation off" behavior. More...
struct	saturation_on_tag
	Gen hardware supports applying saturation to results of certain operations. More...
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	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...
struct	vector_aligned_tag
	vector_aligned_tag type. 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 >
template<int N>
using	simd_mask = detail::simd_mask_type< N >
	Represents a simd mask os size N. More...

Enumerations
enum	rgba_channel : uint8_t { rgba_channel::R, rgba_channel::G, rgba_channel::B, rgba_channel::A }
	Represents a pixel's channel. More...
enum	rgba_channel_mask : uint8_t {   rgba_channel_mask::R = detail::chR, rgba_channel_mask::G = detail::chG, rgba_channel_mask::GR = detail::chG | detail::chR, rgba_channel_mask::B = detail::chB,   rgba_channel_mask::BR = detail::chB | detail::chR, rgba_channel_mask::BG = detail::chB | detail::chG, rgba_channel_mask::BGR = detail::chB | detail::chG | detail::chR, rgba_channel_mask::A = detail::chA,   rgba_channel_mask::AR = detail::chA | detail::chR, rgba_channel_mask::AG = detail::chA | detail::chG, rgba_channel_mask::AGR = detail::chA | detail::chG | detail::chR, rgba_channel_mask::AB = detail::chA | detail::chB,   rgba_channel_mask::ABR = detail::chA | detail::chB | detail::chR, rgba_channel_mask::ABG = detail::chA | detail::chB | detail::chG, rgba_channel_mask::ABGR = detail::chA | detail::chB | detail::chG | detail::chR }
	Represents a pixel's channel mask - all possible combinations of enabled channels. More...
enum	atomic_op : uint8_t {   atomic_op::add = 0x0, atomic_op::sub = 0x1, atomic_op::inc = 0x2, atomic_op::dec = 0x3,   atomic_op::umin = 0x4, atomic_op::umax = 0x5, atomic_op::xchg = 0x6, atomic_op::cmpxchg = 0x7,   atomic_op::bit_and = 0x8, atomic_op::bit_or = 0x9, atomic_op::bit_xor = 0xa, atomic_op::smin = 0xb,   atomic_op::smax = 0xc, atomic_op::__SYCL_DEPRECATED =("fmax" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x10, atomic_op::__SYCL_DEPRECATED =("fmin" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x11, atomic_op::fcmpxchg = 0x12,   atomic_op::__SYCL_DEPRECATED =("fcmpwr" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = fcmpxchg, atomic_op::__SYCL_DEPRECATED =("fadd" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x13, atomic_op::__SYCL_DEPRECATED =("fsub" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x14, atomic_op::load = 0x15,   atomic_op::store = 0x16, atomic_op::predec = 0xff }
	Represents an atomic operation. More...
enum	bfn_t : uint8_t { bfn_t::x = 0xAA, bfn_t::y = 0xCC, bfn_t::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, l3_flush_instructions = 0x2, l3_flush_texture_data = 0x4, l3_flush_constant_data = 0x8,   l3_flush_rw_data = 0x10, local_barrier = 0x20, l1_flush_ro_data = 0x40, sw_barrier = 0x80 }
	Represetns a bit mask to control behavior of esimd::fence. 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, SZ > src0, simd< T, SZ > src1, Sat sat={})
	Selects component-wise the maximum of the two vectors. More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > >	max (simd< T, SZ > src0, T src1, Sat sat={})
	Selects maximums for each element of the input vector and a scalar. More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > >	max (T src0, simd< T, SZ > src1, Sat sat={})
	Selects maximums for each element of the input scalar and a vector. More...
template<typename T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< detail::is_esimd_scalar< T >::value, T >	max (T src0, T src1, Sat sat={})
	Selects maximum between two scalar values. More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API simd< T, SZ >	min (simd< T, SZ > src0, simd< T, SZ > src1, Sat sat={})
	Selects component-wise the minimum of the two vectors. More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > >	min (simd< T, SZ > src0, T src1, Sat sat={})
	Selects minimums for each element of the input vector and a scalar. More...
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > >	min (T src0, simd< T, SZ > src1, Sat sat={})
	Selects minimums for each element of the input scalar and a vector. More...
template<typename T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< detail::is_esimd_scalar< T >::value, T >	min (T src0, T src1, Sat sat={})
	Selects minimum between two scalar values. 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 >	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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_IEEE_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_IEEE_COND >>
__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< __ESIMD_EMATH_IEEE_COND >>
__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 __ESIMD_NS::simd< T, SZ >	rndd (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
	Round-down (also known as floor). More...
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ >	rndu (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
	Round-up (also known as ceil). More...
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ >	rnde (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
	Round-to-even (also known as round). More...
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ >	rndz (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
	Round-to-zero (also known as trunc). More...
template<typename T , class Sat = __ESIMD_NS::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...
template<typename AccessorTy >
__ESIMD_API SurfaceIndex	get_surface_index (AccessorTy acc)
	Get surface index corresponding to a SYCL accessor. More...
template<typename Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N >	gather (const Tx *p, simd< Toffset, N > offsets, simd_mask< N > mask=1)
	Loads ("gathers") elements from different memory locations and returns a vector of them. More...
template<typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N >	gather (const Tx *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1)
	A variation of gather API with offsets represented as simd_view object. 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 Tx , int N, typename Toffset >
__ESIMD_API void	scatter (Tx *p, simd< Toffset, N > offsets, simd< Tx, N > vals, simd_mask< N > mask=1)
	Writes ("scatters") elements of the input vector to different memory locations. More...
template<typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API void	scatter (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > vals, simd_mask< N > mask=1)
	A variation of scatter API with offsets represented as simd_view object. 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 >, 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> && sycl::detail::acc_properties::is_accessor_v<AccessorTy> && !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N >	block_load (AccessorTy acc, uint32_t offset, Flags={})
	Loads a contiguous block of memory from given accessor and offset and returns the loaded data as a vector. More...
template<typename Tx , int N, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API __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 Tx , int N, typename AccessorTy , class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > >	block_store (AccessorTy acc, uint32_t offset, simd< Tx, N > vals)
	Stores elements of a vector to a contiguous block of memory represented by an accessor and an offset within this accessor. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< T, N > >	gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > >	scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
template<typename T , typename AccessorTy >
__ESIMD_API T	scalar_load (AccessorTy acc, uint32_t offset)
	Load a scalar value from an accessor. More...
template<typename T , typename AccessorTy >
__ESIMD_API void	scalar_store (AccessorTy acc, uint32_t 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)
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< T, N *get_num_channels_enabled(RGBAMask)>	gather_rgba (const T *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1)
	A variation of gather_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, simd< T, N *get_num_channels_enabled(RGBAMask)> >	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)
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API void	scatter_rgba (T *p, simd_view< Toffset, RegionTy > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
	A variation of scatter_rgba API with offsets represented as simd_view object. More...
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	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 &&!std::is_pointer_v< AccessorT > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorT >), 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)
	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 &&!std::is_pointer_v< AccessorT > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorT > >	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)
	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>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32), simd< T, N > >	slm_gather (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
	Gather operation over the Shared Local Memory. 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>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32) &&(sizeof(T)<=4)>	slm_scatter (simd< uint32_t, N > offsets, simd< T, N > vals, simd_mask< N > mask=1)
	Scatter operation over the Shared Local Memory. 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 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 Flags = overaligned_tag<detail::OperandSize::OWORD>>
__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<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N >	slm_atomic_update (simd< uint32_t, N > offsets, simd_mask< N > mask)
	Atomic update operation performed on SLM. More...
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N >	slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd_mask< N > mask)
	Atomic update operation performed on SLM. More...
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N >	slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	Atomic update operation performed on SLM. More...
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd< Toffset, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
	Single-argument variant of the atomic update operation. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as simd_view object. 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 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 Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd< Toffset, N > offset, simd_mask< N > mask)
	No-argument variant of the atomic update operation. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1)
	A variation of atomic_update API with offsets represented as simd_view object. 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 offset, simd_mask< N > mask=1)
	A variation of atomic_update API with offset represented as scalar. More...
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd< Toffset, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N >	atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as simd_view object. 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 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 Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
	Single-argument variant of the atomic update operation. More...
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
	Variant of atomic_update that uses local_accessor as a parameter. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as simd_view object. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value &&((Op !=atomic_op::store &&Op !=atomic_op::xchg)||N==1), simd< Tx, N > >	atomic_update (AccessorTy acc, Toffset 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 Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd_mask< N > mask)
	No-argument variant of the atomic update operation. More...
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd_mask< N > mask)
	Variant of atomic_update that uses local_accessor as a parameter. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as simd_view object. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > >	atomic_update (AccessorTy acc, Toffset offset, simd_mask< N > mask)
	A variation of atomic_update API with offset represented as scalar. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > >	atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	Variant of atomic_update that uses local_accessor as a parameter. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
	A variation of atomic_update API with offsets represented as simd_view object. More...
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, 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<uint8_t cntl>
__ESIMD_API void	fence ()
	esimd::fence sets the memory read/write order. More...
__ESIMD_API void	fence (fence_mask cntl)
__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 Tx , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > &&is_simd_flag_type_v< Flags >, simd< Tx, N > >	block_load (AccessorTy acc, uint32_t offset, Flags={})
	Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte-offset offset, then returns the loaded data as a simd object. More...
template<typename Tx , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > &&is_simd_flag_type_v< Flags > >	block_store (AccessorTy acc, uint32_t offset, simd< Tx, N > vals, Flags={})
	Variant of block_store that uses local accessor as a parameter. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< T, N > >	gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
	Variant of gather that uses local accessor as a parameter. More...
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > >	scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
	Variant of scatter 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< sycl::detail::acc_properties::is_local_accessor_v< AccessorT >, 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< sycl::detail::acc_properties::is_local_accessor_v< AccessorT > >	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<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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, 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 __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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API __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_view< Toffset, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0 &&!std::is_pointer< AccessorTy >::value, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&!std::is_pointer< AccessorTy >::value, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2 &&!std::is_pointer< AccessorTy >::value, simd< T, N > >	atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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...

Typedef Documentation

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.

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 && !std::is_pointer<AccessorTy>::value, 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 3872 of file memory.hpp.

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

atomic_update() [2/18]

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

__ESIMD_API __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1 && !std::is_pointer<AccessorTy>::value, 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 3833 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 && !std::is_pointer<AccessorTy>::value, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		simd< Toffset, N >	offset,
		simd_mask< N >	mask
	)

Definition at line 3796 of file memory.hpp.

atomic_update() [4/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >

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

Definition at line 3888 of file memory.hpp.

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

atomic_update() [5/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >

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

Definition at line 3847 of file memory.hpp.

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

atomic_update() [6/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >

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

Definition at line 3808 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 && !std::is_pointer<AccessorTy>::value, 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 3900 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::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 && !std::is_pointer<AccessorTy>::value, 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 3859 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 && !std::is_pointer<AccessorTy>::value, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update	(	AccessorTy	acc,
		Toffset	offset,
		simd_mask< N >	mask
	)

Definition at line 3820 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 3760 of file memory.hpp.

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

atomic_update() [11/18]

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

__ESIMD_API __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 3728 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 3697 of file memory.hpp.

atomic_update() [13/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>

__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_view< Toffset, RegionTy >	offsets,
		simd< T, N >	src0,
		simd< T, N >	src1,
		simd_mask< N >	mask = 1
	)

Definition at line 3774 of file memory.hpp.

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

atomic_update() [14/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>

__ESIMD_API __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_view< Toffset, RegionTy >	offsets,
		simd< T, N >	src0,
		simd_mask< N >	mask = 1
	)

Definition at line 3740 of file memory.hpp.

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

atomic_update() [15/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>

__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_view< Toffset, RegionTy >	offsets,
		simd_mask< N >	mask = 1
	)

Definition at line 3707 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 3784 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::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 3750 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 3717 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.

gather()

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather	(	AccessorTy	acc,
		simd< uint32_t, N >	offsets,
		uint32_t	glob_offset = 0,
		simd_mask< N >	mask = 1
	)

Variant of gather that uses local accessor as a parameter.

Collects elements located at given offsets in an accessor and returns them as a single simd object. 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 gather from.
offsets	Per-element offsets in bytes.
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, their values in the resulting vector are undefined.

Definition at line 2217 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< sycl::detail::acc_properties::is_local_accessor_v<AccessorT>, 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 2281 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 668 of file math.hpp.

References pack_mask(), and sycl::_V1::ext::intel::experimental::esimd::src0.

scatter()

template<typename T , int N, typename AccessorTy >

__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> > sycl::_V1::ext::intel::esimd::scatter	(	AccessorTy	acc,
		simd< uint32_t, N >	offsets,
		simd< T, N >	vals,
		uint32_t	glob_offset = 0,
		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 2244 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< sycl::detail::acc_properties::is_local_accessor_v<AccessorT> > 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 2308 of file memory.hpp.