From eed91c3984748372793e58eb80e90596597f8711 Mon Sep 17 00:00:00 2001
From: Philip Maybank <pmaybank@amd.com>
Date: Fri, 3 Oct 2025 10:50:14 -0400
Subject: [PATCH] change block_sync_lds function to be consistent with gfx12
 path in Old CK

---
 include/ck_tile/core/arch/arch.hpp | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp
index 28ded5439a..f90ec928a8 100644
--- a/include/ck_tile/core/arch/arch.hpp
+++ b/include/ck_tile/core/arch/arch.hpp
@@ -191,7 +191,12 @@ CK_TILE_DEVICE void s_waitcnt_barrier()
 template <index_t lgkmcnt = 0>
 CK_TILE_DEVICE void block_sync_lds()
 {
-    s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
+    // s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
+    asm volatile("\
+    s_wait_dscnt 0x0 \n \
+    s_barrier_signal -1 \n \
+    s_barrier_wait -1 \
+    " ::);
 }
 
 template <index_t vmcnt = 0>