FMHA: zero-init my_p_vals to fix N<128 padding NaN

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.
This commit is contained in:
2026-05-31 23:11:12 +00:00
parent 13be3ad443
commit 4adee1207f

View File

@@ -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;