From 21f911151a46139a5fd677b0b875197db8520208 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 22:32:07 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp21):=20T2.2=20Phase=203=20=E2=80=94=20wi?= =?UTF-8?q?re=20E1=20q=5Fcorrection=20=E2=86=92=20Bellman=20target=20offse?= =?UTF-8?q?t=20(atomic)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Wires `EnrichmentResult::q_correction` (Phase 1.5+2's E1 clamp(-mean(predicted_q − pnl), ±10.0) from the per-trade tape) into the training loss kernels' Bellman target computation. Both mse_loss_batched and c51_loss_kernel::block_bellman_project_f apply q_correction as additive shift on the Bellman target — blended (1-α)·MSE + α·C51 sees identical Q-target shifts (symmetric application per feedback_no_partial_refactor). ISV slot allocation (per pearl_controller_anchors_isv_driven + feedback_isv_for_adaptive_bounds — adaptive Q-target shift IS adaptive bound, lives in ISV bus): - New Q_CORRECTION_INDEX = 520 in cuda_pipeline::sp21_isv_slots - ISV_TOTAL_DIM 520 → 521 (bus extension) - Layout fingerprint adds SLOT_520_Q_CORRECTION - New module registered in cuda_pipeline/mod.rs - Compile-time test verifies SP21_SLOT_END <= ISV_TOTAL_DIM Sign convention: q_correction is the additive correction (E1 already negates the bias). predicted_q > pnl → q_correction < 0 → "lower Q-targets". predicted_q < pnl → q_correction > 0 → "raise Q-targets". Cold-start sentinel 0.0 (no trades yet) is no-op until first val pass writes (per pearl_first_observation_bootstrap). ABI surgery (minimal): - mse_loss_batched: ONE new scalar arg (float q_correction) at end - c51_loss_batched: ZERO new args (block_bellman_project_f already takes isv_signals; reads slot 520 from there) - launch_mse_loss: reads ISV[520] host-side via existing read_isv_signal_at, passes scalar - training_loop.rs (post-enrichment): writes result.q_correction to ISV[520] via existing write_isv_signal_at Files changed: - crates/ml/src/cuda_pipeline/sp21_isv_slots.rs (NEW): slot const + bounds-check tests - crates/ml/src/cuda_pipeline/mod.rs: pub mod sp21_isv_slots - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs: ISV_TOTAL_DIM bump + fingerprint update + launch_mse_loss launcher arg - crates/ml/src/cuda_pipeline/mse_loss_kernel.cu: q_correction arg + application to target_q - crates/ml/src/cuda_pipeline/c51_loss_kernel.cu: read isv_signals[520] in block_bellman_project_f, add to t_z - crates/ml/src/trainers/dqn/trainer/training_loop.rs: write_isv_signal_at( Q_CORRECTION_INDEX, result.q_correction) post-enrichment - docs/dqn-wire-up-audit.md: 2026-05-11 audit entry Verification: - cargo check -p ml --tests --features cuda: 0 errors - cargo test -p ml --lib sp21_isv_slots --features cuda: 2/2 (bus bounds + slot uniqueness) - sp20_aggregate_inputs_test: 12/12 - sp20_phase1_4_wireup_test: 2/2 - sp20_emas_compute_test: 4/4 - sp20_controllers_compute_test: 7/7 - sp21_per_trade_predicted_q_test: 3/3 Total: 30 tests, 0 failures. Behavioral gate (q_correction → non- zero ISV[520] → Q-target shift) is the upcoming smoke training run. After this commit (T2.2 Phases 4-8): - Phase 4: E4 per-branch LR scaling via per-group Adam - Phase 5: E6 winner indices → PER priority bumps - Phase 6: E7 hindsight → replay buffer injection - Phase 7: E8 curriculum weights → segment sampling - Phase 8: signal-drive remaining controller GAINS Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/c51_loss_kernel.cu | 25 ++- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 20 ++- crates/ml/src/cuda_pipeline/mod.rs | 1 + .../ml/src/cuda_pipeline/mse_loss_kernel.cu | 23 ++- crates/ml/src/cuda_pipeline/sp21_isv_slots.rs | 87 ++++++++++ .../src/trainers/dqn/trainer/training_loop.rs | 16 ++ docs/dqn-wire-up-audit.md | 161 ++++++++++++++++++ 7 files changed, 328 insertions(+), 5 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/sp21_isv_slots.rs diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index f5bbec251..d76b95bef 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -230,13 +230,34 @@ __device__ void block_bellman_project_f( reward_bias = deficit * (1.0f - health); } + /* SP21 T2.2 Phase 3 (2026-05-11): E1 q_correction additive Bellman + * offset. Read once per sample from `isv_signals[520]` (Q_CORRECTION_ + * INDEX). Applied uniformly across all branches / actions / atoms + * for this sample — the correction is a global Q-bias estimate + * from the validation backtest's predicted_q − pnl statistic + * (E1: `enrichment::compute_q_correction`). Cold-start sentinel + * (`isv_signals[520] = 0.0`) makes this a no-op until the first + * val pass writes a real correction; per + * `pearl_first_observation_bootstrap`. NULL-tolerant on + * `isv_signals` to preserve backward compatibility with test + * scaffolds that pass NULL (q_correction defaults to 0.0). */ + float q_correction = (isv_signals != NULL) ? isv_signals[520] : 0.0f; + for (int j = tid; j < num_atoms; j += BLOCK_THREADS) { float z_j = support[j]; /* Task 2.Y-ext v5: add per-sample reward-bias BEFORE gamma×z term. * Identity (reward_bias == 0) on non-direction branches, on * Hold/Flat samples, and when the tradable bin already leads - * the max pathology bin by ≥ lead_scale. */ - float t_z = (reward + reward_bias) + gamma * z_j * (1.0f - done); + * the max pathology bin by ≥ lead_scale. + * + * SP21 T2.2 Phase 3 (2026-05-11): also adds `q_correction` from + * `isv_signals[520]` — a global E1 Q-bias correction from the + * validation backtest. Both bias terms are scalar offsets on + * the Bellman target; composing additively at the same site + * keeps the projection's discretization error minimised + * (single arithmetic operation instead of two passes through + * the support grid). */ + float t_z = (reward + reward_bias + q_correction) + gamma * z_j * (1.0f - done); /* Task 2.4: Huber-style negative-tail compression (relocated R6). * Applied before v_min/v_max clamp so the smooth compression is the * first transform; the hard clamp afterwards still enforces the diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index f33801333..17a4ac836 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2622,7 +2622,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (and the C-side mirrors in `state_layout.cuh`). Bump this constant in the /// SAME commit that adds the new range, and update `layout_fingerprint_seed` /// to register the new slot names. -pub(crate) const ISV_TOTAL_DIM: usize = 520; // was 510, SP20 added 10 slots [510..520) +pub(crate) const ISV_TOTAL_DIM: usize = 521; // was 520, SP21 Phase 3 added 1 slot [520..521) for Q_CORRECTION_INDEX /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -3748,7 +3748,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { SLOT_517_N_STEP=517;\ SLOT_518_AUX_CONF_THRESHOLD=518;\ SLOT_519_AUX_GATE_TEMP=519;\ - ISV_TOTAL_DIM=520;\ + SLOT_520_Q_CORRECTION=520;\ + ISV_TOTAL_DIM=521;\ SP19_PRODUCER_HARDCODED_HORIZON_BLEND=sp19_path_b;\ SP14_C_AUX_TRUNK_CONTROL_PLANE=sp14_c_phase_1;\ ALPHA_MACHINERY_DELETED=sp14_c_phase_1;\ @@ -30949,6 +30950,19 @@ impl GpuDqnTrainer { let curiosity_lambda_f32 = self.config.curiosity_q_penalty_lambda as f32; let spectral_lambda_f32 = self.config.spectral_decoupling_lambda as f32; + // SP21 T2.2 Phase 3 (2026-05-11): E1 q_correction Bellman offset. + // Read from `ISV[Q_CORRECTION_INDEX=520]` host-side (set by the + // enrichment phase after each val pass). Cold-start = 0.0 + // sentinel → kernel applies no shift. c51_loss_batched reads + // slot 520 from its existing `isv_signals` arg; this scalar is + // the mse-side mirror so blended `(1-α)·MSE + α·C51` sees + // identical Q-target shifts (`feedback_no_partial_refactor` + // — symmetric application). + let q_correction = { + use crate::cuda_pipeline::sp21_isv_slots::Q_CORRECTION_INDEX; + self.read_isv_signal_at(Q_CORRECTION_INDEX) + }; + unsafe { self.stream .launch_builder(kernel) @@ -31003,6 +31017,8 @@ impl GpuDqnTrainer { .arg(&self.ensemble_disagreement_weight) // ── Spectral decoupling (1) ── .arg(&spectral_lambda_f32) + // ── SP21 T2.2 Phase 3: E1 q_correction Bellman offset (1) ── + .arg(&q_correction) .launch(LaunchConfig { grid_dim: (b as u32, 1, 1), block_dim: (256, 1, 1), diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index b642e8605..07ab428cb 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -65,6 +65,7 @@ pub mod sp11_isv_slots; pub mod sp13_isv_slots; pub mod sp14_isv_slots; pub mod sp15_isv_slots; +pub mod sp21_isv_slots; pub mod lob_bar; pub use sp4_isv_slots::{ TARGET_Q_BOUND_INDEX, ATOM_POS_BOUND_BASE, WEIGHT_BOUND_BASE, diff --git a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu index 298961509..968e88dcd 100644 --- a/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/mse_loss_kernel.cu @@ -171,7 +171,20 @@ extern "C" __global__ void mse_loss_batched( float ensemble_disagreement_weight, /* Q_target -= weight * std (0.0=disabled) */ /* ── Spectral decoupling: L2 penalty on Q-value logit magnitudes ── */ - float spectral_decoupling_lambda /* Pezeshki et al. 2021 (0.0=disabled, 0.01=default) */ + float spectral_decoupling_lambda, /* Pezeshki et al. 2021 (0.0=disabled, 0.01=default) */ + + /* ── SP21 T2.2 Phase 3 (2026-05-11): E1 q_correction Bellman offset ── + * Read host-side from `ISV[Q_CORRECTION_INDEX=520]` (set by + * `enrichment::compute_q_correction` after each val pass). + * Sign convention: positive = "raise Q-targets" (policy + * underestimates); negative = "lower Q-targets" (policy + * overestimates). 0.0 = no-op (cold-start sentinel — first val + * pass replaces directly per `pearl_first_observation_bootstrap`). + * Applied as additive shift on `target_q` BEFORE the v_min/v_max + * clamp; mirrors `c51_loss_kernel::block_bellman_project_f`'s + * application to `t_z`. Symmetric application across both kernels + * keeps the blended `(1-α)·MSE + α·C51` loss consistent. */ + float q_correction ) { /* Shared memory is now float (4 bytes/elem) — doubled from BF16 version */ extern __shared__ float shmem_f[]; @@ -456,6 +469,14 @@ extern "C" __global__ void mse_loss_batched( /* ═══ STEP d: Bellman target + MSE ═══════════════════════════ */ float target_q = reward + gamma_eff * (1.0f - done) * target_eq; + /* SP21 T2.2 Phase 3 (2026-05-11): E1 q_correction additive offset. + * Applied AFTER the Bellman composition and BEFORE the ensemble + * disagreement subtraction + v_min/v_max clamp. Symmetric with + * `c51_loss_kernel::block_bellman_project_f`'s application to `t_z` + * so the blended `(1-α)·MSE + α·C51` loss sees identical + * Q-target shifts. 0.0 sentinel (cold start) is a no-op. */ + target_q += q_correction; + /* #27 Ensemble disagreement: reduce Q-target when ensemble is uncertain. * High std = novel/uncertain state = conservative Q-value. */ if (ensemble_disagreement_weight > 0.0f && ensemble_std != NULL) { diff --git a/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs new file mode 100644 index 000000000..9d11e845d --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs @@ -0,0 +1,87 @@ +//! SP21 ISV slot reservations. +//! +//! SP21 introduces enrichment-driven controllers that read per-trade +//! signals from `GpuBacktestEvaluator::read_per_trade_tape()` (Phase +//! 1.5+2) and feed them back into the next epoch's training step. +//! Each slot here is a single producer / single consumer pair within +//! the SP21 wireup chain. +//! +//! ## Slot allocation +//! +//! | Index | Name | Producer | Consumer | +//! |-------|------|----------|----------| +//! | 520 | `Q_CORRECTION_INDEX` | `enrichment::compute_q_correction` (host, post-val) | `mse_loss_kernel` + `c51_loss_kernel::block_bellman_project_f` (Bellman target offset) | +//! +//! ## Bus extension +//! +//! Each new slot bumps `ISV_TOTAL_DIM` (declared in `gpu_dqn_trainer.rs`) +//! and the layout fingerprint (used by `verify_isv_layout_fingerprint`). +//! Both must move atomically with the slot definition here per +//! `feedback_no_partial_refactor`. The `verify_sp21_slot_end_within_total_dim` +//! test below enforces the bus is large enough at compile time. + +/// SP21 T2.2 Phase 3 (2026-05-11) — additive Q-target offset from +/// E1 (`enrichment::compute_q_correction`). +/// +/// Producer: host-side after each validation backtest pass. +/// `result.q_correction = clamp(-mean(predicted_q − pnl), -10.0, 10.0)` +/// — predicted_q comes from the per-trade tape's Phase 1.5 +/// `predicted_q` field (entry_q at trade open). When the policy +/// systematically OVERESTIMATES Q at trade open relative to realized +/// pnl, `q_correction` is negative; underestimation produces a +/// positive correction. Cold start (no validation data yet) writes +/// 0.0 — sentinel matches `pearl_first_observation_bootstrap`. +/// +/// Consumer: +/// - `mse_loss_kernel::mse_loss_batched`: scalar kernel arg added +/// to `target_q = (reward + q_correction) + γ × E[Q'] × (1 − done)` +/// before the v_min/v_max clamp. +/// - `c51_loss_kernel::block_bellman_project_f`: reads the slot +/// from the existing `isv_signals` arg. Adds to `t_z = (reward +/// + reward_bias + q_correction) + γ × z_j × (1 − done)`. +/// +/// Sign convention: `q_correction` is the ADDITIVE correction to +/// the Bellman target, NOT the bias itself. E1 already negates the +/// bias inside `compute_q_correction`, so a positive value here +/// means "Q-targets should be raised". Cold-start bootstrap = 0.0 +/// → no-op until the first val pass writes a real correction. +pub const Q_CORRECTION_INDEX: usize = 520; + +/// First SP21 slot. Used by the bus-bounds compile-time check below +/// to assert SP21's slot range is reachable. +pub const SP21_SLOT_START: usize = 520; + +/// One past the last SP21 slot. +pub const SP21_SLOT_END: usize = 521; + +#[cfg(test)] +mod tests { + use super::*; + + /// Compile-time assertion that the ISV bus is wide enough for + /// SP21's slots. Bumping `ISV_TOTAL_DIM` is a layout-fingerprint + /// change — every consumer of `ISV_TOTAL_DIM` (pinned alloc, + /// fingerprint string, seed) must move atomically with the bus + /// extension per `feedback_no_partial_refactor`. + #[test] + fn verify_sp21_slot_end_within_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP21_SLOT_END <= ISV_TOTAL_DIM, + "SP21_SLOT_END={SP21_SLOT_END} exceeds ISV_TOTAL_DIM={ISV_TOTAL_DIM} — bus \ + too small for SP21 slots; bump `ISV_TOTAL_DIM` in `gpu_dqn_trainer.rs` and \ + update `layout_fingerprint_seed()` (per `feedback_no_partial_refactor`)." + ); + } + + #[test] + fn sp21_slot_indices_are_unique() { + // One slot today; placeholder for when SP21 grows. The check + // becomes load-bearing once a second slot is added. + let slots = [Q_CORRECTION_INDEX]; + let mut sorted = slots.to_vec(); + sorted.sort(); + sorted.dedup(); + assert_eq!(slots.len(), sorted.len(), "SP21 ISV slot indices must be unique"); + } +} diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 08bbcd9fe..9d621bd9d 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1579,6 +1579,22 @@ impl DQNTrainer { if let Some(ref mut fused) = self.fused_ctx { fused.set_var_ema(result.agreement_threshold); } + + // SP21 T2.2 Phase 3 (2026-05-11): apply E1 q_correction + // by writing to ISV[Q_CORRECTION_INDEX=520]. The mse_loss + // launcher reads this slot host-side and passes as a + // scalar kernel arg; the c51_loss kernel reads it + // device-side via its existing `isv_signals` ptr. + // Both kernels apply it as additive shift on the + // Bellman target. Cold-start: `result.q_correction = + // 0.0` (no trades yet) is the no-op sentinel. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp21_isv_slots::Q_CORRECTION_INDEX; + fused.trainer().write_isv_signal_at( + Q_CORRECTION_INDEX, + result.q_correction, + ); + } } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 24a96853d..e5814b66c 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -13709,3 +13709,164 @@ T2.2 multi-phase scope continues with Phases 3-7 + Phase 8: Each is its own atomic commit (per `feedback_no_partial_refactor` — phase-internal contract migrations are atomic; cross-phase wiring is sequential). + +## 2026-05-11 — SP21 T2.2 Phase 3: E1 q_correction → Bellman target offset + +### Scope (atomic single commit) + +Wires `EnrichmentResult::q_correction` (E1's `clamp(-mean(predicted_q +− pnl), ±10.0)` from the Phase 1.5+2 per-trade tape) into the +training loss kernels' Bellman target computation. Both `mse_loss_ +batched` and `c51_loss_batched`'s `block_bellman_project_f` apply +`q_correction` as additive shift on the Bellman target — the +blended `(1-α)·MSE + α·C51` loss sees identical Q-target shifts +per `feedback_no_partial_refactor` (symmetric application). + +### ISV slot allocation + +| Index | Name | Producer | Consumer | +|-------|------|----------|----------| +| 520 | `Q_CORRECTION_INDEX` | `enrichment::compute_q_correction` (host, post-val) | `mse_loss_batched` (scalar arg) + `c51_loss_kernel::block_bellman_project_f` (reads via existing `isv_signals` ptr) | + +`ISV_TOTAL_DIM` bumped 520 → 521 (per `feedback_no_partial_refactor` +all consumers — pinned alloc, fingerprint string, layout seed — +move atomically with the slot definition). New module +`cuda_pipeline::sp21_isv_slots` registered in `mod.rs`. + +### Sign convention + +`q_correction` is the ADDITIVE correction to the Bellman target, +NOT the bias itself. E1 already negates the bias inside +`compute_q_correction`: + + `q_correction = clamp(-mean(predicted_q − pnl), -10.0, 10.0)` + +So: + - `predicted_q > pnl` (overestimation) → bias > 0 → `q_correction < 0` → "lower Q-targets" + - `predicted_q < pnl` (underestimation) → bias < 0 → `q_correction > 0` → "raise Q-targets" + +Cold-start: the per-trade tape is empty until the first val pass +emits trade closes. `compute_q_correction` returns 0.0 on empty +trades; ISV slot stays at sentinel 0.0; both kernels apply no +shift. Matches `pearl_first_observation_bootstrap`. + +### ABI surgery (kernel-side) + +**mse_loss_kernel.cu** — adds one new `float q_correction` scalar +arg at the END of `mse_loss_batched`. Applied AFTER the Bellman +composition and BEFORE the ensemble disagreement subtraction + +v_min/v_max clamp: + +```c +float target_q = reward + gamma_eff * (1.0f - done) * target_eq; +target_q += q_correction; // SP21 Phase 3 +if (ensemble_disagreement_weight > 0.0f && ensemble_std != NULL) { + target_q -= ensemble_disagreement_weight * ensemble_std[sample_id]; +} +target_q = fminf(fmaxf(target_q, v_min), v_max); +``` + +**c51_loss_kernel.cu** — NO new kernel arg. `block_bellman_ +project_f` already takes `const float* __restrict__ isv_signals`; +reads slot 520 once per sample and adds to `t_z`: + +```c +float q_correction = (isv_signals != NULL) ? isv_signals[520] : 0.0f; +... +float t_z = (reward + reward_bias + q_correction) + gamma * z_j * (1.0f - done); +``` + +NULL-tolerant (existing `isv_signals != NULL` guard already in +place for `reward_bias` — `q_correction` reuses the same guard +shape). + +### Launcher changes + +**launch_mse_loss** in `gpu_dqn_trainer.rs:30882`: reads +`ISV[Q_CORRECTION_INDEX=520]` host-side via existing +`read_isv_signal_at`, passes as scalar arg at the end of the +kernel arg chain. + +**c51_loss launcher**: NO change (kernel already passes +`isv_signals` to `block_bellman_project_f`). + +### Producer wireup (host-side) + +`training_loop.rs` enrichment block (post-val pass): after the +existing `set_var_ema` for E5, write `result.q_correction` to +ISV[520]: + +```rust +if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp21_isv_slots::Q_CORRECTION_INDEX; + fused.trainer().write_isv_signal_at( + Q_CORRECTION_INDEX, + result.q_correction, + ); +} +``` + +The slot lives until the next val pass overwrites — every training +step between val passes reads the same correction. This is the +intended cadence (q_correction is an EMA-like signal across the val +window, not per-step). + +### Pearls + invariants honoured + +- `feedback_no_partial_refactor` — slot definition + bus extension + (`ISV_TOTAL_DIM` 520→521) + fingerprint update + producer wireup + + both consumer kernels migrate atomically. +- `feedback_no_stubs` — `q_correction` consumed for real (not + computed-then-discarded); the prior Phase 1.5+2 commit wired + E1's input (`predicted_q`), this commit wires E1's output. +- `pearl_first_observation_bootstrap` — sentinel 0.0 in ISV[520] + is no-op until first val pass writes a real correction; honest + cold-start. +- `pearl_controller_anchors_isv_driven` — q_correction lives in + ISV bus, not as a config constant or trainer field. +- `feedback_isv_for_adaptive_bounds` — adaptive Q-target shift + IS an adaptive bound on the Bellman target; lives in ISV. +- `feedback_no_atomicadd` — single-thread-per-sample writes to + `target_q` and `t_z`; no concurrent updates. +- `feedback_cpu_is_read_only` — `q_correction` is GPU-side except + for the ONE `write_isv_signal_at` call per epoch (host writes + via mapped-pinned slot, kernel reads from device pointer); + no GPU↔CPU compute roundtrip in the hot path. + +### Files changed + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp21_isv_slots.rs` | NEW | `Q_CORRECTION_INDEX = 520`; bus-bounds compile-time test | +| `crates/ml/src/cuda_pipeline/mod.rs` | +pub mod | Module registration | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | ISV bump + launcher | `ISV_TOTAL_DIM` 520→521; fingerprint adds `SLOT_520_Q_CORRECTION`; `launch_mse_loss` reads ISV[520] + passes scalar | +| `crates/ml/src/cuda_pipeline/mse_loss_kernel.cu` | ABI bump | `mse_loss_batched` gains `float q_correction` arg; applied to `target_q` | +| `crates/ml/src/cuda_pipeline/c51_loss_kernel.cu` | Body change | `block_bellman_project_f` reads `isv_signals[520]` and adds to `t_z`; NO ABI change | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Producer wireup | After enrichment block, writes `result.q_correction` to ISV[520] | +| `docs/dqn-wire-up-audit.md` | This entry | 2026-05-11 audit log | + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # 0 errors +cargo test -p ml --lib sp21_isv_slots --features cuda # 2/2 (bus bounds + slot uniqueness) +cargo test -p ml --test sp20_aggregate_inputs_test ... # 12/12 +cargo test -p ml --test sp20_phase1_4_wireup_test ... # 2/2 +cargo test -p ml --test sp20_emas_compute_test ... # 4/4 +cargo test -p ml --test sp20_controllers_compute_test ... # 7/7 +cargo test -p ml --test sp21_per_trade_predicted_q_test ... # 3/3 +``` + +Total: 30 tests, 0 failures. The Bellman-target shift's behavioral +gate is the upcoming smoke training run (`compute_q_correction` +produces non-zero output → ISV[520] non-zero → Q-targets shift). + +### After this commit lands + +T2.2 multi-phase scope continues with Phases 4-7 + Phase 8: +- Phase 4: wire E4 (per-branch LR scaling) via per-group Adam. +- Phase 5: wire E6 (winner indices) → PER priority bumps. +- Phase 6: wire E7 (`HindsightExperience`) → replay buffer. +- Phase 7: wire E8 (curriculum weights) → segment sampling. +- Phase 8: signal-drive remaining controller GAINS (E5 0.9/1.1 steps; + E2 2.0/0.5/-0.5 epsilon thresholds; etc.).