From 8d30e46ba5431537cd183a3ea450839763d48b29 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Sun, 6 Jul 2025 14:31:36 +0000 Subject: [PATCH] Remove using i_loop and num_loops since seqlen_k_curr and seqlen_k_end is enough --- .../18_hstu_attention/hstu_attention_fwd_pipeline.hpp | 6 +----- 1 file changed, 1 insertion(+), 5 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 73c9b3fc30..501699b277 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 @@ -274,8 +274,6 @@ struct HstuAttentionFwdPipelineQRKSVS } }; - const auto num_loops = integer_divide_ceil(seqlen_k_end - seqlen_k_start, kN0); - const auto bias_origin = bias_dram_block_window_tmp.get_window_origin(); auto bias_dram_window = make_tile_window(bias_dram_block_window_tmp.get_bottom_tensor_view(), @@ -352,8 +350,6 @@ struct HstuAttentionFwdPipelineQRKSVS auto seqlen_k_curr = seqlen_k_start; - index_t i_loop = 0; - // ensure all q_reg_tiles[] have been loaded from LDS, so the LDS can be reused by k_tile __builtin_amdgcn_s_barrier(); @@ -481,7 +477,7 @@ struct HstuAttentionFwdPipelineQRKSVS }; } }); - } while(i_loop++ < num_loops); + } while(seqlen_k_curr < seqlen_k_end); tile_elementwise_inout( [&](auto& x) {