From e80c99b672ed9caacd7467e0fae145bbff6681a3 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Wed, 17 Dec 2025 07:35:19 +0000 Subject: [PATCH] Remove useless call of __builtin_amdgcn_s_waitcnt(0xc07f) --- .../hstu_attention_no_softmax_fwd_pipeline.hpp | 2 -- .../hstu_attention_with_softmax_fwd_pipeline.hpp | 2 -- 2 files changed, 4 deletions(-) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_pipeline.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_pipeline.hpp index e3ac8b0deb..1d3035a43c 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_pipeline.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_pipeline.hpp @@ -336,8 +336,6 @@ struct HstuAttentionNoSoftmaxFwdPipelineQRKSVS q_reg_tiles[i_rep] = load_tile(q_lds_read_window); - __builtin_amdgcn_s_waitcnt(0xc07f); - // the following codes will not generate actual instructions by the compiler set_slice_tile(q_tile, q_reg_tiles[i_rep], diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_pipeline.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_pipeline.hpp index b8f1b7e5be..23c8ed4945 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_pipeline.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_pipeline.hpp @@ -353,8 +353,6 @@ struct HstuAttentionWithSoftmaxFwdPipelineQRKSVS q_reg_tiles[i_rep] = load_tile(q_lds_read_window); - __builtin_amdgcn_s_waitcnt(0xc07f); - // the following codes will not generate actual instructions by the compiler set_slice_tile(q_tile, q_reg_tiles[i_rep],