From 2b007d20088548c48d54e43c6b4938a3f359ae33 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 15:53:39 +0000 Subject: [PATCH] auto: pre-test commit --- tests/unit/test_tmem_zero_pv.cu | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/tests/unit/test_tmem_zero_pv.cu b/tests/unit/test_tmem_zero_pv.cu index bc461c35..69e612a1 100644 --- a/tests/unit/test_tmem_zero_pv.cu +++ b/tests/unit/test_tmem_zero_pv.cu @@ -113,10 +113,16 @@ test_tmem_zero_before_pv(const bf16_t* q, const bf16_t* k, const bf16_t* v, } __syncthreads(); - // ===== ZERO ALL 128 TMEM COLUMNS ===== + // ===== ZERO ALL 128 TMEM COLUMNS using 32x32b.x8 stores ===== if (wid == 0) { - for (int col = 0; col < 128; col++) { - tmem_store(tb + col, 0, 0, 0, 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); + 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) } tmem_fence_store(); }