From a71dc1245f52b5351d9bf4d735219e005650122e Mon Sep 17 00:00:00 2001
From: Po Yen Chen <PoYen.Chen@amd.com>
Date: Tue, 8 Jul 2025 09:17:25 +0800
Subject: [PATCH] Eliminate warning caused by failed to meet occupancy
 requirement (#2389)

Co-authored-by: felix <felix.li@amd.com>

[ROCm/composable_kernel commit: b2dea90116d1060c67db5edddb6d4498188ebac4]
---
 .../block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp       | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp b/include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp
index 6398bf316e..10daea99d1 100644
--- a/include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp
+++ b/include/ck_tile/ops/fmha/pipeline/block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp
@@ -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;
             }