DPC++ Runtime
Runtime libraries for oneAPI DPC++
Shared local memory access functions.
Collaboration diagram for Shared local memory access functions.:

Functions

template<uint32_t SLMSize>
__ESIMD_API void cl::sycl::ext::intel::esimd::slm_init ()
 Declare per-work-group slm size. More...
 
__ESIMD_API void cl::sycl::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 > > cl::sycl::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 cl::sycl::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)> cl::sycl::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 cl::sycl::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)> > cl::sycl::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)> cl::sycl::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 > cl::sycl::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 cl::sycl::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 std::enable_if_t< detail::check_atomic< Op, T, N, 0 >), simd< Tx, N > > cl::sycl::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 std::enable_if_t< detail::check_atomic< Op, T, N, 1 >), simd< Tx, N > > cl::sycl::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 std::enable_if_t< detail::check_atomic< Op, T, N, 2 >), simd< Tx, N > > cl::sycl::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...
 

Detailed Description

Function Documentation

◆ slm_atomic_update() [1/3]

template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t<detail::check_atomic<Op, T, N, 2>), simd<Tx, N> > cl::sycl::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.

Two source operands version. See description of template and function parameters in atomic update operation docs.

Definition at line 1104 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_atomic_update() [2/3]

template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t<detail::check_atomic<Op, T, N, 1>), simd<Tx, N> > cl::sycl::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.

One source operands version. See description of template and function parameters in atomic update operation docs.

Definition at line 1092 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_atomic_update() [3/3]

template<atomic_op Op, typename Tx , int N, class T = detail::__raw_t<Tx>>
__ESIMD_API std::enable_if_t<detail::check_atomic<Op, T, N, 0>), simd<Tx, N> > cl::sycl::ext::intel::esimd::slm_atomic_update ( simd< uint32_t, N >  offsets,
simd_mask< N >  mask 
)

Atomic update operation performed on SLM.

No source operands version. See description of template and function parameters in atomic update operation docs.

Definition at line 1082 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_block_load()

template<typename T , int N>
__ESIMD_API simd<T, N> cl::sycl::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.

Template Parameters
TElement type.
NNumber of elements to load, N * sizeof(Tx) must be 1, 2, 4 or 8 owords long.
Parameters
offsetThe offset to load from in bytes. Must be oword-aligned.
Returns
A vector of loaded elements.

Definition at line 1038 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_block_store()

template<typename T , int N>
__ESIMD_API void cl::sycl::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.

Template Parameters
TElement type.
NNumber of elements to store, N * sizeof(Tx) must be 1, 2, 4 or 8 owords long.
Parameters
offsetThe offset in bytes to store at. Must be oword-aligned.
valsThe vector to store.

Definition at line 1062 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_gather()

template<typename T , int N>
__ESIMD_API std::enable_if_t<(N == 1 || N == 8 || N == 16 || N == 32), simd<T, N> > cl::sycl::ext::intel::esimd::slm_gather ( simd< uint32_t, N >  offsets,
simd_mask< N >  mask = 1 
)

Gather operation over the Shared Local Memory.

This API has almost the same interface as the accessor-based gather, except that it does not have the accessor and the global offset parameters.

Definition at line 948 of file memory.hpp.

◆ slm_gather_rgba()

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)> > cl::sycl::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.

See usm_gather_rgba for information about the operation semantics and parameter restrictions/interdependencies.

Template Parameters
TThe element type of the returned vector.
NThe number of elements to access.
RGBAMaskPixel's channel mask.
Parameters
offsetsByte offsets within the SLM of each element.
maskOperation mask. All-1 by default.
Returns
Gathered data as an N - element vector.

Definition at line 999 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.

◆ slm_init() [1/2]

template<uint32_t SLMSize>
__ESIMD_API void cl::sycl::ext::intel::esimd::slm_init ( )

Declare per-work-group slm size.

Template Parameters
SLMSizeShared Local Memory (SLM) size

Definition at line 931 of file memory.hpp.

◆ slm_init() [2/2]

__ESIMD_API void cl::sycl::ext::intel::esimd::slm_init ( uint32_t  size)

Declare per-work-group slm size.

Non-constant argument version to be used with specialization constants only.

Parameters
sizeShared Local Memory (SLM) size

Definition at line 938 of file memory.hpp.

◆ slm_scalar_load()

template<typename T >
__ESIMD_API T cl::sycl::ext::intel::esimd::slm_scalar_load ( uint32_t  offset)

Load a scalar value from the Shared Local Memory.

Template Parameters
Ttype of the value
Parameters
offsetSLM offset in bytes
Returns
the loaded value

Definition at line 958 of file memory.hpp.

◆ slm_scalar_store()

template<typename T >
__ESIMD_API void cl::sycl::ext::intel::esimd::slm_scalar_store ( uint32_t  offset,
T  val 
)

Store a scalar value into the Shared Local Memory.

Template Parameters
Ttype of the value
Parameters
offsetSLM offset in bytes
valvalue to store

Definition at line 982 of file memory.hpp.

◆ slm_scatter()

template<typename T , int N>
__ESIMD_API std::enable_if_t<(N == 1 || N == 8 || N == 16 || N == 32) && (sizeof(T) <= 4)> cl::sycl::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.

This API has almost the same interface as the accessor-based scatter, except that it does not have the accessor and the global offset parameters.

Definition at line 971 of file memory.hpp.

◆ slm_scatter_rgba()

template<typename T , int N, rgba_channel_mask Mask>
__ESIMD_API std::enable_if_t<(N == 8 || N == 16 || N == 32) && (sizeof(T) == 4)> cl::sycl::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.

See usm_scatter_rgba for information about the operation semantics and parameter restrictions/interdependencies.

Template Parameters
TThe element type of the returned vector.
NThe number of elements to access.
MaskPixel's channel mask.
Parameters
offsetsByte offsets within the SLM of each element.
valsvalues to be written.
maskOperation mask. All-1 by default.

Definition at line 1018 of file memory.hpp.

References __ESIMD_GET_SURF_HANDLE.