test: single-thread MMA (tid==0) for Layout D

This commit is contained in:
2026-05-28 12:56:39 +00:00
parent ba2e390e1e
commit f244c4fdd2

View File

@@ -107,8 +107,8 @@ test_umma_hd64(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k,
uint64_t dq = make_umma_desc_kmajor_none(q_addr, BLOCK_MN);
uint64_t dk = make_umma_desc_kmajor_none(k_addr, BLOCK_MN);
// 4 warp leaders call MMA (same as working HD=16 test)
if (lane == 0) {
// Single thread calls MMA (gau-nernst's elect_one pattern)
if (tid == 0) {
umma_ss_f16(tb, dq, dk, idesc, kt > 0);
}
asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory");