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

test: minimal tmem_store debug (1 column + sentinels)

Browse Source
This commit is contained in:
biondizzle
2026-05-28 09:26:52 +00:00
parent f1aaa50326
commit c3b6c3a5e6
1 changed files with 6 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
tests/unit/test_umma_qk.cu
View File

@@ -71,23 +71,22 @@ test_umma_qk_hd16(
__syncthreads();
uint32_t tmem_base = *sTmemBase;
// Debug: write tmem_base to output
// Debug: write tmem_base to output BEFORE tmem_store
if (tid == 0) {
s_out[138] = 999.0f; // sentinel before tmem ops
s_out[138] = (float)tmem_base; // should be 0.0f if first alloc
s_out[139] = 1.0f; // sentinel: reached this point
}
__syncthreads();
// Zero TMEM — re-enable with debug
// Zero TMEM — only zero 1 column to test
if (wid == 0) {
for (int col = 0; col < 128; col++) {
tmem_store(tmem_base + col, 0, 0, 0, 0);
}
tmem_store(tmem_base, 0, 0, 0, 0);
tmem_fence_store();
}
__syncthreads();
if (tid == 0) {
s_out[138] = (float)tmem_base; // write tmem_base after tmem ops
s_out[140] = 2.0f; // sentinel: survived tmem_store
}
__syncthreads();