DPC++ Runtime
Runtime libraries for oneAPI DPC++
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 L3 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>
__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)
 Transposed SLM gather with 1 channel. More...
 
template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size>
__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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 flags=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 L3H = 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 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 L3H = 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 flags=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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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< uint32_t, 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 L3H = 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>
__ESIMD_API void lsc_slm_block_store (uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals)
 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 flags=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 L3H = 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 L3H = 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 flags=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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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 L3H = 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