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); };