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

• ck_tile
• QuantGemmKernel

Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members

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 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_DEVICE auto	MakeABlockWindow (const ADataType *a_ptr, const QuantGemmKernelArgs &kargs, const index_t k_size, const index_t i_m)
static CK_TILE_DEVICE auto	MakeAQBlockWindow (const AQDataType *aq_ptr, const QuantGemmKernelArgs &kargs, const index_t i_m, const index_t i_n)
static CK_TILE_DEVICE auto	MakeBBlockWindow (const BDataType *b_ptr, const QuantGemmKernelArgs &kargs, const index_t k_size, const index_t i_n)
static CK_TILE_DEVICE auto	MakeBQBlockWindow (const BQDataType *bq_ptr, const QuantGemmKernelArgs &kargs, const index_t i_m, const index_t i_n)
template<memory_operation_enum DstInMemOp = memory_operation_enum::set>
static CK_TILE_DEVICE auto	MakeCBlockWindow (CDataType *c_ptr, const QuantGemmKernelArgs &kargs, const index_t i_m, const index_t i_n)
static CK_TILE_HOST bool	IsSupportedArgument (const QuantGemmKernelArgs &kargs)
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, 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 CK_TILE_HOST auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::BlockSize ( )

inlinestatic

◆ 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

◆ MakeABlockWindow()

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

static CK_TILE_DEVICE auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeABlockWindow ( const ADataType *  a_ptr, const QuantGemmKernelArgs &  kargs, const index_t  k_size, const index_t  i_m  )

inlinestatic

◆ MakeAQBlockWindow()

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

static CK_TILE_DEVICE auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeAQBlockWindow ( const AQDataType *  aq_ptr, const QuantGemmKernelArgs &  kargs, const index_t  i_m, const index_t  i_n  )

inlinestatic

◆ MakeBBlockWindow()

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

static CK_TILE_DEVICE auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeBBlockWindow ( const BDataType *  b_ptr, const QuantGemmKernelArgs &  kargs, const index_t  k_size, const index_t  i_n  )

inlinestatic

◆ MakeBQBlockWindow()

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

static CK_TILE_DEVICE auto ck_tile::QuantGemmKernel< TilePartitioner_, GemmPipeline_, EpiloguePipeline_, QuantType_ >::MakeBQBlockWindow ( const BQDataType *  bq_ptr, const QuantGemmKernelArgs &  kargs, const index_t  i_m, const index_t  i_n  )

inlinestatic

◆ MakeCBlockWindow()

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_ >::MakeCBlockWindow ( CDataType *  c_ptr, const QuantGemmKernelArgs &  kargs, 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_>

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, 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	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.

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