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

Tiny movement in the code lines of the pipeline

Browse Source
This commit is contained in:
Qianfeng Zhang
2025-06-26 10:09:51 +00:00
parent 60d8ffb6cf
commit 3c300d3069
1 changed files with 3 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -368,7 +368,9 @@ struct HstuAttentionFwdPipelineQRKSVS
move_tile_window(v_dram_window, {0, kK1});
__builtin_amdgcn_sched_barrier(0);
// for i_k1 = k1_loop-1, the loading is for next iteration
k_tile = load_tile(k_dram_window);
move_tile_window(k_dram_window, {kK1, 0});
block_sync_lds();
// execute current unroll of gemm_0
@@ -395,14 +397,6 @@ struct HstuAttentionFwdPipelineQRKSVS
tile_elementwise_inout([&scale_s](auto& x) { x = x * scale_s; }, sacc_tile);
}
__builtin_amdgcn_sched_barrier(0);
// for i_k1 = k1_loop-1, the loading is for next iteration
k_tile = load_tile(k_dram_window);
move_tile_window(k_dram_window, {kK1, 0});
__builtin_amdgcn_sched_barrier(0);
if constexpr(HstuMask::IsMasking)
{
if constexpr(HstuMask::kUseLocal)