AQuantGemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference

AQuantGemmPipelineAgBgCrMem&lt; Problem, Policy &gt; Struct Template Reference#

Composable Kernel: ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference
Classes | Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy > Struct Template Reference

#include <gemm_aquant_pipeline_ag_bg_cr_mem.hpp>

Inheritance diagram for ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >:
ck_tile::BaseAQuantGemmPipelineAgBgCrMem< Problem > ck_tile::BaseGemmPipelineAgBgCrMem< Problem >

Classes

struct  PipelineImpl
struct  PipelineImpl< GemmPipelineScheduler::Interwave >

Public Types

using Base = BaseGemmPipelineAgBgCrMem<Problem>
using PipelineImplBase = GemmAQuantPipelineAgBgCrImplBase<Problem, Policy>
using ADataType = remove_cvref_t<typename Problem::ADataType>
using AQDataType = remove_cvref_t<typename Problem::AQDataType>
using BDataType = remove_cvref_t<typename Problem::BDataType>
using CDataType = remove_cvref_t<typename Problem::CDataType>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>
using I0 = number<0>
using I1 = number<1>
using I2 = number<2>
using ALayout = remove_cvref_t<typename Problem::ALayout>
using AQLayout = remove_cvref_t<typename Problem::AQLayout>
using BLayout = remove_cvref_t<typename Problem::BLayout>
using CLayout = remove_cvref_t<typename Problem::CLayout>
using BlockGemm = remove_cvref_t<decltype(Policy::template GetBlockGemm<Problem>())>
Public Types inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
using ADataType = remove_cvref_t<typename Problem::ADataType>
using BDataType = remove_cvref_t<typename Problem::BDataType>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

Public Member Functions

template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AQDramBlockWindowTmp>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BDramBlockWindowTmp &b_dram_block_window_tmp, const AQDramBlockWindowTmp &aq_dram_block_window_tmp, index_t m, index_t num_loop, void *p_smem) const

Static Public Member Functions

static constexpr index_t GetVectorSizeA ()
static constexpr index_t GetVectorSizeB ()
static constexpr index_t GetVectorSizeC ()
static constexpr index_t GetVectorSizeAQ ()
static constexpr index_t GetSmemPackA ()
static constexpr index_t GetSmemPackB ()
static CK_TILE_HOST const std::string GetName ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()
static CK_TILE_HOST std::string Print ()
Static Public Member Functions inherited from ck_tile::BaseAQuantGemmPipelineAgBgCrMem< Problem >
static CK_TILE_HOST_DEVICE constexpr TailNumber GetBlockLoopTailNum (index_t num_loop)
template<typename RunFunction>
static CK_TILE_HOST_DEVICE auto TailHandler (const RunFunction &run_func, bool has_hot_loop, TailNumber tail_number)
Static Public Member Functions inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
static CK_TILE_HOST_DEVICE constexpr bool BlockHasHotloop (index_t num_loop)
static CK_TILE_HOST_DEVICE constexpr TailNumber GetBlockLoopTailNum (index_t num_loop)
template<typename RunFunction>
static CK_TILE_HOST_DEVICE auto TailHandler (const RunFunction &run_func, bool has_hot_loop, TailNumber tail_number)

Static Public Attributes

static constexpr index_t APackedSize
static constexpr index_t BPackedSize
static constexpr index_t AQPackedSize
static constexpr index_t BlockSize = Problem::kBlockSize
static constexpr index_t MPerBlock = BlockGemmShape::kM
static constexpr index_t NPerBlock = BlockGemmShape::kN
static constexpr index_t KPerBlock = BlockGemmShape::kK
static constexpr index_t KPerBlockAQ = BlockGemmShape::kK / QuantGroupSize::kK
static constexpr bool kPadM = Problem::kPadM
static constexpr bool kPadN = Problem::kPadN
static constexpr bool kPadK = Problem::kPadK
static constexpr bool DoubleSmemBuffer = Problem::DoubleSmemBuffer
static constexpr bool PreshuffleQuant = Problem::Traits::PreshuffleQuant
static constexpr bool HasHotLoop = Problem::HasHotLoop
static constexpr auto TailNum = Problem::TailNum
static constexpr auto Scheduler = Problem::Scheduler
static constexpr index_t PrefetchStages
Static Public Attributes inherited from ck_tile::BaseGemmPipelineAgBgCrMem< Problem >
static constexpr index_t APackedSize
static constexpr index_t BPackedSize
static constexpr index_t BlockSize = Problem::kBlockSize
static constexpr index_t MPerBlock = BlockGemmShape::kM
static constexpr index_t NPerBlock = BlockGemmShape::kN
static constexpr index_t KPerBlock = BlockGemmShape::kK
static constexpr index_t MinMemInFlyBytes = 32768
static constexpr index_t WgpPerCU
static constexpr index_t FullMemBandPrefetchStages
static constexpr index_t PrefetchStages
static constexpr index_t LocalPrefillStages = 1
static constexpr index_t GlobalBufferNum = PrefetchStages
static constexpr bool UsePersistentKernel = Problem::Traits::UsePersistentKernel

Member Typedef Documentation

◆ ADataType

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::ADataType = remove_cvref_t<typename Problem::ADataType>

◆ ALayout

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::ALayout = remove_cvref_t<typename Problem::ALayout>

◆ AQDataType

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::AQDataType = remove_cvref_t<typename Problem::AQDataType>

◆ AQLayout

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::AQLayout = remove_cvref_t<typename Problem::AQLayout>

◆ Base

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::Base = BaseGemmPipelineAgBgCrMem<Problem>

◆ BDataType

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BDataType = remove_cvref_t<typename Problem::BDataType>

◆ BLayout

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BLayout = remove_cvref_t<typename Problem::BLayout>

◆ BlockGemm

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BlockGemm = remove_cvref_t<decltype(Policy::template GetBlockGemm<Problem>())>

◆ BlockGemmShape

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ CDataType

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ I0

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::I0 = number<0>

◆ I1

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::I1 = number<1>

◆ I2

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::I2 = number<2>

◆ PipelineImplBase

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::PipelineImplBase = GemmAQuantPipelineAgBgCrImplBase<Problem, Policy>

◆ QuantGroupSize

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
using ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>

Member Function Documentation

◆ GetName()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
CK_TILE_HOST const std::string ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetName ( )
inlinestaticnodiscard

◆ GetSmemPackA()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetSmemPackA ( )
inlinestaticconstexpr

◆ GetSmemPackB()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetSmemPackB ( )
inlinestaticconstexpr

◆ GetSmemSize()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
CK_TILE_HOST_DEVICE constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetSmemSize ( )
inlinestaticconstexpr

◆ GetVectorSizeA()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeA ( )
inlinestaticconstexpr

◆ GetVectorSizeAQ()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeAQ ( )
inlinestaticconstexpr

◆ GetVectorSizeB()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeB ( )
inlinestaticconstexpr

◆ GetVectorSizeC()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
constexpr index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::GetVectorSizeC ( )
inlinestaticconstexpr

◆ operator()()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
template<typename ADramBlockWindowTmp, typename BDramBlockWindowTmp, typename AQDramBlockWindowTmp>
CK_TILE_DEVICE auto ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BDramBlockWindowTmp & b_dram_block_window_tmp,
const AQDramBlockWindowTmp & aq_dram_block_window_tmp,
index_t m,
index_t num_loop,
void * p_smem ) const
inline

◆ Print()

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
CK_TILE_HOST std::string ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::Print ( )
inlinestatic

Member Data Documentation

◆ APackedSize

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::APackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<ADataType>>::PackedSize
ck_tile::numeric_traits
Definition tile/core/numeric/numeric.hpp:81

◆ AQPackedSize

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::AQPackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<AQDataType>>::PackedSize

◆ BlockSize

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BlockSize = Problem::kBlockSize
staticconstexpr

◆ BPackedSize

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::BPackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<BDataType>>::PackedSize

◆ DoubleSmemBuffer

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::DoubleSmemBuffer = Problem::DoubleSmemBuffer
staticconstexpr

◆ HasHotLoop

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::HasHotLoop = Problem::HasHotLoop
staticconstexpr

◆ kPadK

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::kPadK = Problem::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::kPadM = Problem::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::kPadN = Problem::kPadN
staticconstexpr

◆ KPerBlock

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ KPerBlockAQ

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::KPerBlockAQ = BlockGemmShape::kK / QuantGroupSize::kK
staticconstexpr

◆ MPerBlock

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::NPerBlock = BlockGemmShape::kN
staticconstexpr

◆ PrefetchStages

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::PrefetchStages
staticconstexpr

◆ PreshuffleQuant

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
bool ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::PreshuffleQuant = Problem::Traits::PreshuffleQuant
staticconstexpr

◆ Scheduler

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
auto ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::Scheduler = Problem::Scheduler
staticconstexpr

◆ TailNum

template<typename Problem, typename Policy = GemmAQuantPipelineAgBgCrDefaultPolicy>
auto ck_tile::AQuantGemmPipelineAgBgCrMem< Problem, Policy >::TailNum = Problem::TailNum
staticconstexpr
The documentation for this struct was generated from the following file: