From 1ef76a62eaa57a6b0675ce2ccf51887f08d38c99 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Sun, 21 Dec 2025 12:14:39 +0000 Subject: [PATCH] Fix the static_assert expression in the pipeline --- .../pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp | 2 +- .../block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_trload.hpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) 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 0720802d86..ba826196be 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 @@ -167,7 +167,7 @@ struct BlockFmhaPipelineQRKSVSWholeKPrefetch "wrong!"); static_assert(kM0 == QDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && - kK1 == KDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && + kN0Sub == KDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && kQKHeaddim == KDramBlockWindowTmp{}.get_window_lengths()[number<1>{}] && kN1 == VDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && kK1 == VDramBlockWindowTmp{}.get_window_lengths()[number<1>{}] && diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_trload.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_trload.hpp index 3ca4734ac9..65c1af72de 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_trload.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_trload.hpp @@ -171,7 +171,7 @@ struct BlockFmhaPipelineQRKSVSWholeKPrefetchTrLoad "wrong!"); static_assert(kM0 == QDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && - kK1 == KDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && + kN0Sub == KDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && kQKHeaddim == KDramBlockWindowTmp{}.get_window_lengths()[number<1>{}] && kN1 == VDramBlockWindowTmp{}.get_window_lengths()[number<0>{}] && kK1 == VDramBlockWindowTmp{}.get_window_lengths()[number<1>{}] &&