From f244c4fdd232a48bd68f73a0a6c08477d391ba8e Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 12:56:39 +0000 Subject: [PATCH] test: single-thread MMA (tid==0) for Layout D --- tests/unit/test_umma_qk_hd64.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/unit/test_umma_qk_hd64.cu b/tests/unit/test_umma_qk_hd64.cu index 8b3663e0..54fcddc2 100644 --- a/tests/unit/test_umma_qk_hd64.cu +++ b/tests/unit/test_umma_qk_hd64.cu @@ -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");