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.