From 40c83c769ab9086386a617dd3d0fef1002f981b6 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 15:09:57 +0000
Subject: [PATCH] =?UTF-8?q?Fix:=20remove=20=C3=972=20QK=20scale=20correcti?=
 =?UTF-8?q?on=20(MMA=20scale=20is=201.0,=20not=200.5)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

---
 tests/unit/test_fmha_hd64_smem_p.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/tests/unit/test_fmha_hd64_smem_p.cu b/tests/unit/test_fmha_hd64_smem_p.cu
index 06ccb152..43b1dcd6 100644
--- a/tests/unit/test_fmha_hd64_smem_p.cu
+++ b/tests/unit/test_fmha_hd64_smem_p.cu
@@ -114,8 +114,8 @@ test_fmha_hd64_smem_p(const bf16_t* q, const bf16_t* k, const bf16_t* v,
                 : "r"(tb + n*8));
             asm volatile("tcgen05.wait::ld.sync.aligned;");
             if (lane == 0) for (int c=0;c<8;c++) {
-                s_vals[n*8+c] = tmp[c] * scale * 2.0f;  // ×2 to undo QK 0.5 scale
-                row_max = fmaxf(row_max, tmp[c] * scale * 2.0f);
+                s_vals[n*8+c] = tmp[c] * scale;  // No ×2 correction needed for QK MMA
+                row_max = fmaxf(row_max, tmp[c] * scale);
             }
         }
         row_max = wmax(row_max);