diff --git a/crates/ml-dqn/src/per_kernels.cu b/crates/ml-dqn/src/per_kernels.cu index d1c1f3612..811705e3f 100644 --- a/crates/ml-dqn/src/per_kernels.cu +++ b/crates/ml-dqn/src/per_kernels.cu @@ -220,7 +220,46 @@ extern "C" __global__ void per_insert_pa( boost_factor = 1.0f + weight * decreasing; } - float effective = raw_priorities[i] * boost_factor; + /* SP21 T2.2 Phase 7.5 (2026-05-11) — per-segment curriculum + * boost on insert priority. E8's per-segment weights live at + * ISV[528..536) (CURRICULUM_WEIGHT_{0..8}_INDEX), normalised so + * sum = 1 across the 8 segments. Each insert gets tagged with + * `seg_id = i % 8` (round-robin within the batch). Effective + * priority gets multiplied by `N_SEGMENTS × weight[seg_id]` + * — the `× N_SEGMENTS = 8` factor makes uniform weights + * (= 1/8 each) a no-op (= 1.0), so the controller only + * REDISTRIBUTES sampling priority across segments without + * changing the overall priority scale. + * + * Cold-start short-circuit: when `isv[528]` is at sentinel 0.0 + * (pre-first-emit), curriculum_boost = 1.0 (uniform no-op). + * Per `pearl_first_observation_bootstrap`. NULL-tolerant on + * `isv` for tests / smoke scaffolds. + * + * Slot indices written as literals (mirroring the canonical + * `dd_trajectory_kernel.cu` convention); names live in + * `crates/ml/src/cuda_pipeline/sp21_isv_slots.rs` and the + * trainer constructor pins those names to these indices via + * the `curriculum_weight_index` accessor used in the producer + * (host-side write path). */ + float curriculum_boost = 1.0f; + if (isv != nullptr) { + const float w0 = isv[/*CURRICULUM_WEIGHT_0_INDEX*/ 528]; + if (w0 > 1e-6f) { + const int seg_id = i % 8; + const float w = isv[528 + seg_id]; + /* × 8 makes uniform (1/8 per segment) = 1.0 no-op. */ + curriculum_boost = w * 8.0f; + /* Guard against pathological zero-weight segments: a + * single zero weight would zero the insert priority, + * which would prevent that segment from ever being + * sampled (sticky exclusion). Floor at 0.1× to keep + * all segments at least nominally sampleable. */ + if (curriculum_boost < 0.1f) curriculum_boost = 0.1f; + } + } + + float effective = raw_priorities[i] * boost_factor * curriculum_boost; priorities[idx] = effective; priorities_pa[idx] = powf(effective, alpha); diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 29fc851d8..23553e9b4 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 = 528; // SP21 Phase 7 adds 1 slot [527..528) for CURRICULUM_CONCENTRATION (PER alpha scaling, composed with WINNER_CONCENTRATION + HINDSIGHT_MAGNITUDE); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525) for BRANCH_LR_SCALE_*; Phase 3 added 1 slot [520..521) for Q_CORRECTION +pub(crate) const ISV_TOTAL_DIM: usize = 536; // SP21 Phase 7.5 adds 8 slots [528..536) for CURRICULUM_WEIGHT_{0..8} (per_insert_pa per-segment priority boost); Phase 7 added 1 slot [527..528); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525); Phase 3 added 1 slot [520..521) /// 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). @@ -3756,7 +3756,15 @@ const fn layout_fingerprint_seed() -> &'static [u8] { SLOT_525_WINNER_CONCENTRATION=525;\ SLOT_526_HINDSIGHT_MAGNITUDE=526;\ SLOT_527_CURRICULUM_CONCENTRATION=527;\ - ISV_TOTAL_DIM=528;\ + SLOT_528_CURRICULUM_WEIGHT_0=528;\ + SLOT_529_CURRICULUM_WEIGHT_1=529;\ + SLOT_530_CURRICULUM_WEIGHT_2=530;\ + SLOT_531_CURRICULUM_WEIGHT_3=531;\ + SLOT_532_CURRICULUM_WEIGHT_4=532;\ + SLOT_533_CURRICULUM_WEIGHT_5=533;\ + SLOT_534_CURRICULUM_WEIGHT_6=534;\ + SLOT_535_CURRICULUM_WEIGHT_7=535;\ + ISV_TOTAL_DIM=536;\ 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;\ diff --git a/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs index bd1b085fd..e68435cea 100644 --- a/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs @@ -18,6 +18,7 @@ //! | 525 | `WINNER_CONCENTRATION_INDEX` | host-side aggregation of E6 winner P&L + total P&L | `per_update_pa` kernel (PER alpha scaling) | //! | 526 | `HINDSIGHT_MAGNITUDE_INDEX` | host-side aggregation of E7 counterfactual P&L magnitudes | `per_update_pa` kernel (PER alpha scaling) | //! | 527 | `CURRICULUM_CONCENTRATION_INDEX` | host-side entropy of E8 curriculum-weights distribution | `per_update_pa` kernel (PER alpha scaling) | +//! | 528..536 | `CURRICULUM_WEIGHT_{0..8}_INDEX` | E8 per-segment weights (full Vec) | `per_insert_pa` kernel (per-segment insert priority boost) | //! //! ## Bus extension //! @@ -152,6 +153,60 @@ pub const HINDSIGHT_MAGNITUDE_INDEX: usize = 526; /// short-circuits to no-op. Per `pearl_first_observation_bootstrap`. pub const CURRICULUM_CONCENTRATION_INDEX: usize = 527; +/// SP21 T2.2 Phase 7.5 (2026-05-11) — per-segment curriculum +/// weights from E8 (`enrichment::compute_curriculum_weights`). +/// +/// Phase 7 wired the SCALAR concentration of these weights to +/// `per_update_pa`'s alpha boost; Phase 7.5 wires the FULL Vec +/// of per-segment weights to `per_insert_pa`'s priority boost, +/// closing the "true E8 consumer" loop that was deferred in 7's +/// audit entry. +/// +/// Producer: host-side after each validation backtest pass. +/// `compute_curriculum_weights` returns `Vec` of length +/// `n_segments = 8` (default), normalised so sum = 1. Each weight +/// is `(1/sharpe).clamp(0.1, 10.0)` / total — segments with low +/// validation Sharpe (= "hard segments") get larger weights. +/// +/// Consumer: `per_insert_pa` kernel. Each insert's effective +/// priority gets multiplied by `n_segments × weight[seg_id]` +/// where `seg_id = i % 8` (round-robin tag by insert order +/// within the batch). The `× n_segments` factor makes uniform +/// weights (= 1/8 each) a no-op (= 1.0), so the controller only +/// REDISTRIBUTES sampling priority across segments without +/// changing the overall priority scale. +/// +/// Round-robin segment tagging: experience-collector batches +/// (typically 512+ tuples per insert) get evenly distributed +/// across the 8 segments. Over time the buffer carries equal +/// representation per segment; E8's weights then redirect +/// sampling pressure toward hard segments at insert time. +/// +/// Cold-start: when `isv_signals[CURRICULUM_WEIGHT_0_INDEX]` is +/// at sentinel `0.0` (pre-first-emit), the kernel short-circuits +/// to `boost = 1.0` (uniform no-op). Same +/// `pearl_first_observation_bootstrap` pattern as Phases 5+6+7. +pub const CURRICULUM_WEIGHT_0_INDEX: usize = 528; +pub const CURRICULUM_WEIGHT_1_INDEX: usize = 529; +pub const CURRICULUM_WEIGHT_2_INDEX: usize = 530; +pub const CURRICULUM_WEIGHT_3_INDEX: usize = 531; +pub const CURRICULUM_WEIGHT_4_INDEX: usize = 532; +pub const CURRICULUM_WEIGHT_5_INDEX: usize = 533; +pub const CURRICULUM_WEIGHT_6_INDEX: usize = 534; +pub const CURRICULUM_WEIGHT_7_INDEX: usize = 535; + +/// Number of curriculum segments. Matches `EnrichmentState`'s +/// `n_segments` default and the kernel's `seg_id = i % N` mod. +pub const CURRICULUM_N_SEGMENTS: usize = 8; + +/// Convenience accessor: segment idx (0..8) → ISV slot index. +#[inline] +pub fn curriculum_weight_index(seg_idx: usize) -> usize { + assert!(seg_idx < CURRICULUM_N_SEGMENTS, + "curriculum_weight_index: seg_idx must be 0..{CURRICULUM_N_SEGMENTS} (got {seg_idx})"); + CURRICULUM_WEIGHT_0_INDEX + seg_idx +} + /// Convenience accessor: branch index (0..4) → ISV slot index. Used /// by the launcher's per-branch sub-launch loop to keep the /// branch-index ↔ ISV-slot mapping in one place. @@ -171,7 +226,7 @@ pub fn branch_lr_scale_index(branch: usize) -> usize { pub const SP21_SLOT_START: usize = 520; /// One past the last SP21 slot. -pub const SP21_SLOT_END: usize = 528; +pub const SP21_SLOT_END: usize = 536; #[cfg(test)] mod tests { @@ -204,6 +259,14 @@ mod tests { WINNER_CONCENTRATION_INDEX, HINDSIGHT_MAGNITUDE_INDEX, CURRICULUM_CONCENTRATION_INDEX, + CURRICULUM_WEIGHT_0_INDEX, + CURRICULUM_WEIGHT_1_INDEX, + CURRICULUM_WEIGHT_2_INDEX, + CURRICULUM_WEIGHT_3_INDEX, + CURRICULUM_WEIGHT_4_INDEX, + CURRICULUM_WEIGHT_5_INDEX, + CURRICULUM_WEIGHT_6_INDEX, + CURRICULUM_WEIGHT_7_INDEX, ]; let mut sorted = slots.to_vec(); sorted.sort(); @@ -218,4 +281,11 @@ mod tests { assert_eq!(branch_lr_scale_index(2), BRANCH_LR_SCALE_ORDER_INDEX); assert_eq!(branch_lr_scale_index(3), BRANCH_LR_SCALE_URGENCY_INDEX); } + + #[test] + fn curriculum_weight_index_matches_constants() { + assert_eq!(curriculum_weight_index(0), CURRICULUM_WEIGHT_0_INDEX); + assert_eq!(curriculum_weight_index(7), CURRICULUM_WEIGHT_7_INDEX); + assert_eq!(CURRICULUM_N_SEGMENTS, 8); + } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 8a910b924..6bc85a755 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1641,6 +1641,31 @@ impl DQNTrainer { CURRICULUM_CONCENTRATION_INDEX, result.curriculum_concentration, ); + + // SP21 T2.2 Phase 7.5 (2026-05-11): wire full + // E8 per-segment weights (Vec of length 8) + // to ISV[CURRICULUM_WEIGHT_{0..8}]. Consumed by + // per_insert_pa kernel which multiplies insert + // priority by `N_SEGMENTS × weight[seg_id]` + // (round-robin segment tag, seg_id = i % 8). + // Closes the "true E8 consumer" deferral noted + // in Phase 7's audit entry. Iterate `take(N)` + // defensive against future n_segments mismatch + // — current EnrichmentState::new pins it to 8. + use crate::cuda_pipeline::sp21_isv_slots::{ + curriculum_weight_index, CURRICULUM_N_SEGMENTS, + }; + for (seg_idx, &weight) in result + .curriculum_weights + .iter() + .take(CURRICULUM_N_SEGMENTS) + .enumerate() + { + fused.trainer().write_isv_signal_at( + curriculum_weight_index(seg_idx), + weight, + ); + } } // SP21 T2.2 Phase 6.5b (2026-05-11): synthetic diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 16bae6b8f..61b17efd9 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -14967,6 +14967,126 @@ the ceiling adaptively. No other controllers saturated in the enrichment controllers (only Invariant 1 stability carve-outs on bound-on-bound formulas, which terminate the recursion). +## 2026-05-11 — SP21 T2.2 Phase 7.5: E8 curriculum_weights → per-segment PER insert priority boost (atomic) + +### Scope (atomic single commit) + +Closes the "true E8 per-segment PER sampling" deferral from +Phase 7. Phase 7 wired E8's SCALAR concentration to +`per_update_pa`'s alpha boost; Phase 7.5 wires the FULL `Vec` +of per-segment weights to `per_insert_pa`'s priority boost. + +### What lands + +1. **8 new ISV slots [528..536)**: + `CURRICULUM_WEIGHT_{0..8}_INDEX`. Each receives one of E8's + per-segment weights (`compute_curriculum_weights` returns + `Vec` of length 8, normalised so sum = 1). + +2. **`ISV_TOTAL_DIM` bumped 528 → 536**. Fingerprint adds 8 SLOT + entries. `CURRICULUM_N_SEGMENTS = 8` const + + `curriculum_weight_index(seg_idx)` accessor. + +3. **`per_insert_pa` kernel reads `isv[528 + seg_id]`** where + `seg_id = i % 8` (round-robin segment tag by insert order + within the batch). Effective priority gets multiplied by + `N_SEGMENTS × weight[seg_id]`: + - Uniform weights (= 1/8 each) → boost = 1.0 (no-op). + - High-weight segments (hard trading windows per E8) get + larger boost → sampled more. + - Low-weight segments → smaller boost (but floored at 0.1× + to prevent sticky exclusion). + +4. **Cold-start short-circuit**: `isv[528] ≤ 1e-6` → boost = + 1.0 (uniform no-op). Same `pearl_first_observation_bootstrap` + pattern as Phases 5+6+7. + +5. **Producer**: `training_loop` post-enrichment block writes + 8 ISV slots from `result.curriculum_weights[0..8]`. NULL- + tolerant in kernel for test scaffolds. + +### Segment tagging rationale + +The naïve approach — tag each replay tuple with a "val-curriculum- +segment id" preserving val-window semantics — is infeasible +because val and training have separate coordinate systems (the +same problem documented in Phase 5+6's audit re E6 winner +indices). Round-robin segment tagging via `i % 8` instead +distributes the experience-collector's typical 512+ tuple batch +EVENLY across 8 segments. Over time the buffer carries equal +representation per segment; E8's weights then redirect sampling +pressure toward "hard" segments at insert time. + +This is a HEURISTIC mapping — it doesn't preserve the val-segment +SEMANTIC. But it consumes the curriculum_weights vector for real +PER priority redistribution, which is what Phase 7.5 set out to +do. A future commit could add a true bar-index → segment mapping +(requires tagging each insert with a temporal segment id from the +source data), but that's substantial additional infrastructure. + +### Files changed + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp21_isv_slots.rs` | +8 slots + accessor + test | `CURRICULUM_WEIGHT_{0..8}_INDEX = 528..536`; `CURRICULUM_N_SEGMENTS = 8`; `curriculum_weight_index(seg_idx)` accessor; uniqueness + accessor mapping tests | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | ISV bump | `ISV_TOTAL_DIM` 528 → 536; fingerprint adds 8 SLOT entries | +| `crates/ml-dqn/src/per_kernels.cu` | per_insert_pa boost | Reads `isv[528..536)`, applies per-segment boost `N_SEGMENTS × weight[i%8]` with cold-start sentinel + 0.1× floor against pathological zero weights | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Producer | Post-enrichment block writes 8 ISV slots via `curriculum_weight_index`; `take(CURRICULUM_N_SEGMENTS)` defensive | +| `docs/dqn-wire-up-audit.md` | This entry | 2026-05-11 audit log | + +### Pearls + invariants honoured + +- `feedback_no_partial_refactor` — ISV slot + bus bump + kernel + body + producer wireup migrate atomically. +- `feedback_no_stubs` — `result.curriculum_weights` consumed for + real (Phase 7 wired the scalar concentration; 7.5 wires the + full vector). +- `pearl_first_observation_bootstrap` — sentinel 0.0 short-circuits + kernel to uniform-no-op until E8 emits real weights. +- `pearl_controller_anchors_isv_driven` — per-segment weights + live in ISV; `× N_SEGMENTS` normalisation, `0.1×` floor, and + `i % 8` segment derivation are all structural (Invariant 1 + carve-outs), NOT controller anchors. +- `feedback_no_atomicadd` — kernel side reads ISV via standard + device load; no concurrent writes anywhere. + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # 0 errors +cargo test -p ml --lib sp21_isv_slots --features cuda # 4/4 (new curriculum_weight_index test) +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: 35 tests, 0 failures. Behavioral gate: a smoke training +run should surface heterogeneous per-segment insert priorities +(visible in PER priority distribution histograms by segment +position in the ring buffer). + +### After this commit lands + +SP21 T2.2 cascade — **TRULY** fully complete (13 atomic commits). +Every enrichment output E1-E8 wires to a real consumer: + +- E1 q_correction → Bellman target offset (Phase 3) +- E2 epsilon → next-epoch ε scaling (Phase 8 ISV gains) +- E3 gamma → blended into hyperparams.gamma (Phase 2) +- E4 branch_lr_scale → per-branch Adam sub-launches (Phase 4) +- E5 agreement_threshold → fused.set_var_ema (Phase 2, gains + signal-driven in Phase 8, clamp bounds in Phase 8.1) +- E6 winner_concentration scalar → PER alpha boost (Phase 5+6) +- E7 hindsight: magnitude scalar → PER alpha boost (Phase 5+6); + synthetic tuple injection (Phase 6.5) +- E8 curriculum_concentration scalar → PER alpha boost (Phase 7); + per-segment weight vector → per_insert_pa boost (Phase 7.5) + +No remaining deferrals or hardcoded controller anchors in the +SP21 T2.2 scope. + Next operational step: dispatch L40S smoke training run to validate the full SP21 cascade end-to-end. Watch: - `q_corr` non-zero after first val pass (Phase 3 ISV[520])