ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-04 21:51:28 +00:00
Code Issues Packages Projects Releases Wiki Activity

Eliminate warning caused by failed to meet occupancy requirement (#2389)

Browse Source 
Co-authored-by: felix <felix.li@amd.com>
This commit is contained in:
Po Yen Chen
2025-07-08 09:17:25 +08:00
committed by GitHub
parent f240ae3248
commit b2dea90116
1 changed files with 3 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp
View File

@@ -122,6 +122,9 @@ struct BlockFmhaBatchPrefillPipelineQRKSVSAsync
{
if constexpr(kPadSeqLenK && BiasEnum == BlockAttentionBiasEnum::ELEMENTWISE_BIAS)
return 1;
// use larger K/V LDS buffer size will lower the occupancy
else if constexpr(64 <= kK0 || 64 <= kK1)
return 1;
else
return 2;
}