From 23e9a1f78c2a63459c7b108d398bc59e1ce2f837 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 7 May 2026 11:39:34 +0200 Subject: [PATCH] =?UTF-8?q?fix(cuda):=20compute=5Fexpected=5Fq=20stride=20?= =?UTF-8?q?13=20vs=20denoise=5Ftarget=5Fq=5Fbuf=20size=2012=20=E2=80=94=20?= =?UTF-8?q?OOB=20writes=20threads=2060-63?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Pre-existing latent bug surfaced by compute-sanitizer after the SP15 NULL-pointer fix at 2e37af29d unblocked the cascade. Threads 60-63 were writing past denoise_target_q_buf's end every batch. Sanitizer evidence (pre-fix on 2e37af29d): Invalid __global__ write of size 4 bytes at compute_expected_q+0x2480 by thread (60..63, 0, 0) in block (0, 0, 0) Access at is out of bounds (45/97/149/201 bytes after nearest allocation of size B*12*4 bytes) Host backtrace: GpuDqnTrainer::compute_denoise_target_q → submit_post_aux_ops → run_full_step Root cause — semantic mismatch between writer stride and buffer size: - compute_expected_q writes q_values[i*total_actions + a] with total_actions = b0+b1+b2+b3 = 4+3+3+3 = 13 (4-direction factored) - denoise_target_q_buf was allocated b*12 (legacy 3+3+3+3 layout) - threads 60..63 wrote slot 12 of samples (60..63 % batch_size) past the buffer end every step The downstream q_denoise_backward + 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 the identical OOB at threads 60..63); denoise_target_q_buf 8 lines below was missed 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 pattern as q_var_buf_trainer): 1. Widen denoise_target_q_buf from b*12 to b*total_actions (= b*13) to match compute_expected_q's writer stride. 2. Add refined_stride / target_stride / input_stride parameters to q_denoise_backward kernel; the kernel still computes D=12 per- sample (denoiser MLP fixed width) but addresses each input buffer at its own per-sample stride. 3. Add refined_stride / target_stride parameters to denoise_loss_grad kernel (same pattern; used by launch_q_denoise_backward_cublas). 4. Update both Rust launchers (launch_q_denoise_backward, launch_q_denoise_backward_cublas) to pass refined_stride=12 (q_coord and q_input are post-attention narrowed), target_stride=total_actions=13. 5. denoise_q_input_buf STAYS at b*12 — it's a snapshot of q_coord_buf (also b*12) via snapshot_pre_denoise_q's DtoD copy; never written by compute_expected_q. 6. Flat-scan consumers (SP4 target_q_p99_update producer + SP3 slot 46 threshold-check + dqn_clamp_finite_f32) UNCHANGED — they consume the buffer flat via .len(); widening from b*12 to b*13 is monotone (one more valid Q-value per sample in the histogram). Atomic per feedback_no_partial_refactor: 2 kernel signatures + buffer allocation + 2 launch sites + struct doc comments + SP3 slot 46 doc + audit doc — all in one commit. Every consumer of the writer's stride migrates simultaneously. Verification: - SQLX_OFFLINE=true cargo check -p ml --features cuda: clean - compute-sanitizer (RTX 3050 Ti): 0 Invalid __global__ errors at compute_expected_q post-fix (down from 4 per training step in baseline — re-verified on 2e37af29d to confirm the diff) - SQLX_OFFLINE=true cargo test -p ml --features cuda --lib: holds parent baseline 947 pass / 12 fail (same 12 pre-existing failures) 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 — audit entry Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 37 +++++++-- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 78 +++++++++++++++---- docs/dqn-wire-up-audit.md | 2 + 3 files changed, 95 insertions(+), 22 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index fc9f2f61b..73a0e0a34 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -6245,8 +6245,15 @@ extern "C" __global__ void denoise_silu_bwd( /** * denoise_loss_grad — MSE gradient with 0.1x residual scaling. - * Reads row-major [B, D], writes col-major [D, B]. - * d_res[d, b] = 0.1 * 2.0 * (Q_refined[b*D+d] - Q_target[b*D+d]) * inv_B. + * Reads row-major [B, refined_stride] and [B, target_stride], writes col-major [D, B]. + * d_res[d, b] = 0.1 * 2.0 * (Q_refined[b*refined_stride+d] - Q_target[b*target_stride+d]) * inv_B. + * + * refined_stride / target_stride MAY exceed D when the upstream buffers are + * sized per total Q-actions (e.g. 13 for the 4+3+3+3 factored layout) instead + * of the legacy D=12 exposure layout. The denoiser MSE only consumes the + * first D entries per sample; the trailing slot(s) are ignored. Per-sample + * stride must be passed so addressing matches the buffer layout — otherwise + * sample b reads slot 12 of the previous sample. * * Col-major [D, B] output: element (d, b) at offset d + b * D. * Grid: ceil(D*B/256), Block: 256. @@ -6255,14 +6262,16 @@ extern "C" __global__ void denoise_loss_grad( const float* __restrict__ Q_refined, const float* __restrict__ Q_target, float* __restrict__ d_res, - int B, int D, float inv_B + int B, int D, + int refined_stride, int target_stride, + float inv_B ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; int total = D * B; if (idx >= total) return; int d = idx % D; int b = idx / D; - d_res[idx] = 0.2f * (Q_refined[b * D + d] - Q_target[b * D + d]) * inv_B; + d_res[idx] = 0.2f * (Q_refined[b * refined_stride + d] - Q_target[b * target_stride + d]) * inv_B; } /** @@ -6344,7 +6353,16 @@ extern "C" __global__ void denoise_bias_grad_p2( */ /* var_q stride matches the variance buffer's per-sample width (>= D). The * denoiser only uses the first D entries per sample; pass var_stride from the - * Rust launch site to keep alignment when q_var is sized per total Q-actions. */ + * Rust launch site to keep alignment when q_var is sized per total Q-actions. + * + * refined_stride / target_stride / input_stride mirror this pattern for the + * Q buffers themselves: each may be >= D when the upstream buffer is sized + * per total Q-actions (13 for 4+3+3+3 factored layout). The denoiser MSE + * + replay only consumes the first D entries per sample. Compute-sanitizer + * caught the legacy hard-coded `b * D` indexing as 4 OOB writes per batch + * (threads 60..63) when compute_expected_q wrote 13 floats per sample into + * a `b*12` target buffer (commit log: see fix(cuda) compute_expected_q + * stride 13 vs denoise_target_q_buf size 12). */ extern "C" __global__ void q_denoise_backward( const float* __restrict__ Q_refined, const float* __restrict__ Q_target, @@ -6353,7 +6371,10 @@ extern "C" __global__ void q_denoise_backward( const float* __restrict__ Q_input, const float* __restrict__ var_q, int B, - int var_stride) + int var_stride, + int refined_stride, + int target_stride, + int input_stride) { int widx = blockIdx.x * blockDim.x + threadIdx.x; const int D = 12; @@ -6385,7 +6406,7 @@ extern "C" __global__ void q_denoise_backward( /* Compute loss gradient for this sample */ float d_out_b[12]; for (int i = 0; i < D; i++) { - d_out_b[i] = 2.0f * (Q_refined[b * D + i] - Q_target[b * D + i]) * inv_B; + d_out_b[i] = 2.0f * (Q_refined[b * refined_stride + i] - Q_target[b * target_stride + i]) * inv_B; } /* Process steps in reverse. The gradient for the current step is needed. */ @@ -6395,7 +6416,7 @@ extern "C" __global__ void q_denoise_backward( /* Reconstruct Q at start of `step` by running forward from Q_input * through steps 0..(step-1). */ float q_curr[12]; - for (int i = 0; i < D; i++) q_curr[i] = Q_input[b * D + i]; + for (int i = 0; i < D; i++) q_curr[i] = Q_input[b * input_stride + i]; for (int s = 0; s < step; s++) { int soff = s * STEP_PARAMS; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 3f1e6c057..7e25f26c4 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -6203,7 +6203,7 @@ pub struct GpuDqnTrainer { // ── SP4 Layer A Task A5: target_q_p99 producer ───────────────────── /// SP4 first end-to-end producer kernel. Single-block kernel that reads - /// `denoise_target_q_buf` (target-network Q-values [B, 12] flattened), + /// `denoise_target_q_buf` (target-network Q-values [B, total_actions] flattened), /// computes p99(|target_q|) via the `sp4_histogram_p99<256>` header-only /// device function, writes the step observation to /// `producer_step_scratch_buf[0]`. Loaded from @@ -7001,7 +7001,11 @@ pub struct GpuDqnTrainer { q_denoise_backward_kernel: CudaFunction, /// Gradient accumulator for denoise_params [1800]. Zeroed before each backward. denoise_grad: CudaSlice, - /// Target network Q-values [B, 12] for denoiser training signal. + /// Target network Q-values [B, total_actions] for denoiser training signal. + /// Sized at `total_actions = b0+b1+b2+b3` (= 13 for the 4-direction + /// factored layout) to match the `compute_expected_q` writer stride. + /// Denoiser consumers (q_denoise_backward, denoise_loss_grad) read only + /// the first D=12 slots per sample via the `target_stride` parameter. denoise_target_q_buf: CudaSlice, /// Pre-denoise Q-values snapshot [B, 12] — saved before denoiser forward for backward. denoise_q_input_buf: CudaSlice, @@ -11020,15 +11024,23 @@ impl GpuDqnTrainer { (end - start) / std::mem::size_of::() } - /// SP3 slot 46: device pointer to `denoise_target_q_buf` `[B, 12]`, - /// the target-network Q-values written by the C51 target-q compute pass. - /// The threshold-check kernel scans this for post-clip outliers (Task B2 - /// extends the clipping pipeline to feed this buffer). + /// SP3 slot 46: device pointer to `denoise_target_q_buf` `[B, total_actions]`, + /// the target-network Q-values written by the C51 target-q compute pass via + /// `compute_expected_q` (which strides by `total_actions = b0+b1+b2+b3 = 13` + /// in the 4-direction factored layout). The threshold-check kernel scans + /// the full buffer flat — all entries are valid Q-values, so widening from + /// the legacy `b*12` exposure-layout sizing to `b*total_actions` is + /// monotone (more samples in the histogram, no semantic change). Per-sample + /// consumers (q_denoise_backward, denoise_loss_grad) receive + /// `target_stride = total_actions` and read only the first D=12 entries — + /// matching the cross-branch-attention narrowed denoiser pipeline. pub(crate) fn target_q_ptr(&self) -> u64 { self.denoise_target_q_buf.raw_ptr() } - /// Element count of the target-q buffer (= `B * 12`). + /// Element count of the target-q buffer (= `B * total_actions`, e.g. `B * 13` + /// for the 4-direction factored layout). Used by SP4 p99 producer + + /// SP3 threshold-check kernels for flat scans of the full buffer. pub(crate) fn target_q_len(&self) -> usize { self.denoise_target_q_buf.len() } @@ -12777,7 +12789,7 @@ impl GpuDqnTrainer { /// SP4 Layer A Task A5: producer for ISV[TARGET_Q_BOUND_INDEX=131]. /// - /// Reads `denoise_target_q_buf` (target-network Q-values [B, 12]), launches + /// Reads `denoise_target_q_buf` (target-network Q-values [B, total_actions]), launches /// the single-block `target_q_p99_update` kernel which computes /// p99(|target_q|) via the header-only `sp4_histogram_p99<256>` device /// function and writes the step observation to @@ -21670,7 +21682,21 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("q_denoise_backward load: {e}")))?; let denoise_grad = stream.alloc_zeros::(DENOISE_TOTAL_PARAMS) .map_err(|e| MLError::ModelError(format!("alloc denoise_grad: {e}")))?; - let denoise_target_q_buf = alloc_f32(&stream, b * 12, "denoise_target_q_buf")?; + // denoise_target_q_buf is written by `compute_expected_q` which strides + // by `total_actions` (= 4+3+3+3 = 13 in the 4-direction factored layout). + // The legacy `b*12` sizing was the 12-slot exposure layout (3+3+3+3) and + // produced 4 OOB writes per batch (threads 60..63) under compute-sanitizer. + // Sized at `b*total_actions` to match the kernel writer; downstream + // denoiser consumers (q_denoise_backward, denoise_loss_grad) receive + // `target_stride = total_actions` and read only the first D=12 entries + // per sample. Same pattern as q_var_buf_trainer (see comment above + // line 21629). + // denoise_q_input_buf is a snapshot of q_coord_buf which stays at b*12 + // (the cross-branch attention output narrows 13→12 before the denoiser + // consumes it), so its sizing is unchanged. The kernel input_stride + // parameter still threads through for symmetry with target/refined. + let q_var_total_actions_for_target = config.branch_0_size + config.branch_1_size + config.branch_2_size + config.branch_3_size; + let denoise_target_q_buf = alloc_f32(&stream, b * q_var_total_actions_for_target, "denoise_target_q_buf")?; let denoise_q_input_buf = alloc_f32(&stream, b * 12, "denoise_q_input_buf")?; let denoise_norm_partials = stream.alloc_zeros::(1) .map_err(|e| MLError::ModelError(format!("alloc denoise_norm_partials: {e}")))?; @@ -26158,8 +26184,12 @@ impl GpuDqnTrainer { Ok(()) } - /// Compute target network Q-values [B, 12] into denoise_target_q_buf. + /// Compute target network Q-values [B, total_actions] into denoise_target_q_buf. /// Uses compute_expected_q on target logits (tg_v_logits_buf, tg_b_logits_buf). + /// `compute_expected_q` writes `total_actions = b0+b1+b2+b3 = 13` floats per + /// sample (4-direction factored layout); the buffer is sized to match. + /// Per-sample consumers thread `target_stride = total_actions` to read only + /// the first D=12 entries. pub(crate) fn compute_denoise_target_q(&self, batch_size: usize) -> Result<(), MLError> { let n = batch_size as i32; let na = self.config.num_atoms as i32; @@ -26230,10 +26260,19 @@ impl GpuDqnTrainer { let d_params_ptr = self.denoise_grad.raw_ptr(); let q_input_ptr = self.denoise_q_input_buf.raw_ptr(); let var_ptr = self.q_var_buf_trainer.raw_ptr(); - // q_var_buf_trainer is sized [B, total_actions]; denoiser only reads - // the first D=12 slots per sample, but the kernel needs the true - // stride so per-sample addressing matches the buffer layout. + // Per-sample strides for the four [B, ?] inputs. The denoiser MLP only + // reads the first D=12 slots per sample, but each upstream buffer may + // be sized wider — the kernel needs the true stride to align per-sample + // addressing. + // + // q_var_buf_trainer: [B, total_actions=13] (compute_expected_q writes) + // denoise_target_q_buf: [B, total_actions=13] (compute_expected_q writes) + // q_coord_buf: [B, 12] (cross-branch attention output) + // denoise_q_input_buf: [B, 12] (snapshot of q_coord_buf) let var_stride = self.total_actions() as i32; + let target_stride = self.total_actions() as i32; + let refined_stride = 12_i32; + let input_stride = 12_i32; unsafe { self.stream .launch_builder(&self.q_denoise_backward_kernel) @@ -26245,6 +26284,9 @@ impl GpuDqnTrainer { .arg(&var_ptr) .arg(&b) .arg(&var_stride) + .arg(&refined_stride) + .arg(&target_stride) + .arg(&input_stride) .launch(LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (256, 1, 1), @@ -26508,9 +26550,15 @@ impl GpuDqnTrainer { // BACKWARD // ═══════════════════════════════════════════════════════ - // 10. Loss gradient: d_res = 0.1 * 2 * (Q_refined - Q_target) / B + // 10. Loss gradient: d_res = 0.1 * 2 * (Q_refined - Q_target) / B. + // Q_refined = q_coord_buf [B, 12] (post-cross-branch-attention). + // Q_target = denoise_target_q_buf [B, total_actions=13] (post-compute_expected_q). + // Kernel reads only the first D=12 entries from each per sample but + // needs per-sample strides to handle the asymmetric layout. let inv_b: f32 = 1.0 / b as f32; let d_res_ptr = self.denoise_d_res_buf.raw_ptr(); + let refined_stride: i32 = 12; + let target_stride: i32 = self.total_actions() as i32; unsafe { self.stream.launch_builder(&self.denoise_loss_grad_kernel) .arg(&q_refined_ptr) @@ -26518,6 +26566,8 @@ impl GpuDqnTrainer { .arg(&d_res_ptr) .arg(&(b as i32)) .arg(&d_val) + .arg(&refined_stride) + .arg(&target_stride) .arg(&inv_b) .launch(LaunchConfig { grid_dim: (((D * b) as u32 + 255) / 256, 1, 1), diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index beb0ea8e2..4c16fc461 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. +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. SP15 Phase 1.3.b-followup — per-env DD redesign (Path A → Path B atomic migration) (2026-05-07): closes the Phase 1.3.b deferred per-env redesign per `feedback_no_partial_refactor`. The original Phase 1.3.b atomic fix (`132609724`) chose **env 0 as the canonical DD observable** and wrote 6 ISV scalar slots [401..407) directly from env-0's portfolio row. That choice was a documented Path A pragmatic compromise — every downstream consumer (the fused reward composer, plasticity injection trigger, dd_trajectory recovery proxy, HEALTH_DIAG diagnostic) read the env-0-canonical scalars even though each production env has its OWN portfolio + DD trajectory. When env-3 was in 30% DD and env-0 was at ATH, the model learning from env-3's transitions saw `dd_pct=0` and silently skipped the recovery shaping. (1) **`dd_state_kernel.cu` reshape** — grid `[n_envs, 1, 1]` × `[1, 1, 1]`, one thread per env, sequential per-bar walk per env. Each thread computes its env's DD trajectory and writes 6 scalars to a NEW per-env tile buffer `dd_state_per_env[n_envs * 6]` (layout `[env_0_current, env_0_max, env_0_rec, env_0_pers, env_0_calmar, env_0_pct, env_1_current, ...]`); the kernel no longer touches the ISV scalar slots. (2) **New `dd_state_reduce_kernel.cu`** — single-block tree-reduce (no `atomicAdd` per `feedback_no_atomicadd`; BLOCK=256 with strided initial accumulation handles n_envs up to 32768 on H100). Mean-aggregates across envs into the 6 scalar ISV slots [401..407) (smooth across-env summary for HEALTH_DIAG diagnostic — preserves backward-compatible "this is one DD scalar to look at" semantics; only the across-env aggregation rule changed Path A: env-0-canonical → Path B: mean-across-envs); max-aggregates DD_PERSISTENCE into a NEW slot **`DD_PERSISTENCE_MAX_INDEX = 443`** (consumed by `plasticity_injection_kernel` so the global advantage-head reset fires when ANY env's persistence exceeds the threshold — one set of advantage weights → global firing semantics → max-aggregate is the only correct rule; mean-aggregate would silently gate the trigger when only a fraction of envs are in long DD). Aggregation rule rationale per `feedback_isv_for_adaptive_bounds` (the rule itself is structural, not adaptive). `ISV_TOTAL_DIM 443 → 444` to fit the new slot; `SP15_SLOT_END 443 → 444`; `SP15_SLOT_COUNT 46 → 47`. (3) **`compute_sp15_final_reward_kernel.cu` per-env DD lookup** — index→env mapping `env_id = (idx % (N*L)) / L` (the layout is `[on_policy_N*L | cf_N*L]` so `slot_it = idx % (N*L)` is the per-(i,t) flat index, and `env_id = slot_it / L`); each `(i,t)` slot looks up ITS env's `dd_pct` and `dd_current` from the per-env tile (`tile[env_base + 5]` and `tile[env_base + 0]` respectively). On-policy and CF threads at the same (i,t) read the SAME tile entry (one DD trajectory per env, shared across slot kinds). The `__device__` helpers `sp15_dd_asymmetric_reward` / `sp15_dd_penalty` migrated from "read ISV by themselves" to "accept dd_pct + dd_current as scalar parameters" so the per-(i,t) per-env lookup happens at the kernel level, with the helpers staying scalar-clean. Atomic per `feedback_no_partial_refactor` — both gain-side multiplier (Stage 2) and quadratic DD penalty (Stage 3) migrated together. (4) **`plasticity_injection_kernel.cu` consumer migration** — persistence read switched from ISV[DD_PERSISTENCE_INDEX=404] (mean-aggregate across envs) to ISV[DD_PERSISTENCE_MAX_INDEX=443] (max-aggregate across envs). The fired-this-fold flag and warm-bars counter remain global scalars (one set of advantage weights ⇒ one trigger condition ⇒ one fired flag); the kernel itself stays unchanged downstream of the read. (5) **HEALTH_DIAG semantic shift (documented breaking change)** — slots 401-406 now report the cross-env MEAN of DD_CURRENT/DD_MAX/DD_RECOVERY_BARS/DD_PERSISTENCE/CALMAR/DD_PCT, NOT env-0's value. For single-env smoke configs (n_envs=1) the mean equals env-0's value, so the diagnostic stream is bit-stable across the migration. For production (n_envs ≥ 2048) the mean is genuinely different — but it's the right diagnostic: a smoothed across-env summary of "where is the training instance, on average, in its DD trajectory". (6) **Per-env tile owned by `GpuExperienceCollector`** (NOT the trainer) — the collector knows `alloc_episodes` (= `n_envs`); the trainer's `batch_size` is a different quantity (training batch size). The new field `sp15_dd_state_per_env: MappedF32Buffer([alloc_episodes * 6])` is allocated in the constructor next to `sp13_hold_rate_buf` and reset to zero at fold boundary via the `sp15_dd_state_per_env` registry-arm dispatch (`host_slice_mut().fill(0.0)` — mirrors the `sp15_alpha_warm_count` / `sp15_dd_trajectory_prev_dd` non-ISV mapped-pinned reset pattern). (7) **Per-step launch order in `gpu_experience_collector.rs::launch_timestep_loop` step 5b**: `dd_state_kernel` (per-env tile producer) → `dd_state_reduce_kernel` (ISV scalar consumer; gated on `isv_signals_dev_ptr != 0`) → `alpha_split_producer_kernel` (existing) → `compute_sp15_final_reward_kernel` (per-(i,t) consumer; threads `dd_state_per_env_ptr` so each slot's reward shaping uses ITS env's DD context). All four launches on the same stream — CUDA serialises producer→consumer without explicit event sync. Outside the exp-fwd graph capture region per `pearl_no_host_branches_in_captured_graph`. (8) **Layout fingerprint break** — added markers `DD_PERSISTENCE_MAX=443; ISV_TOTAL_DIM=444; DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` to `layout_fingerprint_seed`; pre-followup checkpoints WILL NOT LOAD after this commit (greenfield OK per spec Q1 — same precedent as the Wave 4.1a `TRUNK_INPUT_DD_PCT=sp15_wave_4_1a` fingerprint break). (9) **Migrated 6 oracle tests + added 1 new behavioral test** — `dd_state_kernel_tracks_drawdown_correctly` (now drives both kernels with `n_envs=1` so the existing assertions still hold via the mean-of-one identity); `plasticity_fires_when_persistence_exceeds_threshold` / `plasticity_debounced_within_fold` / `plasticity_no_fire_below_threshold` / `plasticity_fires_force_flat_then_cooldown_holds` (writes seeded ISV[DD_PERSISTENCE_MAX_INDEX=443] instead of the now-stale ISV[DD_PERSISTENCE_INDEX=404]); 5 `final_reward_*` tests (now thread `dd_tile.dev_ptr` through `launch_sp15_final_reward` and write per-env DD into `dd_tile[0..6]` instead of ISV[401..407)). NEW: `dd_state_per_env_diverge_independently` — two-env config where env-0 is in recovery (DD shrinking) and env-1 is deepening into DD; verifies the per-env tile reflects each env's INDEPENDENT trajectory, the mean-aggregated ISV scalars equal the cross-env mean, and the max-aggregated ISV[DD_PERSISTENCE_MAX_INDEX=443] equals the WORSE env's persistence — the canonical "envs in different DD contexts thread their own DD context independently" guarantee that motivated the Path A → Path B migration. (10) **Phase 1.3.b-followup-B (separate split)** — `dd_trajectory_decreasing_kernel` + `per_insert_pa` migration is NOT in scope for this commit. The current Wave 4.3 implementation reads ISV[DD_TRAJECTORY_DECREASING_INDEX=439] at insert-batch time (epoch end), which carries the LAST step's value applied uniformly to ALL inserted transitions — a known pre-existing limitation where ANY env's trajectory could win the boost regardless of the inserting transition's actual env. The proper fix requires a per-(env, t) trajectory buffer `[N*L]` written by `dd_trajectory_kernel` and an `env_id`-aware lookup in `per_insert_pa` (the insert position `j` of the batch maps to `env_id = (j % (N*L)) / L` since the inserted batch has the same `[on_policy_N*L | cf_N*L]` layout). That's a fundamentally different contract change (per-(env, t) tile + env_id-mapped lookup at insert time, NOT just per-env tile + reduction). Splitting it into Phase 1.3.b-followup-B preserves the no-partial-refactor invariant within each migration: contract A (DD state ISV scalar → per-env tile) lands here atomically with all 5 of its consumers (final_reward kernel + plasticity + HEALTH_DIAG + 6 oracle tests + new behavioral test); contract B (DD trajectory ISV scalar → per-(env, t) tile) lands separately with its lone consumer (`per_insert_pa`). Files touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (reshape), `dd_state_reduce_kernel.cu` (NEW), `compute_sp15_final_reward_kernel.cu` (per-env lookup), `plasticity_injection_kernel.cu` (slot 404→443), `sp15_reward_axis_helpers.cuh` (helpers take dd_pct/dd_current as scalars), `sp15_isv_slots.rs` (+ DD_PERSISTENCE_MAX_INDEX), `gpu_dqn_trainer.rs` (launchers + ISV_TOTAL_DIM + fingerprint), `gpu_experience_collector.rs` (per-env tile field + 4-launch sequence), `state_reset_registry.rs` (+ 2 entries), `training_loop.rs` (+ 2 dispatch arms), `build.rs` (+ dd_state_reduce_kernel cubin), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (6 migrated + 1 new test), `docs/isv-slots.md` (slot map), `docs/dqn-wire-up-audit.md` (this entry). Hard rules: `feedback_no_partial_refactor` (5 consumers + 6 oracle tests + 1 new test all atomic in this commit); `feedback_no_atomicadd` (reduction kernel uses block tree-reduce, not atomicAdd); `feedback_no_cpu_compute_strict` (per-env logic + reduction GPU-resident); `feedback_isv_for_adaptive_bounds` (mean/max aggregation rules are structural, documented rationale).