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

 sycl
 ---— Error handling, matching OpenCL plugin semantics.
 
 sycl::_V1
 
 sycl::_V1::ext
 
 sycl::_V1::ext::intel
 
 sycl::_V1::ext::intel::esimd
 
 sycl::_V1::ext::intel::esimd::detail
 

Enumerations

enum  sycl::_V1::ext::intel::esimd::fence_mask : uint8_t {
  sycl::_V1::ext::intel::esimd::global_coherent_fence = 0x1, sycl::_V1::ext::intel::esimd::l3_flush_instructions = 0x2, sycl::_V1::ext::intel::esimd::l3_flush_texture_data = 0x4, sycl::_V1::ext::intel::esimd::l3_flush_constant_data = 0x8,
  sycl::_V1::ext::intel::esimd::l3_flush_rw_data = 0x10, sycl::_V1::ext::intel::esimd::local_barrier = 0x20, sycl::_V1::ext::intel::esimd::l1_flush_ro_data = 0x40, sycl::_V1::ext::intel::esimd::sw_barrier = 0x80
}
 Represetns a bit mask to control behavior of esimd::fence. More...
 

Functions

template<typename AccessorTy >
__ESIMD_API SurfaceIndex sycl::_V1::ext::intel::esimd::get_surface_index (AccessorTy acc)
 Get surface index corresponding to a SYCL accessor. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< detail::isPowerOf2(N, 32), simd< Tx, N > > sycl::_V1::ext::intel::esimd::gather (const Tx *p, simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Loads ("gathers") elements from different memory locations and returns a vector of them. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t< detail::isPowerOf2(N, 32)> sycl::_V1::ext::intel::esimd::scatter (Tx *p, simd< uint32_t, N > offsets, simd< Tx, N > vals, simd_mask< N > mask=1)
 Writes ("scatters") elements of the input vector to different memory locations. More...
 
template<typename Tx , int N, typename Flags = vector_aligned_tag, class T = detail::__raw_t<Tx>, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::block_load (const Tx *addr, Flags={})
 Loads a contiguous block of memory from given memory address and returns the loaded data as a vector. More...
 
template<typename Tx , int N, typename AccessorTy , typename Flags = vector_aligned_tag, typename = std::enable_if_t<is_simd_flag_type_v<Flags>>, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::block_load (AccessorTy acc, uint32_t offset, Flags={})
 Loads a contiguous block of memory from given accessor and offset and returns the loaded data as a vector. More...
 
template<typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API void sycl::_V1::ext::intel::esimd::block_store (Tx *p, simd< Tx, N > vals)
 Stores elements of a vector to a contiguous block of memory at given address. More...
 
template<typename Tx , int N, typename AccessorTy , class T = detail::__raw_t<Tx>>
__ESIMD_API void sycl::_V1::ext::intel::esimd::block_store (AccessorTy acc, uint32_t offset, simd< Tx, N > vals)
 Stores elements of a vector to a contiguous block of memory represented by an accessor and an offset within this accessor. More...
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&!std::is_pointer< AccessorTy >::value, simd< T, N > > sycl::_V1::ext::intel::esimd::gather (AccessorTy acc, simd< uint32_t, N > offsets, uint32_t glob_offset=0, simd_mask< N > mask=1)
 
template<typename T , int N, typename AccessorTy >
__ESIMD_API std::enable_if_t<(sizeof(T)<=4) &&(N==1||N==8||N==16||N==32) &&!std::is_pointer< AccessorTy >::value > sycl::_V1::ext::intel::esimd::scatter (AccessorTy acc, simd< uint32_t, N > offsets, simd< T, N > vals, uint32_t glob_offset=0, simd_mask< N > mask=1)
 
template<typename T , typename AccessorTy >
__ESIMD_API T sycl::_V1::ext::intel::esimd::scalar_load (AccessorTy acc, uint32_t offset)
 Load a scalar value from an accessor. More...
 
template<typename T , typename AccessorTy >
__ESIMD_API void sycl::_V1::ext::intel::esimd::scalar_store (AccessorTy acc, uint32_t offset, T val)
 Store a scalar value into an accessor. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4, simd< T, N *get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba (const T *p, simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 
template<rgba_channel_mask M>
static void sycl::_V1::ext::intel::esimd::detail::validate_rgba_write_channel_mask ()
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename T , int N>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 > sycl::_V1::ext::intel::esimd::scatter_rgba (T *p, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, simd_mask< N > mask=1)
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<((N==8||N==16||N==32) &&sizeof(T)==4 &&!std::is_pointer_v< AccessorT >), simd< T, N *get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::gather_rgba (AccessorT acc, simd< uint32_t, N > offsets, uint32_t global_offset=0, simd_mask< N > mask=1)
 Gather and transpose pixels from the given memory locations defined by the base specified by acc, the global offset global_offset and a vector of offsets offsets. More...
 
template<rgba_channel_mask RGBAMask = rgba_channel_mask::ABGR, typename AccessorT , int N, typename T = typename AccessorT::value_type>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&sizeof(T)==4 &&!std::is_pointer_v< AccessorT > > sycl::_V1::ext::intel::esimd::scatter_rgba (AccessorT acc, simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(RGBAMask)> vals, uint32_t global_offset=0, simd_mask< N > mask=1)
 Gather data from the memory addressed by accessor acc, offset common for all loaded elements global_offset and per-element offsets offsets, and return it as simd vector. More...
 
template<sycl::ext::intel::esimd::atomic_op Op, typename T , int N, unsigned NumSrc>
constexpr void sycl::_V1::ext::intel::esimd::detail::check_atomic ()
 Check the legality of an atomic call in terms of size and type. More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::atomic_update (Tx *p, simd< unsigned, N > offset, simd_mask< N > mask)
 No-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::atomic_update (Tx *p, simd< unsigned, N > offset, simd< Tx, N > src0, simd_mask< N > mask)
 Single-argument variant of the atomic update operation. More...
 
template<atomic_op Op, typename Tx , int N>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::atomic_update (Tx *p, simd< unsigned, N > offset, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 
template<uint8_t cntl>
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence ()
 esimd::fence sets the memory read/write order. More...
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::fence (fence_mask cntl)
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::barrier ()
 Generic work-group barrier. More...
 
template<uint32_t SLMSize>
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_init ()
 Declare per-work-group slm size. More...
 
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_init (uint32_t size)
 Declare per-work-group slm size. More...
 
template<typename T , int N>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32), simd< T, N > > sycl::_V1::ext::intel::esimd::slm_gather (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Gather operation over the Shared Local Memory. More...
 
template<typename T >
__ESIMD_API T sycl::_V1::ext::intel::esimd::slm_scalar_load (uint32_t offset)
 Load a scalar value from the Shared Local Memory. More...
 
template<typename T , int N>
__ESIMD_API std::enable_if_t<(N==1||N==8||N==16||N==32) &&(sizeof(T)<=4)> sycl::_V1::ext::intel::esimd::slm_scatter (simd< uint32_t, N > offsets, simd< T, N > vals, simd_mask< N > mask=1)
 Scatter operation over the Shared Local Memory. More...
 
template<typename T >
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_scalar_store (uint32_t offset, T val)
 Store a scalar value into the Shared Local Memory. More...
 
template<typename T , int N, rgba_channel_mask RGBAMask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4), simd< T, N *get_num_channels_enabled(RGBAMask)> > sycl::_V1::ext::intel::esimd::slm_gather_rgba (simd< uint32_t, N > offsets, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
template<typename T , int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N==8||N==16||N==32) &&(sizeof(T)==4)> sycl::_V1::ext::intel::esimd::slm_scatter_rgba (simd< uint32_t, N > offsets, simd< T, N *get_num_channels_enabled(Mask)> vals, simd_mask< N > mask=1)
 Gather data from the Shared Local Memory at specified offsets and return it as simd vector. More...
 
template<typename T , int N>
__ESIMD_API simd< T, N > sycl::_V1::ext::intel::esimd::slm_block_load (uint32_t offset)
 Loads a contiguous block of memory from the SLM at given offset and returns the loaded data as a vector. More...
 
template<typename T , int N>
__ESIMD_API void sycl::_V1::ext::intel::esimd::slm_block_store (uint32_t offset, simd< T, N > vals)
 Stores elements of a vector to a contiguous block of SLM at given offset. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > offsets, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API simd< Tx, N > sycl::_V1::ext::intel::esimd::slm_atomic_update (simd< uint32_t, N > offsets, simd< Tx, N > src0, simd< Tx, N > src1, simd_mask< N > mask)
 Atomic update operation performed on SLM. More...
 
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API simd< T, m *N > sycl::_V1::ext::intel::esimd::media_block_load (AccessorTy acc, unsigned x, unsigned y)
 Media block load. More...
 
template<typename T , int m, int N, typename AccessorTy , unsigned plane = 0>
__ESIMD_API void sycl::_V1::ext::intel::esimd::media_block_store (AccessorTy acc, unsigned x, unsigned y, simd< T, m *N > vals)
 Media block store. More...