From 08bae35abd66e52a3471b5502f1cfeb20dd9b4e9 Mon Sep 17 00:00:00 2001 From: Feng Shijie Date: Tue, 9 Sep 2025 04:27:33 +0000 Subject: [PATCH] fix backward compatibility --- .../moe_flatmm/kernel/moe_flatmm_kernel.hpp | 46 +++++++++++++++++++ 1 file changed, 46 insertions(+) diff --git a/include/ck_tile/ops/moe_flatmm/kernel/moe_flatmm_kernel.hpp b/include/ck_tile/ops/moe_flatmm/kernel/moe_flatmm_kernel.hpp index 0f798e1088..d840283f0a 100644 --- a/include/ck_tile/ops/moe_flatmm/kernel/moe_flatmm_kernel.hpp +++ b/include/ck_tile/ops/moe_flatmm/kernel/moe_flatmm_kernel.hpp @@ -31,6 +31,52 @@ struct MoeFlatmmHostArgs : ScaleFlatmmHostArgs ExpertBias exp_bias; CK_TILE_HOST MoeFlatmmHostArgs() noexcept = default; + + CK_TILE_HOST MoeFlatmmHostArgs(const ck_tile::index_t* p_sorted_token_ids_, + const void* p_sorted_expert_weights_, + const ck_tile::index_t* p_sorted_expert_ids_, + const ck_tile::index_t* p_max_token_id_, + const void* a_ptr_, + const void* b_ptr_, + void* c_ptr_, + ck_tile::index_t NumTokens_, + ck_tile::index_t NumExperts_, + ck_tile::index_t TopK_, + ck_tile::index_t k_batch_, + ck_tile::index_t M_, + ck_tile::index_t N_, + ck_tile::index_t K_, + ck_tile::index_t stride_A_, + ck_tile::index_t stride_B_, + ck_tile::index_t stride_C_, + ScaleM scale_m_ = {}, + ScaleN scale_n_ = {}, + ExpertBias exp_bias_ = {}) + : MoeFlatmmHostArgs(p_sorted_token_ids_, + p_sorted_expert_weights_, + p_sorted_expert_ids_, + p_max_token_id_, + a_ptr_, + b_ptr_, + c_ptr_, + NumTokens_, + NumExperts_, + TopK_, + k_batch_, + M_, + N_, + K_, + stride_A_, + stride_B_, + stride_C_, + 0, // n_padded_zeros_ + 0, // k_padded_zeros_ + scale_m_, + scale_n_, + exp_bias_) + { + } + CK_TILE_HOST MoeFlatmmHostArgs(const ck_tile::index_t* p_sorted_token_ids_, const void* p_sorted_expert_weights_, const ck_tile::index_t* p_sorted_expert_ids_,