From 43ba672e15ed43a567c71181be8a66b57db41ad4 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 23:13:31 +0000 Subject: [PATCH] fmha_multirow: add fence.sc.gpu after QK GEMM for TMEM visibility --- dsv4/kernels/attention/fmha_6warp_multirow.cuh | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/dsv4/kernels/attention/fmha_6warp_multirow.cuh b/dsv4/kernels/attention/fmha_6warp_multirow.cuh index 4c95504c..b434987e 100644 --- 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 // ================================================================