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

Add s_barrier to sync threads

Browse Source
This commit is contained in:
PoYen, Chen
2024-08-22 09:03:01 +00:00
parent 73378ff95d
commit d3f550f30c
1 changed files with 2 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/ops/fmha/pipeline/block_fmha_fwd_appendkv_pipeline.hpp
View File

@@ -145,6 +145,7 @@ struct BlockFmhaFwdAppendKVPipeline
rotary_sin_window,
rotary_dim,
thread_end);
__builtin_amdgcn_s_barrier();
}
if constexpr(kIsPagedKV)
@@ -222,6 +223,7 @@ struct BlockFmhaFwdAppendKVPipeline
BlockRotaryEmbedding<RotaryEnum>::apply(
q_tile, q_window, rotary_cos_window, rotary_sin_window, rotary_dim, thread_end);
__builtin_amdgcn_s_barrier();
store_tile(q_dram_block_window, q_tile);
}