diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 64fc8753f..e9bf7a3f9 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2584,7 +2584,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 505; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors + Class A audit-fix Batch 4-A adaptive DD saturation floor + Class A audit-fix Batch 4-B adaptive plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE + SP16 Phase 2 adaptive Hold cost scale + SP16 T3 Wiener-optimal adaptive α Welford accumulators + SP17 dueling-Q identifiability diagnostics + SP18 v2 combined Hold-attractor fix (D-leg structural Hold opportunity-cost reward + B-leg TD(λ) Q(s') bootstrap diagnostics + PopArt reset gate). Bumped 483 → 505 by SP18 v2 Pre-Phase Task PP.2 (2026-05-08): added 22 slots [483..505) split across two atomic legs per `feedback_no_partial_refactor`. D-leg [483..493) — 10 slots: HOLD_REWARD_POS_CAP (483) + HOLD_REWARD_NEG_CAP (484) + HOLD_REWARD_DECOMP_DIAG (485) + HOLD_OPP_COST_FIRE_RATE_EMA (486) + 6 HRC_* Welford accumulators [487..493) for Wiener-α (mirrors SP16 T3 HCS_* layout). The two caps replace the SP13 P0a + SP16 P2 + SP16 T3 reactive per-bar Hold-cost chain (DD7=c) with a structural opportunity-cost reward at the 3 Hold-cost subtraction sites in `experience_kernels.cu`: each Hold-action bar pays `compute_sp18_hold_opportunity_cost(-2 × vol_normalized_next_log_return, neg_cap, pos_cap, shaping_scale)` — counterfactual cost in realised next-bar log-return units, asymmetric-capped via the `hold_reward_cap_update_kernel` adaptive producer (Phase 3) reading `p99(|step_ret|)` over Long/Short trade closes (DD3=b) × HOLD_REWARD_CAP_SAFETY_FACTOR=1.5 → POS, with NEG = -2 × POS (DD5=b mirrored 2:1 loss-aversion ratio matching SP14 P0-A REWARD_POS/NEG_CAP). Wiener-optimal α from Welford accumulators in [487..493) per `pearl_wiener_optimal_adaptive_alpha`. Pearl-A first-observation bootstrap sentinels POS=5.0/NEG=-10.0 match SP14 P0-A REWARD_POS/NEG_CAP_ADAPTIVE for bit-identical cold-start. Bounds [HOLD_REWARD_POS_CAP_MIN_BOUND=0.5, HOLD_REWARD_POS_CAP_MAX_BOUND=50.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. SP13/SP16 slots [380..383) and [461..474) remain ALLOCATED but RETIRED — sentinel 0.0, no producer launch — preserving checkpoint compatibility (RESERVED-gap pattern from SP14-C.1). B-leg [493..505) — 12 slots: TD_ERROR_MAG_EMA (493) + Q_NEXT_TARGET_P99 (494) + Q_NEXT_MINUS_REWARD_P99 (495) + V_SHARE_TREND_DIAG (496) + POPART_RESET_FLAG (497) + 6 TDB_* Welford accumulators [498..504) + 1 RESERVED (504). The B-leg replaces the `q_next = rewards` self-bootstrap at `gpu_experience_collector.rs:4143` with a real per-branch DDQN `Q_target(s', argmax_a Q_online(s', a))` computed by a target-net forward pass on `next_states` at experience-collection time (Phase 4-5). Reuses existing `target_params_buf` (SP14 P4.T1B-iv Polyak-tracked) — no new target-net infra (B-DD1). HEALTH_DIAG observability (B-DD12): td_error_mag_ema is the B-DD9 ratio-gate input (`avg(|TD-error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`); q_next_target_p99 is the bootstrap bound check; q_next_minus_reward_p99 sanity-checks magnitude (should be O(1) — proves the fix is doing something AND not blowing up); v_share_trend_diag observes the B-leg synergy hypothesis (bootstrap-bias predicts v_share monotone HD0→HD4). POPART_RESET_FLAG (497, sentinel 1.0) gates the per-fold PopArt slot 63 EMA reset at the SP18 deployment boundary per B-DD11 — switching `q_next` source from rewards-distributed to Q-distributed shifts the PopArt input distribution; cheap insurance is to reset PopArt EMA to identity normalization at the first SP18 epoch. NOT a FoldReset; host writes 1.0 once at deploy, kernel zeroes after consuming. Sentinels for diagnostic slots = 0.0 (Pearl-A direct-replace on first valid observation). 6 TDB_* Welford accumulators mirror HRC_* layout for q_next_target Wiener-α blending. Slot 504 RESERVED for B-leg follow-up. PER priority recomputation NOT migrated (per B-DD4 — already recomputed at training time from current TD-error). Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md. Bumped 474 → 483 by SP17 (2026-05-08): added 9 slots [474..483) for the dueling-Q identifiability diagnostic chain. With mean-zero advantage projection `A_centered[a, z] = A[a, z] − mean_a A[a, z]` (Phase 1-3 of SP17), V and A become uniquely identifiable: V_share ≈ 0.5 + Var(A_centered) > 0 indicates a healthy dueling decomposition; V_share → 1.0 + Var(A_centered) → 0 signals regression to the pre-SP17 V-dominated regime that the optimizer was free to learn under naive `Q = V + A`. Slots: A_VAR_EMA × 4 branches at [474..478) — per-branch EMA of the over-actions variance of the centered advantage; V_SHARE × 4 branches at [478..482) — per-branch |V| / (|V| + |A_centered|) magnitude share; ADVANTAGE_CLIP_BOUND at [482] — adaptive |A_centered| upper bound producer-tracked from p99(|A_centered|) × safety_factor=1.5 (mirrors SP14 P0-A REWARD_POS_CAP producer pattern; α=0.01 slow per-fold cadence). All Pearl-A first-observation bootstraps with sentinels SENTINEL_A_VAR_EMA=0.0, SENTINEL_V_SHARE=0.5 (cold-start 50/50 split), SENTINEL_ADVANTAGE_CLIP_BOUND=1.0 (cold-start "no effective clipping"). Bounds [ADVANTAGE_CLIP_BOUND_MIN=0.1, ADVANTAGE_CLIP_BOUND_MAX=100.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`: floor 0.1 below which centered advantage is essentially noise (clipping would erase the action signal); ceiling 100.0 above which a single adversarial outlier dominates per-batch EMAs. This commit is additive infrastructure — no consumer change yet; the Phase 1 mean-centering producer (atomic across `compute_expected_q` + `c51_loss` + `c51_grad` + `mag_concat_qdir`) lands in the next 4 commits per `feedback_no_partial_refactor`. Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md. Bumped 462 → 474 by SP16 T3 (2026-05-08): added 12 slots [462..474) — 6 Welford accumulators per producer kernel (T1 MIN_HOLD_TEMPERATURE + T2 HOLD_COST_SCALE). Replaces the hardcoded `alpha = 0.05f` EMA constant in both producer kernels (`min_hold_temperature_update_kernel.cu` + `hold_cost_scale_update_kernel.cu`) with a signal-driven Wiener-optimal α per `pearl_wiener_optimal_adaptive_alpha`: α = diff_var / (diff_var + sample_var + ε). Per train-multi-seed-hjzss validation, the SP16 T1+T2 chain was structurally landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n baseline through epoch 3) because α=0.05 needs ~60 epochs from cold start to converge. Wiener-optimal α naturally falls in [0, 1] (clamped to [0.01, 0.95] for numerical safety): cold-start (signal jumps 1.0 → 6.4 → 7.0) → diff_var dominates → α ≈ 0.6+ → near-bootstrap responsiveness in epochs 1-3; steady-state → diff_var drops → α decays naturally → smoothing emerges without hardcoded constant. Each producer maintains 6 Welford accumulators contiguously: [HCS_TARGET_MEAN, HCS_TARGET_M2, HCS_DIFF_MEAN, HCS_DIFF_M2, HCS_PREV_TARGET, HCS_SAMPLE_COUNT] at [462..468) and [MHT_TARGET_MEAN, MHT_TARGET_M2, MHT_DIFF_MEAN, MHT_DIFF_M2, MHT_PREV_TARGET, MHT_SAMPLE_COUNT] at [468..474). FoldReset sentinel 0.0 — combined with the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), the first observation of each fold replaces blended directly, the second seeds diff statistics, the third stabilises Wiener α. Per `feedback_isv_for_adaptive_bounds.md` — α was an adaptive bound being hardcoded; lifting to ISV-driven Wiener-optimal closes that violation. Bumped 461 → 462 by SP16 Phase 2 (2026-05-08): added 1 slot [461..462) — HOLD_COST_SCALE_INDEX (461). Replaces the static `ISV[HOLD_COST_INDEX=380]` per-bar Hold-cost magnitude (~0.006) when over-holding is observed. Producer kernel `hold_cost_scale_update_kernel` reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and `ISV[HOLD_RATE_TARGET_INDEX=381]` (same SP13 hold-pricing chain feeding SP16-P1's MIN_HOLD_TEMPERATURE — same cadence, same input pattern), maps overrun into a [1, 25] multiplier via `target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, and Pearl-A-bootstrap + α=0.05 mid-cadence EMA blends into slot 461. Consumer: `experience_kernels.cu` per-bar Hold cost subtractions (3 sites: segment_complete branch line 3089, per-bar positioned-Hold branch line 3553, per-bar flat-Hold branch line 3617) multiply `isv[ISV_HOLD_COST_IDX=380]` by `isv[HOLD_COST_SCALE_INDEX=461]`. Cold-start fallback: when slot 461 is at sentinel 0.0 OR outside [1, 25], the consumer uses scale=1.0 (bit-identical to pre-Phase-2 v3-P0a Hold-cost magnitude). At observed=2× target → scale=25× → effective cost ~0.15/bar competitive with per-bar reward magnitudes ~0.01-0.1 but well below SP12 asymmetric cap [-10, +5]. Bounds [1, 25] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md`. Bumped 459 → 461 by Class A audit-fix Batch 4-B (2026-05-08): added 2 slots [459..461) — PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) + MIN_HOLD_TEMPERATURE_ADAPTIVE (460). PLAN_THRESHOLD_FLOOR_ADAPTIVE replaces the hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`) — Design Y inline producer (no new kernel; same kernel writes both slot 49 and the slow-EMA-shadow floor at slot 459 in the same launch, then uses the freshly-written floor as the lower bound). Pearl-A bootstrap (sentinel 0.1 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence — the floor is a long-term shadow of `0.5 × readiness_ema` and must move much slower than the readiness EMA itself). Bounds [0.05, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (probability units). MIN_HOLD_TEMPERATURE_ADAPTIVE replaces the deleted epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary directional accuracy), maps to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion` where `confusion = 1 − clamp((short_ema − 0.5)/0.5, 0, 1)` (so dir_acc ≈ 0.5 plateau → temp HIGH/permissive exit ramp; dir_acc → 1.0 → temp LOW/strict informed-commitment), and slow-EMA-blends into slot 460. HIGH T = forgiving (`deficit/(deficit+T)` saturation factor → 0); LOW T = sharp (→ 1) per `compute_min_hold_penalty` at trade_physics.cuh:575. Pearl-A bootstrap (sentinel 50.0 matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes) + α=0.05 mid-cadence EMA. Bounds [5, 50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp) at end of training (max punishment exactly when most needed forgiveness); the signal-driven version reacts to actual realized skill. Per `feedback_no_legacy_aliases.md` the legacy schedule constants + helper function are deleted atomically. Bumped 458 → 459 by Class A audit-fix Batch 4-A (2026-05-08): added 1 adaptive-DD-saturation-floor slot [458..459) — DD_SATURATION_FLOOR_ADAPTIVE (458). Replaces the hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor)). Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). Producer kernel `dd_saturation_floor_update_kernel` aggregates per-env DD_MAX from the existing `dd_state_per_env[env*6 + 1]` tile (same source as the SP15 dd_state_kernel writes to), computes p75 estimator via Welford rule (mean + Z_75 × sigma) × safety_factor=1.5, and slow-EMA blends. Pearl-A bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property, shouldn't track per-epoch noise). Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Atomically deleted alongside this slot addition: legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` and `dd_threshold` kernel args (Item 2 Case A — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting per `feedback_no_partial_refactor`). Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) +pub(crate) const ISV_TOTAL_DIM: usize = 507; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 (post-C.1) + SP15 + SP14-C aux trunk control plane + SP14-C Phase C.4b aux horizon + Class A P0-A adaptive reward caps + Class A P1-Producer adaptive Bayesian Kelly priors + Class A audit-fix Batch 4-A adaptive DD saturation floor + Class A audit-fix Batch 4-B adaptive plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE + SP16 Phase 2 adaptive Hold cost scale + SP16 T3 Wiener-optimal adaptive α Welford accumulators + SP17 dueling-Q identifiability diagnostics + SP18 v2 combined Hold-attractor fix (D-leg structural Hold opportunity-cost reward + B-leg TD(λ) Q(s') bootstrap diagnostics + PopArt reset gate) + SP18 v2 ISV-adaptive shrink-and-perturb at fold transition. Bumped 505 → 507 by SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09): added 2 slots [505..507) replacing the hardcoded `α=0.8 / σ=0.01` constants in `fused_training.rs::reset_for_fold`'s `shrink_and_perturb(alpha, sigma)` call. Per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`, every adaptive bound must be signal-driven from ISV. The driving signal is the relative weight drift `||current_params − best_params||₂ / ||best_params||₂` already computed each epoch by the SP18 v2 weight-drift diagnostic kernel (`weight_drift_diag_kernel.cu`, commit `aab13a83f`) — extended atomically in this commit to also write α and σ into ISV[505] / ISV[506] from the existing two block-tree-reductions (no new kernel — reuses `||best||₂` and `relative_drift` already computed). SHRINK_ALPHA_ADAPTIVE_INDEX (505): drift-driven. Small drift → α near MAX (preserve the converged best-Sharpe weights hard); large drift → α near MIN (shrink hard to escape the divergent region). SHRINK_SIGMA_ADAPTIVE_INDEX (506): scale-aware noise — σ tracks `||best||_RMS` so noise stays scale-relative as weight magnitudes evolve across folds (LayerNorm γ drift, weight-decay relaxation). Pearl-A first-observation bootstrap sentinels α=0.8 / σ=0.01 match the prior hardcoded values for bit-identical cold-start before the first valid weight-drift launch. Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] + [SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bilateral clamp per `pearl_symmetric_clamp_audit`. Cold-start case (first fold, no `best_params_snapshot`): host launcher writes drift = 0.0/0.0 directly and skips the kernel; ISV slots remain at sentinels — bit-identical to prior hardcoded behavior. Producer + consumer wired in same commit per `feedback_no_partial_refactor` and `feedback_wire_everything_up`. Bumped 483 → 505 by SP18 v2 Pre-Phase Task PP.2 (2026-05-08): added 22 slots [483..505) split across two atomic legs per `feedback_no_partial_refactor`. D-leg [483..493) — 10 slots: HOLD_REWARD_POS_CAP (483) + HOLD_REWARD_NEG_CAP (484) + HOLD_REWARD_DECOMP_DIAG (485) + HOLD_OPP_COST_FIRE_RATE_EMA (486) + 6 HRC_* Welford accumulators [487..493) for Wiener-α (mirrors SP16 T3 HCS_* layout). The two caps replace the SP13 P0a + SP16 P2 + SP16 T3 reactive per-bar Hold-cost chain (DD7=c) with a structural opportunity-cost reward at the 3 Hold-cost subtraction sites in `experience_kernels.cu`: each Hold-action bar pays `compute_sp18_hold_opportunity_cost(-2 × vol_normalized_next_log_return, neg_cap, pos_cap, shaping_scale)` — counterfactual cost in realised next-bar log-return units, asymmetric-capped via the `hold_reward_cap_update_kernel` adaptive producer (Phase 3) reading `p99(|step_ret|)` over Long/Short trade closes (DD3=b) × HOLD_REWARD_CAP_SAFETY_FACTOR=1.5 → POS, with NEG = -2 × POS (DD5=b mirrored 2:1 loss-aversion ratio matching SP14 P0-A REWARD_POS/NEG_CAP). Wiener-optimal α from Welford accumulators in [487..493) per `pearl_wiener_optimal_adaptive_alpha`. Pearl-A first-observation bootstrap sentinels POS=5.0/NEG=-10.0 match SP14 P0-A REWARD_POS/NEG_CAP_ADAPTIVE for bit-identical cold-start. Bounds [HOLD_REWARD_POS_CAP_MIN_BOUND=0.5, HOLD_REWARD_POS_CAP_MAX_BOUND=50.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. SP13/SP16 slots [380..383) and [461..474) remain ALLOCATED but RETIRED — sentinel 0.0, no producer launch — preserving checkpoint compatibility (RESERVED-gap pattern from SP14-C.1). B-leg [493..505) — 12 slots: TD_ERROR_MAG_EMA (493) + Q_NEXT_TARGET_P99 (494) + Q_NEXT_MINUS_REWARD_P99 (495) + V_SHARE_TREND_DIAG (496) + POPART_RESET_FLAG (497) + 6 TDB_* Welford accumulators [498..504) + 1 RESERVED (504). The B-leg replaces the `q_next = rewards` self-bootstrap at `gpu_experience_collector.rs:4143` with a real per-branch DDQN `Q_target(s', argmax_a Q_online(s', a))` computed by a target-net forward pass on `next_states` at experience-collection time (Phase 4-5). Reuses existing `target_params_buf` (SP14 P4.T1B-iv Polyak-tracked) — no new target-net infra (B-DD1). HEALTH_DIAG observability (B-DD12): td_error_mag_ema is the B-DD9 ratio-gate input (`avg(|TD-error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`); q_next_target_p99 is the bootstrap bound check; q_next_minus_reward_p99 sanity-checks magnitude (should be O(1) — proves the fix is doing something AND not blowing up); v_share_trend_diag observes the B-leg synergy hypothesis (bootstrap-bias predicts v_share monotone HD0→HD4). POPART_RESET_FLAG (497, sentinel 1.0) gates the per-fold PopArt slot 63 EMA reset at the SP18 deployment boundary per B-DD11 — switching `q_next` source from rewards-distributed to Q-distributed shifts the PopArt input distribution; cheap insurance is to reset PopArt EMA to identity normalization at the first SP18 epoch. NOT a FoldReset; host writes 1.0 once at deploy, kernel zeroes after consuming. Sentinels for diagnostic slots = 0.0 (Pearl-A direct-replace on first valid observation). 6 TDB_* Welford accumulators mirror HRC_* layout for q_next_target Wiener-α blending. Slot 504 RESERVED for B-leg follow-up. PER priority recomputation NOT migrated (per B-DD4 — already recomputed at training time from current TD-error). Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md. Bumped 474 → 483 by SP17 (2026-05-08): added 9 slots [474..483) for the dueling-Q identifiability diagnostic chain. With mean-zero advantage projection `A_centered[a, z] = A[a, z] − mean_a A[a, z]` (Phase 1-3 of SP17), V and A become uniquely identifiable: V_share ≈ 0.5 + Var(A_centered) > 0 indicates a healthy dueling decomposition; V_share → 1.0 + Var(A_centered) → 0 signals regression to the pre-SP17 V-dominated regime that the optimizer was free to learn under naive `Q = V + A`. Slots: A_VAR_EMA × 4 branches at [474..478) — per-branch EMA of the over-actions variance of the centered advantage; V_SHARE × 4 branches at [478..482) — per-branch |V| / (|V| + |A_centered|) magnitude share; ADVANTAGE_CLIP_BOUND at [482] — adaptive |A_centered| upper bound producer-tracked from p99(|A_centered|) × safety_factor=1.5 (mirrors SP14 P0-A REWARD_POS_CAP producer pattern; α=0.01 slow per-fold cadence). All Pearl-A first-observation bootstraps with sentinels SENTINEL_A_VAR_EMA=0.0, SENTINEL_V_SHARE=0.5 (cold-start 50/50 split), SENTINEL_ADVANTAGE_CLIP_BOUND=1.0 (cold-start "no effective clipping"). Bounds [ADVANTAGE_CLIP_BOUND_MIN=0.1, ADVANTAGE_CLIP_BOUND_MAX=100.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`: floor 0.1 below which centered advantage is essentially noise (clipping would erase the action signal); ceiling 100.0 above which a single adversarial outlier dominates per-batch EMAs. This commit is additive infrastructure — no consumer change yet; the Phase 1 mean-centering producer (atomic across `compute_expected_q` + `c51_loss` + `c51_grad` + `mag_concat_qdir`) lands in the next 4 commits per `feedback_no_partial_refactor`. Plan: docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md. Bumped 462 → 474 by SP16 T3 (2026-05-08): added 12 slots [462..474) — 6 Welford accumulators per producer kernel (T1 MIN_HOLD_TEMPERATURE + T2 HOLD_COST_SCALE). Replaces the hardcoded `alpha = 0.05f` EMA constant in both producer kernels (`min_hold_temperature_update_kernel.cu` + `hold_cost_scale_update_kernel.cu`) with a signal-driven Wiener-optimal α per `pearl_wiener_optimal_adaptive_alpha`: α = diff_var / (diff_var + sample_var + ε). Per train-multi-seed-hjzss validation, the SP16 T1+T2 chain was structurally landed but BEHAVIORALLY INERT in 5-epoch smoke (bit-identical to pfh9n baseline through epoch 3) because α=0.05 needs ~60 epochs from cold start to converge. Wiener-optimal α naturally falls in [0, 1] (clamped to [0.01, 0.95] for numerical safety): cold-start (signal jumps 1.0 → 6.4 → 7.0) → diff_var dominates → α ≈ 0.6+ → near-bootstrap responsiveness in epochs 1-3; steady-state → diff_var drops → α decays naturally → smoothing emerges without hardcoded constant. Each producer maintains 6 Welford accumulators contiguously: [HCS_TARGET_MEAN, HCS_TARGET_M2, HCS_DIFF_MEAN, HCS_DIFF_M2, HCS_PREV_TARGET, HCS_SAMPLE_COUNT] at [462..468) and [MHT_TARGET_MEAN, MHT_TARGET_M2, MHT_DIFF_MEAN, MHT_DIFF_M2, MHT_PREV_TARGET, MHT_SAMPLE_COUNT] at [468..474). FoldReset sentinel 0.0 — combined with the cold-start fallback (sample_count < 3 → α = 1.0 REPLACE), the first observation of each fold replaces blended directly, the second seeds diff statistics, the third stabilises Wiener α. Per `feedback_isv_for_adaptive_bounds.md` — α was an adaptive bound being hardcoded; lifting to ISV-driven Wiener-optimal closes that violation. Bumped 461 → 462 by SP16 Phase 2 (2026-05-08): added 1 slot [461..462) — HOLD_COST_SCALE_INDEX (461). Replaces the static `ISV[HOLD_COST_INDEX=380]` per-bar Hold-cost magnitude (~0.006) when over-holding is observed. Producer kernel `hold_cost_scale_update_kernel` reads `ISV[HOLD_RATE_OBSERVED_EMA_INDEX=382]` and `ISV[HOLD_RATE_TARGET_INDEX=381]` (same SP13 hold-pricing chain feeding SP16-P1's MIN_HOLD_TEMPERATURE — same cadence, same input pattern), maps overrun into a [1, 25] multiplier via `target_scale = SCALE_MIN + (SCALE_MAX − SCALE_MIN) × clamp(max(0, observed - target) / max(target, 0.01), 0, 1)`, and Pearl-A-bootstrap + α=0.05 mid-cadence EMA blends into slot 461. Consumer: `experience_kernels.cu` per-bar Hold cost subtractions (3 sites: segment_complete branch line 3089, per-bar positioned-Hold branch line 3553, per-bar flat-Hold branch line 3617) multiply `isv[ISV_HOLD_COST_IDX=380]` by `isv[HOLD_COST_SCALE_INDEX=461]`. Cold-start fallback: when slot 461 is at sentinel 0.0 OR outside [1, 25], the consumer uses scale=1.0 (bit-identical to pre-Phase-2 v3-P0a Hold-cost magnitude). At observed=2× target → scale=25× → effective cost ~0.15/bar competitive with per-bar reward magnitudes ~0.01-0.1 but well below SP12 asymmetric cap [-10, +5]. Bounds [1, 25] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md`. Bumped 459 → 461 by Class A audit-fix Batch 4-B (2026-05-08): added 2 slots [459..461) — PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) + MIN_HOLD_TEMPERATURE_ADAPTIVE (460). PLAN_THRESHOLD_FLOOR_ADAPTIVE replaces the hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`) — Design Y inline producer (no new kernel; same kernel writes both slot 49 and the slow-EMA-shadow floor at slot 459 in the same launch, then uses the freshly-written floor as the lower bound). Pearl-A bootstrap (sentinel 0.1 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence — the floor is a long-term shadow of `0.5 × readiness_ema` and must move much slower than the readiness EMA itself). Bounds [0.05, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (probability units). MIN_HOLD_TEMPERATURE_ADAPTIVE replaces the deleted epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20} + training_loop.rs::min_hold_temperature_for_epoch). Producer kernel `min_hold_temperature_update_kernel` reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary directional accuracy), maps to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × confusion` where `confusion = 1 − clamp((short_ema − 0.5)/0.5, 0, 1)` (so dir_acc ≈ 0.5 plateau → temp HIGH/permissive exit ramp; dir_acc → 1.0 → temp LOW/strict informed-commitment), and slow-EMA-blends into slot 460. HIGH T = forgiving (`deficit/(deficit+T)` saturation factor → 0); LOW T = sharp (→ 1) per `compute_min_hold_penalty` at trade_physics.cuh:575. Pearl-A bootstrap (sentinel 50.0 matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes) + α=0.05 mid-cadence EMA. Bounds [5, 50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp) at end of training (max punishment exactly when most needed forgiveness); the signal-driven version reacts to actual realized skill. Per `feedback_no_legacy_aliases.md` the legacy schedule constants + helper function are deleted atomically. Bumped 458 → 459 by Class A audit-fix Batch 4-A (2026-05-08): added 1 adaptive-DD-saturation-floor slot [458..459) — DD_SATURATION_FLOOR_ADAPTIVE (458). Replaces the hardcoded `0.25f` saturation floor in `trade_physics.cuh::apply_margin_cap` (the upper end of the linear position-size scaling ramp dd_scale = max(0.05, 1.0 − dd_frac / saturation_floor)). Distinct semantic role from `SP15_DD_THRESHOLD_INDEX=421` (the SP15 quadratic DD-penalty *trigger* threshold, a *lower* bound). Producer kernel `dd_saturation_floor_update_kernel` aggregates per-env DD_MAX from the existing `dd_state_per_env[env*6 + 1]` tile (same source as the SP15 dd_state_kernel writes to), computes p75 estimator via Welford rule (mean + Z_75 × sigma) × safety_factor=1.5, and slow-EMA blends. Pearl-A bootstrap (sentinel 0.25 matches pre-fix hardcoded value for bit-identical cold-start) + α=0.01 slow EMA (DD distribution is a slow-moving fold-volatility property, shouldn't track per-epoch noise). Bounds [0.10, 0.50] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Atomically deleted alongside this slot addition: legacy `compute_drawdown_penalty` device function in trade_physics.cuh + its single caller at experience_kernels.cu:3822 + the `w_dd` and `dd_threshold` kernel args (Item 2 Case A — SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting per `feedback_no_partial_refactor`). Bumped 454 → 458 by Class A P1-Producer (2026-05-08): added 4 adaptive-Bayesian-prior slots [454..458) — KELLY_PRIOR_WINS (454) + KELLY_PRIOR_LOSSES (455) + KELLY_PRIOR_SUM_WINS (456) + KELLY_PRIOR_SUM_LOSSES (457). Per Class A audit, hardcoded `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / prior_sum_losses=0.01f` (kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh::kelly_position_cap:304-307) was a static "I don't know" Bayesian prior frozen at sprint-1 defaults — the cold-start of every new fold ignored what prior folds had learned about the realized Kelly distribution. Producer kernel `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats from the existing `portfolio_state[n_envs, PS_STRIDE]` buffer (same source as kelly_cap_update consumes) and slow-EMA blends into the four slots. Pearl-A bootstrap (sentinels 2.0/2.0/0.01/0.01 match pre-P1-Producer hardcoded values for bit-identical cold-start) + α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds). Bounds counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Bumped 452 → 454 by Class A P0-A (2026-05-08): added 2 adaptive-reward-cap slots [452..454) — REWARD_POS_CAP_ADAPTIVE (452) + REWARD_NEG_CAP_ADAPTIVE (453). Per Class A audit, hardcoded `REWARD_POS_CAP=+5.0f` / `REWARD_NEG_CAP=-10.0f` (state_layout.cuh:266-267) was structurally clipping the upper tail of realized alpha — controller signals (sharpe EMA, var_q, q_gap) all derive from the CAPPED buffer, so the controller cannot select for trades it cannot see. The state_layout comment explicitly deferred Phase 2 (ISV-driven) "IF Phase 1 validation reveals adaptive need"; 11 superprojects of WR-stuck-at-46-48% revealed adaptive need. Producer kernel `reward_cap_update_kernel` computes p99(winning_returns) × safety_factor=1.5 → POS cap and NEG = -2 × POS (preserves Kahneman 2:1 loss-aversion asymmetry per `pearl_audit_unboundedness_for_implicit_asymmetry` — moved from hardcoded scalar to producer-time multiplier, single source of truth). Pearl-A bootstrap (sentinels 5.0/-10.0) + Wiener-α=0.01 slow EMA. Bounds POS∈[1, 50], NEG∈[-100, -2] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`, NOT tuning. Bumped 450 → 452 by SP14 Layer C Phase C.4b (2026-05-08): added 2 aux-prediction-horizon slots [450..452) — AUX_PRED_HORIZON_BARS (450) + AVG_WIN_HOLD_TIME_BARS (451). Aux's original label was (p_{t+1} > p_t), pure HFT-scale microstructure noise — pivoted to (p_{t+H} > p_t) with H read from ISV[450], adaptively driven from observed avg winning hold time via Pearl-A bootstrap + Wiener-α EMA. Case B: existing `hold_at_exit_per_sample` + `trade_profitable_per_sample` per-sample buffers, no aggregate slot — added new aggregate slot 451 to expose the EMA target. Bumped 444 → 450 by SP14 Layer C Phase C.1 (2026-05-08): atomically deleted 9 α-machinery slots in [385..396) AND added 6 aux-trunk control-plane slots [444..450) — AUX_TRUNK_LR (444) + AUX_TRUNK_BETA1 (445) + AUX_TRUNK_BETA2 (446) + AUX_TRUNK_EPS (447) + AUX_TRUNK_GRAD_CLIP (448) + H_S2_AUX_RMS_EMA (449). Per `feedback_no_legacy_aliases` and `feedback_no_partial_refactor`. Deleted slots left as RESERVED gap (NOT compacted) to preserve checkpoint layout fingerprint compatibility — the surviving SP14 q_disagreement_* diagnostic slots (383, 384, 389) keep their original indices. Pre-C.1 base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 q_disagreement diagnostic at [383..385), [389] — Q_DISAGREEMENT_SHORT/LONG_EMA + Q_DISAGREEMENT_VARIANCE_EMA (slots [385..389) and [390..396) RESERVED gap from C.1 deletion); SP15 [397..444) per Phase 1; SP14-C [444..450) per Phase C.1. Intentional 5-slot boundary gap at [367..372)) /// 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). @@ -3697,7 +3697,9 @@ const fn layout_fingerprint_seed() -> &'static [u8] { SLOT_502_TDB_PREV_TARGET=502;\ SLOT_503_TDB_SAMPLE_COUNT=503;\ SLOT_504_SP18_B_LEG_RESERVED=504;\ - ISV_TOTAL_DIM=505;\ + SLOT_505_SHRINK_ALPHA_ADAPTIVE=505;\ + SLOT_506_SHRINK_SIGMA_ADAPTIVE=506;\ + ISV_TOTAL_DIM=507;\ SP14_C_AUX_TRUNK_CONTROL_PLANE=sp14_c_phase_1;\ ALPHA_MACHINERY_DELETED=sp14_c_phase_1;\ TRUNK_INPUT_DD_PCT=sp15_wave_4_1a;\ diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 32a5f5d9e..09a073add 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -831,6 +831,101 @@ pub const SENTINEL_POPART_RESET_FLAG: f32 = 1.0; // one-shot ⇒ PopAr pub const SP18_TD_BOOTSTRAP_SLOT_BASE: usize = 493; pub const SP18_TD_BOOTSTRAP_SLOT_END: usize = 505; +// ── SP18 v2 (2026-05-09) — ISV-adaptive shrink-and-perturb at fold transition ── +// +// 2 ISV slots [505..507) drive `reset_for_fold`'s `shrink_and_perturb` +// formula: +// +// params_flat ← α × best_params + (1 − α) × N(0, σ) +// +// Per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned` +// the prior hardcoded `α=0.8 / σ=0.01` constants at the fold-transition +// site (`fused_training.rs::reset_for_fold`) are an adaptive bound that +// must be signal-driven from ISV. The driving signal is the relative +// weight drift `||current_params − best_params||₂ / ||best_params||₂` +// computed by the SP18 v2 weight-drift diagnostic kernel +// (`weight_drift_diag_kernel.cu`, commit `aab13a83f`) — extended in this +// commit to also write α and σ into ISV[505] and ISV[506] from the +// already-computed `relative_drift` and `||best||₂`. +// +// [505] SHRINK_ALPHA_ADAPTIVE shrink factor (preserve fraction). +// Drift-driven: when end-of-fold drift +// was small (model converged near best), +// α high (preserve). When drift large +// (model diverged from best), α low +// (shrink harder to escape divergent +// region). Formula: α = 1.0 − clamp( +// relative_drift, 1.0 − ALPHA_MAX, +// 1.0 − ALPHA_MIN). +// +// [506] SHRINK_SIGMA_ADAPTIVE noise scale. Tracks weight magnitude +// so noise stays scale-relative +// regardless of how |params| evolves +// across folds. Formula: σ = clamp( +// SENTINEL_SIGMA × (||best||_RMS / +// RMS_REFERENCE), SIGMA_MIN, SIGMA_MAX). +// +// Pearl-A first-observation bootstrap: sentinels POS=0.8 / SIGMA=0.01 +// match the prior hardcoded values for bit-identical cold-start before +// the first valid producer launch. After the first kernel launch with a +// valid `best_params_snapshot`, the slots track the adaptive formula. +// +// Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] + +// [SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1 +// dimensional safety floors per `feedback_isv_for_adaptive_bounds`: +// α floor 0.5 — never less than 50% preservation (prevents catastrophic +// forgetting of best weights even on large drift). +// α ceiling 0.95 — never more than 95% (some perturbation always — +// preserves the noise-injection effect that breaks +// overfitting to the prior fold). +// σ floor 1e-4 — prevent zero-noise (gradient bias if α perturbation +// vanishes — the formula degenerates to identity). +// σ ceiling 0.1 — prevent destructive noise (10× the cold-start +// sentinel; ratio caps σ growth on weight-magnitude +// excursions). +// +// Cold-start case (first fold, no `best_params_snapshot`): the host +// launcher (`launch_weight_drift_diag`) writes drift = 0.0 / 0.0 directly +// to the output buffer and skips the kernel. The ISV slots remain at +// sentinels (0.8 / 0.01) — bit-identical to prior hardcoded behavior. +// +// SP14-C.8 / SP17 PP.3 / SP18 PP.3 lineage — single-thread cold-path +// producer per `pearl_no_host_branches_in_captured_graph`. +// +// Spec: this commit; ISV-adaptive shrink-and-perturb close-out of SP18 v2. +pub const SHRINK_ALPHA_ADAPTIVE_INDEX: usize = 505; +pub const SHRINK_SIGMA_ADAPTIVE_INDEX: usize = 506; + +// Sentinels — Pearl-A first-observation bootstrap. Match prior hardcoded +// values exactly so first-fold behavior (where `best_params_snapshot` is +// `None`) is bit-identical to pre-fix. +pub const SENTINEL_SHRINK_ALPHA: f32 = 0.8; +pub const SENTINEL_SHRINK_SIGMA: f32 = 0.01; + +// Category-1 dimensional safety bounds per `feedback_isv_for_adaptive_bounds`. +// NEVER tuned; structural floor/ceiling. Bilateral clamp per +// `pearl_symmetric_clamp_audit`. +pub const SHRINK_ALPHA_MIN: f32 = 0.5; +pub const SHRINK_ALPHA_MAX: f32 = 0.95; +pub const SHRINK_SIGMA_MIN: f32 = 1.0e-4; +pub const SHRINK_SIGMA_MAX: f32 = 0.1; + +// Producer reference scale: σ_target = SENTINEL_SHRINK_SIGMA × (||best||_RMS / +// SHRINK_SIGMA_RMS_REFERENCE). The reference is a Category-1 dimensional +// constant — the "expected RMS scale" for healthy stably-trained DQN +// weights derived from the model's mixed initialisation (He-init weights +// at ~0.05-0.15, LayerNorm γ at 1.0, biases at 0.0; the empirical mean RMS +// across all tensors is ≈ 0.5 for the foxhunt branching DQN). At +// `||best||_RMS = 0.5`, σ = 0.01 = SENTINEL — equivalence with the prior +// hardcoded value for typical healthy weights. As ||best||_RMS rises +// (LayerNorm γ drift, weight-decay relaxation), σ scales up +// proportionally so the noise term remains comparable to the typical +// magnitude of an individual weight. +pub const SHRINK_SIGMA_RMS_REFERENCE: f32 = 0.5; + +pub const SP18_SHRINK_PERTURB_SLOT_BASE: usize = 505; +pub const SP18_SHRINK_PERTURB_SLOT_END: usize = 507; + #[cfg(test)] mod tests { use super::*; @@ -1274,10 +1369,42 @@ mod tests { fn all_sp18_slots_fit_within_isv_total_dim() { use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; assert!( - SP18_TD_BOOTSTRAP_SLOT_END <= ISV_TOTAL_DIM, - "SP18_TD_BOOTSTRAP_SLOT_END={} exceeds ISV_TOTAL_DIM={}; \ + SP18_SHRINK_PERTURB_SLOT_END <= ISV_TOTAL_DIM, + "SP18_SHRINK_PERTURB_SLOT_END={} exceeds ISV_TOTAL_DIM={}; \ bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", - SP18_TD_BOOTSTRAP_SLOT_END, ISV_TOTAL_DIM, + SP18_SHRINK_PERTURB_SLOT_END, ISV_TOTAL_DIM, ); } + + /// Lock SP18 v2 ISV-adaptive shrink-and-perturb slot layout (2026-05-09). + /// 2 contiguous slots [505..507) drive `reset_for_fold`'s + /// `shrink_and_perturb` call replacing the prior hardcoded + /// `α=0.8 / σ=0.01` constants per `feedback_isv_for_adaptive_bounds` + /// and `feedback_adaptive_not_tuned`. + #[test] + fn sp18_shrink_perturb_slot_layout_locked() { + // ── Slot indices ───────────────────────────────────────────── + assert_eq!(SP18_SHRINK_PERTURB_SLOT_BASE, 505); + assert_eq!(SP18_SHRINK_PERTURB_SLOT_END, 507); + assert_eq!(SHRINK_ALPHA_ADAPTIVE_INDEX, 505); + assert_eq!(SHRINK_SIGMA_ADAPTIVE_INDEX, 506); + + // ── Sentinels — match prior hardcoded values for cold-start + // bit-identical equivalence (Pearl-A first-observation + // bootstrap before the first valid weight-drift launch). + assert_eq!(SENTINEL_SHRINK_ALPHA, 0.8); + assert_eq!(SENTINEL_SHRINK_SIGMA, 0.01); + + // ── Category-1 dimensional safety bounds. NEVER tuned; + // structural floor/ceiling per + // `feedback_isv_for_adaptive_bounds`. + assert_eq!(SHRINK_ALPHA_MIN, 0.5); + assert_eq!(SHRINK_ALPHA_MAX, 0.95); + assert_eq!(SHRINK_SIGMA_MIN, 1.0e-4); + assert_eq!(SHRINK_SIGMA_MAX, 0.1); + + // ── Producer reference scale. Category-1 dimensional anchor + // (typical mean RMS of healthy stably-trained DQN weights). + assert_eq!(SHRINK_SIGMA_RMS_REFERENCE, 0.5); + } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index f0bf09e27..b2500dc43 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -555,6 +555,37 @@ #define SENTINEL_V_SHARE_TREND_DIAG 0.0f #define SENTINEL_POPART_RESET_FLAG 1.0f +// ── SP18 v2 (2026-05-09) — ISV-adaptive shrink-and-perturb at fold transition ── +// 2 ISV slots [505..507) drive `reset_for_fold`'s `shrink_and_perturb` +// call replacing prior hardcoded `α=0.8 / σ=0.01` constants per +// `feedback_isv_for_adaptive_bounds`. Mirrors +// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of the same +// names (locked by `sp18_shrink_perturb_slot_layout_locked` test). The +// kernel `weight_drift_diag_kernel.cu` writes both slots in the same +// launch as the existing drift diagnostic — single producer, atomic +// per-epoch. Pearl-A first-observation bootstrap sentinels match prior +// hardcoded values for bit-identical cold-start (first fold, no +// `best_params_snapshot` → host skips kernel → slots stay at +// sentinels). Bounds Category-1 dimensional floors per +// `feedback_isv_for_adaptive_bounds`. Bilateral clamp per +// `pearl_symmetric_clamp_audit`. +#define SHRINK_ALPHA_ADAPTIVE_INDEX 505 +#define SHRINK_SIGMA_ADAPTIVE_INDEX 506 + +#define SENTINEL_SHRINK_ALPHA 0.8f +#define SENTINEL_SHRINK_SIGMA 0.01f + +#define SHRINK_ALPHA_MIN 0.5f +#define SHRINK_ALPHA_MAX 0.95f +#define SHRINK_SIGMA_MIN 1.0e-4f +#define SHRINK_SIGMA_MAX 0.1f + +// Producer reference scale (Category-1 dimensional anchor — typical +// mean RMS of healthy stably-trained DQN weights). At +// `||best||_RMS = 0.5`, σ = 0.01 = SENTINEL — equivalence with prior +// hardcoded value for typical healthy weights. +#define SHRINK_SIGMA_RMS_REFERENCE 0.5f + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu b/crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu index e40efc045..74a3cca68 100644 --- a/crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu +++ b/crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu @@ -1,39 +1,96 @@ /* ══════════════════════════════════════════════════════════════════════════ - * SP18 v2 weight-drift diagnostic kernel (2026-05-09). + * SP18 v2 weight-drift diagnostic + ISV-adaptive shrink-and-perturb producer + * (2026-05-09). * - * Reads `params_flat [n]` and `best_params_snapshot [n]`; emits two - * floats into a `MappedF32Buffer<2>`: + * Original cadence (commit `aab13a83f`): cold-path producer of the per-epoch + * weight-drift diagnostic. Reads `params_flat [n]` and `best_params_snapshot + * [n]`; emits two floats into a `MappedF32Buffer<4>`: * * out[0] = ||params_flat - best_params||₂ (L2 distance) * out[1] = out[0] / max(||best_params||₂, EPS) (relative) * - * Producer cadence: cold-path (epoch boundary) — single launch per - * HEALTH_DIAG emit. Diagnostic only — does NOT modify any consumer - * path. + * Surfaces HD2→HD5 weight drift across epochs so the post-deploy reviewer + * can correlate Sharpe-peak decay with actual parameter trajectory in + * `(W_best, W_curr)` space. * - * Surfaces HD2→HD5 weight drift across epochs so the post-deploy - * reviewer can correlate Sharpe-peak decay with actual parameter - * trajectory in `(W_best, W_curr)` space. + * SP18 v2 ISV-adaptive shrink-and-perturb close-out (this commit, 2026-05-09): + * extended to ALSO write the adaptive α and σ for the next fold's + * `shrink_and_perturb` call into ISV[505] / ISV[506] using the values + * already computed by the two block-tree-reductions. No new kernel — this + * is the single producer for both diagnostics and the controller (fewer + * kernels = same atomicity per the spec's "Watch for: Drift diagnostic + * timing — have the drift kernel ALSO emit the adaptive α/σ as a 4-element + * output — fewer kernels, same atomicity"). Per `feedback_isv_for_adaptive_ + * bounds` and `feedback_adaptive_not_tuned`. Outputs: + * + * out[2] = α_adaptive (clamped to [SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX]) + * out[3] = σ_adaptive (clamped to [SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX]) + * + * AND simultaneously: + * + * isv[SHRINK_ALPHA_ADAPTIVE_INDEX=505] = α_adaptive + * isv[SHRINK_SIGMA_ADAPTIVE_INDEX=506] = σ_adaptive + * + * Formulas (single-thread-0 cold-path arithmetic post-reduction): + * + * relative = ||p − b||₂ / max(||b||₂, EPS) — drift signal + * best_rms = ||b||₂ / sqrt(n) — scale signal + * + * α_target = 1.0 − clamp(relative, + * 1.0 − SHRINK_ALPHA_MAX, + * 1.0 − SHRINK_ALPHA_MIN) + * ≡ small drift → α near MAX (preserve the converged + * best-Sharpe weights hard); large drift → α near MIN + * (shrink hard to escape the divergent region). + * + * σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms / + * SHRINK_SIGMA_RMS_REFERENCE), + * SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX) + * ≡ scale-aware noise — σ tracks `||best||_RMS` so noise + * stays scale-relative as weight magnitudes evolve + * across folds (LayerNorm γ drift, weight-decay + * relaxation). At RMS = REFERENCE = 0.5, σ_target = + * SENTINEL = 0.01 → bit-identical to prior hardcoded + * value for typical healthy weights. + * + * Pearl-A first-observation bootstrap is enforced by the host launcher: + * when `best_params_snapshot` is `None` (first fold pre-best-Sharpe, or + * any fold with no improvement), the launcher writes + * [0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA] directly into + * the mapped-pinned output buffer and skips the kernel — bit-identical + * cold-start to prior hardcoded `α=0.8 / σ=0.01`. The ISV slots are NOT + * touched on that path; consumers fall back to host-side sentinels read + * via `read_isv_signal_at` (which returns the FoldReset-written sentinel + * before the first valid kernel launch lands a real signal). * * ── Pearls + invariants ──────────────────────────────────────────────── * - `feedback_no_atomicadd.md` — block tree-reduce only; no atomicAdd. * Two reductions (sum-of-squares of diff, sum-of-squares of best) * use the same shmem tile sequentially. * - `feedback_no_htod_htoh_only_mapped_pinned.md` — output is a - * `MappedF32Buffer<2>`; kernel writes via `__threadfence_system()` - * for PCIe-visible coherence; host reads after stream sync. + * `MappedF32Buffer<4>`; kernel writes via `__threadfence_system()` + * for PCIe-visible coherence; host reads after stream sync. ISV bus + * write is also mapped-pinned (host-readable directly). * - `pearl_no_host_branches_in_captured_graph.md` — kernel runs * cold-path (epoch boundary, between training-step graph * invocations); no host branches inside. - * - `pearl_first_observation_bootstrap.md` — N/A: this is a single - * instantaneous distance, not an EMA. The host-side launcher - * handles the "no snapshot saved" case before launching the kernel - * (returns 0.0/0.0 directly so the diagnostic line is still emitted). - * - `pearl_symmetric_clamp_audit.md` — relative drift is bounded - * [0, ∞); we floor the denominator at SP18_DRIFT_EPS_F=1e-12 to - * avoid 0/0 NaN when best_params are all-zero (cold-init pre-first - * save). No upper clamp — diagnostic should surface unbounded - * drift faithfully. + * - `pearl_first_observation_bootstrap.md` — first valid observation + * REPLACES sentinel directly. The producer ALWAYS writes a clamped + * target value (no EMA blend), so on the first valid launch, ISV[505] + * and ISV[506] go from sentinel directly to the computed target. + * Subsequent launches REPLACE again — there is no Welford state for + * this controller; the underlying signal (relative drift, best_rms) + * is itself computed afresh each epoch from current/best weights. + * - `pearl_symmetric_clamp_audit.md` — bilateral + * `fmaxf(MIN, fminf(x, MAX))` for both α and σ; never one-sided. + * - `feedback_isv_for_adaptive_bounds.md` — replaces the hardcoded + * `α=0.8 / σ=0.01` constants formerly at + * `fused_training.rs::reset_for_fold` lines 1014-1015. The bounds + * [MIN, MAX] are themselves Category-1 dimensional safety floors — + * structural (never adaptive themselves; their values are the floor + * and ceiling). + * - `feedback_adaptive_not_tuned.md` — α/σ derive from the EMA-tracked + * drift signal + scale-aware best_rms ratio, not tuned constants. * * Algorithm (single block, 256 threads): * @@ -48,13 +105,18 @@ * local_sumsq += best[i] * best[i] * Block tree-reduce into s_reduce[0]; thread 0 stores to local s1. * - * Pass 3 (write): + * Pass 3 (write outputs): * thread 0 only: - * l2 = sqrtf(s0) + * l2 = sqrtf(s0) * norm_best = sqrtf(s1) - * rel = l2 / max(norm_best, SP18_DRIFT_EPS_F) - * out[0] = l2 - * out[1] = rel + * rel = l2 / max(norm_best, SP18_DRIFT_EPS_F) + * best_rms = norm_best / sqrtf(n) [ NEW ] + * α_target = 1.0 − clamp(rel, 1.0 − ALPHA_MAX, 1.0 − ALPHA_MIN) [ NEW ] + * σ_target = clamp(SIGMA_SENTINEL × best_rms / RMS_REF, + * SIGMA_MIN, SIGMA_MAX) [ NEW ] + * out[0..4] = [l2, rel, α_target, σ_target] + * isv[505] = α_target [ NEW ] + * isv[506] = σ_target [ NEW ] * __threadfence_system() * * Launch contract: @@ -65,11 +127,15 @@ * Args: * params — `params_flat [n]` f32 device pointer (online weights). * best — `best_params_snapshot [n]` f32 device pointer. - * out — `MappedF32Buffer<2>` device pointer for [l2, rel]. + * out — `MappedF32Buffer<4>` device pointer for + * [l2, rel, α_target, σ_target]. + * isv — ISV bus device pointer (mapped-pinned). Writes ISV[505] + * and ISV[506]. * n — element count of `params` and `best` (must match). * ══════════════════════════════════════════════════════════════════════════ */ #include +#include "state_layout.cuh" #define SP18_DRIFT_BLOCK 256 #define SP18_DRIFT_EPS_F 1e-12f @@ -78,6 +144,7 @@ extern "C" __global__ void weight_drift_diag_update( const float* __restrict__ params, const float* __restrict__ best, float* __restrict__ out, + float* __restrict__ isv, int n) { /* Single block (gridDim.x == 1). */ @@ -86,11 +153,15 @@ extern "C" __global__ void weight_drift_diag_update( const int tid = (int)threadIdx.x; const int bdim = (int)blockDim.x; - /* Degenerate guard: n <= 0 — write zeros and return. */ + /* Degenerate guard: n <= 0 — write zeros + sentinels and return. */ if (n <= 0) { if (tid == 0) { out[0] = 0.0f; out[1] = 0.0f; + out[2] = SENTINEL_SHRINK_ALPHA; + out[3] = SENTINEL_SHRINK_SIGMA; + isv[SHRINK_ALPHA_ADAPTIVE_INDEX] = SENTINEL_SHRINK_ALPHA; + isv[SHRINK_SIGMA_ADAPTIVE_INDEX] = SENTINEL_SHRINK_SIGMA; __threadfence_system(); } return; @@ -138,8 +209,36 @@ extern "C" __global__ void weight_drift_diag_update( : SP18_DRIFT_EPS_F; const float rel = l2_diff / denom; + /* SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09): + * compute α_target and σ_target from the same reductions. */ + const float best_rms = norm_best / sqrtf((float)n); + + /* α_target: drift-driven preservation factor. Bilateral clamp + * on relative drift, then 1 − clamp gives the preserve fraction. + * `pearl_symmetric_clamp_audit` — `fmaxf(LO, fminf(rel, HI))`. */ + const float rel_lo = 1.0f - SHRINK_ALPHA_MAX; /* 0.05 */ + const float rel_hi = 1.0f - SHRINK_ALPHA_MIN; /* 0.50 */ + const float rel_clamped = fmaxf(rel_lo, fminf(rel, rel_hi)); + const float alpha_target = 1.0f - rel_clamped; + + /* σ_target: scale-aware noise. ratio = best_rms / RMS_REFERENCE + * — when weights have typical RMS, ratio = 1.0 → σ = SENTINEL. + * As RMS grows, σ scales up proportionally, capped at MAX. */ + const float ratio = best_rms / SHRINK_SIGMA_RMS_REFERENCE; + const float sigma_raw = SENTINEL_SHRINK_SIGMA * ratio; + const float sigma_target = fmaxf(SHRINK_SIGMA_MIN, + fminf(sigma_raw, SHRINK_SIGMA_MAX)); + out[0] = l2_diff; out[1] = rel; - __threadfence_system(); /* PCIe-visible write for mapped-pinned output */ + out[2] = alpha_target; + out[3] = sigma_target; + + /* ISV bus write (mapped-pinned). Consumer (`reset_for_fold`) + * reads via `read_isv_signal_at` after stream sync. */ + isv[SHRINK_ALPHA_ADAPTIVE_INDEX] = alpha_target; + isv[SHRINK_SIGMA_ADAPTIVE_INDEX] = sigma_target; + + __threadfence_system(); /* PCIe-visible write for mapped-pinned outputs */ } } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index dce8f3cc0..828d128bc 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -395,17 +395,25 @@ pub(crate) struct FusedTrainingCtx { // ── SP18 v2 weight-drift diagnostic (2026-05-09) ───────────────────── /// Single-block × 256-thread kernel computing - /// `[||params - best||₂, relative_drift]` from `params_flat` and - /// `best_params_snapshot`. Cold-path (epoch boundary). Block tree- - /// reduce per `feedback_no_atomicadd`. Loaded from - /// `weight_drift_diag_kernel.cubin`. Diagnostic-only — does NOT - /// modify any consumer path. Consumed by `read_weight_drift_diag()`. + /// `[||params - best||₂, relative_drift, α_target, σ_target]` from + /// `params_flat` and `best_params_snapshot`. Cold-path (epoch + /// boundary). Block tree-reduce per `feedback_no_atomicadd`. Loaded + /// from `weight_drift_diag_kernel.cubin`. Diagnostic + adaptive- + /// controller producer for ISV slots [505..507). Consumed by + /// `read_weight_drift_diag()` and `reset_for_fold` reading + /// ISV[SHRINK_ALPHA_ADAPTIVE_INDEX] / ISV[SHRINK_SIGMA_ADAPTIVE_INDEX]. + /// Per `feedback_isv_for_adaptive_bounds` — replaces the prior + /// hardcoded `α=0.8 / σ=0.01` constants in `reset_for_fold`. sp18_weight_drift_diag_kernel: cudarc::driver::CudaFunction, - /// Mapped-pinned 2-float output buffer for the weight-drift diagnostic. - /// Layout: `[l2_diff, relative_drift]`. Per + /// Mapped-pinned 4-float output buffer for the weight-drift diagnostic + /// + ISV-adaptive shrink-and-perturb controller. + /// Layout: `[l2_diff, relative_drift, α_target, σ_target]`. Per /// `feedback_no_htod_htoh_only_mapped_pinned`. Lifetime matches the /// fused-training context; written-then-read inside the kernel each /// launch. Host reads via `read_volatile` after stream sync. + /// SP18 v2 ISV-adaptive shrink-and-perturb close-out (2026-05-09): + /// extended from 2 to 4 elements; α/σ are also written to ISV slots + /// 505/506 inside the same kernel launch (single-producer atomicity). sp18_weight_drift_diag_buf: crate::cuda_pipeline::mapped_pinned::MappedF32Buffer, } @@ -879,11 +887,14 @@ impl FusedTrainingCtx { "weight_drift_diag_update load_function: {e}" ))? }; + // 4 floats: [l2_diff, relative_drift, α_target, σ_target] per + // SP18 v2 ISV-adaptive shrink-and-perturb close-out. The kernel + // also writes α/σ into ISV[505]/ISV[506] in the same launch. let sp18_weight_drift_diag_buf = unsafe { - crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(2) + crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(4) } .map_err(|e| anyhow::anyhow!( - "sp18_weight_drift_diag_buf alloc (2 f32): {e}" + "sp18_weight_drift_diag_buf alloc (4 f32): {e}" ))?; info!( @@ -1008,15 +1019,53 @@ impl FusedTrainingCtx { unsafe { cuda_sys::cuGraphExecDestroy(exec); } } // Shrink-and-perturb at fold boundary -- reduce overfit to previous fold. - // alpha=0.8: keep 80% of learned representations, add 20% random noise. - // Without this, the greedy policy from fold 1 produces val_Sharpe=-33 - // on fold 2 data and takes many epochs to recover. - let sp_alpha = 0.8_f32; - let sp_sigma = 0.01_f32; + // + // SP18 v2 ISV-adaptive shrink-and-perturb close-out (2026-05-09): + // α and σ are now read from ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505] + // and ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506] — replaces the prior + // hardcoded `α=0.8 / σ=0.01` constants per + // `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`. + // + // Producer: `weight_drift_diag_update` (extended in this commit + // to also write α/σ alongside the existing l2_diff/relative + // diagnostic outputs). Runs at HEALTH_DIAG cadence (per-epoch); + // by the time `reset_for_fold` executes (fold transition, + // after the last per-epoch HEALTH_DIAG emit), the slots reflect + // the drift OBSERVED at end-of-fold N, applied to fold N+1's + // perturbation. + // + // α: drift-driven preservation factor. + // - Small drift (model converged near best) → α near MAX (preserve hard). + // - Large drift (model diverged from best) → α near MIN (shrink hard). + // σ: scale-aware noise. Tracks `||best||_RMS` so noise stays + // scale-relative as weight magnitudes evolve across folds. + // + // Cold-start (first fold, no `best_params_snapshot` yet): the + // host launcher in `launch_weight_drift_diag` skipped the kernel + // and did NOT write to ISV[505]/ISV[506]; consumers fall back + // to the FoldReset-written sentinels (0.8 / 0.01) — bit-identical + // to prior hardcoded behavior. + // + // Bilateral clamp [SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX] + + // [SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX] is enforced inside the + // producer kernel per `pearl_symmetric_clamp_audit`. The host- + // side defensive clamp here guards against any out-of-range + // value persisted from a prior version's checkpoint or written + // through some unintended path — Category-1 dimensional safety + // floors per `feedback_isv_for_adaptive_bounds`. + use crate::cuda_pipeline::sp14_isv_slots::{ + SHRINK_ALPHA_ADAPTIVE_INDEX, SHRINK_SIGMA_ADAPTIVE_INDEX, + SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX, + SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX, + }; + let sp_alpha = self.trainer.read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX) + .clamp(SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX); + let sp_sigma = self.trainer.read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX) + .clamp(SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX); if let Err(e) = self.trainer.shrink_and_perturb(sp_alpha, sp_sigma) { tracing::warn!("Fold-boundary shrink-and-perturb failed (non-fatal): {e}"); } else { - tracing::info!(alpha = sp_alpha, sigma = sp_sigma, "Fold-boundary shrink-and-perturb applied"); + tracing::info!(alpha = sp_alpha, sigma = sp_sigma, "Fold-boundary shrink-and-perturb applied (ISV-adaptive)"); } // Hard-copy the shrink-and-perturb'd online weights into target params. // Without this, target_params_buf retains end-of-previous-fold weights @@ -4619,30 +4668,49 @@ impl FusedTrainingCtx { Ok(()) } - /// SP18 v2 weight-drift diagnostic (2026-05-09). + /// SP18 v2 weight-drift diagnostic + ISV-adaptive shrink-and-perturb + /// producer (2026-05-09). /// - /// Cold-path launcher for `weight_drift_diag_kernel`. Reads the - /// flat online params buffer + the best-Sharpe snapshot and emits - /// `[||params - best||₂, relative_drift]` into the mapped-pinned - /// 2-float output buffer. Single block × 256 threads; two block - /// tree-reductions (no atomicAdd per `feedback_no_atomicadd`). - /// Returns Ok(()) on success. + /// Cold-path launcher for `weight_drift_diag_kernel`. Reads the flat + /// online params buffer + the best-Sharpe snapshot and emits + /// `[||params - best||₂, relative_drift, α_target, σ_target]` into + /// the mapped-pinned 4-float output buffer. The kernel ALSO writes + /// α_target / σ_target into `ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505]` + /// and `ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506]` in the same launch + /// (single-producer atomicity per `feedback_no_partial_refactor`). + /// Single block × 256 threads; two block tree-reductions + /// (no atomicAdd per `feedback_no_atomicadd`). Returns Ok(()) on success. /// /// Cold-start: if `best_params_snapshot` is `None` (no Sharpe - /// improvement yet), the launcher writes `[0.0, 0.0]` directly into - /// the host-visible buffer and skips the kernel — diagnostic line - /// still emits with zeros so operators distinguish "no snapshot" - /// from "snapshot exists but matches current". This bypass uses the - /// mapped-pinned host slice (no DtoH per - /// `feedback_no_htod_htoh_only_mapped_pinned`). + /// improvement yet), the launcher writes + /// `[0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]` + /// directly into the host-visible buffer and skips the kernel — + /// diagnostic line still emits with zeros so operators distinguish + /// "no snapshot" from "snapshot exists but matches current". The + /// ISV slots 505/506 are NOT touched on this path; consumers fall + /// back to the FoldReset-written sentinels via `read_isv_signal_at`. + /// This bypass uses the mapped-pinned host slice (no DtoH per + /// `feedback_no_htod_htoh_only_mapped_pinned`) and is bit-identical + /// to the prior hardcoded `α=0.8 / σ=0.01` behavior. /// - /// Diagnostic only — does NOT modify any consumer path. + /// Per `feedback_isv_for_adaptive_bounds` — replaces the hardcoded + /// `α=0.8 / σ=0.01` constants formerly at + /// `fused_training.rs::reset_for_fold` lines 1014-1015. The α + /// signal is drift-driven; σ scales with `||best||_RMS`. pub(crate) fn launch_weight_drift_diag(&mut self) -> Result<()> { - // Cold-start: no snapshot saved yet. + // Cold-start: no snapshot saved yet — write zeros + sentinels + // directly into the mapped-pinned buffer and skip the kernel. + // ISV[505]/ISV[506] remain at their FoldReset sentinels (0.8/0.01) + // → bit-identical to prior hardcoded shrink-and-perturb behavior. if self.best_params_snapshot.is_none() { + use crate::cuda_pipeline::sp14_isv_slots::{ + SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA, + }; let host_slice = self.sp18_weight_drift_diag_buf.host_slice_mut(); host_slice[0] = 0.0; host_slice[1] = 0.0; + host_slice[2] = SENTINEL_SHRINK_ALPHA; + host_slice[3] = SENTINEL_SHRINK_SIGMA; return Ok(()); } @@ -4666,6 +4734,7 @@ impl FusedTrainingCtx { } let n_i32 = params_n as i32; let out_ptr = self.sp18_weight_drift_diag_buf.dev_ptr; + let isv_ptr = self.trainer.isv_signals_dev_ptr(); let block_dim: u32 = 256; let shmem_bytes: u32 = block_dim * std::mem::size_of::() as u32; @@ -4676,6 +4745,7 @@ impl FusedTrainingCtx { .arg(¶ms_ptr) .arg(&best_ptr) .arg(&out_ptr) + .arg(&isv_ptr) .arg(&n_i32) .launch(LaunchConfig { grid_dim: (1, 1, 1), @@ -4693,7 +4763,9 @@ impl FusedTrainingCtx { /// /// Convenience wrapper: launches the producer, syncs the stream, /// returns `[l2_diff, relative_drift]` from the mapped-pinned - /// output buffer. Cold-path only (HEALTH_DIAG cadence). + /// output buffer. Cold-path only (HEALTH_DIAG cadence). The α/σ + /// outputs (host_slice[2..4]) are read separately via + /// `read_shrink_perturb_adaptive()` by `reset_for_fold`. pub(crate) fn read_weight_drift_diag(&mut self) -> Result<[f32; 2]> { self.launch_weight_drift_diag()?; // Sync to ensure the kernel write (if any) is visible on host. @@ -4704,6 +4776,26 @@ impl FusedTrainingCtx { Ok([v[0], v[1]]) } + /// SP18 v2 ISV-adaptive shrink-and-perturb readback (2026-05-09). + /// + /// Returns `[α_adaptive, σ_adaptive]` for the next fold's + /// `shrink_and_perturb` call. Reads from the same mapped-pinned + /// 4-float buffer as `read_weight_drift_diag`; the kernel produces + /// both diagnostic and controller outputs in one launch. Caller + /// MUST have invoked `launch_weight_drift_diag` (or + /// `read_weight_drift_diag`) and synced the stream before this read + /// to guarantee visibility — `reset_for_fold` does so via the + /// per-epoch HEALTH_DIAG emit that precedes it. + /// + /// Cold-start (no `best_params_snapshot`): host slice carries + /// `[SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]` from the + /// `launch_weight_drift_diag` cold-start branch — bit-identical to + /// prior hardcoded `α=0.8 / σ=0.01`. + pub(crate) fn read_shrink_perturb_adaptive(&self) -> [f32; 2] { + let v = self.sp18_weight_drift_diag_buf.read_all(); + [v[2], v[3]] + } + /// Raw device pointer to the active loss scalar (pinned device-mapped). /// Returns MSE dev_ptr when c51_alpha ≈ 0, otherwise C51 total_loss dev_ptr. pub(crate) fn loss_gpu_ptr(&self) -> u64 { diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index c3f0021d3..636c8a1ac 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2020,6 +2020,28 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[SP18_B_LEG_RESERVED_INDEX=504] — SP18 B-leg RESERVED slot for follow-up if needed (e.g. distributional flavour of B-DD8 if 5-epoch smoke shows atom-distribution drift). Sentinel 0.0. No producer or consumer in the Pre-Phase scope — preserves the 12-slot B-leg block aligned and the layout fingerprint stable.", }, + // ── SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09) ───── + // 2 ISV slots [505..507) drive `reset_for_fold`'s + // `shrink_and_perturb` call replacing prior hardcoded + // `α=0.8 / σ=0.01` constants per + // `feedback_isv_for_adaptive_bounds` and + // `feedback_adaptive_not_tuned`. Producer kernel is the + // existing `weight_drift_diag_update` (`weight_drift_diag_kernel.cu`, + // commit `aab13a83f`) extended atomically in this commit to + // ALSO write α and σ from the same two block-tree-reductions. + // Pearl-A first-observation bootstrap sentinels match prior + // hardcoded values for bit-identical cold-start before the + // first valid drift-kernel launch. + RegistryEntry { + name: "sp18_shrink_alpha_adaptive", + category: ResetCategory::FoldReset, + description: "ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505] — SP18 v2 adaptive shrink factor for `reset_for_fold`'s `shrink_and_perturb` call. Drift-driven: when end-of-fold relative drift `||current − best||₂ / ||best||₂` is small (model converged near best), α high (preserve); large drift → α low (shrink harder to escape divergent region). Formula: α_target = 1 − clamp(relative_drift, 1 − ALPHA_MAX, 1 − ALPHA_MIN). FoldReset sentinel SENTINEL_SHRINK_ALPHA=0.8 — Pearl-A first-observation bootstrap matching the prior hardcoded `α=0.8` for bit-identical cold-start before the first valid drift-kernel launch (first fold has no `best_params_snapshot` → host launcher skips kernel → slot stays at sentinel). Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` — never less than 50% preservation (catastrophic-forgetting guard) / never more than 95% (some perturbation always — preserves the noise-injection effect that breaks overfitting to the prior fold). Bilateral clamp per `pearl_symmetric_clamp_audit`. Producer + consumer wired in the same commit per `feedback_no_partial_refactor`.", + }, + RegistryEntry { + name: "sp18_shrink_sigma_adaptive", + category: ResetCategory::FoldReset, + description: "ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506] — SP18 v2 adaptive noise scale for `reset_for_fold`'s `shrink_and_perturb` call. Scale-aware: σ tracks `||best||_RMS = ||best||₂ / sqrt(n)` so noise stays scale-relative as weight magnitudes evolve across folds (LayerNorm γ drift, weight-decay relaxation). Formula: σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms / SHRINK_SIGMA_RMS_REFERENCE), SIGMA_MIN, SIGMA_MAX). At RMS = REFERENCE = 0.5 (Category-1 dimensional anchor — typical mean RMS of healthy stably-trained DQN weights), σ = SENTINEL = 0.01 → bit-identical to prior hardcoded value for typical healthy weights. FoldReset sentinel SENTINEL_SHRINK_SIGMA=0.01 — Pearl-A first-observation bootstrap matching the prior hardcoded `σ=0.01` for bit-identical cold-start. Bounds [SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` — floor prevents zero-noise (gradient bias if perturbation vanishes) / ceiling prevents destructive noise (10× sentinel; ratio caps σ growth on weight-magnitude excursions). Bilateral clamp per `pearl_symmetric_clamp_audit`. Same producer kernel as slot 505 (`weight_drift_diag_update` writes both in one launch).", + }, ]; Self { entries } } @@ -2164,18 +2186,18 @@ mod tests { ); } - /// Lock SP18 v2 Pre-Phase Task PP.3: 22 fold-reset entries land for the - /// combined Hold-attractor fix slots [483..505). 10 D-leg + 12 B-leg. + /// Lock SP18 v2 Pre-Phase Task PP.3 + ISV-adaptive shrink-and-perturb + /// close-out: 24 fold-reset entries land for the combined Hold-attractor + /// fix slots [483..507). 10 D-leg + 12 B-leg + 2 shrink-and-perturb. /// Asserts each by name is present with FoldReset category and a - /// description containing "SP18 D-leg" or "SP18 B-leg" (sentinel / - /// description content checked separately for clarity). + /// description marker matching its leg. #[test] fn sp18_fold_reset_entries_present() { let r = StateResetRegistry::new(); let by_name: std::collections::HashMap<&str, &RegistryEntry> = r.all().iter().map(|e| (e.name, e)).collect(); - // 22 expected SP18 entries — D-leg 10 + B-leg 12. + // 24 expected SP18 entries — D-leg 10 + B-leg 12 + shrink-and-perturb 2. let expected: &[(&str, &str)] = &[ // D-leg [483..493) — 10 slots ("sp18_hold_reward_pos_cap", "SP18 D-leg"), @@ -2201,9 +2223,12 @@ mod tests { ("sp18_tdb_prev_target", "SP18 B-leg"), ("sp18_tdb_sample_count", "SP18 B-leg"), ("sp18_b_leg_reserved", "SP18 B-leg"), + // shrink-and-perturb [505..507) — 2 slots + ("sp18_shrink_alpha_adaptive", "SP18 v2 adaptive shrink factor"), + ("sp18_shrink_sigma_adaptive", "SP18 v2 adaptive noise scale"), ]; - assert_eq!(expected.len(), 22, "expected 22 SP18 entries (10 D-leg + 12 B-leg)"); + assert_eq!(expected.len(), 24, "expected 24 SP18 entries (10 D-leg + 12 B-leg + 2 shrink-and-perturb)"); for (name, leg_marker) in expected { let entry = by_name.get(name).unwrap_or_else(|| { diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index acea46d9f..191b2c5fc 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -5424,10 +5424,32 @@ impl DQNTrainer { // (skip_start..skip_end), so trunk drift dominates the // L2 in any case. let branch_max = relative; + // SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09): + // surface the per-epoch α/σ targets the fold-transition + // `shrink_and_perturb` will use. Read directly from ISV + // (the producer kernel that ran inside + // `read_weight_drift_diag` above wrote slots 505/506 + // alongside the diagnostic output). Cold-start (no + // `best_params_snapshot`) → kernel skipped, ISV slots + // remain at FoldReset sentinels (0.8/0.01) — surfaced + // verbatim as the cold-start values. + let (alpha_adaptive, sigma_adaptive) = if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + SHRINK_ALPHA_ADAPTIVE_INDEX, SHRINK_SIGMA_ADAPTIVE_INDEX, + }; + ( + fused.trainer().read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX), + fused.trainer().read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX), + ) + } else { + (0.0, 0.0) + }; tracing::info!( "HEALTH_DIAG[{}]: weight_drift [norm_l2={:.6} \ - relative={:.6} branch_max={:.6}]", + relative={:.6} branch_max={:.6} \ + alpha_adaptive={:.4} sigma_adaptive={:.6}]", epoch, norm_l2, relative, branch_max, + alpha_adaptive, sigma_adaptive, ); } @@ -9412,6 +9434,34 @@ impl DQNTrainer { ); } } + // ── SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09) ───── + // Producer is `weight_drift_diag_update` (extended). Pearl-A + // first-observation bootstrap sentinels match the prior + // hardcoded `α=0.8 / σ=0.01` so first-fold behavior (where + // `best_params_snapshot` is `None` → host launcher skips the + // kernel) is bit-identical to pre-fix. + "sp18_shrink_alpha_adaptive" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + SHRINK_ALPHA_ADAPTIVE_INDEX, SENTINEL_SHRINK_ALPHA, + }; + fused.trainer().write_isv_signal_at( + SHRINK_ALPHA_ADAPTIVE_INDEX, + SENTINEL_SHRINK_ALPHA, + ); + } + } + "sp18_shrink_sigma_adaptive" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + SHRINK_SIGMA_ADAPTIVE_INDEX, SENTINEL_SHRINK_SIGMA, + }; + fused.trainer().write_isv_signal_at( + SHRINK_SIGMA_ADAPTIVE_INDEX, + SENTINEL_SHRINK_SIGMA, + ); + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs b/crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs new file mode 100644 index 000000000..5368ff806 --- /dev/null +++ b/crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs @@ -0,0 +1,240 @@ +//! SP18 v2 ISV-adaptive shrink-and-perturb test (2026-05-09). +//! +//! Pins the math contract for the producer formulas in +//! `weight_drift_diag_kernel.cu` that drive ISV[505] / ISV[506] — +//! replacing the prior hardcoded `α=0.8 / σ=0.01` constants in +//! `fused_training.rs::reset_for_fold`'s `shrink_and_perturb` call. +//! +//! Per `feedback_isv_for_adaptive_bounds` and +//! `feedback_adaptive_not_tuned`, the formulas are: +//! +//! α_target = 1.0 − clamp(relative_drift, +//! 1.0 − SHRINK_ALPHA_MAX, +//! 1.0 − SHRINK_ALPHA_MIN) +//! +//! σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms / +//! SHRINK_SIGMA_RMS_REFERENCE), +//! SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX) +//! +//! ## Why this test is CPU-only +//! +//! `weight_drift_diag_update` is a GPU kernel that reads from device- +//! resident `params_flat` and `best_params_snapshot` buffers. A faithful +//! end-to-end test would require constructing a `GpuDqnTrainer` (heavy: +//! param-buf alloc, cubin loads, CUDA stream, full device init) and +//! exercising the live GPU path — deferred to L40S smoke. +//! +//! Instead, this test pins the MATH CONTRACT that the producer enforces: +//! +//! 1. Small drift → α near MAX (preserve hard). +//! 2. Large drift → α near MIN (clamped — shrink hard to escape). +//! 3. Cold-start (sentinel preserved) → α = SENTINEL exactly. +//! 4. σ scales proportionally with `||best||_RMS`. +//! +//! Real GPU coverage of the kernel-level math is provided by the L40S +//! smoke that gates the SP18 closeout dispatch. + +use ml::cuda_pipeline::sp14_isv_slots::{ + SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA, + SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX, + SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX, + SHRINK_SIGMA_RMS_REFERENCE, +}; + +/// CPU-side oracle for the α formula in `weight_drift_diag_kernel.cu`. +/// +/// Mirrors: +/// `α_target = 1.0 − clamp(relative_drift, 1.0 − ALPHA_MAX, 1.0 − ALPHA_MIN)` +/// +/// Bilateral clamp per `pearl_symmetric_clamp_audit`. +fn alpha_target_cpu(relative_drift: f32) -> f32 { + let rel_lo = 1.0 - SHRINK_ALPHA_MAX; + let rel_hi = 1.0 - SHRINK_ALPHA_MIN; + let rel_clamped = relative_drift.max(rel_lo).min(rel_hi); + 1.0 - rel_clamped +} + +/// CPU-side oracle for the σ formula in `weight_drift_diag_kernel.cu`. +/// +/// Mirrors: +/// `σ_target = clamp(SENTINEL_SIGMA × (best_rms / RMS_REFERENCE), +/// SIGMA_MIN, SIGMA_MAX)` +/// +/// Bilateral clamp per `pearl_symmetric_clamp_audit`. +fn sigma_target_cpu(best_rms: f32) -> f32 { + let ratio = best_rms / SHRINK_SIGMA_RMS_REFERENCE; + let raw = SENTINEL_SHRINK_SIGMA * ratio; + raw.max(SHRINK_SIGMA_MIN).min(SHRINK_SIGMA_MAX) +} + +/// Test 1: small drift → α high (near MAX, preserve hard). +/// +/// drift = 0.02 (model converged within 2% of best) → α should be +/// 1.0 − 0.05 = 0.95 (clamped at the floor of `1 − ALPHA_MAX`). +#[test] +fn small_drift_alpha_near_max() { + let drift = 0.02_f32; + let alpha = alpha_target_cpu(drift); + // Bilateral clamp pins α at exactly SHRINK_ALPHA_MAX when drift is + // below the lower bound `1 − ALPHA_MAX = 0.05`. + assert!( + (alpha - SHRINK_ALPHA_MAX).abs() < 1e-6, + "small drift={} should give α=ALPHA_MAX={}, got {}", + drift, SHRINK_ALPHA_MAX, alpha, + ); + // Sanity: α is in the operating band [MIN, MAX]. + assert!(alpha >= SHRINK_ALPHA_MIN && alpha <= SHRINK_ALPHA_MAX); +} + +/// Test 2: large drift → α low (clamped at MIN). +/// +/// drift = 0.50 (heavy divergence) → α should be 1.0 − 0.50 = 0.50, +/// clamped exactly at SHRINK_ALPHA_MIN. drift > 0.50 stays clamped. +#[test] +fn large_drift_alpha_clamped_min() { + // Exactly at the upper-clamp boundary. + let drift_at = 1.0 - SHRINK_ALPHA_MIN; + let alpha_at = alpha_target_cpu(drift_at); + assert!( + (alpha_at - SHRINK_ALPHA_MIN).abs() < 1e-6, + "drift=1−MIN={} should give α=ALPHA_MIN={}, got {}", + drift_at, SHRINK_ALPHA_MIN, alpha_at, + ); + + // Beyond the upper-clamp boundary — still clamped. + let drift_high = 0.80_f32; + let alpha_high = alpha_target_cpu(drift_high); + assert!( + (alpha_high - SHRINK_ALPHA_MIN).abs() < 1e-6, + "drift=0.80 should give α=ALPHA_MIN={}, got {} (bilateral clamp)", + SHRINK_ALPHA_MIN, alpha_high, + ); +} + +/// Test 3: medium drift → α in operating band (linear interpolation +/// between MIN and MAX inside the unclamped region). +/// +/// drift = 0.20 → α = 0.80 (matches the legacy hardcoded constant — +/// equivalence point for the typical mid-fold drift signal). +#[test] +fn medium_drift_alpha_interpolates() { + let drift = 0.20_f32; + let alpha = alpha_target_cpu(drift); + let expected = 1.0 - 0.20; + assert!( + (alpha - expected).abs() < 1e-6, + "drift=0.20 should give α=0.80, got {}", + alpha, + ); + assert!(alpha > SHRINK_ALPHA_MIN); + assert!(alpha < SHRINK_ALPHA_MAX); +} + +/// Test 4: σ scales with `||best||_RMS` — proportional in the +/// unclamped region. +/// +/// At RMS = REFERENCE = 0.5, σ = SENTINEL = 0.01 (bit-identical to +/// the prior hardcoded value). +/// At RMS = 2 × REFERENCE = 1.0, σ = 2 × SENTINEL = 0.02. +#[test] +fn sigma_scales_with_best_rms() { + let sigma_at_ref = sigma_target_cpu(SHRINK_SIGMA_RMS_REFERENCE); + assert!( + (sigma_at_ref - SENTINEL_SHRINK_SIGMA).abs() < 1e-7, + "RMS=REF should give σ=SENTINEL={}, got {}", + SENTINEL_SHRINK_SIGMA, sigma_at_ref, + ); + + let sigma_at_2x = sigma_target_cpu(2.0 * SHRINK_SIGMA_RMS_REFERENCE); + let expected_2x = 2.0 * SENTINEL_SHRINK_SIGMA; + assert!( + (sigma_at_2x - expected_2x).abs() < 1e-7, + "RMS=2×REF should give σ=2×SENTINEL={}, got {}", + expected_2x, sigma_at_2x, + ); +} + +/// Test 5: σ floor — RMS far below REFERENCE clamps σ at SIGMA_MIN. +#[test] +fn sigma_clamped_at_floor() { + // RMS = 0 → ratio = 0 → raw σ = 0 → clamped at SIGMA_MIN. + let sigma_zero = sigma_target_cpu(0.0); + assert!( + (sigma_zero - SHRINK_SIGMA_MIN).abs() < 1e-9, + "RMS=0 should give σ=SIGMA_MIN={:e}, got {:e}", + SHRINK_SIGMA_MIN, sigma_zero, + ); + + // Tiny RMS → still clamped (ratio = 1e-3, raw = 1e-5 ≪ MIN=1e-4). + let sigma_tiny = sigma_target_cpu(1.0e-3 * SHRINK_SIGMA_RMS_REFERENCE); + assert!( + sigma_tiny >= SHRINK_SIGMA_MIN, + "tiny RMS should clamp σ ≥ SIGMA_MIN, got {:e}", + sigma_tiny, + ); +} + +/// Test 6: σ ceiling — RMS far above REFERENCE clamps σ at SIGMA_MAX. +#[test] +fn sigma_clamped_at_ceiling() { + // RMS = 100 × REF → raw σ = 1.0 ≫ MAX=0.1 → clamped at MAX. + let sigma_huge = sigma_target_cpu(100.0 * SHRINK_SIGMA_RMS_REFERENCE); + assert!( + (sigma_huge - SHRINK_SIGMA_MAX).abs() < 1e-7, + "huge RMS should give σ=SIGMA_MAX={}, got {}", + SHRINK_SIGMA_MAX, sigma_huge, + ); +} + +/// Test 7: cold-start equivalence — at the (drift, RMS) point that +/// corresponds to "model just converged near best with healthy weight +/// magnitudes", α and σ should be in the vicinity of the prior +/// hardcoded values 0.8 / 0.01. This pins the architectural intent of +/// the formulas: typical operating point ≈ legacy constants, with +/// signal-driven excursion when drift or weight scale departs from +/// typical. +#[test] +fn typical_operating_point_near_legacy_constants() { + // drift = 0.20 → α = 0.80 (exactly matches legacy SENTINEL). + let alpha_typical = alpha_target_cpu(0.20); + assert!( + (alpha_typical - SENTINEL_SHRINK_ALPHA).abs() < 1e-6, + "typical drift=0.20 should give α=SENTINEL={}, got {}", + SENTINEL_SHRINK_ALPHA, alpha_typical, + ); + + // RMS = REFERENCE → σ = SENTINEL (exactly matches legacy SENTINEL). + let sigma_typical = sigma_target_cpu(SHRINK_SIGMA_RMS_REFERENCE); + assert!( + (sigma_typical - SENTINEL_SHRINK_SIGMA).abs() < 1e-7, + "typical RMS=REF should give σ=SENTINEL={}, got {}", + SENTINEL_SHRINK_SIGMA, sigma_typical, + ); +} + +/// Test 8: cold-start path (first fold, `best_params_snapshot=None`) +/// — the host launcher writes `[0.0, 0.0, SENTINEL_SHRINK_ALPHA, +/// SENTINEL_SHRINK_SIGMA]` to the mapped-pinned buffer and skips the +/// kernel; ISV[505]/ISV[506] remain at FoldReset sentinels. The +/// effective `shrink_and_perturb(α, σ)` call uses +/// `α=SENTINEL_SHRINK_ALPHA=0.8` and `σ=SENTINEL_SHRINK_SIGMA=0.01` — +/// bit-identical to the prior hardcoded constants. This test verifies +/// the sentinel-value contract. +#[test] +fn cold_start_uses_legacy_constants_via_sentinel() { + assert_eq!( + SENTINEL_SHRINK_ALPHA, 0.8, + "SENTINEL_SHRINK_ALPHA must equal prior hardcoded value 0.8" + ); + assert_eq!( + SENTINEL_SHRINK_SIGMA, 0.01, + "SENTINEL_SHRINK_SIGMA must equal prior hardcoded value 0.01" + ); + // The host-side defensive clamp in `reset_for_fold` (and the + // bilateral kernel clamp) must accept the sentinel values without + // modification — they are inside the operating band. + let alpha_clamped = SENTINEL_SHRINK_ALPHA.clamp(SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX); + let sigma_clamped = SENTINEL_SHRINK_SIGMA.clamp(SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX); + assert_eq!(alpha_clamped, SENTINEL_SHRINK_ALPHA); + assert_eq!(sigma_clamped, SENTINEL_SHRINK_SIGMA); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3e5618389..ab08e0542 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,45 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP18 v2 ISV-adaptive shrink-and-perturb at fold transition + +Replaces the hardcoded `α=0.8 / σ=0.01` constants in `fused_training.rs::reset_for_fold`'s `shrink_and_perturb` call with ISV-driven adaptive bounds per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`. The driving signal is the relative weight drift `||current − best||₂ / ||best||₂` already computed each epoch by the SP18 v2 weight-drift diagnostic kernel (commit `aab13a83f`). + +**Files added/modified:** + +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — 2 new ISV slot constants `SHRINK_ALPHA_ADAPTIVE_INDEX=505` + `SHRINK_SIGMA_ADAPTIVE_INDEX=506`; sentinels `SENTINEL_SHRINK_ALPHA=0.8` + `SENTINEL_SHRINK_SIGMA=0.01` (match prior hardcoded values for bit-identical cold-start); Category-1 dimensional safety bounds `[SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95]` + `[SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1]`; producer reference scale `SHRINK_SIGMA_RMS_REFERENCE=0.5` (Category-1 dimensional anchor — typical mean RMS of healthy DQN weights). Lock test `sp18_shrink_perturb_slot_layout_locked`. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `ISV_TOTAL_DIM` 505 → 507 with full SP18 close-out commentary; `layout_fingerprint_seed` extended with `SLOT_505_SHRINK_ALPHA_ADAPTIVE` + `SLOT_506_SHRINK_SIGMA_ADAPTIVE`. +- `crates/ml/src/cuda_pipeline/state_layout.cuh` — mirror `#define`s for the 2 new slots, sentinels, bounds, and `SHRINK_SIGMA_RMS_REFERENCE`. +- `crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu` — extended from 2-element output `[l2_diff, rel]` to 4-element `[l2_diff, rel, α_target, σ_target]`. Added 5th `isv` argument; same single-block × 256-thread launch + two block-tree-reductions. Thread-0 cold-path arithmetic computes `α_target = 1 − clamp(rel, 1−ALPHA_MAX, 1−ALPHA_MIN)` and `σ_target = clamp(SENTINEL_SIGMA × (best_rms / RMS_REF), SIGMA_MIN, SIGMA_MAX)` with bilateral clamp per `pearl_symmetric_clamp_audit`, writes to both the mapped-pinned output buffer AND `isv[505] / isv[506]` with `__threadfence_system()`. No new kernel — this is the sole producer of both diagnostic and controller outputs (single-launch atomicity per the spec's "Watch for: Drift diagnostic timing — have the drift kernel ALSO emit the adaptive α/σ as a 4-element output — fewer kernels, same atomicity"). +- `crates/ml/src/trainers/dqn/fused_training.rs` — `sp18_weight_drift_diag_buf` resized 2 → 4 floats; `launch_weight_drift_diag` plumbs `isv_signals_dev_ptr` as a new arg + cold-start path writes `[0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]` to the host-pinned slice (skip kernel; ISV slots remain at FoldReset sentinels — bit-identical to prior hardcoded behavior). New convenience reader `read_shrink_perturb_adaptive`. `reset_for_fold` migrated: replaces hardcoded `let sp_alpha = 0.8; let sp_sigma = 0.01;` with ISV reads via `read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX)` / `read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX)` + defensive host-side clamp at the same MIN/MAX bounds the kernel enforces (Category-1 dimensional safety). 4-arg `shrink_and_perturb(alpha, sigma)` signature unchanged — no ripple to the 3 other call sites in `training_loop.rs` (those continue to use `hyperparams.shrink_perturb_alpha/sigma` — different driving signal: unhealthy-streak rescue vs fold-transition plasticity refresh). +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — 2 new `RegistryEntry` records (`sp18_shrink_alpha_adaptive`, `sp18_shrink_sigma_adaptive`) with `FoldReset` category. Sentinels match prior hardcoded values for cold-start bit-identical equivalence. `sp18_fold_reset_entries_present` lock test bumped to assert 24 SP18 entries (was 22). +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — 2 dispatch arms in `reset_named_state` writing the sentinels at fold boundary (so the cold-start fallback in `launch_weight_drift_diag` reads the sentinel rather than stale state). Per-epoch HEALTH_DIAG `weight_drift` line extended: `[norm_l2={} relative={} branch_max={} alpha_adaptive={} sigma_adaptive={}]` — surfaces the per-epoch α/σ targets the next fold-transition `shrink_and_perturb` will use. +- `crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs` (new) — 8 CPU-only oracle tests pinning the math contract: small drift → α near MAX, large drift → α clamped MIN, medium drift → linear interpolation, σ scales with `||best||_RMS`, σ floor/ceiling clamps, typical operating point ≡ legacy constants, cold-start = sentinel = legacy. + +**Atomicity envelope** (per `feedback_no_partial_refactor` and `feedback_wire_everything_up`): ISV slots + producer + consumer + HEALTH_DIAG + tests all land in the same commit. The producer kernel's outputs are wired to two consumers in this commit — the existing diagnostic readback (`read_weight_drift_diag` returns `[l2, rel]` from `out[0..2]`) and the new fold-transition controller (`reset_for_fold` reads `ISV[505] / ISV[506]`). + +**Cold-start handling** (per `pearl_first_observation_bootstrap`): on the first fold or any fold where `best_params_snapshot` is `None` (no Sharpe improvement yet), the host launcher writes zeros + sentinels to the mapped-pinned buffer and skips the kernel. ISV slots stay at FoldReset sentinels (`0.8 / 0.01`) — `reset_for_fold`'s `shrink_and_perturb(0.8, 0.01)` call is bit-identical to the prior hardcoded behavior. + +**Producer cadence**: HEALTH_DIAG (per-epoch). The drift kernel runs at `read_weight_drift_diag()` inside the per-epoch HEALTH_DIAG block, which sync's the stream — host visibility of `ISV[505]/ISV[506]` is guaranteed before the next fold-transition. `reset_for_fold` reads the values written by the LAST epoch's drift launch — i.e., drift OBSERVED at end-of-fold N, applied to fold N+1's perturbation (per the spec). + +**Concerns surfaced** (DONE_WITH_CONCERNS): + +1. **3 other `shrink_and_perturb` call sites unchanged**: `training_loop.rs:1096` (Phase 3 boundary), `training_loop.rs:3549` (D3/N3 health-triggered), `training_loop.rs:3609` (D6/N6 ensemble-collapse-triggered), and `training_loop.rs:7320` (PerturbationStrategy::ShrinkPerturbBranches with hardcoded `0.9, 0.1`). These all read from `self.hyperparams.shrink_perturb_alpha/sigma` (or hardcoded `0.9, 0.1` for backtracking). They're a different driving-signal regime — health-driven rescue interventions, not fold-transition plasticity refresh. Migrating those to ISV-adaptive is a separate concern (different signal: health collapse + ensemble disagreement, not weight drift) — surfaced for follow-up but explicitly out of scope for this commit per `feedback_no_partial_refactor` (the contract change is the fold-transition site only; the others have their own driving signal). + +2. **Cold-start drift = 0.0 → α = 1.0?**: Per the spec's "Watch for" section, when drift = 0 and we're not in cold-start, `α_target = 1.0 − clamp(0, 0.05, 0.5) = 1.0 − 0.05 = 0.95` — clamped at MAX, not 1.0. The bilateral clamp's lower bound `1 − ALPHA_MAX = 0.05` ensures we never get α=1.0 (full preservation, no perturbation). The cold-start case (no `best_params_snapshot`) is handled via the host-side bypass writing sentinels — drift never appears as 0.0 from the kernel side because the kernel doesn't run in cold-start. + +**Verification**: +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace` — clean. +- `cargo test -p ml --lib sp18` — 4 SP18 lock tests pass: `all_sp18_slots_fit_within_isv_total_dim`, `sp18_combined_slot_layout_locked`, `sp18_shrink_perturb_slot_layout_locked`, `sp18_fold_reset_entries_present`. +- `cargo test -p ml --lib every_fold_and_soft_reset_entry_has_dispatch_arm` — pass (24 SP18 dispatch arms verified). +- `cargo test -p ml --test sp18_shrink_perturb_adaptive_test` — 8/8 pass (CPU oracle math contract). +- `cargo test -p ml --test sp18_fold_transition_test` — 3/3 pass (existing fold-transition restore-best test still validates). + +**HEALTH_DIAG line format** (extended): +``` +HEALTH_DIAG[N]: weight_drift [norm_l2=0.123456 relative=0.234567 branch_max=0.234567 alpha_adaptive=0.7654 sigma_adaptive=0.012345] +``` + ## 2026-05-08 — SP18 v2 Phase 0 Task 0.2: B-leg V_SHARE trajectory + TD-error magnitude diagnostic Phase 0 observability: kernel + Rust launcher + V_SHARE history ring buffer + per-epoch HEALTH_DIAG emit + GPU oracle tests. All in the same atomic commit per `feedback_no_partial_refactor`.