/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_uk.hpp Source File#
fused_moegemm_pipeline_flatmm_uk.hpp
Go to the documentation of this file.
126 CK_TILE_DEVICE auto GetRowID(const ROW_COORDS coords, const IndexDataType* sorted_token_ids_ptr)
Definition: cluster_descriptor.hpp:13
CK_TILE_DEVICE auto cmp_lt_to_exec(const X &x, const Y &y)
Definition: utility.hpp:133
CK_TILE_DEVICE auto tile_elementwise_in(const InElementFunc &in_element_func, const InTensor &... in_dstr_tensors)
Definition: tile_elementwise.hpp:40
remove_cv_t< std::remove_reference_t< T > > remove_cvref_t
Definition: type_traits.hpp:21
CK_TILE_DEVICE auto make_tile_window_linear_raw(const TensorView_ &tensor_view, const WindowLengths_ &window_lengths, const multi_index< TensorView_::get_num_of_dimension()> &origin, const StaticTileDistribution_ &tile_distribution, LinearBottomDims_={})
Definition: tile_window_linear.hpp:1029
constexpr CK_TILE_DEVICE auto make_tile_window_linear(const TensorView_ &tensor_view, const WindowLengths_ &window_lengths, const multi_index< TensorView_::get_num_of_dimension()> &origin, const StaticTileDistribution_ &tile_distribution, LinearBottomDims_={})
Definition: tile_window_linear.hpp:993
constexpr CK_TILE_HOST_DEVICE auto generate_tuple(F &&f, number< N >)
Definition: tuple.hpp:429
constexpr CK_TILE_HOST_DEVICE auto make_tuple(Xs &&... xs)
Definition: tuple.hpp:360
CK_TILE_DEVICE void store_tile(tile_window_with_static_lengths< BottomTensorView_, WindowLengths_ > &tile_window_tmp, const static_distributed_tensor< DataType_, TileDistribution_ > &dstr_tensor)
Definition: store_tile.hpp:23
constexpr CK_TILE_HOST_DEVICE void sweep_tile(const F &f, UnpacksPerXDim={})
Definition: sweep_tile.hpp:231
CK_TILE_DEVICE int32x4_t make_wave_buffer_resource(const void *ptr, uint32_t size=0xffffffff)
Definition: amd_buffer_addressing.hpp:40
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:23
typename Problem::IndexDataType IndexDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:39
typename Problem::ADataType ADataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:29
static constexpr index_t kAlignmentO
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:52
constexpr CK_TILE_DEVICE auto GetNumRowCoords_A()
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:101
typename Problem::DDataType DDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:31
static constexpr bool PadIntermediateSize
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:47
CK_TILE_DEVICE auto operator()(const Karg &kargs, CK_TILE_LDS_ADDR void *smem, index_t sorted_tile_id, index_t intermediate_tile_id)
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:170
static constexpr const char * name
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:69
CK_TILE_DEVICE auto GetRowCoords_A(index_t base_offset)
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:111
static CK_TILE_HOST_DEVICE auto GetOCoord()
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:94
typename Problem::DScaleDataType DScaleDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:36
static constexpr CK_TILE_HOST_DEVICE ck_tile::index_t GetSmemSize()
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:71
static constexpr index_t kAlignmentD
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:51
static constexpr index_t kAlignmentA
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:49
static constexpr index_t GLD_B
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:56
static constexpr index_t kBlockPerCu
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:59
typename Problem::GScaleDataType GScaleDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:35
typename Problem::TopkWeightDataType TopkWeightDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:38
static constexpr bool IsGateOnly
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:44
typename Problem::ODataType ODataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:33
typename Problem::BlockShape BlockShape
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:27
static constexpr index_t kAlignmentG
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:50
static constexpr bool PadHiddenSize
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:46
static constexpr index_t GST_O
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:57
static constexpr index_t SLD_A
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:54
CK_TILE_DEVICE auto GetWeightScale(const ROW_COORDS coords, const TopkWeightDataType *sorted_weight_ptr)
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:142
static constexpr index_t GLD_A
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:55
remove_cvref_t< Problem_ > Problem
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:24
typename Problem::YSmoothScaleDataType YSmoothScaleDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:37
CK_TILE_DEVICE auto GetRowID(const ROW_COORDS coords, const IndexDataType *sorted_token_ids_ptr)
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:126
typename Problem::GDataType GDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:30
typename Problem::YDataType YDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:40
CK_TILE_DEVICE auto GetRowCoords_O(index_t base_offset)
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:156
remove_cvref_t< Policy_ > Policy
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:25
static constexpr bool UseSmoothQuant
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:45
typename Problem::Traits Traits
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:42
typename Problem::AccDataType AccDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:32
typename Problem::AScaleDataType AScaleDataType
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:34
static CK_TILE_HOST_DEVICE auto GetACoord()
Definition: fused_moegemm_pipeline_flatmm_uk.hpp:86
A fixed-size array container similar to std::array with additional utilities.
Definition: array.hpp:43
Definition: integral_constant.hpp:13
Definition: sequence.hpp:49
Definition: functional.hpp:43