auto: pre-test commit

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();
     }