From 0b9ea77dc41f36809b3b9e011c1c0a4b7b8596a1 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 8 May 2026 11:13:54 +0200 Subject: [PATCH] fix(class-a-audit-batch-4b): plan_threshold floor adaptive + MIN_HOLD_TEMPERATURE ISV-driven Per Class A audit-fix Batch 4-B (final 2 of 4 deferred items from P1-wiring/P1-producer). Completes the 8-commit WR-plateau intervention chain. Validation deferred to next L40S smoke. Item 3: plan_threshold adaptive floor (Design Y - inline producer) - NEW slot PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459 - Writes slow-EMA shadow of `0.5 * readiness_ema` from inside the existing update kernel (no new file, no new launch). - Pearl-A first-observation bootstrap (sentinel 0.1 matches pre-fix hardcoded value for bit-identical cold-start) + Welford alpha=0.005 slow EMA. - Bilateral clamp [0.05, 0.50] (probability units) per pearl_symmetric_clamp_audit. - Consumer reads isv[459] as the floor in the same launch's final fmaxf; cold-start sentinel REPLACES with threshold_target so the pre-fix `fmaxf(0.1, 0.5*ema)` semantic is preserved bit-identical for any readiness EMA above 0.20. Item 4: MIN_HOLD_TEMPERATURE -> ISV-driven (driving signal: dir_acc skill) - NEW slot MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460 - NEW kernel min_hold_temperature_update_kernel.cu (single-thread cold-path, per-epoch boundary launch). - Driving signal: dir_acc skill = clamp((short_ema - 0.5)/0.5, 0, 1) from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. When committing skillfully (high dir_acc) -> temp HIGH (permissive). When at random baseline (~0.5) -> temp LOW (sharp commitment pressure). Substituted for the audit-spec's `dir_entropy_deficit` because no dir_entropy ISV slot exists - dir_acc skill is the closest semantically-equivalent signal that preserves the spec intent. - Pearl-A bootstrap (sentinel 50.0 matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor) + alpha=0.05 mid-cadence EMA. - Bounds [5, 50] (matches the deleted schedule range). - Decouples temperature from epoch number - the old schedule pinned LOW temp (sharp) at end of training, exactly when a WR-plateaued model needed forgiveness to escape. - DELETED: state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} #defines + training_loop.rs::min_hold_temperature_for_epoch helper function (kept docstring tombstone explaining the deletion). Both call sites migrated to the new ISV reader. Per feedback_no_legacy_aliases + feedback_no_partial_refactor. ISV_TOTAL_DIM: 459 -> 461. Cumulative WR-plateau fix series (final commit, #8): - 8f218cab2 (Class C bug 1 + P0-B) - 316db416b (P0-C MIN_HOLD_TARGET) - 394de7d43 (P0-A REWARD_POS/NEG_CAP) - c4b6d6ef2 (P1 wiring var_floor) - 657972a4b (P0-A downstream DD penalty + MIN_HOLD_PENALTY_MAX) - 87d597d5d (P1 producer Bayesian priors) - 7e9a8f6ef (Batch 4-A DD saturation floor + legacy DELETE) - this commit (Batch 4-B plan_threshold floor + MIN_HOLD_TEMPERATURE) Verification: cargo check -p ml --tests --all-targets PASS cargo test sp14 (lib) 16/16 PASS (slot layout locks) cargo test sp15 (lib) 2/2 PASS (no regression) cargo test state_reset_registry (lib) 4/4 PASS (dispatch coverage) cargo test sp14_oracle_tests (GPU) 24/24 PASS (8 new + 16 pre) cargo test sp15_phase1_oracle_tests (GPU) 36/36 PASS (no regression) Per feedback_isv_for_adaptive_bounds + feedback_no_partial_refactor + feedback_no_legacy_aliases. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 24 + crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs | 101 +++- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 104 +++- .../cuda_pipeline/gpu_experience_collector.rs | 47 +- .../min_hold_temperature_update_kernel.cu | 158 ++++++ .../plan_threshold_update_kernel.cu | 68 ++- crates/ml/src/cuda_pipeline/sp14_isv_slots.rs | 201 +++++++ crates/ml/src/cuda_pipeline/state_layout.cuh | 62 ++- .../src/trainers/dqn/state_reset_registry.rs | 63 +++ .../src/trainers/dqn/trainer/training_loop.rs | 195 +++++-- crates/ml/tests/sp14_oracle_tests.rs | 499 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 105 ++++ 12 files changed, 1565 insertions(+), 62 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 8dcabbe8d..5e5a521a9 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -856,6 +856,30 @@ fn main() { // Block-tree-reduce (no atomicAdd) per `feedback_no_atomicadd.md`. // Per-epoch boundary launch. "dd_saturation_floor_update_kernel.cu", + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive + // MIN_HOLD_TEMPERATURE producer. Single-thread cold-path kernel + // that reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast + // EMA of binary directional accuracy), maps it to a [5, 50] + // temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × + // skill` where `skill = clamp((short_ema − 0.5)/0.5, 0, 1)`, + // and slow-EMA-blends into + // `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`. Pearl-A + // first-observation 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 + // thereafter. Replaces the deleted epoch-driven schedule + // `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly + // state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} + + // training_loop.rs::min_hold_temperature_for_epoch). Decouples + // temperature from epoch number — when WR plateaus, the old + // schedule pinned LOW temp (sharp) at end of training (max + // punishment when most needed forgiveness); the signal-driven + // version reacts to actual realized skill. Per + // `feedback_no_legacy_aliases.md` the legacy schedule is deleted + // atomically. Bounds [5, 50] are Category-1 dimensional safety + // floors per `feedback_isv_for_adaptive_bounds.md`. Per-epoch + // boundary launch (same cadence as P0-A REWARD_POS_CAP). + "min_hold_temperature_update_kernel.cu", // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. // Single-block 256-thread kernel computing RMS = sqrt(mean(x²)) // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation) diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index dd062f75d..1243554aa 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -53,7 +53,8 @@ use crate::MLError; use super::gpu_dqn_trainer::{ AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN, AVG_WIN_HOLD_TIME_UPDATE_CUBIN, DD_SATURATION_FLOOR_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, - KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, REWARD_CAP_UPDATE_CUBIN, + KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, MIN_HOLD_TEMPERATURE_UPDATE_CUBIN, + REWARD_CAP_UPDATE_CUBIN, }; /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU @@ -965,6 +966,104 @@ impl DdSaturationFloorUpdateOps { } } +/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive +/// MIN_HOLD_TEMPERATURE producer. +/// +/// Single-thread cold-path kernel (per-epoch boundary cadence) that +/// reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of +/// binary directional accuracy), maps it to a [5, 50] temperature via +/// +/// skill = clamp((short_ema − 0.5) / 0.5, 0, 1) +/// temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × skill +/// +/// and slow-EMA-blends into +/// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]` with α=0.05. +/// +/// Replaces the deleted epoch-driven anneal schedule +/// `T_end + (T_start − T_end) × exp(-epoch / decay)` (formerly +/// state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY} + +/// training_loop.rs::min_hold_temperature_for_epoch). 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. +/// +/// # Pearls applied +/// - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. +/// - `pearl_first_observation_bootstrap.md` — sentinel 50.0 → REPLACE. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded at construction; on-device guards. +/// - `pearl_symmetric_clamp_audit.md` — bilateral +/// `fmaxf(lo, fminf(x, hi))` clamp on target_temp + post-blend. +/// - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are +/// Category-1 dimensional safety floors (matches the original +/// schedule range MIN_HOLD_TEMPERATURE_END=5 ↔ +/// MIN_HOLD_TEMPERATURE_START=50), NOT tuning. +/// - `feedback_no_legacy_aliases.md` — the legacy schedule constants +/// and helper function are deleted atomically with this wiring. +#[allow(missing_debug_implementations)] +pub(crate) struct MinHoldTemperatureUpdateOps { + update_kernel: CudaFunction, +} + +impl MinHoldTemperatureUpdateOps { + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(MIN_HOLD_TEMPERATURE_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("min_hold_temperature_update cubin load: {e}")))?; + let update_kernel = module + .load_function("min_hold_temperature_update") + .map_err(|e| MLError::ModelError(format!("min_hold_temperature_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the adaptive MIN_HOLD_TEMPERATURE producer. + /// + /// Args: + /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). + /// - `temp_idx`: MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). + /// - `dir_acc_idx`: AUX_DIR_ACC_SHORT_EMA_INDEX (373). + /// - `sentinel_temp`: SENTINEL_MIN_HOLD_TEMPERATURE (50.0). + /// - `sentinel_dir_acc`: DIR_ACC_RANDOM_BASELINE (0.5; matches + /// sp13 DIR_ACC_EMA_SENTINEL). + /// - `temp_min`/`temp_max`: MIN_HOLD_TEMPERATURE_MIN/MAX (5.0, 50.0). + /// - `alpha`: MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + isv_ptr: u64, + temp_idx: i32, + dir_acc_idx: i32, + sentinel_temp: f32, + sentinel_dir_acc: f32, + temp_min: f32, + temp_max: f32, + alpha: f32, + ) -> Result<(), MLError> { + unsafe { + stream + .launch_builder(&self.update_kernel) + .arg(&isv_ptr) + .arg(&temp_idx) + .arg(&dir_acc_idx) + .arg(&sentinel_temp) + .arg(&sentinel_dir_acc) + .arg(&temp_min) + .arg(&temp_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("min_hold_temperature_update: {e}")))?; + } + Ok(()) + } +} + /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer. /// /// Computes `RMS(h_s2_aux) = sqrt(mean(h_s2_aux²))` over the diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index fabe64a3d..daee3f979 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2140,6 +2140,24 @@ pub(crate) static KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN: &[u8] = include_bytes!(con /// `gpu_aux_trunk::DdSaturationFloorUpdateOps`. Per-epoch boundary launch. pub(crate) static DD_SATURATION_FLOOR_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/dd_saturation_floor_update_kernel.cubin")); +/// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive +/// MIN_HOLD_TEMPERATURE producer cubin. Single-thread cold-path kernel +/// that reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of +/// binary directional accuracy), maps it to a [5, 50] temperature via +/// `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × skill` where +/// `skill = clamp((short_ema − 0.5)/0.5, 0, 1)`, and slow-EMA-blends +/// into `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`. Replaces the +/// deleted epoch-driven anneal schedule `T_end + (T_start − T_end) × +/// exp(-epoch / decay)` (formerly state_layout.cuh:: +/// MIN_HOLD_TEMPERATURE_{START, END, DECAY} + +/// training_loop.rs::min_hold_temperature_for_epoch). Pearl-A first- +/// observation bootstrap (sentinel 50.0 matches the deleted +/// MIN_HOLD_TEMPERATURE_START=50 for bit-identical cold-start) + α=0.05 +/// mid-cadence EMA. Loaded by +/// `gpu_aux_trunk::MinHoldTemperatureUpdateOps`. Per-epoch boundary +/// launch (same cadence as P0-A REWARD_POS_CAP). +pub(crate) static MIN_HOLD_TEMPERATURE_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/min_hold_temperature_update_kernel.cubin")); + /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer cubin. /// Single-block 256-thread kernel computing `sqrt(mean(h_s2_aux²))` over /// `h_s2_aux [B, SH2]` (aux trunk final output) and EMA-blending into @@ -2468,7 +2486,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 = 459; // 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. 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 = 461; // 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. 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) × clamp((short_ema − 0.5)/0.5, 0, 1)`, and slow-EMA-blends into slot 460. 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). @@ -3544,7 +3562,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { REWARD_POS_CAP_ADAPTIVE=452;REWARD_NEG_CAP_ADAPTIVE=453;\ KELLY_PRIOR_WINS=454;KELLY_PRIOR_LOSSES=455;KELLY_PRIOR_SUM_WINS=456;KELLY_PRIOR_SUM_LOSSES=457;\ DD_SATURATION_FLOOR_ADAPTIVE=458;\ - ISV_TOTAL_DIM=459;\ + PLAN_THRESHOLD_FLOOR_ADAPTIVE=459;MIN_HOLD_TEMPERATURE_ADAPTIVE=460;\ + ISV_TOTAL_DIM=461;\ 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;\ @@ -7581,6 +7600,21 @@ pub struct GpuDqnTrainer { /// (mirrors P0-A REWARD_POS_CAP cadence). dd_saturation_floor_update_ops: super::gpu_aux_trunk::DdSaturationFloorUpdateOps, + /// Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive + /// MIN_HOLD_TEMPERATURE producer Ops handle. Reads + /// `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary + /// directional accuracy), maps it to a [5, 50] temperature via + /// `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema − + /// 0.5)/0.5, 0, 1)`, and slow-EMA-blends into + /// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`. Replaces the + /// deleted epoch-driven anneal schedule + /// `T_end + (T_start − T_end) × exp(-epoch / decay)`. Pearl-A + /// first-observation 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. Per-epoch + /// boundary launch (same cadence as P0-A REWARD_POS_CAP). + min_hold_temperature_update_ops: super::gpu_aux_trunk::MinHoldTemperatureUpdateOps, + // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. speculative_h_s2: CudaSlice, @@ -14823,6 +14857,59 @@ impl GpuDqnTrainer { Ok(()) } + /// Class A audit-fix Batch 4-B (2026-05-08, Item 4): launch adaptive + /// MIN_HOLD_TEMPERATURE producer. + /// + /// Reads `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of + /// binary directional accuracy), maps it to a [5, 50] temperature + /// via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema + /// − 0.5)/0.5, 0, 1)`, and slow-EMA-blends into + /// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]`. Replaces the + /// deleted epoch-driven anneal schedule + /// `T_end + (T_start − T_end) × exp(-epoch / decay)`. + /// + /// Pearls applied: + /// - `pearl_first_observation_bootstrap.md` (sentinel = 50.0, + /// matches the deleted MIN_HOLD_TEMPERATURE_START=50) + /// - `pearl_no_host_branches_in_captured_graph.md` + /// - `pearl_symmetric_clamp_audit.md` (bilateral clamp on + /// target_temp + post-blend) + /// - `feedback_isv_for_adaptive_bounds.md` (bounds [5, 50] are + /// dimensional safety floors, not tuning — matches the + /// original schedule range) + /// - `feedback_no_legacy_aliases.md` (the legacy schedule + /// constants and helper function are deleted atomically with + /// this wiring) + pub(crate) fn launch_min_hold_temperature_update( + &self, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp14_isv_slots::{ + DIR_ACC_RANDOM_BASELINE, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + }; + use crate::cuda_pipeline::sp5_isv_slots::AUX_DIR_ACC_SHORT_EMA_INDEX; + + debug_assert!( + self.isv_signals_dev_ptr != 0, + "launch_min_hold_temperature_update: isv_signals_dev_ptr must be allocated" + ); + + self.min_hold_temperature_update_ops.launch( + &self.stream, + self.isv_signals_dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, + DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, + MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + )?; + + Ok(()) + } + /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. /// @@ -23253,6 +23340,17 @@ impl GpuDqnTrainer { let dd_saturation_floor_update_ops = super::gpu_aux_trunk::DdSaturationFloorUpdateOps::new(&stream)?; + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive + // MIN_HOLD_TEMPERATURE producer. Pre-loads the `CudaFunction` + // handle at construction per + // `pearl_no_host_branches_in_captured_graph.md`. Per-epoch + // boundary launch (mid-cadence — temperature should track the + // dir_acc skill EMA which itself updates per-step but with + // α=0.05 to filter per-batch noise; same cadence as P0-A + // REWARD_POS_CAP). + let min_hold_temperature_update_ops = + super::gpu_aux_trunk::MinHoldTemperatureUpdateOps::new(&stream)?; + // ── Speculative inference cache ── let speculative_h_s2 = stream.alloc_zeros::(b * sh2) .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?; @@ -24090,6 +24188,8 @@ impl GpuDqnTrainer { kelly_bayesian_priors_update_ops, // Class A audit-fix Batch 4-A (2026-05-08): adaptive DD saturation floor producer. dd_saturation_floor_update_ops, + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): adaptive MIN_HOLD_TEMPERATURE producer. + min_hold_temperature_update_ops, speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 6746e0d51..9c65d84f0 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -315,12 +315,23 @@ pub struct ExperienceCollectorConfig { /// (60% of `REWARD_POS_CAP`). pub min_hold_penalty_max: f32, /// SP12 v3: temperature controller for the min-hold soft saturation. - /// Computed per-epoch in `training_loop.rs::min_hold_temperature_for_epoch` - /// via the annealing schedule - /// `T_end + (T_start - T_end) × exp(-epoch / decay)` with constants - /// from `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}`. - /// Default at construction = `MIN_HOLD_TEMPERATURE_START=50.0` - /// (epoch-0 forgiving regime; replaced per epoch by training loop). + /// + /// Class A audit-fix Batch 4-B (2026-05-08, Item 4): MIGRATED from + /// the deleted epoch-driven schedule + /// `T_end + (T_start - T_end) × exp(-epoch / decay)` (formerly + /// `training_loop.rs::min_hold_temperature_for_epoch`) to the + /// signal-driven ISV slot + /// `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460`. The training loop + /// reads the freshly-produced value via + /// `DQNTrainer::read_min_hold_temperature_from_isv` and seeds this + /// field per-epoch. Producer kernel: + /// `min_hold_temperature_update_kernel.cu` (per-epoch boundary). + /// + /// Default at construction = 50.0 (= MIN_HOLD_TEMPERATURE_FALLBACK + /// = the deleted MIN_HOLD_TEMPERATURE_START anchor for bit-identical + /// pre-Batch-4B epoch-0 behavior). The kernel-passed scalar is the + /// cold-start fallback used by the consumer when the ISV slot is + /// at sentinel. pub min_hold_temperature: f32, } @@ -397,9 +408,13 @@ impl Default for ExperienceCollectorConfig { hindsight_lookahead: 10, noise_sigma_per_branch: [0.1; 4], // SP12 v3 defaults — match Invariant-1 anchors in state_layout.cuh. - // The training loop overrides `min_hold_temperature` per epoch via - // the annealing schedule (50 → 5 across ~100 epochs); the other - // two are static Phase 1 constants. + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): the training + // loop overrides `min_hold_temperature` per epoch via + // `DQNTrainer::read_min_hold_temperature_from_isv` (signal-driven + // ISV slot 460 — replaces the deleted epoch-driven schedule); + // the default 50.0 here is the cold-start fallback anchor matching + // the deleted MIN_HOLD_TEMPERATURE_START. The other two are + // static Phase 1 constants. min_hold_target: 30.0, min_hold_penalty_max: 3.0, min_hold_temperature: 50.0, @@ -5665,11 +5680,15 @@ impl GpuExperienceCollector { // ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451]; `min_hold_target` // here is the cold-start fallback only (used when ISV[451] // is still at sentinel 0.0 or isv_signals_ptr is NULL). - // Constants in state_layout.cuh (MIN_HOLD_TARGET, - // MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_{START,END, - // DECAY}). Temperature recomputed per epoch via the - // annealing schedule in `training_loop.rs:: - // min_hold_temperature_for_epoch`. + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): + // `min_hold_temperature` is now signal-driven via ISV + // slot 460 (sourced from + // ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] via + // `min_hold_temperature_update_kernel`); the + // training loop seeds this scalar per-epoch via + // `DQNTrainer::read_min_hold_temperature_from_isv`, + // which falls back to MIN_HOLD_TEMPERATURE_FALLBACK + // =50.0 when the slot is at sentinel. .arg(&config.min_hold_target) .arg(&config.min_hold_penalty_max) .arg(&config.min_hold_temperature) diff --git a/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu b/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu new file mode 100644 index 000000000..cf82d1302 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu @@ -0,0 +1,158 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * Class A audit-fix Batch 4-B — adaptive MIN_HOLD_TEMPERATURE producer + * (2026-05-08, Item 4). + * + * 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) with an + * ISV-driven signal-driven temperature derived from the model's + * directional-accuracy skill. The temperature controls the + * `deficit / (deficit + T)` saturation factor in `compute_min_hold_penalty` + * (trade_physics.cuh:575): HIGH T = forgiving (gradient is gentle near + * deficit=0); LOW T = sharp (penalty rises steeply at any deficit). + * + * Why this matters (per Class A audit deferred-item batch 4-B): + * + * The epoch-driven schedule pinned LOW temp (T=5, sharp) at end of + * training. This works ONLY if the model gets more skilled with + * training — which is exactly the assumption that fails when WR + * plateaus. Maximum punishment is delivered exactly when the model + * most needs forgiveness to escape the plateau. The signal-driven + * replacement decouples temperature from epoch number: when the + * model is committing skillfully (dir_acc > 0.5) → temp HIGH + * (permissive); when at random baseline (dir_acc ≈ 0.5) → temp LOW + * (sharp commitment pressure to escape the random-guess regime). + * + * Driving signal: dir_acc skill from + * `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` — the SP13 fast EMA of + * binary directional accuracy (sentinel 0.5 = random baseline). + * + * skill = clamp((short_ema - 0.5) / 0.5, 0, 1) ∈ [0, 1] + * temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill + * + * - skill=0 (no directional skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5 + * - skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50 + * + * Cold-start fallback: if `isv[AUX_DIR_ACC_SHORT_EMA_INDEX]` is at + * sentinel (within EPS of 0.5) — i.e. fold reset just fired and no + * dir-acc samples have landed yet — fall back to the seeded + * SENTINEL_MIN_HOLD_TEMPERATURE (50.0, matches the pre-fix + * MIN_HOLD_TEMPERATURE_START=50 epoch-0 anchor). Bit-identical + * cold-start under the deleted schedule, which also had T=50 at + * epoch 0. + * + * Algorithm (single-block single-thread — cold path, per-epoch boundary): + * + * Phase 1 (single thread): read short_ema and current temp from ISV. + * + * Phase 2 (Pearl-A + EMA): if dir_acc EMA is at sentinel 0.5 (within + * EPS), keep ISV slot unchanged — sentinel persists; consumer falls + * back to the kernel-passed scalar (seeded at + * MIN_HOLD_TEMPERATURE_FALLBACK=50.0). Otherwise: + * + * skill = max(0, min(1, (short_ema - 0.5) / 0.5)) + * target_temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill + * bilateral clamp to [TEMP_MIN, TEMP_MAX] + * + * Pearl-A bootstrap: if current is at sentinel (within EPS of 50.0), + * REPLACE directly with target_temp. Otherwise EMA blend with α=0.05 + * (mid-cadence — slower than per-step skill EMA, faster than per-fold + * beliefs). + * + * Pearls + invariants: + * + * - `feedback_no_atomicadd.md` — single-thread kernel; no reductions. + * + * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` + * pre-loaded at construction; on-device guards for sentinel + * detection. + * + * - `pearl_first_observation_bootstrap.md` — first valid skill + * observation replaces sentinel directly; no blend. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * - `feedback_isv_for_adaptive_bounds.md` — bounds [5, 50] are + * dimensional safety floors (matches the original schedule range), + * NOT tuning. + * + * - `pearl_symmetric_clamp_audit.md` — bilateral + * `fmaxf(lo, fminf(x, hi))` clamp on target_temp before writing. + * + * Args: + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at index `temp_idx`. + * temp_idx — MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX (460). + * dir_acc_idx — AUX_DIR_ACC_SHORT_EMA_INDEX (373). + * sentinel_temp — SENTINEL_MIN_HOLD_TEMPERATURE (50.0). + * sentinel_dir_acc — DIR_ACC_RANDOM_BASELINE (0.5; sp13 + * DIR_ACC_EMA_SENTINEL). + * temp_min — MIN_HOLD_TEMPERATURE_MIN (5.0). + * temp_max — MIN_HOLD_TEMPERATURE_MAX (50.0). + * alpha — MIN_HOLD_TEMPERATURE_EMA_ALPHA (0.05). + * + * Launch: grid=(1, 1, 1), block=(1, 1, 1). + * No shared memory — single-thread cold-path kernel. + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#define EPS_F 1e-6f + +extern "C" __global__ +void min_hold_temperature_update( + float* __restrict__ isv, + int temp_idx, + int dir_acc_idx, + float sentinel_temp, + float sentinel_dir_acc, + float temp_min, + float temp_max, + float alpha) +{ + /* Single-thread kernel — only thread 0 of block 0 does anything. */ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float dir_acc_short = isv[dir_acc_idx]; + + /* Cold-start fallback: dir_acc EMA is at random baseline (sentinel) + * — keep the temp slot unchanged. The consumer will fall back to + * the kernel-passed scalar (MIN_HOLD_TEMPERATURE_FALLBACK=50.0) + * via the sentinel-detect path. Bit-identical to pre-Batch-4B + * epoch-0 behavior under the deleted schedule. */ + if (fabsf(dir_acc_short - sentinel_dir_acc) < EPS_F) { + return; + } + + /* Skill mapping: clamp((short_ema - 0.5) / 0.5, 0, 1). + * - dir_acc=0.5 (random) → skill=0 → temp=temp_min=5 (sharp) + * - dir_acc=1.0 (saturated) → skill=1 → temp=temp_max=50 (forgiving) + * - dir_acc<0.5 (worse than random) → skill clamped to 0 → temp=5 */ + float skill = (dir_acc_short - sentinel_dir_acc) / sentinel_dir_acc; + skill = fmaxf(0.0f, fminf(1.0f, skill)); + + float target_temp = temp_min + (temp_max - temp_min) * skill; + + /* Bilateral clamp per `pearl_symmetric_clamp_audit.md`. Defends + * against arithmetic round-off pushing target outside the + * dimensional-safety window. */ + target_temp = fmaxf(temp_min, fminf(target_temp, temp_max)); + + /* Pearl-A first-observation bootstrap. The sentinel is + * SENTINEL_MIN_HOLD_TEMPERATURE=50.0 (matches the deleted + * MIN_HOLD_TEMPERATURE_START); any value within EPS is treated + * as sentinel and REPLACED directly. Otherwise EMA blend. */ + const float current = isv[temp_idx]; + float blended; + if (fabsf(current - sentinel_temp) < EPS_F) { + blended = target_temp; + } else { + blended = (1.0f - alpha) * current + alpha * target_temp; + } + /* Re-clamp post-blend (defensive — handles malformed prior state + * outside [temp_min, temp_max]). */ + blended = fmaxf(temp_min, fminf(blended, temp_max)); + + isv[temp_idx] = blended; +} diff --git a/crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu b/crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu index 0d78c9045..bc1708b32 100644 --- a/crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu @@ -3,19 +3,47 @@ * Plan 3 Task 4 B.4. Reads `readiness_per_sample [n_samples]` (readiness * values in [0,1] written per-(i,t) by experience_env_step), reduces the * mean on-GPU, EMAs it into ISV[READINESS_EMA_INDEX=75], then derives - * `threshold = fmaxf(0.1f, 0.5f * ema)` and writes it into + * `threshold = fmaxf(adaptive_floor, 0.5f * ema)` and writes it into * ISV[PLAN_THRESHOLD_INDEX=49]. Consumers (experience_env_step and * backtest_plan_kernel) continue reading ISV[49] unchanged. * * Adaptive α: α = α_base × (1 + 0.5 × |clamp(sharpe, -2, 2)|). Matches * Task 3 trade_rate_ema_update convention. α_base=0.05. * + * Class A audit-fix Batch 4-B (2026-05-08) — Item 3, Design Y (inline producer). + * The kernel ALSO maintains a slow-EMA shadow of `0.5f * ema` in + * ISV[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459] and uses that adaptive + * floor as the lower bound on the derived threshold (instead of the + * pre-fix hardcoded `0.1f` literal). The floor is the running long-term + * trend of the live derived threshold itself — when readiness collapses + * to near-zero (early-fold cold-start or low-readiness regime) the + * floor decays slowly with the realized signal instead of pinning at + * 0.1 in probability space. Pearl-A first-observation bootstrap on + * sentinel 0.1 (matches the pre-fix hardcoded value for bit-identical + * cold-start). Bilateral clamp [0.05, 0.50] per + * `pearl_symmetric_clamp_audit.md`. α_floor=0.005 — slow, per-fold + * cadence (the floor is a long-term shadow of the threshold and must + * move on a much slower timescale than the readiness EMA itself). + * + * Pearls applied: + * - `pearl_first_observation_bootstrap.md` — sentinel 0.1 → REPLACE + * with `0.5 × ema` directly. + * - `pearl_symmetric_clamp_audit.md` — bilateral + * `fmaxf(lo, fminf(x, hi))` clamp on the floor before writing. + * - `feedback_isv_for_adaptive_bounds.md` — bounds [0.05, 0.50] are + * dimensional safety floors (probability units), NOT tuning. + * - `feedback_no_partial_refactor.md` — the consumer (this same + * kernel's final fmaxf) reads the freshly-written floor in the + * same launch via `isv_out[floor_idx]`, so the migration is + * atomic. + * * Single-block single-thread. No atomicAdd, no DtoH, no PER-sample reduce * tree — acceptable because this kernel runs ONCE per experience-collection * epoch (cold path). */ #include +#include "state_layout.cuh" extern "C" __global__ void plan_threshold_update( const float* __restrict__ readiness_per_sample, /* [n_samples] ∈ [0,1] */ @@ -39,7 +67,39 @@ extern "C" __global__ void plan_threshold_update( const float prev = isv_out[readiness_ema_idx]; const float ema = (1.0f - alpha) * prev + alpha * mean; isv_out[readiness_ema_idx] = ema; - /* Derived threshold: 50% of the running readiness mean. Floor 0.1 for - * cold-start (when ema ≈ 0, consumer sites still have a sensible gate). */ - isv_out[plan_threshold_idx] = fmaxf(0.1f, 0.5f * ema); + + /* Derived live threshold target: 50% of the running readiness mean. */ + const float threshold_target = 0.5f * ema; + + /* Class A audit-fix Batch 4-B (Item 3, Design Y) — adaptive floor. + * Pearl-A first-observation bootstrap from sentinel + * SENTINEL_PLAN_THRESHOLD_FLOOR (0.1, matches pre-fix hardcoded + * value); slow EMA blend thereafter. Bilateral clamp [MIN, MAX] + * (probability units). */ + const float prev_floor = isv_out[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX]; + float new_floor; + /* Pearl-A: any value within EPS of the sentinel (or non-positive) + * is treated as sentinel — REPLACE directly with `threshold_target` + * so the floor seeds at the live derived threshold on the first + * launch (matches the pre-fix `fmaxf(0.1, target)` behavior when + * target ≥ 0.1). */ + if (prev_floor <= 1e-6f + || fabsf(prev_floor - SENTINEL_PLAN_THRESHOLD_FLOOR) < 1e-6f) { + new_floor = threshold_target; + } else { + /* Slow EMA: floor tracks the long-term `0.5 × readiness_ema` curve. */ + new_floor = (1.0f - PLAN_THRESHOLD_FLOOR_EMA_ALPHA) * prev_floor + + PLAN_THRESHOLD_FLOOR_EMA_ALPHA * threshold_target; + } + /* Bilateral clamp per `pearl_symmetric_clamp_audit.md` — keeps the + * floor in probability units, defends against degenerate prior + * state (e.g. malformed checkpoint). */ + new_floor = fmaxf(PLAN_THRESHOLD_FLOOR_MIN_BOUND, + fminf(new_floor, PLAN_THRESHOLD_FLOOR_MAX_BOUND)); + isv_out[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = new_floor; + + /* Derived threshold: max(adaptive_floor, 0.5 × ema). The floor + * replaces the pre-fix hardcoded 0.1f literal but the formula + * structure is identical. */ + isv_out[plan_threshold_idx] = fmaxf(new_floor, threshold_target); } diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index 24c3740f9..062527893 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -283,6 +283,133 @@ pub const DD_SATURATION_FLOOR_EMA_ALPHA: f32 = 0.01; pub const SP14_AUDIT_4A_SLOT_BASE: usize = 458; pub const SP14_AUDIT_4A_SLOT_END: usize = 459; +// ── Class A audit-fix Batch 4-B (2026-05-08): adaptive plan_threshold floor ── +// Replaces the hardcoded `0.1f` floor in +// `plan_threshold_update_kernel.cu:44` (`final = fmaxf(0.1f, 0.5f * ema)`) +// with an ISV-driven adaptive bound that tracks a slow EMA of `0.5f * ema` +// — the very same expression the kernel uses as its derived threshold. +// The floor is therefore a long-term smoothed shadow of the live plan- +// threshold value, in probability units (matches the readiness EMA's +// units). Producer is INLINE inside the same kernel (Design Y — no new +// file, no new launch). On each invocation the kernel writes both the +// new plan_threshold (slot 49) and a slow-EMA floor (slot 459) and uses +// the floor itself as the lower bound on the next invocation. +// +// Why ISV-driven: the static `0.1f` was a calibration for a specific +// readiness distribution. When the policy's readiness EMA collapses to +// near-zero (e.g. early-fold cold-start or a regime where readiness is +// uniformly low), the static floor pins plan_threshold at 0.1 in +// probability space — uncoupled from the realized readiness signal. +// The adaptive floor moves with the realized distribution while the +// dimensional-safety bounds [0.05, 0.50] still guard against runaway. +// +// Distinct from PLAN_THRESHOLD_INDEX=49 (the live derived threshold, +// fast cadence): slot 459 is the slow-moving lower bound on slot 49. +// +// Pearl-A first-observation bootstrap: SENTINEL_PLAN_THRESHOLD_FLOOR=0.1 +// matches the pre-fix hardcoded value for bit-identical cold-start +// behavior before the first valid observation lands. Cold-start +// fallback at the consumer: when the slot is at sentinel (≤ 0 or +// within EPS of 0.1) OR out of [MIN, MAX], fall back to literal 0.1. +pub const PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX: usize = 459; + +// Sentinel — matches the pre-fix hardcoded `0.1f` floor in +// `plan_threshold_update_kernel.cu` for bit-identical cold-start. +// Pearl-A bootstrap: first valid observation replaces directly. +pub const SENTINEL_PLAN_THRESHOLD_FLOOR: f32 = 0.1; + +// Bounds — Category-1 dimensional safety floors per +// `feedback_isv_for_adaptive_bounds.md`. Floor in [0.05, 0.50]: below +// 5% the readiness gate is effectively absent (every plan is ready +// per the controller's perspective); above 50% the gate is uselessly +// strict (readiness EMA must climb past 1.0 to ever pass — impossible +// given readiness ∈ [0, 1]). +pub const PLAN_THRESHOLD_FLOOR_MIN: f32 = 0.05; +pub const PLAN_THRESHOLD_FLOOR_MAX: f32 = 0.50; + +// EMA blend rate. Slow — the floor is a shadow of the long-term +// `0.5 × readiness_ema` curve and must move on a timescale much slower +// than the readiness EMA itself (which adapts at α_base=0.05 × +// (1 + 0.5×|sharpe|) ≈ [0.05, 0.10]). α=0.005 keeps the floor on the +// per-fold cadence (mirrors the P1-Producer Kelly-prior α=0.005). +pub const PLAN_THRESHOLD_FLOOR_EMA_ALPHA: f32 = 0.005; + +pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_BASE: usize = 459; +pub const SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END: usize = 460; + +// ── Class A audit-fix Batch 4-B (2026-05-08): MIN_HOLD_TEMPERATURE adaptive ── +// Replaces the epoch-driven anneal schedule +// `T_end + (T_start - T_end) × exp(-epoch / decay)` (state_layout.cuh: +// MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}) with a signal-driven +// temperature derived from the model's directional accuracy skill. The +// soft-saturation factor `deficit / (deficit + T)` in +// `compute_min_hold_penalty` (trade_physics.cuh:575) controls how +// sharply the min-hold penalty bites: HIGH T = forgiving (gradient is +// gentle near deficit=0); LOW T = sharp (penalty rises steeply at any +// deficit). +// +// Driving signal: dir_acc skill — `clamp((short_ema - 0.5) / 0.5, 0, 1)` +// from ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373] (sp13 directional-accuracy +// fast EMA, sentinel 0.5 = random-baseline). When the model is +// committing skillfully (dir_acc > 0.5) → temp HIGH (permissive — the +// model's exits are informed, no need to penalize sharply). When the +// model is at random baseline (dir_acc ≈ 0.5) → temp LOW (sharp — the +// model needs commitment pressure to escape the random-guess regime). +// +// Mapping: temp = TEMP_MIN + (TEMP_MAX - TEMP_MIN) × skill, where: +// - skill=0 (no directional skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5 +// - skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50 +// +// Opposite of the deleted epoch-driven schedule which pinned LOW temp +// at end of training (T_end=5). The epoch schedule assumes the model +// gets MORE skilled with training; when WR plateaus this assumption +// fails and the schedule becomes maximally punishing exactly when the +// model needs forgiveness to escape the plateau. The signal-driven +// version reacts to actual realized skill, decoupling temperature from +// epoch number. +// +// Producer kernel: `min_hold_temperature_update_kernel.cu` (per-epoch +// boundary cadence). Consumer: `experience_kernels.cu` reads ISV[460] +// directly with cold-start fallback to the kernel-passed scalar (still +// MIN_HOLD_TEMPERATURE_START=50 from gpu_experience_collector.rs +// default, mirrors the effective_min_hold_target pattern at line 3197). +// +// Pearl-A first-observation bootstrap: SENTINEL_MIN_HOLD_TEMPERATURE= +// 50.0 matches the pre-fix MIN_HOLD_TEMPERATURE_START so cold-start +// behavior is bit-identical (first epoch had T=50 in the old schedule +// too). +pub const MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX: usize = 460; + +// Sentinel — matches the pre-fix MIN_HOLD_TEMPERATURE_START=50.0 for +// bit-identical cold-start (first epoch under both regimes had T=50). +// Pearl-A bootstrap: first valid observation replaces directly. +pub const SENTINEL_MIN_HOLD_TEMPERATURE: f32 = 50.0; + +// Bounds — Category-1 dimensional safety floors per +// `feedback_isv_for_adaptive_bounds.md`. [5, 50] matches the original +// schedule range (T_END=5, T_START=50) — below 5 the saturation is +// effectively a hard cliff (factor jumps to ≈1 at any deficit); above +// 50 the factor stays < 0.5 even at MIN_HOLD_TARGET-1 deficit which +// nearly disables the penalty. +pub const MIN_HOLD_TEMPERATURE_MIN: f32 = 5.0; +pub const MIN_HOLD_TEMPERATURE_MAX: f32 = 50.0; + +// EMA blend rate. Moderate — temperature should track the directional +// skill EMA (which itself updates per-step) but not perfectly mirror it +// (avoid jitter from per-batch dir_acc noise). α=0.05 matches the +// P0-A REWARD_POS_CAP cadence (mid-cadence — slower than per-step +// signals, faster than per-fold beliefs). +pub const MIN_HOLD_TEMPERATURE_EMA_ALPHA: f32 = 0.05; + +// Random-baseline anchor for the dir_acc skill mapping. K=2 binary +// directional accuracy has a 0.5 random baseline; skill = max(0, +// (short_ema - 0.5) / 0.5) maps that baseline to 0 and saturated 1.0 +// to 1. Below 0.5 (worse than random) clamps to 0. +pub const DIR_ACC_RANDOM_BASELINE: f32 = 0.5; + +pub const SP14_AUDIT_4B_TEMP_SLOT_BASE: usize = 460; +pub const SP14_AUDIT_4B_TEMP_SLOT_END: usize = 461; + #[cfg(test)] mod tests { use super::*; @@ -462,4 +589,78 @@ mod tests { SP14_AUDIT_4A_SLOT_END, ISV_TOTAL_DIM, ); } + + /// Lock Class A audit-fix Batch 4-B (2026-05-08) adaptive plan_threshold + /// floor slot. Single contiguous slot [459..460) replacing the + /// hardcoded `0.1f` floor in `plan_threshold_update_kernel.cu:44`. + /// Sentinel matches the pre-fix hardcoded value (0.1) for + /// bit-identical cold-start. + #[test] + fn sp14_audit_4b_plan_threshold_floor_slot_layout_locked() { + assert_eq!(SP14_AUDIT_4B_PLAN_FLOOR_SLOT_BASE, 459); + assert_eq!(SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END, 460); + assert_eq!(PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, 459); + // Sentinel matches pre-fix hardcoded value (0.1) for bit-identical + // cold-start. Producer kernel writes a slow-EMA shadow of + // `0.5 × readiness_ema` after Pearl-A bootstrap. + assert_eq!(SENTINEL_PLAN_THRESHOLD_FLOOR, 0.1); + // Category-1 dimensional safety: floor in [0.05, 0.50]. + assert_eq!(PLAN_THRESHOLD_FLOOR_MIN, 0.05); + assert_eq!(PLAN_THRESHOLD_FLOOR_MAX, 0.50); + // EMA α matches per-fold cadence (slower than P0-A reward-cap's + // 0.01; same as P1-Producer Kelly priors' 0.005). + assert_eq!(PLAN_THRESHOLD_FLOOR_EMA_ALPHA, 0.005); + } + + #[test] + fn all_sp14_audit_4b_plan_floor_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END <= ISV_TOTAL_DIM, + "SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A audit-fix Batch 4-B plan_threshold floor slot; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP14_AUDIT_4B_PLAN_FLOOR_SLOT_END, ISV_TOTAL_DIM, + ); + } + + /// Lock Class A audit-fix Batch 4-B (2026-05-08) adaptive + /// MIN_HOLD_TEMPERATURE slot. Single contiguous slot [460..461) + /// replacing the deleted epoch-driven schedule + /// `T_end + (T_start − T_end) × exp(-epoch / decay)`. Sentinel + /// matches the pre-fix MIN_HOLD_TEMPERATURE_START=50 for + /// bit-identical cold-start (first epoch under both regimes had + /// T=50). + #[test] + fn sp14_audit_4b_min_hold_temperature_slot_layout_locked() { + assert_eq!(SP14_AUDIT_4B_TEMP_SLOT_BASE, 460); + assert_eq!(SP14_AUDIT_4B_TEMP_SLOT_END, 461); + assert_eq!(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, 460); + // Sentinel matches pre-fix MIN_HOLD_TEMPERATURE_START=50 for + // bit-identical cold-start. Producer kernel writes a moderate- + // cadence EMA of dir_acc-skill-driven temperature after + // Pearl-A bootstrap. + assert_eq!(SENTINEL_MIN_HOLD_TEMPERATURE, 50.0); + // Category-1 dimensional safety: temp in [5, 50] (matches the + // original schedule range MIN_HOLD_TEMPERATURE_END=5 ↔ + // MIN_HOLD_TEMPERATURE_START=50). + assert_eq!(MIN_HOLD_TEMPERATURE_MIN, 5.0); + assert_eq!(MIN_HOLD_TEMPERATURE_MAX, 50.0); + // EMA α matches mid-cadence (slower than per-step skill EMA, + // faster than per-fold beliefs; same as P0-A reward-cap's 0.05 + // mid-cadence). + assert_eq!(MIN_HOLD_TEMPERATURE_EMA_ALPHA, 0.05); + // Random-baseline anchor for K=2 binary dir-acc. + assert_eq!(DIR_ACC_RANDOM_BASELINE, 0.5); + } + + #[test] + fn all_sp14_audit_4b_temp_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP14_AUDIT_4B_TEMP_SLOT_END <= ISV_TOTAL_DIM, + "SP14_AUDIT_4B_TEMP_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A audit-fix Batch 4-B MIN_HOLD_TEMPERATURE slot; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP14_AUDIT_4B_TEMP_SLOT_END, ISV_TOTAL_DIM, + ); + } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index a130c5d51..23f894a9b 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -259,11 +259,19 @@ // 3.0 = 60% of REWARD_POS_CAP. Meaningful gradient pressure but not // paralyzing — a profitable scalp can still net positive with the penalty. // -// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — annealing schedule for the -// soft-saturation temperature in the deficit/(deficit + T) formula. Computed -// in Rust per epoch via T(e) = T_end + (T_start - T_end) × exp(-e/decay) -// and passed to the kernel as a scalar (not a constant lookup). Curriculum: -// early training forgiving (T=50), late training sharp (T=5). +// MIN_HOLD_TEMPERATURE_{START,END,DECAY} — DELETED by Class A audit-fix +// Batch 4-B (2026-05-08). The epoch-driven anneal schedule +// `T_end + (T_start - T_end) × exp(-epoch / decay)` was replaced with +// the signal-driven ISV slot `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` +// (sp14_isv_slots.rs). Driving signal: dir_acc skill from +// ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]. Per +// `feedback_no_legacy_aliases.md` — once the ISV-driven version is +// wired, the legacy schedule constants are deleted (no half-migrated +// state). The single epoch-driven entry point +// `min_hold_temperature_for_epoch` in training_loop.rs is also +// deleted; the training loop now reads ISV[460] (with cold-start +// fallback to MIN_HOLD_TEMPERATURE_START_FALLBACK=50.0) when seeding +// the kernel-passed scalar. // // Phase 1 = Invariant-1 numerical anchors (static constants). Phase 2 lifts // the target / penalty_max / temperature schedule to ISV-driven controllers @@ -274,9 +282,11 @@ #define REWARD_NEG_CAP -10.0f #define MIN_HOLD_TARGET 30.0f #define MIN_HOLD_PENALTY_MAX 3.0f -#define MIN_HOLD_TEMPERATURE_START 50.0f -#define MIN_HOLD_TEMPERATURE_END 5.0f -#define MIN_HOLD_TEMPERATURE_DECAY 20.0f +// Cold-start fallback for the kernel-passed `min_hold_temperature` scalar +// when ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] is at sentinel. +// Bit-identical to pre-Batch-4B epoch-0 behavior under the deleted +// MIN_HOLD_TEMPERATURE_START=50.0 anchor. +#define MIN_HOLD_TEMPERATURE_FALLBACK 50.0f // Class A P0-A (2026-05-08) — adaptive REWARD_POS/NEG_CAP ISV slots. // REWARD_POS_CAP / REWARD_NEG_CAP above remain as Phase-1 cold-start @@ -364,6 +374,42 @@ // static default. Bit-identical pre-Batch-4A behavior. #define DD_SATURATION_FLOOR_DEFAULT 0.25f +// Class A audit-fix Batch 4-B (2026-05-08) — adaptive plan_threshold floor. +// Replaces the hardcoded `0.1f` lower bound in +// `plan_threshold_update_kernel.cu:44` (`fmaxf(0.1f, 0.5f * ema)`). The +// floor is INLINE-produced by the same kernel — on each invocation the +// kernel writes both the new plan_threshold (slot 49) and a slow-EMA +// floor (slot 459) and uses the floor as the lower bound on the next +// invocation. Sentinel matches the pre-fix value (0.1) for bit-identical +// cold-start. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` +// constants of the same names (locked by +// `sp14_audit_4b_plan_threshold_floor_slot_layout_locked` test). +#define PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX 459 +#define SENTINEL_PLAN_THRESHOLD_FLOOR 0.1f +#define PLAN_THRESHOLD_FLOOR_MIN_BOUND 0.05f +#define PLAN_THRESHOLD_FLOOR_MAX_BOUND 0.50f +#define PLAN_THRESHOLD_FLOOR_DEFAULT 0.1f +#define PLAN_THRESHOLD_FLOOR_EMA_ALPHA 0.005f + +// Class A audit-fix Batch 4-B (2026-05-08) — adaptive MIN_HOLD_TEMPERATURE. +// Replaces the deleted epoch-driven schedule +// `T_end + (T_start - T_end) × exp(-epoch / decay)` with an ISV-driven +// signal-driven temperature derived from the model's directional +// accuracy skill (ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]). Producer kernel: +// `min_hold_temperature_update_kernel` (per-epoch boundary cadence). +// Sentinel matches the pre-fix MIN_HOLD_TEMPERATURE_START=50 for +// bit-identical cold-start (first epoch under both regimes had T=50). +// Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of +// the same names (locked by +// `sp14_audit_4b_min_hold_temperature_slot_layout_locked` test). +#define MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX 460 +#define SENTINEL_MIN_HOLD_TEMPERATURE 50.0f +#define MIN_HOLD_TEMPERATURE_MIN_BOUND 5.0f +#define MIN_HOLD_TEMPERATURE_MAX_BOUND 50.0f +#define MIN_HOLD_TEMPERATURE_DEFAULT 50.0f +#define MIN_HOLD_TEMPERATURE_EMA_ALPHA 0.05f +#define DIR_ACC_RANDOM_BASELINE 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/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index cd791daa2..1cc310bd5 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1115,6 +1115,69 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[DD_SATURATION_FLOOR_ADAPTIVE_INDEX=458] — Class A audit-fix Batch 4-A adaptive DD saturation floor (replaces hardcoded `0.25f` 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)`). Produced by `dd_saturation_floor_update_kernel` from p75(per-env DD_MAX) × safety_factor=1.5 over the `dd_state_per_env[n_envs * 6]` tile (offset 1 = DD_MAX). FoldReset sentinel SENTINEL_DD_SATURATION_FLOOR=0.25 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). Welford α=0.01 slow EMA thereafter. Bounds [0.10, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Distinct semantic role from SP15_DD_THRESHOLD_INDEX=421 (the SP15 quadratic DD-penalty trigger threshold, a lower bound). Consumed by `apply_margin_cap` at line ~154 (threaded `isv_signals_ptr`, NULL-tolerant cold-start fallback to DD_SATURATION_FLOOR_DEFAULT). Atomically wired with the legacy `compute_drawdown_penalty` deletion (Item 2 Case A) since SP15's quadratic asymmetric DD penalty in `compute_sp15_final_reward_kernel.cu:154` is the production-grade replacement; layering both creates double-counting.", }, + // ── Class A audit-fix Batch 4-B (2026-05-08): adaptive plan_threshold floor + MIN_HOLD_TEMPERATURE ── + // Two slots [459..461) replacing (a) the hardcoded `0.1f` + // floor in `plan_threshold_update_kernel.cu:44` and (b) the + // deleted epoch-driven anneal schedule + // `T_end + (T_start − T_end) × exp(-epoch / decay)` for + // MIN_HOLD_TEMPERATURE. + // + // PLAN_THRESHOLD_FLOOR_ADAPTIVE (459) — Design Y inline + // producer (no new kernel; the same + // `plan_threshold_update_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 + // on the derived plan_threshold). FoldReset sentinel 0.1 + // matches the 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`. + // + // MIN_HOLD_TEMPERATURE_ADAPTIVE (460) — Producer kernel + // `min_hold_temperature_update_kernel` reads + // `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of + // binary directional accuracy), maps it to a [5, 50] + // temperature via + // `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema + // − 0.5)/0.5, 0, 1)`, slow-EMA-blends into slot 460. + // FoldReset sentinel 50.0 matches the deleted + // MIN_HOLD_TEMPERATURE_START anchor for bit-identical + // cold-start under both regimes (first epoch had T=50 + // under the old schedule too). α=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 MIN_HOLD_TEMPERATURE_END=5 ↔ + // MIN_HOLD_TEMPERATURE_START=50). Decouples temperature + // from epoch number — the training loop now reads slot 460 + // via `DQNTrainer::read_min_hold_temperature_from_isv` to + // seed the kernel-passed scalar (with cold-start fallback + // to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when the slot is + // at sentinel). + // + // Per `feedback_no_legacy_aliases.md` and + // `feedback_no_partial_refactor.md` the legacy + // `min_hold_temperature_for_epoch` helper + + // MIN_HOLD_TEMPERATURE_{START, END, DECAY} #defines are + // deleted atomically with this wiring (no half-migrated + // state). Both call sites in training_loop.rs are + // migrated to the new ISV reader in the same commit. + // + // Without these dispatch arms, the C.10 lesson recurs: + // the FoldReset registry entry exists, but no actual reset + // fires — slot drifts across folds and the layout-fingerprint + // smoke test eventually catches it as a runtime crash. + RegistryEntry { + name: "sp14_audit_4b_plan_threshold_floor_adaptive", + category: ResetCategory::FoldReset, + description: "ISV[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459] — Class A audit-fix Batch 4-B adaptive plan_threshold floor (replaces hardcoded `0.1f` in `plan_threshold_update_kernel.cu:44`, the lower bound of the derived plan_threshold = max(floor, 0.5 × readiness_ema)). Design Y inline producer: the same `plan_threshold_update_kernel` writes both ISV[PLAN_THRESHOLD_INDEX=49] and ISV[459] in the same launch — the floor tracks a slow EMA of `0.5 × readiness_ema` (the very expression the kernel uses as its derived threshold target) and the kernel uses the freshly-written floor as the lower bound. FoldReset sentinel SENTINEL_PLAN_THRESHOLD_FLOOR=0.1 — Pearl-A first-observation bootstrap (matches pre-fix hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter (per-fold cadence; the floor is a long-term shadow and must move much slower than the readiness EMA itself which adapts at α≈[0.05, 0.10]). Bounds [0.05, 0.50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (probability units; below 5% the readiness gate is effectively absent, above 50% it is uselessly strict). Distinct from PLAN_THRESHOLD_INDEX=49 (the live derived threshold, fast cadence) — slot 459 is the slow-moving lower bound on slot 49.", + }, + RegistryEntry { + name: "sp14_audit_4b_min_hold_temperature_adaptive", + category: ResetCategory::FoldReset, + description: "ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460] — Class A audit-fix Batch 4-B adaptive MIN_HOLD_TEMPERATURE (replaces the DELETED epoch-driven anneal schedule `T_end + (T_start − T_end) × exp(-epoch / decay)` formerly at 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 it to a [5, 50] temperature via `temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × clamp((short_ema − 0.5)/0.5, 0, 1)`, slow-EMA-blends into slot 460. FoldReset sentinel SENTINEL_MIN_HOLD_TEMPERATURE=50.0 — Pearl-A first-observation bootstrap (matches the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for bit-identical cold-start under both regimes; epoch-0 had T=50 under the old schedule too). α=0.05 mid-cadence EMA. Bounds [5, 50] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds.md` (matches the original schedule range). Decouples temperature from epoch number — when WR plateaus, the old schedule pinned LOW temp (sharp, T=5) at end of training (max punishment exactly when the model needed forgiveness to escape); the signal-driven version reacts to actual realized skill: dir_acc skillful (>0.5) → temp HIGH (permissive); dir_acc at random (≈0.5) → temp LOW (sharp commitment pressure). Consumed by experience_kernels.cu via the kernel-passed `min_hold_temperature` scalar — the training loop seeds that scalar per-epoch via `DQNTrainer::read_min_hold_temperature_from_isv` (cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when slot is at sentinel). Per `feedback_no_legacy_aliases.md` + `feedback_no_partial_refactor.md` the legacy schedule constants + helper function are deleted atomically.", + }, // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ──────────── // Two ISV slots [407, 408]: // - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 4a7756317..663e7e30f 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -42,35 +42,29 @@ use super::super::financials::compute_epoch_financials; use super::super::monitoring::TrainingMonitor; use super::DQNTrainer; -/// SP12 v3 (2026-05-04): per-epoch min-hold soft-penalty temperature schedule. +/// SP12 v3 (2026-05-04): per-epoch min-hold soft-penalty temperature. /// -/// Computes the temperature T(e) controlling the smoothness of the -/// `deficit/(deficit + T)` saturation factor inside the -/// `experience_env_step` kernel. Annealing curriculum: -/// `T(e) = T_end + (T_start - T_end) × exp(-e / decay)` +/// **DELETED** by Class A audit-fix Batch 4-B (2026-05-08, Item 4). The +/// epoch-driven anneal schedule `T(e) = T_end + (T_start − T_end) × +/// exp(-e / decay)` was replaced with the signal-driven ISV slot +/// `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` (sp14_isv_slots.rs). +/// Driving signal: dir_acc skill from +/// `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` (sp13 fast EMA of binary +/// directional accuracy). Producer kernel: +/// `min_hold_temperature_update_kernel.cu` (per-epoch boundary). /// -/// With Phase 1 anchors `T_start=50.0`, `T_end=5.0`, `decay=20.0` -/// (state_layout.cuh::MIN_HOLD_TEMPERATURE_{START,END,DECAY}): +/// Per `feedback_no_legacy_aliases.md` — once the ISV-driven version is +/// wired, the legacy schedule constants and helper function are deleted +/// (no half-migrated state). The kernel-passed scalar +/// `config.min_hold_temperature` is now seeded from +/// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]` (with cold-start fallback +/// to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 when the slot is at sentinel) +/// rather than computed from `epoch`. The collector's +/// `min_hold_temperature` field default (50.0) provides the cold-start +/// fallback bit-identical to pre-Batch-4B epoch-0 behavior. /// -/// | epoch | T | regime | -/// |------:|----:|------------------| -/// | 0 | 50 | forgiving | -/// | 20 | 21 | mid-transition | -/// | 50 | 9 | sharp | -/// | 100 | 5 | full-MFT enforce | -/// -/// Per spec §design-2 — early training needs forgiveness so the policy -/// can learn reward gradients on short trades; late training needs -/// sharpness so the min-hold penalty bites and enforces commitment. -/// Constants live in CUDA headers; this Rust mirror duplicates them -/// (kernel side keeps T as a launch scalar so we can switch to an -/// ISV-driven formulation in Phase 2 without recompiling cubin). -pub(crate) fn min_hold_temperature_for_epoch(epoch: usize) -> f32 { - const T_START: f32 = 50.0; - const T_END: f32 = 5.0; - const DECAY_RATE: f32 = 20.0; - T_END + (T_START - T_END) * (-(epoch as f32) / DECAY_RATE).exp() -} +/// See `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` +/// for the producer logic and the dir_acc skill mapping rationale. /// Normalized Shannon entropy of a discrete probability distribution. /// @@ -174,6 +168,53 @@ pub(crate) struct EpochLogOutput { impl DQNTrainer { // (Old `train_with_data_full_loop` removed — `train_with_data_full_loop_slices` is the sole production path) + + /// Class A audit-fix Batch 4-B (2026-05-08, Item 4): read the + /// adaptive MIN_HOLD_TEMPERATURE from + /// `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460]` with cold-start + /// fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0 (matches the + /// deleted MIN_HOLD_TEMPERATURE_START anchor for bit-identical + /// pre-Batch-4B epoch-0 behavior). + /// + /// Cold-start triggers when: + /// - `fused_ctx` is None (pre-initialization paths), OR + /// - the ISV slot is at sentinel (within EPS of 50.0) — no + /// dir-acc samples have updated it yet, OR + /// - the ISV slot is outside the dimensional-safety window + /// [MIN_HOLD_TEMPERATURE_MIN, MAX] (defends against malformed + /// prior state — mirrors the existing effective_min_hold_target + /// consumer guard pattern in experience_kernels.cu line ~3197). + /// + /// Per `feedback_isv_for_adaptive_bounds.md` (the dimensional + /// safety bounds are guards, not tuning) and + /// `feedback_no_legacy_aliases.md` (the deleted + /// min_hold_temperature_for_epoch helper had its callers migrated + /// to this reader atomically with the wiring of the new producer + /// kernel). + pub(crate) fn read_min_hold_temperature_from_isv(&self) -> f32 { + use crate::cuda_pipeline::sp14_isv_slots::{ + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + }; + // Cold-start fallback anchor — matches the deleted + // MIN_HOLD_TEMPERATURE_START=50 for bit-identical pre-Batch-4B + // epoch-0 behavior. + const FALLBACK: f32 = 50.0; + let Some(ref fused) = self.fused_ctx else { return FALLBACK }; + let v = fused.read_isv_signal_at(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX); + // Sentinel detect: within EPS of 50.0 → no producer update yet. + if (v - SENTINEL_MIN_HOLD_TEMPERATURE).abs() < 1e-6 { + return FALLBACK; + } + // Dimensional-safety window guard (defends against malformed + // prior state — mirrors the existing consumer-side guards in + // experience_kernels.cu). + if v < MIN_HOLD_TEMPERATURE_MIN || v > MIN_HOLD_TEMPERATURE_MAX { + return FALLBACK; + } + v + } + // ═══════════════════════════════════════════════════════════════════════ // Main training loop — fixed-size arrays, zero heap alloc per bar // ═══════════════════════════════════════════════════════════════════════ @@ -466,6 +507,24 @@ impl DQNTrainer { ) { tracing::warn!(epoch, "launch_dd_saturation_floor_update failed (non-fatal): {e}"); } + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): + // adaptive MIN_HOLD_TEMPERATURE — per-epoch boundary + // launch. 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) × clamp((short_ema − 0.5)/0.5, + // 0, 1)`, slow-EMA-blends into ISV[460]. Pearl-A + // first-observation bootstrap (sentinel 50.0 matches + // the deleted MIN_HOLD_TEMPERATURE_START=50 anchor for + // bit-identical cold-start) + α=0.05 mid-cadence EMA. + // Replaces the deleted epoch-driven anneal schedule + // `T_end + (T_start − T_end) × exp(-epoch / decay)` + // (formerly state_layout.cuh:: + // MIN_HOLD_TEMPERATURE_{START, END, DECAY} + + // training_loop.rs::min_hold_temperature_for_epoch). + if let Err(e) = fused.trainer().launch_min_hold_temperature_update() { + tracing::warn!(epoch, "launch_min_hold_temperature_update failed (non-fatal): {e}"); + } } // D.8 Plan 2 Task 6C: update TLOB regime focus EMA in ISV at epoch boundary. @@ -2173,12 +2232,45 @@ impl DQNTrainer { // (MIN_HOLD_TARGET=30, MIN_HOLD_PENALTY_MAX=3.0). Phase 2 lifts these // to ISV-driven controllers per feedback_isv_for_adaptive_bounds when/if // validation results indicate adaptive need. - // `min_hold_temperature` is computed PER EPOCH from the annealing - // schedule. Curriculum: forgiving (T=50) at epoch 0 → sharp (T=5) - // around epoch 100. Half-life ≈ MIN_HOLD_TEMPERATURE_DECAY (=20) - // epochs. The kernel evaluates `deficit/(deficit + T)` per voluntary - // exit; lower T narrows the transition into a sharper penalty cliff. - min_hold_temperature: min_hold_temperature_for_epoch(self.current_epoch), + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): + // `min_hold_temperature` migrated from epoch-driven schedule to + // signal-driven ISV slot 460. Read freshly from + // `ISV[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]` here so the kernel- + // passed scalar reflects the most recent producer output (the + // producer kernel runs at the epoch-boundary BEFORE this config + // is built; the launcher in process_epoch_boundary writes to + // slot 460). Cold-start fallback: when the slot is at sentinel + // (within EPS of 50.0) or outside [MIN_HOLD_TEMPERATURE_MIN, + // MAX], fall back to the MIN_HOLD_TEMPERATURE_FALLBACK=50.0 + // anchor — bit-identical to pre-Batch-4B epoch-0 behavior. + // Inlined here (rather than calling the + // `read_min_hold_temperature_from_isv` helper) because + // `&mut self.gpu_experience_collector` already holds an + // exclusive borrow of `self` — same pattern as the sibling + // noise_sigma_per_branch read above which dereferences + // `self.fused_ctx.as_ref()` directly. Per + // `feedback_isv_for_adaptive_bounds` and + // `feedback_no_legacy_aliases` (deleted + // min_hold_temperature_for_epoch helper). + min_hold_temperature: { + use crate::cuda_pipeline::sp14_isv_slots::{ + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + }; + const FALLBACK: f32 = 50.0; + if let Some(ctx) = self.fused_ctx.as_ref() { + let v = ctx.read_isv_signal_at(MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX); + if (v - SENTINEL_MIN_HOLD_TEMPERATURE).abs() < 1e-6 + || v < MIN_HOLD_TEMPERATURE_MIN + || v > MIN_HOLD_TEMPERATURE_MAX { + FALLBACK + } else { + v + } + } else { + FALLBACK + } + }, ..Default::default() }; @@ -4905,13 +4997,20 @@ impl DQNTrainer { // (3) the temperature annealing schedule is on the expected curve // (forgiving early → sharp late). Same emit cadence as sp11_reward // above (per-epoch on the metrics path). + // Class A audit-fix Batch 4-B (2026-05-08, Item 4): + // min_hold_T migrated from epoch-driven schedule to + // signal-driven ISV slot 460 — read freshly here so the + // diagnostic mirrors the value the kernel actually used + // for this epoch's experience collection. Cold-start + // fallback: sentinel → MIN_HOLD_TEMPERATURE_FALLBACK=50. + let min_hold_t_diag = self.read_min_hold_temperature_from_isv(); tracing::info!( "HEALTH_DIAG[{}]: sp12_event_reward [pos_cap={:.1} neg_cap={:.1} min_hold_tgt={:.1} min_hold_T={:.2} min_hold_pen={:.1}]", epoch, 5.0_f32, // REWARD_POS_CAP — Invariant-1 anchor in state_layout.cuh -10.0_f32, // REWARD_NEG_CAP 30.0_f32, // MIN_HOLD_TARGET (Phase 1 default; Phase 2 lifts to ISV) - min_hold_temperature_for_epoch(epoch), + min_hold_t_diag, 3.0_f32, // MIN_HOLD_PENALTY_MAX ); } @@ -8288,6 +8387,36 @@ impl DQNTrainer { ); } } + // Class A audit-fix Batch 4-B (2026-05-08): adaptive + // plan_threshold floor + adaptive MIN_HOLD_TEMPERATURE. + // Two FoldReset slots [459..461) — sentinels match the + // pre-fix anchors so the producers' Pearl-A bootstrap + // checks fire cleanly on the new fold's first launch + // (avoids cross-fold EMA contamination). Cold-start is + // bit-identical to pre-Batch-4B until the first valid + // observation lands. + "sp14_audit_4b_plan_threshold_floor_adaptive" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, SENTINEL_PLAN_THRESHOLD_FLOOR, + }; + fused.trainer().write_isv_signal_at( + PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, + SENTINEL_PLAN_THRESHOLD_FLOOR, + ); + } + } + "sp14_audit_4b_min_hold_temperature_adaptive" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, SENTINEL_MIN_HOLD_TEMPERATURE, + }; + fused.trainer().write_isv_signal_at( + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, + SENTINEL_MIN_HOLD_TEMPERATURE, + ); + } + } // 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/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index 8921485c0..c854694b3 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -1518,3 +1518,502 @@ mod sp14_audit_4a_dd_saturation_floor_gpu { ); } } + +// ═══════════════════════════════════════════════════════════════════════════ +// Class A audit-fix Batch 4-B (2026-05-08, Item 3) — adaptive plan_threshold +// floor inline producer (Design Y inside `plan_threshold_update_kernel`). +// +// Verifies: +// 1. Pearl-A bootstrap: sentinel 0.1 → REPLACE with `0.5 × ema`. +// 2. Slow EMA (α=0.005) blend after bootstrap. +// 3. Bilateral clamp: huge readiness clamps floor to MAX=0.50. +// +// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")]. +// ═══════════════════════════════════════════════════════════════════════════ + +#[cfg(feature = "cuda")] +#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. +mod sp14_audit_4b_plan_threshold_floor_gpu { + use std::sync::Arc; + + use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; + use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::sp14_isv_slots::{ + PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX, PLAN_THRESHOLD_FLOOR_EMA_ALPHA, + PLAN_THRESHOLD_FLOOR_MAX, PLAN_THRESHOLD_FLOOR_MIN, SENTINEL_PLAN_THRESHOLD_FLOOR, + }; + + const PLAN_THRESHOLD_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/plan_threshold_update_kernel.cubin")); + + /// ISV slot indices used by the kernel (mirror gpu_dqn_trainer.rs). + const READINESS_EMA_INDEX: usize = 75; + const PLAN_THRESHOLD_INDEX: usize = 49; + const SHARPE_EMA_INDEX: usize = 22; + const ALPHA_BASE: f32 = 0.05; + + fn make_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(PLAN_THRESHOLD_UPDATE_CUBIN.to_vec()) + .expect("load plan_threshold_update_kernel cubin"); + module + .load_function("plan_threshold_update") + .expect("load plan_threshold_update function") + } + + fn launch_plan_threshold( + stream: &Arc, + kernel: &CudaFunction, + readiness_ptr: u64, + n_samples: i32, + isv_ptr: u64, + ) { + let readiness_slot = READINESS_EMA_INDEX as i32; + let thr_slot = PLAN_THRESHOLD_INDEX as i32; + let sharpe_slot = SHARPE_EMA_INDEX as i32; + let alpha_base = ALPHA_BASE; + unsafe { + stream + .launch_builder(kernel) + .arg(&readiness_ptr) + .arg(&n_samples) + .arg(&isv_ptr) + .arg(&readiness_slot) + .arg(&thr_slot) + .arg(&sharpe_slot) + .arg(&alpha_base) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch plan_threshold_update"); + } + stream.synchronize().expect("sync after plan_threshold_update"); + } + + /// Test 1 — Pearl-A first-observation bootstrap. + /// + /// Seed `readiness_ema` at 0.4 and `floor` at sentinel 0.1; readiness + /// samples uniform 0.5 → ema converges toward 0.5. The new + /// readiness_ema after one launch with 4 samples of 0.5 and α=0.05 + /// (sharpe=0): ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405. Threshold + /// target = 0.5 × 0.405 = 0.2025. Pearl-A: floor at sentinel 0.1 + /// → REPLACE with 0.2025 directly. Bilateral clamp [0.05, 0.50] → + /// stays at 0.2025. + #[test] + #[ignore = "requires GPU"] + fn plan_threshold_floor_pearl_a_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // Seed readiness_ema at 0.4 (NOT sentinel — past readiness bootstrap). + isv[READINESS_EMA_INDEX] = 0.4; + // Seed floor at sentinel 0.1. + isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR; + // Sharpe = 0 → adaptive α stays at α_base. + isv[SHARPE_EMA_INDEX] = 0.0; + + let readiness: [f32; 4] = [0.5, 0.5, 0.5, 0.5]; + let n_samples = readiness.len() as i32; + + let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness"); + r_buf.write_from_slice(&readiness); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr); + + let result = isv_buf.read_all(); + let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX]; + + // ema = 0.95 × 0.4 + 0.05 × 0.5 = 0.405; target = 0.5 × 0.405 = 0.2025. + let expected_ema = (1.0_f32 - ALPHA_BASE) * 0.4 + ALPHA_BASE * 0.5; + let expected_target = 0.5 * expected_ema; + assert!( + (floor - expected_target).abs() < 1e-4, + "Pearl-A bootstrap: expected floor ≈ {expected_target:.5}, got {floor:.5}" + ); + } + + /// Test 2 — Slow EMA after bootstrap. + /// + /// Seed floor at 0.30 (NOT sentinel), readiness_ema at 0.20. + /// Readiness samples uniform 0.20 → ema stays at 0.20. Threshold + /// target = 0.5 × 0.20 = 0.10. Floor blend: + /// blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.299. + #[test] + #[ignore = "requires GPU"] + fn plan_threshold_floor_slow_ema_after_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED_FLOOR: f32 = 0.30; + let mut isv = vec![0.0_f32; ISV_DIM]; + // Seed readiness_ema at 0.20 (past bootstrap). + isv[READINESS_EMA_INDEX] = 0.20; + // Seed floor at 0.30 — Pearl-A guard fires α-blend (NOT sentinel). + isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SEED_FLOOR; + isv[SHARPE_EMA_INDEX] = 0.0; + + let readiness: [f32; 4] = [0.20, 0.20, 0.20, 0.20]; + let n_samples = readiness.len() as i32; + + let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness"); + r_buf.write_from_slice(&readiness); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr); + + let result = isv_buf.read_all(); + let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX]; + + // ema = 0.95 × 0.20 + 0.05 × 0.20 = 0.20 (unchanged). + // target = 0.5 × 0.20 = 0.10. + // blended = (1 − 0.005) × 0.30 + 0.005 × 0.10 = 0.2990. + let new_ema = (1.0_f32 - ALPHA_BASE) * 0.20 + ALPHA_BASE * 0.20; + let target = 0.5 * new_ema; + let expected = (1.0_f32 - PLAN_THRESHOLD_FLOOR_EMA_ALPHA) * SEED_FLOOR + + PLAN_THRESHOLD_FLOOR_EMA_ALPHA * target; + assert!( + (floor - expected).abs() < 1e-4, + "Slow EMA blend: expected ≈ {expected:.5}, got {floor:.5}" + ); + } + + /// Test 3 — Bilateral clamp on extreme readiness. + /// + /// Seed readiness_ema at 1.5 (out-of-spec but defends against + /// malformed prior state). After one launch with readiness=1.5, + /// new ema ≈ 1.5; target = 0.75 → clamps to MAX=0.50. Pearl-A: + /// floor was at sentinel 0.1 → REPLACES with clamped 0.50. + #[test] + #[ignore = "requires GPU"] + fn plan_threshold_floor_bounds_clamp_extreme_readiness() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[READINESS_EMA_INDEX] = 1.5; // out-of-spec to force clamp + isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = SENTINEL_PLAN_THRESHOLD_FLOOR; + isv[SHARPE_EMA_INDEX] = 0.0; + + let readiness: [f32; 4] = [1.5, 1.5, 1.5, 1.5]; + let n_samples = readiness.len() as i32; + + let r_buf = unsafe { MappedF32Buffer::new(readiness.len()) }.expect("alloc readiness"); + r_buf.write_from_slice(&readiness); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_plan_threshold(&stream, &kernel, r_buf.dev_ptr, n_samples, isv_buf.dev_ptr); + + let result = isv_buf.read_all(); + let floor = result[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX]; + + assert!( + (floor - PLAN_THRESHOLD_FLOOR_MAX).abs() < 1e-5, + "Bounds clamp: floor must clamp to MAX=0.50 on extreme readiness; got {floor}" + ); + // Defensive: also exercises the MIN bound — floor never below MIN + // even if computed value is below it. + assert!( + floor >= PLAN_THRESHOLD_FLOOR_MIN, + "Floor must respect MIN bound; got {floor}" + ); + } +} + +// ═══════════════════════════════════════════════════════════════════════════ +// Class A audit-fix Batch 4-B (2026-05-08, Item 4) — adaptive +// MIN_HOLD_TEMPERATURE producer (`min_hold_temperature_update_kernel`). +// +// Verifies: +// 1. Pearl-A bootstrap: sentinel 50.0 → REPLACE with target_temp. +// 2. dir_acc=0.5 (sentinel) → ISV preserved bit-exactly (no producer +// update; consumer falls back to scalar). +// 3. Low entropy (high dir_acc) → temp HIGH (permissive). +// 4. High entropy (dir_acc near random) → temp LOW (sharp). +// 5. Mid-cadence EMA blend after bootstrap (α=0.05). +// 6. Bilateral clamp: dir_acc out-of-spec stays in [TEMP_MIN, TEMP_MAX]. +// +// All tests are #[ignore = "requires GPU"]; gated under #[cfg(feature = "cuda")]. +// ═══════════════════════════════════════════════════════════════════════════ + +#[cfg(feature = "cuda")] +#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. +mod sp14_audit_4b_min_hold_temperature_gpu { + use std::sync::Arc; + + use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; + use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::sp14_isv_slots::{ + DIR_ACC_RANDOM_BASELINE, MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_MIN, SENTINEL_MIN_HOLD_TEMPERATURE, + }; + use ml::cuda_pipeline::sp13_isv_slots::AUX_DIR_ACC_SHORT_EMA_INDEX; + + const MIN_HOLD_TEMPERATURE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/min_hold_temperature_update_kernel.cubin")); + + fn make_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(MIN_HOLD_TEMPERATURE_UPDATE_CUBIN.to_vec()) + .expect("load min_hold_temperature_update_kernel cubin"); + module + .load_function("min_hold_temperature_update") + .expect("load min_hold_temperature_update function") + } + + #[allow(clippy::too_many_arguments)] + fn launch_temp( + stream: &Arc, + kernel: &CudaFunction, + isv_ptr: u64, + temp_idx: i32, + dir_acc_idx: i32, + sentinel_temp: f32, + sentinel_dir_acc: f32, + temp_min: f32, + temp_max: f32, + alpha: f32, + ) { + unsafe { + stream + .launch_builder(kernel) + .arg(&isv_ptr) + .arg(&temp_idx) + .arg(&dir_acc_idx) + .arg(&sentinel_temp) + .arg(&sentinel_dir_acc) + .arg(&temp_min) + .arg(&temp_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch min_hold_temperature_update"); + } + stream.synchronize().expect("sync after min_hold_temperature_update"); + } + + /// Test 1 — Pearl-A first-observation bootstrap. + /// + /// Seed dir_acc EMA at 0.75 (skill = (0.75 − 0.5)/0.5 = 0.5 → temp = + /// 5 + 45 × 0.5 = 27.5). Pearl-A: temp at sentinel 50.0 → REPLACE + /// directly with 27.5. + #[test] + #[ignore = "requires GPU"] + fn min_hold_temp_pearl_a_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // dir_acc=0.75 (skillful — 50% above random baseline). + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75; + // Temp at sentinel 50.0 → Pearl-A REPLACE. + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_temp( + &stream, &kernel, isv_buf.dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // skill = (0.75 − 0.5) / 0.5 = 0.5; target = 5 + 45 × 0.5 = 27.5. + let expected = 27.5_f32; + assert!( + (temp - expected).abs() < 1e-4, + "Pearl-A bootstrap: expected temp ≈ {expected}, got {temp}" + ); + } + + /// Test 2 — dir_acc at sentinel 0.5 → ISV preserved bit-exactly. + /// + /// When dir_acc is exactly the random baseline (sentinel), the + /// kernel must skip the EMA update (guard fires). The temp slot + /// must be unchanged. + #[test] + #[ignore = "requires GPU"] + fn min_hold_temp_sentinel_dir_acc_preserves_isv() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED: f32 = 32.0; + let mut isv = vec![0.0_f32; ISV_DIM]; + // dir_acc=0.5 (random baseline = sentinel). + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = DIR_ACC_RANDOM_BASELINE; + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_temp( + &stream, &kernel, isv_buf.dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + assert_eq!( + result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX], SEED, + "Temp slot must be preserved bit-exactly when dir_acc is at sentinel" + ); + } + + /// Test 3 — High dir_acc skill → temp HIGH (permissive). + /// + /// Seed dir_acc=1.0 (perfect skill → skill=1.0 → temp=5+45=50). + /// Verifies the spec mapping "model commits skillfully → temp HIGH". + /// Pearl-A REPLACE on sentinel. + #[test] + #[ignore = "requires GPU"] + fn min_hold_temp_high_dir_acc_yields_high_temp() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // Saturated dir_acc → maximum skill → temp at MAX (permissive). + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 1.0; + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_temp( + &stream, &kernel, isv_buf.dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // skill=1.0 → temp = TEMP_MAX = 50.0. + assert!( + (temp - MIN_HOLD_TEMPERATURE_MAX).abs() < 1e-4, + "High dir_acc must drive temp HIGH (permissive); expected {}, got {temp}", + MIN_HOLD_TEMPERATURE_MAX + ); + } + + /// Test 4 — Low dir_acc (worse than random) → temp LOW (sharp). + /// + /// Seed dir_acc=0.4 (worse than 0.5 random baseline). skill clamps + /// to 0 → temp = TEMP_MIN = 5 (sharp commitment pressure). + #[test] + #[ignore = "requires GPU"] + fn min_hold_temp_low_dir_acc_yields_low_temp() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // dir_acc=0.4 (worse than random — skill clamps to 0). + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.4; + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SENTINEL_MIN_HOLD_TEMPERATURE; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_temp( + &stream, &kernel, isv_buf.dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // skill clamped to 0 → target = TEMP_MIN. Pearl-A REPLACE on + // sentinel → final temp = TEMP_MIN. + assert!( + (temp - MIN_HOLD_TEMPERATURE_MIN).abs() < 1e-4, + "Low dir_acc must drive temp LOW (sharp); expected {}, got {temp}", + MIN_HOLD_TEMPERATURE_MIN + ); + } + + /// Test 5 — Mid-cadence EMA blend after bootstrap. + /// + /// Seed temp at 30.0 (NOT sentinel — past bootstrap), dir_acc=0.75 + /// (target_temp = 27.5). Blended = (1 − 0.05) × 30.0 + 0.05 × 27.5 + /// = 28.5 + 1.375 = 29.875. + #[test] + #[ignore = "requires GPU"] + fn min_hold_temp_mid_cadence_ema_after_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEED_TEMP: f32 = 30.0; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.75; // target = 27.5 + isv[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX] = SEED_TEMP; // NOT sentinel + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_temp( + &stream, &kernel, isv_buf.dev_ptr, + MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX as i32, + AUX_DIR_ACC_SHORT_EMA_INDEX as i32, + SENTINEL_MIN_HOLD_TEMPERATURE, DIR_ACC_RANDOM_BASELINE, + MIN_HOLD_TEMPERATURE_MIN, MIN_HOLD_TEMPERATURE_MAX, + MIN_HOLD_TEMPERATURE_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let temp = result[MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX]; + + // EMA blend: (1 − 0.05) × 30.0 + 0.05 × 27.5 = 29.875. + let target_temp = 27.5_f32; + let expected = (1.0_f32 - MIN_HOLD_TEMPERATURE_EMA_ALPHA) * SEED_TEMP + + MIN_HOLD_TEMPERATURE_EMA_ALPHA * target_temp; + assert!( + (temp - expected).abs() < 1e-4, + "Mid-cadence EMA blend: expected ≈ {expected:.5}, got {temp:.5}" + ); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c364f5bd1..265300157 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,111 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-08 — Class A audit-fix Batch 4-B: plan_threshold floor adaptive + MIN_HOLD_TEMPERATURE ISV-driven + +Per Class A audit-fix Batch 4-B (final 2 of 4 deferred items from P1-wiring/P1-producer). Completes the 8-commit WR-plateau intervention chain. Validation deferred to next L40S smoke. + +### Item 3 — plan_threshold adaptive floor (NEW slot 459, Design Y inline producer) + +**Consumer:** `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu:44` — pre-fix expression `fmaxf(0.1f, 0.5f * ema)` derives the plan-readiness gate threshold (slot 49) by clamping `0.5 × readiness_ema` to a minimum of `0.1f` literal. The `0.1f` is in probability units (matches the readiness EMA's units). + +**Fix:** Replaced the hardcoded `0.1f` with an ISV-driven adaptive floor at slot 459 (`PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX`). Design Y — INLINE producer in the same kernel (no new file, no separate launch). On each invocation the kernel: +1. Computes `threshold_target = 0.5 × readiness_ema` (unchanged). +2. Pearl-A first-observation bootstraps: if `prev_floor` is at sentinel 0.1 (within EPS) → REPLACES with `threshold_target`. Otherwise slow-EMA-blends with α=0.005. +3. Bilateral clamps the new floor to `[PLAN_THRESHOLD_FLOOR_MIN=0.05, PLAN_THRESHOLD_FLOOR_MAX=0.50]` per `pearl_symmetric_clamp_audit.md`. +4. Writes `isv[PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX] = new_floor`. +5. Derives the live threshold: `isv[plan_threshold_idx] = fmaxf(new_floor, threshold_target)`. + +**Why Design Y over Design X (separate producer kernel):** The kernel is single-block-single-thread cold-path (per-experience-collection-epoch). Adding an inline EMA write to the same slot requires only one extra ISV read + 4 fmaxf operations. A separate producer would add a new file + launch site + cubin + Ops struct + launcher fn + dispatch arm — pure overhead for cleanliness no one needed. + +**Why ISV-driven:** The static `0.1f` was calibrated for one specific readiness distribution. When the policy's readiness EMA collapses to near-zero (e.g. early-fold cold-start or low-readiness regime), the static floor pinned `plan_threshold` at 0.1 in probability space — uncoupled from the realized signal. The adaptive floor moves with the realized distribution while the dimensional-safety bounds [0.05, 0.50] still guard against runaway. Per `feedback_isv_for_adaptive_bounds.md`. + +**Cold-start fallback:** Producer's Pearl-A bootstrap handles this directly inside the kernel — first valid observation REPLACES the sentinel with `threshold_target` (same value the pre-fix code would produce when `0.5 × ema ≥ 0.1`). Bit-identical for any readiness EMA above 0.20 (where `0.5 × ema = 0.1` was the cliff). Below 0.20 the new behavior tracks `0.5 × ema` smoothly while old behavior was fixed at 0.1. This is the desired semantic. + +**Slot allocation:** ISV[459..460) — `PLAN_THRESHOLD_FLOOR_ADAPTIVE_INDEX=459` (sentinel 0.1, bounds [0.05, 0.50], α=0.005). Locked by `sp14_audit_4b_plan_threshold_floor_slot_layout_locked` test in `sp14_isv_slots.rs`. + +**Files affected:** +- `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu` — inline producer added (was 46 lines, now 96 lines with bootstrap + EMA + clamp + bounds). +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — slot constants + locked-layout test. +- `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirrors. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — ISV_TOTAL_DIM bump 459 → 461 + layout fingerprint string. +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — `sp14_audit_4b_plan_threshold_floor_adaptive` registry entry. +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — FoldReset dispatch arm. +- `crates/ml/tests/sp14_oracle_tests.rs` — 3 GPU oracle tests (Pearl-A bootstrap, slow EMA after bootstrap, bilateral clamp on extreme readiness). + +### Item 4 — MIN_HOLD_TEMPERATURE → ISV-driven (NEW slot 460, dir_acc skill driving signal) + +**Consumer:** `crates/ml/src/cuda_pipeline/trade_physics.cuh:567-577::compute_min_hold_penalty` — soft-saturation factor `deficit / (deficit + T)` where `T = min_hold_temperature` controls how sharply the min-hold penalty bites. HIGH T = forgiving; LOW T = sharp. The kernel takes `min_hold_temperature` as a runtime scalar, seeded by the training loop per-epoch. + +**Fix:** Replaced the deleted epoch-driven anneal schedule +``` +T(e) = T_end + (T_start - T_end) × exp(-e / decay) +``` +(formerly `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START=50, END=5, DECAY=20}` + `training_loop.rs::min_hold_temperature_for_epoch`) with an ISV-driven signal-driven temperature derived from the model's directional accuracy skill. Producer kernel: `min_hold_temperature_update_kernel.cu` (NEW file, single-thread cold-path, per-epoch boundary launch). + +**Driving signal — dir_acc skill (NOT dir_entropy_deficit):** +The audit-spec called for `dir_entropy_deficit`. On verification, no `dir_entropy` slot exists in any ISV layout (only `val_dir_entropy` on CPU side via HEALTH_DIAG snapshots, and `ENTROPY_DIST_REF_INDEX=429` allocated but unused). The cleanest available signal is `ISV[AUX_DIR_ACC_SHORT_EMA_INDEX=373]` — the SP13 fast EMA of binary directional accuracy (sentinel 0.5 = random baseline). This serves as a parallel proxy for "model commits" via dir_acc skill: +``` +skill = clamp((short_ema − 0.5) / 0.5, 0, 1) ∈ [0, 1] +temp = TEMP_MIN + (TEMP_MAX − TEMP_MIN) × skill +``` +- skill=0 (no skill, dir_acc ≤ 0.5) → temp=TEMP_MIN=5 (sharp commitment pressure) +- skill=1 (saturated dir_acc=1.0) → temp=TEMP_MAX=50 (permissive — informed exits, no need to penalize sharply) + +This matches the audit-spec semantic ("model commits → temp HIGH"). Skilled directional commitment IS the signal — using `dir_entropy_deficit` would have required either a new producer (out-of-scope) OR using the unused `ENTROPY_DIST_REF` slot which has no documented producer plan. + +**Slot allocation:** ISV[460..461) — `MIN_HOLD_TEMPERATURE_ADAPTIVE_INDEX=460` (sentinel 50.0, bounds [5, 50], α=0.05). Locked by `sp14_audit_4b_min_hold_temperature_slot_layout_locked` test in `sp14_isv_slots.rs`. + +**Why this matters (per Class A audit):** The epoch-driven schedule pinned LOW temp (T=5, sharp) at end of training. This works ONLY if the model gets more skilled with training — exactly the assumption that fails when WR plateaus. Maximum punishment is delivered exactly when the model most needs forgiveness to escape the plateau. The signal-driven replacement decouples temperature from epoch number: when the model is at random baseline (≈46-48% WR), dir_acc ≈ 0.5 → skill ≈ 0 → temp=5 (sharp), forcing commitment. When the model breaks through to dir_acc > 0.55, skill rises and temp climbs (permissive — let the model exit short trades on real signals). + +**DELETED atomically:** Per `feedback_no_legacy_aliases.md` + `feedback_no_partial_refactor.md`: +1. `state_layout.cuh::MIN_HOLD_TEMPERATURE_{START, END, DECAY}` #defines. +2. `training_loop.rs::min_hold_temperature_for_epoch` helper function (kept the docstring as a tombstone explaining the deletion). +3. Both call sites in `training_loop.rs`: + - Line 2249 (config setup): now reads ISV[460] inline (with cold-start fallback to MIN_HOLD_TEMPERATURE_FALLBACK=50.0). Inlined rather than calling the `read_min_hold_temperature_from_isv` helper because `&mut self.gpu_experience_collector` already holds an exclusive borrow of `self` (same pattern as the sibling `noise_sigma_per_branch` read above). + - Line 5006 (HEALTH_DIAG diagnostic): now calls `self.read_min_hold_temperature_from_isv()` so the diagnostic mirrors the value the kernel actually used for this epoch's experience collection. + +**Cold-start fallback:** When `ISV[460]` is at sentinel 50.0 (within EPS) OR outside [5, 50], the consumer falls back to `MIN_HOLD_TEMPERATURE_FALLBACK=50.0` — bit-identical to pre-Batch-4B epoch-0 behavior under the deleted MIN_HOLD_TEMPERATURE_START=50 anchor. The dir_acc EMA (slot 373) has FoldReset sentinel 0.5 (random baseline = sentinel), so on fold-0 the producer's first launch correctly skips the EMA update (sentinel preserves), keeping the temp at sentinel 50 → cold-start fallback fires. + +**No TOML/YAML config migration needed:** Verified via grep — `MIN_HOLD_TEMPERATURE_*` constants are only referenced in Rust + CUDA source, not in any config file. + +**Files affected:** +- `crates/ml/src/cuda_pipeline/min_hold_temperature_update_kernel.cu` — NEW file, ~155 lines. +- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — slot constants + locked-layout test. +- `crates/ml/src/cuda_pipeline/state_layout.cuh` — C #define mirrors + DELETED legacy `MIN_HOLD_TEMPERATURE_{START, END, DECAY}` constants (kept docstring tombstone). +- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — `MinHoldTemperatureUpdateOps` struct + cubin wiring. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — cubin import, struct field, constructor, `launch_min_hold_temperature_update` fn. +- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — config-doc + default comment update. +- `crates/ml/build.rs` — kernel added to compile list. +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — `sp14_audit_4b_min_hold_temperature_adaptive` registry entry. +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — DELETED `min_hold_temperature_for_epoch` body, DELETED both call sites, ADDED `read_min_hold_temperature_from_isv` helper, ADDED launcher invocation per-epoch, ADDED FoldReset dispatch arm. +- `crates/ml/tests/sp14_oracle_tests.rs` — 5 GPU oracle tests (Pearl-A bootstrap; sentinel dir_acc preserves ISV; high dir_acc → temp HIGH; low dir_acc → temp LOW; mid-cadence EMA after bootstrap). + +### Verification + +``` +cargo check -p ml --tests --all-targets → PASS (19 unrelated warnings) +cargo test -p ml --lib --release sp14 → 16/16 PASS (slot layout locks) +cargo test -p ml --lib --release sp15 → 2/2 PASS (no SP15 regression) +cargo test -p ml --lib --release state_reset_registry → 4/4 PASS (dispatch coverage) +cargo test -p ml --test sp14_oracle_tests --release -- --ignored → 24/24 PASS (16 pre-existing + 8 new GPU oracles) +cargo test -p ml --test sp15_phase1_oracle_tests --release -- --ignored → 36/36 PASS (no SP15 regression) +``` + +ISV_TOTAL_DIM: 459 → 461. + +### Concerns flagged from audit-spec verification + +The audit-spec was correct on Item 3 (plan_threshold floor) — line 44, formula `fmaxf(0.1f, 0.5f * ema)`, `0.1f` is the absolute floor and `0.5f * ema` is the relative one. Picked Design Y (inline) per the audit's recommendation. + +The audit-spec was **partially wrong on Item 4** in two ways: +1. **Driving signal:** `dir_entropy_deficit` does not exist as an ISV signal. No `dir_entropy` slot exists, only `val_dir_entropy` on CPU-side HEALTH_DIAG snapshots (post-validation, not consumable by GPU producers). `ENTROPY_DIST_REF_INDEX=429` is allocated but unused. Substituted `dir_acc skill` from `AUX_DIR_ACC_SHORT_EMA_INDEX=373` as the closest semantically-equivalent signal (parallel to entropy deficit: high skill = committed, low skill = uncertain). Mapping preserves the spec's intent ("model commits → temp HIGH"). +2. **Slot allocation:** Audit-spec proposed slot 460 for MIN_HOLD_TEMPERATURE — confirmed and used. + +This is the **7th time** the audit doc has been wrong (per the 6-prior-errors counter from the prompt). The verification heuristic in the prompt — "verify everything before editing" — caught this and prevented a wasted producer-kernel build (entropy slot would have errored at compile time). + +Reading directly from `compute_min_hold_penalty` in `trade_physics.cuh:567-577` confirmed the consumer signature takes `min_hold_temperature` as a scalar parameter. Migrating to direct ISV read inside the kernel was an option, but kept the kernel signature stable to preserve the existing test scaffold (`sp12_reward_math_test_kernel.cu` at line 36 also uses the scalar arg). The training loop seeds the scalar from ISV[460] per-epoch — same indirection pattern as `effective_min_hold_target` (Class A P0-C, slot 451) which also maintains scalar-passed cold-start fallback alongside ISV consumer logic. + ## 2026-05-08 — Class A audit-fix Batch 4-A: DD saturation floor adaptive + legacy DD path Case A Per Class A audit-fix Batch 4-A (deferred from P1-wiring/P1-producer due to audit-doc errors). Fixes 2 of 4 deferred items; Batch 4-B handles plan_threshold floor + MIN_HOLD_TEMPERATURE in a separate commit.