GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference#

Composable Kernel: ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

#include <gemm_kernel.hpp>

Inheritance diagram for ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >:

Classes
struct	GemmKernelArgs

struct	SplitKBatchOffset

Public Types
using	TilePartitioner = remove_cvref_t< TilePartitioner_ >

using	GemmPipeline = remove_cvref_t< GemmPipeline_ >

using	EpiloguePipeline = remove_cvref_t< EpiloguePipeline_ >

using	ALayout = remove_cvref_t< typename GemmPipeline::ALayout >

using	BLayout = remove_cvref_t< typename GemmPipeline::BLayout >

using	CLayout = remove_cvref_t< typename GemmPipeline::CLayout >

using	ADataType = remove_cvref_t< typename GemmPipeline::ADataType >

using	BDataType = remove_cvref_t< typename GemmPipeline::BDataType >

using	CDataType = remove_cvref_t< typename EpiloguePipeline::ODataType >

Public Member Functions
CK_TILE_DEVICE void	operator() (GemmKernelArgs kargs) const

Static Public Member Functions
static constexpr CK_TILE_HOST auto	GridSize (index_t M, index_t N, index_t KBatch)

static constexpr CK_TILE_HOST auto	BlockSize ()

static constexpr CK_TILE_HOST GemmKernelArgs	MakeKernelArgs (const GemmHostArgs &hostArgs)

static constexpr CK_TILE_HOST_DEVICE index_t	GetSmemSize ()

static CK_TILE_HOST bool	IsSupportedArgument (const GemmKernelArgs &kargs)

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
static CK_TILE_DEVICE auto	MakeGemmTensorViews (const ADataType a_ptr, const BDataType b_ptr, CDataType *c_ptr, const GemmKernelArgs &kargs, const SplitKBatchOffset &splitk_batch_offset)

template<typename TensorView >
static CK_TILE_DEVICE auto	MakeGemmPadViews (const TensorView &views)

template<typename PadView >
static CK_TILE_DEVICE auto	MakeGemmTileWindows (const PadView &views, const index_t i_m, const index_t i_n)

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
static CK_TILE_DEVICE void	RunGemm (const ADataType a_ptr, const BDataType b_ptr, CDataType c_ptr, void smem_ptr, const GemmKernelArgs &kargs, const SplitKBatchOffset &splitk_batch_offset, const index_t block_idx_m, const index_t block_idx_n)
	Runs single GEMM problem cooperatively by whole workgroup. More...

Static Public Attributes
static constexpr index_t	KernelBlockSize = GemmPipeline::BlockSize

static constexpr auto	I0 = number<0>()

static constexpr auto	I1 = number<1>()

static constexpr auto	I2 = number<2>()

Member Typedef Documentation

◆ ADataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ADataType = remove_cvref_t<typename GemmPipeline::ADataType>

◆ ALayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::ALayout = remove_cvref_t<typename GemmPipeline::ALayout>

◆ BDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BDataType = remove_cvref_t<typename GemmPipeline::BDataType>

◆ BLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BLayout = remove_cvref_t<typename GemmPipeline::BLayout>

◆ CDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::CDataType = remove_cvref_t<typename EpiloguePipeline::ODataType>

◆ CLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::CLayout = remove_cvref_t<typename GemmPipeline::CLayout>

◆ EpiloguePipeline

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::EpiloguePipeline = remove_cvref_t<EpiloguePipeline_>

◆ GemmPipeline

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GemmPipeline = remove_cvref_t<GemmPipeline_>

◆ TilePartitioner

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

using ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::TilePartitioner = remove_cvref_t<TilePartitioner_>

Member Function Documentation

◆ BlockSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

static constexpr CK_TILE_HOST auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::BlockSize ( )

inlinestaticconstexpr

◆ GetSmemSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

static constexpr CK_TILE_HOST_DEVICE index_t ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GetSmemSize ( )

inlinestaticconstexpr

◆ GridSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

static constexpr CK_TILE_HOST auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::GridSize	(	index_t	M,
		index_t	N,
		index_t	KBatch
	)

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

static CK_TILE_HOST bool ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::IsSupportedArgument ( const GemmKernelArgs & kargs )

inlinestatic

◆ MakeGemmPadViews()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

template<typename TensorView >

static CK_TILE_DEVICE auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeGemmPadViews ( const TensorView & views )

inlinestatic

◆ MakeGemmTensorViews()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>

static CK_TILE_DEVICE auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeGemmTensorViews	(	const ADataType *	a_ptr,
		const BDataType *	b_ptr,
		CDataType *	c_ptr,
		const GemmKernelArgs &	kargs,
		const SplitKBatchOffset &	splitk_batch_offset
	)

inlinestatic

◆ MakeGemmTileWindows()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

template<typename PadView >

static CK_TILE_DEVICE auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeGemmTileWindows	(	const PadView &	views,
		const index_t	i_m,
		const index_t	i_n
	)

inlinestatic

◆ MakeKernelArgs()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

static constexpr CK_TILE_HOST GemmKernelArgs ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::MakeKernelArgs ( const GemmHostArgs & hostArgs )

inlinestaticconstexpr

◆ operator()()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

CK_TILE_DEVICE void ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::operator() ( GemmKernelArgs kargs ) const

inline

◆ RunGemm()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>

static CK_TILE_DEVICE void ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::RunGemm	(	const ADataType *	a_ptr,
		const BDataType *	b_ptr,
		CDataType *	c_ptr,
		void *	smem_ptr,
		const GemmKernelArgs &	kargs,
		const SplitKBatchOffset &	splitk_batch_offset,
		const index_t	block_idx_m,
		const index_t	block_idx_n
	)

inlinestatic

Runs single GEMM problem cooperatively by whole workgroup.

Parameters

a_ptr	input A pointer
b_ptr	input B pointer
c_ptr	output C pointer
kargs	GEMM kernel arguments
block_idx_m	The GEMM's output M dimension tile index processed by this workgroup.
block_idx_n	The GEMM's output N dimension tile index processed by this workgroup.

Template Parameters

DstInMemOp Destination memory operation (default: set).

Member Data Documentation

◆ I0

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

constexpr auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::I0 = number<0>()

staticconstexpr

◆ I1

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

constexpr auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::I1 = number<1>()

staticconstexpr

◆ I2

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

constexpr auto ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::I2 = number<2>()

staticconstexpr

◆ KernelBlockSize

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ >

constexpr index_t ck_tile::GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ >::KernelBlockSize = GemmPipeline::BlockSize

staticconstexpr

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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp

GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference

GemmKernel&lt; TilePartitioner_, GemmPipeline_, EpiloguePipeline_ &gt; Struct Template Reference#

Classes

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ ADataType

◆ ALayout

◆ BDataType

◆ BLayout

◆ CDataType

◆ CLayout

◆ EpiloguePipeline

◆ GemmPipeline

◆ TilePartitioner

Member Function Documentation

◆ BlockSize()

◆ GetSmemSize()

◆ GridSize()

◆ IsSupportedArgument()

◆ MakeGemmPadViews()

◆ MakeGemmTensorViews()

◆ MakeGemmTileWindows()

◆ MakeKernelArgs()

◆ operator()()

◆ RunGemm()

Member Data Documentation

◆ I0

◆ I1

◆ I2

◆ KernelBlockSize

GemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_ > Struct Template Reference#