feat(sp21): T2.2 Phase 5+6 combined — E6/E7 signal-driven PER alpha scaling (atomic)
Wires E6 (compute_winner_indices) + E7 (compute_hindsight_labels) outputs into the PER priority update kernel via signal-driven ISV slots. Phases 5 and 6 combined into one atomic commit per feedback_no_deferrals_for_complementary_fixes — both attack the same consumer kernel (per_update_pa) with non-overlapping refactor scopes (E6 contributes one ISV-mediated scalar, E7 contributes another, kernel composes both into alpha_boost). Plan revision (briefed and accepted 2026-05-11): - E6's literal Vec<usize> of val bar indices CANNOT directly bump training-PER priorities — val and training have separate coord systems (no bar-index → buffer-slot mapping). - E7's true synthetic injection requires val state retention infrastructure (n_windows × max_len × state_dim scratch buffer) — significant scope expansion deferred to Phase 6.5. - The MEANINGFUL signal in both phases is scalar aggregations of the per-trade tape: E6 winner concentration (top-decile-mean P&L / all-mean P&L) and E7 hindsight magnitude (mean |counterfactual_pnl|). Both compose multiplicatively into per_update_pa's alpha_eff. ISV slot allocation: - WINNER_CONCENTRATION_INDEX = 525 - HINDSIGHT_MAGNITUDE_INDEX = 526 - ISV_TOTAL_DIM 525 → 527 (bus extension) - Layout fingerprint adds 2 SLOT entries - Compile-time test asserts SP21_SLOT_END (527) ≤ ISV_TOTAL_DIM Signal semantics: - winner_concentration: top_decile_mean_pnl / max(all_mean_pnl, EPS). > 1.0 = top decile dominates; ≈ 1.0 = uniform; ≤ 0 → 0.0 sentinel (losing strategy, signal ill-defined). - hindsight_magnitude: mean(|counterfactual_pnl|). Higher = larger missed opportunities. - Cold start: both at 0.0 sentinel; kernel short-circuits to alpha_eff = alpha (no-op). ABI surgery (minimal): - per_update_pa: ONE new arg `const float* __restrict__ isv_signals` at end. NULL-tolerant. Reads slots 525/526 device-side, composes boost_delta (bounded [0, 0.4]), applies alpha_eff = alpha + delta. - update_priorities_gpu launcher: passes existing self.isv_signals_dev_ptr (already on struct, settable via set_isv_signals_ptr — same infra per_insert_pa uses for recovery_oversample). Zero new public API. Producer wireup: - enrichment.rs: 2 new helpers (compute_winner_concentration, compute_hindsight_magnitude); EnrichmentResult gains 2 fields (winner_concentration, hindsight_magnitude); run_enrichments populates both; log line extended. - training_loop.rs: post-enrichment block writes both ISV slots via fused.trainer().write_isv_signal_at. Files changed: - crates/ml/src/cuda_pipeline/sp21_isv_slots.rs: +2 slot consts - crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs: ISV_TOTAL_DIM bump + fingerprint - crates/ml-dqn/src/per_kernels.cu: per_update_pa ABI + boost logic - crates/ml-dqn/src/gpu_replay_buffer.rs: launcher passes ISV ptr - crates/ml/src/trainers/dqn/trainer/enrichment.rs: 2 helpers + 2 fields + populate - crates/ml/src/trainers/dqn/trainer/training_loop.rs: write 2 slots - docs/dqn-wire-up-audit.md: 2026-05-11 audit entry Verification (passing): - cargo check -p ml --tests --features cuda: 0 errors - cargo test -p ml --lib sp21_isv_slots: 3/3 - 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: 34 tests, 0 failures. Behavioral gate (alpha_eff lift on post-warmup val passes) is the upcoming smoke training run. Deferred follow-up (Phase 6.5): true E7 hindsight synthetic experience injection via val state retention + per_insert_pa extension. Significant infrastructure (n_windows × max_len × state_dim scratch + insert API extension). Documented in audit. After this commit (T2.2 Phases 7-8 + 6.5 deferred): - Phase 7: E8 curriculum weights → segment sampling - Phase 8: signal-drive remaining controller GAINS - Phase 6.5: synthetic hindsight injection (deferred — needs val state buffer infrastructure) Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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}")))?;
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
|
||||
@@ -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;\
|
||||
|
||||
@@ -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<usize>` 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();
|
||||
|
||||
@@ -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<usize>,
|
||||
/// 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<HindsightExperience>,
|
||||
/// 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<f32>,
|
||||
}
|
||||
@@ -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<f32> = 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::<f32>() / top_n as f32;
|
||||
let all_mean: f32 = sorted.iter().sum::<f32>() / 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::<f32>()
|
||||
/ hindsight.len() as f32
|
||||
}
|
||||
|
||||
/// E6: Winner distillation — bar indices of top-10% trades by P&L.
|
||||
fn compute_winner_indices(trades: &[EvalTrade]) -> Vec<usize> {
|
||||
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,
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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,
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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<usize>` 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.
|
||||
|
||||
Reference in New Issue
Block a user