From a9d71ff6ab843ef7eb5f3912ca97043703686e51 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 08:38:08 +0000
Subject: [PATCH] debug: print TMEM values after MMA

---
 dsv4/kernels/attention/fmha_qk_verify.cuh | 14 +++++++++++++-
 1 file changed, 13 insertions(+), 1 deletion(-)

diff --git a/dsv4/kernels/attention/fmha_qk_verify.cuh b/dsv4/kernels/attention/fmha_qk_verify.cuh
index b260c392..db146b42 100644
--- a/dsv4/kernels/attention/fmha_qk_verify.cuh
+++ b/dsv4/kernels/attention/fmha_qk_verify.cuh
@@ -82,7 +82,19 @@ fmha_qk_verify(
     if (wid == 0 && lane == 0) tmem_fence_store();
     __syncthreads();
 
-    // Read S row 0 from TMEM
+    // Read S row 0 from TMEM — try multiple column offsets
+    if (tid == 0) {
+        printf("[qk] tmem_base=%u, reading S...\n", tmem_base);
+        for (int col = 0; col < 4; col++) {
+            uint32_t u0, u1, u2, u3;
+            tmem_load(tmem_base + col, u0, u1, u2, u3);
+            printf("[qk] S[0,%d] raw: %f %f %f %f\n", col,
+                   u32_to_f32(u0), u32_to_f32(u1), u32_to_f32(u2), u32_to_f32(u3));
+        }
+    }
+    __syncthreads();
+
+    // Also read using the output array (same as before)
     if (wid == 0) {
         for (int col = lane; col < 128; col += WARP) {
             uint32_t u0, u1, u2, u3;