diff --git a/dsv4/kernels/attention/fmha_qk_verify.cuh b/dsv4/kernels/attention/fmha_qk_verify.cuh
index 52217a09..b260c392 100644
--- a/dsv4/kernels/attention/fmha_qk_verify.cuh
+++ b/dsv4/kernels/attention/fmha_qk_verify.cuh
@@ -68,7 +68,14 @@ fmha_qk_verify(
     uint64_t desc_k = make_umma_desc_k_none(sK_smem, HD);
 
     // QK GEMM: S = Q @ K^T (SS, both SMEM → TMEM)
-    if (wid == 0 && lane == 0) {
+    // MMA: called by ONE lane (elect_one_sync pattern)
+    if (tid == 0) printf("[qk] tmem_base=%u sQ_smem=0x%x sK_smem=0x%x\n", tmem_base, sQ_smem, sK_smem);
+    __syncthreads();
+
+    // Try: ALL 32 lanes of warp 0 call MMA (not just lane 0)
+    // The CUTLASS code uses elect_one_sync which selects lane 0,
+    // but maybe the PTX requires all lanes to participate?
+    if (wid == 0) {
         umma_ss_f16(tmem_base, desc_q, desc_k, false);
     }
     __syncwarp();