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 // ================================================================