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");