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) {