|
DPC++ Runtime
Runtime libraries for oneAPI DPC++
|
|
|
DPC++ Runtime
Runtime libraries for oneAPI DPC++
|
|
#include <sycl/ext/intel/esimd/common.hpp>#include <sycl/ext/intel/esimd/memory.hpp>#include <sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp>#include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
Go to the source code of this file.
Classes | |
| class | sycl::_V1::ext::intel::experimental::esimd::config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > |
Container class to hold parameters for load2d/store2d functions More... | |
| class | sycl::_V1::ext::intel::esimd::slm_allocator< SLMAmount > |
| RAII-style class used to implement "semi-dynamic" SLM allocation. More... | |
Namespaces | |
| sycl | |
| ---— Error handling, matching OpenCL plugin semantics. | |
| sycl::_V1 | |
| sycl::_V1::ext | |
| sycl::_V1::ext::intel | |
| sycl::_V1::ext::intel::experimental | |
| sycl::_V1::ext::intel::experimental::esimd | |
| sycl::_V1::ext::intel::experimental::esimd::detail | |
| sycl::_V1::ext::intel::esimd | |
Functions | |
| template<split_barrier_action flag> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::split_barrier () |
| Generic work-group split barrier. More... | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::split_barrier (split_barrier_action flag) |
| template<typename T1 , int n1, typename T2 , int n2, typename T3 , int n3, int N = 16> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > | sycl::_V1::ext::intel::experimental::esimd::raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, sycl::ext::intel::esimd::simd< T3, n3 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| Raw sends. More... | |
| template<typename T1 , int n1, typename T2 , int n2, typename T3 , int n3, int N = 16> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > | sycl::_V1::ext::intel::experimental::esimd::raw_sends_load (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, sycl::ext::intel::esimd::simd< T3, n3 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| template<typename T1 , int n1, typename T2 , int n2, int N = 16> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > | sycl::_V1::ext::intel::experimental::esimd::raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| Raw send. More... | |
| template<typename T1 , int n1, typename T2 , int n2, int N = 16> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > | sycl::_V1::ext::intel::experimental::esimd::raw_send_load (sycl::ext::intel::esimd::simd< T1, n1 > msgDst, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| template<typename T1 , int n1, typename T2 , int n2, int N = 16> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::raw_sends (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| Raw sends. More... | |
| template<typename T1 , int n1, typename T2 , int n2, int N = 16> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::raw_sends_store (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, sycl::ext::intel::esimd::simd< T2, n2 > msgSrc1, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| template<typename T1 , int n1, int N = 16> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::raw_send (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| Raw send. More... | |
| template<typename T1 , int n1, int N = 16> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::raw_send_store (sycl::ext::intel::esimd::simd< T1, n1 > msgSrc0, uint32_t exDesc, uint32_t msgDesc, uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT=0, uint8_t isSendc=0, sycl::ext::intel::esimd::simd_mask< N > mask=1) |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::named_barrier_wait (uint8_t id) |
| Wait on a named barrier Available only on PVC. More... | |
| template<uint8_t NbarCount> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::named_barrier_init () |
| Initialize number of named barriers for a kernel Available only on PVC. More... | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::named_barrier_signal (uint8_t barrier_id, uint8_t producer_consumer_mode, uint32_t num_producers, uint32_t num_consumers) |
| Perform signal operation for the given named barrier Available only on PVC. More... | |
| template<typename T , int N> | |
| __ESIMD_API std::enable_if_t<(sizeof(T) *N >=2)> | sycl::_V1::ext::intel::experimental::esimd::wait (sycl::ext::intel::esimd::simd< T, N > value) |
Create explicit scoreboard dependency to avoid device code motion across this call and preserve the value computation even if it is unused. More... | |
| template<typename T , typename RegionT > | |
| __ESIMD_API std::enable_if_t<(RegionT::length *sizeof(typename RegionT::element_type) >=2)> | sycl::_V1::ext::intel::experimental::esimd::wait (sycl::ext::intel::esimd::simd_view< T, RegionT > value) |
Create explicit scoreboard dependency to avoid device code motion across this call and preserve the value computation even if it is unused. More... | |
| template<typename T , int NBlocks, int Height, int Width, bool Transposed, bool Transformed> | |
| constexpr int | sycl::_V1::ext::intel::experimental::esimd::detail::get_lsc_block_2d_data_size () |
| template<typename RT , typename T , int N> | |
| ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, N > | sycl::_V1::ext::intel::experimental::esimd::detail::lsc_format_input (sycl::ext::intel::esimd::simd< T, N > Vals) |
| template<typename T , typename T1 , int N> | |
| ESIMD_INLINE sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::detail::lsc_format_ret (sycl::ext::intel::esimd::simd< T1, N > Vals) |
| template<sycl::ext::intel::esimd::native::lsc::atomic_op Op, typename T , int N, unsigned NumSrc> | |
| constexpr void | sycl::_V1::ext::intel::experimental::esimd::detail::check_lsc_atomic () |
| Check the legality of lsc atomic call in terms of size and type. More... | |
| template<cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none> | |
| constexpr uint32_t | sycl::_V1::ext::intel::experimental::esimd::detail::get_lsc_load_cache_mask () |
| template<cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none> | |
| constexpr uint32_t | sycl::_V1::ext::intel::experimental::esimd::detail::get_lsc_store_cache_mask () |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| SLM gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_gather (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values) |
| SLM gather. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1) |
| Transposed SLM gather with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_load (uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > old_values) |
| Transposed SLM gather with 1 channel. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| USM pointer gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values) |
| USM pointer gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset >, sycl::ext::intel::esimd::simd< T, N *NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N *NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| Accessor-based gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer_v< AccessorTy >, sycl::ext::intel::esimd::simd< T, N *NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_gather (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred, sycl::ext::intel::esimd::simd< T, N *NElts > old_values) |
| Accessor-based gather. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{}) |
| USM pointer transposed gather with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, FlagsT flags) |
| A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (const T *p, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > old_values, FlagsT flags=FlagsT{}) |
| USM pointer transposed gather with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{}) |
| Accessor-based transposed gather with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, FlagsT flags) |
| A variation of lsc_block_load without predicate parameter to simplify use of alignment parameter. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT >, sycl::ext::intel::esimd::simd< T, NElts > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd_mask< 1 > pred, sycl::ext::intel::esimd::simd< T, NElts > old_values, FlagsT flags=FlagsT{}) |
| Accessor-based transposed gather with 1 channel. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| USM pointer prefetch gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > > | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (const T *p) |
| USM pointer prefetch transposed gather with 1 channel. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value > | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| Accessor-based prefetch gather. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value > | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch (AccessorTy acc, uint32_t offset) |
| Accessor-based transposed prefetch gather with 1 channel. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_scatter (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| SLM scatter. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_block_store (uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals) |
| Transposed SLM scatter with 1 channel. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| USM pointer scatter. More... | |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&N==1 > | sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<typename T , int NElts = 1, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value > | sycl::_V1::ext::intel::experimental::esimd::lsc_scatter (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N *NElts > vals, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| Accessor-based scatter. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{}) |
| USM pointer transposed scatter with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t< sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags) |
| A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, sycl::ext::intel::esimd::simd_mask< 1 > pred=1, FlagsT flags=FlagsT{}) |
| Accessor-based transposed scatter with 1 channel. More... | |
| template<typename T , int NElts, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy , typename FlagsT = __ESIMD_DNS::dqword_element_aligned_tag> | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value &&sycl::ext::intel::esimd::is_simd_flag_type_v< FlagsT > > | sycl::_V1::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals, FlagsT flags) |
| A variation of lsc_block_store without predicate parameter to simplify use of alignment parameter. More... | |
| template<typename T , int BlockWidth, int BlockHeight, int NBlocks, bool Transposed, bool Transformed, bool IsStore = false> | |
| constexpr void | sycl::_V1::ext::intel::experimental::esimd::detail::check_lsc_block_2d_restrictions () |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y) |
| 2D USM pointer block load. More... | |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_load2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y) |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y) |
| 2D USM pointer block prefetch. More... | |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch2d (const T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y) |
| template<typename T , int BlockWidth, int BlockHeight = 1, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false, false>()> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, sycl::ext::intel::esimd::simd< T, N > Vals) |
| 2D USM pointer block store. More... | |
| template<typename T , int BlockWidth, int BlockHeight = 1, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, 1u, BlockHeight, BlockWidth, false, false>()> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_store2d (T *Ptr, unsigned SurfaceWidth, unsigned SurfaceHeight, unsigned SurfacePitch, int X, int Y, sycl::ext::intel::esimd::simd< T, N > Vals) |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()> | |
| ESIMD_INLINE SYCL_ESIMD_FUNCTION sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_load_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload) |
A variation of 2D stateless block load with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More... | |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, bool Transposed = false, bool Transformed = false, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, Transposed, Transformed>()> | |
| ESIMD_INLINE SYCL_ESIMD_FUNCTION void | sycl::_V1::ext::intel::experimental::esimd::lsc_prefetch_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload) |
A variation of 2D stateless block prefetch with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More... | |
| template<typename T , int BlockWidth, int BlockHeight = 1, int NBlocks = 1, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, int N = detail::get_lsc_block_2d_data_size< T, NBlocks, BlockHeight, BlockWidth, false, false>()> | |
| ESIMD_INLINE SYCL_ESIMD_FUNCTION void | sycl::_V1::ext::intel::experimental::esimd::lsc_store_2d (config_2d_mem_access< T, BlockWidth, BlockHeight, NBlocks > &payload, sycl::ext::intel::esimd::simd< T, N > Data) |
A variation of 2D stateless block store with parameters passed as config_2d_mem_access object Note: Compatibility with future hardware versions is not guaranteed. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred) |
| SLM atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred) |
| SLM atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size> | |
| __ESIMD_API sycl::ext::intel::esimd::simd< T, N > | sycl::_V1::ext::intel::experimental::esimd::lsc_slm_atomic_update (sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred) |
| SLM atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==0, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred) |
| USM pointer atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==0, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==0, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==1, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred) |
| USM pointer atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==1, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==1 &&((Op !=sycl::ext::intel::esimd::atomic_op::store &&Op !=sycl::ext::intel::esimd::atomic_op::xchg)||N==1), sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==2, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< Toffset, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred) |
| USM pointer atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==2, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd_view< Toffset, RegionTy > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< __ESIMD_DNS::to_lsc_atomic_op< Op >)>)==2, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (T *p, Toffset offset, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred) |
| Accessor-based atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd_mask< N > pred) |
| Accessor-based atomic. More... | |
| template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, lsc_data_size DS = lsc_data_size::default_size, cache_hint L1H = cache_hint::none, cache_hint L3H = cache_hint::none, typename AccessorTy > | |
| __ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::experimental::esimd::lsc_atomic_update (AccessorTy acc, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd_mask< N > pred) |
| Accessor-based atomic. More... | |
| template<lsc_memory_kind Kind = lsc_memory_kind::untyped_global, lsc_fence_op FenceOp = lsc_fence_op::none, lsc_scope Scope = lsc_scope::group, int N = 16> | |
| __ESIMD_API void | sycl::_V1::ext::intel::experimental::esimd::lsc_fence (sycl::ext::intel::esimd::simd_mask< N > pred=1) |
| Memory fence. More... | |
| __ESIMD_API int32_t | sycl::_V1::ext::intel::experimental::esimd::get_hw_thread_id () |
| Get HW Thread ID. More... | |
| __ESIMD_API int32_t | sycl::_V1::ext::intel::experimental::esimd::get_subdevice_id () |
| Get subdevice ID. More... | |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > offset, simd_mask< N > mask) |
LSC version of no argument variant of the atomic_update - accepts native::lsc::atomic_op instead of atomic_op as atomic operation template argument. More... | |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< Toffset, RegionTy > offsets, simd_mask< N > mask=1) |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==0, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, Toffset offset, simd_mask< N > mask=1) |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > offset, simd< T, N > src0, simd_mask< N > mask) |
| LSC version of the single-argument atomic update. More... | |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< Toffset, RegionTy > offsets, simd< T, N > src0, simd_mask< N > mask=1) |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==1, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, Toffset offset, simd< T, N > src0, simd_mask< N > mask=1) |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd< Toffset, N > offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask) |
| LSC version of the two-argument atomic update. More... | |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset , typename RegionTy = sycl::ext::intel::esimd::region1d_t<Toffset, N, 1>> | |
| __ESIMD_API std::enable_if_t< __ESIMD_DNS::get_num_args< Op >)==2, simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, simd_view< Toffset, RegionTy > offsets, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1) |
| template<native::lsc::atomic_op Op, typename T , int N, typename Toffset > | |
| __ESIMD_API std::enable_if_t< std::is_integral_v< Toffset > &&__ESIMD_DNS::get_num_args< Op >)==2, sycl::ext::intel::esimd::simd< T, N > > | sycl::_V1::ext::intel::esimd::atomic_update (T *p, Toffset offset, simd< T, N > src0, simd< T, N > src1, simd_mask< N > mask=1) |
1.8.17