From a2d5fd01eb05d2aa477fb618e69bb2b7d7e31428 Mon Sep 17 00:00:00 2001
From: Qianfeng <qianfeng.zhang@amd.com>
Date: Fri, 14 Feb 2025 12:44:32 +0800
Subject: [PATCH] Tiny Fix (#1888)

[ROCm/composable_kernel commit: 4cfb24feb67602d38b60a1568492c6313bf25a82]
---
 .../block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp
index 3726cd433c..3d53535b28 100644
--- a/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp
+++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp
@@ -343,6 +343,8 @@ struct BlockFmhaFwdSplitKVPipelineNWarpSShuffleQRKSVS
         // moving k_dram_window is an in-page-block operation, so there is
         // no need to invoke k_page_block_navigator.move_tile_window() here.
         move_tile_window(k_dram_window, {0, kK0});
+        // ensure LDS access by Q is done before the over-writting by K
+        block_sync_lds();
         store_tile(k_lds_window, tile_elementwise_in(k_element_func, k_block_tile));
 
         do