feat(sp21): T2.2 Phase 3 — wire E1 q_correction → Bellman target offset (atomic)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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),
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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) {
|
||||
|
||||
87
crates/ml/src/cuda_pipeline/sp21_isv_slots.rs
Normal file
87
crates/ml/src/cuda_pipeline/sp21_isv_slots.rs
Normal file
@@ -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");
|
||||
}
|
||||
}
|
||||
@@ -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,
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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.).
|
||||
|
||||
Reference in New Issue
Block a user