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

test: skip MMA, just test descriptor values

Browse Source
This commit is contained in:
biondizzle
2026-05-28 09:20:59 +00:00
parent ea6b42e649
commit a6c0ce51a2
1 changed files with 11 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

22
tests/unit/test_umma_qk.cu
View File

@@ -106,22 +106,22 @@ test_umma_qk_hd16(
s_out[133] = (float)sQ_smem;
s_out[134] = (float)sK_smem;
s_out[135] = (float)__cvta_generic_to_shared(sQ);
s_out[136] = (float)(sQ_smem >> 4); // start_address for desc
s_out[137] = (float)(sQ_smem & 0xF); // alignment check
}
__syncthreads();
// ================================================================
// Call tcgen05.mma SS
// SKIP MMA for now — just test descriptor construction
// ================================================================
if (tid == 0) {
umma_ss_f16(tmem_base, desc_q, desc_k, idesc, /*accumulate=*/false);
}
__syncwarp();
// Wait for MMA
if (wid == 0 && lane == 0) {
tmem_fence_store();
}
__syncthreads();
// if (tid == 0) {
// umma_ss_f16(tmem_base, desc_q, desc_k, idesc, /*accumulate=*/false);
// }
// __syncwarp();
// if (wid == 0 && lane == 0) {
// tmem_fence_store();
// }
// __syncthreads();
// ================================================================
// Read S from TMEM