From babff1f40239ab47c0de74db0c639765fa6ef519 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 15:54:05 +0000 Subject: [PATCH] auto: pre-test commit --- tests/unit/test_tmem_zero_pv.cu | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) diff --git a/tests/unit/test_tmem_zero_pv.cu b/tests/unit/test_tmem_zero_pv.cu index 69e612a1..992116f7 100644 --- a/tests/unit/test_tmem_zero_pv.cu +++ b/tests/unit/test_tmem_zero_pv.cu @@ -114,15 +114,13 @@ test_tmem_zero_before_pv(const bf16_t* q, const bf16_t* k, const bf16_t* v, __syncthreads(); // ===== ZERO ALL 128 TMEM COLUMNS using 32x32b.x8 stores ===== + // All 32 lanes must call. Each lane writes 8 zero FP32 values. if (wid == 0) { - for (int n = 0; n < 16; n++) { // 16 groups of 8 columns - // 32x32b.x8 store: all 32 lanes must participate - // Each lane writes 8 zero FP32 values (one per column) - float z = 0.0f; - uint32_t uz = f32_to_u32(z); + float z0 = 0.0f, z1 = 0.0f, z2 = 0.0f, z3 = 0.0f; + float z4 = 0.0f, z5 = 0.0f, z6 = 0.0f, z7 = 0.0f; + for (int n = 0; n < 16; n++) { asm volatile("tcgen05.st.sync.aligned.32x32b.x8.b32 [%0], {%1, %2, %3, %4, %5, %6, %7, %8};" - :: "r"(tb + n*8), "f"(z), "f"(z), "f"(z), "f"(z), "f"(z), "f"(z), "f"(z)); - // wait not needed between stores (same warp, same format) + :: "r"(tb + n*8), "f"(z0), "f"(z1), "f"(z2), "f"(z3), "f"(z4), "f"(z5), "f"(z6), "f"(z7)); } tmem_fence_store(); }