WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy > Struct Template Reference

WPQuantBPipelineAgBgCrV2&lt; Problem, PipelinePolicy &gt; Struct Template Reference#

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

#include <gemm_wp_bquant_pipeline_ag_bg_cr_v2.hpp>

Inheritance diagram for ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >:
ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy > ck_tile::BaseWeightPreshufflePipelineAGmemBGmemCRegV2< Problem >

Public Types

using Base = WeightPreshufflePipelineAGmemBGmemCRegV2<Problem>
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 ALayout = remove_cvref_t<typename Problem::ALayout>
using BLayout = remove_cvref_t<typename Problem::BLayout>
using BQLayout = remove_cvref_t<typename Problem::BQLayout>
using CLayout = remove_cvref_t<typename Problem::CLayout>
using BlockWeightPreshuffle
using WG = remove_cvref_t<decltype(config.template at<0>())>
Public Types inherited from ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >
using Base = BaseWeightPreshufflePipelineAGmemBGmemCRegV2<Problem>
using AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>
using BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>
using CDataType = remove_cvref_t<typename Problem::CDataType>
using AElementWise = remove_cvref_t<typename Problem::AElementWise>
using BElementWise = remove_cvref_t<typename Problem::BElementWise>
using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>
using AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>
using BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>
using CLayout = remove_cvref_t<typename Problem::CLayout>
using ALayout = remove_cvref_t<std::tuple_element_t<0, AsLayout>>
using BLayout = remove_cvref_t<std::tuple_element_t<0, BsLayout>>
using ADataType = remove_cvref_t<std::tuple_element_t<0, AsDataType>>
using BDataType = remove_cvref_t<std::tuple_element_t<0, BsDataType>>
using BlockWeightPreshuffle
using WG = remove_cvref_t<decltype(config.template at<0>())>
using BlockTile = remove_cvref_t<typename BlockGemmShape::BlockTile>
using BlockWarps = remove_cvref_t<typename BlockGemmShape::BlockWarps>
using WarpTile = remove_cvref_t<typename BlockGemmShape::WarpTile>

Public Member Functions

template<TailNumber TailNum, typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp, typename AElementFunction, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, const BQDramBlockWindowTmp &bq_dram_block_window_tmp, index_t num_loop, void *p_smem_ping, void *p_smem_pong) const
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, const BQDramBlockWindowTmp &bq_dram_block_window_tmp, index_t num_loop, void *p_smem_ping, void *p_smem_pong) const
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, const BQDramBlockWindowTmp &bq_dram_block_window_tmp, index_t num_loop, TailNumber tail_number, void *p_smem_ping, void *p_smem_pong) const
Public Member Functions inherited from ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >
template<TailNumber TailNum, typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename AElementFunction, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BFlatBlockWindowTmp >::value, bool > * = nullptr, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, index_t num_loop, void *p_smem_ping, void *p_smem_pong) const
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename AElementFunction, typename BElementFunction, typename std::enable_if_t< is_detected< is_tuple, ADramBlockWindowTmp >::value &&is_detected< is_tuple, BFlatBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const AElementFunction &a_element_func, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, const BElementFunction &b_element_func, index_t num_loop, void *p_smem_ping, void *p_smem_pong) const
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BFlatBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, index_t num_loop, void *p_smem_ping, void *p_smem_pong) const
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename std::enable_if_t<!is_detected< is_tuple, ADramBlockWindowTmp >::value &&!is_detected< is_tuple, BFlatBlockWindowTmp >::value, bool > * = nullptr>
CK_TILE_DEVICE auto operator() (const ADramBlockWindowTmp &a_dram_block_window_tmp, const BFlatBlockWindowTmp &b_flat_dram_block_window_tmp, index_t num_loop, TailNumber tail_number, void *__restrict__ p_smem_0, void *__restrict__ p_smem_1) const

Static Public Member Functions

static constexpr index_t GetVectorSizeBQ ()
static CK_TILE_HOST const std::string GetName ()
Static Public Member Functions inherited from ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >
template<bool IsWave32Host = false>
static constexpr index_t GetVectorSizeA ()
template<bool IsWave32Host = false>
static constexpr index_t GetVectorSizeB ()
static constexpr index_t GetVectorSizeC ()
static CK_TILE_HOST const std::string GetName ()
static CK_TILE_HOST_DEVICE constexpr auto TransposeC ()
static CK_TILE_HOST_DEVICE constexpr index_t GetSmemSize ()
static CK_TILE_HOST_DEVICE constexpr auto SchedulerPerM (index_t dsread_perM, index_t dswrite_perM, index_t load_perM)
static CK_TILE_HOST_DEVICE constexpr auto HotLoopScheduler ()
static CK_TILE_HOST_DEVICE constexpr auto Last2ndHotLoopScheduler ()
static CK_TILE_HOST_DEVICE constexpr auto LastHotLoopScheduler ()
Static Public Member Functions inherited from ck_tile::BaseWeightPreshufflePipelineAGmemBGmemCRegV2< 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, TailNumber tail_number)

Static Public Attributes

static constexpr auto config
static constexpr index_t KPerBlockBQ
static constexpr index_t QScalesPerBlockRow
static constexpr index_t KIterPerQScale = KIterPerWarp / QScalesPerBlockRow
static constexpr bool PreshuffleB = Problem::PreshuffleB
static constexpr auto TailNum = Problem::TailNum
static constexpr index_t kKPerBlock
static constexpr index_t kMPerBlock
static constexpr index_t kNPerBlock
static constexpr index_t KIterPerWarp
static constexpr index_t MIterPerWarp
static constexpr index_t NIterPerWarp
static constexpr index_t BlockSize
static constexpr bool kPadK
static constexpr bool kPadM
static constexpr bool kPadN
static constexpr auto I0
static constexpr auto I1
static constexpr auto I2
static constexpr index_t MWarp
static constexpr index_t NWarp
static constexpr index_t KPerBlockPerIter
static constexpr index_t MPerBlockPerIter
static constexpr index_t flatKPerWarp
static constexpr index_t flatNPerWarp
static constexpr index_t m_preload
Static Public Attributes inherited from ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >
static constexpr auto config
static constexpr index_t DsWritePreIssue = 3
static constexpr index_t DsReadPreload = 2
static constexpr index_t BlockSize = Problem::kBlockSize
static constexpr index_t WaveSize = get_warp_size()
static constexpr index_t kMPerBlock = BlockGemmShape::kM
static constexpr index_t kNPerBlock = BlockGemmShape::kN
static constexpr index_t kKPerBlock = BlockGemmShape::kK
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 flatKPerWarp = BlockGemmShape::flatKPerWarp
static constexpr index_t flatNPerWarp = BlockGemmShape::flatNPerWarp
static constexpr bool kPadM = Problem::kPadM
static constexpr bool kPadN = Problem::kPadN
static constexpr bool kPadK = Problem::kPadK
static constexpr index_t kLdsAlignmentInBytes = 16
static constexpr index_t NumWaveGroups = Problem::NumWaveGroups
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 index_t MWarp = config.template at<1>()
static constexpr index_t NWarp = config.template at<2>()
static constexpr index_t MIterPerWarp = kMPerBlock / (MWarp * WG::kM)
static constexpr index_t NIterPerWarp = kNPerBlock / (NWarp * WG::kN)
static constexpr index_t KIterPerWarp = kKPerBlock / WG::kK
static constexpr index_t KFlatPerBlockPerIter = flatKPerWarp
static constexpr index_t NFlatPerBlockPerIter = flatNPerWarp
static constexpr index_t MPerBlockPerIter = kMPerBlock / MIterPerWarp
static constexpr index_t KPerBlockPerIter = kKPerBlock / KIterPerWarp
static constexpr index_t K1 = Problem::VectorLoadSize / sizeof(ADataType)
static constexpr index_t m_preload
static constexpr auto TailNum = Problem::TailNum
static constexpr index_t mfma_per_wg = 1
static constexpr index_t dsread_per_wg
static constexpr index_t dsread_num_perK
static constexpr index_t dswrite_num_perK = dsread_num_perK / (MWarp * NWarp)
static constexpr index_t dswrite_rep = (dswrite_num_perK + MIterPerWarp - 1) / MIterPerWarp
static constexpr index_t Aload_num_perK = dswrite_num_perK
static constexpr index_t Aload_rep = dswrite_rep
static constexpr index_t Bload_num_perK = kNPerBlock * WG::kK / NWarp / K1 / WaveSize
static constexpr index_t HalfMIter = (MIterPerWarp + 1) / 2
static constexpr index_t Bload_rep = (Bload_num_perK + HalfMIter - 1) / HalfMIter
static constexpr index_t mfma_perM_perK = NIterPerWarp * mfma_per_wg
static constexpr index_t dswrite_mIter = (DsWritePreIssue - 1) % MIterPerWarp
static constexpr index_t dswrite_kIter = (DsWritePreIssue - 1) / MIterPerWarp
static constexpr bool DoubleSmemBuffer = Problem::DoubleSmemBuffer
static constexpr index_t Preshuffle = Problem::Preshuffle
static constexpr bool UsePersistentKernel
Static Public Attributes inherited from ck_tile::BaseWeightPreshufflePipelineAGmemBGmemCRegV2< Problem >
static constexpr index_t PrefetchStages = 2
static constexpr index_t PrefillStages = 1
static constexpr index_t GlobalBufferNum = 1
static constexpr bool UsePersistentKernel = Problem::Traits::UsePersistentKernel

Member Typedef Documentation

◆ ADataType

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::ADataType = remove_cvref_t<typename Problem::ADataType>

◆ ALayout

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::ALayout = remove_cvref_t<typename Problem::ALayout>

◆ Base

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::Base = WeightPreshufflePipelineAGmemBGmemCRegV2<Problem>

◆ BDataType

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BDataType = remove_cvref_t<typename Problem::BDataType>

◆ BLayout

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BLayout = remove_cvref_t<typename Problem::BLayout>

◆ BlockGemmShape

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>

◆ BlockWeightPreshuffle

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BlockWeightPreshuffle
Initial value:
remove_cvref_t<
decltype(PipelinePolicy::template GetBlockWeightPreshuffleBQuant<Problem>())>
ck_tile::remove_cvref_t
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition type_traits.hpp:21

◆ BQDataType

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BQDataType = remove_cvref_t<typename Problem::BQDataType>

◆ BQLayout

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::BQLayout = remove_cvref_t<typename Problem::BQLayout>

◆ CDataType

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::CDataType = remove_cvref_t<typename Problem::CDataType>

◆ CLayout

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::CLayout = remove_cvref_t<typename Problem::CLayout>

◆ ComputeDataType

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>

◆ QuantGroupSize

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::QuantGroupSize = remove_cvref_t<typename Problem::QuantGroupSize>

◆ WG

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
using ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::WG = remove_cvref_t<decltype(config.template at<0>())>

Member Function Documentation

◆ GetName()

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
CK_TILE_HOST const std::string ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::GetName ( )
inlinestaticnodiscard

◆ GetVectorSizeBQ()

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
constexpr index_t ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::GetVectorSizeBQ ( )
inlinestaticconstexpr

◆ operator()() [1/3]

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
template<TailNumber TailNum, typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp, typename AElementFunction, index_t UnaryOpSize_ = 8>
CK_TILE_DEVICE auto ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const AElementFunction & a_element_func,
const BFlatBlockWindowTmp & b_flat_dram_block_window_tmp,
const BQDramBlockWindowTmp & bq_dram_block_window_tmp,
index_t num_loop,
void * p_smem_ping,
void * p_smem_pong ) const
inline

◆ operator()() [2/3]

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp>
CK_TILE_DEVICE auto ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BFlatBlockWindowTmp & b_flat_dram_block_window_tmp,
const BQDramBlockWindowTmp & bq_dram_block_window_tmp,
index_t num_loop,
TailNumber tail_number,
void * p_smem_ping,
void * p_smem_pong ) const
inline

◆ operator()() [3/3]

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
template<typename ADramBlockWindowTmp, typename BFlatBlockWindowTmp, typename BQDramBlockWindowTmp>
CK_TILE_DEVICE auto ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::operator() ( const ADramBlockWindowTmp & a_dram_block_window_tmp,
const BFlatBlockWindowTmp & b_flat_dram_block_window_tmp,
const BQDramBlockWindowTmp & bq_dram_block_window_tmp,
index_t num_loop,
void * p_smem_ping,
void * p_smem_pong ) const
inline

Member Data Documentation

◆ BlockSize

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::BlockSize
staticconstexpr

◆ config

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
auto ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::config
staticconstexpr
Initial value:
=
BlockWeightPreshuffle::BlockPolicy::template GetWarpGemmMWarpNWarp<Problem>()

◆ flatKPerWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::flatKPerWarp
staticconstexpr

◆ flatNPerWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::flatNPerWarp
staticconstexpr

◆ I0

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
auto ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::I0
staticconstexpr

◆ I1

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
auto ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::I1
staticconstexpr

◆ I2

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
auto ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::I2
staticconstexpr

◆ KIterPerQScale

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::KIterPerQScale = KIterPerWarp / QScalesPerBlockRow
staticconstexpr

◆ KIterPerWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::KIterPerWarp
staticconstexpr

◆ kKPerBlock

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kKPerBlock
staticconstexpr

◆ kMPerBlock

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kMPerBlock
staticconstexpr

◆ kNPerBlock

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kNPerBlock
staticconstexpr

◆ kPadK

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
bool ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kPadK
staticconstexpr

◆ kPadM

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
bool ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kPadM
staticconstexpr

◆ kPadN

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
bool ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::kPadN
staticconstexpr

◆ KPerBlockBQ

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::KPerBlockBQ
staticconstexpr
Initial value:
=
integer_divide_ceil(BlockGemmShape::kK, 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

◆ KPerBlockPerIter

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::KPerBlockPerIter
staticconstexpr

◆ m_preload

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::m_preload
staticconstexpr

◆ MIterPerWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::MIterPerWarp
staticconstexpr

◆ MPerBlockPerIter

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::MPerBlockPerIter
staticconstexpr

◆ MWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::MWarp
staticconstexpr

◆ NIterPerWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::NIterPerWarp
staticconstexpr

◆ NWarp

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WeightPreshufflePipelineAGmemBGmemCRegV2< Problem, PipelinePolicy >::NWarp
staticconstexpr

◆ PreshuffleB

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
bool ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::PreshuffleB = Problem::PreshuffleB
staticconstexpr

◆ QScalesPerBlockRow

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
index_t ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::QScalesPerBlockRow
staticconstexpr
Initial value:
=
integer_divide_ceil(kKPerBlock, QuantGroupSize::kK)
ck_tile::WPQuantBPipelineAgBgCrV2::kKPerBlock
static constexpr index_t kKPerBlock
Definition wp_pipeline_agmem_bgmem_creg_v2.hpp:91

◆ TailNum

template<typename Problem, typename PipelinePolicy = GemmWPQuantPipelineAgBgCrPolicy>
auto ck_tile::WPQuantBPipelineAgBgCrV2< Problem, PipelinePolicy >::TailNum = Problem::TailNum
staticconstexpr
The documentation for this struct was generated from the following file: