Functions

◆ gemm_direct()

template<typename T>

BATCHLAS_INTERNAL_API Event batchlas::sycl_gemm::gemm_direct	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		Transpose	transA,
		Transpose	transB )

◆ gemm_tiled()

template<typename T>

BATCHLAS_INTERNAL_API Event batchlas::sycl_gemm::gemm_tiled	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		Transpose	transA,
		Transpose	transB )

◆ gemm_small()

template<typename T>

BATCHLAS_INTERNAL_API Event batchlas::sycl_gemm::gemm_small	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		Transpose	transA,
		Transpose	transB )

◆ gemm_reg()

BATCHLAS_INTERNAL_API Event batchlas::sycl_gemm::gemm_reg	(	Queue &	ctx,
		int	tm,
		int	tn,
		int	tk,
		int	u,
		const MatrixView< float, MatrixFormat::Dense > &	A,
		const MatrixView< float, MatrixFormat::Dense > &	B,
		const MatrixView< float, MatrixFormat::Dense > &	C,
		float	alpha,
		float	beta,
		Transpose	transA,
		Transpose	transB )

◆ gemm_wide()

template<typename T>

BATCHLAS_INTERNAL_API Event batchlas::sycl_gemm::gemm_wide	(	Queue &	ctx,
		int	tm,
		int	tn,
		int	tk,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		Transpose	transA,
		Transpose	transB )

◆ maybe_conj() [1/2]

template<typename T>

T batchlas::sycl_gemm::maybe_conj	(	const T &	value,
		bool	conjugate )

inline

Definition at line 10 of file accessors.hh.

◆ maybe_conj() [2/2]

template<typename T>

std::complex< T > batchlas::sycl_gemm::maybe_conj	(	const std::complex< T > &	value,
		bool	conjugate )

inline

Definition at line 16 of file accessors.hh.

◆ operand_value()

template<typename T>

T batchlas::sycl_gemm::operand_value	(	const T *	ptr,
		int	ld,
		int	batch_offset,
		int	row,
		int	col,
		Transpose	trans )

inline

Definition at line 21 of file accessors.hh.

◆ supports_aligned_packet_loads()

template<typename T, int Width>

bool batchlas::sycl_gemm::supports_aligned_packet_loads	(	const T *	ptr,
		int	ld,
		int	stride )

inline

Definition at line 19 of file load_policies.hh.

◆ packet_load_aligned()

template<typename T, int Width>

sycl::vec< T, Width > batchlas::sycl_gemm::packet_load_aligned	(	const T *	ptr,
		int	offset )

inline

Definition at line 32 of file load_policies.hh.

◆ vec4_ref() [1/2]

template<typename T>

const Vec4< T > & batchlas::sycl_gemm::vec4_ref ( const T * p )

inline

Definition at line 37 of file register_128x128.hh.

◆ vec4_ref() [2/2]

template<typename T>

Vec4< T > & batchlas::sycl_gemm::vec4_ref ( T * p )

inline

Definition at line 42 of file register_128x128.hh.

◆ load4_l2_line()

template<typename T>

Vec4< T > batchlas::sycl_gemm::load4_l2_line ( const T * p )

inline

Definition at line 49 of file register_128x128.hh.

◆ load1_l2_line()

template<typename T>

T batchlas::sycl_gemm::load1_l2_line ( const T * p )

inline

Definition at line 79 of file register_128x128.hh.

◆ can_use_128x128_fast_path()

template<typename T>

bool batchlas::sycl_gemm::can_use_128x128_fast_path	(	const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C )

inline

Definition at line 100 of file register_128x128.hh.

◆ launch_register_128x128_k8_leg()

template<typename T, bool AlignedFastPath, bool StagedEpilogue>

Event batchlas::sycl_gemm::launch_register_128x128_k8_leg	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	kernel_trace_name (*)(KernelVariant) )

Definition at line 344 of file register_128x128.hh.

◆ launch_register_128x128_k8()

template<typename T, bool AlignedFastPath = false>

Event batchlas::sycl_gemm::launch_register_128x128_k8	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	kernel_trace_name (*)(KernelVariant) )

Definition at line 381 of file register_128x128.hh.

◆ can_use_64x64_k16_wide_fast_path()

template<typename T>

bool batchlas::sycl_gemm::can_use_64x64_k16_wide_fast_path	(	const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C )

inline

Definition at line 107 of file register_64x64_k16_wide.hh.

◆ launch_register_64x64_k16_wide()

template<typename T, bool AlignedFastPath = false>

Event batchlas::sycl_gemm::launch_register_64x64_k16_wide	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	kernel_trace_name (*)(KernelVariant) )

Definition at line 132 of file register_64x64_k16_wide.hh.

◆ launch_reg()

template<typename T, RegTile P>

Event batchlas::sycl_gemm::launch_reg	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	trace,
		const char *	trace_aligned = nullptr )

Definition at line 35 of file register_launchers.hh.

◆ static_for_impl()

template<int Count, typename Fn, std::size_t... I>

void batchlas::sycl_gemm::static_for_impl	(	Fn &&	fn,
		std::index_sequence< I... >	)

inline

Definition at line 18 of file register_tiled_common.hh.

◆ static_for()

template<int Count, typename Fn>

void batchlas::sycl_gemm::static_for ( Fn && fn )

inline

Definition at line 23 of file register_tiled_common.hh.

◆ register_ceil_div()

template<typename T>

int batchlas::sycl_gemm::register_ceil_div	(	int	value,
		int	divisor )

inline

Definition at line 43 of file register_tiled_common.hh.

◆ is_contiguous_dense_matrix()

template<typename T>

bool batchlas::sycl_gemm::is_contiguous_dense_matrix ( const MatrixView< T, MatrixFormat::Dense > & M )

inline

Definition at line 76 of file register_tiled_common.hh.

◆ can_use_aligned_nn_fast_path()

template<typename T, int TileM, int TileN, int TileK, int VecA, int VecB>

bool batchlas::sycl_gemm::can_use_aligned_nn_fast_path	(	const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C )

inline

Definition at line 81 of file register_tiled_common.hh.

◆ launch_register_tiled()

template<typename T, int TileM, int TileN, int TileK, int ThreadTileRows, int ThreadTileCols = ThreadTileRows, int VecA = 4, int VecB = 4, int UnrollK = 1, int Stages = 1, bool AlignedFastPath = false, Transpose OpA = Transpose::NoTrans, Transpose OpB = Transpose::NoTrans>

Event batchlas::sycl_gemm::launch_register_tiled	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	kernel_trace_name )

Definition at line 114 of file register_tiled_common.hh.

◆ wide_trans_matches()

template<typename T>

bool batchlas::sycl_gemm::wide_trans_matches	(	Transpose	requested,
		Transpose	instantiated )

inline

Definition at line 80 of file register_wide_transposed.hh.

◆ launch_wide_transposed()

template<typename T, WideTile P>

Event batchlas::sycl_gemm::launch_wide_transposed	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	trace )

Definition at line 94 of file register_wide_transposed.hh.

◆ launch_tiled_general()

template<typename T, int Tile, Transpose OpA, Transpose OpB>

Event batchlas::sycl_gemm::launch_tiled_general	(	Queue &	ctx,
		const MatrixView< T, MatrixFormat::Dense > &	A,
		const MatrixView< T, MatrixFormat::Dense > &	B,
		const MatrixView< T, MatrixFormat::Dense > &	C,
		T	alpha,
		T	beta,
		const char *	kernel_trace_name (*)(KernelVariant) )

Definition at line 15 of file tiled_general.hh.

Types

◆ Vec4

template<typename T>

using batchlas::sycl_gemm::Vec4 = T __attribute__((ext_vector_type(4)))

Definition at line 33 of file vec4.hh.

Enumerations

◆ KernelVariant

enum class batchlas::sycl_gemm::KernelVariant

strong

Enumerator
Direct
Tiled16
Tiled32x32Register
Tiled64x64Register
Tiled64x64RegisterK16
Tiled64x64RegisterK16TN
Tiled64x64RegisterK16NT
Tiled64x64RegisterK16TT
Tiled128x32RegisterK16
Tiled128x32RegisterK16TN
Tiled128x32RegisterK16NT
Tiled128x32RegisterK16TT
Tiled128x32RegisterK32TN
Tiled128x32RegisterK32NT
Tiled128x32RegisterK32TT
Tiled128x64RegisterK16TN
Tiled128x64RegisterK16NT
Tiled128x64RegisterK16TT
Tiled128x32RegisterK32S2U1Aligned
Tiled128x32RegisterK32S2U1Generic
Tiled128x64RegisterK32Large
Tiled128x64RegisterK32LargeU2
Tiled128x128RegisterK8
Tiled64x64RegisterK16Wide
Tiled64x64RegisterK16WideCN
Tiled64x64RegisterK16WideNC
Tiled128x32RegisterK16WideNC
Tiled32x128RegisterK16WideCN
Tiled32x128RegisterK16
Tiled32x128RegisterK16TN
Tiled32x128RegisterK16TT
Tiled32x32RegisterK16Wide
Tiled16x16RegisterK16Wide

Definition at line 12 of file gemm_kernels.hh.

Variables

◆ supports_packet4_v

template<typename T>

bool batchlas::sycl_gemm::supports_packet4_v = std::is_same_v<T, float>

inlineconstexpr

Definition at line 10 of file load_policies.hh.

◆ supports_packet_v

template<typename T, int Width>

bool batchlas::sycl_gemm::supports_packet_v = false

inlineconstexpr

Definition at line 13 of file load_policies.hh.

◆ supports_packet_v< T, 4 >

template<typename T>

bool batchlas::sycl_gemm::supports_packet_v< T, 4 > = supports_packet4_v<T>

inlineconstexpr

Definition at line 16 of file load_policies.hh.

◆ kStagedEpilogueMaxK

int batchlas::sycl_gemm::kStagedEpilogueMaxK = 64

inlineconstexpr

Definition at line 75 of file register_128x128.hh.

Namespaces
namespace	wide_scalar
namespace	wide_transposed

Classes
struct	Gemm128x128Body
class	GemmRegister128x128Kernel
class	GemmRegister64x64K16WideKernel
class	GemmRegisterTiledKernel
class	GemmTiledGeneralKernel
class	GemmWideTransposedKernel
struct	LinearEpilogue
struct	OperandAccessor
struct	RegisterTilePolicy
struct	RegTile
struct	WideTile

Functions

◆ gemm_direct()

◆ gemm_tiled()

◆ gemm_small()

◆ gemm_reg()

◆ gemm_wide()

◆ maybe_conj() [1/2]

◆ maybe_conj() [2/2]

◆ operand_value()

◆ supports_aligned_packet_loads()

◆ packet_load_aligned()

◆ vec4_ref() [1/2]

◆ vec4_ref() [2/2]

◆ load4_l2_line()

◆ load1_l2_line()

◆ can_use_128x128_fast_path()

◆ launch_register_128x128_k8_leg()

◆ launch_register_128x128_k8()

◆ can_use_64x64_k16_wide_fast_path()

◆ launch_register_64x64_k16_wide()

◆ launch_reg()

◆ static_for_impl()

◆ static_for()

◆ register_ceil_div()

◆ is_contiguous_dense_matrix()

◆ can_use_aligned_nn_fast_path()

◆ launch_register_tiled()

◆ wide_trans_matches()

◆ launch_wide_transposed()

◆ launch_tiled_general()

Types

◆ Vec4

Enumerations

◆ KernelVariant

Variables

◆ supports_packet4_v

◆ supports_packet_v

◆ supports_packet_v< T, 4 >

◆ kStagedEpilogueMaxK

Namespaces

Classes