FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ > Struct Template Reference

FusedMoeGemmPipeline_FlatmmEx&lt; Problem_, Policy_ &gt; Struct Template Reference#

Composable Kernel: ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ > Struct Template Reference
Public Types | Public Member Functions | Static Public Member Functions | Static Public Attributes | List of all members
ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ > Struct Template Reference

#include <fused_moegemm_pipeline_flatmm_ex.hpp>

Public Types

using Problem = remove_cvref_t<Problem_>
using Policy = remove_cvref_t<Policy_>
using BlockShape = typename Problem::BlockShape
using ADataType = typename Problem::ADataType
using GDataType = typename Problem::GDataType
using DDataType = typename Problem::DDataType
using AccDataType = typename Problem::AccDataType
using ODataType = typename Problem::ODataType
using AScaleDataType = typename Problem::AScaleDataType
using GScaleDataType = typename Problem::GScaleDataType
using DScaleDataType = typename Problem::DScaleDataType
using YSmoothScaleDataType = typename Problem::YSmoothScaleDataType
using TopkWeightDataType = typename Problem::TopkWeightDataType
using IndexDataType = typename Problem::IndexDataType
using YDataType = typename Problem::YDataType
using Traits = typename Problem::Traits

Public Member Functions

template<typename AWindow, typename GWindow, typename DWindow, typename OWindow>
CK_TILE_DEVICE auto operator() (const AWindow &a_window_, const GWindow &g_window_, const DWindow &d_window_, OWindow &o_window_, TopkWeightDataType, CK_TILE_LDS_ADDR void *smem, index_t hidden_size, index_t intermediate_size)

Static Public Member Functions

static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize_A ()
static CK_TILE_HOST_DEVICE constexpr ck_tile::index_t GetSmemSize ()
static CK_TILE_HOST_DEVICE auto GetACoord ()
static CK_TILE_HOST_DEVICE auto GetOCoord ()

Static Public Attributes

static constexpr bool IsGateOnly = Traits::IsGateOnly
static constexpr bool UseSmoothQuant = Traits::UseSmoothQuant
static constexpr bool PadHiddenSize = Traits::PadHiddenSize
static constexpr bool PadIntermediateSize = Traits::PadIntermediateSize
static constexpr index_t kAlignmentA = Policy::template GetAlignment_A<Problem>()
static constexpr index_t kAlignmentG = Policy::template GetAlignment_G<Problem>()
static constexpr index_t kAlignmentD = Policy::template GetAlignment_D<Problem>()
static constexpr index_t kAlignmentO = Policy::template GetAlignment_O<Problem>()
static constexpr index_t SLD_A = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::SLD_A)
static constexpr index_t GLD_A = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GLD_A)
static constexpr index_t GLD_B = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GLD_B)
static constexpr index_t GST_O = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GST_O)
static constexpr index_t kBlockPerCu
static constexpr const char * name = "fused_moe_flatmm"

Member Typedef Documentation

◆ AccDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::AccDataType = typename Problem::AccDataType

◆ ADataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::ADataType = typename Problem::ADataType

◆ AScaleDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::AScaleDataType = typename Problem::AScaleDataType

◆ BlockShape

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::BlockShape = typename Problem::BlockShape

◆ DDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::DDataType = typename Problem::DDataType

◆ DScaleDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::DScaleDataType = typename Problem::DScaleDataType

◆ GDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GDataType = typename Problem::GDataType

◆ GScaleDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GScaleDataType = typename Problem::GScaleDataType

◆ IndexDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::IndexDataType = typename Problem::IndexDataType

◆ ODataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::ODataType = typename Problem::ODataType

◆ Policy

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::Policy = remove_cvref_t<Policy_>

◆ Problem

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::Problem = remove_cvref_t<Problem_>

◆ TopkWeightDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::TopkWeightDataType = typename Problem::TopkWeightDataType

◆ Traits

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::Traits = typename Problem::Traits

◆ YDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::YDataType = typename Problem::YDataType

◆ YSmoothScaleDataType

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
using ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::YSmoothScaleDataType = typename Problem::YSmoothScaleDataType

Member Function Documentation

◆ GetACoord()

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
CK_TILE_HOST_DEVICE auto ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GetACoord ( )
inlinestatic

◆ GetOCoord()

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
CK_TILE_HOST_DEVICE auto ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GetOCoord ( )
inlinestatic

◆ GetSmemSize()

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GetSmemSize ( )
inlinestaticconstexpr

◆ GetSmemSize_A()

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
CK_TILE_HOST_DEVICE constexpr ck_tile::index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GetSmemSize_A ( )
inlinestaticconstexpr

◆ operator()()

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
template<typename AWindow, typename GWindow, typename DWindow, typename OWindow>
CK_TILE_DEVICE auto ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::operator() ( const AWindow & a_window_,
const GWindow & g_window_,
const DWindow & d_window_,
OWindow & o_window_,
TopkWeightDataType ,
CK_TILE_LDS_ADDR void * smem,
index_t hidden_size,
index_t intermediate_size )
inline

Member Data Documentation

◆ GLD_A

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GLD_A = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GLD_A)
staticconstexpr

◆ GLD_B

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GLD_B = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GLD_B)
staticconstexpr

◆ GST_O

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::GST_O = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::GST_O)
staticconstexpr

◆ IsGateOnly

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
bool ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::IsGateOnly = Traits::IsGateOnly
staticconstexpr

◆ kAlignmentA

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::kAlignmentA = Policy::template GetAlignment_A<Problem>()
staticconstexpr

◆ kAlignmentD

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::kAlignmentD = Policy::template GetAlignment_D<Problem>()
staticconstexpr

◆ kAlignmentG

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::kAlignmentG = Policy::template GetAlignment_G<Problem>()
staticconstexpr

◆ kAlignmentO

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::kAlignmentO = Policy::template GetAlignment_O<Problem>()
staticconstexpr

◆ kBlockPerCu

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::kBlockPerCu
staticconstexpr
Initial value:
= []() {
if constexpr(Problem::kBlockPerCu != -1)
return Problem::kBlockPerCu;
else
{
return 2;
}
}()

◆ name

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
const char* ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::name = "fused_moe_flatmm"
staticconstexpr

◆ PadHiddenSize

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
bool ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::PadHiddenSize = Traits::PadHiddenSize
staticconstexpr

◆ PadIntermediateSize

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
bool ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::PadIntermediateSize = Traits::PadIntermediateSize
staticconstexpr

◆ SLD_A

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
index_t ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::SLD_A = static_cast<index_t>(FusedMoeGemmPipelineSequencerEnum::SLD_A)
staticconstexpr

◆ UseSmoothQuant

template<typename Problem_, typename Policy_ = FusedMoeGemmPipelineFlatmmPolicy>
bool ck_tile::FusedMoeGemmPipeline_FlatmmEx< Problem_, Policy_ >::UseSmoothQuant = Traits::UseSmoothQuant
staticconstexpr
The documentation for this struct was generated from the following file: