diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 5d53fbde9..c40e94105 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -806,6 +806,27 @@ fn main() { // WR-stuck-at-46-48% plateau across 11 superprojects. Per-epoch // boundary launch. "reward_cap_update_kernel.cu", + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors + // producer. Single-block kernel that sweeps the per-fold-end + // `portfolio_state[n_envs, PS_STRIDE]` buffer (same source the + // `kelly_cap_update_kernel` reads from at the same boundary), + // aggregates the realized PS_KELLY_{WIN_COUNT, LOSS_COUNT, + // SUM_WINS, SUM_LOSSES} fields across envs, and slow-EMA-blends + // into ISV[454..458). Pearl-A first-observation 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 + // thereafter (per-fold cadence — the Bayesian prior is a *prior + // belief* and should change slowly across folds; a fast + // "prior" would just be a noisy Kelly estimate). Bounds + // counts∈[0.5, 100], sums∈[0.001, 1.0] are Category-1 dimensional + // safety floors per `feedback_isv_for_adaptive_bounds`. Replaces + // hardcoded `prior_wins=2.0f / prior_losses=2.0f / + // prior_sum_wins=0.01f / prior_sum_losses=0.01f` in + // kelly_cap_update_kernel.cu:39-42 + trade_physics.cuh:: + // kelly_position_cap:304-307 — Class A audit P1-Producer batch. + // Block-tree-reduce (no atomicAdd) per + // `feedback_no_atomicadd.md`. Per-epoch boundary launch. + "kelly_bayesian_priors_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 816b8fd72..27de97a95 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs @@ -52,7 +52,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, H_S2_AUX_RMS_EMA_CUBIN, REWARD_CAP_UPDATE_CUBIN, + AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN, KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN, + REWARD_CAP_UPDATE_CUBIN, }; /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU @@ -734,6 +735,128 @@ impl RewardCapUpdateOps { } } +/// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer. +/// +/// Replaces the previously-hardcoded `prior_wins=2.0f, prior_losses=2.0f, +/// prior_sum_wins=0.01f, prior_sum_losses=0.01f` constants used in +/// `kelly_cap_update_kernel.cu` and `trade_physics.cuh::kelly_position_cap` +/// with ISV-driven slow-EMA values fed by the +/// `kelly_bayesian_priors_update_kernel` producer at per-fold-end (per-epoch +/// boundary in the current scheduling). Aggregates the realized +/// PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, SUM_LOSSES} fields across all +/// envs from the same `portfolio_state[n_envs, PS_STRIDE]` buffer that +/// `kelly_cap_update_kernel` reads from at the same boundary, then +/// slow-EMA-blends into ISV[KELLY_PRIOR_*_INDEX] = ISV[454..458). +/// +/// Single-block 256-thread kernel; shmem block-tree-reduce +/// (no atomicAdd per `feedback_no_atomicadd.md`). `CudaFunction` pre-loaded +/// at construction per `pearl_no_host_branches_in_captured_graph.md`. +/// +/// Pearls applied: +/// - `feedback_no_atomicadd.md` — block-tree-reduce in shmem, single +/// global write per slot from thread 0. +/// - `pearl_first_observation_bootstrap.md` — sentinel match (within 1e-6 +/// of pre-P1-Producer hardcoded value) → REPLACE; otherwise EMA blend. +/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` +/// pre-loaded at construction; on-device guard for "no realized trades". +/// - `feedback_no_stubs.md` — full body, no placeholder. +/// - `feedback_isv_for_adaptive_bounds.md` — count bounds [0.5, 100], +/// sum bounds [0.001, 1.0] are dimensional safety floors, NOT tuning. +/// - `pearl_symmetric_clamp_audit.md` — bilateral +/// `fmaxf(lo, fminf(x, hi))` clamp on each slot before writing. +/// - `pearl_controller_anchors_isv_driven.md` — every controller anchor +/// (priors are anchors for the Kelly cap formula) is ISV-driven. +#[allow(missing_debug_implementations)] +pub(crate) struct KellyBayesianPriorsUpdateOps { + update_kernel: CudaFunction, +} + +impl KellyBayesianPriorsUpdateOps { + /// Block dim used by the producer kernel — must match `BLK_DIM` in + /// `kelly_bayesian_priors_update_kernel.cu`. + const BLK_DIM: u32 = 256; + + pub(crate) fn new(stream: &Arc) -> Result { + let context = stream.context(); + let module = context + .load_cubin(KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_update cubin load: {e}")))?; + let update_kernel = module + .load_function("kelly_bayesian_priors_update") + .map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_update load: {e}")))?; + Ok(Self { update_kernel }) + } + + /// Launch the adaptive Bayesian-priors producer. + /// + /// Args: + /// - `portfolio_state_ptr`: f32 device ptr `[n_envs * ps_stride]` — + /// the same buffer `kelly_cap_update_kernel` reads from. We use + /// the four PS_KELLY_* fields per env. + /// - `n_envs`: number of envs (alloc_episodes from the collector). + /// - `ps_stride`: PS_STRIDE from state_layout.cuh (43 post-Plan-3-D.4c). + /// - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned). + /// - `wins_idx`/`losses_idx`/`sum_wins_idx`/`sum_losses_idx`: + /// KELLY_PRIOR_*_INDEX (454..458). + /// - `sentinel_*`: SENTINEL_KELLY_PRIOR_* (2.0/2.0/0.01/0.01). + /// - `count_min`/`count_max`: KELLY_PRIOR_COUNT_MIN/MAX (0.5, 100). + /// - `sum_min`/`sum_max`: KELLY_PRIOR_SUM_MIN/MAX (0.001, 1.0). + /// - `alpha`: KELLY_PRIOR_EMA_ALPHA (0.005, slow per-fold cadence). + #[allow(clippy::too_many_arguments)] + pub(crate) fn launch( + &self, + stream: &Arc, + portfolio_state_ptr: u64, + n_envs: i32, + ps_stride: i32, + isv_ptr: u64, + wins_idx: i32, + losses_idx: i32, + sum_wins_idx: i32, + sum_losses_idx: i32, + sentinel_wins: f32, + sentinel_losses: f32, + sentinel_sum_wins: f32, + sentinel_sum_losses: f32, + count_min: f32, + count_max: f32, + sum_min: f32, + sum_max: f32, + alpha: f32, + ) -> Result<(), MLError> { + // 4 arrays × BLK_DIM × sizeof(f32) — see the kernel comment block. + let smem_bytes = 4 * Self::BLK_DIM * std::mem::size_of::() as u32; + unsafe { + stream + .launch_builder(&self.update_kernel) + .arg(&portfolio_state_ptr) + .arg(&n_envs) + .arg(&ps_stride) + .arg(&isv_ptr) + .arg(&wins_idx) + .arg(&losses_idx) + .arg(&sum_wins_idx) + .arg(&sum_losses_idx) + .arg(&sentinel_wins) + .arg(&sentinel_losses) + .arg(&sentinel_sum_wins) + .arg(&sentinel_sum_losses) + .arg(&count_min) + .arg(&count_max) + .arg(&sum_min) + .arg(&sum_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (Self::BLK_DIM, 1, 1), + shared_mem_bytes: smem_bytes, + }) + .map_err(|e| MLError::ModelError(format!("kelly_bayesian_priors_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 4bb49025e..4e65a42b6 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2110,6 +2110,20 @@ pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat! /// `gpu_aux_trunk::RewardCapUpdateOps`. Per-epoch boundary launch. pub(crate) static REWARD_CAP_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_cap_update_kernel.cubin")); +/// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer cubin. +/// Single-block 256-thread kernel that sweeps the per-fold-end +/// `portfolio_state[n_envs, PS_STRIDE]` buffer (same source the +/// `kelly_cap_update_kernel` reads from at the same boundary), aggregates +/// the realized PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, SUM_LOSSES} +/// fields across envs, and slow-EMA-blends into ISV[454..458). Pearl-A +/// first-observation 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 thereafter. Loaded by +/// `gpu_aux_trunk::KellyBayesianPriorsUpdateOps`. Per-epoch boundary launch +/// (RIGHT BEFORE `kelly_cap_update` so that kernel sees the freshly-blended +/// priors). +pub(crate) static KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/kelly_bayesian_priors_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 @@ -2438,7 +2452,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 = 454; // 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. 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 = 458; // 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. 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). @@ -3512,7 +3526,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { AUX_TRUNK_EPS=447;AUX_TRUNK_GRAD_CLIP=448;H_S2_AUX_RMS_EMA=449;\ AUX_PRED_HORIZON_BARS=450;AVG_WIN_HOLD_TIME_BARS=451;\ REWARD_POS_CAP_ADAPTIVE=452;REWARD_NEG_CAP_ADAPTIVE=453;\ - ISV_TOTAL_DIM=454;\ + KELLY_PRIOR_WINS=454;KELLY_PRIOR_LOSSES=455;KELLY_PRIOR_SUM_WINS=456;KELLY_PRIOR_SUM_LOSSES=457;\ + ISV_TOTAL_DIM=458;\ 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;\ @@ -7525,6 +7540,20 @@ pub struct GpuDqnTrainer { /// foundation of training and shouldn't move fast. reward_cap_update_ops: super::gpu_aux_trunk::RewardCapUpdateOps, + /// Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors + /// producer. Sweeps `portfolio_state[n_envs, PS_STRIDE]` (same source + /// `kelly_cap_update_kernel` reads), aggregates the realized + /// PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, SUM_LOSSES} fields + /// across envs, and slow-EMA-blends into ISV[454..458). Pearl-A + /// first-observation bootstrap (sentinels 2.0/2.0/0.01/0.01); α=0.005 + /// slow EMA. Replaces hardcoded `prior_wins/prior_losses/ + /// prior_sum_wins/prior_sum_losses` constants in + /// `kelly_cap_update_kernel.cu` and + /// `trade_physics.cuh::kelly_position_cap`. Per-epoch boundary launch + /// (RIGHT BEFORE `kelly_cap_update` so that kernel sees the + /// freshly-blended priors). + kelly_bayesian_priors_update_ops: super::gpu_aux_trunk::KellyBayesianPriorsUpdateOps, + // ── Speculative inference cache ── /// Cached trunk output [B, SH2] from between-bar speculative forward. speculative_h_s2: CudaSlice, @@ -14643,6 +14672,72 @@ impl GpuDqnTrainer { Ok(()) } + /// Class A P1-Producer (2026-05-08): launch the adaptive Bayesian Kelly + /// priors producer. Per-fold-end (per-epoch boundary in current + /// scheduling) launch — Bayesian prior is a *prior belief* and should + /// change slowly across folds (α=0.005 EMA). + /// + /// Aggregates realized PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, + /// SUM_LOSSES} across envs from the same `portfolio_state` buffer + /// `kelly_cap_update_kernel` reads from at the same boundary and + /// slow-EMA-blends into ISV[454..458). MUST be called RIGHT BEFORE + /// `launch_kelly_cap_update` so that kernel sees the freshly-blended + /// priors. + /// + /// Pearls applied: + /// - `pearl_first_observation_bootstrap.md` (sentinels 2.0/2.0/0.01/0.01) + /// - `pearl_no_host_branches_in_captured_graph.md` + /// - `feedback_no_atomicadd.md` (block-tree-reduce in shmem only) + /// - `pearl_symmetric_clamp_audit.md` (bilateral clamp) + /// - `feedback_isv_for_adaptive_bounds.md` (count [0.5, 100], sum + /// [0.001, 1.0] are dimensional safety floors, not tuning) + /// - `pearl_controller_anchors_isv_driven.md` (priors are + /// controller anchors) + pub(crate) fn launch_kelly_bayesian_priors_update( + &self, + portfolio_state_dev_ptr: u64, + n_envs: i32, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp14_isv_slots::{ + KELLY_PRIOR_COUNT_MAX, KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_EMA_ALPHA, + KELLY_PRIOR_LOSSES_INDEX, KELLY_PRIOR_SUM_LOSSES_INDEX, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_WINS_INDEX, KELLY_PRIOR_WINS_INDEX, + SENTINEL_KELLY_PRIOR_LOSSES, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_WINS, + }; + // PS_STRIDE from state_layout.cuh — matches the value used by + // `launch_kelly_cap_update`. Grown 41→43 by Plan 3 D.4c. + const PS_STRIDE: i32 = 43; + + debug_assert!( + self.isv_signals_dev_ptr != 0, + "launch_kelly_bayesian_priors_update: isv_signals_dev_ptr must be allocated" + ); + + self.kelly_bayesian_priors_update_ops.launch( + &self.stream, + portfolio_state_dev_ptr, + n_envs, + PS_STRIDE, + self.isv_signals_dev_ptr, + KELLY_PRIOR_WINS_INDEX as i32, + KELLY_PRIOR_LOSSES_INDEX as i32, + KELLY_PRIOR_SUM_WINS_INDEX as i32, + KELLY_PRIOR_SUM_LOSSES_INDEX as i32, + SENTINEL_KELLY_PRIOR_WINS, + SENTINEL_KELLY_PRIOR_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, + SENTINEL_KELLY_PRIOR_SUM_LOSSES, + KELLY_PRIOR_COUNT_MIN, + KELLY_PRIOR_COUNT_MAX, + KELLY_PRIOR_SUM_MIN, + KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_EMA_ALPHA, + )?; + + Ok(()) + } + /// SP8 (Fix 36, 2026-05-03): launch GPU-only `train_active_frac` canary /// producer. /// @@ -23057,6 +23152,14 @@ impl GpuDqnTrainer { let reward_cap_update_ops = super::gpu_aux_trunk::RewardCapUpdateOps::new(&stream)?; + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors + // producer. Pre-loads the `CudaFunction` handle at construction per + // `pearl_no_host_branches_in_captured_graph.md`. Per-fold-end (per-epoch + // boundary in current scheduling) launch — Bayesian prior is a *prior + // belief* and should change slowly across folds (α=0.005 EMA). + let kelly_bayesian_priors_update_ops = + super::gpu_aux_trunk::KellyBayesianPriorsUpdateOps::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}")))?; @@ -23890,6 +23993,8 @@ impl GpuDqnTrainer { avg_win_hold_time_update_ops, // Class A P0-A (2026-05-08): adaptive reward-cap producer. reward_cap_update_ops, + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors producer. + kelly_bayesian_priors_update_ops, speculative_h_s2, speculative_features, speculative_valid: false, diff --git a/crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu new file mode 100644 index 000000000..423893015 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu @@ -0,0 +1,250 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * Class A P1-Producer — adaptive Bayesian Kelly priors producer (2026-05-08). + * + * Replaces the hardcoded Bayesian priors used in Kelly cap calculations: + * - `kelly_cap_update_kernel.cu:39-42` (epoch-boundary effective Kelly cap) + * - `trade_physics.cuh::kelly_position_cap:304-307` (per-step env cap) + * + * Both consumers ran with `prior_wins=2.0f, prior_losses=2.0f, + * prior_sum_wins=0.01f, prior_sum_losses=0.01f` — 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; the prior was *literally* the prior (in the Bayesian sense), + * but it was a fixed scalar, not an adaptive belief. + * + * This producer aggregates realized Kelly stats from the existing + * `portfolio_state[n_envs, PS_STRIDE]` buffer — same source the + * `kelly_cap_update_kernel` already reads from at the same per-epoch + * boundary. The four PS slots (`PS_KELLY_WIN_COUNT`, + * `PS_KELLY_LOSS_COUNT`, `PS_KELLY_SUM_WINS`, `PS_KELLY_SUM_LOSSES`) are + * realized counts/sums per env; we aggregate across envs and slow-EMA-blend + * into ISV[454..458). On the next fold's cold-start path, the consumer + * reads the EMA prior (a learned belief about the distribution shape) instead + * of the fixed sprint-1 default. + * + * Ordering: launched per-fold-end (currently per-epoch boundary in the + * scheduling chain in `training_loop.rs`) BEFORE `kelly_cap_update` so the + * effective-Kelly-cap kernel sees the freshly-blended priors. The launch + * site lives at the same epoch-boundary block as `kelly_cap_update` and + * `reward_cap_update`. + * + * Algorithm: + * + * Phase 1 (block-tree-reduce): each thread of the single block sweeps + * its tile of `n_envs` portfolio_state slots, accumulating + * local_win_count, local_loss_count, local_sum_wins, local_sum_losses + * from the four PS_KELLY_* fields. Block-tree-reduce in shmem (no + * atomicAdd per `feedback_no_atomicadd.md`) yields the four sums in + * thread 0. + * + * Phase 2 (single thread, EMA blend): if no realized trades observed + * (total_win_count + total_loss_count == 0), keep ISV slots unchanged + * (cold-start sentinel persists; consumers fall back to the static + * defaults via the NULL-tolerant guard in `kelly_position_cap` / + * `kelly_cap_update_kernel`). Otherwise: + * + * target_wins = total_win_count + * target_losses = total_loss_count + * target_sum_wins = total_sum_wins + * target_sum_losses = total_sum_losses + * + * Pearl-A first-observation bootstrap: each slot independently checks + * `|current - sentinel| < EPS_F`; if so, REPLACE with target. Otherwise + * slow EMA blend with α=0.005 (per-fold cadence — the Bayesian prior + * is a *prior belief* and should change slowly across folds; a fast + * "prior" would just be a noisy Kelly estimate). Per + * `pearl_first_observation_bootstrap.md` and + * `pearl_controller_anchors_isv_driven.md`. + * + * Bilateral clamp on each slot per `pearl_symmetric_clamp_audit.md`: + * - counts in [0.5, 100]: below 0.5 the prior is effectively absent + * (kelly_f swings on first real trade); above 100 the prior + * dominates 10+ real trades (the maturity threshold for warmup_floor + * blend in kelly_position_cap). + * - sums in [0.001, 1.0]: below 1e-3 hits the same numerical floor + * as `fmaxf(avg_loss, 0.0001f)` in the Kelly formula; above 1.0 + * implies a per-trade win/loss > 100% of equity which would be + * risk-management-failure territory anyway. + * + * Pearls + invariants: + * + * - `feedback_no_atomicadd.md` — single-block kernel; block-tree-reduce + * in shmem; no atomics across blocks. + * + * - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction` + * pre-loaded at construction; on-device guards for "no realized + * trades" and Pearl-A sentinel detection. + * + * - `pearl_first_observation_bootstrap.md` — first valid observation + * replaces sentinel directly; no blend. + * + * - `feedback_no_stubs.md` — full body, no return-zero placeholders. + * + * - `feedback_isv_for_adaptive_bounds.md` — count bounds [0.5, 100], + * sum bounds [0.001, 1.0] are dimensional safety floors, NOT tuning. + * + * - `pearl_symmetric_clamp_audit.md` — bilateral + * `fmaxf(lo, fminf(x, hi))` clamp on each slot before writing. + * + * - `pearl_controller_anchors_isv_driven.md` — every controller anchor + * (the priors are anchors for the Kelly cap formula) is ISV-driven. + * + * Args: + * portfolio_state — `[n_envs * ps_stride]` device f32; the + * same buffer `kelly_cap_update_kernel` + * already reads. We use the four + * PS_KELLY_* fields per env. + * n_envs — number of envs (alloc_episodes from the + * collector). + * ps_stride — PS_STRIDE from state_layout.cuh (43 + * post-Plan-3-D.4c). + * isv — `[ISV_TOTAL_DIM]` device f32, modified + * in place at indices `*_idx`. + * wins_idx — KELLY_PRIOR_WINS_INDEX (454). + * losses_idx — KELLY_PRIOR_LOSSES_INDEX (455). + * sum_wins_idx — KELLY_PRIOR_SUM_WINS_INDEX (456). + * sum_losses_idx — KELLY_PRIOR_SUM_LOSSES_INDEX (457). + * sentinel_wins — SENTINEL_KELLY_PRIOR_WINS (2.0). + * sentinel_losses — SENTINEL_KELLY_PRIOR_LOSSES (2.0). + * sentinel_sum_wins — SENTINEL_KELLY_PRIOR_SUM_WINS (0.01). + * sentinel_sum_losses — SENTINEL_KELLY_PRIOR_SUM_LOSSES (0.01). + * count_min, count_max — KELLY_PRIOR_COUNT_MIN/MAX (0.5, 100). + * sum_min, sum_max — KELLY_PRIOR_SUM_MIN/MAX (0.001, 1.0). + * alpha — KELLY_PRIOR_EMA_ALPHA (0.005, slow). + * + * Launch: grid=(1, 1, 1), block=(BLK_DIM=256, 1, 1). + * Shared memory: 4 * BLK_DIM floats: + * sh_win_count — sum of PS_KELLY_WIN_COUNT across envs + * sh_loss_count — sum of PS_KELLY_LOSS_COUNT across envs + * sh_sum_wins — sum of PS_KELLY_SUM_WINS across envs + * sh_sum_losses — sum of PS_KELLY_SUM_LOSSES across envs + * ══════════════════════════════════════════════════════════════════════════ */ + +#include +#include "state_layout.cuh" /* PS_KELLY_WIN_COUNT, PS_KELLY_LOSS_COUNT, + * PS_KELLY_SUM_WINS, PS_KELLY_SUM_LOSSES */ + +#define BLK_DIM 256 +#define EPS_F 1e-6f + +extern "C" __global__ +void kelly_bayesian_priors_update( + const float* __restrict__ portfolio_state, + int n_envs, + int ps_stride, + float* __restrict__ isv, + int wins_idx, + int losses_idx, + int sum_wins_idx, + int sum_losses_idx, + float sentinel_wins, + float sentinel_losses, + float sentinel_sum_wins, + float sentinel_sum_losses, + float count_min, + float count_max, + float sum_min, + float sum_max, + float alpha) +{ + /* Single-block kernel — only block 0 does anything. */ + if (blockIdx.x != 0) return; + + extern __shared__ float shmem[]; + float* sh_win_count = shmem; /* [BLK_DIM] */ + float* sh_loss_count = shmem + BLK_DIM; /* [BLK_DIM] */ + float* sh_sum_wins = shmem + 2 * BLK_DIM; /* [BLK_DIM] */ + float* sh_sum_losses = shmem + 3 * BLK_DIM; /* [BLK_DIM] */ + + int tid = threadIdx.x; + + /* Phase 1a: stride-based sweep — accumulate the four Kelly stats + * across all envs. PS_KELLY_* indices come from state_layout.cuh + * (the same constants kelly_cap_update_kernel uses). Counts are + * stored as floats in PS slots (matches Kelly stat accumulator + * convention in apply_kelly_cap). */ + float local_wins = 0.0f; + float local_losses = 0.0f; + float local_sumw = 0.0f; + float local_suml = 0.0f; + for (int e = tid; e < n_envs; e += BLK_DIM) { + const float* ps = portfolio_state + (long long)e * ps_stride; + local_wins += ps[PS_KELLY_WIN_COUNT]; + local_losses += ps[PS_KELLY_LOSS_COUNT]; + local_sumw += ps[PS_KELLY_SUM_WINS]; + local_suml += ps[PS_KELLY_SUM_LOSSES]; + } + sh_win_count[tid] = local_wins; + sh_loss_count[tid] = local_losses; + sh_sum_wins[tid] = local_sumw; + sh_sum_losses[tid] = local_suml; + __syncthreads(); + + /* Phase 1b: block-tree-reduce (no atomicAdd) — halve every iter. + * All four streams fold via addition. */ + for (int s = BLK_DIM >> 1; s > 0; s >>= 1) { + if (tid < s) { + sh_win_count[tid] += sh_win_count[tid + s]; + sh_loss_count[tid] += sh_loss_count[tid + s]; + sh_sum_wins[tid] += sh_sum_wins[tid + s]; + sh_sum_losses[tid] += sh_sum_losses[tid + s]; + } + __syncthreads(); + } + + /* Phase 2: thread 0 finalises EMA + Pearl-A + clamp + writes the four slots. */ + if (tid != 0) return; + + float total_wins = sh_win_count[0]; + float total_losses = sh_loss_count[0]; + float total_sumw = sh_sum_wins[0]; + float total_suml = sh_sum_losses[0]; + + /* Guard: no realized trades observed across all envs — keep ISV slots + * unchanged (cold-start sentinel persists; consumers fall back to + * the static defaults via the NULL-tolerant guard pattern). */ + if (total_wins + total_losses <= 0.0f) { + return; + } + + /* Bilateral pre-EMA clamp on the targets so the blended value stays in + * the dimensional-safety bounds even before the post-EMA defensive + * clamp. Per `pearl_symmetric_clamp_audit.md`. */ + float target_wins = fmaxf(count_min, fminf(total_wins, count_max)); + float target_losses = fmaxf(count_min, fminf(total_losses, count_max)); + float target_sumw = fmaxf(sum_min, fminf(total_sumw, sum_max)); + float target_suml = fmaxf(sum_min, fminf(total_suml, sum_max)); + + /* Pearl-A first-observation bootstrap on each slot independently. + * The sentinel matches the pre-P1-Producer hardcoded constants + * (2.0/2.0/0.01/0.01) for bit-identical cold-start behavior. */ + float cur_wins = isv[wins_idx]; + float cur_losses = isv[losses_idx]; + float cur_sumw = isv[sum_wins_idx]; + float cur_suml = isv[sum_losses_idx]; + + float blended_wins = (fabsf(cur_wins - sentinel_wins) < EPS_F) + ? target_wins + : ((1.0f - alpha) * cur_wins + alpha * target_wins); + float blended_losses = (fabsf(cur_losses - sentinel_losses) < EPS_F) + ? target_losses + : ((1.0f - alpha) * cur_losses + alpha * target_losses); + float blended_sumw = (fabsf(cur_sumw - sentinel_sum_wins) < EPS_F) + ? target_sumw + : ((1.0f - alpha) * cur_sumw + alpha * target_sumw); + float blended_suml = (fabsf(cur_suml - sentinel_sum_losses) < EPS_F) + ? target_suml + : ((1.0f - alpha) * cur_suml + alpha * target_suml); + + /* Defensive post-EMA bilateral clamp — handles malformed prior state + * outside the dimensional-safety range. */ + blended_wins = fmaxf(count_min, fminf(blended_wins, count_max)); + blended_losses = fmaxf(count_min, fminf(blended_losses, count_max)); + blended_sumw = fmaxf(sum_min, fminf(blended_sumw, sum_max)); + blended_suml = fmaxf(sum_min, fminf(blended_suml, sum_max)); + + isv[wins_idx] = blended_wins; + isv[losses_idx] = blended_losses; + isv[sum_wins_idx] = blended_sumw; + isv[sum_losses_idx] = blended_suml; +} diff --git a/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu index ddc8f3b87..7806ffd93 100644 --- a/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu @@ -1,5 +1,7 @@ /* kelly_cap_update — effective Kelly cap, GPU-driven (Plan 1 Task 11). * Reads: ISV[LEARNING_HEALTH_INDEX=12] (safety multiplier coupled to health) + * ISV[KELLY_PRIOR_*_INDEX=454..458) (Class A P1-Producer adaptive + * Bayesian priors — slow-EMA from realized fold trades) * portfolio_state[n_envs, PS_STRIDE] — Kelly win/loss stats per env * Writes: ISV[KELLY_CAP_EFF_INDEX=47] (shifted +3 from 44 by D.2 per-branch gamma, Plan 2 Task 3) * Cold path (per-epoch boundary). Single-thread kernel. @@ -9,9 +11,14 @@ * health-coupled safety multiplier. Result is a scalar cap in [0, 1]. * Downstream kernel readers use this cap to limit position sizing. * - * Bayesian priors (matching trade_physics.cuh): - * prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01 - * This ensures the cap never collapses to zero on cold start. + * Bayesian priors (Class A P1-Producer 2026-05-08): now read from + * ISV[454..458) where they are produced by `kelly_bayesian_priors_update_kernel` + * at the SAME per-epoch boundary, launched RIGHT BEFORE this kernel so the + * cold-path Kelly cap sees the freshly-blended priors. Sentinels match + * the pre-P1-Producer hardcoded constants + * (2.0 / 2.0 / 0.01 / 0.01) for bit-identical cold-start behavior. The + * fallback path below applies when an ISV slot is at sentinel (Pearl-A + * pre-bootstrap) or out of the dimensional-safety range — defensive only. */ #include "state_layout.cuh" @@ -27,6 +34,21 @@ * Slot constant mirrors `sp5_isv_slots.rs::KELLY_WARMUP_FLOOR_INDEX = 330`. */ #define SP9_KELLY_WARMUP_FLOOR_INDEX 330 +/* Class A P1-Producer cold-start fallback helper. Returns the ISV slot value + * if it is within the dimensional-safety range, else the static default + * (the pre-P1-Producer hardcoded value for bit-identical cold-start). */ +__device__ __forceinline__ float kelly_prior_or_default( + float isv_val, + float fallback_default, + float lo_bound, + float hi_bound +) { + if (isv_val >= lo_bound && isv_val <= hi_bound) { + return isv_val; + } + return fallback_default; +} + extern "C" __global__ void kelly_cap_update( const float* __restrict__ isv, float* __restrict__ isv_out, @@ -36,10 +58,20 @@ extern "C" __global__ void kelly_cap_update( ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; - const float prior_wins = 2.0f; - const float prior_losses = 2.0f; - const float prior_sum_wins = 0.01f; - const float prior_sum_losses = 0.01f; + /* Class A P1-Producer (2026-05-08): adaptive Bayesian priors from + * ISV[454..458). Producer kernel `kelly_bayesian_priors_update_kernel` + * runs at the same per-epoch boundary RIGHT BEFORE this kernel. + * Cold-start path: ISV slot at sentinel (Pearl-A pre-bootstrap) or + * out-of-range → fall back to the static default (matches pre-P1-Producer + * hardcoded values for bit-identical cold-start behavior). */ + const float prior_wins = kelly_prior_or_default( + isv[KELLY_PRIOR_WINS_INDEX], KELLY_PRIOR_WINS_DEFAULT, 0.5f, 100.0f); + const float prior_losses = kelly_prior_or_default( + isv[KELLY_PRIOR_LOSSES_INDEX], KELLY_PRIOR_LOSSES_DEFAULT, 0.5f, 100.0f); + const float prior_sum_wins = kelly_prior_or_default( + isv[KELLY_PRIOR_SUM_WINS_INDEX], KELLY_PRIOR_SUM_WINS_DEFAULT, 0.001f, 1.0f); + const float prior_sum_losses = kelly_prior_or_default( + isv[KELLY_PRIOR_SUM_LOSSES_INDEX], KELLY_PRIOR_SUM_LOSSES_DEFAULT, 0.001f, 1.0f); float health = fminf(fmaxf(isv[12], 0.0f), 1.0f); /* Safety multiplier: 0.5 + 0.5*health — collapses cap during degradation */ diff --git a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs index d0d6e79c1..029bdb04d 100644 --- a/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp14_isv_slots.rs @@ -164,6 +164,65 @@ pub const REWARD_CAP_EMA_ALPHA: f32 = 0.01; pub const SP14_P0A_SLOT_BASE: usize = 452; pub const SP14_P0A_SLOT_END: usize = 454; +// ── Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors ───── +// Replaces the hardcoded +// `prior_wins=2.0, prior_losses=2.0, prior_sum_wins=0.01, prior_sum_losses=0.01` +// in `kelly_cap_update_kernel.cu:39-42` AND +// `trade_physics.cuh::kelly_position_cap:304-307` with ISV-driven slow-EMA +// values produced by `kelly_bayesian_priors_update_kernel.cu` at per-fold +// boundary. The producer aggregates realized Kelly stats from +// `portfolio_state[n_envs, PS_STRIDE]` (the same buffer +// `kelly_cap_update_kernel` already reads from `PS_KELLY_*` indices) and +// EMA-blends into the slot. The static defaults remain as Pearl-A +// sentinels so cold-start (pre-first-observation) is bit-identical to the +// pre-P1-Producer behavior. +// +// Why slow EMA at fold cadence (α=0.005): the Bayesian prior is a *prior +// belief* about the Kelly distribution before observing the new fold's +// trades. It should change slowly across folds — a fast-moving "prior" +// would just be a noisy Kelly estimate. Per +// `pearl_controller_anchors_isv_driven.md` and +// `feedback_isv_for_adaptive_bounds.md` (the prior is a controller +// anchor, not a tuning constant). +// +// Producer kernel: `kelly_bayesian_priors_update_kernel.cu`. +// Consumers (atomic migration in same commit): +// - `kelly_cap_update_kernel.cu:39-42` → ISV[454..458) +// - `trade_physics.cuh::kelly_position_cap:304-307` → threaded +// `isv_signals_ptr` argument with NULL-tolerant cold-start fallback. +pub const KELLY_PRIOR_WINS_INDEX: usize = 454; // adaptive prior_wins (replaces 2.0f) +pub const KELLY_PRIOR_LOSSES_INDEX: usize = 455; // adaptive prior_losses (replaces 2.0f) +pub const KELLY_PRIOR_SUM_WINS_INDEX: usize = 456; // adaptive prior_sum_wins (replaces 0.01f) +pub const KELLY_PRIOR_SUM_LOSSES_INDEX: usize = 457; // adaptive prior_sum_losses (replaces 0.01f) + +// Sentinels — match pre-P1-Producer hardcoded values for bit-identical +// cold-start behavior before the first valid observation lands. +pub const SENTINEL_KELLY_PRIOR_WINS: f32 = 2.0; +pub const SENTINEL_KELLY_PRIOR_LOSSES: f32 = 2.0; +pub const SENTINEL_KELLY_PRIOR_SUM_WINS: f32 = 0.01; +pub const SENTINEL_KELLY_PRIOR_SUM_LOSSES: f32 = 0.01; + +// Bounds — Category-1 dimensional safety floors per +// `feedback_isv_for_adaptive_bounds.md`. Counts in [0.5, 100]: below 0.5 +// the prior is effectively absent (kelly_f swings on first real trade); +// above 100 the prior dominates 10+ real trades (the maturity threshold +// for warmup_floor blend in kelly_position_cap). Sums in [0.001, 1.0]: +// below 1e-3 hits the same numerical floor as `fmaxf(avg_loss, 0.0001f)` +// in the Kelly formula; above 1.0 implies a per-trade win/loss > 100% of +// equity which would be risk-management-failure territory anyway. +pub const KELLY_PRIOR_COUNT_MIN: f32 = 0.5; +pub const KELLY_PRIOR_COUNT_MAX: f32 = 100.0; +pub const KELLY_PRIOR_SUM_MIN: f32 = 0.001; +pub const KELLY_PRIOR_SUM_MAX: f32 = 1.0; + +// EMA blend rate — per-fold cadence, slower than reward_cap's α=0.01. +// The Bayesian prior is a prior belief; it should change slowly across +// folds (a fast-moving "prior" would just be a noisy Kelly estimate). +pub const KELLY_PRIOR_EMA_ALPHA: f32 = 0.005; + +pub const SP14_P1_SLOT_BASE: usize = 454; +pub const SP14_P1_SLOT_END: usize = 458; + #[cfg(test)] mod tests { use super::*; @@ -273,4 +332,42 @@ mod tests { SP14_P0A_SLOT_END, ISV_TOTAL_DIM, ); } + + /// Lock Class A P1-Producer (2026-05-08) adaptive Bayesian-prior slot + /// layout. Four contiguous slots [454..458) replacing the hardcoded + /// `prior_wins=2.0 / prior_losses=2.0 / prior_sum_wins=0.01 / + /// prior_sum_losses=0.01` constants in `kelly_cap_update_kernel.cu` and + /// `trade_physics.cuh::kelly_position_cap`. Sentinels match the + /// pre-P1-Producer hardcoded values for bit-identical cold-start. + #[test] + fn sp14_p1_kelly_prior_slot_layout_locked() { + assert_eq!(SP14_P1_SLOT_BASE, 454); + assert_eq!(SP14_P1_SLOT_END, 458); + assert_eq!(KELLY_PRIOR_WINS_INDEX, 454); + assert_eq!(KELLY_PRIOR_LOSSES_INDEX, 455); + assert_eq!(KELLY_PRIOR_SUM_WINS_INDEX, 456); + assert_eq!(KELLY_PRIOR_SUM_LOSSES_INDEX, 457); + // Sentinels match pre-P1-Producer hardcoded constants for + // bit-identical cold-start behavior before first valid observation. + assert_eq!(SENTINEL_KELLY_PRIOR_WINS, 2.0); + assert_eq!(SENTINEL_KELLY_PRIOR_LOSSES, 2.0); + assert_eq!(SENTINEL_KELLY_PRIOR_SUM_WINS, 0.01); + assert_eq!(SENTINEL_KELLY_PRIOR_SUM_LOSSES, 0.01); + // Category-1 dimensional safety: counts [0.5, 100], sums [0.001, 1.0]. + assert_eq!(KELLY_PRIOR_COUNT_MIN, 0.5); + assert_eq!(KELLY_PRIOR_COUNT_MAX, 100.0); + assert_eq!(KELLY_PRIOR_SUM_MIN, 0.001); + assert_eq!(KELLY_PRIOR_SUM_MAX, 1.0); + } + + #[test] + fn all_sp14_p1_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP14_P1_SLOT_END <= ISV_TOTAL_DIM, + "SP14_P1_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for Class A P1-Producer Bayesian-prior slots; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP14_P1_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 b2932ed07..3ee15d7e3 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -201,6 +201,13 @@ #define ISV_HOLD_RATE_TARGET_IDX 381 // SP13 v3 target Hold-pick rate (static, default 0.20) #define ISV_HOLD_RATE_OBSERVED_EMA_IDX 382 // SP13 v3 observed Hold-pick rate EMA (per-fold, sentinel 0.0) +// NOTE: this macro is kept at 383 for the historic SP13-era boundary marker +// and is referenced only in commentary inside .cu/.cuh files. The actual +// ISV bus dimension is the Rust `ISV_TOTAL_DIM` constant in +// `gpu_dqn_trainer.rs` (currently 458, post Class A P1-Producer +// 2026-05-08). Kernels take `const float*` pointers and use raw indices — +// no .cu/.cuh file dimensions an array with this macro. The Rust constant +// is the single source of truth for the bus size. #define ISV_TOTAL_DIM 383 // ──────────────────────────────────────────────────────────────────────────── @@ -299,6 +306,36 @@ #define REWARD_POS_CAP_MIN_BOUND 1.0f #define REWARD_POS_CAP_MAX_BOUND 50.0f +// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors. +// Replaces hardcoded `prior_wins=2.0f / prior_losses=2.0f / +// prior_sum_wins=0.01f / prior_sum_losses=0.01f` in kelly_cap_update_kernel +// and trade_physics.cuh::kelly_position_cap with ISV-driven slow-EMA values +// fed by the `kelly_bayesian_priors_update_kernel` producer at per-fold +// boundary (per-epoch boundary in current scheduling — same cadence as +// kelly_cap_update). Sentinels match the pre-P1-Producer hardcoded values +// so cold-start (pre-first-observation) is bit-identical to the prior +// behavior. Mirrors `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` +// constants of the same names (locked by +// `sp14_p1_kelly_prior_slot_layout_locked` test). +#define KELLY_PRIOR_WINS_INDEX 454 +#define KELLY_PRIOR_LOSSES_INDEX 455 +#define KELLY_PRIOR_SUM_WINS_INDEX 456 +#define KELLY_PRIOR_SUM_LOSSES_INDEX 457 +#define SENTINEL_KELLY_PRIOR_WINS 2.0f +#define SENTINEL_KELLY_PRIOR_LOSSES 2.0f +#define SENTINEL_KELLY_PRIOR_SUM_WINS 0.01f +#define SENTINEL_KELLY_PRIOR_SUM_LOSSES 0.01f +// Cold-start fallback for the consumer-side ISV-NULL guard (mirrors the +// existing `kelly_warmup_floor_sp9` and `kelly_f_smooth` patterns in +// trade_physics.cuh). When `isv_signals_ptr == NULL` (validation envs that +// don't pass the bus through) OR the slot value is at sentinel, the +// consumer falls back to the static defaults below. Bit-identical +// pre-P1-Producer behavior. +#define KELLY_PRIOR_WINS_DEFAULT 2.0f +#define KELLY_PRIOR_LOSSES_DEFAULT 2.0f +#define KELLY_PRIOR_SUM_WINS_DEFAULT 0.01f +#define KELLY_PRIOR_SUM_LOSSES_DEFAULT 0.01f + // ── Compile-time checks ── static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index 931f08399..e693ee266 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -263,10 +263,21 @@ __device__ __forceinline__ float compute_drawdown( * the environment refuses to let the agent over-lever, rather than * scoring the agent for matching a formula. * - * Priors (prior_wins=2, prior_losses=2, prior_sum_wins=0.01, - * prior_sum_losses=0.01) prevent the cap from being degenerate on cold-start - * runs where no trades have completed yet — without priors Kelly would - * be undefined for the first ~10 trades. + * Priors (Class A P1-Producer 2026-05-08): ISV-driven adaptive Bayesian + * priors at ISV[454..458) replace the previously-hardcoded + * `prior_wins=2.0f, prior_losses=2.0f, prior_sum_wins=0.01f, + * prior_sum_losses=0.01f` constants. Producer kernel + * `kelly_bayesian_priors_update_kernel` aggregates realized Kelly stats + * across envs and slow-EMA-blends (α=0.005) into the four slots at the + * same per-epoch boundary as `kelly_cap_update_kernel`. Sentinels match + * the pre-P1-Producer hardcoded values for bit-identical cold-start. + * NULL-tolerant (`isv_signals_ptr == NULL`) and out-of-range guards + * fall back to the static defaults — the helper is safe to call from + * validation envs that don't pass the bus through (matches the existing + * `kelly_f_smooth` / `kelly_warmup_floor_sp9` pattern in + * `unified_env_step_core`). The priors prevent the cap from being + * degenerate on cold-start runs where no trades have completed yet — + * without priors Kelly would be undefined for the first ~10 trades. * * The safety_multiplier is composed by the caller from two orthogonal * adaptive signals: @@ -299,12 +310,36 @@ __device__ __forceinline__ float kelly_position_cap( float max_position, float safety_multiplier, float conviction, - float health_floor + float health_floor, + /* ── Class A P1-Producer (2026-05-08): adaptive Bayesian priors ── + * NULL-tolerant: when NULL OR slot at sentinel/out-of-range, falls + * back to the static defaults (matches pre-P1-Producer hardcoded + * values for bit-identical cold-start). The fallback path mirrors + * the existing `kelly_f_smooth` / `kelly_warmup_floor_sp9` pattern + * already in use elsewhere in this file. */ + const float* __restrict__ isv_signals_ptr /* [ISV_TOTAL_DIM]; NULL = static defaults */ ) { - const float prior_wins = 2.0f; - const float prior_losses = 2.0f; - const float prior_sum_wins = 0.01f; - const float prior_sum_losses = 0.01f; + /* Class A P1-Producer (2026-05-08): adaptive Bayesian priors. NULL + * guard handles validation envs that don't pass the bus through; + * range guard handles cold-start (sentinel) and malformed prior + * state. Both fall back to the pre-P1-Producer static defaults + * for bit-identical cold-start behavior. */ + float prior_wins, prior_losses, prior_sum_wins, prior_sum_losses; + if (isv_signals_ptr != NULL) { + float pw = isv_signals_ptr[KELLY_PRIOR_WINS_INDEX]; + float pl = isv_signals_ptr[KELLY_PRIOR_LOSSES_INDEX]; + float psw = isv_signals_ptr[KELLY_PRIOR_SUM_WINS_INDEX]; + float psl = isv_signals_ptr[KELLY_PRIOR_SUM_LOSSES_INDEX]; + prior_wins = (pw >= 0.5f && pw <= 100.0f) ? pw : KELLY_PRIOR_WINS_DEFAULT; + prior_losses = (pl >= 0.5f && pl <= 100.0f) ? pl : KELLY_PRIOR_LOSSES_DEFAULT; + prior_sum_wins = (psw >= 0.001f && psw <= 1.0f) ? psw : KELLY_PRIOR_SUM_WINS_DEFAULT; + prior_sum_losses = (psl >= 0.001f && psl <= 1.0f) ? psl : KELLY_PRIOR_SUM_LOSSES_DEFAULT; + } else { + prior_wins = KELLY_PRIOR_WINS_DEFAULT; + prior_losses = KELLY_PRIOR_LOSSES_DEFAULT; + prior_sum_wins = KELLY_PRIOR_SUM_WINS_DEFAULT; + prior_sum_losses = KELLY_PRIOR_SUM_LOSSES_DEFAULT; + } float eff_wins = win_count + prior_wins; float eff_losses = loss_count + prior_losses; @@ -389,7 +424,11 @@ __device__ __forceinline__ float kelly_position_cap( /* Apply Kelly cap: clamp a signed target_position to [-cap, +cap]. * Convenience for call sites that want a single clamp call. * `conviction` ∈ [0, 1] threads through to kelly_position_cap's adaptive - * warmup_floor — see docstring there. */ + * warmup_floor — see docstring there. + * `isv_signals_ptr` threads through to kelly_position_cap's adaptive + * Bayesian priors (Class A P1-Producer 2026-05-08). NULL-tolerant — + * validation envs that don't pass the bus through fall back to the + * pre-P1-Producer static defaults for bit-identical cold-start. */ __device__ __forceinline__ float apply_kelly_cap( float target_position, float win_count, @@ -399,11 +438,12 @@ __device__ __forceinline__ float apply_kelly_cap( float max_position, float safety_multiplier, float conviction, - float health_floor + float health_floor, + const float* __restrict__ isv_signals_ptr /* [ISV_TOTAL_DIM]; NULL = static defaults */ ) { float cap = kelly_position_cap(win_count, loss_count, sum_wins, sum_losses, max_position, safety_multiplier, conviction, - health_floor); + health_floor, isv_signals_ptr); return fmaxf(-cap, fminf(cap, target_position)); } @@ -861,7 +901,13 @@ __device__ __forceinline__ void unified_env_step_core( *sum_wins, *sum_losses, max_position_physics, safety, policy_conviction, - health_safety_sp9 + health_safety_sp9, + /* Class A P1-Producer (2026-05-08): adaptive Bayesian priors. + * `isv_signals_ptr` is the existing parameter to + * `unified_env_step_core` (already used for SP5 Pearl 6 + * `kelly_f_smooth` and SP9 Fix 37 `kelly_warmup_floor_sp9` + * just above). NULL-tolerant on the consumer side. */ + isv_signals_ptr ); } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index a1f93e933..88f14ab57 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1046,6 +1046,48 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[REWARD_NEG_CAP_ADAPTIVE_INDEX=453] — Class A P0-A adaptive negative reward cap (replaces hardcoded REWARD_NEG_CAP=-10.0f from state_layout.cuh:267). Produced by `reward_cap_update_kernel` as NEG = −REWARD_NEG_TO_POS_RATIO=2.0 × POS_CAP. The 2:1 ratio (Kahneman/Tversky meta-analysis ~2.0-2.25) is preserved at producer-time inside the kernel (single source of truth — no consumer applies the multiplier itself), per `pearl_audit_unboundedness_for_implicit_asymmetry`. FoldReset sentinel SENTINEL_REWARD_NEG_CAP=-10.0 (matches pre-P0-A hardcoded value for bit-identical cold-start). Resulting bounds [-100, -2] derived from POS bounds × ratio. Consumed alongside POS slot at the same 3 sites.", }, + // ── Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly priors ── + // Four slots [454..458) replacing the hardcoded + // `prior_wins=2.0f / prior_losses=2.0f / prior_sum_wins=0.01f / + // prior_sum_losses=0.01f` constants in + // `kelly_cap_update_kernel.cu:39-42` and + // `trade_physics.cuh::kelly_position_cap:304-307`. Producer + // kernel `kelly_bayesian_priors_update_kernel` aggregates + // realized PS_KELLY_{WIN_COUNT, LOSS_COUNT, SUM_WINS, + // SUM_LOSSES} across envs from the same `portfolio_state` buffer + // `kelly_cap_update_kernel` reads from at the same per-epoch + // boundary, then slow-EMA-blends into the four slots. Pearl-A + // bootstrap sentinels (2.0/2.0/0.01/0.01) match the + // pre-P1-Producer hardcoded values for bit-identical cold-start + // behavior — consumers fall back to the static defaults via the + // NULL-tolerant + range guard pattern (mirrors the existing + // `kelly_f_smooth` / `kelly_warmup_floor_sp9` patterns in + // `unified_env_step_core`). Once the first valid observation + // lands the adaptive path takes over via slow EMA (α=0.005; + // per-fold cadence — the Bayesian prior is a *prior belief* + // and should change slowly across folds). Per + // `feedback_isv_for_adaptive_bounds.md` and + // `pearl_controller_anchors_isv_driven.md`. + RegistryEntry { + name: "sp14_p1_kelly_prior_wins", + category: ResetCategory::FoldReset, + description: "ISV[KELLY_PRIOR_WINS_INDEX=454] — Class A P1-Producer adaptive Bayesian prior_wins (replaces hardcoded prior_wins=2.0f from kelly_cap_update_kernel.cu:39 + trade_physics.cuh::kelly_position_cap:304). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_WIN_COUNT field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_WINS=2.0 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.5, 100] — Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds`. Consumed by kelly_cap_update_kernel (cold-path epoch-boundary effective Kelly cap) and trade_physics.cuh::kelly_position_cap (per-step env Kelly cap, NULL-tolerant via threaded `isv_signals_ptr`).", + }, + RegistryEntry { + name: "sp14_p1_kelly_prior_losses", + category: ResetCategory::FoldReset, + description: "ISV[KELLY_PRIOR_LOSSES_INDEX=455] — Class A P1-Producer adaptive Bayesian prior_losses (replaces hardcoded prior_losses=2.0f from kelly_cap_update_kernel.cu:40 + trade_physics.cuh::kelly_position_cap:305). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_LOSS_COUNT field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_LOSSES=2.0 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.5, 100] same as prior_wins. Consumed at the same 2 sites as prior_wins.", + }, + RegistryEntry { + name: "sp14_p1_kelly_prior_sum_wins", + category: ResetCategory::FoldReset, + description: "ISV[KELLY_PRIOR_SUM_WINS_INDEX=456] — Class A P1-Producer adaptive Bayesian prior_sum_wins (replaces hardcoded prior_sum_wins=0.01f from kelly_cap_update_kernel.cu:41 + trade_physics.cuh::kelly_position_cap:306). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_SUM_WINS field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_SUM_WINS=0.01 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.001, 1.0] — sum below 1e-3 hits the `fmaxf(avg_loss, 0.0001f)` numerical floor; sum above 1.0 implies > 100% per-trade equity which is risk-management-failure territory. Consumed at the same 2 sites as prior_wins.", + }, + RegistryEntry { + name: "sp14_p1_kelly_prior_sum_losses", + category: ResetCategory::FoldReset, + description: "ISV[KELLY_PRIOR_SUM_LOSSES_INDEX=457] — Class A P1-Producer adaptive Bayesian prior_sum_losses (replaces hardcoded prior_sum_losses=0.01f from kelly_cap_update_kernel.cu:42 + trade_physics.cuh::kelly_position_cap:307). Produced by `kelly_bayesian_priors_update_kernel` from the realized PS_KELLY_SUM_LOSSES field aggregated across envs in the per-epoch `portfolio_state` buffer. FoldReset sentinel SENTINEL_KELLY_PRIOR_SUM_LOSSES=0.01 — Pearl-A first-observation bootstrap (matches pre-P1-Producer hardcoded value for bit-identical cold-start). α=0.005 slow EMA thereafter. Bounds [0.001, 1.0] same as prior_sum_wins. Consumed at the same 2 sites as prior_wins.", + }, // ── 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 9510e9613..deb7ca54e 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -383,6 +383,18 @@ impl DQNTrainer { { let ps_dev_ptr = collector.portfolio_states_dev_ptr(); let n_envs = collector.alloc_episodes() as i32; + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly + // priors update. MUST run BEFORE `launch_kelly_cap_update` + // so that kernel sees the freshly-blended priors. Sweeps + // the same `portfolio_state` buffer with the same n_envs + // and writes ISV[454..458). Pearl-A bootstrap + α=0.005 + // slow EMA. Replaces hardcoded + // prior_wins/prior_losses/prior_sum_wins/prior_sum_losses + // constants in kelly_cap_update_kernel.cu and + // trade_physics.cuh::kelly_position_cap. + if let Err(e) = fused.trainer().launch_kelly_bayesian_priors_update(ps_dev_ptr, n_envs) { + tracing::warn!(epoch, "launch_kelly_bayesian_priors_update failed (non-fatal): {e}"); + } if let Err(e) = fused.launch_kelly_cap_update(ps_dev_ptr, n_envs) { tracing::warn!(epoch, "launch_kelly_cap_update failed (non-fatal): {e}"); } @@ -8183,6 +8195,57 @@ impl DQNTrainer { ); } } + // Class A P1-Producer (2026-05-08): adaptive Bayesian Kelly + // priors. Four FoldReset slots [454..458) — sentinels match + // pre-P1-Producer hardcoded values (2.0/2.0/0.01/0.01) so + // the kernel's "first observation" check fires cleanly on + // the new fold's first launch (avoids cross-fold EMA + // contamination). Cold-start is bit-identical to + // pre-P1-Producer until the first valid observation lands. + "sp14_p1_kelly_prior_wins" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + KELLY_PRIOR_WINS_INDEX, SENTINEL_KELLY_PRIOR_WINS, + }; + fused.trainer().write_isv_signal_at( + KELLY_PRIOR_WINS_INDEX, + SENTINEL_KELLY_PRIOR_WINS, + ); + } + } + "sp14_p1_kelly_prior_losses" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + KELLY_PRIOR_LOSSES_INDEX, SENTINEL_KELLY_PRIOR_LOSSES, + }; + fused.trainer().write_isv_signal_at( + KELLY_PRIOR_LOSSES_INDEX, + SENTINEL_KELLY_PRIOR_LOSSES, + ); + } + } + "sp14_p1_kelly_prior_sum_wins" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + KELLY_PRIOR_SUM_WINS_INDEX, SENTINEL_KELLY_PRIOR_SUM_WINS, + }; + fused.trainer().write_isv_signal_at( + KELLY_PRIOR_SUM_WINS_INDEX, + SENTINEL_KELLY_PRIOR_SUM_WINS, + ); + } + } + "sp14_p1_kelly_prior_sum_losses" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp14_isv_slots::{ + KELLY_PRIOR_SUM_LOSSES_INDEX, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + }; + fused.trainer().write_isv_signal_at( + KELLY_PRIOR_SUM_LOSSES_INDEX, + SENTINEL_KELLY_PRIOR_SUM_LOSSES, + ); + } + } // 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 bbb68b739..31d170493 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -787,6 +787,396 @@ mod sp14_p0a_reward_cap_gpu { } } +// ═══════════════════════════════════════════════════════════════════════════ +// Class A P1-Producer (2026-05-08) — adaptive Bayesian Kelly priors producer +// tests. +// +// Verifies the `kelly_bayesian_priors_update_kernel` producer: +// 1. Pearl-A first-observation bootstrap: ISV at sentinel → REPLACES with +// the realized aggregated stats (no blend). +// 2. Slow EMA (α=0.005) blend after bootstrap. +// 3. Bounds enforced: counts in [0.5, 100], sums in [0.001, 1.0]. +// 4. No realized trades (zero PS_KELLY_* fields) → ISV slots preserved +// bit-exactly. +// +// 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_p1_kelly_priors_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::{ + KELLY_PRIOR_COUNT_MAX, KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_EMA_ALPHA, + KELLY_PRIOR_LOSSES_INDEX, KELLY_PRIOR_SUM_LOSSES_INDEX, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_WINS_INDEX, KELLY_PRIOR_WINS_INDEX, + SENTINEL_KELLY_PRIOR_LOSSES, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_WINS, + }; + + // PS_STRIDE constant — matches state_layout.cuh::PS_STRIDE and the + // value used by `launch_kelly_cap_update`. Grown 41→43 by Plan 3 D.4c. + const PS_STRIDE: usize = 43; + // PS_KELLY_* slot offsets within a single env's portfolio_state row. + const PS_KELLY_WIN_COUNT: usize = 14; + const PS_KELLY_LOSS_COUNT: usize = 15; + const PS_KELLY_SUM_WINS: usize = 16; + const PS_KELLY_SUM_LOSSES: usize = 17; + + const KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/kelly_bayesian_priors_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(KELLY_BAYESIAN_PRIORS_UPDATE_CUBIN.to_vec()) + .expect("load kelly_bayesian_priors_update_kernel cubin"); + module + .load_function("kelly_bayesian_priors_update") + .expect("load kelly_bayesian_priors_update function") + } + + /// 4 arrays × 256 × 4 bytes = 4096 bytes shmem. + const BLK_DIM: u32 = 256; + const SMEM_BYTES: u32 = 4 * BLK_DIM * std::mem::size_of::() as u32; + + /// Helper: launch with a fixed config. + #[allow(clippy::too_many_arguments)] + fn launch_kelly_priors( + stream: &Arc, + kernel: &CudaFunction, + portfolio_state_ptr: u64, + n_envs: i32, + ps_stride: i32, + isv_ptr: u64, + wins_idx: i32, + losses_idx: i32, + sum_wins_idx: i32, + sum_losses_idx: i32, + sentinel_wins: f32, + sentinel_losses: f32, + sentinel_sum_wins: f32, + sentinel_sum_losses: f32, + count_min: f32, + count_max: f32, + sum_min: f32, + sum_max: f32, + alpha: f32, + ) { + unsafe { + stream + .launch_builder(kernel) + .arg(&portfolio_state_ptr) + .arg(&n_envs) + .arg(&ps_stride) + .arg(&isv_ptr) + .arg(&wins_idx) + .arg(&losses_idx) + .arg(&sum_wins_idx) + .arg(&sum_losses_idx) + .arg(&sentinel_wins) + .arg(&sentinel_losses) + .arg(&sentinel_sum_wins) + .arg(&sentinel_sum_losses) + .arg(&count_min) + .arg(&count_max) + .arg(&sum_min) + .arg(&sum_max) + .arg(&alpha) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (BLK_DIM, 1, 1), + shared_mem_bytes: SMEM_BYTES, + }) + .expect("launch kelly_bayesian_priors_update"); + } + stream.synchronize().expect("sync after kelly_bayesian_priors_update"); + } + + /// Build a portfolio_state buffer of `n_envs * PS_STRIDE` entries with + /// the four PS_KELLY_* fields populated per env from the parallel slices. + /// All other PS_* fields are zero. + fn build_portfolio_state( + wins_per_env: &[f32], + losses_per_env: &[f32], + sum_wins_per_env: &[f32], + sum_losses_per_env: &[f32], + ) -> Vec { + let n = wins_per_env.len(); + assert_eq!(losses_per_env.len(), n); + assert_eq!(sum_wins_per_env.len(), n); + assert_eq!(sum_losses_per_env.len(), n); + let mut ps = vec![0.0_f32; n * PS_STRIDE]; + for e in 0..n { + let row_off = e * PS_STRIDE; + ps[row_off + PS_KELLY_WIN_COUNT] = wins_per_env[e]; + ps[row_off + PS_KELLY_LOSS_COUNT] = losses_per_env[e]; + ps[row_off + PS_KELLY_SUM_WINS] = sum_wins_per_env[e]; + ps[row_off + PS_KELLY_SUM_LOSSES] = sum_losses_per_env[e]; + } + ps + } + + /// Test 1 — Pearl-A first-observation bootstrap. + /// + /// 4 envs with realized stats. Total wins = 1+2+3+4 = 10, losses = 0+1+0+1 = 2, + /// sum_wins = 0.05+0.10+0.15+0.20 = 0.50, sum_losses = 0+0.03+0+0.05 = 0.08. + /// Pre-EMA clamp: counts to [0.5, 100], sums to [0.001, 1.0] — + /// 10/2/0.5/0.08 all in range, pass through unchanged. Cold-start + /// (sentinels 2.0/2.0/0.01/0.01): Pearl-A REPLACES → ISV = + /// (10, 2, 0.5, 0.08). + #[test] + #[ignore = "requires GPU"] + fn kelly_priors_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]; + // Cold-start: sentinels (matches pre-P1-Producer constants). + isv[KELLY_PRIOR_WINS_INDEX] = SENTINEL_KELLY_PRIOR_WINS; + isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES; + isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS; + isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES; + + let wins_per_env = [1.0_f32, 2.0, 3.0, 4.0]; + let losses_per_env = [0.0_f32, 1.0, 0.0, 1.0]; + let sum_wins_per_env = [0.05_f32, 0.10, 0.15, 0.20]; + let sum_losses_per_env = [0.0_f32, 0.03, 0.0, 0.05]; + let n_envs = wins_per_env.len(); + + let ps = build_portfolio_state(&wins_per_env, &losses_per_env, + &sum_wins_per_env, &sum_losses_per_env); + + let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps"); + ps_buf.write_from_slice(&ps); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_kelly_priors( + &stream, &kernel, + ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr, + KELLY_PRIOR_WINS_INDEX as i32, + KELLY_PRIOR_LOSSES_INDEX as i32, + KELLY_PRIOR_SUM_WINS_INDEX as i32, + KELLY_PRIOR_SUM_LOSSES_INDEX as i32, + SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let pw = result[KELLY_PRIOR_WINS_INDEX]; + let pl = result[KELLY_PRIOR_LOSSES_INDEX]; + let psw = result[KELLY_PRIOR_SUM_WINS_INDEX]; + let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX]; + + // Pearl-A: sentinel → REPLACE with target. Targets are the + // aggregated totals clamped to dimensional-safety bounds (all + // within range here, no clamping applied). + assert!((pw - 10.0).abs() < 1e-4, + "Pearl-A bootstrap: prior_wins expected 10.0, got {pw}"); + assert!((pl - 2.0).abs() < 1e-4, + "Pearl-A bootstrap: prior_losses expected 2.0, got {pl}"); + assert!((psw - 0.50).abs() < 1e-4, + "Pearl-A bootstrap: prior_sum_wins expected 0.50, got {psw}"); + assert!((psl - 0.08).abs() < 1e-4, + "Pearl-A bootstrap: prior_sum_losses expected 0.08, got {psl}"); + } + + /// Test 2 — No realized trades → ISV slots preserved bit-exactly. + /// + /// All envs have zero PS_KELLY_* fields → total_wins + total_losses == 0. + /// Kernel guard skips the EMA update; both slots stay at the seeded + /// values (NOT sentinels — to verify the guard preserves arbitrary + /// pre-existing state). + #[test] + #[ignore = "requires GPU"] + fn kelly_priors_no_realized_trades_preserves_isv() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + // Seed with non-sentinel values to verify they survive the guard. + const SEED_WINS: f32 = 7.0; + const SEED_LOSSES: f32 = 5.0; + const SEED_SUM_WINS: f32 = 0.25; + const SEED_SUM_LOSSES: f32 = 0.15; + + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[KELLY_PRIOR_WINS_INDEX] = SEED_WINS; + isv[KELLY_PRIOR_LOSSES_INDEX] = SEED_LOSSES; + isv[KELLY_PRIOR_SUM_WINS_INDEX] = SEED_SUM_WINS; + isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SEED_SUM_LOSSES; + + // 3 envs with no realized trades. + let n_envs: usize = 3; + let ps = vec![0.0_f32; n_envs * PS_STRIDE]; + + let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps"); + ps_buf.write_from_slice(&ps); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_kelly_priors( + &stream, &kernel, + ps_buf.dev_ptr, n_envs as i32, PS_STRIDE as i32, isv_buf.dev_ptr, + KELLY_PRIOR_WINS_INDEX as i32, + KELLY_PRIOR_LOSSES_INDEX as i32, + KELLY_PRIOR_SUM_WINS_INDEX as i32, + KELLY_PRIOR_SUM_LOSSES_INDEX as i32, + SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + assert_eq!(result[KELLY_PRIOR_WINS_INDEX], SEED_WINS, + "prior_wins must be preserved bit-exactly when no realized trades"); + assert_eq!(result[KELLY_PRIOR_LOSSES_INDEX], SEED_LOSSES, + "prior_losses must be preserved bit-exactly when no realized trades"); + assert_eq!(result[KELLY_PRIOR_SUM_WINS_INDEX], SEED_SUM_WINS, + "prior_sum_wins must be preserved bit-exactly when no realized trades"); + assert_eq!(result[KELLY_PRIOR_SUM_LOSSES_INDEX], SEED_SUM_LOSSES, + "prior_sum_losses must be preserved bit-exactly when no realized trades"); + } + + /// Test 3 — Bounds enforcement: extreme aggregate clamps to upper/lower bounds. + /// + /// One env with huge counts (1e6 wins, 1e6 losses) and tiny sums (1e-9). + /// Pre-EMA clamp: wins/losses → 100 (count_max), sums → 0.001 (sum_min). + /// Cold-start (sentinels): Pearl-A REPLACES with clamped targets. + #[test] + #[ignore = "requires GPU"] + fn kelly_priors_bounds_clamp_extreme() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[KELLY_PRIOR_WINS_INDEX] = SENTINEL_KELLY_PRIOR_WINS; + isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES; + isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS; + isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES; + + // Single env with extreme aggregates. + let wins_per_env = [1.0e6_f32]; + let losses_per_env = [1.0e6_f32]; + let sum_wins_per_env = [1.0e-9_f32]; + let sum_losses_per_env = [1.0e-9_f32]; + + let ps = build_portfolio_state(&wins_per_env, &losses_per_env, + &sum_wins_per_env, &sum_losses_per_env); + + let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps"); + ps_buf.write_from_slice(&ps); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_kelly_priors( + &stream, &kernel, + ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr, + KELLY_PRIOR_WINS_INDEX as i32, + KELLY_PRIOR_LOSSES_INDEX as i32, + KELLY_PRIOR_SUM_WINS_INDEX as i32, + KELLY_PRIOR_SUM_LOSSES_INDEX as i32, + SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let pw = result[KELLY_PRIOR_WINS_INDEX]; + let pl = result[KELLY_PRIOR_LOSSES_INDEX]; + let psw = result[KELLY_PRIOR_SUM_WINS_INDEX]; + let psl = result[KELLY_PRIOR_SUM_LOSSES_INDEX]; + + // 1e6 wins → clamped to count_max=100. + assert!((pw - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4, + "prior_wins must clamp to count_max=100 on extreme; got {pw}"); + assert!((pl - KELLY_PRIOR_COUNT_MAX).abs() < 1e-4, + "prior_losses must clamp to count_max=100 on extreme; got {pl}"); + // 1e-9 sums → clamped to sum_min=0.001. + assert!((psw - KELLY_PRIOR_SUM_MIN).abs() < 1e-6, + "prior_sum_wins must clamp to sum_min=0.001 on tiny; got {psw}"); + assert!((psl - KELLY_PRIOR_SUM_MIN).abs() < 1e-6, + "prior_sum_losses must clamp to sum_min=0.001 on tiny; got {psl}"); + } + + /// Test 4 — Slow EMA blend after bootstrap (α=0.005). + /// + /// Pre-seed prior_wins=10.0 (NOT sentinel — bootstrap path NOT taken). + /// Single env with 5 wins. Target = 5 (in range). EMA blend: + /// (1 - 0.005) × 10 + 0.005 × 5 = 0.995 × 10 + 0.025 = 9.975. + #[test] + #[ignore = "requires GPU"] + fn kelly_priors_slow_ema_after_bootstrap() { + let stream = make_stream(); + let kernel = load_kernel(&stream); + + const ISV_DIM: usize = 1024; + const SEEDED_WINS: f32 = 10.0; + // Other slots stay at sentinel — only test EMA on prior_wins. + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[KELLY_PRIOR_WINS_INDEX] = SEEDED_WINS; + isv[KELLY_PRIOR_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_LOSSES; + isv[KELLY_PRIOR_SUM_WINS_INDEX] = SENTINEL_KELLY_PRIOR_SUM_WINS; + isv[KELLY_PRIOR_SUM_LOSSES_INDEX] = SENTINEL_KELLY_PRIOR_SUM_LOSSES; + + // Single env with 5 wins. + let wins_per_env = [5.0_f32]; + let losses_per_env = [1.0_f32]; + let sum_wins_per_env = [0.10_f32]; + let sum_losses_per_env = [0.02_f32]; + + let ps = build_portfolio_state(&wins_per_env, &losses_per_env, + &sum_wins_per_env, &sum_losses_per_env); + + let ps_buf = unsafe { MappedF32Buffer::new(ps.len()) }.expect("alloc ps"); + ps_buf.write_from_slice(&ps); + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) }.expect("alloc isv"); + isv_buf.write_from_slice(&isv); + + launch_kelly_priors( + &stream, &kernel, + ps_buf.dev_ptr, 1, PS_STRIDE as i32, isv_buf.dev_ptr, + KELLY_PRIOR_WINS_INDEX as i32, + KELLY_PRIOR_LOSSES_INDEX as i32, + KELLY_PRIOR_SUM_WINS_INDEX as i32, + KELLY_PRIOR_SUM_LOSSES_INDEX as i32, + SENTINEL_KELLY_PRIOR_WINS, SENTINEL_KELLY_PRIOR_LOSSES, + SENTINEL_KELLY_PRIOR_SUM_WINS, SENTINEL_KELLY_PRIOR_SUM_LOSSES, + KELLY_PRIOR_COUNT_MIN, KELLY_PRIOR_COUNT_MAX, + KELLY_PRIOR_SUM_MIN, KELLY_PRIOR_SUM_MAX, + KELLY_PRIOR_EMA_ALPHA, + ); + + let result = isv_buf.read_all(); + let pw = result[KELLY_PRIOR_WINS_INDEX]; + // EMA: 0.995 × 10 + 0.005 × 5 = 9.975. + let expected = (1.0_f32 - KELLY_PRIOR_EMA_ALPHA) * SEEDED_WINS + + KELLY_PRIOR_EMA_ALPHA * 5.0_f32; + assert!( + (pw - expected).abs() < 1e-4, + "Slow EMA blend (α={KELLY_PRIOR_EMA_ALPHA}): expected {expected}, got {pw}" + ); + } +} + // ═══════════════════════════════════════════════════════════════════════════ // Test B.8: Layout-fingerprint regression for direction Q-head input bump. // diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 0e7d3bc84..c26435324 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -7992,3 +7992,150 @@ P0-A made `REWARD_POS_CAP` adaptive via `ISV[REWARD_POS_CAP_ADAPTIVE_INDEX=452]` 1. **The legacy `compute_drawdown_penalty` is now partially adaptive but the SP15 `sp15_dd_penalty` quadratic path is fully ISV-driven via slots 420-421**. The two paths still coexist with different shapes (legacy: linear ramp scaled by POS_CAP; SP15: quadratic gated by λ_dd). The P1 batch already flagged this as future work. P0-A-downstream does not retire the legacy path — that's a separate refactor. 2. **`MIN_HOLD_PENALTY_RATIO = 0.6f` is a magic number at the call site** rather than a `#define` in `state_layout.cuh`. Kept local because it's strictly a relationship between two existing constants (3.0 / 5.0 = 0.6), not a new tunable signal. If a future batch lifts the ratio itself to ISV, the constant should move to `state_layout.cuh` first. + +## Class A P1-Producer batch — adaptive Bayesian Kelly priors (2026-05-08) + +The Class A P1 batch (commit `c4b6d6ef2`) wired Item 4 (var_floor q_gap-only adaptive) and explicitly deferred Items 1-3 to "a producer batch" because they required new ISV slot allocations. This entry lands one of those deferred items: the adaptive Bayesian Kelly priors. Item 2 (MIN_HOLD_TEMPERATURE EMA) was deferred separately — the audit-spec was wrong about the consumer site (see Concerns below). + +### Background + +`kelly_cap_update_kernel.cu:39-42` and `trade_physics.cuh::kelly_position_cap:304-307` both ran with hardcoded Bayesian priors: + +```c +const float prior_wins = 2.0f; +const float prior_losses = 2.0f; +const float prior_sum_wins = 0.01f; +const float prior_sum_losses = 0.01f; +``` + +These are static "I don't know" priors frozen at sprint-1 defaults. Conceptually, Bayesian priors encode *prior belief* about the Kelly distribution before observing real trades. A static prior IS literally the prior — but it's also constant across folds, ignoring everything prior folds learned about the realized win/loss distribution shape. + +The fix is an ISV-driven slow-EMA prior fed from the realized Kelly stats already aggregated in `portfolio_state[n_envs, PS_STRIDE]` (the same buffer `kelly_cap_update_kernel` reads from at the same per-epoch boundary). On the first fold, the sentinels match the pre-P1-Producer hardcoded values for bit-identical cold-start; once the first valid observation lands, the slow EMA blends in learned beliefs from realized distribution shape; on subsequent folds, the cold-start prior is no longer "I don't know" but "what we learned from the prior fold's distribution shape, slow-discounted." + +### Slots allocated + +| Slot | Constant | Sentinel | Bounds | Replaces | +|-----:|----------|---------:|--------|----------| +| 454 | `KELLY_PRIOR_WINS_INDEX` | 2.0 | [0.5, 100] | `prior_wins=2.0f` | +| 455 | `KELLY_PRIOR_LOSSES_INDEX` | 2.0 | [0.5, 100] | `prior_losses=2.0f` | +| 456 | `KELLY_PRIOR_SUM_WINS_INDEX` | 0.01 | [0.001, 1.0] | `prior_sum_wins=0.01f` | +| 457 | `KELLY_PRIOR_SUM_LOSSES_INDEX` | 0.01 | [0.001, 1.0] | `prior_sum_losses=0.01f` | + +`ISV_TOTAL_DIM` bumped 454 → 458. Layout fingerprint regenerated. + +Bounds rationale (Category-1 dimensional safety per `feedback_isv_for_adaptive_bounds.md`): +- Counts in [0.5, 100]: below 0.5 the prior is effectively absent (kelly_f swings on first real trade); above 100 the prior dominates 10+ real trades (the maturity threshold for warmup_floor blend in `kelly_position_cap`). +- Sums in [0.001, 1.0]: below 1e-3 hits the same numerical floor as `fmaxf(avg_loss, 0.0001f)` in the Kelly formula; above 1.0 implies a per-trade win/loss > 100% of equity which would be risk-management-failure territory. + +### Producer + +`crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu` (NEW, 250 LOC): + +- Single-block 256-thread kernel; block-tree-reduce in shmem (no `atomicAdd` per `feedback_no_atomicadd.md`). +- Stride-based sweep accumulates four sums across all envs from the four `PS_KELLY_*` fields per env's portfolio_state row (PS_STRIDE=43 per state_layout.cuh). +- Phase 2 (thread 0): pre-EMA bilateral clamp on each target slot, then per-slot Pearl-A first-observation bootstrap (sentinel match within 1e-6 → REPLACE; otherwise EMA blend), then defensive post-EMA bilateral clamp. +- α=0.005 slow EMA (per-fold cadence; the Bayesian prior is a *prior belief* and should change slowly across folds — a fast-moving "prior" would just be a noisy Kelly estimate). Slower than `reward_cap_update`'s α=0.01. +- Cold-start fallback: `total_wins + total_losses == 0` → kernel guard skips the EMA update and ISV slots remain unchanged (sentinel persists; consumers fall back to the `state_layout.cuh::KELLY_PRIOR_*_DEFAULT` macros via the NULL-tolerant + range guard pattern). + +### Consumers migrated atomically + +**Site A** — `kelly_cap_update_kernel.cu:39-42`: + +```c +// Before: +const float prior_wins = 2.0f; /* + 3 more */ + +// After: +const float prior_wins = kelly_prior_or_default( + isv[KELLY_PRIOR_WINS_INDEX], KELLY_PRIOR_WINS_DEFAULT, 0.5f, 100.0f); +/* + 3 more — same pattern */ +``` + +The `kelly_prior_or_default` device helper returns the ISV slot value if it's within the dimensional-safety range, else the `state_layout.cuh::KELLY_PRIOR_*_DEFAULT` macro (which equals the pre-P1-Producer hardcoded value). Cold-start (sentinel) is in-range so the macro path only fires for malformed prior state. + +**Site B** — `trade_physics.cuh::kelly_position_cap:304-307`: + +`kelly_position_cap` gained an `isv_signals_ptr` parameter (NULL-tolerant). The function reads the four ISV slots and falls back to `KELLY_PRIOR_*_DEFAULT` macros under either `isv_signals_ptr == NULL` (validation envs that don't pass the bus through) OR out-of-range (cold-start sentinel + malformed state). `apply_kelly_cap` (the only caller of `kelly_position_cap`) gained the same parameter and passes it through. The single call site in `unified_env_step_core` (line 898) already had `isv_signals_ptr` — it's the same parameter already used for SP5 Pearl 6 `kelly_f_smooth` and SP9 Fix 37 `kelly_warmup_floor_sp9` just above. No new parameter threading needed. + +### Producer launch site + +Wired in `training_loop.rs:393` RIGHT BEFORE `launch_kelly_cap_update`: + +```rust +if let Err(e) = fused.trainer().launch_kelly_bayesian_priors_update(ps_dev_ptr, n_envs) { + tracing::warn!(epoch, "launch_kelly_bayesian_priors_update failed (non-fatal): {e}"); +} +if let Err(e) = fused.launch_kelly_cap_update(ps_dev_ptr, n_envs) { + /* ... */ +} +``` + +Order matters: the priors producer MUST run before the cap update so the cap kernel sees the freshly-blended priors. + +### State reset registry + dispatch arms (C.10 lesson) + +Four FoldReset entries (`sp14_p1_kelly_prior_*`) added to `state_reset_registry.rs`, plus four matching dispatch arms in `training_loop.rs::reset_named_state` that write the sentinels (2.0/2.0/0.01/0.01) at fold boundary. Without these dispatch arms, the FoldReset entry exists but no actual reset fires — the C.10 lesson where slot drifts across folds and the layout-fingerprint smoke test eventually catches it as a runtime crash. + +The `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test confirms all four new entries have arms. + +### Sites modified + +| File | LOC delta | Change | +|------|-----------|--------| +| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | +97 / -0 | 4 new slot constants + sentinels + bounds + EMA α + 2 new tests | +| `crates/ml/src/cuda_pipeline/state_layout.cuh` | +37 / -0 | C #define mirrors for slots + sentinels + defaults | +| `crates/ml/src/cuda_pipeline/kelly_bayesian_priors_update_kernel.cu` | +250 / -0 | NEW producer kernel | +| `crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu` | +37 / -8 | Consumer migration: ISV reads + helper for cold-start fallback | +| `crates/ml/src/cuda_pipeline/trade_physics.cuh` | +63 / -9 | `kelly_position_cap` + `apply_kelly_cap` gain `isv_signals_ptr` parameter; consumer migration | +| `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` | +120 / -1 | `KellyBayesianPriorsUpdateOps` struct + `new()` + `launch()` | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | +90 / -0 | Cubin static + struct field + constructor + `launch_kelly_bayesian_priors_update` wrapper | +| `crates/ml/build.rs` | +21 / -0 | Manifest entry for new cubin | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +35 / -0 | 4 FoldReset entries | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +63 / -0 | 4 dispatch arms + producer launch wiring | +| `crates/ml/tests/sp14_oracle_tests.rs` | +390 / -0 | 4 GPU-gated oracle tests | +| `docs/dqn-wire-up-audit.md` | +90 | This entry | + +### Architectural decisions + +1. **No new buffers**: producer reads from existing `portfolio_state[n_envs, PS_STRIDE]`, the same buffer `kelly_cap_update_kernel` already aggregates Kelly stats from at the same boundary. Symmetry preserved; no buffer-allocation surface area added. +2. **Per-fold-end cadence (= per-epoch boundary in current scheduling)**: matches `kelly_cap_update`. The producer runs RIGHT BEFORE `kelly_cap_update` so that consumer sees the freshly-blended priors. Same ordering pattern as `aux_horizon_update_chain` runs before downstream consumers. +3. **α=0.005 (slower than reward_cap's α=0.01)**: Bayesian prior is a *prior belief* and should change slowly across folds. A fast-moving "prior" would just be a noisy Kelly estimate, defeating the purpose. +4. **NULL-tolerant + range guard on consumer side**: mirrors the existing `kelly_f_smooth` / `kelly_warmup_floor_sp9` patterns in `unified_env_step_core`. Validation envs that don't pass the bus through fall back to static defaults; in-range sentinels also fall back (within tolerance of the pre-P1-Producer hardcoded values). Single source of truth for the adaptive priors; consumers never re-derive bounds. +5. **Atomic per-batch commit (Item 1 of audit) — Item 2 deferred**: the audit asked for two items combined; Item 1 is independently completable, Item 2's audit-spec was wrong (see Concerns). Per `feedback_no_partial_refactor`, Item 1 is committed atomically with all consumers migrated. + +### Verification + +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` — clean (warnings unchanged from prior baseline). +- `cargo build -p ml --release` — clean. +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --lib --release sp14` — 10/10 pass (including new `sp14_p1_kelly_prior_slot_layout_locked` + `all_sp14_p1_slots_fit_within_isv_total_dim`). +- `cargo test -p ml --lib --release dispatch` — `every_fold_and_soft_reset_entry_has_dispatch_arm` passes (C.10 regression closed for the 4 new entries). +- `cargo test -p ml --test sp14_oracle_tests --release layout_fingerprint` — `layout_fingerprint_bumps_after_sp14_wire` passes. +- 4 new oracle tests in `sp14_oracle_tests.rs::sp14_p1_kelly_priors_gpu` (#[ignore = "requires GPU"]): Pearl-A bootstrap, no-realized-trades preserves ISV, bounds clamp, slow EMA after bootstrap. GPU smoke runs at deploy time. + +### Cumulative WR-plateau fix series (continued) + +- Class C bug 1 + P0-B (`8f218cab2`): replay buffer intent→realized + Kelly warmup floor wiring. +- P0-C (`316db416b`): MIN_HOLD_TARGET adaptive from AVG_WIN_HOLD_TIME. +- P0-A (`394de7d43`): REWARD_POS/NEG_CAP adaptive producer (slots 452-453). +- P1 wiring (`c4b6d6ef2`): var_floor q_gap-only adaptive (1 of 4 wireable, 3 deferred for slot allocation). +- P0-A-downstream (`657972a4b`): DD penalty + MIN_HOLD_PENALTY_MAX scale to POS_CAP_ADAPTIVE. +- **P1-Producer (this commit): adaptive Bayesian Kelly priors (slots 454-457).** + +### Concerns + +1. **Item 2 (MIN_HOLD_TEMPERATURE EMA) DEFERRED — audit-spec error**: The prompt described "MIN_HOLD_TEMPERATURE = 0.5f hardcoded somewhere" but the actual code has `MIN_HOLD_TEMPERATURE_{START=50.0f, END=5.0f, DECAY=20.0f}` defined in `state_layout.cuh:270-272` and consumed via the **per-epoch annealing schedule** `min_hold_temperature_for_epoch(epoch)` in `training_loop.rs:68-73`: + + ```rust + 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() + } + ``` + + This is already an epoch-driven adaptive schedule (50→5 over training) per the SP12 v3 design. The kernel takes T as a runtime scalar specifically to enable Phase 2 ISV-driven lift "without recompiling cubin" per the design comment in the same function — Phase 2 is anticipated, but the *signal* that should drive it is a separate spec decision. Per `feedback_no_quickfixes.md` (every issue gets a proper fix per established patterns), guessing at the right Phase 2 signal would not be principled — the right move is to defer until a follow-up spec defines the signal (likely a feedback control on observed hold-length distribution vs target band, but that's a design decision, not a wiring fix). Reporting back per the prompt's hard rule #8. + +2. **Layout-fingerprint bump invalidates all pre-P1-Producer DQN checkpoints**: `LAYOUT_FINGERPRINT_CURRENT` regenerated from the new seed string with `KELLY_PRIOR_*=454..457` and `ISV_TOTAL_DIM=458`. Old checkpoints fail-fast on load (intentional, per spec §4.A.2 — no migration path). New training runs only. + +3. **Validation envs receive the bus through `isv_signals_ptr`**: the migration extends `apply_kelly_cap` / `kelly_position_cap` signatures with `isv_signals_ptr`. NULL-tolerant for safety, but production code should pass the bus through so val and train see the same priors. `unified_env_step_core` (line 898) already does — same parameter as `kelly_f_smooth` / `kelly_warmup_floor_sp9` paths. No `backtest_env_kernel.cu` direct callers verified — the only caller in the codebase is `unified_env_step_core` (verified by `grep -nE "apply_kelly_cap\\("`).