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

Namespaces

 detail
 
 native
 
 xmx
 

Classes

struct  element_aligned_tag
 element_aligned_tag type. More...
 
struct  is_simd_flag_type
 Checks if type is a simd load/store flag. More...
 
struct  is_simd_flag_type< detail::dqword_element_aligned_tag >
 
struct  is_simd_flag_type< element_aligned_tag >
 
struct  is_simd_flag_type< overaligned_tag< N > >
 
struct  is_simd_flag_type< vector_aligned_tag >
 
struct  overaligned_tag
 overaligned_tag type. More...
 
struct  saturation_off_tag
 This type tag represents "saturation off" behavior. More...
 
struct  saturation_on_tag
 Gen hardware supports applying saturation to results of certain operations. More...
 
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...
 
class  slm_allocator
 RAII-style class used to implement "semi-dynamic" SLM allocation. More...
 
struct  vector_aligned_tag
 vector_aligned_tag type. 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  rgba_channel : uint8_t { rgba_channel::R, rgba_channel::G, rgba_channel::B, rgba_channel::A }
 Represents a pixel's channel. More...
 
enum  rgba_channel_mask : uint8_t {
  rgba_channel_mask::R = detail::chR, rgba_channel_mask::G = detail::chG, rgba_channel_mask::GR = detail::chG | detail::chR, rgba_channel_mask::B = detail::chB,
  rgba_channel_mask::BR = detail::chB | detail::chR, rgba_channel_mask::BG = detail::chB | detail::chG, rgba_channel_mask::BGR = detail::chB | detail::chG | detail::chR, rgba_channel_mask::A = detail::chA,
  rgba_channel_mask::AR = detail::chA | detail::chR, rgba_channel_mask::AG = detail::chA | detail::chG, rgba_channel_mask::AGR = detail::chA | detail::chG | detail::chR, rgba_channel_mask::AB = detail::chA | detail::chB,
  rgba_channel_mask::ABR = detail::chA | detail::chB | detail::chR, rgba_channel_mask::ABG = detail::chA | detail::chB | detail::chG, rgba_channel_mask::ABGR = detail::chA | detail::chB | detail::chG | detail::chR
}
 Represents a pixel's channel mask - all possible combinations of enabled channels. More...
 
enum  atomic_op : uint8_t {
  atomic_op::add = 0x0, atomic_op::sub = 0x1, atomic_op::inc = 0x2, atomic_op::dec = 0x3,
  atomic_op::umin = 0x4, atomic_op::umax = 0x5, atomic_op::xchg = 0x6, atomic_op::cmpxchg = 0x7,
  atomic_op::bit_and = 0x8, atomic_op::bit_or = 0x9, atomic_op::bit_xor = 0xa, atomic_op::smin = 0xb,
  atomic_op::smax = 0xc, atomic_op::__SYCL_DEPRECATED =("fmax" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x10, atomic_op::__SYCL_DEPRECATED =("fmin" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x11, atomic_op::fcmpxchg = 0x12,
  atomic_op::__SYCL_DEPRECATED =("fcmpwr" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = fcmpxchg, atomic_op::__SYCL_DEPRECATED =("fadd" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x13, atomic_op::__SYCL_DEPRECATED =("fsub" __ESIMD_USM_DWORD_ATOMIC_TO_LSC) = 0x14, atomic_op::load = 0x15,
  atomic_op::store = 0x16, atomic_op::predec = 0xff
}
 Represents an atomic operation. More...
 
enum  bfn_t : uint8_t { bfn_t::x = 0xAA, bfn_t::y = 0xCC, bfn_t::z = 0xF0 }
 This enum is used to encode all possible logical operations performed on the 3 input operands. More...
 
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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __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< 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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_COND >>
__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< __ESIMD_EMATH_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_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< __ESIMD_EMATH_IEEE_COND >>
__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< __ESIMD_EMATH_IEEE_COND >>
__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 __ESIMD_NS::simd< T, SZ > rndd (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
 Round-down (also known as floor). More...
 
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ > rndu (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
 Round-up (also known as ceil). More...
 
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ > rnde (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
 Round-to-even (also known as round). More...
 
template<typename T , class Sat = __ESIMD_NS::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 __ESIMD_NS::simd< T, SZ > rndz (__ESIMD_NS::simd< float, SZ > src0, Sat sat={})
 Round-to-zero (also known as trunc). More...
 
template<typename T , class Sat = __ESIMD_NS::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_typefbl (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_typefbh (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...
 
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...
 
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, typename Toffset >
__ESIMD_API simd< Tx, N > gather (const Tx *p, simd< Toffset, 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, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N > gather (const Tx *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1)
 A variation of gather API with offsets represented as simd_view object. 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 Tx , int N, typename Toffset >
__ESIMD_API void scatter (Tx *p, simd< Toffset, 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 Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API void scatter (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > vals, simd_mask< N > mask=1)
 A variation of scatter API with offsets represented as simd_view object. 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 >, 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> && sycl::detail::acc_properties::is_accessor_v<AccessorTy> && !sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>>, 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, typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API __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 Tx , int N, typename AccessorTy , class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > > 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) &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< T, N > > gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > > scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
 
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, 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)
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< T, N *get_num_channels_enabled(RGBAMask)> gather_rgba (const T *p, simd_view< Toffset, RegionTy > 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)
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API void scatter_rgba (T *p, simd_view< Toffset, RegionTy > 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 &&!std::is_pointer_v< AccessorT > &&!sycl::detail::acc_properties::is_local_accessor_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 > &&!sycl::detail::acc_properties::is_local_accessor_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<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, 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 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 Flags = overaligned_tag<detail::OperandSize::OWORD>>
__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<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<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< Toffset, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
 Single-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as simd_view object. 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 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 Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< Toffset, N > offset, simd_mask< N > mask)
 No-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1)
 A variation of atomic_update API with offsets represented as simd_view object. 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 offset, simd_mask< N > mask=1)
 A variation of atomic_update API with offset represented as scalar. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset >
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd< Toffset, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API simd< Tx, N > atomic_update (Tx *p, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as simd_view object. 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 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 Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
 Single-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
 Variant of atomic_update that uses local_accessor as a parameter. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > > atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as simd_view object. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value &&((Op !=atomic_op::store &&Op !=atomic_op::xchg)||N==1), simd< Tx, N > > atomic_update (AccessorTy acc, Toffset 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 Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd_mask< N > mask)
 No-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd_mask< N > mask)
 Variant of atomic_update that uses local_accessor as a parameter. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > > atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as simd_view object. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > > atomic_update (AccessorTy acc, Toffset offset, simd_mask< N > mask)
 A variation of atomic_update API with offset represented as scalar. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&sycl::detail::acc_properties::is_accessor_v< AccessorTy > &&!sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< Toffset, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 
template<atomic_op Op, typename Tx , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< Tx, N > > atomic_update (AccessorTy acc, simd< uint32_t, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 Variant of atomic_update that uses local_accessor as a parameter. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy , typename RegionTy = region1d_t<Toffset, N, 1>>
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, simd< Tx, N > > atomic_update (AccessorTy acc, simd_view< Toffset, RegionTy > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 A variation of atomic_update API with offsets represented as simd_view object. More...
 
template<atomic_op Op, typename Tx , int N, typename Toffset , typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&!std::is_pointer< AccessorTy >::value, 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<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<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 Tx , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > &&is_simd_flag_type_v< Flags >, simd< Tx, N > > block_load (AccessorTy acc, uint32_t offset, Flags={})
 Loads a contiguous block of SLM memory referenced by the given local-accessor acc and byte-offset offset, then returns the loaded data as a simd object. More...
 
template<typename Tx , int N, typename AccessorTy , typename Flags = overaligned_tag<detail::OperandSize::OWORD>>
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > &&is_simd_flag_type_v< Flags > > block_store (AccessorTy acc, uint32_t offset, simd< Tx, N > vals, Flags={})
 Variant of block_store that uses local accessor as a parameter. More...
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy >, simd< T, N > > gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
 Variant of gather that uses local accessor as a parameter. More...
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v< AccessorTy > > scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
 Variant of scatter 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< sycl::detail::acc_properties::is_local_accessor_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)
 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< sycl::detail::acc_properties::is_local_accessor_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)
 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<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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, 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 __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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API __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_view< Toffset, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0 &&!std::is_pointer< AccessorTy >::value, simd< T, N > > atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1 &&!std::is_pointer< AccessorTy >::value, simd< T, N > > atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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 &&!std::is_pointer< AccessorTy >::value, 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 Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>, typename AccessorTy >
__ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2 &&!std::is_pointer< AccessorTy >::value, simd< T, N > > atomic_update (AccessorTy acc, simd_view< Toffset, 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 &&!std::is_pointer< AccessorTy >::value, 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...
 

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 391 of file types.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 && !std::is_pointer<AccessorTy>::value, 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 3872 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ atomic_update() [2/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename AccessorTy >
__ESIMD_API __ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 1 && !std::is_pointer<AccessorTy>::value, 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 3833 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 && !std::is_pointer<AccessorTy>::value, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update ( AccessorTy  acc,
simd< Toffset, N >  offset,
simd_mask< N >  mask 
)

Definition at line 3796 of file memory.hpp.

◆ atomic_update() [4/18]

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

◆ atomic_update() [5/18]

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

Definition at line 3847 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ atomic_update() [6/18]

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

Definition at line 3808 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 && !std::is_pointer<AccessorTy>::value, 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 
)

◆ 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 && !std::is_pointer<AccessorTy>::value, 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 3859 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 && !std::is_pointer<AccessorTy>::value, simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update ( AccessorTy  acc,
Toffset  offset,
simd_mask< N >  mask 
)

Definition at line 3820 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 3760 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.

◆ atomic_update() [11/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset >
__ESIMD_API __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 3728 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 3697 of file memory.hpp.

◆ atomic_update() [13/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, RegionTy >  offsets,
simd< T, N >  src0,
simd< T, N >  src1,
simd_mask< N >  mask = 1 
)

◆ atomic_update() [14/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__ESIMD_API __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_view< Toffset, RegionTy >  offsets,
simd< T, N >  src0,
simd_mask< N >  mask = 1 
)

Definition at line 3740 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::src0.

◆ atomic_update() [15/18]

template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>>
__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_view< Toffset, RegionTy >  offsets,
simd_mask< N >  mask = 1 
)

Definition at line 3707 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 
)

◆ 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 3750 of file memory.hpp.

References sycl::_V1::ext::intel::experimental::esimd::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 3717 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.

◆ gather()

template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::gather ( AccessorTy  acc,
simd< uint32_t, N >  offsets,
uint32_t  glob_offset = 0,
simd_mask< N >  mask = 1 
)

Variant of gather that uses local accessor as a parameter.

Collects elements located at given offsets in an accessor and returns them as a single simd object. 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 gather from.
offsetsPer-element offsets in bytes.
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, their values in the resulting vector are undefined.

Definition at line 2217 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< sycl::detail::acc_properties::is_local_accessor_v<AccessorT>, 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 2281 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 668 of file math.hpp.

References pack_mask(), and sycl::_V1::ext::intel::experimental::esimd::src0.

◆ scatter()

template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t< sycl::detail::acc_properties::is_local_accessor_v<AccessorTy> > sycl::_V1::ext::intel::esimd::scatter ( AccessorTy  acc,
simd< uint32_t, N >  offsets,
simd< T, N >  vals,
uint32_t  glob_offset = 0,
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 2244 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< sycl::detail::acc_properties::is_local_accessor_v<AccessorT> > 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 2308 of file memory.hpp.