From 5c9d4711629129bbb1e6ab2f7026bef3480d5206 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 15:13:47 +0000
Subject: [PATCH] Add register-math PV reference for HD=64 debug

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

diff --git a/tests/unit/test_fmha_hd64_smem_p.cu b/tests/unit/test_fmha_hd64_smem_p.cu
index c19882bc..934a40ab 100644
--- a/tests/unit/test_fmha_hd64_smem_p.cu
+++ b/tests/unit/test_fmha_hd64_smem_p.cu
@@ -190,6 +190,21 @@ test_fmha_hd64_smem_p(const bf16_t* q, const bf16_t* k, const bf16_t* v,
         }
     }
 
+    // ===== Register-math PV for verification =====
+    float* s_ref_o = s_p_vals;  // Reuse s_p_vals (already consumed by PV)
+    if (tid == 0) {
+        float s_ref[SK];
+        for (int j=0;j<SK;j++) s_ref[j] = s_p_vals[j];
+        for (int d=0;d<HD;d++) {
+            float ov = 0.0f;
+            for (int j=0;j<SK;j++) ov += s_ref[j] * bf16_to_f32(v[d*SK+j]);
+            s_ref_o[d] = ov;
+        }
+    }
+    __syncthreads();
+    if (tid == 0) { printf("RegPV[0..7]: "); for(int d=0;d<8;d++) printf("%.6f ", s_ref_o[d]); printf("\n"); }
+    __syncthreads();
+
     if (wid == 0) tmem_dealloc(tb, 128);
 }