From efc786f6a3e7dda72ca36d9b7d89b3073765a0aa Mon Sep 17 00:00:00 2001
From: Qianfeng Zhang <Qianfeng.Zhang@amd.com>
Date: Fri, 18 Apr 2025 10:05:57 +0000
Subject: [PATCH] Remove un-needed __builtin_amdgcn_sched_barrier(0)

---
 .../ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp  | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp
index 9f5a890002..c458054010 100644
--- a/example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp
+++ b/example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp
@@ -481,10 +481,7 @@ struct HstuAttentionFwdPipelineQRKSVS
 
             // the over-lap only occurs when k1_loops is 3/5/7, NumVLdsBuffers is 2
             if constexpr(Policy::template IsFirstKLdsBufferOverlapLastVLdsBuffer<Problem>())
-            {
-                __builtin_amdgcn_sched_barrier(0);
                 __builtin_amdgcn_s_barrier();
-            };
         } while(++i_loop < num_loops);
 
         o_acc = tile_elementwise_in(o_acc_element_func, o_acc);