From ecf6a86d884bbf1429c4f7241bd538ce174c8e88 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Mon, 21 Jul 2025 09:04:49 +0000 Subject: [PATCH] Correct some comments --- .../18_hstu_attention/hstu_attention_fwd_pipeline.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 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 97b0656f5c..1d9691da65 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 @@ -207,7 +207,7 @@ struct HstuAttentionFwdPipelineQRKSVS __builtin_amdgcn_sched_barrier(0); - // K tile in LDS + // Q tile in LDS QKVDataType* q_lds_ptr = static_cast(smem_ptr); auto q_lds = make_tensor_view( q_lds_ptr, Policy::template MakeQLdsBlockDescriptor()); @@ -392,9 +392,8 @@ struct HstuAttentionFwdPipelineQRKSVS // ds_load for K __builtin_amdgcn_sched_group_barrier(0x00000100, 1, 0); - // buffer_load for V & K - __builtin_amdgcn_sched_group_barrier( - 0x00000020, 1, 0); // buffer_load for K & V + // buffer_load for V & K + __builtin_amdgcn_sched_group_barrier(0x00000020, 1, 0); }); static_for<0, K_LDS_READ_INST - (V_VMEM_LOAD_INST + K_VMEM_LOAD_INST), 1>{}(