BaseGemmPipelineAgBgCrMem< Problem > Struct Template Reference

BaseGemmPipelineAgBgCrMem&lt; Problem &gt; Struct Template Reference#

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

#include <gemm_pipeline_ag_bg_cr_mem.hpp>

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

Public Types

using ADataType = remove_cvref_t<typename Problem::ADataType>
using BDataType = remove_cvref_t<typename Problem::BDataType>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

Static Public Member Functions

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 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>
using ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::ADataType = remove_cvref_t<typename Problem::ADataType>

◆ BDataType

template<typename Problem>
using ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::BDataType = remove_cvref_t<typename Problem::BDataType>

◆ BlockGemmShape

template<typename Problem>
using ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

Member Function Documentation

◆ BlockHasHotloop()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr bool ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::BlockHasHotloop ( index_t num_loop)
inlinestaticconstexpr

◆ GetBlockLoopTailNum()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr TailNumber ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::GetBlockLoopTailNum ( index_t num_loop)
inlinestaticconstexpr

◆ TailHandler()

template<typename Problem>
template<typename RunFunction>
CK_TILE_HOST_DEVICE auto ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::TailHandler ( const RunFunction & run_func,
bool has_hot_loop,
TailNumber tail_number )
inlinestatic

◆ TransposeC()

template<typename Problem>
CK_TILE_HOST_DEVICE constexpr auto ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::TransposeC ( )
inlinestaticconstexpr

Member Data Documentation

◆ APackedSize

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

◆ BlockSize

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::BlockSize = Problem::kBlockSize
staticconstexpr

◆ BPackedSize

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::BPackedSize
staticconstexpr
Initial value:
=
ck_tile::numeric_traits<remove_cvref_t<BDataType>>::PackedSize

◆ FullMemBandPrefetchStages

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::FullMemBandPrefetchStages
staticconstexpr
Initial value:
=
integer_divide_ceil(MinMemInFlyBytes / WgpPerCU,
(MPerBlock * sizeof(ADataType) / APackedSize +
NPerBlock * sizeof(BDataType) / BPackedSize) *
KPerBlock)
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::BlockwiseGemmXdlops_pipeline_v2< BlockGemmPipelineScheduler::Intrawave, BlockSize, ADataType, BDataType, ComputeDataType, AccDataType, ATileDesc, BTileDesc, AMmaTileDesc, BMmaTileDesc, ABlockTransferSrcScalarPerVector, BBlockTransferSrcScalarPerVector, MPerBlock, NPerBlock, KPerBlock, MPerXDL, NPerXDL, MRepeat, NRepeat, KPack >::WgpPerCU
static constexpr index_t WgpPerCU
Definition blockwise_gemm_pipeline_xdlops_v2.hpp:146
ck_tile::BaseGemmPipelineAgBgCrMem::MPerBlock
static constexpr index_t MPerBlock
Definition gemm_pipeline_ag_bg_cr_mem.hpp:32
ck_tile::BaseGemmPipelineAgBgCrMem::BDataType
remove_cvref_t< typename Problem::BDataType > BDataType
Definition gemm_pipeline_ag_bg_cr_mem.hpp:21
ck_tile::BaseGemmPipelineAgBgCrMem::BPackedSize
static constexpr index_t BPackedSize
Definition gemm_pipeline_ag_bg_cr_mem.hpp:26
ck_tile::BaseGemmPipelineAgBgCrMem::ADataType
remove_cvref_t< typename Problem::ADataType > ADataType
Definition gemm_pipeline_ag_bg_cr_mem.hpp:20
ck_tile::BaseGemmPipelineAgBgCrMem::MinMemInFlyBytes
static constexpr index_t MinMemInFlyBytes
Definition gemm_pipeline_ag_bg_cr_mem.hpp:37
ck_tile::BaseGemmPipelineAgBgCrMem::KPerBlock
static constexpr index_t KPerBlock
Definition gemm_pipeline_ag_bg_cr_mem.hpp:34
ck_tile::BaseGemmPipelineAgBgCrMem::APackedSize
static constexpr index_t APackedSize
Definition gemm_pipeline_ag_bg_cr_mem.hpp:24
ck_tile::BaseGemmPipelineAgBgCrMem::NPerBlock
static constexpr index_t NPerBlock
Definition gemm_pipeline_ag_bg_cr_mem.hpp:33

◆ GlobalBufferNum

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::GlobalBufferNum = PrefetchStages
staticconstexpr

◆ KPerBlock

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::KPerBlock = BlockGemmShape::kK
staticconstexpr

◆ LocalPrefillStages

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::LocalPrefillStages = 1
staticconstexpr

◆ MinMemInFlyBytes

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::MinMemInFlyBytes = 32768
staticconstexpr

◆ MPerBlock

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::MPerBlock = BlockGemmShape::kM
staticconstexpr

◆ NPerBlock

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::NPerBlock = BlockGemmShape::kN
staticconstexpr

◆ PrefetchStages

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::PrefetchStages
staticconstexpr
Initial value:
=
FullMemBandPrefetchStages >= 2
? FullMemBandPrefetchStages <= 8 ? FullMemBandPrefetchStages : 8
: 2
ck::BlockwiseGemmXdlops_pipeline_v2< BlockGemmPipelineScheduler::Intrawave, BlockSize, ADataType, BDataType, ComputeDataType, AccDataType, ATileDesc, BTileDesc, AMmaTileDesc, BMmaTileDesc, ABlockTransferSrcScalarPerVector, BBlockTransferSrcScalarPerVector, MPerBlock, NPerBlock, KPerBlock, MPerXDL, NPerXDL, MRepeat, NRepeat, KPack >::FullMemBandPrefetchStages
static constexpr index_t FullMemBandPrefetchStages
Definition blockwise_gemm_pipeline_xdlops_v2.hpp:148

◆ UsePersistentKernel

template<typename Problem>
bool ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::UsePersistentKernel = Problem::Traits::UsePersistentKernel
staticconstexpr

◆ WgpPerCU

template<typename Problem>
index_t ck_tile::BaseGemmPipelineAgBgCrMem< Problem >::WgpPerCU
staticconstexpr
Initial value:
=
(4 * get_warp_size() / BlockSize) >= 1 ? 4 * get_warp_size() / BlockSize : 1
ck_tile::get_warp_size
CK_TILE_HOST_DEVICE constexpr index_t get_warp_size()
Definition arch.hpp:63
ck_tile::BaseGemmPipelineAgBgCrMem::BlockSize
static constexpr index_t BlockSize
Definition gemm_pipeline_ag_bg_cr_mem.hpp:31
The documentation for this struct was generated from the following file: