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	remove_alignment_property

struct	remove_alignment_property< properties< std::tuple< alignment_key::value_t< Alignment >, LastTs... > > >

struct	remove_alignment_property< properties< std::tuple< FirstT, alignment_key::value_t< Alignment >, LastTs... > > >

struct	add_or_replace_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<typename PropertyListT >
using	remove_alignment_property_t = typename remove_alignment_property< PropertyListT >::type

template<typename PropertyListT , size_t Alignment>
using	add_or_replace_alignment_property_t = typename add_or_replace_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 ()

template<cache_hint Val>
constexpr bool	are_all (cache_hint First, cache_hint Second)

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 PropertyListT , cache_level Level>
constexpr cache_hint	getCacheHintForIntrin ()
	Extracts a cache hint with the given 'Level' to pass it to ESIMD/GENX intrinsics. More...

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 231 of file memory_properties.hpp.

◆ add_or_replace_alignment_property_t

template<typename PropertyListT , size_t Alignment>

using sycl::_V1::ext::intel::esimd::detail::add_or_replace_alignment_property_t = typedef typename add_or_replace_alignment_property<PropertyListT, Alignment>::type

Definition at line 263 of file memory_properties.hpp.

◆ DeviceAccessorOffsetT

using sycl::_V1::ext::intel::esimd::detail::DeviceAccessorOffsetT = typedef uint32_t

Definition at line 1263 of file memory.hpp.

◆ is_property_list

template<typename PropsT >

using sycl::_V1::ext::intel::esimd::detail::is_property_list = typedef ext::oneapi::experimental::is_property_list<PropsT>

Definition at line 160 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 274 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 283 of file memory_properties.hpp.

◆ remove_alignment_property_t

template<typename PropertyListT >

using sycl::_V1::ext::intel::esimd::detail::remove_alignment_property_t = typedef typename remove_alignment_property<PropertyListT>::type

Definition at line 250 of file memory_properties.hpp.

Enumeration Type Documentation

◆ cache_action

enum sycl::_V1::ext::intel::esimd::detail::cache_action

strong

Enumerator
prefetch
load
store
atomic

Definition at line 548 of file common.hpp.

◆ lsc_data_order

enum sycl::_V1::ext::intel::esimd::detail::lsc_data_order : uint8_t

strong

Enumerator
nontranspose
transpose

Definition at line 521 of file common.hpp.

◆ lsc_data_size

enum sycl::_V1::ext::intel::esimd::detail::lsc_data_size : uint8_t

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 402 of file common.hpp.

◆ lsc_vector_size

enum sycl::_V1::ext::intel::esimd::detail::lsc_vector_size : uint8_t

strong

Enumerator
n1
n2
n3
n4
n8
n16
n32
n64

Definition at line 448 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_all()

template<cache_hint Val>

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

constexpr

Definition at line 544 of file common.hpp.

◆ 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

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.

Parameters

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.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 8988 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(), 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

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.

Parameters

acc	is the SYCL accessor.
byte_offset	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 8927 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(), 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

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.

Parameters

acc	is the SYCL accessor.
byte_offsets	is the zero-based offsets.
pred	is predicates.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 8874 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(), 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

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.

Parameters

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 8818 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(), 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

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.

Parameters

p	is the base pointer.
offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicates.

Definition at line 8768 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(), 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

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.

Parameters

p	is the base pointer.
offsets	is the zero-based offsets.
pred	is predicates.

Definition at line 8727 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(), 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

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.

Parameters

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.

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 1392 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(), 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

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.

Parameters

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.

Returns: is a vector of type T and size NElts

Definition at line 1491 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(), 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

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.

Parameters

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.

Returns: is a vector of type T and size NElts.

Definition at line 1299 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(), 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
	)

Definition at line 1610 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(), transpose, u32, and u64.

◆ 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
	)

Definition at line 1554 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(), transpose, u32, and u64.

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

References __ESIMD_FP_ATOMIC_OP_TYPE_CHECK, sycl::_V1::detail::device_global_map::add(), sycl::_V1::ext::oneapi::fmax(), sycl::_V1::ext::oneapi::fmin(), isPowerOf2(), sycl::_V1::ext::intel::math::umax(), and sycl::_V1::ext::intel::math::umin().

Referenced by atomic_update_impl(), and slm_atomic_update_impl().

◆ check_cache_hints()

template<cache_action Action, typename PropertyListT >

void sycl::_V1::ext::intel::esimd::detail::check_cache_hints ( )

Definition at line 561 of file common.hpp.

References 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().

◆ 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 413 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 459 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 603 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 432 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

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).

Parameters

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.

Returns: is a vector of type T and size N * NElts

Definition at line 133 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(), 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 307 of file common.hpp.

References sycl::_V1::detail::device_global_map::add(), sycl::_V1::ext::oneapi::fmax(), sycl::_V1::ext::oneapi::fmin(), sycl::_V1::ext::intel::math::umax(), and sycl::_V1::ext::intel::math::umin().

◆ getCacheHintForIntrin()

template<typename PropertyListT , cache_level Level>

constexpr cache_hint sycl::_V1::ext::intel::esimd::detail::getCacheHintForIntrin ( )

constexpr

Extracts a cache hint with the given 'Level' to pass it to ESIMD/GENX intrinsics.

If PropertyListT does not have the requested cache-hint, then 'cache_hint::none' is returned.

Definition at line 102 of file memory.hpp.

References sycl::_V1::ext::intel::esimd::L1, sycl::_V1::ext::intel::esimd::L2, and sycl::_V1::ext::intel::esimd::none.

◆ 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 170 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 550 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 208 of file common.hpp.

◆ isMaskedGatherScatterLLVMAvailable()

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

constexpr

Definition at line 203 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(), sycl::_V1::ext::intel::esimd::media_block_load(), 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 7763 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 1556 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 1533 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
	)

Definition at line 1509 of file math.hpp.

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

◆ 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 1493 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

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).

Parameters

p	is the base pointer.
offsets	is the zero-based offsets in bytes.
vals	is values to store.
pred	is predicates.

Definition at line 176 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(), 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

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

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.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 7861 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

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

offsets	is the zero-based offsets.
src0	is the first atomic operand.
pred	is predicate.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 7818 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

Op	is operation type.
T	is element type.
N	is the number of channels (platform dependent).
DS	is the data size.

Parameters

offsets	is the zero-based offsets.
pred	is predicate.

Returns: A vector of the old values at the memory locations before the update.

Definition at line 7785 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 1548 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 264 of file common.hpp.

References sycl::_V1::detail::device_global_map::add(), sycl::_V1::ext::oneapi::fmax(), sycl::_V1::ext::oneapi::fmin(), sycl::_V1::ext::intel::math::umax(), and sycl::_V1::ext::intel::math::umin().

◆ to_int()

template<lsc_vector_size VS>

constexpr uint8_t sycl::_V1::ext::intel::esimd::detail::to_int ( )

constexpr

Definition at line 473 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 218 of file common.hpp.

References sycl::_V1::detail::device_global_map::add(), sycl::_V1::ext::oneapi::fmax(), sycl::_V1::ext::oneapi::fmin(), sycl::_V1::ext::intel::math::umax(), and sycl::_V1::ext::intel::math::umin().

◆ 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 497 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 5331 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)

sycl::_V1::ext::intel::esimd::SurfaceIndex

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 163 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().

Classes

Typedefs

Enumerations

Functions

Variables

Typedef Documentation

◆ add_alignment_property_t

◆ add_or_replace_alignment_property_t

◆ DeviceAccessorOffsetT

◆ is_property_list

◆ make_L1_L2_alignment_properties_t

◆ make_L1_L2_properties_t

◆ remove_alignment_property_t

Enumeration Type Documentation

◆ cache_action

◆ lsc_data_order

◆ lsc_data_size

◆ lsc_vector_size

Function Documentation

◆ __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP() [1/2]

◆ __ESIMD_DEF_SIMD_OBJ_IMPL_CMP_OP() [2/2]

◆ are_all()

◆ atomic_update_impl() [1/6]

◆ atomic_update_impl() [2/6]

◆ atomic_update_impl() [3/6]

◆ atomic_update_impl() [4/6]

◆ atomic_update_impl() [5/6]

◆ atomic_update_impl() [6/6]

◆ block_load_impl() [1/3]

◆ block_load_impl() [2/3]

◆ block_load_impl() [3/3]

◆ block_store_impl() [1/2]

◆ block_store_impl() [2/2]

◆ check_atomic()

◆ check_cache_hints()

◆ check_lsc_data_size()

◆ check_lsc_vector_size()

◆ expand_data_size()

◆ finalize_data_size()

◆ gather_impl()

◆ get_num_args()

◆ getCacheHintForIntrin()

◆ getPropertyValue()

◆ has_cache_hints()

◆ has_lsc_equivalent()

◆ isMaskedGatherScatterLLVMAvailable()

◆ isPowerOf2() [1/2]

◆ isPowerOf2() [2/2]

◆ lsc_format_input()

◆ lsc_format_ret()

◆ lsc_to_internal_atomic_op()

◆ operator%() [1/3]

◆ operator%() [2/3]

◆ operator%() [3/3]

◆ operator&() [1/3]

◆ operator&() [2/3]

◆ operator&() [3/3]

◆ operator*() [1/3]

◆ operator*() [2/3]

◆ operator*() [3/3]

◆ operator+() [1/3]

◆ operator+() [2/3]

◆ operator+() [3/3]

◆ operator-() [1/3]

◆ operator-() [2/3]

◆ operator-() [3/3]

◆ operator/() [1/3]

◆ operator/() [2/3]

◆ operator/() [3/3]

◆ operator<<() [1/3]

◆ operator<<() [2/3]

◆ operator<<() [3/3]

◆ operator>>() [1/3]

◆ operator>>() [2/3]

◆ operator>>() [3/3]

◆ operator^() [1/3]

◆ operator^() [2/3]

◆ operator^() [3/3]

◆ operator|() [1/3]

◆ operator|() [2/3]