From c32949b28537020e5b1dd822ec4e5ca359da4401 Mon Sep 17 00:00:00 2001 From: Qianfeng Zhang Date: Fri, 5 Dec 2025 02:04:27 +0000 Subject: [PATCH] Change in GetKVBlockGemm to let gemm1 to use WarpTile-16x16x16/32x32x8 on mi350 --- ...ha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp | 6 ------ 1 file changed, 6 deletions(-) diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp index e5a45afeea..f15ac2f08d 100644 --- a/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp +++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp @@ -525,15 +525,9 @@ struct BlockFmhaPipelineQRKSVSWholeKPrefetchDefaultPolicy constexpr index_t WarpGemmK = Problem::BlockFmhaShape::Gemm1WarpTile::at(number<2>{}); -#ifdef __gfx950__ - static_assert((WarpGemmM == 16 && WarpGemmK == 32) || - (WarpGemmM == 32 && WarpGemmK == 16), - "Not supported WarpGemm sizes!"); -#else static_assert((WarpGemmM == 16 && (WarpGemmK == 16 || WarpGemmK == 32)) || (WarpGemmM == 32 && (WarpGemmK == 8 || WarpGemmK == 16)), "Not supported WarpGemm sizes!"); -#endif if constexpr((WarpGemmM == 16 && WarpGemmK == 32) || (WarpGemmM == 32 && WarpGemmK == 16))