From b69a538ab19aa708a02b7116f05ea4861cbe433d Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 09:53:10 +0000
Subject: [PATCH] test: add fence+sync between 2 tmem_stores

---
 tests/unit/test_tmem_cols.cu | 5 +++++
 1 file changed, 5 insertions(+)

diff --git a/tests/unit/test_tmem_cols.cu b/tests/unit/test_tmem_cols.cu
index 71dd550b..1b73d4fd 100644
--- a/tests/unit/test_tmem_cols.cu
+++ b/tests/unit/test_tmem_cols.cu
@@ -45,8 +45,13 @@ __global__ void test_tmem_loop(float* out) {
         memcpy(&u2, &v2, 4); memcpy(&u3, &v3, 4);
         
         tmem_store(tb + 0, u0, u1, u2, u3);
+    }
+    asm volatile("tcgen05.wait::st.sync.aligned;" ::: "memory");
+    __syncthreads();
+    if (threadIdx.x < 32) {
         tmem_store(tb + 1, u0, u1, u2, u3);
     }
+    }
     asm volatile("tcgen05.wait::st.sync.aligned;" ::: "memory");
     __syncthreads();