From d20ed807e42ae720d1d1a01315585f3c2ded3591 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Tue, 24 Feb 2026 07:53:50 +0000 Subject: [PATCH] Fix in comments --- include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp | 2 +- .../block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp | 2 +- include/ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp | 11 ++++++----- 3 files changed, 8 insertions(+), 7 deletions(-) diff --git a/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp b/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp index 2dd844e18b..871273cca8 100644 --- a/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp +++ b/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp @@ -87,7 +87,7 @@ struct has_naive_hdim_load_flag< template static inline constexpr bool is_naive_hdim_load_v = has_naive_hdim_load_flag::value; -// A helper struct for detechting kUseTrLoad +// A helper struct for detecting kUseTrLoad template struct has_use_trload_flag : std::false_type { diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp index 9ca077f643..de52b84335 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp @@ -318,7 +318,7 @@ struct BlockFmhaPipelineQRKSVSWholeKPrefetch Policy::template MakeBiasDramTileDistribution()); // assuming no random values need be saved, this is true when the pipeline is called from - // xformers, since we have a separate kernel to generated randomm values + // xformers, since we have a separate kernel to generatd random values auto null_randval_window = [&]() { if constexpr(kHasDropout) { diff --git a/include/ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp b/include/ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp index 230aa1fa33..4f07c3da27 100644 --- a/include/ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp +++ b/include/ck_tile/ops/fmha/pipeline/tile_fmha_shape.hpp @@ -52,9 +52,10 @@ struct TileFmhaShape static constexpr index_t kM0 = BlockTile::at(number<0>{}); // tile size along q seqlen static constexpr index_t kN0 = BlockTile::at(number<1>{}); // tile size along k seqlen static constexpr index_t kK0 = BlockTile::at(number<2>{}); // tile size along qk gemm unroll - static constexpr index_t kN0Sub = BlockTile::at(number<2>{}); // tile size for dividing kN0 - static constexpr index_t kN1 = BlockTile::at(number<3>{}); // tile size along v head_dim - static constexpr index_t kK1 = BlockTile::at(number<4>{}); // tile size along kv gemm unroll + static constexpr index_t kN0Sub = BlockTile::at( + number<2>{}); // same index as kK0; used as subdivision factor when dividing kN0 + static constexpr index_t kN1 = BlockTile::at(number<3>{}); // tile size along v head_dim + static constexpr index_t kK1 = BlockTile::at(number<4>{}); // tile size along kv gemm unroll static constexpr index_t kQKHeaddim = BlockTile::at(number<5>{}); // total length of K0, used for pipeline that need load Q at // once (or repeately load Q as a whole tile) @@ -65,8 +66,8 @@ struct TileFmhaShape // v, rowmajor : seqlen*hdim, colmajor : hdim*seqlen static constexpr bool IsVLayoutRowMajor = IsVLayoutRowMajor_; using VLayout = std::conditional_t; + ck_tile::tensor_layout::gemm::RowMajor, + ck_tile::tensor_layout::gemm::ColumnMajor>; }; template