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

debug: print TMEM values after MMA

Browse Source
This commit is contained in:
biondizzle
2026-05-28 08:38:08 +00:00
parent bfb1e177ce
commit a9d71ff6ab
1 changed files with 13 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
dsv4/kernels/attention/fmha_qk_verify.cuh
View File

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