add fence.sc.gpu between PV and epilogue for TMEM visibility

2026-05-28 23:21:53 +00:00
parent 52809b0ec6
commit 55c0604a71
1 changed files with 4 additions and 0 deletions
--- a/dsv4/kernels/attention/fmha_6warp_multirow.cuh
+++ b/dsv4/kernels/attention/fmha_6warp_multirow.cuh
@@ -259,6 +259,10 @@ fmha_6warp_multirow_kernel(FmhaMultiRowParams params) {
        }
    }

+    // Ensure PV output is visible to all warps before epilogue
+    asm volatile("fence.sc.gpu;" ::: "memory");
+    __syncthreads();
+
    // ================================================================
    // EPILOGUE: TMEM → regs → normalize → BF16 → GMEM
    // ================================================================