From fa6c124163c6b83a9a980f4a278305f7187e65fa Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 14:15:18 +0000 Subject: [PATCH] Debug: QK only, skip softmax+PV --- tests/unit/test_fmha_ts_full.cu | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/tests/unit/test_fmha_ts_full.cu b/tests/unit/test_fmha_ts_full.cu index 474491bf..1a876999 100644 --- a/tests/unit/test_fmha_ts_full.cu +++ b/tests/unit/test_fmha_ts_full.cu @@ -90,7 +90,8 @@ test_fmha_ts_full(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k, asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); __syncthreads(); - // ===== STEP 2: Softmax ===== + // ===== STEP 2: Softmax — SKIPPED FOR DEBUG ===== + /* if (wid == 0) { float s_vals[SK], row_max = -INFINITY; for (int n = 0; n < SK / 8; n++) { @@ -127,6 +128,7 @@ test_fmha_ts_full(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k, tmem_fence_store(); } __syncthreads(); + */ // ===== STEP 3: PV GEMM (TS) — SKIPPED FOR DEBUG ===== /* @@ -148,8 +150,8 @@ test_fmha_ts_full(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k, __syncthreads(); */ - // ===== STEP 4: Epilogue — read O from TMEM ===== - // MMA output is scaled by 0.5, so multiply by 2.0 + // ===== STEP 4: Epilogue — read S from TMEM (QK output) ===== + // Just verify QK works by reading first 16 values if (wid == 0) { float o_vals[HD]; for (int n = 0; n < HD / 8; n++) { @@ -157,9 +159,9 @@ test_fmha_ts_full(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k, asm volatile("tcgen05.ld.sync.aligned.32x32b.x8.b32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8];" : "=f"(tmp[0]),"=f"(tmp[1]),"=f"(tmp[2]),"=f"(tmp[3]), "=f"(tmp[4]),"=f"(tmp[5]),"=f"(tmp[6]),"=f"(tmp[7]) - : "r"(tb_o + n*8)); + : "r"(tb + n*8)); asm volatile("tcgen05.wait::ld.sync.aligned;"); - if (lane == 0) for (int c=0;c<8;c++) o_vals[n*8+c] = tmp[c] * 2.0f; + if (lane == 0) for (int c=0;c<8;c++) o_vals[n*8+c] = tmp[c]; } if (lane == 0) for (int d=0;d