From b75077475bd18c2437a5d1e4c6bd0d6d9fde4e80 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Sat, 15 Nov 2025 13:48:00 +0000 Subject: [PATCH] Remove useless codes in the two trload pipelines --- .../hstu_attention_no_softmax_fwd_trload_pipeline.hpp | 10 +--------- ...hstu_attention_with_softmax_fwd_trload_pipeline.hpp | 10 +--------- 2 files changed, 2 insertions(+), 18 deletions(-) diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_trload_pipeline.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_trload_pipeline.hpp index cd360041b1..b887dd832d 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_trload_pipeline.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_no_softmax_fwd_trload_pipeline.hpp @@ -70,9 +70,6 @@ struct HstuAttentionNoSoftmaxFwdPipelineQRKSVSTrLoad static constexpr index_t kAlignmentBias = kPadSeqLenK ? 1 : Policy::template GetAlignmentBias(); - static constexpr index_t kGemmSingleRepM = Policy::template GetQKBlockGemmSingleRepM(); - static constexpr index_t kGemmNumRepM = kM0 / kGemmSingleRepM; - // used by NRepetitions2DEpilogue static constexpr index_t kGemm1SingleRepN = Policy::template GetKVBlockGemmSingleRepN(); @@ -318,11 +315,6 @@ struct HstuAttentionNoSoftmaxFwdPipelineQRKSVSTrLoad return make_null_tile_window(make_tuple(number<1>{}, number<1>{})); }(); - using q_tile_type = decltype(make_static_distributed_tensor( - Policy::template MakeQRegTileDistribution())); - - q_tile_type q_tile; - store_tile(q_lds_write_window, q_dram_tile); clear_tile(o_acc); @@ -331,7 +323,7 @@ struct HstuAttentionNoSoftmaxFwdPipelineQRKSVSTrLoad block_sync_lds(); - q_tile = load_tile(q_lds_read_window); + auto q_tile = load_tile(q_lds_read_window); q_tile = tile_elementwise_in(q_element_func, q_tile); diff --git a/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_trload_pipeline.hpp b/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_trload_pipeline.hpp index c4647f25d5..94135d9e84 100644 --- a/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_trload_pipeline.hpp +++ b/example/ck_tile/18_hstu_attention/hstu_attention_with_softmax_fwd_trload_pipeline.hpp @@ -70,9 +70,6 @@ struct HstuAttentionWithSoftmaxFwdPipelineQRKSVSTrLoad static constexpr index_t kAlignmentBias = kPadSeqLenK ? 1 : Policy::template GetAlignmentBias(); - static constexpr index_t kGemmSingleRepM = Policy::template GetQKBlockGemmSingleRepM(); - static constexpr index_t kGemmNumRepM = kM0 / kGemmSingleRepM; - // used by NRepetitions2DEpilogue static constexpr index_t kGemm1SingleRepN = Policy::template GetKVBlockGemmSingleRepN(); @@ -332,11 +329,6 @@ struct HstuAttentionWithSoftmaxFwdPipelineQRKSVSTrLoad return make_null_tile_window(make_tuple(number<1>{}, number<1>{})); }(); - using q_tile_type = decltype(make_static_distributed_tensor( - Policy::template MakeQRegTileDistribution())); - - q_tile_type q_tile; - store_tile(q_lds_write_window, q_dram_tile); clear_tile(o_acc); @@ -345,7 +337,7 @@ struct HstuAttentionWithSoftmaxFwdPipelineQRKSVSTrLoad block_sync_lds(); - q_tile = load_tile(q_lds_read_window); + auto q_tile = load_tile(q_lds_read_window); q_tile = tile_elementwise_in(q_element_func, q_tile);