ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 18:42:06 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove using i_loop and num_loops since seqlen_k_curr and seqlen_k_end is enough

Browse Source
This commit is contained in:
Qianfeng Zhang
2025-07-06 14:31:36 +00:00
parent 5451912526
commit 8d30e46ba5
1 changed files with 1 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
example/ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp
View File

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