From 0672373e51d62ebb6d52e5b4ee342f9332e503b9 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 13:08:06 +0000 Subject: [PATCH] =?UTF-8?q?test:=20debug=20=E2=80=94=20just=20QK+softmax+P?= =?UTF-8?q?=20read=20(no=20PV)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- tests/unit/test_fmha_hd16.cu | 15 ++------------- 1 file changed, 2 insertions(+), 13 deletions(-) diff --git a/tests/unit/test_fmha_hd16.cu b/tests/unit/test_fmha_hd16.cu index e290a0a2..efa452fa 100644 --- a/tests/unit/test_fmha_hd16.cu +++ b/tests/unit/test_fmha_hd16.cu @@ -86,11 +86,8 @@ test_fmha_hd16(const bf16_t* q, const bf16_t* k, const bf16_t* v, } __syncthreads(); - // STEP 3: PV — register math (decode T=1) - // O[d] = Σ P[0,j] × V[d,j] for j=0..SK-1 - // P is already in s_vals (warp 0, lane 0). Read V from GMEM. + // STEP 3: Just read P from TMEM and write to output (debug) if (wid == 0) { - // Re-read P from TMEM (already written back) float p_vals[SK]; for (int n = 0; n < SK / 8; n++) { float tmp[8]; @@ -98,15 +95,7 @@ test_fmha_hd16(const bf16_t* q, const bf16_t* k, const bf16_t* v, asm volatile("tcgen05.wait::ld.sync.aligned;"); if (lane == 0) for (int c=0;c<8;c++) p_vals[n*8+c] = tmp[c]; } - // Compute O[d] = Σ p[j] × V[d,j] - if (lane == 0) { - for (int d = 0; d < HD; d++) { - float ov = 0.0f; - for (int j = 0; j < SK; j++) - ov += p_vals[j] * bf16_to_f32(v[d * SK + j]); - o_out[d] = f32_to_bf16(ov); - } - } + if (lane == 0) for (int j=0;j