BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ > Struct Template Reference

BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ > Struct Template Reference#

Composable Kernel: ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ > Struct Template Reference

#include <block_universal_gemm_ar_flatbr_bquant_cr.hpp>

Public Types
using	Problem = remove_cvref_t<Problem_>
using	BlockPolicy = remove_cvref_t<BlockPolicy_>
using	ADataType = remove_cvref_t<typename Problem::ADataType>
using	BDataType = remove_cvref_t<typename Problem::BDataType>
using	BQDataType = remove_cvref_t<typename Problem::BQDataType>
using	CDataType = remove_cvref_t<typename Problem::CDataType>
using	ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>
using	BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using	QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>
using	BlockTile = remove_cvref_t<typename BlockGemmShape::BlockTile>
using	BlockWarps = remove_cvref_t<typename BlockGemmShape::BlockWarps>
using	WarpTile = remove_cvref_t<typename BlockGemmShape::WarpTile>
using	WG = remove_cvref_t<decltype(config.template at<0>())>

Public Member Functions
template<typename CBlockTensor, typename ABlockTensor, typename BFlatBlockTensor, typename BQBlockTensor, typename ABlockWindow>
CK_TILE_DEVICE void	operator() (CBlockTensor &c_block_tensor, ABlockTensor &a_warp_tensor, BFlatBlockTensor &b_warp_tensor, BQBlockTensor &bq_block_tensor, ABlockWindow &a_warp_windows) const

Static Public Member Functions
template<typename T>
static CK_TILE_DEVICE float	cvt_scale_to_fp32 (T &scale)
static CK_TILE_DEVICE constexpr auto	MakeCBlockTile ()

Static Public Attributes
static constexpr auto	I0 = number<0>()
static constexpr auto	I1 = number<1>()
static constexpr auto	I2 = number<2>()
static constexpr auto	idxM = I0
static constexpr auto	idxN = I1
static constexpr auto	idxK = I2
static constexpr auto	config = BlockPolicy::template GetWarpGemmMWarpNWarp<Problem>()
static constexpr auto	warp_size = get_warp_size()
static constexpr index_t	MWarp = config.template at<1>()
static constexpr index_t	NWarp = config.template at<2>()
static constexpr index_t	MPerBlock = BlockGemmShape::kM
static constexpr index_t	KPerBlock = BlockGemmShape::kK
static constexpr index_t	kBlockSize = Problem::kBlockSize
static constexpr index_t	MIterPerWarp = MPerBlock / (MWarp * WG::kM)
static constexpr index_t	NIterPerWarp
static constexpr index_t	KIterPerWarp = KPerBlock / WG::kK
static constexpr auto	MIter_2nd_last
static constexpr index_t	KPerBlockBQ = KPerBlock / QuantGroupSize::kK
static constexpr index_t	QScalesPerBlockRow
static constexpr index_t	QScalesPerWarpGemmRow
static constexpr index_t	KIterPerQScale = KIterPerWarp / QScalesPerBlockRow
static constexpr index_t	DsReadPreload = 2
static constexpr index_t	m_preload

Member Typedef Documentation

◆ ADataType

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::ADataType = remove_cvref_t<typename Problem::ADataType>

◆ BDataType

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BDataType = remove_cvref_t<typename Problem::BDataType>

◆ BlockGemmShape

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BlockPolicy

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BlockPolicy = remove_cvref_t<BlockPolicy_>

◆ BlockTile

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BlockTile = remove_cvref_t<typename BlockGemmShape::BlockTile>

◆ BlockWarps

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BlockWarps = remove_cvref_t<typename BlockGemmShape::BlockWarps>

◆ BQDataType

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::BQDataType = remove_cvref_t<typename Problem::BQDataType>

◆ CDataType

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ ComputeDataType

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>

◆ Problem

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::Problem = remove_cvref_t<Problem_>

◆ QuantGroupSize

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>

◆ WarpTile

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::WarpTile = remove_cvref_t<typename BlockGemmShape::WarpTile>

◆ WG

template<typename Problem_, typename BlockPolicy_>

using ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::WG = remove_cvref_t<decltype(config.template at<0>())>

Member Function Documentation

◆ cvt_scale_to_fp32()

template<typename Problem_, typename BlockPolicy_>

template<typename T>

CK_TILE_DEVICE float ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::cvt_scale_to_fp32 ( T & scale )

inlinestatic

◆ MakeCBlockTile()

template<typename Problem_, typename BlockPolicy_>

CK_TILE_DEVICE constexpr auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::MakeCBlockTile ( )

inlinestaticconstexpr

◆ operator()()

template<typename Problem_, typename BlockPolicy_>

template<typename CBlockTensor, typename ABlockTensor, typename BFlatBlockTensor, typename BQBlockTensor, typename ABlockWindow>

CK_TILE_DEVICE void ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::operator()	(	CBlockTensor &	c_block_tensor,
		ABlockTensor &	a_warp_tensor,
		BFlatBlockTensor &	b_warp_tensor,
		BQBlockTensor &	bq_block_tensor,
		ABlockWindow &	a_warp_windows ) const

inline

Member Data Documentation

◆ config

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::config = BlockPolicy::template GetWarpGemmMWarpNWarp<Problem>()

staticconstexpr

◆ DsReadPreload

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::DsReadPreload = 2

staticconstexpr

◆ I0

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::I0 = number<0>()

staticconstexpr

◆ I1

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::I1 = number<1>()

staticconstexpr

◆ I2

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::I2 = number<2>()

staticconstexpr

◆ idxK

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::idxK = I2

staticconstexpr

◆ idxM

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::idxM = I0

staticconstexpr

◆ idxN

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::idxN = I1

staticconstexpr

◆ kBlockSize

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::kBlockSize = Problem::kBlockSize

staticconstexpr

◆ KIterPerQScale

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::KIterPerQScale = KIterPerWarp / QScalesPerBlockRow

staticconstexpr

◆ KIterPerWarp

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::KIterPerWarp = KPerBlock / WG::kK

staticconstexpr

◆ KPerBlock

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::KPerBlock = BlockGemmShape::kK

staticconstexpr

◆ KPerBlockBQ

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::KPerBlockBQ = KPerBlock / QuantGroupSize::kK

staticconstexpr

◆ m_preload

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::m_preload

staticconstexpr

Initial value:

= (MIterPerWarp * KIterPerWarp >= DsReadPreload)
                                             ? DsReadPreload
                                             : MIterPerWarp * KIterPerWarp

◆ MIter_2nd_last

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::MIter_2nd_last

staticconstexpr

Initial value:

=

(MIterPerWarp >= 2) ? MIterPerWarp - 2 : MIterPerWarp - 1

ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg::MIterPerWarp

static constexpr index_t MIterPerWarp

Definition block_universal_gemm_ar_flatbr_bquant_cr.hpp:56

◆ MIterPerWarp

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::MIterPerWarp = MPerBlock / (MWarp * WG::kM)

staticconstexpr

◆ MPerBlock

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::MPerBlock = BlockGemmShape::kM

staticconstexpr

◆ MWarp

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::MWarp = config.template at<1>()

staticconstexpr

◆ NIterPerWarp

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::NIterPerWarp

staticconstexpr

Initial value:

=

BlockTile::at(idxN) / (WarpTile::at(idxN) * BlockWarps::at(idxN))

ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg::idxN

static constexpr auto idxN

Definition block_universal_gemm_ar_flatbr_bquant_cr.hpp:36

◆ NWarp

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::NWarp = config.template at<2>()

staticconstexpr

◆ QScalesPerBlockRow

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::QScalesPerBlockRow

staticconstexpr

Initial value:

=

integer_divide_ceil(KPerBlock, QuantGroupSize::kK)

ck_tile::integer_divide_ceil

CK_TILE_HOST_DEVICE constexpr auto integer_divide_ceil(X x, Y y)

Definition tile/core/numeric/math.hpp:149

ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg::KPerBlock

static constexpr index_t KPerBlock

Definition block_universal_gemm_ar_flatbr_bquant_cr.hpp:52

◆ QScalesPerWarpGemmRow

template<typename Problem_, typename BlockPolicy_>

index_t ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::QScalesPerWarpGemmRow

staticconstexpr

Initial value:

=

integer_divide_ceil(WG::kK, QuantGroupSize::kK)

◆ warp_size

template<typename Problem_, typename BlockPolicy_>

auto ck_tile::BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ >::warp_size = get_warp_size()

staticconstexpr

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

block_universal_gemm_ar_flatbr_bquant_cr.hpp

BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ > Struct Template Reference

BlockGemmWeightPreshuffleBQuantARegBRegCReg&lt; Problem_, BlockPolicy_ &gt; Struct Template Reference#

Public Types

Public Member Functions

Static Public Member Functions

Static Public Attributes

Member Typedef Documentation

◆ ADataType

◆ BDataType

◆ BlockGemmShape

◆ BlockPolicy

◆ BlockTile

◆ BlockWarps

◆ BQDataType

◆ CDataType

◆ ComputeDataType

◆ Problem

◆ QuantGroupSize

◆ WarpTile

◆ WG

Member Function Documentation

◆ cvt_scale_to_fp32()

◆ MakeCBlockTile()

◆ operator()()

Member Data Documentation

◆ config

◆ DsReadPreload

◆ I0

◆ I1

◆ I2

◆ idxK

◆ idxM

◆ idxN

◆ kBlockSize

◆ KIterPerQScale

◆ KIterPerWarp

◆ KPerBlock

◆ KPerBlockBQ

◆ m_preload

◆ MIter_2nd_last

◆ MIterPerWarp

◆ MPerBlock

◆ MWarp

◆ NIterPerWarp

◆ NWarp

◆ QScalesPerBlockRow

◆ QScalesPerWarpGemmRow

◆ warp_size

BlockGemmWeightPreshuffleBQuantARegBRegCReg< Problem_, BlockPolicy_ > Struct Template Reference#