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), uint > | pack_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), uint > | pack_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), uint > | ballot (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... | |
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... | |
using sycl::_V1::ext::intel::esimd::mask_type_t = typedef detail::simd_mask_storage_t<N> |
__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.
__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.
__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.
__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 | ||
) |
Definition at line 3888 of file memory.hpp.
References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.
__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.
__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.
__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 | ||
) |
Definition at line 3900 of file memory.hpp.
References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.
__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.
__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.
__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.
__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.
__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.
__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 |
||
) |
Definition at line 3774 of file memory.hpp.
References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.
__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.
__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.
__ESIMD_API std::enable_if_t<std::is_integral_v<Toffset> && __ESIMD_DNS::get_num_args<Op>) == 2, sycl::ext::intel::esimd::simd<T, N> > sycl::_V1::ext::intel::esimd::atomic_update | ( | T * | p, |
Toffset | offset, | ||
simd< T, N > | src0, | ||
simd< T, N > | src1, | ||
simd_mask< N > | mask = 1 |
||
) |
Definition at line 3784 of file memory.hpp.
References sycl::_V1::ext::intel::experimental::esimd::src0, and sycl::_V1::ext::intel::experimental::esimd::src1.
__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.
__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.
__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).
src0 | the source operand to count bits in. |
0xFFFFffff
is returned for an element equal to 0
. __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.
T | Element type; can only be a 1,2,4-byte integer, sycl::half or float . |
N | The number of vector elements. Can be 1 , 8 , 16 or 32 . |
AccessorTy | The accessor type. |
acc | The accessor to gather from. |
offsets | Per-element offsets in bytes. |
glob_offset | Offset in bytes added to each individual element's offset to compute actual memory access offset for that element. |
mask | Memory 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.
__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.
RGBAMask | A pixel's channel mask. |
AccessorT | The 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. |
N | Number of pixels to access (matches the size of the offsets vector). Must be 8, 16 or 32. |
acc | The accessor representing memory address of the access. |
offsets | Byte offsets of the pixels relative to the base pointer. |
global_offset | Byte offset of the pixels relative to the base pointer. |
mask | Memory access mask. Pixels with zero corresponding mask's predicate are not accessed. Their values in the resulting vector are undefined. |
Tx
. Definition at line 2281 of file memory.hpp.
__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.
__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.
T | Element type; can only be a 1,2,4-byte integer, sycl::half or float . |
N | The number of vector elements. Can be 1 , 8 , 16 or 32 . |
AccessorTy | The accessor type. |
acc | The accessor to scatter to. |
offsets | Per-element offsets in bytes. |
vals | Values to write. |
glob_offset | Offset in bytes added to each individual element's offset to compute actual memory access offset for that element. |
mask | Memory access mask. Elements with zero corresponding mask's predicate are not accessed. |
Definition at line 2244 of file memory.hpp.
__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.
RGBAMask | Pixel's channel mask. |
AccessorT | The 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. |
N | The number of elements to access. |
offsets | Byte offsets of each element. |
vals | values to be written. |
global_offset | Byte offset of the pixels relative to the base pointer. |
mask | Operation mask. All-1 by default. |
Definition at line 2308 of file memory.hpp.