From 09a2e029b96bd94459db7b6ea5666f71c1667a2c Mon Sep 17 00:00:00 2001
From: joye <joye@amd.com>
Date: Wed, 11 Jun 2025 12:44:03 +0800
Subject: [PATCH] add lds barrier

---
 .../device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp  | 1 +
 1 file changed, 1 insertion(+)

diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
index db4852c9e0..b153202fcd 100644
--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
+++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp
@@ -511,6 +511,7 @@ __global__ void kernel_grouped_conv_bwd_data_optimized_v2(const ABDataType* __re
 #pragma unroll
     for(int batch_id = 0; batch_id < BatchPerBlock; batch_id += WaveNum)
     {
+        block_sync_lds();
         for(int i = tid; i < InLoopNum; i += BlockSize)
         {
             int group_id          = i % GroupPerBlockInFP4;