From d408192bfefd1cce723c083b15cae5e91d359f98 Mon Sep 17 00:00:00 2001 From: Enrico Degregori Date: Wed, 1 Jul 2026 13:34:22 +0000 Subject: [PATCH] Instruction scheduling in block gemm instead of scheduler Compiler generates better assembly decreasing spilling for large tiles AB Quant is unchanged --- .../block/block_gemm_areg_breg_creg_eight_waves_v1.hpp | 6 ++++++ .../gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp | 8 -------- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_eight_waves_v1.hpp b/include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_eight_waves_v1.hpp index 9f91c06e8e..0c5fee5c67 100644 --- a/include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_eight_waves_v1.hpp +++ b/include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_eight_waves_v1.hpp @@ -248,6 +248,12 @@ struct BlockGemmARegBRegCRegEightWavesV1 merge_sequences(c_iter_idx{}, c_warp_y_index_zeros), merge_sequences(sequence<1, 1>{}, c_warp_y_lengths), c_warp_tensor.get_thread_buffer()); + __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); + if constexpr(nIter == 0 && mIter == MIterPerWarp - 1 && kIter == 0) + { + s_waitcnt_lgkm<4>(); + __builtin_amdgcn_sched_barrier(0); + } }); }); } diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp index 5e46be1aaa..9ac4cc2042 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp @@ -183,14 +183,6 @@ struct GemmPipelineAgBgCrCompAsyncEightWaves : public BaseGemmPipelineAgBgCrComp // Hot loop scheduler // ------------------ auto hot_loop_scheduler = [&]() { - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - s_waitcnt_lgkm<4>(); - __builtin_amdgcn_sched_group_barrier(0x004, 1, 0); // lgkmcnt / SALU - static_for<0, MFMA_INST - 3, 1>{}([&](auto) { - __builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA - }); __builtin_amdgcn_sched_barrier(0); };