ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-02 21:27:45 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix in using sched_group_barrier()

Browse Source
This commit is contained in:
Qianfeng Zhang
2025-07-21 07:34:54 +00:00
parent 1caef1fb89
commit 34edc4391c
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -387,13 +387,13 @@ struct HstuAttentionFwdPipelineQRKSVS
__builtin_amdgcn_sched_group_barrier(0x00000020, V_VMEM_LOAD_INST, 0);
__builtin_amdgcn_sched_group_barrier(0x00000100, K_LDS_READ_INST, 0);
__builtin_amdgcn_sched_group_barrier(0x00000100, 1, 0);
__builtin_amdgcn_sched_group_barrier(0x00000020, K_VMEM_LOAD_INST, 0);
static_for<0, K_LDS_READ_INST - 1, 1>{}([&](auto i) {
ignore = i;
__builtin_amdgcn_sched_group_barrier(0x00000100, K_LDS_READ_INST, 0);
__builtin_amdgcn_sched_group_barrier(0x00000100, 1, 0);
__builtin_amdgcn_sched_group_barrier(0x00000008, kGemmNumRepM, 0);
});