auto: pre-test commit

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