DPC++ Runtime
Runtime libraries for oneAPI DPC++
Namespaces | Classes | Enumerations | Functions
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
 

Enumerations

enum  lsc_scope : uint8_t {
  lsc_scope::group = 0, lsc_scope::local = 1, lsc_scope::tile = 2, lsc_scope::gpu = 3,
  lsc_scope::gpus = 4, lsc_scope::system = 5, lsc_scope::sysacq = 6
}
 The scope that lsc_fence operation should apply to Supported platforms: DG2, PVC. More...
 
enum  lsc_fence_op : uint8_t {
  lsc_fence_op::none = 0, lsc_fence_op::evict = 1, lsc_fence_op::invalidate = 2, lsc_fence_op::discard = 3,
  lsc_fence_op::clean = 4, lsc_fence_op::flushl3 = 5
}
 The lsc_fence operation to apply to caches Supported platforms: DG2, PVC. More...
 
enum  lsc_memory_kind : uint8_t { lsc_memory_kind::untyped_global = 0, lsc_memory_kind::untyped_global_low_pri = 1, lsc_memory_kind::typed_global = 2, lsc_memory_kind::shared_local = 3 }
 The specific LSC shared function to fence with lsc_fence Supported platforms: DG2, PVC. More...
 
enum  lsc_data_size : uint8_t {
  lsc_data_size::default_size = 0, lsc_data_size::u8 = 1, lsc_data_size::u16 = 2, lsc_data_size::u32 = 3,
  lsc_data_size::u64 = 4, lsc_data_size::u8u32 = 5, lsc_data_size::u16u32 = 6, lsc_data_size::u16u32h = 7
}
 Data size or format to read or store. More...
 
enum  cache_hint : uint8_t {
  cache_hint::none = 0, cache_hint::uncached = 1, cache_hint::cached = 2, cache_hint::write_back = 3,
  cache_hint::write_through = 4, cache_hint::streaming = 5, cache_hint::read_invalidate = 6
}
 L1 or L3 cache hint kinds. More...
 
enum  split_barrier_action : uint8_t { split_barrier_action::wait = 0, split_barrier_action::signal = 1 }
 Represents a split barrier action. More...
 
enum  bfn_t : uint8_t { bfn_t::x = 0xAA, bfn_t::y = 0xCC, bfn_t::z = 0xF0 }
 This enum is used to encode all possible logical operations performed on the 3 input operands. More...
 

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<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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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_NODEBUG ESIMD_INLINE 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<argument_type src1_precision, argument_type src2_precision, typename T , int systolic_depth, int repeat_count, typename T0 , typename T1 , typename T2 , int N, int N1, int N2, typename Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpas (sycl::ext::intel::esimd::simd< T0, N > src0, sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
 
template<argument_type src1_precision, argument_type src2_precision, int systolic_depth, int repeat_count, typename T , typename T1 , typename T2 , int N, int N1, int N2, typename Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpas (sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
 DPAS. More...
 
template<argument_type src1_precision, argument_type src2_precision, int systolic_depth, int repeat_count, typename T , typename T1 , typename T2 , int N, int N1, int N2, typename Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpas (sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
 DPAS. More...
 
template<argument_type src1_precision, argument_type src2_precision, int systolic_depth, int repeat_count, typename T , typename T1 , typename T2 , int N, int N1, int N2, typename Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpasw (sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
 DPASW. More...
 
template<argument_type src1_precision, argument_type src2_precision, int systolic_depth, int repeat_count, typename T , typename T1 , typename T2 , int N, int N1, int N2, typename Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpasw2 (sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
 DPASW2. More...
 
static constexpr bfn_t operator~ (bfn_t x)
 
static constexpr bfn_t operator| (bfn_t x, bfn_t y)
 
static constexpr bfn_t operator& (bfn_t x, bfn_t y)
 
static constexpr bfn_t operator^ (bfn_t x, bfn_t y)
 
template<bfn_t FuncControl, typename T , int N>
__ESIMD_API std::enable_if_t< std::is_integral_v< T >, sycl::ext::intel::esimd::simd< T, N > > bfn (sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd< T, N > src2)
 Performs binary function computation with three vector operands. More...
 
template<bfn_t FuncControl, typename T >
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T >::value &&std::is_integral_v< T >, T > bfn (T src0, T src1, T src2)
 Performs binary function computation with three scalar operands. More...
 
template<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<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_load (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)
 
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<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > raw_send_load (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)
 
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<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API void raw_sends_store (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)
 
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<typename T1 , int n1, int N = 16>
__ESIMD_API void raw_send_store (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)
 
__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 > old_values)
 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 > old_values)
 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 > old_values)
 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 , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, 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 , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values)
 
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 > old_values)
 
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<!std::is_pointer_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=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<!std::is_pointer_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 > old_values)
 Accessor-based gather. 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=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 > old_values, 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<!std::is_pointer< AccessorTy >::value &&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{})
 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<!std::is_pointer< AccessorTy >::value &&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)
 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<!std::is_pointer< AccessorTy >::value &&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 > old_values, FlagsT flags=FlagsT{})
 Accessor-based 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 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API void lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, 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<!std::is_pointer< AccessorTy >::value > 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<!std::is_pointer< AccessorTy >::value > lsc_prefetch (AccessorTy acc, uint32_t 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API void lsc_scatter (T *p, sycl::ext::intel::esimd::simd_view< Toffset, 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<!std::is_pointer< AccessorTy >::value > 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)
 Accessor-based scatter. 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, 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<!std::is_pointer< AccessorTy >::value &&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, 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<!std::is_pointer< AccessorTy >::value &&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)
 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, 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_load2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
 
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, 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_prefetch2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y)
 
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, 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_store2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, sycl::ext::intel::esimd::simd< T, N > Vals)
 
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< __ESIMD_DNS::to_lsc_atomic_op< 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 , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==0, sycl::ext::intel::esimd::simd< T, N > > lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, 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< __ESIMD_DNS::to_lsc_atomic_op< 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< __ESIMD_DNS::to_lsc_atomic_op< 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==1, sycl::ext::intel::esimd::simd< T, N > > lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, 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< __ESIMD_DNS::to_lsc_atomic_op< 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< __ESIMD_DNS::to_lsc_atomic_op< 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==2, sycl::ext::intel::esimd::simd< T, N > > lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, 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< __ESIMD_DNS::to_lsc_atomic_op< 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 >
__ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value, 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)
 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<!std::is_pointer< AccessorTy >::value, 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)
 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<!std::is_pointer< AccessorTy >::value, 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)
 Accessor-based atomic. 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...