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 > >
 
class  simd
 The main simd vector class. More...
 
class  simd_view
 This class represents a reference to a sub-region of a base simd object. More...
 
class  simd_view< BaseTy, region1d_scalar_t< ViewedElemT > >
 This is a specialization of simd_view class with a single element. More...
 
class  simd_view< BaseTy, std::pair< region1d_scalar_t< ViewedElemT >, NestedRegion > >
 This is a specialization of nested simd_view class with a single element. More...
 

Typedefs

using uchar = unsigned char
 
using ushort = unsigned short
 
using uint = unsigned int
 
using SurfaceIndex = unsigned int
 Surface index type. More...
 
template<int N>
using mask_type_t = detail::simd_mask_storage_t< N >
 
template<int N>
using simd_mask = detail::simd_mask_type< N >
 Represents a simd mask os size N. More...
 

Enumerations

enum class  rgba_channel : uint8_t { R , G , B , A }
 Represents a pixel's channel. 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 , __SYCL_DEPRECATED =("use umin") = umin , umax = 0x5 , __SYCL_DEPRECATED =("use smax") = umax ,
  xchg = 0x6 , cmpxchg = 0x7 , bit_and = 0x8 , bit_or = 0x9 ,
  bit_xor = 0xa , smin = 0xb , __SYCL_DEPRECATED =("use smin") = smin , smax = 0xc ,
  __SYCL_DEPRECATED =("use smax") = 0xc , __SYCL_DEPRECATED =("fmax" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x10 , __SYCL_DEPRECATED =("fmin" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x11 , fcmpxchg = 0x12 ,
  __SYCL_DEPRECATED =("fcmpwr" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = fcmpxchg , __SYCL_DEPRECATED =("fadd" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x13 , __SYCL_DEPRECATED =("fsub" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x14 , load = 0x15 ,
  store = 0x16 , predec = 0xff
}
 Represents an atomic operation. More...
 
enum  fence_mask : uint8_t {
  global_coherent_fence = 0x1 , l3_flush_instructions = 0x2 , l3_flush_texture_data = 0x4 , l3_flush_constant_data = 0x8 ,
  l3_flush_rw_data = 0x10 , local_barrier = 0x20 , l1_flush_ro_data = 0x40 , sw_barrier = 0x80
}
 Represetns a bit mask to control behavior of esimd::fence. More...
 

Functions

template<class T , int N>
__ESIMD_API simd< T, N > merge (simd< T, N > a, simd< T, N > b, simd_mask< N > m)
 "Merges" elements of the input simd object according to the merge mask. More...
 
template<int N>
__ESIMD_API simd_mask< N > merge (simd_mask< N > a, simd_mask< N > b, simd_mask< N > m)
 "Merges" elements of the input masks according to the merge mask. More...
 
template<class BaseT1 , class BaseT2 , class RegionT1 , class RegionT2 , class = std::enable_if_t< (shape_type<RegionT1>::length == shape_type<RegionT2>::length) && std::is_same_v<detail::element_type_t<BaseT1>, detail::element_type_t<BaseT2>>>>
__ESIMD_API auto merge (simd_view< BaseT1, RegionT1 > v1, simd_view< BaseT2, RegionT2 > v2, simd_mask< shape_type< RegionT1 >::length > m)
 "Merges" elements of vectors referenced by the input views. More...
 
constexpr int is_channel_enabled (rgba_channel_mask M, rgba_channel Ch)
 
constexpr int get_num_channels_enabled (rgba_channel_mask M)
 
template<typename T0 , typename T1 , int SZ>
__ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v< T0 >||std::is_same_v< T1, T0 >, simd< T0, SZ > > saturate (simd< T1, SZ > src)
 Conversion of input vector elements of type T1 into vector of elements of type T0 with saturation. More...
 
template<typename TRes , typename TArg , int SZ>
ESIMD_DETAIL __ESIMD_API std::enable_if_t< !std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value, simd< TRes, SZ > > abs (simd< TArg, SZ > src0)
 Get absolute value (vector version) More...
 
template<typename TRes , typename TArg >
__ESIMD_API std::enable_if_t<!std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value &&detail::is_esimd_scalar< TRes >::value &&detail::is_esimd_scalar< TArg >::value, std::remove_const_t< TRes > > abs (TArg src0)
 Get absolute value (scalar version) More...
 
template<typename T1 , int SZ>
__ESIMD_API simd< T1, SZ > abs (simd< T1, SZ > src0)
 Get absolute value (vector version). More...
 
template<typename T1 >
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T1 >::value, std::remove_const_t< T1 > > abs (T1 src0)
 Get absolute value (scalar version). More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API simd< T, SZ > max (simd< T, SZ > src0, simd< T, SZ > src1, Sat sat={})
 Selects component-wise the maximum of the two vectors. More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > > max (simd< T, SZ > src0, T src1, Sat sat={})
 Selects maximums for each element of the input vector and a scalar. More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > > max (T src0, simd< T, SZ > src1, Sat sat={})
 Selects maximums for each element of the input scalar and a vector. More...
 
template<typename T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< detail::is_esimd_scalar< T >::value, T > max (T src0, T src1, Sat sat={})
 Selects maximum between two scalar values. More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API simd< T, SZ > min (simd< T, SZ > src0, simd< T, SZ > src1, Sat sat={})
 Selects component-wise the minimum of the two vectors. More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > > min (simd< T, SZ > src0, T src1, Sat sat={})
 Selects minimums for each element of the input vector and a scalar. More...
 
template<typename T , int SZ, class Sat = saturation_off_tag>
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T >::value, simd< T, SZ > > min (T src0, simd< T, SZ > src1, Sat sat={})
 Selects minimums for each element of the input scalar and a vector. More...
 
template<typename T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< detail::is_esimd_scalar< T >::value, T > min (T src0, T src1, Sat sat={})
 Selects minimum between two scalar values. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > inv (simd< T, N > src, Sat sat={})
 Inversion - calculates (1/x). More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T inv (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > log2 (simd< T, N > src, Sat sat={})
 Logarithm base 2. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T log2 (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > exp2 (simd< T, N > src, Sat sat={})
 Exponent base 2. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T exp2 (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > sqrt (simd< T, N > src, Sat sat={})
 Square root. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T sqrt (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API simd< T, N > sqrt_ieee (simd< T, N > src, Sat sat={})
 IEEE754-compliant square root. Supports float and double. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4) >>
__ESIMD_API T sqrt_ieee (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > rsqrt (simd< T, N > src, Sat sat={})
 Square root reciprocal - calculates 1/sqrt(x). More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T rsqrt (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > sin (simd< T, N > src, Sat sat={})
 Sine. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T sin (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > cos (simd< T, N > src, Sat sat={})
 Cosine. More...
 
template<typename T , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API T cos (T src, Sat sat={})
 Scalar version. More...
 
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__ESIMD_API simd< T, N > pow (simd< T, N > src0, simd< U, N > src1, Sat sat={})
 Power - calculates src0 in power of src1. More...
 
template<class T , int N, class U , class Sat = saturation_off_tag, class = std::enable_if_t< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__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< detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4) >>
__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, SZ > log (simd< T, SZ > src0, Sat sat={})
 Computes the natural logarithm of the given argument. More...
 
template<class T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T log (T src0, Sat sat={})
 
template<class T , int SZ, class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > exp (simd< T, SZ > src0, Sat sat={})
 Computes e raised to the power of the given argument. More...
 
template<class T , class Sat = saturation_off_tag>
ESIMD_NODEBUG ESIMD_INLINE T exp (T src0, Sat sat={})
 
template<typename T , int SZ, class Sat = sycl::ext::intel::esimd ::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ > rndd (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, SZ > rndu (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, SZ > rnde (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, SZ > rndz (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, SZ > floor (const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
 "Floor" operation, vector version - alias of rndd. More...
 
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE RT floor (float src0, Sat sat={})
 "Floor" operation, scalar version - alias of rndd. More...
 
template<typename RT , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > ceil (const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
 "Ceiling" operation, vector version - alias of rndu. More...
 
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
ESIMD_INLINE RT ceil (float src0, Sat sat={})
 "Ceiling" operation, scalar version - alias of rndu. More...
 
template<typename RT , int SZ, class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ > trunc (const sycl::ext::intel::esimd::simd< float, SZ > &src0, Sat sat={})
 Round to integral value using the round to zero rounding mode (vector version). More...
 
template<typename RT , class Sat = sycl::ext::intel::esimd::saturation_off_tag>
__ESIMD_API RT trunc (float src0, Sat sat={})
 Round to integral value using the round to zero rounding mode (scalar version). More...
 
template<int N>
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<(N==8||N==16||N==32), 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 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...
 
template<typename AccessorTy >
__ESIMD_API SurfaceIndex get_surface_index (AccessorTy acc)
 Get surface index corresponding to a SYCL accessor. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< detail::isPowerOf2(N, 32), simd< Tx, N > > gather (const Tx *p, simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Loads ("gathers") elements from different memory locations and returns a vector of them. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< detail::isPowerOf2(N, 32)> scatter (Tx *p, simd< uint32_t, N > offsets, simd< Tx, N > vals, simd_mask< N > mask=1)
 Writes ("scatters") elements of the input vector to different memory locations. More...
 
template<typename Tx , int N, typename Flags = vector_aligned_tag, class T = detail::__raw_t<Tx>, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
__ESIMD_API simd< Tx, N > block_load (const Tx *addr, Flags={})
 Loads a contiguous block of memory from given memory address 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>>, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > block_load (AccessorTy acc, uint32_t offset, Flags={})
 Loads a contiguous block of memory from given accessor and offset and returns the loaded data as a vector. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API void block_store (Tx *p, simd< Tx, N > vals)
 Stores elements of a vector to a contiguous block of memory at given address. More...
 
template<typename Tx , int N, typename AccessorTy , class T = detail::__raw_t<Tx>>
__ESIMD_API void block_store (AccessorTy acc, uint32_t offset, simd< Tx, N > vals)
 Stores elements of a vector to a contiguous block of memory represented by an accessor and an offset within this accessor. More...
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&!std::is_pointer< AccessorTy >::value, simd< T, N > > gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
 Accessor-based gather. More...
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&!std::is_pointer< AccessorTy >::value > scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
 Accessor-based scatter. More...
 
template<typename T , typename AccessorTy >
__ESIMD_API T scalar_load (AccessorTy acc, uint32_t offset)
 Load a scalar value from an accessor. More...
 
template<typename T , typename AccessorTy >
__ESIMD_API void scalar_store (AccessorTy acc, uint32_t offset, T val)
 Store a scalar value into an accessor. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4, simd< T, N *get_num_channels_enabled(RGBAMask)> > gather_rgba (const T *p, simd< uint32_t, 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>
__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)
 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 AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<((N==8||N==16||N==32) &&sizeof(T)==4 &&!std::is_pointer_v< AccessorT >), simd< T, N *get_num_channels_enabled(RGBAMask)> > gather_rgba (AccessorT acc, simd< uint32_t, N > offsets, uint32_t global_offset=0, simd_mask< N > mask=1)
 Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 &&!std::is_pointer_v< AccessorT > > scatter_rgba (AccessorT acc, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, uint32_t global_offset=0, simd_mask< N > mask=1)
 Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< unsigned, N > offset, simd_mask< N > mask)
  More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< unsigned, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
  More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< unsigned, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 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<uint8_t cntl>
__ESIMD_API void fence ()
 esimd::fence sets the memory read/write order. More...
 
__ESIMD_API void fence (fence_mask cntl)
 
__ESIMD_API void barrier ()
 Generic work-group barrier. More...
 
template<uint32_t SLMSize>
__ESIMD_API void slm_init ()
 Declare per-work-group slm size. More...
 
__ESIMD_API void slm_init (uint32_t size)
 Declare per-work-group slm size. More...
 
template<typename T , int N>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32), simd< T, N > > slm_gather (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Gather operation over the Shared Local Memory. More...
 
template<typename T >
__ESIMD_API T slm_scalar_load (uint32_t offset)
 Load a scalar value from the Shared Local Memory. More...
 
template<typename T , int N>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32) &&(sizeof(T)<=4)> slm_scatter (simd< uint32_t, N > offsets, simd< T, N > vals, simd_mask< N > mask=1)
 Scatter operation over the Shared Local Memory. More...
 
template<typename T >
__ESIMD_API void slm_scalar_store (uint32_t offset, T val)
 Store a scalar value into the Shared Local Memory. More...
 
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4), simd< T, N *get_num_channels_enabled(RGBAMask)> > slm_gather_rgba (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
template<typename T , int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4)> slm_scatter_rgba (simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(Mask)> vals, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
template<typename T , int N>
__ESIMD_API simd< T, N > slm_block_load (uint32_t offset)
 Loads a contiguous block of memory from the SLM at given offset and returns the loaded data as a vector. More...
 
template<typename T , int N>
__ESIMD_API void slm_block_store (uint32_t offset, simd< T, N > vals)
 Stores elements of a vector to a contiguous block of SLM at given offset. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > slm_atomic_update (simd< uint32_t, N > offsets, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<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 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>
__ESIMD_API simd< T, N > atomic_update (T *p, simd< unsigned, 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>
__ESIMD_API simd< T, N > atomic_update (T *p, simd< unsigned, 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>
__ESIMD_API simd< T, N > atomic_update (T *p, simd< unsigned, N > offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask)
 LSC version of the two-argument atomic update. More...
 

Variables

static constexpr saturation_off_tag saturation_off {}
 Type tag object representing "saturation off" behavior. More...
 
static constexpr saturation_on_tag saturation_on {}
 Type tag object representing "saturation on" behavior. More...
 
constexpr element_aligned_tag element_aligned = {}
 
constexpr vector_aligned_tag vector_aligned = {}
 
template<unsigned N>
constexpr overaligned_tag< N > overaligned = {}
 
template<typename T >
static constexpr bool is_simd_flag_type_v = is_simd_flag_type<T>::value
 Checks if given type is a simd load/store flag. More...
 

Typedef Documentation

◆ mask_type_t

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

Definition at line 392 of file types.hpp.

Function Documentation

◆ atomic_update() [1/3]

template<native::lsc::atomic_op Op, typename T , int N>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::atomic_update ( T *  p,
simd< unsigned, 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 1751 of file memory.hpp.

◆ atomic_update() [2/3]

template<native::lsc::atomic_op Op, typename T , int N>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::atomic_update ( T *  p,
simd< unsigned, N >  offset,
simd< T, N >  src0,
simd_mask< N >  mask 
)

LSC version of the single-argument atomic update.

Definition at line 1743 of file memory.hpp.

◆ atomic_update() [3/3]

template<native::lsc::atomic_op Op, typename T , int N>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::atomic_update ( T *  p,
simd< unsigned, 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 1735 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 804 of file math.hpp.

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

◆ pack_mask()

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

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

Definition at line 668 of file math.hpp.