From d313917e702f18e761e2f893238278c803fe32fe Mon Sep 17 00:00:00 2001
From: "cheheng.ch" <haoc80996@gmail.com>
Date: Fri, 31 Jul 2026 06:25:30 +0000
Subject: [PATCH] k1: fence generic smem writes before the TMA workspace stores

bar.sync orders threads inside the generic proxy; it does not order
those writes against the TMA engine's reads. The C++ kernel this was
ported from has the same fence (FlashKDA fwd_kernel1.cuh:515).
---
 cula/ops/kda/sm90/k1.py | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/cula/ops/kda/sm90/k1.py b/cula/ops/kda/sm90/k1.py
index de44036..8b2fdd1 100644
--- a/cula/ops/kda/sm90/k1.py
+++ b/cula/ops/kda/sm90/k1.py
@@ -529,6 +529,8 @@ def k1_kernel(
             smem_thr_store_C.retile(tCrInvC_bf16),
             smem_thr_store_C.partition_D(sINV_bf16),
         )
+    # Generic-proxy smem writes -> visible to the async proxy.
+    cute.arch.fence_view_async_shared()
     cute.arch.barrier()
 
     # TMA bulk store all 5 workspace tensors (one elect_one, one thread).