ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-13 17:55:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

Change in GetKVBlockGemm to let gemm1 to use WarpTile-16x16x16/32x32x8 on mi350

Browse Source
This commit is contained in:
Qianfeng Zhang
2025-12-05 02:04:27 +00:00
parent 98f9b4a47b
commit c32949b285
1 changed files with 0 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp
View File

@@ -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))