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 |
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.
using sycl::_V1::ext::intel::esimd::detail::DeviceAccessorOffsetT = typedef uint32_t |
Definition at line 836 of file memory.hpp.
using sycl::_V1::ext::intel::esimd::detail::is_property_list = typedef ext::oneapi::experimental::is_property_list<PropsT> |
Definition at line 159 of file memory_properties.hpp.
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.
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.
Enumerator | |
---|---|
prefetch | |
load | |
store | |
atomic |
Definition at line 552 of file common.hpp.
|
strong |
Enumerator | |
---|---|
nontranspose | |
transpose |
Definition at line 526 of file common.hpp.
|
strong |
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.
|
strong |
Enumerator | |
---|---|
n1 | |
n2 | |
n3 | |
n4 | |
n8 | |
n16 | |
n32 | |
n64 |
Definition at line 453 of file common.hpp.
sycl::_V1::ext::intel::esimd::detail::__ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP | ( | ) |
__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 > | |||
) |
|
constexpr |
Definition at line 548 of file common.hpp.
Referenced by check_cache_hints().
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
AccessorTy | is the sycl::accessor type. |
acc | is the SYCL accessor. |
byte_offset | is the zero-based offsets. |
src0 | is the first atomic operand (expected value). |
src1 | is the second atomic operand (new value). |
pred | is predicates. |
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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
AccessorTy | is the sycl::accessor type. |
acc | is the SYCL accessor. |
byte_offset | is the zero-based offsets. |
src0 | is the first atomic operand. |
pred | is predicates. |
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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
AccessorTy | is the sycl::accessor type. |
acc | is the SYCL accessor. |
byte_offsets | is the zero-based offsets. |
pred | is predicates. |
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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
p | is the base pointer. |
offsets | is the zero-based offsets. |
src0 | is the first atomic operand (expected value). |
src1 | is the second atomic operand (new value). |
pred | predicates. |
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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
p | is the base pointer. |
offsets | is the zero-based offsets. |
src0 | is the first atomic operand. |
pred | is 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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
p | is the base pointer. |
offsets | is the zero-based offsets. |
pred | is 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.
__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.
T | is element type. |
NElts | is the number of elements to load per address. |
PropertyListT | is the list of optional cache-hint properties and the required alignment property. |
AccessorT | is the sycl::accessor type. |
acc | is the SYCL accessor. |
offset | is the zero-based offset in bytes. |
pred | is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. The default is '1' - perform the operation. |
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.
__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.
T | is element type. |
NElts | is the number of elements to load per address. |
PropertyListT | is the list of optional cache-hint properties and the required alignment property. |
AccessorTy | is the sycl::accessor type. |
acc | is the SYCL accessor. |
offset | is the zero-based offset in bytes. |
pred | is 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_thru | contains the values copied to the result if pred is 0. |
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.
__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.
T | is element type. |
NElts | is the number of elements to load per address. |
PropertyListT | is the list of optional cache-hint properties and the required alignment property. |
p | is the base pointer. |
pred | is operation predicate. Zero means operation is skipped entirely, non-zero - operation is performed. |
pass_thru | contains the vector which elements are copied to the returned result when the corresponding element of pred is 0. |
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.
__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 | ||
) |
Definition at line 1191 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.
__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 | ||
) |
|
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().
void sycl::_V1::ext::intel::esimd::detail::check_cache_hints | ( | ) |
Definition at line 566 of file common.hpp.
References are_both(), atomic, sycl::_V1::ext::intel::esimd::cached, load, sycl::_V1::ext::intel::esimd::none, prefetch, sycl::_V1::ext::intel::esimd::read_invalidate, store, sycl::_V1::ext::intel::esimd::streaming, sycl::_V1::ext::intel::esimd::uncached, sycl::_V1::ext::intel::esimd::write_back, and sycl::_V1::ext::intel::esimd::write_through.
Referenced by sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d(), sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d(), and sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d().
|
constexpr |
Definition at line 418 of file common.hpp.
References default_size, u16, u16u32, u16u32h, u32, u64, u8, and u8u32.
|
constexpr |
Definition at line 464 of file common.hpp.
|
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().
|
constexpr |
Definition at line 437 of file common.hpp.
References default_size, u16, u32, u64, and u8.
__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.
T | is element type. |
NElts | is the number of elements to load per address. |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
N | is the number of channels (platform dependent). |
p | is the base pointer. |
offsets | is the zero-based offsets in bytes. |
pred | is predicates. |
pass_thru | contains the vector which elements are copied to the returned result when the corresponding element of pred is 0. |
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().
|
constexpr |
Definition at line 311 of file common.hpp.
References sycl::_V1::detail::device_global_map::add().
|
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.
|
constexpr |
Definition at line 554 of file common.hpp.
References sycl::_V1::ext::intel::esimd::none.
|
constexpr |
Definition at line 211 of file common.hpp.
|
constexpr |
Definition at line 193 of file memory.hpp.
|
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().
|
constexpr |
Check at compile time if given 32 bit positive integer is both:
Definition at line 103 of file common.hpp.
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.
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.
|
constexpr |
Definition at line 5767 of file memory.hpp.
|
inline |
Definition at line 177 of file operators.hpp.
|
inline |
Definition at line 177 of file operators.hpp.
|
inline |
Definition at line 177 of file operators.hpp.
|
inline |
Definition at line 171 of file operators.hpp.
|
inline |
Definition at line 171 of file operators.hpp.
|
inline |
Definition at line 171 of file operators.hpp.
|
inline |
Definition at line 189 of file operators.hpp.
|
inline |
Definition at line 189 of file operators.hpp.
|
inline |
Definition at line 189 of file operators.hpp.
|
inline |
Definition at line 187 of file operators.hpp.
|
inline |
Definition at line 187 of file operators.hpp.
|
inline |
Definition at line 187 of file operators.hpp.
|
inline |
Definition at line 188 of file operators.hpp.
|
inline |
Definition at line 188 of file operators.hpp.
|
inline |
Definition at line 188 of file operators.hpp.
|
inline |
Definition at line 190 of file operators.hpp.
|
inline |
Definition at line 190 of file operators.hpp.
|
inline |
Definition at line 190 of file operators.hpp.
|
inline |
Definition at line 178 of file operators.hpp.
|
inline |
Definition at line 178 of file operators.hpp.
|
inline |
Definition at line 178 of file operators.hpp.
|
inline |
Definition at line 179 of file operators.hpp.
|
inline |
Definition at line 179 of file operators.hpp.
|
inline |
Definition at line 179 of file operators.hpp.
|
inline |
Definition at line 169 of file operators.hpp.
|
inline |
Definition at line 169 of file operators.hpp.
|
inline |
Definition at line 169 of file operators.hpp.
|
inline |
Definition at line 170 of file operators.hpp.
|
inline |
Definition at line 170 of file operators.hpp.
|
inline |
Definition at line 170 of file operators.hpp.
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.
T0 sycl::_V1::ext::intel::esimd::detail::reduce_pair | ( | simd< T1, N1 > | v1, |
simd< T1, N2 > | v2 | ||
) |
Definition at line 969 of file math.hpp.
References isPowerOf2(), and sycl::_V1::ext::intel::esimd::detail::simd_obj_impl< RawTy, N, Derived, SFINAE >::read().
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.
__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.
T | is element type. |
NElts | is the number of elements to store per address. |
DS | is the data size. |
PropertyListT | is the properties with optional cache hints. |
N | is the number of channels (platform dependent). |
p | is the base pointer. |
offsets | is the zero-based offsets in bytes. |
vals | is values to store. |
pred | is 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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
offsets | is the zero-based offsets. |
src0 | is the first atomic operand (expected value). |
src1 | is the second atomic operand (new value). |
pred | is predicates. |
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().
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
offsets | is the zero-based offsets. |
src0 | is the first atomic operand. |
pred | is predicate. |
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.
__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
Op | is operation type. |
T | is element type. |
N | is the number of channels (platform dependent). |
DS | is the data size. |
offsets | is the zero-based offsets. |
pred | is predicate. |
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().
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().
|
constexpr |
Definition at line 268 of file common.hpp.
References sycl::_V1::detail::device_global_map::add().
|
constexpr |
|
constexpr |
Definition at line 222 of file common.hpp.
References sycl::_V1::detail::device_global_map::add().
|
constexpr |
|
static |
Definition at line 4037 of file memory.hpp.
|
inlinestaticconstexpr |
Definition at line 108 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 112 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 111 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 110 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 109 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 116 of file common.hpp.
|
inlineconstexpr |
Definition at line 162 of file memory_properties.hpp.
|
inlineconstexpr |
Definition at line 93 of file common.hpp.
|
inlinestaticconstexpr |
Definition at line 115 of file common.hpp.
Referenced by sycl::_V1::ext::intel::esimd::get_surface_index().