diff --git a/crates/ml-dqn/src/gpu_replay_buffer.rs b/crates/ml-dqn/src/gpu_replay_buffer.rs index a497db2cc..08e2badfe 100644 --- a/crates/ml-dqn/src/gpu_replay_buffer.rs +++ b/crates/ml-dqn/src/gpu_replay_buffer.rs @@ -1302,7 +1302,16 @@ impl GpuReplayBuffer { .map_err(|e| MLError::ModelError(format!("pow_alpha_diverse_f32: {e}")))?; } } else { - // Standard PER (unchanged): reads f32 td_errors directly, writes priorities + priorities_pa + // Standard PER: reads f32 td_errors directly, writes priorities + priorities_pa. + // SP21 Phase 5+6 (2026-05-11): pass `isv_signals_dev_ptr` so the + // kernel can read ISV[525] (E6 winner concentration) + ISV[526] + // (E7 hindsight magnitude) device-side and lift `alpha_eff` per + // sample. NULL-tolerant: when the trainer hasn't wired the ISV + // bus yet (smoke tests / oracle scaffolds), `isv_ptr = 0` and + // the kernel short-circuits to `alpha_eff = alpha` (no-op). + // Same null-degradation contract as the established + // `per_insert_pa` ISV[440] recovery_oversample wireup. + let isv_ptr = self.isv_signals_dev_ptr; unsafe { stream.launch_builder(&self.kernels.per_update_pa) .arg(&self.priorities_pa) @@ -1311,6 +1320,7 @@ impl GpuReplayBuffer { .arg(indices) .arg(td_errors) .arg(&al).arg(&ep).arg(&cap_i).arg(&bsi) + .arg(&isv_ptr) .launch(lcfg(bs)) .map_err(|e| MLError::ModelError(format!("per_update_pa: {e}")))?; } diff --git a/crates/ml-dqn/src/per_kernels.cu b/crates/ml-dqn/src/per_kernels.cu index 689d20965..f4daf6013 100644 --- a/crates/ml-dqn/src/per_kernels.cu +++ b/crates/ml-dqn/src/per_kernels.cu @@ -38,7 +38,23 @@ extern "C" __global__ void per_update_pa( const unsigned int* __restrict__ indices, const float* __restrict__ td_errors_bf16, float alpha, float epsilon, - int capacity, int batch_size) + int capacity, int batch_size, + /* SP21 T2.2 Phase 5+6 (2026-05-11) — combined E6/E7 signal-driven + * PER alpha scaling. ISV[525] = winner P&L concentration from + * E6 (top-decile mean P&L / all-trade mean P&L); ISV[526] = + * hindsight counterfactual-P&L magnitude from E7. Both + * normalised host-side; this kernel composes them + * multiplicatively into `alpha_boost` and applies as + * `alpha_eff = alpha + boost_delta` where boost_delta is + * scaled to [0, 0.4] (so alpha_eff stays in + * [alpha_baseline, alpha_baseline + 0.4]). + * + * Cold-start sentinel: when BOTH ISV[525] AND ISV[526] are at + * sentinel 0.0 (pre-first-val-pass), kernel short-circuits + * `alpha_eff = alpha` — no-op. Per + * `pearl_first_observation_bootstrap`. NULL-tolerant on + * `isv_signals` for tests / smoke scaffolds. */ + const float* __restrict__ isv_signals) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= batch_size) return; @@ -60,10 +76,33 @@ extern "C" __global__ void per_update_pa( int ival = __float_as_int(new_prio); atomicMax((int*)batch_max, ival); - /* Sampling weight = priority^alpha. Sum-tree (per_prefix_scan + + /* SP21 Phase 5+6: compose alpha_boost from E6 (winner concentration) + * and E7 (hindsight magnitude). Both signals at sentinel 0.0 → + * boost_delta = 0 → alpha_eff = alpha (cold-start no-op). Otherwise + * the additive delta is bounded to keep alpha_eff in a sane range: + * - winner_term: (concentration - 1.0) clipped to [0, 2.0], scaled + * by 0.1 → [0, 0.2]. Concentration > 1.0 means top-decile drives + * more profit per trade than average; we want PER to focus more + * on hard cases (alpha up). + * - hindsight_term: magnitude clipped to [0, 0.02] (typical + * counterfactual P&L scale ≈ 1-2%), scaled by 10.0 → [0, 0.2]. + * Higher magnitude = more missed opportunities = need more + * PER focus. + * - Total delta clipped at 0.4 so alpha_eff ≤ alpha + 0.4. */ + float winner_conc = (isv_signals != NULL) ? isv_signals[525] : 0.0f; + float hindsight_mag = (isv_signals != NULL) ? isv_signals[526] : 0.0f; + float boost_delta = 0.0f; + if (winner_conc > 1e-6f || hindsight_mag > 1e-6f) { + float winner_term = fmaxf(0.0f, fminf(2.0f, winner_conc - 1.0f)) * 0.1f; + float hindsight_term = fminf(0.02f, hindsight_mag) * 10.0f; + boost_delta = fminf(0.4f, winner_term + hindsight_term); + } + float alpha_eff = alpha + boost_delta; + + /* Sampling weight = priority^alpha_eff. Sum-tree (per_prefix_scan + * per_sample) walks this buffer; IS-weight (is_weights_f32) reads * the same column for normalisation. */ - float pa = powf(new_prio, alpha); + float pa = powf(new_prio, alpha_eff); priorities_pa[idx] = pa; } diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 12fa6a182..938d3647c 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 = 525; // SP21 Phase 4 adds 4 slots [521..525) for BRANCH_LR_SCALE_{DIR,MAG,ORDER,URGENCY}; Phase 3 added 1 slot [520..521) for Q_CORRECTION +pub(crate) const ISV_TOTAL_DIM: usize = 527; // SP21 Phase 5+6 adds 2 slots [525..527) for WINNER_CONCENTRATION + HINDSIGHT_MAGNITUDE (PER alpha scaling); Phase 4 added 4 slots [521..525) for BRANCH_LR_SCALE_*; Phase 3 added 1 slot [520..521) for Q_CORRECTION /// 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). @@ -3753,7 +3753,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] { SLOT_522_BRANCH_LR_SCALE_MAG=522;\ SLOT_523_BRANCH_LR_SCALE_ORDER=523;\ SLOT_524_BRANCH_LR_SCALE_URGENCY=524;\ - ISV_TOTAL_DIM=525;\ + SLOT_525_WINNER_CONCENTRATION=525;\ + SLOT_526_HINDSIGHT_MAGNITUDE=526;\ + ISV_TOTAL_DIM=527;\ 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 19e5c83e5..75897351b 100644 --- a/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp21_isv_slots.rs @@ -15,6 +15,8 @@ //! | 522 | `BRANCH_LR_SCALE_MAG_INDEX` | `enrichment::compute_branch_lr_scale[1]` | `dqn_adam_update_kernel` per-branch sub-launch (mag tensors 21-24) | //! | 523 | `BRANCH_LR_SCALE_ORDER_INDEX` | `enrichment::compute_branch_lr_scale[2]` | `dqn_adam_update_kernel` per-branch sub-launch (order tensors 25-28) | //! | 524 | `BRANCH_LR_SCALE_URGENCY_INDEX` | `enrichment::compute_branch_lr_scale[3]` | `dqn_adam_update_kernel` per-branch sub-launch (urgency tensors 29-32) | +//! | 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) | //! //! ## Bus extension //! @@ -82,6 +84,49 @@ pub const BRANCH_LR_SCALE_MAG_INDEX: usize = 522; pub const BRANCH_LR_SCALE_ORDER_INDEX: usize = 523; pub const BRANCH_LR_SCALE_URGENCY_INDEX: usize = 524; +/// SP21 T2.2 Phase 5+6 combined (2026-05-11) — winner P&L +/// concentration from E6 (`enrichment::compute_winner_indices`). +/// +/// Producer: host-side after each validation backtest pass. The raw +/// `result.winner_indices: Vec` returns top-decile bar +/// indices, but the literal indices are val-space (no mapping to +/// the training replay buffer). The MEANINGFUL signal in E6 is the +/// CONCENTRATION of profitability — top-decile mean P&L divided by +/// all-trade mean P&L. Concentration > 1.0 means winners +/// disproportionately drive total profit (a few big trades carry +/// the strategy); concentration ≈ 1.0 means profit is evenly +/// distributed (more uniform policy quality). +/// +/// Consumer: `per_update_pa` kernel (PER priority alpha scaling). +/// Higher concentration → higher alpha multiplier → sharper PER +/// sampling on high-TD-error transitions (the "hard cases" that +/// the policy is missing). Lower concentration → uniform sampling. +/// +/// Cold-start: 0.0 sentinel before first val pass. Kernel checks +/// `(winner_conc < 1e-6 && hindsight_mag < 1e-6)` and short-circuits +/// to alpha unchanged (no-op). Per `pearl_first_observation_bootstrap`. +pub const WINNER_CONCENTRATION_INDEX: usize = 525; + +/// SP21 T2.2 Phase 5+6 combined (2026-05-11) — hindsight +/// counterfactual-P&L magnitude from E7 +/// (`enrichment::compute_hindsight_labels`). +/// +/// Producer: host-side after each validation backtest pass. Mean +/// of `|counterfactual_pnl|` across the `result.hindsight: Vec< +/// HindsightExperience>` items. Higher magnitude means the policy +/// missed larger opportunities relative to optimal counterfactuals +/// — more reason to focus PER on under-trained samples. +/// +/// Consumer: `per_update_pa` kernel (PER priority alpha scaling), +/// composed multiplicatively with `WINNER_CONCENTRATION_INDEX` to +/// form `alpha_boost`. The combined boost lifts alpha when EITHER +/// E6 (winner concentration) OR E7 (missed opportunities) signal +/// that the policy needs more focus on hard cases. +/// +/// Cold-start: 0.0 sentinel; kernel short-circuits to alpha +/// unchanged when both signals are at sentinel. +pub const HINDSIGHT_MAGNITUDE_INDEX: usize = 526; + /// 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. @@ -101,7 +146,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 = 525; +pub const SP21_SLOT_END: usize = 527; #[cfg(test)] mod tests { @@ -131,6 +176,8 @@ mod tests { BRANCH_LR_SCALE_MAG_INDEX, BRANCH_LR_SCALE_ORDER_INDEX, BRANCH_LR_SCALE_URGENCY_INDEX, + WINNER_CONCENTRATION_INDEX, + HINDSIGHT_MAGNITUDE_INDEX, ]; let mut sorted = slots.to_vec(); sorted.sort(); diff --git a/crates/ml/src/trainers/dqn/trainer/enrichment.rs b/crates/ml/src/trainers/dqn/trainer/enrichment.rs index 1e7321147..cbc71a3a6 100644 --- a/crates/ml/src/trainers/dqn/trainer/enrichment.rs +++ b/crates/ml/src/trainers/dqn/trainer/enrichment.rs @@ -54,8 +54,21 @@ pub(crate) struct EnrichmentResult { pub agreement_threshold: f32, /// E6: bar indices of top-10% trades (priority replay). pub winner_indices: Vec, + /// SP21 T2.2 Phase 5+6 (2026-05-11) — winner P&L concentration: + /// `top_decile_mean_pnl / max(all_trade_mean_pnl, EPS)`. > 1.0 when + /// the top decile drives more profit per trade than average. Cold + /// start: trades.is_empty() returns 0.0 sentinel. Consumed by + /// `per_update_pa` via ISV[WINNER_CONCENTRATION_INDEX=525] for + /// per-sample alpha boost. + pub winner_concentration: f32, /// E7: hindsight labels for worst losers. pub hindsight: Vec, + /// SP21 T2.2 Phase 5+6 (2026-05-11) — mean |counterfactual_pnl| + /// across hindsight experiences. Higher magnitude = larger missed + /// opportunities. Cold start: hindsight.is_empty() returns 0.0 + /// sentinel. Consumed by `per_update_pa` via + /// ISV[HINDSIGHT_MAGNITUDE_INDEX=526] for per-sample alpha boost. + pub hindsight_magnitude: f32, /// E8: curriculum weights (one per segment). pub curriculum_weights: Vec, } @@ -242,6 +255,53 @@ fn compute_agreement_threshold( new.clamp(0.01, 10.0) } +/// SP21 T2.2 Phase 5+6 (2026-05-11) — winner P&L concentration. +/// +/// `top_decile_mean_pnl / max(all_trade_mean_pnl, EPS)`. Captures +/// how concentrated the policy's profitability is: > 1.0 means top- +/// decile trades carry disproportionate profit (a few big wins); +/// ≈ 1.0 means uniform distribution; < 1.0 indicates the top decile +/// underperforms the average (rare — typically signals top-decile +/// includes near-zero-pnl trades while the bulk of trades skew +/// negative). +/// +/// Used by `per_update_pa` (PER alpha scaling) — see +/// `WINNER_CONCENTRATION_INDEX` in `cuda_pipeline::sp21_isv_slots`. +/// Negative or zero `all_mean_pnl` (i.e., losing strategy) +/// short-circuits to 0.0 sentinel — the signal is ill-defined when +/// the policy isn't profitable. +fn compute_winner_concentration(trades: &[EvalTrade]) -> f32 { + if trades.is_empty() { + return 0.0; + } + let mut sorted: Vec = trades.iter().map(|t| t.pnl).collect(); + sorted.sort_by(|a, b| b.partial_cmp(a).unwrap_or(std::cmp::Ordering::Equal)); + let top_n = (sorted.len() / 10).max(1); + let top_mean: f32 = sorted[..top_n].iter().sum::() / top_n as f32; + let all_mean: f32 = sorted.iter().sum::() / sorted.len() as f32; + if all_mean <= 1e-6 { + // Losing or break-even strategy — concentration is ill-defined. + // Return 0.0 sentinel so the kernel short-circuits to no boost. + return 0.0; + } + top_mean / all_mean +} + +/// SP21 T2.2 Phase 5+6 (2026-05-11) — hindsight magnitude. +/// +/// Mean of `|counterfactual_pnl|` across all hindsight experiences. +/// Higher magnitude = larger missed opportunities (the policy +/// produced losing trades whose optimal counterfactual would have +/// been substantially profitable). Used by `per_update_pa` — +/// see `HINDSIGHT_MAGNITUDE_INDEX` in `cuda_pipeline::sp21_isv_slots`. +fn compute_hindsight_magnitude(hindsight: &[HindsightExperience]) -> f32 { + if hindsight.is_empty() { + return 0.0; + } + hindsight.iter().map(|h| h.counterfactual_pnl.abs()).sum::() + / hindsight.len() as f32 +} + /// E6: Winner distillation — bar indices of top-10% trades by P&L. fn compute_winner_indices(trades: &[EvalTrade]) -> Vec { if trades.is_empty() { @@ -357,17 +417,21 @@ pub(crate) fn run_enrichments( // E6: Winner distillation let winner_indices = compute_winner_indices(eval_trades); + // SP21 T2.2 Phase 5+6 (2026-05-11) — scalar signals for PER alpha + // scaling, derived from the same per-trade tape that feeds E6/E7. + let winner_concentration = compute_winner_concentration(eval_trades); // E7: Hindsight labels let hindsight = compute_hindsight_labels(eval_trades); + let hindsight_magnitude = compute_hindsight_magnitude(&hindsight); // E8: Curriculum weights let curriculum_weights = compute_curriculum_weights(eval_trades, state.n_segments); info!( "Enrichment cycle {} complete: q_corr={:.4}, eps={:.4}, gamma={:?}, \ - branch_lr={:?}, agree_thr={:.4}, winners={}, hindsight={}, \ - curriculum_segs={}", + branch_lr={:?}, agree_thr={:.4}, winners={}, win_conc={:.4}, \ + hindsight={}, hindsight_mag={:.4}, curriculum_segs={}", state.cycle_count, q_correction, epsilon, @@ -375,7 +439,9 @@ pub(crate) fn run_enrichments( branch_lr_scale, agreement_threshold, winner_indices.len(), + winner_concentration, hindsight.len(), + hindsight_magnitude, curriculum_weights.len(), ); @@ -386,7 +452,9 @@ pub(crate) fn run_enrichments( branch_lr_scale, agreement_threshold, winner_indices, + winner_concentration, hindsight, + hindsight_magnitude, curriculum_weights, } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index e6adf3dd5..0ead35915 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1600,6 +1600,7 @@ impl DQNTrainer { if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp21_isv_slots::{ Q_CORRECTION_INDEX, branch_lr_scale_index, + WINNER_CONCENTRATION_INDEX, HINDSIGHT_MAGNITUDE_INDEX, }; fused.trainer().write_isv_signal_at( Q_CORRECTION_INDEX, @@ -1613,6 +1614,23 @@ impl DQNTrainer { scale, ); } + // SP21 T2.2 Phase 5+6 (2026-05-11): wire E6 winner + // concentration + E7 hindsight magnitude into PER + // alpha scaling. Both are scalar summaries of the + // per-trade tape (computed once per val pass). + // `per_update_pa` reads slots 525/526 device-side + // and lifts alpha multiplicatively per-batch + // (cold-start sentinel 0.0 → kernel short-circuits + // to no-op). See pearl_first_observation_bootstrap + // discipline. + fused.trainer().write_isv_signal_at( + WINNER_CONCENTRATION_INDEX, + result.winner_concentration, + ); + fused.trainer().write_isv_signal_at( + HINDSIGHT_MAGNITUDE_INDEX, + result.hindsight_magnitude, + ); } } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index a3b5e986b..c89edaede 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -14155,3 +14155,219 @@ will surface per-branch engagement-rate-deficit EMAs in HEALTH_DIAG (emit threshold `> 0.005`) for the first time — Phase 4.5 re-instates the per-group Pearl C signal that Phase 4 temporarily disabled. + +## 2026-05-11 — SP21 T2.2 Phase 5+6 (combined): E6 winner concentration + E7 hindsight magnitude → PER alpha scaling + +### Scope (atomic single commit, Phases 5 + 6 fused) + +Wires E6 (`enrichment::compute_winner_indices`) + E7 +(`enrichment::compute_hindsight_labels`) outputs into the PER +priority update kernel via signal-driven ISV slots. The original +plan listed Phase 5 and Phase 6 as separate atomic commits, but +the design audit (briefed and accepted by jgrusewski 2026-05-11) +identified that: + +1. E6's literal `Vec` of val bar indices CANNOT directly bump + training-PER priorities — val and training have separate + coordinate systems (no bar-index → buffer-slot mapping). +2. E7's hindsight injection (counterfactual experience tuples + inserted into PER) requires retaining val state vectors at each + bar_index — significant new infrastructure (n_windows × max_len + × state_dim per-eval scratch buffer). +3. Both phases' MEANINGFUL signal collapses to **scalar + aggregations** of the per-trade tape: E6's winner concentration + (top-decile mean P&L / all-trade mean P&L) and E7's hindsight + magnitude (mean |counterfactual_pnl|). + +Per `feedback_no_deferrals_for_complementary_fixes`, when two fixes +attack distinct mechanisms with non-overlapping refactor scopes, +combine into one plan rather than sequencing. The two scalar +signals compose multiplicatively into one `alpha_boost` consumed +by the same kernel (`per_update_pa`) — exactly the canonical +"complementary fixes" pattern. + +### Plan revision from on-paper design + +**Original plan (2026-05-11)**: +- Phase 5: "Wire E6 (winner indices) → PER priority bumps." +- Phase 6: "Wire E7 (`HindsightExperience`) → replay buffer + synthetic experience injection." + +**Resolution (this commit)**: +- Phase 5: signal-driven via E6 winner concentration scalar. +- Phase 6: signal-driven via E7 hindsight magnitude scalar. +- Both feed `per_update_pa` alpha boost via ISV[525] + ISV[526]. +- Synthetic experience injection (true E7 hindsight replay) is + documented as a future Phase 6.5 follow-up requiring val state + retention infrastructure. + +### Signal semantics + +**Winner concentration** (`compute_winner_concentration`): +- `top_decile_mean_pnl / max(all_trade_mean_pnl, EPS)`. +- > 1.0 means top-decile drives more profit per trade than average + — "a few big wins carry the strategy". +- ≈ 1.0 means uniform distribution. +- Negative or ≤ 0 `all_mean_pnl` (losing/break-even strategy) → + short-circuits to 0.0 sentinel (signal ill-defined). +- Cold start (empty trades): 0.0 sentinel. + +**Hindsight magnitude** (`compute_hindsight_magnitude`): +- `mean(|counterfactual_pnl|)` across hindsight experiences. +- Higher = larger missed opportunities (the policy made losing + trades whose optimal counterfactual would have been substantially + profitable). +- Cold start (empty hindsight): 0.0 sentinel. + +### ISV slot allocation + +| Index | Name | Producer | Consumer | +|-------|------|----------|----------| +| 525 | `WINNER_CONCENTRATION_INDEX` | `enrichment::compute_winner_concentration` (host, post-val) | `per_update_pa` (PER alpha boost) | +| 526 | `HINDSIGHT_MAGNITUDE_INDEX` | `enrichment::compute_hindsight_magnitude` (host, post-val) | `per_update_pa` (PER alpha boost) | + +`ISV_TOTAL_DIM` 525 → 527. Layout fingerprint adds +`SLOT_525_WINNER_CONCENTRATION` + `SLOT_526_HINDSIGHT_MAGNITUDE`. + +### ABI surgery (kernel-side) + +**`per_update_pa` in `crates/ml-dqn/src/per_kernels.cu`**: ONE new +arg `const float* __restrict__ isv_signals` at the END. NULL-tolerant +(NULL → kernel applies `alpha_eff = alpha`, identical to pre-Phase-5 +behaviour). Inside kernel: + +```c +float winner_conc = (isv_signals != NULL) ? isv_signals[525] : 0.0f; +float hindsight_mag = (isv_signals != NULL) ? isv_signals[526] : 0.0f; +float boost_delta = 0.0f; +if (winner_conc > 1e-6f || hindsight_mag > 1e-6f) { + float winner_term = fmaxf(0.0f, fminf(2.0f, winner_conc - 1.0f)) * 0.1f; + float hindsight_term = fminf(0.02f, hindsight_mag) * 10.0f; + boost_delta = fminf(0.4f, winner_term + hindsight_term); +} +float alpha_eff = alpha + boost_delta; +priorities_pa[idx] = powf(new_prio, alpha_eff); +``` + +`boost_delta` is bounded `[0, 0.4]` so `alpha_eff ≤ alpha + 0.4` +(prevents runaway sharpening of PER sampling). Cold-start +short-circuit guarantees no behavioural change before E6/E7 emit +their first values. + +### Launcher change + +`crates/ml-dqn/src/gpu_replay_buffer.rs::update_priorities_gpu`: +the standard PER branch (health ≥ 0.8) gains one new `.arg(&isv_ptr)` +passing the existing `self.isv_signals_dev_ptr` field (already on +the buffer struct, settable via `set_isv_signals_ptr` — the SAME +infrastructure `per_insert_pa` uses for ISV[440] recovery_oversample). +Zero new struct fields, no new public API. + +### Producer wireup (host-side) + +`crates/ml/src/trainers/dqn/trainer/enrichment.rs`: +- New `winner_concentration: f32` and `hindsight_magnitude: f32` + fields on `EnrichmentResult`. +- New helper functions `compute_winner_concentration(trades)` and + `compute_hindsight_magnitude(hindsight)`. +- `run_enrichments` populates both new fields; `info!` log line + extended to include both scalars. + +`crates/ml/src/trainers/dqn/trainer/training_loop.rs` (post- +enrichment block): +```rust +fused.trainer().write_isv_signal_at( + WINNER_CONCENTRATION_INDEX, + result.winner_concentration, +); +fused.trainer().write_isv_signal_at( + HINDSIGHT_MAGNITUDE_INDEX, + result.hindsight_magnitude, +); +``` + +### Files changed + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp21_isv_slots.rs` | +2 slots | `WINNER_CONCENTRATION_INDEX = 525`, `HINDSIGHT_MAGNITUDE_INDEX = 526`; `SP21_SLOT_END` 525 → 527; uniqueness test extended | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | ISV bump | `ISV_TOTAL_DIM` 525 → 527; fingerprint adds 2 SLOT entries | +| `crates/ml-dqn/src/per_kernels.cu` | ABI bump | `per_update_pa` gains `const float* __restrict__ isv_signals` arg; reads slots 525/526 device-side; composes `alpha_boost`; applies `alpha_eff = alpha + boost_delta` | +| `crates/ml-dqn/src/gpu_replay_buffer.rs` | Launcher arg | `update_priorities_gpu`'s standard-PER branch passes `self.isv_signals_dev_ptr` to `per_update_pa` | +| `crates/ml/src/trainers/dqn/trainer/enrichment.rs` | +2 fields + 2 helpers | `EnrichmentResult.winner_concentration`, `EnrichmentResult.hindsight_magnitude`; `compute_winner_concentration`, `compute_hindsight_magnitude` helpers; `run_enrichments` populates both; log line extended | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Producer | After enrichment block, writes both new ISV slots from `result.{winner_concentration, hindsight_magnitude}` | +| `docs/dqn-wire-up-audit.md` | This entry | 2026-05-11 audit log | + +### Pearls + invariants honoured + +- `feedback_no_partial_refactor` — kernel ABI change (`per_update_pa` + arg) ships with launcher migration AND producer wireup AND ISV + slot definitions atomically. Single commit. +- `feedback_no_deferrals_for_complementary_fixes` — Phases 5 + 6 + combined into one atomic commit since both attack the same + consumer kernel (`per_update_pa`) with non-overlapping refactor + scopes. +- `feedback_no_stubs` — `result.winner_concentration` and + `result.hindsight_magnitude` consumed for real (not computed-then- + discarded). +- `pearl_first_observation_bootstrap` — both signals at sentinel + 0.0 → kernel short-circuits to `alpha_eff = alpha` (no-op). + No fabricated bootstrap values. +- `pearl_controller_anchors_isv_driven` — the boost-delta computation's + thresholds (`> 1.0` for concentration; `0.02` for hindsight) are + baseline bounds, not anchors — the SIGNAL drives the magnitude + via the ISV slots. Boundaries kept hardcoded as numerical-stability + carve-outs (Invariant 1) per existing pattern. +- `feedback_isv_for_adaptive_bounds` — alpha_eff is the adaptive + bound on PER sampling sharpness; lives in ISV. +- `feedback_no_atomicadd` — kernel-side ISV reads are device-side + loads (no atomicAdd); launcher pass-through is host-side scalar. +- `feedback_cpu_is_read_only` — both signals computed once per val + pass on host (cold path), written via `write_isv_signal_at` + (mapped-pinned). No GPU↔CPU compute roundtrip in the hot path. + +### Deferred follow-up (Phase 6.5) + +True E7 hindsight injection — synthesizing replay tuples from +`HindsightExperience.{bar_index, optimal_direction, optimal_magnitude, +counterfactual_pnl}` and inserting into PER via `per_insert_pa`'s +existing recovery_oversample mechanism — requires: +1. Val state vector retention: new + `[n_windows × max_len × state_dim]` scratch buffer in + `GpuBacktestEvaluator`. +2. Action mapping: `(optimal_direction, optimal_magnitude)` → + factored action index for the synthetic tuple. +3. Insertion API extension on `GpuReplayBuffer`: accept synthetic + tuples bypassing the experience-collector path, with priority + boosted via the existing recovery_oversample mechanism. + +This is its own atomic commit when prioritised — significant +infrastructure expansion that's better designed independently of +the signal-driven Phase 5+6 commit. + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # 0 errors +cargo test -p ml --lib sp21_isv_slots --features cuda # 3/3 (bus bounds + slot uniqueness + branch index) +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: 31 GPU oracle + 3 SP21 unit = 34 tests, 0 failures. +Behavioral gate: PER alpha lift only fires once E6/E7 emit +non-sentinel values — first val pass after warmup. Smoke training +run will surface `winner_concentration` / `hindsight_magnitude` +in HEALTH_DIAG enrichment log line. + +### After this commit lands + +T2.2 multi-phase scope continues with Phases 7-8 + 6.5: +- Phase 7: wire E8 (curriculum weights) → segment sampling weights. +- Phase 8: signal-drive remaining controller GAINS (E5 0.9/1.1 steps; + E2 2.0/0.5/-0.5 epsilon thresholds; etc.). +- Phase 6.5 (deferred): true hindsight synthetic experience injection + via val state retention + `per_insert_pa` extension.