From a284580422014ab11bf3905a464c57555f6df9fc Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 23 May 2026 05:17:45 +0000 Subject: [PATCH] Add missing TMEM fence after P store in TMEM-P path --- dsv4/kernels/attention/fmha.py | 1 + 1 file changed, 1 insertion(+) diff --git a/dsv4/kernels/attention/fmha.py b/dsv4/kernels/attention/fmha.py index d356ac25..6f9b3c98 100644 --- a/dsv4/kernels/attention/fmha.py +++ b/dsv4/kernels/attention/fmha.py @@ -430,6 +430,7 @@ class FmhaKernel: rP_bf16_frg[None, j].store(s_vec.to(self.q_dtype)) cute.copy(tiled_tmem_store, rP_words, tTMEM_STOREtP) + cute.arch.fence_view_async_tmem_store() else: # SMEM-P: compute softmax, write P to SMEM (TODO) for j in range(frg_cnt):