From 4adee1207fde53c517264ea9ce1e14a715c96e37 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sun, 31 May 2026 23:11:12 +0000 Subject: [PATCH] FMHA: zero-init my_p_vals to fix N<128 padding NaN MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit When N<128, padded KV positions have my_p_vals[col] uninitialized for col >= kv_len. The PV GEMM then computes garbage_P × zero_V, which can produce NaN on tensor cores (0 × NaN = NaN). Fix: zero-initialize my_p_vals so padded positions contribute 0. --- dsv4/kernels/attention/fmha_6warp_tma_multirow_multitile.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dsv4/kernels/attention/fmha_6warp_tma_multirow_multitile.cuh b/dsv4/kernels/attention/fmha_6warp_tma_multirow_multitile.cuh index 917ef877..86e9f53b 100644 --- a/dsv4/kernels/attention/fmha_6warp_tma_multirow_multitile.cuh +++ b/dsv4/kernels/attention/fmha_6warp_tma_multirow_multitile.cuh @@ -211,7 +211,7 @@ fmha_6warp_tma_multirow_multitile_kernel(FmhaTmaMultiRowMultiTileParams params) if (my_row_active) sTileRowMax[my_row] = my_row_max; __syncthreads(); - float my_p_vals[SK_TILE]; + float my_p_vals[SK_TILE] = {}; // Zero-init: padded positions contribute 0 to PV float my_row_sum = 0.0f; if (my_warp_active) { float rm = my_row_active ? sTileRowMax[my_row] : 0.0f;