DPC++ Runtime
Runtime libraries for oneAPI DPC++
memory.hpp File Reference
Include dependency graph for memory.hpp:
This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Namespaces

 cl
 We provide new interfaces for matrix muliply in this patch:
 
 cl::sycl
 
 cl::sycl::ext
 
 cl::sycl::ext::intel
 
 cl::sycl::ext::intel::experimental
 
 cl::sycl::ext::intel::experimental::esimd
 
 cl::sycl::ext::intel::experimental::esimd::detail
 

Macros

#define __ESIMD_GET_SURF_HANDLE(acc)   __ESIMD_NS::get_surface_index(acc)
 

Functions

template<split_barrier_action flag>
__ESIMD_API void cl::sycl::ext::intel::experimental::esimd::split_barrier ()
 Generic work-group split barrier. More...
 
__ESIMD_API void cl::sycl::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 > cl::sycl::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)
 Raw sends load. More...
 
template<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API sycl::ext::intel::esimd::simd< T1, n1 > cl::sycl::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)
 Raw send load. More...
 
template<typename T1 , int n1, typename T2 , int n2, int N = 16>
__ESIMD_API void cl::sycl::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)
 Raw sends store. More...
 
template<typename T1 , int n1, int N = 16>
__ESIMD_API void cl::sycl::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)
 Raw send store. More...
 
__ESIMD_API void cl::sycl::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 cl::sycl::ext::intel::experimental::esimd::named_barrier_init ()
 Initialize number of named barriers for a kernel Available only on PVC. More...
 
__ESIMD_API void cl::sycl::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 NBlocks, int Height, int Width, bool Transposed, bool Transformed>
constexpr int cl::sycl::ext::intel::experimental::esimd::detail::get_lsc_block_2d_data_size ()
 
template<typename T , typename T1 , int N>
ESIMD_INLINE sycl::ext::intel::esimd::simd< T, N > cl::sycl::ext::intel::experimental::esimd::detail::lsc_format_ret (sycl::ext::intel::esimd::simd< T1, N > Vals)
 
template<typename T , uint8_t NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > cl::sycl::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 , uint8_t NElts = 1, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API sycl::ext::intel::esimd::simd< T, NElts > cl::sycl::ext::intel::experimental::esimd::lsc_slm_block_load (uint32_t offset)
 Transposed SLM gather with 1 channel. More...
 
template<typename T , uint8_t 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::ext::intel::esimd::simd< T, N *NElts > > cl::sycl::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 , uint8_t 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::ext::intel::esimd::simd< T, NElts > > cl::sycl::ext::intel::experimental::esimd::lsc_block_load (AccessorTy acc, uint32_t offset)
 Accessor-based transposed gather with 1 channel. More...
 
template<typename T , uint8_t 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>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N *NElts > cl::sycl::ext::intel::experimental::esimd::lsc_gather (const T *p, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
 USM pointer gather. More...
 
template<typename T , uint8_t 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 sycl::ext::intel::esimd::simd< T, NElts > cl::sycl::ext::intel::experimental::esimd::lsc_block_load (const T *p)
 USM pointer transposed gather with 1 channel. More...
 
template<typename T , uint8_t 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 > cl::sycl::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 , uint8_t 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 > cl::sycl::ext::intel::experimental::esimd::lsc_prefetch (AccessorTy acc, uint32_t offset)
 Accessor-based transposed prefetch gather with 1 channel. More...
 
template<typename T , uint8_t 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>
__ESIMD_API void cl::sycl::ext::intel::experimental::esimd::lsc_prefetch (const T *p, sycl::ext::intel::esimd::simd< uint32_t, N > offsets, sycl::ext::intel::esimd::simd_mask< N > pred=1)
 USM pointer prefetch gather. More...
 
template<typename T , uint8_t 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 cl::sycl::ext::intel::experimental::esimd::lsc_prefetch (const T *p)
 USM pointer prefetch transposed gather with 1 channel. More...
 
template<typename T , uint8_t NElts = 1, lsc_data_size DS = lsc_data_size::default_size, int N>
__ESIMD_API void cl::sycl::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 , uint8_t NElts = 1, lsc_data_size DS = lsc_data_size::default_size>
__ESIMD_API void cl::sycl::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 , uint8_t 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 > cl::sycl::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 , uint8_t 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 > cl::sycl::ext::intel::experimental::esimd::lsc_block_store (AccessorTy acc, uint32_t offset, sycl::ext::intel::esimd::simd< T, NElts > vals)
 Accessor-based transposed scatter with 1 channel. More...
 
template<typename T , uint8_t 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>
__ESIMD_API void cl::sycl::ext::intel::experimental::esimd::lsc_scatter (T *p, 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)
 USM pointer scatter. More...
 
template<typename T , uint8_t 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 cl::sycl::ext::intel::experimental::esimd::lsc_block_store (T *p, sycl::ext::intel::esimd::simd< T, NElts > vals)
 USM pointer transposed scatter with 1 channel. More...
 
template<typename T , int BlockWidth, int BlockHeight, int NBlocks, bool Transposed, bool Transformed, bool IsStore = false>
constexpr void cl::sycl::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 > cl::sycl::ext::intel::experimental::esimd::lsc_load2d (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, 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 cl::sycl::ext::intel::experimental::esimd::lsc_prefetch2d (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, 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 cl::sycl::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)
 2D USM pointer block store. 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 > cl::sycl::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 > cl::sycl::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 > cl::sycl::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 AccessorTy >
__ESIMD_API std::enable_if_t<!std::is_pointer< AccessorTy >::value, sycl::ext::intel::esimd::simd< T, N > > cl::sycl::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 > > cl::sycl::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 > > cl::sycl::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<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>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > cl::sycl::ext::intel::experimental::esimd::lsc_atomic_update (T *p, sycl::ext::intel::esimd::simd< uint32_t, 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>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > cl::sycl::ext::intel::experimental::esimd::lsc_atomic_update (T *p, 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)
 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>
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > cl::sycl::ext::intel::experimental::esimd::lsc_atomic_update (T *p, 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)
 USM pointer 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 cl::sycl::ext::intel::experimental::esimd::lsc_fence (sycl::ext::intel::esimd::simd_mask< N > pred=1)
 Memory fence. More...
 

Macro Definition Documentation

◆ __ESIMD_GET_SURF_HANDLE

#define __ESIMD_GET_SURF_HANDLE (   acc)    __ESIMD_NS::get_surface_index(acc)

Definition at line 21 of file memory.hpp.