From e9d9afbd61f544cf153ea76fad02256f511311fa Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 7 May 2026 12:21:12 +0200 Subject: [PATCH] diag(sp15-wave5): stderr checkpoints in run_full_step step-0 ungraphed path MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit L40S smoke train-jfbzr (commit 23e9a1f78) segfaults at exit 139 in the first run_full_step invocation, after "GPU training guard initialized (epoch loop)" and before any HEALTH_DIAG output. Static analysis debugger couldn't reproduce locally (test_data/futures-baseline/ lacks .fxcache). This commit adds eprintln! checkpoints at each phase boundary in the ungraphed step-0 path of FusedTrainer::run_full_step (PER sample, PopArt, counters, spectral norm, TLOB forward, forward_main, DDQN, aux_ops, post_aux, NaN checks). stderr is line-flushed (vs stdout buffered through tracing JSON formatter), so the last printed checkpoint identifies the SEGV site. Each fires once per fold's first step (graph capture absorbs subsequent steps), so log overhead is minimal. Diagnostic-only commit — no behavior change. Will be reverted after the next L40S smoke localizes the SEGV and the root-cause fix lands. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/src/trainers/dqn/fused_training.rs | 17 +++++++++++++++++ docs/dqn-wire-up-audit.md | 2 ++ 2 files changed, 19 insertions(+) diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 7ccad66b4..9c62c897c 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -1489,6 +1489,11 @@ impl FusedTrainingCtx { tracing::debug!("parent graph launch DONE (async — GPU working)"); } else { // Ungraphed step 0: run everything, then capture + compose. + // SP15 Wave 5 SEGV diagnostics: stderr checkpoints survive + // host-side crashes that don't get emitted to the tracing + // JSON formatter (stdout is buffered; stderr is line-flushed). + // Each checkpoint prints once per fold's first step. + eprintln!("STEP0_PHASE_BEGIN"); // Ensure params are initialized for the first step. if !self.trainer.params_initialized { self.trainer.flatten_online_weights( @@ -1496,6 +1501,7 @@ impl FusedTrainingCtx { ).map_err(|e| anyhow::anyhow!("flatten weights: {e}"))?; self.trainer.params_initialized = true; } + eprintln!("STEP0_PHASE_FLATTEN_ONLINE_DONE"); if !self.trainer.target_params_initialized { let n_bytes = self.trainer.total_params() * std::mem::size_of::(); let src = self.trainer.params_buf_ptr(); @@ -1506,10 +1512,12 @@ impl FusedTrainingCtx { } self.trainer.target_params_initialized = true; } + eprintln!("STEP0_PHASE_TARGET_INIT_DONE"); // Phase -1: PER sampling (writes directly to trainer buffers via gather_padded) let ptrs = agent.primary_dqn_mut().memory.gpu.sample_proportional(self.batch_size) .map_err(|e| anyhow::anyhow!("PER sample step 0: {e}"))?; + eprintln!("STEP0_PHASE_PER_SAMPLE_DONE"); let gpu_batch = crate::dqn::replay_buffer_type::GpuBatch { states_ptr: ptrs.states_ptr, actions_ptr: ptrs.actions_ptr, @@ -1535,9 +1543,11 @@ impl FusedTrainingCtx { self.trainer.normalize_rewards_popart_inplace(self.batch_size) .map_err(|e| anyhow::anyhow!("PopArt normalize: {e}"))?; } + eprintln!("STEP0_PHASE_POPART_DONE"); // Phase 0b: Counter increments + stochastic depth mask (ungraphed) self.submit_counters_ops()?; + eprintln!("STEP0_PHASE_COUNTERS_DONE"); // HER donor computation (outside graph -- indices vary per step) if let Some(ref mut her) = self.gpu_her { @@ -1563,6 +1573,7 @@ impl FusedTrainingCtx { // Phase 1: Spectral norm self.trainer.apply_spectral_norm(&self.online_dueling, &self.online_branching) .map_err(|e| anyhow::anyhow!("spectral norm: {e}"))?; + eprintln!("STEP0_PHASE_SPECTRAL_NORM_DONE"); // D.8 TLOB forward: OFI[32]→TLOB[16] written into states_buf before trunk forward. { @@ -1571,21 +1582,26 @@ impl FusedTrainingCtx { .forward(states, self.batch_size) .map_err(|e| anyhow::anyhow!("TLOB forward: {e}"))?; } + eprintln!("STEP0_PHASE_TLOB_FORWARD_DONE"); // Phase 2: Forward (cuBLAS trunk + ISV + loss + backward) self.trainer.submit_forward_ops_main() .map_err(|e| anyhow::anyhow!("forward ops: {e}"))?; + eprintln!("STEP0_PHASE_FORWARD_MAIN_DONE"); // Phase 3: DDQN Pass 3 self.trainer.submit_forward_ops_ddqn() .map_err(|e| anyhow::anyhow!("ddqn ops: {e}"))?; + eprintln!("STEP0_PHASE_FORWARD_DDQN_DONE"); // Phase 4: Aux ops (HER + EMA + IQL + IQN + attention + CQL) self.submit_aux_ops(agent, &gpu_batch)?; + eprintln!("STEP0_PHASE_AUX_OPS_DONE"); // Phase 5b: Post-aux ops (selectivity + denoise + Q-stats + risk_sgd) self.trainer.submit_post_aux_ops(self.batch_size) .map_err(|e| anyhow::anyhow!("post_aux ops: {e}"))?; + eprintln!("STEP0_PHASE_POST_AUX_DONE"); // SP1 Phase B (Task 4) + SP2 (Task A4): single fused NaN-check // launch covers all 12 backward-path slots (24-35) in 12 blocks. @@ -1603,6 +1619,7 @@ impl FusedTrainingCtx { // are skipped by the kernel's null-pointer guard. self.trainer.run_nan_checks_post_backward() .map_err(|e| anyhow::anyhow!("post_backward nan checks: {e}"))?; + eprintln!("STEP0_PHASE_NAN_CHECKS_DONE"); // Phase 6: TLOB bwd + Mamba2 bwd + OFI embed bwd + pruning + grad_norm + Adam + ISV update // TLOB backward: reads TLOB-slice gradient from bn_d_concat_buf, updates W_Q/K/V/O. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 4c16fc461..475c04ec3 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 (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 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::()` 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). SP15 Phase 1.3.b-followup-B — per-(env, t) dd_trajectory + PER sampler migration (2026-05-07): closes the deferred Phase 1.3.b-followup point (10) per `feedback_no_partial_refactor`. Phase 1.3.b-followup (`5b394f103`) landed the `dd_state_kernel.cu` per-env redesign + 5 of its 6 consumers (final_reward + plasticity + HEALTH_DIAG + 6 oracle tests + new behavioral test) but explicitly deferred the `dd_trajectory_decreasing_kernel` + `per_insert_pa` migration as a separate split because the contract shape is fundamentally different: dd_trajectory's output needs a per-(env, t)-capable buffer (read at insert time by `per_insert_pa` via `env_id = (j % (n_envs * lookback)) / lookback` lookup), and the previous Wave 4.3 implementation was a documented statistical-bias bug — `ISV[DD_TRAJECTORY_DECREASING_INDEX=439]` was read ONCE per insert call and applied uniformly to ALL `n_envs * lookback * 2` transitions in the batch (whichever env's trajectory write landed last determined the boost for ALL inserted transitions, regardless of each transition's actual env). (1) **`dd_trajectory_kernel.cu` reshape** — grid `[n_envs, 1, 1] × [1, 1, 1]` (mirrors `dd_state_kernel`'s per-env shape; one thread per env, no reductions or atomics inside the kernel). Each thread reads its env's DD_PCT from the shared `dd_state_per_env` tile (offset 5 within the 6-stat sub-array — written by `dd_state_kernel` in the launch immediately before this one on the same stream), reads its persistent prev_dd_pct from a per-env scratch buffer, and writes its trajectory boolean to a per-env output tile. ISV[DD_TRAJECTORY_FLOOR_INDEX=441] remains a global ISV scalar (structural cold-start anchor; the same threshold applies across all envs by design — the gate's job is to reject "0.001 → 0.0005" PnL flutter which is a fixed magnitude concept, not a per-env quantile). (2) **Two new collector-owned mapped-pinned tiles**: `sp15_dd_trajectory_per_env: MappedF32Buffer([alloc_episodes])` (per-env trajectory output — read by `per_insert_pa` per-transition) and `sp15_dd_trajectory_prev_dd_per_env: MappedF32Buffer([alloc_episodes])` (per-env persistent prev_dd_pct scratch; replaces the deleted trainer-owned `[1]` `sp15_dd_trajectory_prev_dd` from Wave 4.3). Both allocated in `GpuExperienceCollector::new` next to `sp15_dd_state_per_env`. The trainer-owned `[1]` scratch buffer + its ctor allocation + struct field + destructure in `Self { ... }` were DELETED — the `set_sp15_dd_trajectory_prev_dd_ptr` collector setter + the trainer→collector wiring block in `training_loop.rs` were also DELETED (production wiring is now unconditional via the collector-owned tiles). Mirrors the `sp15_dd_state_per_env` collector-owned ownership pattern. (3) **`dd_state_reduce_kernel.cu` extension** — added `dd_trajectory_per_env` parameter (nullable; null-skips the slot 439 mean-aggregate). When provided, the kernel mean-aggregates the per-env trajectory tile into ISV[DD_TRAJECTORY_DECREASING_INDEX=439] for HEALTH_DIAG diagnostic preservation in the same single-block tree-reduce that produces ISV[401..407) (mean) + ISV[443] (max-persistence). The shared-memory tile grew from `[BLOCK_SIZE × 7]` to `[BLOCK_SIZE × 8]` (one new lane for trajectory sum); the tree-reduce loop got one extra `+=`. The reduce launcher signature gained a 4th parameter (`dd_trajectory_per_env`) — single launch site updated in `gpu_experience_collector.rs::launch_timestep_loop` step 5b + 2 oracle test sites pass `0` (null) to keep their dd_state-only assertions bit-stable. (4) **`per_insert_pa` migration** (`crates/ml-dqn/src/per_kernels.cu`): kernel signature gained 3 new parameters — `dd_trajectory_per_env` (nullable per-env tile pointer), `n_envs` (i32), `lookback` (i32). Old `isv[DD_TRAJECTORY_DECREASING_INDEX=439]` read DELETED; per-transition lookup added: `env_id = (i % (n_envs * lookback)) / lookback` → `decreasing = dd_trajectory_per_env[env_id]`. ISV[440]=RECOVERY_OVERSAMPLE_WEIGHT remains a global ISV scalar read (the oversample weight is structural, not per-env). Boost gate degrades together: `boost_factor = 1.0` when ANY of `(isv == nullptr, dd_trajectory_per_env == nullptr, n_envs <= 0, lookback <= 0)`. The env_id mapping is consistent across BOTH halves of the cf_mult=2 batch layout (`[on_policy: N*L | counterfactual: N*L]` per `experience_kernels.cu::out_off = i*L+t` and `cf_off = N*L + i*L+t`) because the modulo `i % (N*L)` collapses both halves into the same `slot_it = i*L+t` semantics, and the divide `slot_it / L` recovers the env. The on-policy and CF transitions for the same (env, t) pair share the same env's DD trajectory because the env's portfolio state machine is unitary across the cf-sibling pair — the same shared semantic that `compute_sp15_final_reward_kernel` already uses for its per-(i,t) DD lookup (Phase 1.3.b-followup point 3). (5) **`GpuReplayBuffer` setters** — added `set_sp15_dd_trajectory_per_env_ptr(dev_ptr: u64)` + `set_sp15_per_env_dims(n_envs: i32, lookback: i32)` (mirrors the existing `set_isv_signals_ptr` pattern). Three new fields: `sp15_dd_trajectory_per_env_dev_ptr: u64`, `sp15_n_envs: i32`, `sp15_lookback: i32` (all 0 / null until trainer wires them; kernel falls back to boost_factor=1.0 when unwired). Public accessors `sp15_dd_trajectory_per_env_dev_ptr()` + `sp15_per_env_dims()` mirror `isv_signals_dev_ptr()` for oracle-test wiring verification. (6) **Trainer plumbing in `training_loop.rs`** — alongside the existing `set_isv_signals_ptr` call to the GPU PER replay buffer, added a new wiring block that reads `collector.sp15_dd_trajectory_per_env.dev_ptr` + `collector.alloc_episodes()` + `collector.alloc_timesteps()` and calls `set_sp15_dd_trajectory_per_env_ptr` + `set_sp15_per_env_dims` on the buffer. The deleted `set_sp15_dd_trajectory_prev_dd_ptr` block was replaced with a comment explaining the migration (the buffer it pointed at no longer exists). (7) **Per-step launch order in `gpu_experience_collector.rs::launch_timestep_loop` step 5b** — moved `launch_sp15_dd_trajectory_decreasing` from a separate "step 5b" gate (after `dd_state_reduce`) INTO the dd_state block, between `launch_sp15_dd_state` (per-env tile producer) and `launch_sp15_dd_state_reduce` (now also reads the per-env trajectory tile) so the reduce kernel can mean-aggregate the per-env trajectory output into ISV[439] in the same launch. New launch sequence (all on the same stream — CUDA serializes producer→consumer without explicit event sync): `dd_state_kernel` (per-env tile producer) → `dd_trajectory_decreasing_kernel` (per-env trajectory producer; gated on `isv_signals_dev_ptr != 0` because it reads ISV[441]) → `dd_state_reduce_kernel` (ISV scalar consumer for both per-env tiles; gated on `isv_signals_dev_ptr != 0`) → `alpha_split_producer_kernel` (existing) → `compute_sp15_final_reward_kernel` (per-(i,t) consumer). Outside the exp-fwd graph capture region per `pearl_no_host_branches_in_captured_graph`. (8) **State-reset registry** — replaced the single `sp15_dd_trajectory_prev_dd` entry (trainer-owned `[1]` scratch) with TWO collector-owned `[alloc_episodes]` per-env tile entries: `sp15_dd_trajectory_per_env` + `sp15_dd_trajectory_prev_dd_per_env`. Both reset arms `host_slice_mut().fill(0.0)` mirror the `sp15_dd_state_per_env` pattern. The `sp15_dd_trajectory_decreasing` ISV-slot reset arm (slot 439) was UNCHANGED — slot 439 remains a HEALTH_DIAG diagnostic scalar (now mean-aggregate from the per-env tile via the reduce kernel) so its FoldReset 0 sentinel still applies. (9) **Layout fingerprint break** — added marker `DD_TRAJECTORY_PER_ENV=sp15_phase_1_3_b_followup_B` to `layout_fingerprint_seed`. Pre-followup-B checkpoints WILL NOT LOAD after this commit (greenfield OK per spec Q1 — same precedent as the Phase 1.3.b-followup `DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` fingerprint break). (10) **Migrated 4 oracle tests + 1 new behavioral test** — `dd_trajectory_decreasing_fires_during_recovery` / `dd_trajectory_no_fire_when_dd_increasing` / `dd_trajectory_no_fire_when_prev_below_floor` (all 3 set up per-env tiles with `n_envs=1`; tile slot 5 holds DD_PCT; the assertions read the per-env trajectory output rather than ISV[439] — preserves bit-identical semantics for single-env scenarios while exercising the per-env contract). `per_sampler_weights_recovery_transitions_higher` (Wave 4.3 test; migrated to set per-env tile `dd_trajectory_per_env[0]` instead of ISV[439] across the two batches; preserves the recovery=1 vs recovery=0 toggle assertion). NEW: `per_sampler_weights_per_transition_not_uniform_across_batch` — two-env config (n_envs=2, lookback=2, batch_size=8) with env-0's trajectory=1 and env-1's trajectory=0, BOTH inserted in the SAME batch. The test directly fails the Wave 4.3 implementation (which would produce uniform 3.0× OR uniform 1.0× across all 8 priorities) and passes the followup-B per-transition implementation: priorities[0..2]=3.0 (env-0 main), priorities[2..4]=1.0 (env-1 main), priorities[4..6]=3.0 (env-0 cf), priorities[6..8]=1.0 (env-1 cf). The boost pattern verifies that each transition's env_id mapping `env_id = (j % (n_envs * lookback)) / lookback` correctly threads its env's flag — env_0_mean=3.0, env_1_mean=1.0, env_0_mean - env_1_mean ≥ 1.0. The 2 oracle tests that exercise `dd_state_reduce_kernel` without the per-env trajectory tile (`dd_state_kernel_tracks_drawdown_correctly` / `dd_state_per_env_diverge_independently`) pass `0` (null) for the new `dd_trajectory_per_env` parameter — preserves their dd_state-only assertions bit-identically (the kernel skips the slot 439 write when null). (11) **Closes the per-env DD redesign chain end-to-end** per `feedback_wire_everything_up`: every consumer of "current per-env DD context" (final reward shaping, plasticity-injection trigger, HEALTH_DIAG, recovery-curriculum PER oversample) now reads ITS env's context via per-env tile + per-(i,t) lookup; no more "whichever env's write landed last wins" paths. The Wave 4.3 statistical-bias bug is fixed end-to-end. (12) **Atomicity** per `feedback_no_partial_refactor`: kernel reshape (`dd_trajectory_decreasing_kernel`) + buffer alloc (2 collector-owned per-env tiles) + trainer struct cleanup (delete `[1]` scratch + ctor + destructure) + reduction kernel extension (`dd_state_reduce_kernel` shared-mem grow + tree-reduce extra lane) + reduction launcher signature change + per_insert_pa kernel signature change + 3 new GPU PER replay buffer fields/setters/accessors + per_insert_pa launch site update + collector launch sequence update (move dd_trajectory into dd_state block) + state-reset registry rename + 2 dispatch arm renames + training_loop wiring block delete + new training_loop wiring block (per-env tile + dims into PER buffer) + 2 dd_state_reduce oracle test callsite updates (pass null) + 4 dd_trajectory oracle test migrations + 1 PER oracle test migration + 1 NEW behavioral oracle test + audit doc entry all in one commit; pre-existing tests preserve their pre-followup-B semantics bit-identically (the dd_state_reduce null-trajectory and dd_trajectory single-env paths). Files touched: `crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu` (reshape to per-env grid), `crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu` (extend with dd_trajectory_per_env mean-aggregate to ISV[439]), `crates/ml-dqn/src/per_kernels.cu` (per-transition lookup signature change), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (delete `sp15_dd_trajectory_prev_dd` field + ctor + destructure; update `launch_sp15_dd_trajectory_decreasing` signature; update `launch_sp15_dd_state_reduce` signature; layout fingerprint marker), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (add `sp15_dd_trajectory_per_env` + `sp15_dd_trajectory_prev_dd_per_env` fields + ctor + struct literal; delete `sp15_dd_trajectory_prev_dd_dev_ptr` field + setter; merge dd_trajectory launch into dd_state block; update reduce launch with new arg), `crates/ml-dqn/src/gpu_replay_buffer.rs` (add 3 new fields + 2 setters + 2 accessors; update per_insert_pa launch with 3 new args), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (replace 1 entry with 2 collector-owned per-env tile entries), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (delete `set_sp15_dd_trajectory_prev_dd_ptr` wiring; replace 1 dispatch arm with 2; add new PER buffer wiring block for per-env tile + dims), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (4 dd_trajectory test migrations + 1 PER test migration + 1 new behavioral test + 2 dd_state_reduce callsite null updates), `docs/dqn-wire-up-audit.md` (this entry). Hard rules: `feedback_no_partial_refactor` (kernel + buffers + PER sampler + tests + audit doc atomic — every consumer of the kernel signature change migrates simultaneously), `feedback_no_atomicadd` (per-env grid one thread per env, no atomics; reduction uses block tree-reduce), `feedback_no_cpu_compute_strict` (per-env logic + reduction GPU-resident; host-side wiring is cold-path setter calls only), `feedback_isv_for_adaptive_bounds` (the mean aggregation rule for ISV[439] is structural, NOT adaptive — same as the existing slot 401-406 mean-aggregates), `feedback_no_legacy_aliases` (the trainer-owned `[1]` `sp15_dd_trajectory_prev_dd` was DELETED rather than kept as a legacy alias — production paths use the collector-owned per-env tiles unconditionally), `feedback_wire_everything_up` (the deferred Phase 1.3.b-followup point (10) is closed end-to-end; the Phase 1.3.b-followup audit doc entry's open paragraph is now historical context for THIS commit). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean; `SQLX_OFFLINE=true cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored dd_state` 2 of 2 dd_state oracle tests green; `... --ignored dd_trajectory` 3 of 3 dd_trajectory oracle tests green (post-migration to per-env contract); `... --ignored per_sampler` 2 of 2 PER oracle tests green (Wave 4.3 migration + new behavioral test); `... --ignored final_reward` 6 of 6 final-reward oracle tests green (per-env DD lookup chain); `... --ignored plasticity` 5 of 5 plasticity oracle tests green (DD_PERSISTENCE_MAX_INDEX=443 consumption unchanged); `cargo test -p ml --features cuda --lib` HOLDS the Phase 1.3.b-followup baseline (947 pass / 12 fail) on RTX 3050 Ti — no new regressions, the 12 failing tests are the same pre-existing infra failures.