From 1b3d08bd1b2b79571d2693cda810d8158c7066d6 Mon Sep 17 00:00:00 2001 From: "BingYuan.Zhou" Date: Wed, 14 May 2025 09:31:26 +0800 Subject: [PATCH] fix moe sorting build fail (#2190) * fix moe sorting build fail * refile code --------- Co-authored-by: solin [ROCm/composable_kernel commit: 41c17d0a953a5c399a3cf15ff283d1b57992f06d] --- .../flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp | 3 ++- .../pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp index 2ff9d1ebf0..cbd20a6ea3 100644 --- a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp +++ b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp @@ -75,6 +75,7 @@ struct FlatmmPipelineAGmemBGmemCRegV1 CK_TILE_HOST_DEVICE static constexpr auto HotLoopScheduler() { +#if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8) || defined(USING_MFMA_32x32x16) constexpr auto config = BlockFlatmm::BlockPolicy::template GetWarpGemmMWarpNWarp(); using WG = remove_cvref_t())>; @@ -90,7 +91,7 @@ struct FlatmmPipelineAGmemBGmemCRegV1 constexpr index_t A_Buffer_Load_Inst_Num = kMPerBlock * kKPerBlock / BlockSize / KPerLoad; constexpr index_t A_LDS_Read_Inst_Num = MIterPerWarp * KIterPerWarp; constexpr index_t B_Buffer_Load_Inst_Num = NIterPerWarp * KIterPerWarp; - // constexpr index_t A_LDS_Read_Inst_Remain = A_LDS_Read_Inst_Num - A_Buffer_Load_Inst_Num; +#endif #if defined(USING_MFMA_16x16x32) && defined(ENABLE_FP8) static_for<0, A_Buffer_Load_Inst_Num, 1>{}([&](auto i) { ignore = i; diff --git a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp index 474924ec84..1a1b729394 100644 --- a/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp +++ b/include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1_policy.hpp @@ -112,8 +112,8 @@ struct UniversalFlatmmPipelineAgBgCrPolicy make_tuple(number{}, number{}))), make_tuple(sequence<1, 0>{}, sequence<2, 3>{}), make_tuple(sequence<0>{}, sequence<1>{})); + return a_lds_block_desc; #endif - return a_lds_block_desc; } template