fmha_multirow: add fence.sc.gpu after QK GEMM for TMEM visibility

2026-05-28 23:13:31 +00:00
parent d840fbbf85
commit 43ba672e15
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
@@ -164,6 +164,10 @@ fmha_6warp_multirow_kernel(FmhaMultiRowParams params) {
        __syncthreads();
    }

+    // Ensure TMEM visibility across all warps after QK GEMM
+    asm volatile("fence.sc.gpu;" ::: "memory");
+    __syncthreads();
+
    // ================================================================
    // SOFTMAX PASS 1: compute row_max for each row
    // ================================================================