fix performance regression on blockgemm v3 pipe

031ddf35 · coderfeli · 400cac28 · 031ddf35
Commit 031ddf35 authored Dec 27, 2024 by coderfeli
Show whitespace changes
Inline Side-by-side

Showing with 6 additions and 0 deletions

include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp ...operation/gpu/block/blockwise_gemm_pipeline_xdlops_v3.hpp +6 -0

No files found.
--- 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) {