biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add register-math PV reference for HD=64 debug

Browse Source
This commit is contained in:
biondizzle
2026-05-28 15:13:47 +00:00
parent 43e9efbc2b
commit 5c9d471162
1 changed files with 15 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
tests/unit/test_fmha_hd64_smem_p.cu
View File

@@ -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);
}