diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs
index 9c62c897c..87e7ac8e3 100644
--- a/crates/ml/src/trainers/dqn/fused_training.rs
+++ b/crates/ml/src/trainers/dqn/fused_training.rs
@@ -1663,12 +1663,18 @@ impl FusedTrainingCtx {
                     )
                     .map_err(|e| anyhow::anyhow!("TLOB Adam: {e}"))?;
             }
+            eprintln!("STEP0_PHASE_TLOB_BWD_ADAM_DONE");
             self.mamba2_backward(self.batch_size)?;
+            eprintln!("STEP0_PHASE_MAMBA2_BWD_DONE");
             self.step_mamba2_adam()?;
+            eprintln!("STEP0_PHASE_MAMBA2_ADAM_DONE");
             self.launch_ofi_embed_backward(self.batch_size)?;
+            eprintln!("STEP0_PHASE_OFI_EMBED_BWD_DONE");
             self.step_ofi_embed_adam()?;
+            eprintln!("STEP0_PHASE_OFI_EMBED_ADAM_DONE");
             self.trainer.apply_pruning_mask()
                 .map_err(|e| anyhow::anyhow!("pruning mask: {e}"))?;
+            eprintln!("STEP0_PHASE_PRUNING_DONE");
             // Adaptive per-branch gradient-norm balancer — caps each
             // branch's weight-gradient L2 norm at
             // `num_branches × median_branch_norm`.  Must run AFTER all
@@ -1680,10 +1686,13 @@ impl FusedTrainingCtx {
             // tuned knobs).
             self.trainer.launch_branch_grad_balance()
                 .map_err(|e| anyhow::anyhow!("branch_grad_balance: {e}"))?;
+            eprintln!("STEP0_PHASE_BRANCH_GRAD_BALANCE_DONE");
             self.trainer.compute_grad_norm_for_adam()
                 .map_err(|e| anyhow::anyhow!("grad_norm: {e}"))?;
+            eprintln!("STEP0_PHASE_GRAD_NORM_DONE");
             self.trainer.submit_adam_ops(&self.online_dueling, &self.online_branching)
                 .map_err(|e| anyhow::anyhow!("adam ops: {e}"))?;
+            eprintln!("STEP0_PHASE_ADAM_OPS_DONE");
             // Plan C Phase 2 follow-up F (2026-04-29): per-step bin_means_reduce
             // producers feed ISV slots [13..16] (magnitude) and [17..21]
             // (direction) via the captured `update_isv_signals` below. Without
@@ -1701,23 +1710,30 @@ impl FusedTrainingCtx {
             // no extra sync, captures cleanly into the adam_update child.
             self.trainer.launch_q_mag_bin_means_reduce(self.batch_size)
                 .map_err(|e| anyhow::anyhow!("q_mag_bin_means_reduce: {e}"))?;
+            eprintln!("STEP0_PHASE_Q_MAG_BIN_DONE");
             self.trainer.launch_q_dir_bin_means_reduce(self.batch_size)
                 .map_err(|e| anyhow::anyhow!("q_dir_bin_means_reduce: {e}"))?;
+            eprintln!("STEP0_PHASE_Q_DIR_BIN_DONE");
             self.trainer.update_isv_signals()
                 .map_err(|e| anyhow::anyhow!("ISV update: {e}"))?;
+            eprintln!("STEP0_PHASE_ISV_UPDATE_DONE");
 
             // Phase 7: Maintenance ops (causal + vaccine)
             self.trainer.submit_maintenance_ops()
                 .map_err(|e| anyhow::anyhow!("maintenance ops: {e}"))?;
+            eprintln!("STEP0_PHASE_MAINTENANCE_DONE");
 
             // Phase 8: IQL modulate + PER priority update (ungraphed)
             self.submit_iql_modulate_ops(agent, &gpu_batch)?;
+            eprintln!("STEP0_PHASE_IQL_MODULATE_DONE");
             self.submit_per_priority_ops(agent)?;
+            eprintln!("STEP0_PHASE_PER_PRIORITY_DONE");
 
             // Capture all children + compose parent after ungraphed step succeeds.
             if let Err(e) = self.capture_training_graph(agent, &gpu_batch) {
                 tracing::warn!("parent graph capture failed (continuing ungraphed): {e}");
             }
+            eprintln!("STEP0_PHASE_CAPTURE_DONE");
         }
 
         // ── Outside graph: kernel-free host ops only ─────────────────
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 475c04ec3..6f5d41349 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,8 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+SP15 Wave 5 SEGV diagnostic — extend checkpoints into Phase 6/7/8 (2026-05-07): smoke `train-xq9hg` (commit `e9d9afbd6`) got past all 13 original checkpoints (`BEGIN` → `NAN_CHECKS_DONE`) cleanly, confirming the SEGV is downstream in Phase 6 (TLOB bwd + Adam, Mamba2 bwd + Adam, OFI embed bwd + Adam, pruning_mask, branch_grad_balance, grad_norm, submit_adam_ops, q_mag/dir_bin_means_reduce, update_isv_signals), Phase 7 (submit_maintenance_ops), or Phase 8 (submit_iql_modulate_ops, submit_per_priority_ops, capture_training_graph). 14 more `STEP0_PHASE_*` checkpoints added across these phases — the next L40S smoke will localize the SEGV to a single operation. Same diagnostic-only contract; will be reverted with the root-cause fix.
+
 SP15 Wave 5 SEGV diagnostic (2026-05-07): L40S smoke `train-jfbzr` on commit `23e9a1f78` segfaulted at exit 139 in the first `FusedTrainer::run_full_step` invocation after `"GPU training guard initialized (epoch loop)"` and before any HEALTH_DIAG output. Static analysis at commit `23e9a1f78` found no defects in the suspect SP15 changes (`2e37af29d` pointer wiring, `23e9a1f78` denoise stride, Wave 4.1b bn_concat_dim, Phase 1.3.b-followup per-env DD); local repro impossible because `test_data/futures-baseline/` lacks the `.fxcache` the production path requires. This commit adds `eprintln!` checkpoints at each phase boundary of `run_full_step`'s ungraphed step-0 path (between PER sample / PopArt / counters / spectral norm / TLOB forward / forward_main / DDQN / aux_ops / post_aux / NaN checks). stderr is line-flushed (vs the tracing JSON stdout formatter which buffers), so the last printed checkpoint pinpoints the SEGV site. Diagnostic-only — no behavior change. Will be reverted after the next L40S smoke localizes the SEGV and the root-cause fix lands.
 
 CUDA OOB fix — `compute_expected_q` writer stride 13 vs `denoise_target_q_buf` size 12 (2026-05-07): pre-existing latent bug surfaced by `compute-sanitizer --tool memcheck` after the SP15 NULL-pointer fix at `2e37af29d` unblocked the cascade and made the next downstream OOB visible. **Sanitizer evidence (pre-fix)**: 4 `Invalid __global__ write of size 4 bytes at compute_expected_q+0x2480` errors per training step at threads (60..63, 0, 0) of block (0, 0, 0), with "Access at <addr> is out of bounds" landing 45 / 97 / 149 / 201 bytes after `denoise_target_q_buf`'s allocation — exactly slots 12 of samples b ∈ {b_oob_0, b_oob_1, b_oob_2, b_oob_3} where the kernel wrote `q_values[i*total_actions + 12]` past the buffer end. Host backtrace pointed at `GpuDqnTrainer::compute_denoise_target_q` (which launches `compute_expected_q` with `q_out_ptr = denoise_target_q_buf.raw_ptr()`). **Root cause** — semantic mismatch between writer stride and buffer size: `compute_expected_q` strides per sample by `total_actions = b0_size + b1_size + b2_size + b3_size = 4 + 3 + 3 + 3 = 13` (the 4-direction factored layout, `q_values[i*total_actions + action_idx]` with `action_idx ∈ [0, total_actions)`); but `denoise_target_q_buf` was allocated at `b * 12` (the legacy 12-slot exposure-layout sizing 3+3+3+3, dating from before the 4-direction factored migration). The downstream `q_denoise_backward` and `denoise_loss_grad` kernels read `Q_target[b * D + i]` with hardcoded `D = 12` because the diffusion denoiser MLP itself only has 12 output slots (W2[12, 24] + b2[12]) — its 12 are the `q_coord_buf`'s post-cross-branch-attention narrowed output, NOT a clean subset of the 13 raw Q-actions. The exact same fix pattern was already applied to the sibling `q_var_buf_trainer` allocation in the prior SP4 audit (see comment block at `gpu_dqn_trainer.rs:21620-21630` referencing "compute-sanitizer caught the b*12 sizing as 4 OOB writes at threads 60..63"); the `denoise_target_q_buf` allocation 8 lines below was missed during that fix because it was guarded by the SP15 NULL-pointer ILLEGAL_ADDRESS that fault-stopped the cascade before this OOB could fire — `2e37af29d` removed the upstream ILLEGAL_ADDRESS, surfacing the latent OOB. **Fix architecture** (Option A — pad buffer to total_actions, pass stride to consumer; same as `q_var_buf_trainer`): (1) **Allocation widened** in `gpu_dqn_trainer.rs:~21673` — `denoise_target_q_buf` now sized `b * total_actions` (= b * 13) instead of `b * 12`. The sibling `denoise_q_input_buf` STAYS at `b * 12` because it's a snapshot of `q_coord_buf` (= 12 wide, post-cross-branch-attention) via `snapshot_pre_denoise_q`'s `n_bytes = batch_size * 12 * sizeof::<f32>()` DtoD copy; it's NEVER written by `compute_expected_q`. (2) **`compute_expected_q` kernel UNCHANGED** — its 13-stride write pattern is correct (it produces all 13 Q-actions per sample); the bug was downstream sizing, not the writer. (3) **`q_denoise_backward` kernel** (`experience_kernels.cu:6348`) — added 3 new stride parameters: `refined_stride`, `target_stride`, `input_stride`. The kernel still operates on D=12 per-sample (the denoiser MLP's fixed output width) but addresses each input buffer at its own per-sample stride: `Q_refined[b*refined_stride+i]` (q_coord_buf, stride=12), `Q_target[b*target_stride+i]` (denoise_target_q_buf, stride=total_actions=13), `Q_input[b*input_stride+i]` (denoise_q_input_buf, stride=12). The hard-coded `b*D` indexing was the smoking gun — for a stride-13 buffer, sample b=1's slot 0 lives at byte offset 13*4=52, but `b*D=12` reads byte offset 12*4=48 (= sample 0's slot 12 = OOB tail of buf[b*12] sizing) — corruption masked when buffer was sized for stride 12 (the same hardcoded D coincidentally matched). With buffer sized 13, the same hardcoded `b*D` would now read sample (b-1)'s slot 12 — semantically incorrect — hence stride parameters thread the truth from the launcher. (4) **`denoise_loss_grad` kernel** (`experience_kernels.cu:6254`, used by `launch_q_denoise_backward_cublas`) — same pattern: added `refined_stride` and `target_stride` parameters; reads `Q_refined[b*refined_stride+d]` (q_coord_buf, stride=12) and `Q_target[b*target_stride+d]` (denoise_target_q_buf, stride=13). (5) **Rust launchers updated atomically** — `launch_q_denoise_backward` (`gpu_dqn_trainer.rs:~26253`) passes `refined_stride=12, target_stride=total_actions(), input_stride=12`; `launch_q_denoise_backward_cublas` (`gpu_dqn_trainer.rs:~26511`) passes `refined_stride=12, target_stride=total_actions()` to `denoise_loss_grad`. (6) **Flat-scan consumers UNCHANGED** — `target_q_p99_update` (SP4 producer at `target_q_p99_kernel.cu`) and `dqn_clamp_finite_f32_kernel` both consume the buffer flat via `denoise_target_q_buf.len()` for histogram p99 and clamp; widening from `b*12` to `b*13` is monotone (one more value per sample in the histogram, all valid Q-values; clamp covers all 13 instead of 12) — no semantic change, no kernel modification needed. The SP3 threshold-check entry at slot 46 (`target_q_ptr` + `target_q_len`) likewise consumes the full buffer flat; doc updated to reflect new size. **Atomic per `feedback_no_partial_refactor`**: kernel signatures (2 kernels) + buffer allocation + 2 launch sites + struct doc comments + SP3 slot doc + audit doc all in one commit. **Verification**: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean; `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 compute-sanitizer --tool memcheck --error-exitcode=42 ./target/debug/deps/ml-* --ignored test_no_hang_single_epoch --nocapture --test-threads=1` shows **0 `Invalid __global__` errors** at `compute_expected_q` (down from 4 per training step in baseline — verified by re-running sanitizer on the parent commit `2e37af29d` which produced 4 OOB writes at threads 60-63 of block 0); remaining sanitizer messages are `CUDA_ERROR_STREAM_CAPTURE_INVALIDATED` cascades from a sanitizer-infrastructure VRAM exhaustion on the RTX 3050 Ti (4GB) under sanitizer overhead — not memory errors, and confirmed pre-existing because the test runs to the same point in baseline before sanitizer-induced OOM. `SQLX_OFFLINE=true cargo test -p ml --features cuda --lib` HOLDS the parent baseline at **947 pass / 12 fail** (same 12 pre-existing failures: `cuda_pipeline::tests::test_ppo_gpu_data_upload`, `monitors::state_kl_monitor::tests::observe_tracks_fire_rate_on_change`, 2 spectral-norm tests, 4 batched-action-selection tests, 1 ppo reward test, 1 training_profile test). Hard rules: `feedback_no_partial_refactor` (kernel + buffer + launchers + tests + audit doc atomic — every consumer of the writer's stride change migrates simultaneously), `feedback_no_quickfixes` (real fix that addresses the semantic mismatch via stride parameter, NOT a one-line patch like dropping threads 12 from the kernel grid which would silently lose action 12's Q-value), `feedback_no_legacy_aliases` (no padding-to-12-and-discarding-13 wrapper kept; the writer's 13-action contract is the truth, downstream consumers thread the stride to read their D=12 from the right base), `feedback_wire_everything_up` (the parallel `q_var_buf_trainer` fix from the prior SP4 audit is now matched by `denoise_target_q_buf`; the same `var_stride`/`target_stride` pattern is applied uniformly to all per-sample buffers fed by `compute_expected_q`). Files touched: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (q_denoise_backward + denoise_loss_grad kernel signatures), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (alloc widen + 2 launcher updates + 4 doc comment updates), `docs/dqn-wire-up-audit.md` (this entry).