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

Classes

struct  LocalAccessorMarker
 
struct  is_saturation_tag
 
class  cache_hint_wrap
 
struct  lsc_expand_type
 
class  simd_mask_impl
 This class is a simd_obj_impl specialization representing a simd mask, which is basically a simd_obj_impl with fixed element type and limited set of APIs. More...
 
class  simd_obj_impl
 This is a base class for all ESIMD simd classes with real storage (simd, simd_mask_impl). More...
 
class  simd_view_impl
 Base class for "simd view" types. More...
 
struct  esimd_apply_sum
 
struct  esimd_apply_prod
 
struct  esimd_apply_reduced_max
 
struct  esimd_apply_reduced_min
 
struct  get_ext_oneapi_properties
 This helper returns the ext::oneapi::experimental::properties class for ext::oneapi::experimental::properties and it's child in esimd namespace. More...
 
struct  get_ext_oneapi_properties< ext::oneapi::experimental::properties< PropertiesT > >
 
struct  get_ext_oneapi_properties< properties< PropertiesT > >
 
struct  add_alignment_property_helper
 Simply returns 'PropertyListT' as it already has the alignment property. More...
 
struct  add_alignment_property_helper< PropertyListT, Alignment, false >
 Returns a new property list type that contains the properties from 'PropertyListT' and the newly added alignment property. More...
 
class  add_alignment_property
 
struct  make_L1_L2_alignment_properties
 
struct  make_L1_L2_properties
 

Typedefs

using DeviceAccessorOffsetT = uint32_t
 
template<typename PropsT >
using is_property_list = ext::oneapi::experimental::is_property_list< PropsT >
 
template<typename PropertyListT , size_t Alignment>
using add_alignment_property_t = typename add_alignment_property< PropertyListT, Alignment >::type
 
template<cache_hint L1H, cache_hint L2H, size_t Alignment>
using make_L1_L2_alignment_properties_t = typename make_L1_L2_alignment_properties< L1H, L2H, Alignment >::type
 
template<cache_hint L1H, cache_hint L2H>
using make_L1_L2_properties_t = typename make_L1_L2_properties< L1H, L2H >::type
 

Enumerations

enum class  lsc_data_size : uint8_t {
  default_size = 0 , u8 = 1 , u16 = 2 , u32 = 3 ,
  u64 = 4 , u8u32 = 5 , u16u32 = 6 , u16u32h = 7
}
 Data size or format to read or store. More...
 
enum class  lsc_vector_size : uint8_t {
  n1 = 1 , n2 = 2 , n3 = 3 , n4 = 4 ,
  n8 = 5 , n16 = 6 , n32 = 7 , n64 = 8
}
 
enum class  lsc_data_order : uint8_t { nontranspose = 1 , transpose = 2 }
 
enum class  cache_action { prefetch , load , store , atomic }
 

Functions

constexpr ESIMD_INLINE bool isPowerOf2 (unsigned int n)
 Check if a given 32 bit positive integer is a power of 2 at compile time. More...
 
constexpr ESIMD_INLINE bool isPowerOf2 (unsigned int n, unsigned int limit)
 Check at compile time if given 32 bit positive integer is both: More...
 
template<sycl::ext::intel::esimd::atomic_op Op>
constexpr bool has_lsc_equivalent ()
 
template<sycl::ext::intel::esimd::atomic_op Op>
constexpr sycl::ext::intel::esimd::native::lsc::atomic_op to_lsc_atomic_op ()
 
template<sycl::ext::intel::esimd::native::lsc::atomic_op Op>
constexpr sycl::ext::intel::esimd::atomic_op to_atomic_op ()
 
template<sycl::ext::intel::esimd::atomic_op Op>
constexpr int get_num_args ()
 
template<typename T , lsc_data_size DS>
constexpr void check_lsc_data_size ()
 
template<typename T , lsc_data_size DS>
constexpr lsc_data_size finalize_data_size ()
 
template<int VS>
constexpr void check_lsc_vector_size ()
 
template<lsc_vector_size VS>
constexpr uint8_t to_int ()
 
template<int VS>
constexpr lsc_vector_size to_lsc_vector_size ()
 
constexpr bool are_both (cache_hint First, cache_hint Second, cache_hint Val)
 
template<typename PropertyListT >
constexpr bool has_cache_hints ()
 
template<cache_action Action, typename PropertyListT >
void check_cache_hints ()
 
constexpr lsc_data_size expand_data_size (lsc_data_size DS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator^ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator^ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator^ (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator| (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator| (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator| (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator& (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator& (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto operator& (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator% (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator% (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator% (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator<< (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator<< (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator<< (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator>> (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator>> (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator>> (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator+ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator+ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator+ (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator- (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator- (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator- (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator* (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator* (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator* (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator/ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &RHS)
 
template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator/ (const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &LHS, T2 RHS)
 
template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto operator/ (T1 LHS, const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &RHS)
 
 __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP (<, CmpOp::lt, __ESIMD_DNS::is_simd_type_v< SimdTx >) __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(>
 
__ESIMD_DNS::is_simd_type_v< SimdTx > __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP (<=, CmpOp::lte, __ESIMD_DNS::is_simd_type_v< SimdTx >) __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP(>
 
template<typename T0 , typename T1 , int SZ, template< typename RT, typename T, int N > class OpType>
T0 reduce_single (simd< T1, SZ > v)
 
template<typename T0 , typename T1 , int N1, int N2, template< typename RT, typename T, int N > class OpType>
T0 reduce_pair (simd< T1, N1 > v1, simd< T1, N2 > v2)
 
template<typename T0 , typename T1 , int SZ, template< typename RT, typename T, int N > class OpType>
T0 reduce (simd< T1, SZ > v)
 
template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 sum (simd< T1, SZ > v)
 
template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 prod (simd< T1, SZ > v)
 
template<typename RT , typename T , int N>
ESIMD_INLINE simd< RT, N > lsc_format_input (simd< T, N > Vals)
 
template<typename T , typename T1 , int N>
ESIMD_INLINE simd< T, N > lsc_format_ret (simd< T1, N > Vals)
 
template<typename T , int NElts, lsc_data_size DS, typename PropertyListT , int N, typename OffsetT >
__ESIMD_API simd< T, N *NElts > gather_impl (const T *p, simd< OffsetT, N > offsets, simd_mask< N > pred, simd< T, N *NElts > pass_thru)
 USM pointer gather. More...
 
template<typename T , int NElts, lsc_data_size DS, typename PropertyListT , int N, typename Toffset >
__ESIMD_API void scatter_impl (T *p, simd< Toffset, N > offsets, simd< T, N *NElts > vals, simd_mask< N > pred)
 USM pointer scatter. More...
 
constexpr bool isMaskedGatherScatterLLVMAvailable ()
 
template<typename T , int NElts, typename PropertyListT >
__ESIMD_API std::enable_if_t< is_property_list_v< PropertyListT >, simd< T, NElts > > block_load_impl (const T *p, simd_mask< 1 > pred, simd< T, NElts > pass_thru)
 USM pointer transposed gather with 1 channel. More...
 
template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&is_property_list_v< PropertyListT >, simd< T, NElts > > block_load_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd_mask< 1 > pred)
 Accessor-based transposed gather with 1 channel. More...
 
template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read > &&is_property_list_v< PropertyListT >, simd< T, NElts > > block_load_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd_mask< 1 > pred, simd< T, NElts > pass_thru)
 Accessor-based transposed gather with 1 channel. More...
 
template<typename T , int NElts, typename PropertyListT >
__ESIMD_API std::enable_if_t< detail::is_property_list_v< PropertyListT > > block_store_impl (T *p, simd< T, NElts > vals, simd_mask< 1 > pred)
 
template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t< detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write > &&detail::is_property_list_v< PropertyListT > > block_store_impl (AccessorT acc, DeviceAccessorOffsetT offset, simd< T, NElts > vals, simd_mask< 1 > pred)
 
template<rgba_channel_mask M>
static void validate_rgba_write_channel_mask ()
 
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, unsigned NumSrc, bool IsLSC = false>
constexpr void check_atomic ()
 Check the legality of an atomic call in terms of size and type. More...
 
template<typename T , sycl::ext::intel::esimd::atomic_op Op>
constexpr int lsc_to_internal_atomic_op ()
 
template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t< get_num_args< Op >)==0, simd< T, N > > slm_atomic_update_impl (simd< uint32_t, N > offsets, simd_mask< N > pred)
 SLM atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t< get_num_args< Op >)==1, simd< T, N > > slm_atomic_update_impl (simd< uint32_t, N > offsets, simd< T, N > src0, simd_mask< N > pred)
 SLM atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API simd< T, N > slm_atomic_update_impl (simd< uint32_t, N > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 SLM atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==0, simd< T, N > > atomic_update_impl (T *p, simd< Toffset, N > offsets, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==1, simd< T, N > > atomic_update_impl (T *p, simd< Toffset, N > offsets, simd< T, N > src0, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==2, simd< T, N > > atomic_update_impl (T *p, simd< Toffset, N > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 USM pointer atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==0 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, simd< T, N > > atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offsets, simd_mask< N > pred)
 Accessor-based atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==1 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, simd< T, N > > atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd_mask< N > pred)
 Accessor-based atomic. More...
 
template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t< get_num_args< Op >)==2 &&__ESIMD_DNS::is_rw_device_accessor_v< AccessorTy >, simd< T, N > > atomic_update_impl (AccessorTy acc, simd< Toffset, N > byte_offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > pred)
 Accessor-based atomic. More...
 
template<typename PropertiesT , typename KeyT , typename KeyValueT , typename = std::enable_if_t<is_property_list_v<PropertiesT>>>
constexpr auto getPropertyValue (KeyValueT DefaultValue)
 Helper-function that returns the value of the compile time property KeyT if PropertiesT includes it. More...
 

Variables

template<class T >
constexpr bool is_saturation_tag_v = is_saturation_tag<T>::value
 
template<rgba_channel Ch>
static constexpr uint8_t ch = 1 << static_cast<int>(Ch)
 
static constexpr uint8_t chR = ch<rgba_channel::R>
 
static constexpr uint8_t chG = ch<rgba_channel::G>
 
static constexpr uint8_t chB = ch<rgba_channel::B>
 
static constexpr uint8_t chA = ch<rgba_channel::A>
 
static constexpr SurfaceIndex SLM_BTI = 254
 
static constexpr SurfaceIndex INVALID_BTI
 
template<typename PropsT >
constexpr bool is_property_list_v = is_property_list<PropsT>::value
 

Typedef Documentation

◆ add_alignment_property_t

template<typename PropertyListT , size_t Alignment>
using sycl::_V1::ext::intel::esimd::detail::add_alignment_property_t = typedef typename add_alignment_property<PropertyListT, Alignment>::type

Definition at line 230 of file memory_properties.hpp.

◆ DeviceAccessorOffsetT

Definition at line 836 of file memory.hpp.

◆ is_property_list

Definition at line 159 of file memory_properties.hpp.

◆ make_L1_L2_alignment_properties_t

template<cache_hint L1H, cache_hint L2H, size_t Alignment>
using sycl::_V1::ext::intel::esimd::detail::make_L1_L2_alignment_properties_t = typedef typename make_L1_L2_alignment_properties<L1H, L2H, Alignment>::type

Definition at line 241 of file memory_properties.hpp.

◆ make_L1_L2_properties_t

template<cache_hint L1H, cache_hint L2H>
using sycl::_V1::ext::intel::esimd::detail::make_L1_L2_properties_t = typedef typename make_L1_L2_properties<L1H, L2H>::type

Definition at line 250 of file memory_properties.hpp.

Enumeration Type Documentation

◆ cache_action

Enumerator
prefetch 
load 
store 
atomic 

Definition at line 552 of file common.hpp.

◆ lsc_data_order

Enumerator
nontranspose 
transpose 

Definition at line 526 of file common.hpp.

◆ lsc_data_size

Data size or format to read or store.

Enumerator
default_size 
u8 
u16 
u32 
u64 
u8u32 
u16u32 

load 8b, zero extend to 32b; store the opposite

u16u32h 

load 16b, zero extend to 32b; store the opposite

Definition at line 407 of file common.hpp.

◆ lsc_vector_size

Enumerator
n1 
n2 
n3 
n4 
n8 
n16 
n32 
n64 

Definition at line 453 of file common.hpp.

Function Documentation

◆ __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP() [1/2]

sycl::_V1::ext::intel::esimd::detail::__ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP ( )

◆ __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP() [2/2]

__ESIMD_DNS::is_simd_type_v<SimdTx> sycl::_V1::ext::intel::esimd::detail::__ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP ( <=  ,
CmpOp::lte  ,
__ESIMD_DNS::is_simd_type_v< SimdTx >   
)

◆ are_both()

constexpr bool sycl::_V1::ext::intel::esimd::detail::are_both ( cache_hint  First,
cache_hint  Second,
cache_hint  Val 
)
constexpr

Definition at line 548 of file common.hpp.

Referenced by check_cache_hints().

◆ atomic_update_impl() [1/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 2 && __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( AccessorTy  acc,
simd< Toffset, N >  byte_offset,
simd< T, N >  src0,
simd< T, N >  src1,
simd_mask< N >  pred 
)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
AccessorTyis the sycl::accessor type.
Parameters
accis the SYCL accessor.
byte_offsetis the zero-based offsets.
src0is the first atomic operand (expected value).
src1is the second atomic operand (new value).
predis predicates.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 6411 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::get_surface_index(), sycl::_V1::ext::intel::esimd::none, nontranspose, sycl::_V1::ext::intel::esimd::src0, and sycl::_V1::ext::intel::esimd::src1.

◆ atomic_update_impl() [2/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 1 && __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( AccessorTy  acc,
simd< Toffset, N >  byte_offset,
simd< T, N >  src0,
simd_mask< N >  pred 
)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
AccessorTyis the sycl::accessor type.
Parameters
accis the SYCL accessor.
byte_offsetis the zero-based offsets.
src0is the first atomic operand.
predis predicates.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 6354 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::get_surface_index(), sycl::_V1::ext::intel::esimd::none, nontranspose, and sycl::_V1::ext::intel::esimd::src0.

◆ atomic_update_impl() [3/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, typename PropertyListT , typename AccessorTy , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 0 && __ESIMD_DNS::is_rw_device_accessor_v<AccessorTy>, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( AccessorTy  acc,
simd< Toffset, N >  byte_offsets,
simd_mask< N >  pred 
)

Accessor-based atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
AccessorTyis the sycl::accessor type.
Parameters
accis the SYCL accessor.
byte_offsetsis the zero-based offsets.
predis predicates.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 6299 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::get_surface_index(), sycl::_V1::ext::intel::esimd::none, and nontranspose.

◆ atomic_update_impl() [4/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 2, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( T *  p,
simd< Toffset, N >  offsets,
simd< T, N >  src0,
simd< T, N >  src1,
simd_mask< N >  pred 
)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
Parameters
pis the base pointer.
offsetsis the zero-based offsets.
src0is the first atomic operand (expected value).
src1is the second atomic operand (new value).
predpredicates.

Definition at line 6248 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, nontranspose, sycl::_V1::ext::intel::esimd::src0, and sycl::_V1::ext::intel::esimd::src1.

◆ atomic_update_impl() [5/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 1, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( T *  p,
simd< Toffset, N >  offsets,
simd< T, N >  src0,
simd_mask< N >  pred 
)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
Parameters
pis the base pointer.
offsetsis the zero-based offsets.
src0is the first atomic operand.
predis predicates.

Definition at line 6202 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, nontranspose, and sycl::_V1::ext::intel::esimd::src0.

◆ atomic_update_impl() [6/6]

template<atomic_op Op, typename T , int N, lsc_data_size DS, typename PropertyListT , typename Toffset >
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 0, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::atomic_update_impl ( T *  p,
simd< Toffset, N >  offsets,
simd_mask< N >  pred 
)

USM pointer atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.ugm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
PropertyListTis the properties with optional cache hints.
Parameters
pis the base pointer.
offsetsis the zero-based offsets.
predis predicates.

Definition at line 6159 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, and nontranspose.

◆ block_load_impl() [1/3]

template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t<detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read> && is_property_list_v<PropertyListT>, simd<T, NElts> > sycl::_V1::ext::intel::esimd::detail::block_load_impl ( AccessorT  acc,
DeviceAccessorOffsetT  offset,
simd_mask< 1 >  pred 
)

Accessor-based transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm Instruction can load max: DG2(64xD32 or 32xD64), PVC(64xD32 or 64xD64).

Collects elements located at surface and returns them as a single simd object. When sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-byte aligned. When T is 1- or 2-byte type, the memory block is loaded with DWORDs or QWORDs depending on the alignment. Allowed NElts values for 8-byte data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 4-byte data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 2-byte data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 1-byte data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8-byte alignment is required for 8-byte data, or if sizeof(T) * NElts > 256. Otherwise, 4-byte alignment is required.

Template Parameters
Tis element type.
NEltsis the number of elements to load per address.
PropertyListTis the list of optional cache-hint properties and the required alignment property.
AccessorTis the sycl::accessor type.
Parameters
accis the SYCL accessor.
offsetis the zero-based offset in bytes.
predis operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation.
Returns
is a vector of type T and size NElts. The elements of the returned vector for which the corresponding element in pred is 0 are undefined.

Definition at line 967 of file memory.hpp.

References sycl::_V1::ext::oneapi::experimental::detail::Alignment, sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), sycl::_V1::ext::intel::esimd::get_surface_index(), sycl::_V1::ext::intel::esimd::none, transpose, u32, and u64.

◆ block_load_impl() [2/3]

template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t<detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_read> && is_property_list_v<PropertyListT>, simd<T, NElts> > sycl::_V1::ext::intel::esimd::detail::block_load_impl ( AccessorT  acc,
DeviceAccessorOffsetT  offset,
simd_mask< 1 >  pred,
simd< T, NElts >  pass_thru 
)

Accessor-based transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm Instruction can load max: DG2(64xD32 or 32xD64), PVC(64xD32 or 64xD64).

Collects elements located at surface and returns them as a single simd object. When sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-byte aligned. When T is 1- or 2-byte type, the memory block is loaded with DWORDs or QWORDs depending on the alignment. Allowed NElts values for 8-byte data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 4-byte data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 2-byte data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 1-byte data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8-byte alignment is required for 8-byte data, or if sizeof(T) * NElts > 256. Otherwise, 4-byte alignment is required.

Template Parameters
Tis element type.
NEltsis the number of elements to load per address.
PropertyListTis the list of optional cache-hint properties and the required alignment property.
AccessorTyis the sycl::accessor type.
Parameters
accis the SYCL accessor.
offsetis the zero-based offset in bytes.
predis operation predicate. Operation is skipped for index 'i' if pred[0] == 0 the result element is taken from pass_thru[i]. Otherwise, the operation is performed and the result if it copied to the result.
pass_thrucontains the values copied to the result if pred is 0.
Returns
is a vector of type T and size NElts

Definition at line 1068 of file memory.hpp.

References sycl::_V1::ext::oneapi::experimental::detail::Alignment, sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), sycl::_V1::ext::intel::esimd::get_surface_index(), sycl::_V1::ext::intel::esimd::none, transpose, u32, and u64.

◆ block_load_impl() [3/3]

template<typename T , int NElts, typename PropertyListT >
__ESIMD_API std::enable_if_t<is_property_list_v<PropertyListT>, simd<T, NElts> > sycl::_V1::ext::intel::esimd::detail::block_load_impl ( const T *  p,
simd_mask< 1 >  pred,
simd< T, NElts >  pass_thru 
)

USM pointer transposed gather with 1 channel.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm Instruction can load max: DG2(64xD32 or 32xD64), PVC(64xD32 or 64xD64).

Accesses contiguous block of memory of NElts * sizeof(T) bytes starting from the given address p. The maximum size of accessed block is 512 bytes for PVC and 256 bytes for ACM (DG2). When sizeof(T) equal to 8 the address must be 8-byte aligned, otherwise - 4-byte aligned. When T is 1- or 2-byte type, the memory block is loaded with DWORDs or QWORDs depending on the alignment. Allowed NElts values for 8-byte data are 1, 2, 3, 4, 8, 16, 32, 64. Allowed NElts values for 4-byte data are 1, 2, 3, 4, 8, 16, 32, 64, 128. Allowed NElts values for 2-byte data are 2, 4, 8, 16, 32, 64, 128, 256. Allowed NElts values for 1-byte data are 4, 8, 12, 16, 32, 64, 128, 256, 512. 8-byte alignment is required for 8-byte data, or if sizeof(T) * NElts > 256. Otherwise, 4-byte alignment is required.

Template Parameters
Tis element type.
NEltsis the number of elements to load per address.
PropertyListTis the list of optional cache-hint properties and the required alignment property.
Parameters
pis the base pointer.
predis operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed.
pass_thrucontains the vector which elements are copied to the returned result when the corresponding element of pred is 0.
Returns
is a vector of type T and size NElts.

Definition at line 872 of file memory.hpp.

References sycl::_V1::ext::oneapi::experimental::detail::Alignment, sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), sycl::_V1::ext::intel::esimd::none, transpose, u32, and u64.

◆ block_store_impl() [1/2]

template<typename T , int NElts, typename PropertyListT , typename AccessorT >
__ESIMD_API std::enable_if_t<detail::is_device_accessor_with_v< AccessorT, detail::accessor_mode_cap::can_write> && detail::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::detail::block_store_impl ( AccessorT  acc,
DeviceAccessorOffsetT  offset,
simd< T, NElts >  vals,
simd_mask< 1 >  pred 
)

◆ block_store_impl() [2/2]

template<typename T , int NElts, typename PropertyListT >
__ESIMD_API std::enable_if_t<detail::is_property_list_v<PropertyListT> > sycl::_V1::ext::intel::esimd::detail::block_store_impl ( T *  p,
simd< T, NElts >  vals,
simd_mask< 1 >  pred 
)

◆ check_atomic()

template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, unsigned NumSrc, bool IsLSC = false>
constexpr void sycl::_V1::ext::intel::esimd::detail::check_atomic ( )
constexpr

Check the legality of an atomic call in terms of size and type.

Definition at line 4272 of file memory.hpp.

References __ESIMD_FP_ATOMIC_OP_TYPE_CHECK, sycl::_V1::detail::device_global_map::add(), and isPowerOf2().

Referenced by atomic_update_impl(), and slm_atomic_update_impl().

◆ check_cache_hints()

◆ check_lsc_data_size()

template<typename T , lsc_data_size DS>
constexpr void sycl::_V1::ext::intel::esimd::detail::check_lsc_data_size ( )
constexpr

Definition at line 418 of file common.hpp.

References default_size, u16, u16u32, u16u32h, u32, u64, u8, and u8u32.

◆ check_lsc_vector_size()

template<int VS>
constexpr void sycl::_V1::ext::intel::esimd::detail::check_lsc_vector_size ( )
constexpr

Definition at line 464 of file common.hpp.

◆ expand_data_size()

constexpr lsc_data_size sycl::_V1::ext::intel::esimd::detail::expand_data_size ( lsc_data_size  DS)
constexpr

Definition at line 608 of file common.hpp.

References u16, u16u32, u8, and u8u32.

Referenced by atomic_update_impl(), gather_impl(), scatter_impl(), and slm_atomic_update_impl().

◆ finalize_data_size()

template<typename T , lsc_data_size DS>
constexpr lsc_data_size sycl::_V1::ext::intel::esimd::detail::finalize_data_size ( )
constexpr

Definition at line 437 of file common.hpp.

References default_size, u16, u32, u64, and u8.

◆ gather_impl()

template<typename T , int NElts, lsc_data_size DS, typename PropertyListT , int N, typename OffsetT >
__ESIMD_API simd<T, N * NElts> sycl::_V1::ext::intel::esimd::detail::gather_impl ( const T *  p,
simd< OffsetT, N >  offsets,
simd_mask< N >  pred,
simd< T, N *NElts >  pass_thru 
)

USM pointer gather.

Supported platforms: DG2, PVC VISA instruction: lsc_load.ugm

Collects elements located at specified address and returns them as a single simd object.

Template Parameters
Tis element type.
NEltsis the number of elements to load per address.
DSis the data size.
PropertyListTis the properties with optional cache hints.
Nis the number of channels (platform dependent).
Parameters
pis the base pointer.
offsetsis the zero-based offsets in bytes.
predis predicates.
pass_thrucontains the vector which elements are copied to the returned result when the corresponding element of pred is 0.
Returns
is a vector of type T and size N * NElts

Definition at line 119 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, and nontranspose.

Referenced by sycl::_V1::ext::intel::esimd::gather().

◆ get_num_args()

template<sycl::ext::intel::esimd::atomic_op Op>
constexpr int sycl::_V1::ext::intel::esimd::detail::get_num_args ( )
constexpr

Definition at line 311 of file common.hpp.

References sycl::_V1::detail::device_global_map::add().

◆ getPropertyValue()

template<typename PropertiesT , typename KeyT , typename KeyValueT , typename = std::enable_if_t<is_property_list_v<PropertiesT>>>
constexpr auto sycl::_V1::ext::intel::esimd::detail::getPropertyValue ( KeyValueT  DefaultValue)
constexpr

Helper-function that returns the value of the compile time property KeyT if PropertiesT includes it.

If it does not then the default value DefaultValue is returned.

Definition at line 169 of file memory_properties.hpp.

◆ has_cache_hints()

template<typename PropertyListT >
constexpr bool sycl::_V1::ext::intel::esimd::detail::has_cache_hints ( )
constexpr

Definition at line 554 of file common.hpp.

References sycl::_V1::ext::intel::esimd::none.

◆ has_lsc_equivalent()

template<sycl::ext::intel::esimd::atomic_op Op>
constexpr bool sycl::_V1::ext::intel::esimd::detail::has_lsc_equivalent ( )
constexpr

Definition at line 211 of file common.hpp.

◆ isMaskedGatherScatterLLVMAvailable()

constexpr bool sycl::_V1::ext::intel::esimd::detail::isMaskedGatherScatterLLVMAvailable ( )
constexpr

Definition at line 193 of file memory.hpp.

◆ isPowerOf2() [1/2]

constexpr ESIMD_INLINE bool sycl::_V1::ext::intel::esimd::detail::isPowerOf2 ( unsigned int  n)
constexpr

Check if a given 32 bit positive integer is a power of 2 at compile time.

Definition at line 96 of file common.hpp.

Referenced by sycl::_V1::ext::intel::esimd::block_load(), check_atomic(), sycl::_V1::ext::intel::esimd::gather(), reduce(), reduce_pair(), reduce_single(), and sycl::_V1::ext::intel::esimd::slm_atomic_update().

◆ isPowerOf2() [2/2]

constexpr ESIMD_INLINE bool sycl::_V1::ext::intel::esimd::detail::isPowerOf2 ( unsigned int  n,
unsigned int  limit 
)
constexpr

Check at compile time if given 32 bit positive integer is both:

  • a power of 2
  • less or equal to given limit

Definition at line 103 of file common.hpp.

◆ lsc_format_input()

template<typename RT , typename T , int N>
ESIMD_INLINE simd<RT, N> sycl::_V1::ext::intel::esimd::detail::lsc_format_input ( simd< T, N >  Vals)

Definition at line 74 of file memory.hpp.

◆ lsc_format_ret()

template<typename T , typename T1 , int N>
ESIMD_INLINE simd<T, N> sycl::_V1::ext::intel::esimd::detail::lsc_format_ret ( simd< T1, N >  Vals)

Definition at line 88 of file memory.hpp.

◆ lsc_to_internal_atomic_op()

template<typename T , sycl::ext::intel::esimd::atomic_op Op>
constexpr int sycl::_V1::ext::intel::esimd::detail::lsc_to_internal_atomic_op ( )
constexpr

Definition at line 5767 of file memory.hpp.

◆ operator%() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator% ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 177 of file operators.hpp.

◆ operator%() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator% ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 177 of file operators.hpp.

◆ operator%() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator% ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 177 of file operators.hpp.

◆ operator&() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator& ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 171 of file operators.hpp.

◆ operator&() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator& ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 171 of file operators.hpp.

◆ operator&() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator& ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 171 of file operators.hpp.

◆ operator*() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator* ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 189 of file operators.hpp.

◆ operator*() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator* ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 189 of file operators.hpp.

◆ operator*() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator* ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 189 of file operators.hpp.

◆ operator+() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator+ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 187 of file operators.hpp.

◆ operator+() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator+ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 187 of file operators.hpp.

◆ operator+() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator+ ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 187 of file operators.hpp.

◆ operator-() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator- ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 188 of file operators.hpp.

◆ operator-() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator- ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 188 of file operators.hpp.

◆ operator-() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator- ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 188 of file operators.hpp.

◆ operator/() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator/ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 190 of file operators.hpp.

◆ operator/() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator/ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 190 of file operators.hpp.

◆ operator/() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< __ESIMD_DNS::is_valid_simd_elem_type_v<T1> &&__ESIMD_DNS::is_valid_simd_elem_type_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator/ ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 190 of file operators.hpp.

◆ operator<<() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator<< ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 178 of file operators.hpp.

◆ operator<<() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator<< ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 178 of file operators.hpp.

◆ operator<<() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator<< ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 178 of file operators.hpp.

◆ operator>>() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator>> ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 179 of file operators.hpp.

◆ operator>>() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator>> ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 179 of file operators.hpp.

◆ operator>>() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> &&__ESIMD_DNS::is_simd_type_v<SimdTx> >>
auto sycl::_V1::ext::intel::esimd::detail::operator>> ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 179 of file operators.hpp.

◆ operator^() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator^ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 169 of file operators.hpp.

◆ operator^() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator^ ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 169 of file operators.hpp.

◆ operator^() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator^ ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 169 of file operators.hpp.

◆ operator|() [1/3]

template<class T1 , class T2 , int N, template< class, int > class SimdT, class SimdTx = SimdT<T1, N>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator| ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N, SimdT< T1, N >> &  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N, SimdT< T2, N >> &  RHS 
)
inline

Definition at line 170 of file operators.hpp.

◆ operator|() [2/3]

template<class T1 , int N1, template< class, int > class SimdT1, class T2 , class SimdTx = SimdT1<T1, N1>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator| ( const __ESIMD_DNS::simd_obj_impl< __raw_t< T1 >, N1, SimdT1< T1, N1 >> &  LHS,
T2  RHS 
)
inline

Definition at line 170 of file operators.hpp.

◆ operator|() [3/3]

template<class T1 , class T2 , int N2, template< class, int > class SimdT2, class SimdTx = SimdT2<T2, N2>, class = std::enable_if_t< std::is_integral_v<T1> &&std::is_integral_v<T2> >>
auto sycl::_V1::ext::intel::esimd::detail::operator| ( T1  LHS,
const __ESIMD_DNS::simd_obj_impl< __raw_t< T2 >, N2, SimdT2< T2, N2 >> &  RHS 
)
inline

Definition at line 170 of file operators.hpp.

◆ prod()

template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 sycl::_V1::ext::intel::esimd::detail::prod ( simd< T1, SZ v)

Definition at line 1016 of file math.hpp.

◆ reduce()

template<typename T0 , typename T1 , int SZ, template< typename RT, typename T, int N > class OpType>
T0 sycl::_V1::ext::intel::esimd::detail::reduce ( simd< T1, SZ v)

Definition at line 993 of file math.hpp.

References isPowerOf2(), and sycl::_V1::ext::intel::esimd::SZ.

◆ reduce_pair()

template<typename T0 , typename T1 , int N1, int N2, template< typename RT, typename T, int N > class OpType>
T0 sycl::_V1::ext::intel::esimd::detail::reduce_pair ( simd< T1, N1 >  v1,
simd< T1, N2 >  v2 
)

◆ reduce_single()

template<typename T0 , typename T1 , int SZ, template< typename RT, typename T, int N > class OpType>
T0 sycl::_V1::ext::intel::esimd::detail::reduce_single ( simd< T1, SZ v)

Definition at line 953 of file math.hpp.

References isPowerOf2(), and sycl::_V1::ext::intel::esimd::SZ.

◆ scatter_impl()

template<typename T , int NElts, lsc_data_size DS, typename PropertyListT , int N, typename Toffset >
__ESIMD_API void sycl::_V1::ext::intel::esimd::detail::scatter_impl ( T *  p,
simd< Toffset, N >  offsets,
simd< T, N *NElts >  vals,
simd_mask< N >  pred 
)

USM pointer scatter.

Supported platforms: DG2, PVC VISA instruction: lsc_store.ugm

Scatters elements to specific address.

Template Parameters
Tis element type.
NEltsis the number of elements to store per address.
DSis the data size.
PropertyListTis the properties with optional cache hints.
Nis the number of channels (platform dependent).
Parameters
pis the base pointer.
offsetsis the zero-based offsets in bytes.
valsis values to store.
predis predicates.

Definition at line 164 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, and nontranspose.

◆ slm_atomic_update_impl() [1/3]

template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API simd<T, N> sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl ( simd< uint32_t, N >  offsets,
simd< T, N >  src0,
simd< T, N >  src1,
simd_mask< N >  pred 
)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
Parameters
offsetsis the zero-based offsets.
src0is the first atomic operand (expected value).
src1is the second atomic operand (new value).
predis predicates.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 5865 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, nontranspose, sycl::_V1::ext::intel::esimd::src0, and sycl::_V1::ext::intel::esimd::src1.

Referenced by sycl::_V1::ext::intel::esimd::slm_atomic_update().

◆ slm_atomic_update_impl() [2/3]

template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 1, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl ( simd< uint32_t, N >  offsets,
simd< T, N >  src0,
simd_mask< N >  pred 
)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
Parameters
offsetsis the zero-based offsets.
src0is the first atomic operand.
predis predicate.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 5822 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, nontranspose, and sycl::_V1::ext::intel::esimd::src0.

◆ slm_atomic_update_impl() [3/3]

template<atomic_op Op, typename T , int N, lsc_data_size DS>
__ESIMD_API std::enable_if_t<get_num_args<Op>) == 0, simd<T, N> > sycl::_V1::ext::intel::esimd::detail::slm_atomic_update_impl ( simd< uint32_t, N >  offsets,
simd_mask< N >  pred 
)

SLM atomic.

Supported platforms: DG2, PVC VISA instruction: lsc_atomic_<OP>.slm

Template Parameters
Opis operation type.
Tis element type.
Nis the number of channels (platform dependent).
DSis the data size.
Parameters
offsetsis the zero-based offsets.
predis predicate.
Returns
A vector of the old values at the memory locations before the update.

Definition at line 5789 of file memory.hpp.

References check_atomic(), sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::data(), simd_mask< _Tp, _Abi >::data(), expand_data_size(), sycl::_V1::ext::intel::esimd::none, and nontranspose.

Referenced by sycl::_V1::ext::intel::esimd::slm_atomic_update().

◆ sum()

template<typename T0 , typename T1 , int SZ>
ESIMD_INLINE ESIMD_NODEBUG T0 sycl::_V1::ext::intel::esimd::detail::sum ( simd< T1, SZ v)

Definition at line 1008 of file math.hpp.

Referenced by syclcompat::vectorized_sum_abs_diff().

◆ to_atomic_op()

template<sycl::ext::intel::esimd::native::lsc::atomic_op Op>
constexpr sycl::ext::intel::esimd::atomic_op sycl::_V1::ext::intel::esimd::detail::to_atomic_op ( )
constexpr

Definition at line 268 of file common.hpp.

References sycl::_V1::detail::device_global_map::add().

◆ to_int()

template<lsc_vector_size VS>
constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::to_int ( )
constexpr

Definition at line 478 of file common.hpp.

References n1, n16, n2, n3, n32, n4, n64, and n8.

◆ to_lsc_atomic_op()

template<sycl::ext::intel::esimd::atomic_op Op>
constexpr sycl::ext::intel::esimd::native::lsc::atomic_op sycl::_V1::ext::intel::esimd::detail::to_lsc_atomic_op ( )
constexpr

Definition at line 222 of file common.hpp.

References sycl::_V1::detail::device_global_map::add().

◆ to_lsc_vector_size()

template<int VS>
constexpr lsc_vector_size sycl::_V1::ext::intel::esimd::detail::to_lsc_vector_size ( )
constexpr

Definition at line 502 of file common.hpp.

References n1, n16, n2, n3, n32, n4, n64, and n8.

◆ validate_rgba_write_channel_mask()

template<rgba_channel_mask M>
static void sycl::_V1::ext::intel::esimd::detail::validate_rgba_write_channel_mask ( )
static

Definition at line 4037 of file memory.hpp.

Variable Documentation

◆ ch

template<rgba_channel Ch>
constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::ch = 1 << static_cast<int>(Ch)
inlinestaticconstexpr

Definition at line 108 of file common.hpp.

◆ chA

constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::chA = ch<rgba_channel::A>
inlinestaticconstexpr

Definition at line 112 of file common.hpp.

◆ chB

constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::chB = ch<rgba_channel::B>
inlinestaticconstexpr

Definition at line 111 of file common.hpp.

◆ chG

constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::chG = ch<rgba_channel::G>
inlinestaticconstexpr

Definition at line 110 of file common.hpp.

◆ chR

constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::chR = ch<rgba_channel::R>
inlinestaticconstexpr

Definition at line 109 of file common.hpp.

◆ INVALID_BTI

constexpr SurfaceIndex sycl::_V1::ext::intel::esimd::detail::INVALID_BTI
inlinestaticconstexpr
Initial value:
=
static_cast<SurfaceIndex>(-1)
unsigned int SurfaceIndex
Surface index type.
Definition: common.hpp:64

Definition at line 116 of file common.hpp.

◆ is_property_list_v

template<typename PropsT >
constexpr bool sycl::_V1::ext::intel::esimd::detail::is_property_list_v = is_property_list<PropsT>::value
inlineconstexpr

Definition at line 162 of file memory_properties.hpp.

◆ is_saturation_tag_v

template<class T >
constexpr bool sycl::_V1::ext::intel::esimd::detail::is_saturation_tag_v = is_saturation_tag<T>::value
inlineconstexpr

Definition at line 93 of file common.hpp.

◆ SLM_BTI

constexpr SurfaceIndex sycl::_V1::ext::intel::esimd::detail::SLM_BTI = 254
inlinestaticconstexpr

Definition at line 115 of file common.hpp.

Referenced by sycl::_V1::ext::intel::esimd::get_surface_index().