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

Namespaces
	detail

Classes
class	config_2d_mem_access
	Container class to hold parameters for load2d/store2d functions More...
class	tfloat32

Typedefs
using	lsc_data_size = ("use sycl::ext::intel::esimd::memory_kind") lsc_memory_kind __ESIMD_DNS::lsc_data_size
	The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC. More...
using	cache_hint = sycl::ext::intel::esimd::cache_hint
	L1 or L2 cache hint kinds. More...

Enumerations
enum class	split_barrier_action : uint8_t { wait = 0 , signal = 1 }
	Represents a split barrier action. More...

Functions
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	shl (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Shift left operation (vector version) More...
template<typename T0 , typename T1 , typename T2 , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > >	shl (T1 src0, T2 src1, Sat sat={})
	Shift left operation (scalar version) More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	lsr (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Logical Shift Right (vector version) More...
template<typename T0 , typename T1 , typename T2 , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > >	lsr (T1 src0, T2 src1, Sat sat={})
	Logical Shift Right (scalar version) More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	asr (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Arithmetical Shift Right (vector version) More...
template<typename T0 , typename T1 , typename T2 , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > >	asr (T1 src0, T2 src1, Sat sat={})
	Arithmetical Shift Right (scalar version) More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	shr (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Shift right operation (vector version) More...
template<typename T0 , typename T1 , typename T2 , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > >	shr (T1 src0, T2 src1, Sat sat={})
	Shift right operation (scalar version) More...
template<typename T0 , typename T1 , int SZ>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), sycl::ext::intel::esimd::simd< T0, SZ > >	rol (sycl::ext::intel::esimd::simd< T1, SZ > src0, sycl::ext::intel::esimd::simd< T1, SZ > src1)
	Rotate left operation with two vector inputs. More...
template<typename T0 , typename T1 , int SZ, typename U >
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< U, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), sycl::ext::intel::esimd::simd< T0, SZ > >	rol (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1)
	Rotate left operation with a vector and a scalar inputs. More...
template<typename T0 , typename T1 , typename T2 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T2, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), std::remove_const_t< T0 > >	rol (T1 src0, T2 src1)
	Rotate left operation with two scalar inputs. More...
template<typename T0 , typename T1 , int SZ>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), sycl::ext::intel::esimd::simd< T0, SZ > >	ror (sycl::ext::intel::esimd::simd< T1, SZ > src0, sycl::ext::intel::esimd::simd< T1, SZ > src1)
	Rotate right operation with two vector inputs. More...
template<typename T0 , typename T1 , int SZ, typename U >
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< U, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), sycl::ext::intel::esimd::simd< T0, SZ > >	ror (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1)
	Rotate right operation with a vector and a scalar inputs. More...
template<typename T0 , typename T1 , typename T2 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T2, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), std::remove_const_t< T0 > >	ror (T1 src0, T2 src1)
	Rotate right operation with two scalar inputs. More...
template<typename T , typename T0 , typename T1 , int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	imul_impl (sycl::ext::intel::esimd::simd< T, N > &rmd, sycl::ext::intel::esimd::simd< T0, N > src0, sycl::ext::intel::esimd::simd< T1, N > src1)
	Computes the 64-bit result of two 32-bit element vectors src0 and src1 multiplication. More...
template<typename T , typename T0 , typename T1 , int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	imul (sycl::ext::intel::esimd::simd< T, N > &rmd, sycl::ext::intel::esimd::simd< T0, N > src0, sycl::ext::intel::esimd::simd< T1, N > src1)
	Computes the 64-bit multiply result of two 32-bit integer vectors src0 and src1. More...
template<typename T , typename T0 , typename T1 , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_dword_type< T1 >::value, sycl::ext::intel::esimd::simd< T, N > >	imul (sycl::ext::intel::esimd::simd< T, N > &rmd, sycl::ext::intel::esimd::simd< T0, N > src0, T1 src1)
	Computes the 64-bit multiply result of 32-bit integer vector src0 and 32-bit integer scalar src1. More...
template<typename T , typename T0 , typename T1 , int N>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_dword_type< T0 >::value, sycl::ext::intel::esimd::simd< T, N > >	imul (sycl::ext::intel::esimd::simd< T, N > &rmd, T0 src0, sycl::ext::intel::esimd::simd< T1, N > src1)
	Computes the 64-bit multiply result of a scalar 32-bit integer src0 and 32-bit integer vector src1. More...
template<typename T , typename T0 , typename T1 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_dword_type< T >::value &&__ESIMD_DNS::is_dword_type< T0 >::value &&__ESIMD_DNS::is_dword_type< T1 >::value, T >	imul (T &rmd, T0 src0, T1 src1)
	Computes the 64-bit multiply result of two scalar 32-bit integer values src0 and src1. More...
template<int N>
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::addc(carry, src0, src1);") __ESIMD_API sycl
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::addc(carry, src0, src1);") __ESIMD_API uint32_t addc(uint32_t &carry
template<int N>
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);") __ESIMD_API sycl
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::subb(borrow, src0, src1);") __ESIMD_API uint32_t subb(uint32_t &borrow
template<typename T , int SZ, typename U >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	quot (sycl::ext::intel::esimd::simd< T, SZ > src0, U src1)
	Integral quotient (vector version) More...
template<typename T0 , typename T1 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value, std::remove_const_t< T0 > >	quot (T0 src0, T1 src1)
	Integral quotient (scalar version) More...
template<typename T , int SZ, typename U >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	mod (sycl::ext::intel::esimd::simd< T, SZ > src0, U src1)
	Modulo (vector version) More...
template<typename T0 , typename T1 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value, std::remove_const_t< T0 > >	mod (T0 src0, T1 src1)
	Modulo (scalar version) More...
template<typename T , int SZ, typename U >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	div (sycl::ext::intel::esimd::simd< T, SZ > &remainder, sycl::ext::intel::esimd::simd< T, SZ > src0, U src1)
	Integral division with a vector dividend and a scalar divisor. More...
template<typename T , int SZ, typename U >
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_integral< U >::value &&__ESIMD_DNS::is_esimd_scalar< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	div (sycl::ext::intel::esimd::simd< T, SZ > &remainder, U src0, sycl::ext::intel::esimd::simd< T, SZ > src1)
	Integral division with a scalar dividend and a vector divisor. More...
template<typename RT , typename T0 , typename T1 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< RT >::value &&__ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value, std::remove_const_t< RT > >	div (sycl::ext::intel::esimd::simd< std::remove_const_t< RT >, 1 > &remainder, T0 src0, T1 src1)
	Integral division (scalar version). More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T1 >::value &&std::is_floating_point< T1 >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	dp2 (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Dot product on groups of 4 elements. More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T1 >::value &&std::is_floating_point< T1 >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	dp3 (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Dot product on groups of 4 elements. More...
template<typename T0 , typename T1 , int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T1 >::value &&std::is_floating_point< T1 >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	dp4 (sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
	Dot product on groups of 4 elements. More...
template<typename T , typename U , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	dph (sycl::ext::intel::esimd::simd< T, SZ > src0, U src1, Sat sat={})
	Dot product on groups of 4 elements. More...
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > >	line (sycl::ext::intel::esimd::simd< T, 4 > src0, sycl::ext::intel::esimd::simd< T, SZ > src1, Sat sat={})
	Linear equation. More...
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > >	line (float P, float Q, sycl::ext::intel::esimd::simd< T, SZ > src1, Sat sat={})
	Linear equation. More...
template<typename T , int SZ>
__ESIMD_API sycl::ext::intel::esimd::simd< T, SZ >	frc (sycl::ext::intel::esimd::simd< T, SZ > src0)
	Performs component-wise truncate-to-minus-infinity fraction operation of src0. More...
template<typename T >
__ESIMD_API T	frc (T src0)
	Performs truncate-to-minus-infinity fraction operation of src0. More...
template<typename RT , typename T0 , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ >	lzd (sycl::ext::intel::esimd::simd< T0, SZ > src0, Sat sat={})
template<typename RT , typename T0 , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< RT >::value &&__ESIMD_DNS::is_esimd_scalar< T0 >::value, std::remove_const_t< RT > >	lzd (T0 src0, Sat sat={})
template<typename T , int SZ, typename U , typename V , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T, SZ > >	lrp (sycl::ext::intel::esimd::simd< T, SZ > src0, U src1, V src2, Sat sat={})
template<typename T0 , typename T1 , int SZ>
__ESIMD_API sycl::ext::intel::esimd::simd< T0, SZ >	bf_reverse (sycl::ext::intel::esimd::simd< T1, SZ > src0)
	bf_reverse More...
template<typename T0 , typename T1 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value, std::remove_const_t< T0 > >	bf_reverse (T1 src0)
	bf_reverse More...
template<typename T0 , typename T1 , int SZ, typename U , typename V , typename W >
__ESIMD_API std::enable_if_t< std::is_integral< T1 >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	bf_insert (U src0, V src1, W src2, sycl::ext::intel::esimd::simd< T1, SZ > src3)
	bf_insert More...
template<typename T0 , typename T1 , typename T2 , typename T3 , typename T4 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T4 >::value, std::remove_const_t< T0 > >	bf_insert (T1 src0, T2 src1, T3 src2, T4 src3)
	bf_insert More...
template<typename T0 , typename T1 , int SZ, typename U , typename V >
__ESIMD_API std::enable_if_t< std::is_integral< T1 >::value, sycl::ext::intel::esimd::simd< T0, SZ > >	bf_extract (U src0, V src1, sycl::ext::intel::esimd::simd< T1, SZ > src2)
	bf_extract More...
template<typename T0 , typename T1 , typename T2 , typename T3 >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T3 >::value, std::remove_const_t< T0 > >	bf_extract (T1 src0, T2 src1, T3 src2)
	bf_extract More...
template<int SZ, typename U , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< float, SZ >	sincos (sycl::ext::intel::esimd::simd< float, SZ > &dstcos, U src0, Sat sat={})
template<typename T , int SZ>
ESIMD_DETAIL __ESIMD_API sycl::ext::intel::esimd::simd< T, SZ >	atan (sycl::ext::intel::esimd::simd< T, SZ > src0)
template<typename T >
__ESIMD_API T	atan (T src0)
template<typename T , int SZ>
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > >	acos (sycl::ext::intel::esimd::simd< T, SZ > src0)
template<typename T >
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T >	acos (T src0)
template<typename T , int SZ>
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > >	asin (sycl::ext::intel::esimd::simd< T, SZ > src0)
template<typename T >
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T >	asin (T src0)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	atan2_fast (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<typename T >
float	atan2_fast (T y, T x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	atan2 (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<typename T >
float	atan2 (T y, T x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	fmod (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<typename T >
float	fmod (T y, T x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	sin_emu (sycl::ext::intel::esimd::simd< float, N > x)
template<typename T >
float	sin_emu (T x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	cos_emu (sycl::ext::intel::esimd::simd< float, N > x)
template<typename T >
float	cos_emu (T x)
float	tanh_cody_waite (float x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	tanh_cody_waite (sycl::ext::intel::esimd::simd< float, N > x)
float	tanh (float x)
template<int N>
sycl::ext::intel::esimd::simd< float, N >	tanh (sycl::ext::intel::esimd::simd< float, N > x)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	atan2_fast (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<>
ESIMD_INLINE float	atan2_fast (float y, float x)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	atan2 (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<>
ESIMD_INLINE float	atan2 (float y, float x)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	fmod (sycl::ext::intel::esimd::simd< float, N > y, sycl::ext::intel::esimd::simd< float, N > x)
template<>
ESIMD_INLINE float	fmod (float y, float x)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	sin_emu (sycl::ext::intel::esimd::simd< float, N > x)
template<>
ESIMD_INLINE float	sin_emu (float x0)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	cos_emu (sycl::ext::intel::esimd::simd< float, N > x)
template<>
ESIMD_INLINE float	cos_emu (float x0)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	tanh_cody_waite (sycl::ext::intel::esimd::simd< float, N > x)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N >	tanh (sycl::ext::intel::esimd::simd< float, N > x)
template<typename T , int N>
sycl::ext::intel::esimd::simd< T, N >	dp4 (sycl::ext::intel::esimd::simd< T, N > v1, sycl::ext::intel::esimd::simd< T, N > v2)
template<int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< sycl::half, N >	srnd (sycl::ext::intel::esimd::simd< float, N > src0, sycl::ext::intel::esimd::simd< uint16_t, N > src1)
	srnd - perform stochastic rounding. More...
template<bfn_t FuncControl, typename T , int N>
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);") __ESIMD_API std
	Performs binary function computation with three vector operands. More...
template<bfn_t FuncControl, typename T >
	__SYCL_DEPRECATED ("Please use sycl::ext::intel::esimd::bfn<FuncControl>(src0, src1, src2);") __ESIMD_API std
	Performs binary function computation with three scalar operands. More...
ESIMD_INLINE uint64_t	rdtsc ()
	rdtsc - get the value of timestamp counter. More...
template<split_barrier_action flag>
__ESIMD_API void	split_barrier ()
	Generic work-group split barrier. More...
__ESIMD_API void	split_barrier (split_barrier_action flag)
template<typename T1 , int n1, typename T2 , int n2, typename T3 , int n3, int N = 16>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 >	raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, sycl::ext::intel::esimd::simd< T3, n3 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1)
	Raw sends. More...
template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, 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 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, sycl::ext::intel::esimd::simd< T3, n3 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw sends. More...
template<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 >	raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1)
	Raw send. More...
template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, 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 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw send. More...
template<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API void	raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1)
	Raw sends. More...
template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2>
__ESIMD_API void	raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw sends. More...
template<typename T1 , int n1, int N = 16>
__ESIMD_API void	raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1)
	Raw send. More...
template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1>
__ESIMD_API void	raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw send. More...
__ESIMD_API void	named_barrier_wait (uint8_t id)
	Wait on a named barrier Available only on PVC. More...
template<uint8_t NbarCount>
__ESIMD_API void	named_barrier_init ()
	Initialize number of named barriers for a kernel Available only on PVC. More...
__ESIMD_API void	named_barrier_signal (uint8_t barrier_id, uint8_t producer_consumer_mode, uint32_t num_producers, uint32_t num_consumers)
	Perform signal operation for the given named barrier Available only on PVC. More...
template<typename T , int N>
__ESIMD_API std::enable_if_t<(sizeof(T) *N >=2)>	wait (sycl::ext::intel::esimd::simd< T, N > value)
	Create explicit scoreboard dependency to avoid device code motion across this call and preserve the value computation even if it is unused. More...
template<typename T , typename RegionT >
__ESIMD_API std::enable_if_t<(RegionT::length *sizeof(typename RegionT::element_type) >=2)>	wait (sycl::ext::intel::esimd::simd_view< T, RegionT > value)
	Create explicit scoreboard dependency to avoid device code motion across this call and preserve the value computation even if it is unused. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	SLM gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
	SLM gather. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, NElts >	lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})
	Transposed SLM gather with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, NElts >	lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru)
	Transposed SLM gather with 1 channel. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_gather (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_gather (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
	USM pointer gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts >	lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
	Accessor-based gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N *NElts > >	lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > pass_thru)
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	USM pointer transposed gather with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (const T *p, FlagsT)
	A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT={})
	USM pointer transposed gather with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})
	Accessor-based transposed gather with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{})
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, FlagsT flags)
	A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, uint32_t offset, FlagsT flags)
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT={})
	Accessor-based transposed gather with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > >	lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > pass_thru, FlagsT flags=FlagsT{})
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API void	lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer prefetch gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API void	lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > >	lsc_prefetch (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none>
__ESIMD_API void	lsc_prefetch (const T *p)
	USM pointer prefetch transposed gather with 1 channel. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > >	lsc_prefetch (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based prefetch gather. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > >	lsc_prefetch (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset)
	Accessor-based transposed prefetch gather with 1 channel. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API void	lsc_slm_scatter (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	SLM scatter. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API void	lsc_slm_block_store (uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags=FlagsT{})
	Transposed SLM scatter with 1 channel. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API void	lsc_scatter (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	USM pointer scatter. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename OffsetObjT , typename RegionTy >
__ESIMD_API void	lsc_scatter (T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 >	lsc_scatter (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > >	lsc_scatter (AccessorTy acc, sycl::ext::intel::esimd::simd< __ESIMD_DNS::DeviceAccessorOffsetT, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Accessor-based scatter. More...
template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > >	lsc_scatter (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	USM pointer transposed scatter with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags)
	A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	lsc_block_store (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT={})
	Accessor-based transposed scatter with 1 channel. More...
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_local_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	lsc_block_store (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags=FlagsT{})
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write > &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > >	lsc_block_store (AccessorTy acc, __ESIMD_DNS::DeviceAccessorOffsetT offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags)
	A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	lsc_load_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
	2D USM pointer block load. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>
__ESIMD_API void	lsc_prefetch_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
	2D USM pointer block prefetch. More...
template<typename T , int BlockWidth, int BlockHeight = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false, false>()>
__ESIMD_API void	lsc_store_2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, sycl::ext::intel::esimd::simd< T, N > Vals)
	2D USM pointer block store. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION sycl::ext::intel::esimd::simd< T, N >	lsc_load_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload)
	A variation of 2D stateless block load with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION void	lsc_prefetch_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload)
	A variation of 2D stateless block prefetch with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More...
template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()>
ESIMD_INLINE SYCL_ESIMD_FUNCTION void	lsc_store_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload, sycl::ext::intel::esimd::simd< T, N > Data)
	A variation of 2D stateless block store with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N >	lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	SLM atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&((Op !=sycl::ext::intel::esimd::atomic_op::store &&Op !=sycl::ext::intel::esimd::atomic_op::xchg)||N==1), sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	USM pointer atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename OffsetObjT , typename RegionTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< OffsetObjT, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename Toffset >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1)
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_read > &&(Op==sycl::ext::intel::esimd::atomic_op::load||__ESIMD_DNS::is_device_accessor_with_v< AccessorTy, __ESIMD_DNS::accessor_mode_cap::can_write >), sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of lsc_atomic_update that uses local_accessor as a parameter. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of lsc_atomic_update that uses local_accessor as a parameter. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	Accessor-based atomic. More...
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L2H = cache_hint::none, typename AccessorTy >
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_rw_local_accessor_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N > >	lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred)
	Variant of lsc_atomic_update that uses local_accessor as a parameter. More...
template<lsc_memory_kind Kind = lsc_memory_kind::untyped_global, lsc_fence_op FenceOp = lsc_fence_op::none, lsc_scope Scope = lsc_scope::group, int N = 16>
__ESIMD_API void	lsc_fence (sycl::ext::intel::esimd::simd_mask< N > pred=1)
	Memory fence. More...
__ESIMD_API int32_t	get_hw_thread_id ()
	Get HW Thread ID. More...
__ESIMD_API int32_t	get_subdevice_id ()
	Get subdevice ID. More...

Variables
uint32_t	src0
uint32_t uint32_t	src1