TensorContractionKernel is a template class that provides Tensor -Tensor contraction operation. More...

#include <TensorContractionSycl.h>

Classes
struct	MemHolder
	MemHolder this is a place holder struct for creating memory hierarchy in SYCL. Inside SYCL kernel it is not allowed to have dynamic memory allocation. While the local memory is created outside of the kernel and passed to the kernel as an accessor, the private memory can only allowed to be allocated statically. Since we are abstracting the TiledMemory for both local and private memory, the MemHolder structs is used as a helper to abstract out different type of memory needed when local/no_local memory computation is called. More...

struct	MemHolder< contraction_type::no_local, MemSize >
	specialization of memHolder class when no local memory kernel is used. More...

struct	TiledMemory
	TiledMemory: contains required memory pointer for loading each tile of the TensorContraction panel from global memory to local/private memory when local/no_local algorithm used. More...

Public Types
typedef BlockProperties< is_lhs_transposed, false, input_mapper_properties::is_lhs_matrix &&Vectorizable, PacketReturnType >	LHSBlockProperties

typedef cl::sycl::multi_ptr< OutScalar, cl::sycl::access::address_space::local_space >	local_ptr

typedef Eigen::TensorSycl::internal::Vectorise< OutScalar, Eigen::SyclDevice, Vectorizable >::PacketReturnType	PacketReturnType

typedef OutScalar *	private_ptr

typedef BlockProperties< is_rhs_transposed, true, input_mapper_properties::is_rhs_matrix &&Vectorizable, PacketReturnType >	RHSBlockProperties

typedef cl::sycl::accessor< OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local >	Scratch

typedef typename ::Eigen::internal::conditional< contraction_tp==contraction_type::local, local_ptr, private_ptr >::type	tile_ptr

Public Member Functions
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	compute_block_per_tile (OutScalar lhs_block_ptr, OutScalar rhs_block_ptr, PacketReturnType *privateRes)

template<bool is_internal_block, typename OutPtr >
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	compute_panel (const cl::sycl::nd_item< 1 > &itemID, ThreadProperties< StorageIndex > &thread_properties, OutPtr out_ptr)

template<bool is_internal_block>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	compute_tile_per_panel (const cl::sycl::nd_item< 1 > &itemID, ThreadProperties< StorageIndex > &thread_properties, TiledMemory &tiled_input_block, PacketReturnType *privateRes, bool &db_offset)

template<typename InputBlockProperties , bool is_internal_block, typename Input , typename PrivateReg , contraction_type contract_tp = contraction_tp>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< contract_tp==contraction_type::no_local >::type	extract_block (const Input &inpt, PrivateReg private_ptr, const std::pair< StorageIndex, StorageIndex > &, const StorageIndex &ncOffset, const StorageIndex cOffset)

template<typename InputBlockProperties , bool is_internal_block, typename Input , typename Local , contraction_type contract_tp = contraction_tp>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< contract_tp==contraction_type::local >::type	extract_block (const Input &inpt, Local local_ptr, const std::pair< StorageIndex, StorageIndex > &local_index, const StorageIndex &ncOffset, const StorageIndex cOffset)

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	operator() (cl::sycl::nd_item< 1 > itemID)

template<bool is_internal_block, StorageIndex PrivateNStride, typename OutPtr >
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void	store (OutPtr out_ptr, PacketReturnType privateRes, StorageIndex mGlobalOffset, StorageIndex nGlobalOffset)

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	TensorContractionKernel (Scratch scratch_, const LhsMapper lhs_, const RhsMapper rhs_, OutAccessor out_res_, const StorageIndex groupSizeM_, const StorageIndex groupSizeN_, const StorageIndex numTiles_, const TripleDim triple_dim_)

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE	TensorContractionKernel (Scratch scratch_, const LhsMapper lhs_, const RhsMapper rhs_, OutAccessor out_res_, const StorageIndex groupSizeM_, const StorageIndex numTiles_, const TripleDim triple_dim_)

Static Public Member Functions
template<typename InputBlockProperties , StorageIndex TileSizeDimNC>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::pair< StorageIndex, StorageIndex >	local_id_extract (const StorageIndex &linearLocalThreadId)

template<bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< db &&ctp==contraction_type::local >::type	sync_mem (const cl::sycl::nd_item< 1 > &, bool &db_offset) noexcept

template<bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<!db &&ctp==contraction_type::local >::type	sync_mem (const cl::sycl::nd_item< 1 > &itemID, bool &) noexcept

template<contraction_type ctp = contraction_tp>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< ctp==contraction_type::no_local >::type	sync_mem (const cl::sycl::nd_item< 1 > &, bool &) noexcept

template<bool need_sync, contraction_type ctp = contraction_tp>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< need_sync &&ctp==contraction_type::no_local >::type	sync_thread (const cl::sycl::nd_item< 1 > &) noexcept

template<bool need_sync, contraction_type ctp = contraction_tp>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if< need_sync &&ctp==contraction_type::local >::type	sync_thread (const cl::sycl::nd_item< 1 > &itemID)

template<bool need_sync>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ::Eigen::internal::enable_if<!need_sync >::type	sync_thread (const cl::sycl::nd_item< 1 > &)

Public Attributes
const StorageIndex	groupSizeM

const StorageIndex	groupSizeN

const LhsMapper	lhs

const StorageIndex	numTiles

OutAccessor	out_res

const RhsMapper	rhs

Scratch	scratch

const TripleDim	triple_dim

Static Public Attributes
static EIGEN_CONSTEXPR bool	is_lhs_transposed

static EIGEN_CONSTEXPR bool	is_rhs_transposed

static EIGEN_CONSTEXPR StorageIndex	LocalOffset = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN

static EIGEN_CONSTEXPR StorageIndex	LSDL

static EIGEN_CONSTEXPR StorageIndex	LSDR

static EIGEN_CONSTEXPR StorageIndex	NStride

static EIGEN_CONSTEXPR int	PacketSize

Detailed Description

template<typename OutScalar, typename LhsScalar, typename RhsScalar, typename OutAccessor, typename LhsMapper, typename RhsMapper, typename StorageIndex, typename Properties, typename TripleDim, bool Vectorizable, typename input_mapper_properties, bool IsFinal, contraction_type contraction_tp>
class Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >

TensorContractionKernel is a template class that provides Tensor -Tensor contraction operation.

Template Parameters

OutScalar	determines the output scalar type
LhsScalar	determines the left-hand-side scalar type
RhsScalar	determines the right-hand-side scalar type
OutAccessor	determines the sycl accessor type for out put (please see the sycl-1.2.1 specification (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for accessor definition)
LhsMapper	determines the tensor contraction mapper type for left-hand-side matrix
RhsMapper	determines the tensor contraction mapper type for right-hand-side matrix
StorageIndex	determines the StorageIndex Type
Properties	determines the Contraction Panel properties
TripleDim	determines the M, K, N dimensions for the flatten tensors in order to treat them as a matrix
Vectorizable	determines whether or not the vectorization is enabled for the Eigen expression.
input_mapper_properties	: determine if the input tensors are matrix. If they are matrix, special memory access is used to guarantee that always the memory access are coalesced.

IsFinal : determine if this is the final kernel. If so, the result will be written in a final output. Otherwise, the result of contraction will be written iin a temporary buffer. This is the case when Tall/Skinny contraction is used. So in this case, a final reduction step is required to compute final output.

Template Parameters

contraction_tp it is an enum value representing whether the local memroy/no local memory implementation of the algorithm to be used

Parameters

scratch	local memory containing tiles of LHS and RHS tensors for each work-group
lhs	determines the left-hand-side flattened tensor (tensor mapper)
rhs	determines the right-hand-side flattened tensor (tensor mapper)
out_res	determines the output tensor containing the contraction result
groupSizeM	a logical number determining the number of work-group for m dimension
groupSizeN	a logical number determining the number of work-group for n dimension
numTiles	determines total number of tiles on the k dimension
TripleDim	determines the M, K, N dimensions for the flatten tensors in order to treat them as a matrix

Definition at line 453 of file TensorContractionSycl.h.

Member Typedef Documentation

◆ LHSBlockProperties

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef BlockProperties<is_lhs_transposed, false, input_mapper_properties::is_lhs_matrix && Vectorizable, PacketReturnType> Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::LHSBlockProperties

Definition at line 466 of file TensorContractionSycl.h.

◆ local_ptr

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef cl::sycl::multi_ptr<OutScalar, cl::sycl::access::address_space::local_space> Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::local_ptr

Definition at line 476 of file TensorContractionSycl.h.

◆ PacketReturnType

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketReturnType Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::PacketReturnType

Definition at line 456 of file TensorContractionSycl.h.

◆ private_ptr

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef OutScalar* Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::private_ptr

Definition at line 477 of file TensorContractionSycl.h.

◆ RHSBlockProperties

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef BlockProperties<is_rhs_transposed, true, input_mapper_properties::is_rhs_matrix && Vectorizable, PacketReturnType> Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::RHSBlockProperties

Definition at line 470 of file TensorContractionSycl.h.

◆ Scratch

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::Scratch

Definition at line 475 of file TensorContractionSycl.h.

◆ tile_ptr

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

typedef typename ::Eigen::internal::conditional<contraction_tp == contraction_type::local, local_ptr, private_ptr>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::tile_ptr

Definition at line 480 of file TensorContractionSycl.h.

Constructor & Destructor Documentation

◆ TensorContractionKernel() [1/2]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::TensorContractionKernel	(	Scratch	scratch_,
		const LhsMapper	lhs_,
		const RhsMapper	rhs_,
		OutAccessor	out_res_,
		const StorageIndex	groupSizeM_,
		const StorageIndex	groupSizeN_,
		const StorageIndex	numTiles_,
		const TripleDim	triple_dim_
	)

inline

Definition at line 577 of file TensorContractionSycl.h.

◆ TensorContractionKernel() [2/2]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::TensorContractionKernel	(	Scratch	scratch_,
		const LhsMapper	lhs_,
		const RhsMapper	rhs_,
		OutAccessor	out_res_,
		const StorageIndex	groupSizeM_,
		const StorageIndex	numTiles_,
		const TripleDim	triple_dim_
	)

inline

Definition at line 592 of file TensorContractionSycl.h.

Member Function Documentation

◆ compute_block_per_tile()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::compute_block_per_tile	(	OutScalar *	lhs_block_ptr,
		OutScalar *	rhs_block_ptr,
		PacketReturnType *	privateRes
	)

inline

Definition at line 637 of file TensorContractionSycl.h.

◆ compute_panel()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool is_internal_block, typename OutPtr >

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::compute_panel	(	const cl::sycl::nd_item< 1 > &	itemID,
		ThreadProperties< StorageIndex > &	thread_properties,
		OutPtr	out_ptr
	)

inline

Definition at line 871 of file TensorContractionSycl.h.

◆ compute_tile_per_panel()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool is_internal_block>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::compute_tile_per_panel	(	const cl::sycl::nd_item< 1 > &	itemID,
		ThreadProperties< StorageIndex > &	thread_properties,
		TiledMemory &	tiled_input_block,
		PacketReturnType *	privateRes,
		bool &	db_offset
	)

inline

Definition at line 832 of file TensorContractionSycl.h.

◆ extract_block() [1/2]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<typename InputBlockProperties , bool is_internal_block, typename Input , typename PrivateReg , contraction_type contract_tp = contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<contract_tp == contraction_type::no_local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::extract_block	(	const Input &	inpt,
		PrivateReg	private_ptr,
		const std::pair< StorageIndex, StorageIndex > &	,
		const StorageIndex &	ncOffset,
		const StorageIndex	cOffset
	)

inline

Definition at line 714 of file TensorContractionSycl.h.

◆ extract_block() [2/2]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<typename InputBlockProperties , bool is_internal_block, typename Input , typename Local , contraction_type contract_tp = contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<contract_tp == contraction_type::local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::extract_block	(	const Input &	inpt,
		Local	local_ptr,
		const std::pair< StorageIndex, StorageIndex > &	local_index,
		const StorageIndex &	ncOffset,
		const StorageIndex	cOffset
	)

inline

Definition at line 898 of file TensorContractionSycl.h.

◆ local_id_extract()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<typename InputBlockProperties , StorageIndex TileSizeDimNC>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::pair<StorageIndex, StorageIndex> Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::local_id_extract ( const StorageIndex & linearLocalThreadId )

inlinestatic

Definition at line 771 of file TensorContractionSycl.h.

◆ operator()()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::operator() ( cl::sycl::nd_item< 1 > itemID )

inline

Definition at line 599 of file TensorContractionSycl.h.

◆ store()

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool is_internal_block, StorageIndex PrivateNStride, typename OutPtr >

EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::store	(	OutPtr *	out_ptr,
		PacketReturnType *	privateRes,
		StorageIndex	mGlobalOffset,
		StorageIndex	nGlobalOffset
	)

inline

Definition at line 662 of file TensorContractionSycl.h.

◆ sync_mem() [1/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<db && ctp == contraction_type::local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_mem	(	const cl::sycl::nd_item< 1 > &	,
		bool &	db_offset
	)

inlinestaticnoexcept

Definition at line 787 of file TensorContractionSycl.h.

◆ sync_mem() [2/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<!db && ctp == contraction_type::local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_mem	(	const cl::sycl::nd_item< 1 > &	itemID,
		bool &
	)

inlinestaticnoexcept

Definition at line 794 of file TensorContractionSycl.h.

◆ sync_mem() [3/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<contraction_type ctp = contraction_tp>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<ctp == contraction_type::no_local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_mem	(	const cl::sycl::nd_item< 1 > &	,
		bool &
	)

inlinestaticnoexcept

Definition at line 801 of file TensorContractionSycl.h.

◆ sync_thread() [1/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool need_sync, contraction_type ctp = contraction_tp>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<need_sync && ctp == contraction_type::no_local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_thread ( const cl::sycl::nd_item< 1 > & )

inlinestaticnoexcept

Definition at line 808 of file TensorContractionSycl.h.

◆ sync_thread() [2/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool need_sync, contraction_type ctp = contraction_tp>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE typename ::Eigen::internal::enable_if<need_sync && ctp == contraction_type::local>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_thread ( const cl::sycl::nd_item< 1 > & itemID )

inlinestatic

Definition at line 822 of file TensorContractionSycl.h.

◆ sync_thread() [3/3]

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

template<bool need_sync>

static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ::Eigen::internal::enable_if<!need_sync>::type Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::sync_thread ( const cl::sycl::nd_item< 1 > & )

inlinestatic

Definition at line 826 of file TensorContractionSycl.h.

Member Data Documentation

◆ groupSizeM

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::groupSizeM

Definition at line 572 of file TensorContractionSycl.h.

◆ groupSizeN

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::groupSizeN

Definition at line 573 of file TensorContractionSycl.h.

◆ is_lhs_transposed

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR bool Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::is_lhs_transposed

static

Initial value:

=

!::Eigen::internal::TensorContractionInputMapperTrait<LhsMapper>::inner_dim_contiguous

Definition at line 459 of file TensorContractionSycl.h.

◆ is_rhs_transposed

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR bool Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::is_rhs_transposed

static

Initial value:

=

!::Eigen::internal::TensorContractionInputMapperTrait<RhsMapper>::inner_dim_contiguous

Definition at line 461 of file TensorContractionSycl.h.

◆ lhs

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const LhsMapper Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::lhs

Definition at line 569 of file TensorContractionSycl.h.

◆ LocalOffset

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::LocalOffset = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN

static

Definition at line 487 of file TensorContractionSycl.h.

◆ LSDL

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::LSDL

static

Initial value:

= contraction_tp == contraction_type::local
                                                 ? Properties::TileSizeDimM + Properties::BC
                                                 : Properties::WorkLoadPerThreadM

Definition at line 481 of file TensorContractionSycl.h.

◆ LSDR

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::LSDR

static

Initial value:

= contraction_tp == contraction_type::local
                                                 ? Properties::TileSizeDimN + Properties::BC
                                                 : Properties::WorkLoadPerThreadN

Definition at line 484 of file TensorContractionSycl.h.

◆ NStride

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::NStride

static

Initial value:

=

contraction_tp == contraction_type::local ? Properties::WorkLoadPerThreadN : RHSBlockProperties::nc_stride

Definition at line 472 of file TensorContractionSycl.h.

◆ numTiles

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const StorageIndex Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::numTiles

Definition at line 574 of file TensorContractionSycl.h.

◆ out_res

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

OutAccessor Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::out_res

Definition at line 571 of file TensorContractionSycl.h.

◆ PacketSize

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

EIGEN_CONSTEXPR int Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::PacketSize

static

Initial value:

=

Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketSize

Definition at line 457 of file TensorContractionSycl.h.

◆ rhs

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const RhsMapper Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::rhs

Definition at line 570 of file TensorContractionSycl.h.

◆ scratch

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

Scratch Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::scratch

Definition at line 568 of file TensorContractionSycl.h.

◆ triple_dim

template<typename OutScalar , typename LhsScalar , typename RhsScalar , typename OutAccessor , typename LhsMapper , typename RhsMapper , typename StorageIndex , typename Properties , typename TripleDim , bool Vectorizable, typename input_mapper_properties , bool IsFinal, contraction_type contraction_tp>

const TripleDim Eigen::TensorSycl::internal::TensorContractionKernel< OutScalar, LhsScalar, RhsScalar, OutAccessor, LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim, Vectorizable, input_mapper_properties, IsFinal, contraction_tp >::triple_dim

Definition at line 575 of file TensorContractionSycl.h.

The documentation for this class was generated from the following file:

TensorContractionSycl.h

Classes

Public Types

Public Member Functions

Static Public Member Functions

Public Attributes

Static Public Attributes

Detailed Description

Member Typedef Documentation

◆ LHSBlockProperties

◆ local_ptr

◆ PacketReturnType

◆ private_ptr

◆ RHSBlockProperties

◆ Scratch

◆ tile_ptr

Constructor & Destructor Documentation

◆ TensorContractionKernel() [1/2]

◆ TensorContractionKernel() [2/2]

Member Function Documentation

◆ compute_block_per_tile()

◆ compute_panel()

◆ compute_tile_per_panel()

◆ extract_block() [1/2]

◆ extract_block() [2/2]

◆ local_id_extract()

◆ operator()()

◆ store()

◆ sync_mem() [1/3]

◆ sync_mem() [2/3]

◆ sync_mem() [3/3]

◆ sync_thread() [1/3]

◆ sync_thread() [2/3]

◆ sync_thread() [3/3]

Member Data Documentation

◆ groupSizeM

◆ groupSizeN

◆ is_lhs_transposed

◆ is_rhs_transposed

◆ lhs

◆ LocalOffset

◆ LSDL

◆ LSDR

◆ NStride

◆ numTiles

◆ out_res

◆ PacketSize

◆ rhs

◆ scratch

◆ triple_dim