diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
index b63d22703..94beef9df 100644
--- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
@@ -1700,3 +1700,37 @@ extern "C" __global__ void popart_normalize_robust(
     float safe_iqr = fmaxf(iqr, 1e-4f);
     rewards[i] = (rewards[i] - median) / safe_iqr;
 }
+
+// ============================================================================
+// SP2: Fused multi-buffer NaN check
+//
+// Replaces 8 individual `dqn_nan_check_f32` launches in run_nan_checks_post_backward
+// (slots 24-30 + 32) with a single launch. Each block (blockIdx.x = 0..N_SLOTS)
+// scans its assigned buffer end-to-end, sets the corresponding flag if any
+// non-finite value found. buf_ptrs[i] = nullptr → block i no-ops (deferred slot).
+//
+// Invariants preserved from dqn_nan_check_f32:
+//   - sticky-flag semantics (writes 1 only; never clears)
+//   - block-local reduce (no atomicAdd)
+//   - graph-capture safe (pure kernel launch, no host ops)
+// ============================================================================
+extern "C" __global__ void dqn_nan_check_fused_f32_kernel(
+    const float* const* buf_ptrs,    // [N_SLOTS]
+    const int* buf_lens,             // [N_SLOTS]
+    int base_flag_idx,               // first slot index (24 for backward checks)
+    int* nan_flags_buf
+) {
+    const int slot = blockIdx.x;
+    const float* buf = buf_ptrs[slot];
+    const int n = buf_lens[slot];
+    if (buf == nullptr || n == 0) return;  // deferred slot or unused entry
+
+    int has_nan = 0;
+    for (int i = threadIdx.x; i < n; i += blockDim.x) {
+        if (!isfinite(buf[i])) { has_nan = 1; }
+    }
+    has_nan = __syncthreads_or(has_nan);
+    if (threadIdx.x == 0 && has_nan) {
+        nan_flags_buf[base_flag_idx + slot] = 1;
+    }
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index a46ae33a8..8c116a3f9 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2264,3 +2264,5 @@ SP1 Phase C quality fix-up (2026-04-29): commit-quality follow-up to `97f1d25f5`
 SP1 Phase C fix-up #2 (2026-04-29): tightened ε floor formula at the two ISV-driven max_abs sites in `apply_iqn_trunk_gradient` (line ~6967) and `launch_cublas_backward_to` (line ~18631). Original `(1e6 × isv).max(1e3)` clipped legitimate F1 startup gradients to ±1e3 when fold-boundary ISV reset put `H_S2_RMS_EMA_INDEX` or `Q_DIR_ABS_REF_INDEX` near 0 — verified by smoke `smoke-test-dr2bn` (commit 19b008e1c) which F1-NaN'd at step 240 (vs step 890 in pre-fix smoke `smoke-test-xvzgk`). New formula `1e6 × isv.max(1.0)` guarantees `max_abs ≥ 1e6` in all ISV states (cold-start = 1e6, warm = 1e6 × isv), removing cold-start clamp pathology. F0 no-op intent preserved (F0 inputs ≪ 1e6 in all states). The 1.0 ε floor is on the ISV multiplier (Invariant 1 carve-out per `feedback_isv_for_adaptive_bounds`), not on the bound itself — so the bound is still ISV-driven when ISV is meaningful.
 
 SP1 closure (2026-04-29): F1 cold-start clamp pathology fixed at commit `ab2133463` (ε floor formula `1e6 × isv.max(1.0)`). F1 NaN moved from step 240 → step 3540 (15× later), validating the diagnosis. Remaining structural F1 NaN at step 3540 classified as SP3 scope (Adam weight pathology — clamps cannot prevent NaN/Inf weights from accumulating via Q-target inflation over training). F0 regression ~35 (vs baseline 55.87) classified as SP2 scope (consolidate 11 per-step `check_nan_f32` launches into a single fused reduce kernel). Permanent diagnostic infrastructure (nan_flags_buf 24→48 with slots 24-35 backward coverage) verified working across 3 L40S smokes — sentinel correctly fires on cuBLAS GEMM accumulator overflow at slots 26 (`iqn_trunk_m`) + 32 (`bn_d_concat_buf`); IQN-internal buffers (slots 27, 28, 35) stayed clean throughout, ruling out IQN backward chain as the seed. Audit doc `docs/dqn-backward-nan-audit.md` becomes durable artifact for SP2/SP3. Memory entry `project_sp1_f1_nan_root_cause_resolved.md` documents handoff including ε-floor fix-up rationale (ε on multiplier, not bound) and NaN diagnostic ordering pearl (inline check before each clamp to preserve sentinel).
+
+SP2 Phase A1 — fused NaN-check kernel (2026-04-29): foundational kernel-only commit appending `dqn_nan_check_fused_f32_kernel` to `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`. Single-launch replacement target for the 8 per-step `dqn_nan_check_f32` calls in `run_nan_checks_post_backward` (slots 24-30 + 32). Each block (`blockIdx.x = 0..N_SLOTS`) scans its assigned buffer end-to-end via grid-strided loop, block-local `__syncthreads_or` reduce (no atomicAdd per `feedback_no_atomicadd`), sticky-flag write `nan_flags_buf[base_flag_idx + slot] = 1` only when `has_nan && threadIdx.x == 0`. Nullptr bounds check `if (buf == nullptr || n == 0) return;` makes deferred slots (e.g. slot 31 ensemble) and unused entries no-op cleanly — same call site can pass null for cross-struct slots without spurious flag writes. Invariants preserved from `dqn_nan_check_f32`: sticky-flag (writes 1, never clears), graph-capture safe (pure kernel launch), no DtoD/HtoD/HtoH copies. `extern "C"` linkage matches existing `dqn_*` symbol convention for cudarc `module.load_function` lookup. Unused yet — Rust wrapper, metadata buffer (per-slot ptr/len arrays), and call-site replacement land in subsequent A2/A3/A4 commits. Drives F0 regression remediation: 8 per-step kernel launches collapse to 1 fused launch (~7× fewer graph nodes for backward NaN coverage) while preserving identical diagnostic semantics. Pattern reference for future fused diagnostic launches if SP3 expands the slot range.