feat: wire temporal pipeline + fix hold + add aux_frequency

1. Temporal pipeline in training forward (submit_forward_ops_main):
   - mamba2_step: temporal scan enriches h_s2 with history
   - compute_predictive_coding_loss: temporal smoothness
   - apply_regime_dropout: regime-conditioned dropout
   - launch_isv_temporal_route: per-feature temporal weights
   - risk_budget_forward: risk budget from h_s2
   These were only in the monitoring path (reduce_current_q_stats),
   never in the actual training forward. The model trained without
   any temporal context.

2. ISV signal update after each adam step:
   update_isv_signals() called after mamba2 backward, reads pinned
   loss/grad_norm/Q-mean and updates the 12-element ISV vector.
   Was never called during training — ISV signals stayed at zero.

3. Backtest min_hold fixed:
   eval_min_hold was hardcoded 0 — no hold enforcement in validation.
   Now uses config.min_hold_bars.

4. Backtest ISV signals:
   Passes frozen ISV signals from last training step to validation
   backtest for adaptive hold enforcement.

5. aux_frequency parameter (default 4):
   IQL, IQN, attention, CQL run every 4th step instead of every step.
   ~4x faster epochs. graph_forward still runs every step.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

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<i32>,
@@ -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}"

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) ─

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)

crates/ml/src/trainers/dqn/fused_training.rs

@@ -184,6 +184,8 @@ pub(crate) struct FusedTrainingCtx {
     stream: Arc<cudarc::driver::CudaStream>,
     /// 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.

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.