DPC++ Runtime
Runtime libraries for oneAPI DPC++

Implements the send instruction to send messages to variaous components of the Intel(R) processor graphics, as defined in the documentation at https://01.org/sites/default/files/documentation/intel-gfx-prm-osrc-icllp-vol02a-commandreference-instructions_2.pdf. More...

Collaboration diagram for Raw send APIs.:

Functions

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)
 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 > 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)
 Raw send load. 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)
 Raw sends store. 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)
 Raw send store. More...
 

Detailed Description

Implements the send instruction to send messages to variaous components of the Intel(R) processor graphics, as defined in the documentation at https://01.org/sites/default/files/documentation/intel-gfx-prm-osrc-icllp-vol02a-commandreference-instructions_2.pdf.

Function Documentation

◆ raw_send_load()

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 
)

Raw send load.

Parameters
msgDstis the old value of the destination operand.
msgSrc0is the first source operand of send message.
exDescis the extended message descriptor.
msgDescis the message descriptor.
execSizeis the execution size, which must be a compile time constant.
sfidis the shared function ID, which must be a compile time constant.
numSrc0is the number of GRFs for source-0, which must be a compile time constant.
numDstis the number of GRFs for destination, which must be a compile time constant.
isEOTis the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendcis the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
maskis the predicate to specify enabled channels (optional - default to on).
Returns
the vector value read from memory.

Definition at line 116 of file memory.hpp.

◆ raw_send_store()

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 
)

Raw send store.

Generates a send or sendc instruction for the message gateway.

Parameters
msgSrc0is the first source operand of send message.
exDescis the extended message descriptor.
msgDescis the message descriptor.
execSizeis the execution size, which must be a compile time constant.
sfidis the shared function ID, which must be a compile time constant.
numSrc0is the number of GRFs for source-0, which must be a compile time constant.
isEOTis the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendcis the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
maskis the predicate to specify enabled channels (optional - default to on).

Definition at line 188 of file memory.hpp.

◆ raw_sends_load()

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 
)

Raw sends load.

"s" suffix designates "split" variant - i.e. two sources.

Parameters
msgDstis the old value of the destination operand.
msgSrc0is the first source operand of send message.
msgSrc1is the second source operand of send message.
exDescis the extended message descriptor.
msgDescis the message descriptor.
execSizeis the execution size, which must be a compile time constant.
sfidis the shared function ID, which must be a compile time constant.
numSrc0is the number of GRFs for source-0, which must be a compile time constant.
numSrc1is the number of GRFs for source-1, which must be a compile constant.
numDstis the number of GRFs for destination, which must be a compile time constant.
isEOTis the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendcis the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
maskis the predicate to specify enabled channels (optional - default to on).
Returns
the vector value read from memory.

Definition at line 74 of file memory.hpp.

◆ raw_sends_store()

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 
)

Raw sends store.

"s" suffix designates "split" variant - i.e. two sources.

Parameters
msgSrc0is the first source operand of send message.
msgSrc1is the second source operand of send message.
exDescis the extended message descriptor.
msgDescis the message descriptor.
execSizeis the execution size, which must be a compile time constant.
sfidis the shared function ID, which must be a compile time constant.
numSrc0is the number of GRFs for source-0, which must be a compile time constant.
numSrc1is the number of GRFs for source-1, which must be a compile time constant.
isEOTis the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendcis the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
maskis the predicate to specify enabled channels (optional - default to on).

Definition at line 153 of file memory.hpp.