From 37fdb18e0619523ed8fcfe77fe170c4f726f536e Mon Sep 17 00:00:00 2001 From: felix Date: Sat, 25 Oct 2025 13:43:51 +0000 Subject: [PATCH] fix build --- .../67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp | 2 +- ...e_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp b/example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp index 8a3d3c30f7..e0838c221f 100644 --- a/example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp +++ b/example/67_gemm_microscaling/moe_gemm1_xdl_mx_fp4_bpreshuffle.cpp @@ -178,7 +178,7 @@ static constexpr auto GemmSpec = ck::tensor_operation::device::GemmSpecializatio constexpr ck::index_t DataPackedSize = 2; // Packed representation of data constexpr ck::index_t ScaleBlockSize = 32; // scaling block size -constexpr ck::index_t KPerBlock = 128 / DataPackedSize; // 256 f4 = 128 fp4x2 +constexpr ck::index_t KPerBlock = 256 / DataPackedSize; // 256 f4 = 128 fp4x2 static constexpr ck::index_t Nswizzle = false; static constexpr ck::index_t ActOP = 0; // 0: gelu_and_mul, 1: silu_and_mul static constexpr ck::index_t MPerBlock = 32; diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp index b3b3d312c7..6e25af0fcd 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_preshuffle_mx_moe_gufusion_v3.hpp @@ -726,8 +726,8 @@ struct BlockwiseGemmXdlops_pipeline_bpreshuffle_mx_moe_gufusion_v3< }); }); }); - - HotLoopScheduler(); + if constexpr(MPerBlock >= 64) + HotLoopScheduler(); __builtin_amdgcn_sched_barrier(0); };