From 4f7be0f2b591ab451c28001cb5f651463a7d67f0 Mon Sep 17 00:00:00 2001
From: aska-0096 <haocwang@amd.com>
Date: Tue, 13 May 2025 08:37:11 +0000
Subject: [PATCH] disable the scheduler in v3; bring it back when compiler
 feature ready.

---
 example/65_gemm_multiply_multiply/CMakeLists.txt             | 4 ++--
 ...kwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp | 5 +++--
 2 files changed, 5 insertions(+), 4 deletions(-)

diff --git a/example/65_gemm_multiply_multiply/CMakeLists.txt b/example/65_gemm_multiply_multiply/CMakeLists.txt
index f38ac6fe4a..8d2cbb1c70 100644
--- a/example/65_gemm_multiply_multiply/CMakeLists.txt
+++ b/example/65_gemm_multiply_multiply/CMakeLists.txt
@@ -31,8 +31,8 @@ set(GEMM_OPTIONS)
 list(APPEND GEMM_OPTIONS "SHELL: -mllvm -greedy-reverse-local-assignment=1 -mllvm --slp-threshold=-32")
 list(APPEND GEMM_OPTIONS -v --save-temps -Wno-gnu-line-marker)
 set(BLOCKSCALE_GEMM_OPTIONS)
-list(APPEND BLOCKSCALE_GEMM_OPTIONS "SHELL: -mllvm -greedy-reverse-local-assignment=1 -mllvm --slp-threshold=-32 -mllvm --disable-schedmodel-in-sched-mi=1 -mllvm --amdgpu-sched-strategy=gcn-iterative-max-occupancy-experimental -mllvm --misched-bottomup=1")
-# list(APPEND BLOCKSCALE_GEMM_OPTIONS "SHELL: -mllvm -greedy-reverse-local-assignment=1 -mllvm --slp-threshold=-32")
+# list(APPEND BLOCKSCALE_GEMM_OPTIONS "SHELL: -mllvm -greedy-reverse-local-assignment=1 -mllvm --slp-threshold=-32 -mllvm --disable-schedmodel-in-sched-mi=1 -mllvm --amdgpu-sched-strategy=gcn-iterative-max-occupancy-experimental -mllvm --misched-bottomup=1")
+list(APPEND BLOCKSCALE_GEMM_OPTIONS "SHELL: -mllvm -greedy-reverse-local-assignment=1 -mllvm --slp-threshold=-32 -mllvm --misched-bottomup=1")
 list(APPEND BLOCKSCALE_GEMM_OPTIONS -v --save-temps -Wno-gnu-line-marker)
 target_compile_options(example_gemm_multiply_multiply_xdl_fp8_bpreshuffle PRIVATE ${GEMM_OPTIONS})
 target_compile_options(example_moe_gemm1_xdl_fp8 PRIVATE ${GEMM_OPTIONS})
diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
index 4b7e255934..365f91767d 100644
--- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
+++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_blockscale_b_preshuffle_v3.hpp
@@ -784,7 +784,8 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                                                  b_scale_thread_bufs[mfma_reg_buf][I0];
                     });
 
-                    HotLoopScheduler();
+                    // We need new compiler to enable this feature
+                    // HotLoopScheduler();
                     __builtin_amdgcn_sched_barrier(0);
                 };
 
@@ -933,7 +934,7 @@ struct BlockwiseGemmXdlops_pipeline_blockscale_bpreshuffle_v3<BlockGemmPipelineS
                 }
             });
 
-            HotLoopScheduler();
+            // HotLoopScheduler();
 
             static_for<0, MRepeat, 1>{}([&](auto m0) {
                 c_scale_thread_buf(m0) = a_scale_thread_bufs[I0][m0] * b_scale_thread_bufs[I0][I0];