From 031ddf356d992c14dd81ce5394fb193eca5dd61e Mon Sep 17 00:00:00 2001
From: coderfeli <coderfeli@163.com>
Date: Fri, 27 Dec 2024 06:40:43 +0000
Subject: [PATCH] fix performance regression on blockgemm v3 pipe

---
 .../gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp         | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp
index a2ba7c70c3..e87616efe1 100644
--- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp
+++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp
@@ -477,6 +477,9 @@ struct BlockwiseGemmXdlops_pipeline_v3<BlockGemmPipelineScheduler::Intrawave,
                                        b_thread_buf_tail);
                 });
             });
+
+            HotLoopScheduler();
+            __builtin_amdgcn_sched_barrier(0);
         }
     }
 
@@ -692,6 +695,9 @@ struct BlockwiseGemmXdlops_pipeline_v3<BlockGemmPipelineScheduler::Intrawave,
                 });
             });
 
+            HotLoopScheduler();
+            __builtin_amdgcn_sched_barrier(0);
+
             static_for<0, KRepeat, 1>{}([&](auto k0) {
                 static_for<0, MRepeat, 1>{}([&](auto m0) {
                     static_for<0, NRepeat, 1>{}([&](auto n0) {