DPC++ Runtime
Runtime libraries for oneAPI DPC++
sycl::_V1::ext::intel::esimd Namespace Reference

Namespaces

 detail
 
 native
 
 xmx
 

Classes

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

Typedefs

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

Enumerations

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

Functions

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

Variables

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

Typedef Documentation

◆ alignment_key

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

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

Definition at line 96 of file memory_properties.hpp.

◆ default_cache_hint_L1

◆ default_cache_hint_L2

◆ default_cache_hint_L3

◆ mask_type_t

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

Definition at line 391 of file types.hpp.

Enumeration Type Documentation

◆ cache_hint

L1, L2 or L3 cache hints.

Enumerator
none 
uncached 

load/store/atomic: do not cache data to cache;

cached 
write_back 

store: write data into cache level and mark the cache line as "dirty".

Upon eviction, the "dirty" data will be written into the furthest subsequent cache;

write_through 

store: immediately write data to the subsequent furthest cache, marking the cache line in the current cache as "not dirty";

streaming 

load: cache data to cache using the evict-first policy to minimize cache pollution caused by temporary streaming data that may only be accessed once or twice; store/atomic: same as write-through, but use the evict-first policy to limit cache pollution by streaming;

read_invalidate 

load: asserts that the cache line containing the data will not be read again until it’s overwritten, therefore the load operation can invalidate the cache line and discard "dirty" data.

If the assertion is violated (the cache line is read again) then behavior is undefined.

const_cached 

load, L2 cache only, next gen GPU after Xe required: asserts that the L2 cache line containing the data will not be written until all invocations of the shader or kernel execution are finished.

If the assertion is violated (the cache line is written), the behavior is undefined.

Definition at line 28 of file memory_properties.hpp.

◆ cache_level

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

Enumerator
L1 
L2 
L3 

Definition at line 25 of file memory_properties.hpp.

◆ fence_flush_op

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

Supported platforms: DG2, PVC

Enumerator
none 
evict 

no operation;

invalidate 

R/W: evict dirty lines; R/W and RO: invalidate clean lines.

clean 

R/W and RO: invalidate all clean lines;.

Definition at line 384 of file common.hpp.

◆ fence_scope

The scope that fence() operation should apply to.

Supported platforms: DG2, PVC

Enumerator
group 

Wait until all previous memory transactions from this thread are observed within the local thread-group.

local 

Wait until all previous memory transactions from this thread are observed within the local sub-slice.

tile 

Wait until all previous memory transactions from this thread are observed in the local tile.

gpu 

Wait until all previous memory transactions from this thread are observed in the local GPU.

gpus 

Wait until all previous memory transactions from this thread are observed across all GPUs in the system.

system 

Global memory data-port only: wait until all previous memory transactions from this thread are observed at the "system" level.

system_acquire 

Global memory data-port only: for GPUs that do not follow PCIe Write ordering for downstream writes targeting device memory, this op will commit to device memory all downstream and peer writes that have reached the device.

Definition at line 350 of file common.hpp.

◆ memory_kind

The target memory kind for fence() operation.

Supported platforms: DG2, PVC

Enumerator
global 
image 

untyped global memory

local 

image (also known as typed global memory)

Definition at line 397 of file common.hpp.

Function Documentation

◆ atomic_update() [1/18]

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

LSC version of the two-argument atomic update.

Definition at line 2909 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [2/18]

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

LSC version of the single-argument atomic update.

Definition at line 2873 of file memory.hpp.

References src0.

◆ atomic_update() [3/18]

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

Definition at line 2839 of file memory.hpp.

◆ atomic_update() [4/18]

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

Definition at line 2923 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [5/18]

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

Definition at line 2884 of file memory.hpp.

References src0.

◆ atomic_update() [6/18]

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

Definition at line 2849 of file memory.hpp.

◆ atomic_update() [7/18]

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

Definition at line 2935 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [8/18]

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

Definition at line 2896 of file memory.hpp.

References src0.

◆ atomic_update() [9/18]

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

Definition at line 2861 of file memory.hpp.

◆ atomic_update() [10/18]

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

LSC version of the two-argument atomic update.

Definition at line 2805 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [11/18]

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

LSC version of the single-argument atomic update.

Definition at line 2775 of file memory.hpp.

References src0.

◆ atomic_update() [12/18]

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

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

Definition at line 2747 of file memory.hpp.

◆ atomic_update() [13/18]

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

Definition at line 2817 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [14/18]

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

Definition at line 2785 of file memory.hpp.

References src0.

◆ atomic_update() [15/18]

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

Definition at line 2755 of file memory.hpp.

◆ atomic_update() [16/18]

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

Definition at line 2827 of file memory.hpp.

References src0, and src1.

◆ atomic_update() [17/18]

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

Definition at line 2795 of file memory.hpp.

References src0.

◆ atomic_update() [18/18]

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

Definition at line 2765 of file memory.hpp.

◆ fbh()

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

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

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

Definition at line 903 of file math.hpp.

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

◆ gather() [1/4]

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

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

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

Definition at line 12646 of file memory.hpp.

◆ gather() [2/4]

template<int VS = 1, typename AccessorT , typename OffsetSimdViewT , typename PassThruSimdViewT , int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< (detail::is_local_accessor_with_v<AccessorT, detail::accessor_mode_cap::can_read> && detail::is_simd_view_type_v<OffsetSimdViewT> && detail::is_simd_view_type_v<PassThruSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::gather ( AccessorT  acc,
OffsetSimdViewT  byte_offsets,
simd_mask< N/VS >  mask,
PassThruSimdViewT  pass_thru,
PropertyListT  props = {} 
)

template <int VS = 1, typename AccessorT, typename OffsetSimdViewT, typename PassThruSimdViewT, int N = PassThruSimdViewT::getSizeX() * PassThruSimdViewT::getSizeY(), typename T = PassThruSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>> simd<T, N> gather(AccessorT acc, OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, PassThruSimdViewT pass_thru, PropertyListT props = {}); This function is identical to (lacc-ga-1) except that the byte_offsets and pass_thru are represented as simd_view.

Definition at line 12571 of file memory.hpp.

◆ gather() [3/4]

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

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

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

Definition at line 12625 of file memory.hpp.

◆ gather() [4/4]

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

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

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

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

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

Definition at line 12518 of file memory.hpp.

◆ gather_rgba()

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

Variant of gather_rgba that uses local accessor as a parameter.

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

Template Parameters
RGBAMaskA pixel's channel mask.
AccessorTThe accessor type for the memory to be loaded/gathered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
NNumber of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32.
Parameters
accThe accessor representing memory address of the access.
offsetsByte offsets of the pixels relative to the base pointer.
global_offsetByte offset of the pixels relative to the base pointer.
maskMemory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined.
Returns
Read data - up to N*4 values of type Tx.

Definition at line 14058 of file memory.hpp.

◆ pack_mask()

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

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

Definition at line 766 of file math.hpp.

References pack_mask(), and src0.

◆ scatter() [1/7]

template<int VS, typename AccessorTy , typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
OffsetSimdViewT  byte_offsets,
simd< T, N >  vals,
PropertyListT  props = {} 
)

template <int VS, typename AccessorTy, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 12898 of file memory.hpp.

◆ scatter() [2/7]

template<int VS, typename AccessorTy , typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
OffsetSimdViewT  byte_offsets,
simd< T, N >  vals,
simd_mask< N/VS >  mask,
PropertyListT  props = {} 
)

template <int VS, typename AccessorTy, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, simd<T, N> vals, simd_mask<N / VS> mask, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the store to (acc + byte_offsets[i]) is skipped.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
maskThe access mask.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 12861 of file memory.hpp.

◆ scatter() [3/7]

template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetSimdViewT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
OffsetSimdViewT  byte_offsets,
ValuesSimdViewT  vals,
PropertyListT  props = {} 
)

template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 12991 of file memory.hpp.

◆ scatter() [4/7]

template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetSimdViewT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<OffsetSimdViewT> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
OffsetSimdViewT  byte_offsets,
ValuesSimdViewT  vals,
simd_mask< N/VS >  mask,
PropertyListT  props = {} 
)

template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetSimdViewT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, OffsetSimdViewT byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the store to (acc + byte_offsets[i]) is skipped.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
maskThe access mask.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 12947 of file memory.hpp.

◆ scatter() [5/7]

template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
simd< OffsetT, N/VS >  byte_offsets,
ValuesSimdViewT  vals,
PropertyListT  props = {} 
)

template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 13077 of file memory.hpp.

◆ scatter() [6/7]

template<int VS = 1, typename AccessorTy , typename ValuesSimdViewT , typename OffsetT , int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< detail::is_local_accessor_with_v<AccessorTy, detail::accessor_mode_cap::can_write> && detail::is_simd_view_type_v<ValuesSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
simd< OffsetT, N/VS >  byte_offsets,
ValuesSimdViewT  vals,
simd_mask< N/VS >  mask,
PropertyListT  props = {} 
)

template <int VS = 1, typename AccessorTy, typename ValuesSimdViewT, typename OffsetT, int N = ValuesSimdViewT::getSizeX() * ValuesSimdViewT::getSizeY(), typename T = ValuesSimdViewT::value_type::element_type, typename PropertyListT = empty_properties_t> void scatter(AccessorTy acc, simd<OffsetT, N / VS> byte_offsets, ValuesSimdViewT vals, simd_mask<N / VS> mask, PropertyListT props = {});

Variation of the API that allows using simd_view without specifying T and N template parameters. Stores ("scatters") elements of the type 'T' to memory locations addressed by the local accessor acc and byte offsets byte_offsets. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the store to (acc + byte_offsets[i]) is skipped.

Template Parameters
VSVector size. It can also be read as the number of writes per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC and only for 4- and 8-byte element vectors.
Parameters
accAccessor referencing the data to store.
byte_offsetsthe vector of 32-bit or 64-bit offsets in bytes represented as a 'simd_view' object. For each i, (acc + byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
maskThe access mask.
propsThe optional compile-time properties. Only 'alignment' and cache hint properties are used.

Definition at line 13038 of file memory.hpp.

◆ scatter() [7/7]

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

Variant of scatter that uses local accessor as a parameter.

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

Template Parameters
TElement type; can only be a 1,2,4-byte integer, sycl::half or float.
NThe number of vector elements. Can be 1, 8, 16 or 32.
AccessorTyThe accessor type.
Parameters
accThe accessor to scatter to.
offsetsPer-element offsets in bytes.
valsValues to write.
glob_offsetOffset in bytes added to each individual element's offset to compute actual memory access offset for that element.
maskMemory access mask. Elements with zero corresponding mask's predicate are not accessed.

Definition at line 13141 of file memory.hpp.

◆ scatter_rgba()

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

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

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

Template Parameters
RGBAMaskPixel's channel mask.
AccessorTThe accessor type for the memory to be stored/scattered. The returned vector elements must match the accessor data type. The loaded elements must be 4 bytes in size.
NThe number of elements to access.
Parameters
offsetsByte offsets of each element.
valsvalues to be written.
global_offsetByte offset of the pixels relative to the base pointer.
maskOperation mask. All-1 by default.

Definition at line 14085 of file memory.hpp.

◆ slm_gather()

template<int VS, typename T , int N, typename OffsetSimdViewT , typename PropertyListT = ext::oneapi::experimental::empty_properties_t>
__ESIMD_API std::enable_if_t< (detail::is_simd_view_type_v<OffsetSimdViewT> && ext::oneapi::experimental::is_property_list_v<PropertyListT>), simd<T, N> > sycl::_V1::ext::intel::esimd::slm_gather ( OffsetSimdViewT  byte_offsets,
simd_mask< N/VS >  mask,
simd< T, N >  pass_thru,
PropertyListT  props = {} 
)

template <int VS, typename T, int N, typename OffsetSimdViewT, typename PropertyListT = empty_props_t> simd <T, N> slm_gather( OffsetSimdViewT byte_offsets, simd_mask<N / VS> mask, simd<T, N> pass_thru, PropertyListT props = {}); Variation of the API that allows using simd_view without specifying T and N template parameters.

Loads ("gathers") elements of the type 'T' from Shared Local Memory locations addressed by byte offsets byte_offsets, and returns the loaded elements. Access to any element's memory location can be disabled via the input vector of predicates mask. If mask[i] is unset, then the load from (byte_offsets[i]) is skipped and the corresponding i-th element from pass_thru operand is returned.

Template Parameters
VSVector size. It can also be read as the number of reads per each address. The parameter 'N' must be divisible by 'VS'. (VS > 1) is supported only on DG2 and PVC.
Parameters
byte_offsetsthe vector of 32-bit offsets in bytes. For each i, (byte_offsets[i]) must be element size aligned. If the alignment property is not passed, then it is assumed that each accessed address is aligned by element-size.
maskThe access mask, defaults to all 1s.
pass_thruThe vector pass through values.
propsThe optional compile-time properties. Only 'alignment' property is used.
Returns
A vector of elements read.

Definition at line 6079 of file memory.hpp.

Variable Documentation

◆ alignment

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

◆ cache_hint_L1

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

Definition at line 130 of file memory_properties.hpp.

◆ cache_hint_L2

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

Definition at line 132 of file memory_properties.hpp.

◆ cache_hint_L3

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

Definition at line 134 of file memory_properties.hpp.