diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index a5b8c80f6..392d283e5 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -282,6 +282,8 @@ pub struct GpuBacktestEvaluator { config: GpuBacktestConfig, /// Precomputed annualization factor: sqrt(bars_per_day * 252). annualization_factor: f32, + /// ISV signals device pointer for adaptive hold. 0 = NULL (static hold). + isv_signals_ptr: u64, /// Actions buffer for forward kernel output (reused across steps). forward_actions_buf: CudaSlice, @@ -532,6 +534,7 @@ impl GpuBacktestEvaluator { portfolio_dim: PORTFOLIO_AND_MTF_DIM, state_dim, annualization_factor: (config.bars_per_day * config.trading_days_per_year).sqrt(), + isv_signals_ptr: 0, config, forward_actions_buf, action_select_kernel: None, @@ -548,6 +551,12 @@ impl GpuBacktestEvaluator { }) } + /// Set ISV signals pointer for adaptive hold in validation backtest. + /// Pass the frozen ISV signals from the last training step. + pub fn set_isv_signals_ptr(&mut self, ptr: u64) { + self.isv_signals_ptr = ptr; + } + // ── Private helpers ─────────────────────────────────────────────────────── /// Build the initial portfolio state vector. @@ -955,7 +964,7 @@ impl GpuBacktestEvaluator { shared_mem_bytes: 0, }; let null_portfolio: u64 = 0; - let eval_min_hold: i32 = 0; + let eval_min_hold: i32 = self.config.min_hold_bars; let eval_max_pos: f32 = 0.0; unsafe { self.stream @@ -981,7 +990,7 @@ impl GpuBacktestEvaluator { .arg(&1.0_f32) // eps_urg_mult (no-op at eps=0) .arg(&(chunk_start as i32)) // timestep for stateless Philox RNG .arg(&0u64) // per_sample_epsilon: NULL → use cosine schedule - .arg(&0u64) // isv_signals_ptr: NULL → static hold (backtest) + .arg(&self.isv_signals_ptr) // ISV signals for adaptive hold .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "chunked action_select chunk_start={chunk_start}: {e}" diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 210379cfa..a684de44b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -7874,14 +7874,36 @@ impl GpuDqnTrainer { // ── 1. Forward (cuBLAS SGEMM — Pass 1 + 2, no Pass 3) ────────── self.launch_cublas_forward()?; - // ── 1b. ISV + plan heads inside graph capture (same as working 11b1a1ca9) ── + // ── 1b. Temporal + ISV + plan pipeline (full training forward) ── { let batch_size = self.config.batch_size; + + // Mamba2: temporal scan enriches h_s2 with history + self.mamba2_step(batch_size)?; + + // Predictive coding: temporal smoothness loss on enriched trunk + self.compute_predictive_coding_loss(batch_size)?; + + // Regime-conditioned dropout on h_s2 + self.apply_regime_dropout(batch_size, true)?; + + // ISV forward: encoder MLP → branch gate + gamma mod self.launch_isv_forward()?; + + // ISV temporal routing: per-feature temporal_weight for next mamba2 + self.launch_isv_temporal_route()?; + + // ISV feature gate: modulate h_s2 features by regime embedding self.launch_isv_feature_gate(batch_size)?; + + // Recursive confidence: predict own TD-error from h_s2 self.launch_recursive_confidence_forward(batch_size)?; + + // Trade plan: h_s2 → 6 plan parameters self.launch_trade_plan_forward(batch_size)?; - // plan_noise_inject runs outside graph (pre-replay) — see fused_training.rs + + // Risk budget: h_s2 → R ∈ (0,1) before Q-value computation + self.risk_budget_forward(batch_size)?; } // ── 2+3. Loss + gradient (blended MSE + C51 via c51_alpha ramp) ─ diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs index 5740766ff..9a3c45482 100644 --- a/crates/ml/src/trainers/dqn/config.rs +++ b/crates/ml/src/trainers/dqn/config.rs @@ -565,6 +565,11 @@ pub struct DQNHyperparameters { /// Default: 5 (about 5 minutes for 1-min bars). Range: [2, 20]. pub min_hold_bars: usize, + /// Auxiliary ops (IQL, IQN, attention, CQL) run every Nth training step. + /// Default: 4 (aux ops every 4th step, ~4× faster epochs). + /// 1 = every step (full quality), higher = faster but less frequent aux updates. + pub aux_frequency: usize, + // P2-B Enhancement: Cash reserve requirement /// Cash reserve requirement as percentage of portfolio value (0-100) /// Default: 0.0 (no reserve, backward compatible) @@ -1344,6 +1349,7 @@ impl DQNHyperparameters { initial_capital: 100_000.0, // $100K default bars_per_day: 390.0, // 1-minute bars (6.5h × 60min) min_hold_bars: 5, // 5-bar minimum prevents coin-flip exits (1-2 bar churn) + aux_frequency: 4, // aux ops every 4th step (~4× faster epochs) // P2-B Enhancement cash_reserve_percent: 0.0, // Default: no reserve (backward compatible) diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 822a34cb1..9f3b2b81d 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -184,6 +184,8 @@ pub(crate) struct FusedTrainingCtx { stream: Arc, /// Batch size at creation time -- must match `current_batch_size` to reuse CUDA Graph. batch_size: usize, + /// Aux ops run every Nth step (IQL, IQN, attention, CQL). + aux_frequency: usize, /// Steps since last flat buffer sync (deferred to epoch boundary). steps_since_varmap_sync: usize, /// GPU-resident copy of best-epoch params. Pure DtoD — no Candle, no safetensors. @@ -731,6 +733,7 @@ impl FusedTrainingCtx { } }, phase_batch_count: 0, + aux_frequency: hyperparams.aux_frequency.max(1), graph_aux: None, graph_spectral: None, graph_mega: None, @@ -966,13 +969,16 @@ impl FusedTrainingCtx { } } - // ── Step 3: Auxiliary ops (graph_aux via cudarc wrapper) ──────── - if self.graph_aux.is_some() { - self.pre_replay_state_update(agent); - self.graph_aux.as_ref().unwrap().0.launch() - .map_err(|e| anyhow::anyhow!("graph_aux replay: {e}"))?; - } else { - self.submit_aux_ops(agent, gpu_batch)?; + // ── Step 3: Auxiliary ops (every Nth step for speed) ──────── + let aux_freq = self.aux_frequency; + if step % aux_freq == 0 { + if self.graph_aux.is_some() { + self.pre_replay_state_update(agent); + self.graph_aux.as_ref().unwrap().0.launch() + .map_err(|e| anyhow::anyhow!("graph_aux replay: {e}"))?; + } else { + self.submit_aux_ops(agent, gpu_batch)?; + } } if let Some(ref pe) = self.phase_events { PhaseEvents::record(pe.fwd_bwd_end, cu_stream); @@ -1031,6 +1037,13 @@ impl FusedTrainingCtx { self.mamba2_backward(self.batch_size)?; self.step_mamba2_adam()?; + // ── Step 5a-isv: Update ISV signals from training scalars ──────── + // Reads pinned loss/grad_norm/Q-mean (written by GPU during adam), + // updates the 12-element ISV signal vector + history buffer. + // Must run AFTER adam (scalars valid) and BEFORE next forward (ISV used). + self.trainer.update_isv_signals() + .map_err(|e| anyhow::anyhow!("ISV signal update: {e}"))?; + // ── Step 5a-risk: Risk branch gentle weight decay ─────────────── // MVP training signal: decay toward initial values to prevent R // collapse. Full BPTT backward deferred. diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index a441e5c87..a9635889c 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -501,6 +501,11 @@ impl DQNTrainer { self.gpu_evaluator = Some(evaluator); } + // Wire ISV signals to evaluator for adaptive hold in validation + if let (Some(ref mut eval), Some(ref fused)) = (&mut self.gpu_evaluator, &self.fused_ctx) { + eval.set_isv_signals_ptr(fused.isv_signals_dev_ptr()); + } + // ── Extract weight pointers + v_range, then split borrow for fused_ctx ── // DuelingWeightSet / BranchingWeightSet contain raw u64 device pointers. // We extract stable pointers to them, then take &mut fused_ctx for QValueProvider.