QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ > Struct Template Reference

QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ > Struct Template Reference#

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

#include <gemm_quant_kernel.hpp>

Classes
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	AQLayout = remove_cvref_t< typename detail::get_aq_layout_or< GemmPipeline, typename GemmPipeline::ALayout >::type >

using	BQLayout = remove_cvref_t< typename detail::get_bq_layout_or< GemmPipeline, typename GemmPipeline::BLayout >::type >

using	ADataType = remove_cvref_t< typename GemmPipeline::ADataType >

using	BDataType = remove_cvref_t< typename GemmPipeline::BDataType >

using	CDataType = remove_cvref_t< typename EpiloguePipeline::ODataType >

using	AccDataType = remove_cvref_t< typename EpiloguePipeline::AccDataType >

using	AQDataType = remove_cvref_t< typename detail::get_aq_data_type_or< GemmPipeline, AccDataType >::type >

using	BQDataType = remove_cvref_t< typename detail::get_bq_data_type_or< GemmPipeline, AccDataType >::type >

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

Static Public Member Functions
static CK_TILE_HOST const std::string	GetName ()

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 QuantGemmKernelArgs	MakeKernelArgs (const QuantGemmHostArgs &hostArgs)

static constexpr CK_TILE_HOST_DEVICE index_t	GetSmemSize ()

static CK_TILE_HOST bool	IsSupportedArgument (const QuantGemmKernelArgs &kargs)

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
static CK_TILE_DEVICE auto	MakeGemmTensorViews (const ADataType a_ptr, const BDataType b_ptr, const AQDataType aq_ptr, const BQDataType bq_ptr, CDataType *c_ptr, const QuantGemmKernelArgs &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, const AQDataType aq_ptr, const BQDataType bq_ptr, CDataType c_ptr, void smem_ptr_0, const QuantGemmKernelArgs &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...

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
static CK_TILE_DEVICE void	RunGemm2LDS (const ADataType a_ptr, const BDataType b_ptr, const AQDataType aq_ptr, const BQDataType bq_ptr, CDataType c_ptr, void smem_ptr_0, void *smem_ptr_1, const QuantGemmKernelArgs &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	kBlockSize = GemmPipeline::BlockSize

static constexpr bool	PreshuffleQuant

static constexpr bool	PreshuffleB = detail::is_preshuffleB_enabled<GemmPipeline_>::value

static constexpr auto	I0 = number<0>()

static constexpr auto	I1 = number<1>()

static constexpr auto	I2 = number<2>()

static constexpr auto	I3 = number<3>()

static constexpr auto	I4 = number<4>()

static constexpr auto	kQuantType = QuantType_

Member Typedef Documentation

◆ AccDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

using ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::AccDataType = remove_cvref_t<typename EpiloguePipeline::AccDataType>

◆ ADataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ ALayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ AQDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

using ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::AQDataType = remove_cvref_t<typename detail::get_aq_data_type_or<GemmPipeline, AccDataType>::type>

◆ AQLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

using ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::AQLayout = remove_cvref_t< typename detail::get_aq_layout_or<GemmPipeline, typename GemmPipeline::ALayout>::type>

◆ BDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ BLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ BQDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

using ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::BQDataType = remove_cvref_t<typename detail::get_bq_data_type_or<GemmPipeline, AccDataType>::type>

◆ BQLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

using ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::BQLayout = remove_cvref_t< typename detail::get_bq_layout_or<GemmPipeline, typename GemmPipeline::BLayout>::type>

◆ CDataType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ CLayout

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ EpiloguePipeline

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ GemmPipeline

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

◆ TilePartitioner

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

Member Function Documentation

◆ BlockSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

inlinestaticconstexpr

◆ GetName()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

static CK_TILE_HOST const std::string ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::GetName ( )

inlinestatic

◆ GetSmemSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

inlinestaticconstexpr

◆ GridSize()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

inlinestaticconstexpr

◆ IsSupportedArgument()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

static CK_TILE_HOST bool ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::IsSupportedArgument ( const QuantGemmKernelArgs & kargs )

inlinestatic

◆ MakeGemmPadViews()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

template<typename TensorView >

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

inlinestatic

◆ MakeGemmTensorViews()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>

static CK_TILE_DEVICE auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeGemmTensorViews	(	const ADataType *	a_ptr,
		const BDataType *	b_ptr,
		const AQDataType *	aq_ptr,
		const BQDataType *	bq_ptr,
		CDataType *	c_ptr,
		const QuantGemmKernelArgs &	kargs,
		const SplitKBatchOffset &	splitk_batch_offset
	)

inlinestatic

◆ MakeGemmTileWindows()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

template<typename PadView >

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

inlinestatic

◆ MakeKernelArgs()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

static constexpr CK_TILE_HOST QuantGemmKernelArgs ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeKernelArgs ( const QuantGemmHostArgs & hostArgs )

inlinestaticconstexpr

◆ operator()()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

CK_TILE_DEVICE void ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::operator() ( QuantGemmKernelArgs kargs ) const

inline

◆ RunGemm()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>

static CK_TILE_DEVICE void ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::RunGemm	(	const ADataType *	a_ptr,
		const BDataType *	b_ptr,
		const AQDataType *	aq_ptr,
		const BQDataType *	bq_ptr,
		CDataType *	c_ptr,
		void *	smem_ptr_0,
		const QuantGemmKernelArgs &	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
aq_ptr	input AQ pointer
bq_ptr	input BQ pointer
c_ptr	output C pointer
smem_ptr_0	The start memory pointer of the shared memory block.
kargs	GEMM kernel arguments
splitk_batch_offset	splitk_batch_offset Utility structure used to calculate k batch.
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).

◆ RunGemm2LDS()

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

template<memory_operation_enum DstInMemOp = memory_operation_enum::set>

static CK_TILE_DEVICE void ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::RunGemm2LDS	(	const ADataType *	a_ptr,
		const BDataType *	b_ptr,
		const AQDataType *	aq_ptr,
		const BQDataType *	bq_ptr,
		CDataType *	c_ptr,
		void *	smem_ptr_0,
		void *	smem_ptr_1,
		const QuantGemmKernelArgs &	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
aq_ptr	input AQ pointer
c_ptr	output C pointer
smem_ptr_0	The start memory pointer of the shared memory block.
kargs	GEMM kernel arguments
splitk_batch_offset	splitk_batch_offset Utility structure used to calculate k batch.
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_ , QuantType QuantType_>

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

staticconstexpr

◆ I1

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

staticconstexpr

◆ I2

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

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

staticconstexpr

◆ I3

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::I3 = number<3>()

staticconstexpr

◆ I4

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::I4 = number<4>()

staticconstexpr

◆ kBlockSize

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr index_t ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::kBlockSize = GemmPipeline::BlockSize

staticconstexpr

◆ kQuantType

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::kQuantType = QuantType_

staticconstexpr

◆ PreshuffleB

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr bool ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::PreshuffleB = detail::is_preshuffleB_enabled<GemmPipeline_>::value

staticconstexpr

◆ PreshuffleQuant

template<typename TilePartitioner_ , typename GemmPipeline_ , typename EpiloguePipeline_ , QuantType QuantType_>

constexpr bool ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::PreshuffleQuant

staticconstexpr

Initial value:

=

detail::is_quantpreshuffle_enabled<GemmPipeline_>::value

ck_tile::detail::is_quantpreshuffle_enabled::value

static constexpr bool value

Definition: gemm_quant_kernel.hpp:72

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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp

QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ > Struct Template Reference

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

Classes

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ AccDataType

◆ ADataType

◆ ALayout

◆ AQDataType

◆ AQLayout

◆ BDataType

◆ BLayout

◆ BQDataType

◆ BQLayout

◆ CDataType

◆ CLayout

◆ EpiloguePipeline

◆ GemmPipeline

◆ TilePartitioner

Member Function Documentation

◆ BlockSize()

◆ GetName()

◆ GetSmemSize()

◆ GridSize()

◆ IsSupportedArgument()

◆ MakeGemmPadViews()

◆ MakeGemmTensorViews()

◆ MakeGemmTileWindows()

◆ MakeKernelArgs()

◆ operator()()

◆ RunGemm()

◆ RunGemm2LDS()

Member Data Documentation

◆ I0

◆ I1

◆ I2

◆ I3

◆ I4

◆ kBlockSize

◆ kQuantType

◆ PreshuffleB

◆ PreshuffleQuant

QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ > Struct Template Reference#