From 34edc4391c752c5a9d1b205d9f2cd514abdcf2d2 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Mon, 21 Jul 2025 07:34:54 +0000 Subject: [PATCH] Fix in using sched_group_barrier() --- .../ck_tile/18_hstu_attention/hstu_attention_fwd_pipeline.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 01ef3cc496..6667019f78 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 @@ -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); });