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 (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<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2, typename T3 , int n3>
__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, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw sends. 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 (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<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2>
__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, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw send. More...

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<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2>
__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, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw sends. More...

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<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1>
__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, sycl::ext::intel::esimd::simd_mask< execSize > mask=1)
	Raw send. 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() [1/4]

template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2>

__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,
		sycl::ext::intel::esimd::simd_mask< execSize >	mask = `1`
	)

Raw send.

Template Parameters

execSize	is the execution size.
sfid	is the shared function ID.
numSrc0	is the number of GRFs for source-0.
numDst	is the number of GRFs for destination.
isEOT	is the flag that indicates whether this is an EOT message (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used (optional - default to 0).

Parameters

msgDst	is the old value of the destination operand.
msgSrc0	is the first source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
mask	is the predicate to specify enabled channels (optional - default to on).

Returns: the vector value read from memory

Definition at line 197 of file memory.hpp.

◆ raw_send() [2/4]

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.

Parameters

msgDst	is the old value of the destination operand.
msgSrc0	is the first source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
execSize	is the execution size, which must be a compile time constant.
sfid	is the shared function ID, which must be a compile time constant.
numSrc0	is the number of GRFs for source-0, which must be a compile time constant.
numDst	is the number of GRFs for destination, which must be a compile time constant.
isEOT	is the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
mask	is the predicate to specify enabled channels (optional - default to on).

Returns: the vector value read from memory.

Definition at line 158 of file memory.hpp.

◆ raw_send() [3/4]

template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1>

__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,
		sycl::ext::intel::esimd::simd_mask< execSize >	mask = `1`
	)

Raw send.

Generates a send or sendc instruction for the message gateway.

Template Parameters

execSize	is the execution size.
sfid	is the shared function ID.
numSrc0	is the number of GRFs for source-0.
isEOT	is the flag that indicates whether this is an EOT message (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used (optional - default to 0).

Parameters

msgSrc0	is the first source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
mask	is the predicate to specify enabled channels (optional - default to on).

Definition at line 340 of file memory.hpp.

◆ raw_send() [4/4]

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.

Generates a send or sendc instruction for the message gateway.

Parameters

msgSrc0	is the first source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
execSize	is the execution size, which must be a compile time constant.
sfid	is the shared function ID, which must be a compile time constant.
numSrc0	is the number of GRFs for source-0, which must be a compile time constant.
isEOT	is the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
mask	is the predicate to specify enabled channels (optional - default to on).

Definition at line 311 of file memory.hpp.

◆ raw_sends() [1/4]

template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t numDst, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2, typename T3 , int n3>

__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,
		sycl::ext::intel::esimd::simd_mask< execSize >	mask = `1`
	)

Raw sends.

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

Template Parameters

execSize	is the execution size.
sfid	is the shared function ID.
numSrc0	is the number of GRFs for source-0.
numSrc1	is the number of GRFs for source-1.
numDst	is the number of GRFs for destination.
isEOT	is the flag that indicates whether this is an EOT message (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used (optional - default to 0).

Parameters

msgDst	is the old value of the destination operand.
msgSrc0	is the first source operand of send message.
msgSrc1	is the second source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
mask	is the predicate to specify enabled channels (optional - default to on).

Returns: the vector value read from memory.

Definition at line 114 of file memory.hpp.

◆ raw_sends() [2/4]

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.

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

Parameters

msgDst	is the old value of the destination operand.
msgSrc0	is the first source operand of send message.
msgSrc1	is the second source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
execSize	is the execution size, which must be a compile time constant.
sfid	is the shared function ID, which must be a compile time constant.
numSrc0	is the number of GRFs for source-0, which must be a compile time constant.
numSrc1	is the number of GRFs for source-1, which must be a compile constant.
numDst	is the number of GRFs for destination, which must be a compile time constant.
isEOT	is the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
mask	is the predicate to specify enabled channels (optional - default to on).

Returns: the vector value read from memory.

Definition at line 69 of file memory.hpp.

◆ raw_sends() [3/4]

template<uint8_t execSize, uint8_t sfid, uint8_t numSrc0, uint8_t numSrc1, uint8_t isEOT = 0, uint8_t isSendc = 0, typename T1 , int n1, typename T2 , int n2>

__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,
		sycl::ext::intel::esimd::simd_mask< execSize >	mask = `1`
	)

Raw sends.

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

Template Parameters

execSize	is the execution size.
sfid	is the shared function ID.
numSrc0	is the number of GRFs for source-0.
numSrc1	is the number of GRFs for source-1.
isEOT	is the flag that indicates whether this is an EOT message (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used (optional - default to 0).

Parameters

msgSrc0	is the first source operand of send message.
msgSrc1	is the second source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
mask	is the predicate to specify enabled channels (optional - default to on).

Definition at line 273 of file memory.hpp.

◆ raw_sends() [4/4]

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.

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

Parameters

msgSrc0	is the first source operand of send message.
msgSrc1	is the second source operand of send message.
exDesc	is the extended message descriptor.
msgDesc	is the message descriptor.
execSize	is the execution size, which must be a compile time constant.
sfid	is the shared function ID, which must be a compile time constant.
numSrc0	is the number of GRFs for source-0, which must be a compile time constant.
numSrc1	is the number of GRFs for source-1, which must be a compile time constant.
isEOT	is the flag that indicates whether this is an EOT message, which must be a compile time constant (optional - default to 0).
isSendc	is the flag that indicates whether sendc should be used, which must be a compile time constant (optional - default to 0).
mask	is the predicate to specify enabled channels (optional - default to on).

Definition at line 236 of file memory.hpp.

Functions

Detailed Description

Function Documentation

◆ raw_send() [1/4]

◆ raw_send() [2/4]

◆ raw_send() [3/4]

◆ raw_send() [4/4]

◆ raw_sends() [1/4]

◆ raw_sends() [2/4]

◆ raw_sends() [3/4]

◆ raw_sends() [4/4]