diff --git a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp index 7544ad8120..629f0ee8f1 100644 --- a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp +++ b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_pipeline_flatmm_policy.hpp @@ -826,7 +826,7 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Block_M1 == 32 && S_::Block_N1 == 128 && S_::Block_K1 == 512 && S_::Warp_M0 == 16 && S_::Warp_N0 == 16 && S_::Warp_K0 == 32 && - T_::PipeInterleave == false) + T_::PipeInterleave == false) { return FlatmmSn_32x128x512_1x4x1_16x16x32_FP16{}; // return FlatmmSn_32x128x512_1x4x1_16x16x32_FP16_itl{}; @@ -836,17 +836,17 @@ struct FusedMoeGemmPipelineFlatmmPolicy std::is_same_v && S_::Block_M1 == 32 && S_::Block_N1 == 128 && S_::Block_K1 == 512 && S_::Warp_M0 == 16 && S_::Warp_N0 == 16 && S_::Warp_K0 == 32 && - T_::PipeInterleave == true) + T_::PipeInterleave == true) { // return FlatmmSn_32x128x512_1x4x1_16x16x32_FP16{}; return FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl{}; } - else if constexpr(std::is_same_v && + else if constexpr(std::is_same_v && std::is_same_v && std::is_same_v && S_::Block_M1 == 32 && S_::Block_N1 == 128 && S_::Block_K1 == 512 && S_::Warp_M0 == 16 && S_::Warp_N0 == 16 && S_::Warp_K0 == 32 && - T_::PipeInterleave == true) + T_::PipeInterleave == true) { // return FlatmmSn_32x128x512_1x4x1_16x16x32_FP16{}; return FlatmmSn_32x128x512_1x4x1_16x16x32_FP16_itl{}; diff --git a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_traits.hpp b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_traits.hpp index bb29fb23b3..3fb82bc099 100644 --- a/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_traits.hpp +++ b/include/ck_tile/ops/fused_moe/pipeline/fused_moegemm_traits.hpp @@ -23,7 +23,7 @@ template + bool PipeInterleave_ = true> struct FusedMoeGemmTraits { // Gate+Up or Gate only