metal: bring v2 f16 attention to MLX-faithful exp/sentinel pattern

Per `feedback_ferrite_metal_mlx_only.md` rule 7, both v2 attention
variants must match MLX `sdpa_vector` bit-exact:
  - `exp` → `metal::fast::exp`
  - `-INFINITY` → `Limits<U>::finite_min` (= `-FLT_MAX` for f32 accum)

The bf16 sibling (`attention_via_cache_v2_bf16_specialized`) was
already updated; the f16 variant lagged with bare `exp` + `-INFINITY`
at three sites (online softmax + per-simdgroup combine).

Drift on the f16 path is small in steady-state usage but compounds
across long sequences and degrades agreement with the bf16 sibling
at mixed-dtype model swaps. Patching the f16 kernel restores the
contract spelled out in the kernel header comment ("port of MLX
sdpa_vector").

Goldens unchanged: 68/68 ferrite-forward metal lib tests pass,
including the f16 numerical-correctness check
`attention_via_cache_matches_cpu_golden`.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

Author: starpit

vllm-rs/crates/ferrite-metal-kernels/shaders/attention.metal

@@ -148,12 +148,16 @@ kernel void attention_via_cache_v2_f16_specialized(
     device       half* o_row = output + (seq_idx * num_q + q_head_idx) * head_dim;
     device const uint* row_block_table = block_table + seq_idx * max_blocks;
 
+    // Pre-multiply Q by scale (MLX `sdpa_vector`: `q[i] = scale * queries[i]`).
     for (uint i = 0; i < qk_per_thread; ++i) {
         q_reg[i] = U(scale) * U(q_row[simd_lid * qk_per_thread + i]);
         o_reg[i] = 0;
     }
 
-    U max_score = -INFINITY;
+    // Initialize per-thread max with finite minimum (MLX uses
+    // `Limits<U>::finite_min`; -FLT_MAX is the f32 equivalent).
+    // fast::exp doesn't handle -INFINITY safely so we avoid it.
+    U max_score = -FLT_MAX;
     U sum_exp_score = 0;
 
     // For each key, simdgroup `simd_gid` handles tokens at indices
@@ -185,10 +189,11 @@ kernel void attention_via_cache_v2_f16_specialized(
         }
         score = simd_sum(score);
 
-        // Online softmax update.
+        // Online softmax update. Match MLX `sdpa_vector`: fast::exp
+        // for both factor + exp_score.
         U new_max = max(max_score, score);
-        U factor = exp(max_score - new_max);
-        U exp_score = exp(score - new_max);
+        U factor = metal::fast::exp(max_score - new_max);
+        U exp_score = metal::fast::exp(score - new_max);
 
         max_score = new_max;
         sum_exp_score = sum_exp_score * factor + exp_score;
@@ -214,7 +219,7 @@ kernel void attention_via_cache_v2_f16_specialized(
     // and (factor-rescaled) global sum_exp.
     U other_max = tg_max[simd_lid];
     U global_max = simd_max(other_max);
-    U factor = exp(other_max - global_max);
+    U factor = metal::fast::exp(other_max - global_max);
     U global_sum = simd_sum(tg_sum[simd_lid] * factor);
 
     // Combine output partials. Each simdgroup wrote o_reg[j] for