From 11b1a1ca9f520d1718a53477abd5f87ab600b796 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 17 Apr 2026 02:05:20 +0200 Subject: [PATCH] =?UTF-8?q?fix:=20zen=20precommit=20=E2=80=94=202=20critic?= =?UTF-8?q?als=20+=203=20lows=20resolved?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CRITICAL: ISV+plan heads wired into submit_forward_ops_main (training). CRITICAL: backtest single-step 3-branch→4-branch action encoding. LOW: Remove dead bar_start_ns/bar_duration_ns from MicrostructureState. LOW: Pre-allocated eps_host_buf replaces Vec allocation in hot path. LOW: training_sharpe + val_loss added to metrics additional_metrics. Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml-features/src/ofi_calculator.rs | 7 +- .../src/cuda_pipeline/gpu_action_selector.rs | 8 +- pal_precommit.changeset | 4742 +++++++++++++++++ 3 files changed, 4749 insertions(+), 8 deletions(-) create mode 100644 pal_precommit.changeset diff --git a/crates/ml-features/src/ofi_calculator.rs b/crates/ml-features/src/ofi_calculator.rs index 25f87e12e..295423cd5 100644 --- a/crates/ml-features/src/ofi_calculator.rs +++ b/crates/ml-features/src/ofi_calculator.rs @@ -684,11 +684,8 @@ const fn safe_clip(value: f64, min: f64, max: f64) -> f64 { /// | 18 | OFI Acceleration | /// | 19 | Toxicity Gradient | #[derive(Debug, Clone)] -#[allow(dead_code)] pub struct MicrostructureState { - // --- bar timing --- - bar_start_ns: u64, - bar_duration_ns: u64, + /// Bar midpoint for intra-bar momentum half-split (used by update_snapshot). bar_mid_ns: u64, // --- [8] OFI Trajectory: online linear regression of OFI_L1 over time --- @@ -758,8 +755,6 @@ impl MicrostructureState { /// Create a new state for a bar starting at `bar_start_ns` with given duration. pub fn new(bar_start_ns: u64, bar_duration_ns: u64) -> Self { Self { - bar_start_ns, - bar_duration_ns, bar_mid_ns: bar_start_ns + bar_duration_ns / 2, ofi_linreg_n: 0.0, diff --git a/crates/ml/src/cuda_pipeline/gpu_action_selector.rs b/crates/ml/src/cuda_pipeline/gpu_action_selector.rs index 9da092462..4221ff913 100644 --- a/crates/ml/src/cuda_pipeline/gpu_action_selector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_action_selector.rs @@ -46,6 +46,8 @@ pub struct GpuActionSelector { /// Per-sample epsilon buffer for branching kernel: [max_batch_size] on GPU. /// Filled with uniform epsilon before each branching launch. epsilon_buf: CudaSlice, + /// Pre-allocated host buffer for epsilon fill — avoids Vec allocation in hot path. + eps_host_buf: Vec, } impl GpuActionSelector { @@ -71,6 +73,7 @@ impl GpuActionSelector { bonus_exposure_ptr: 0, bonus_order_ptr: 0, bonus_urgency_ptr: 0, bonus_exposure_buf: None, bonus_order_buf: None, bonus_urgency_buf: None, epsilon_buf, + eps_host_buf: vec![0.0_f32; max_batch_size], }) } @@ -151,8 +154,9 @@ impl GpuActionSelector { if batch_size > self.max_batch_size { return Err(MLError::ModelError(format!("batch_size {} exceeds max {}", batch_size, self.max_batch_size))); } // Fill per-sample epsilon buffer with uniform value (kernel reads per_sample_epsilon[idx]) // Slice to batch_size to match eps_host length (epsilon_buf allocated for max_batch_size) - let eps_host = vec![epsilon; batch_size]; - self.stream.memcpy_htod(&eps_host, &mut self.epsilon_buf.slice_mut(..batch_size)) + // Fill pre-allocated host buffer (no heap allocation in hot path) + self.eps_host_buf[..batch_size].fill(epsilon); + self.stream.memcpy_htod(&self.eps_host_buf[..batch_size], &mut self.epsilon_buf.slice_mut(..batch_size)) .map_err(|e| MLError::ModelError(format!("upload epsilon_buf: {e}")))?; let config = launch_config_1d(batch_size); let bs_i32 = batch_size as i32; diff --git a/pal_precommit.changeset b/pal_precommit.changeset new file mode 100644 index 000000000..4b69d1e41 --- /dev/null +++ b/pal_precommit.changeset @@ -0,0 +1,4742 @@ +diff --git a/config/training/dqn-localdev.toml b/config/training/dqn-localdev.toml +index 8d1a95cc7..28dddd711 100644 +--- a/config/training/dqn-localdev.toml ++++ b/config/training/dqn-localdev.toml +@@ -19,7 +19,8 @@ lr_min = 0.00001 + # Data source: "ohlcv" for local dev (may not have MBP-10 data) + data_source = "ohlcv" + # OFI feature enrichment from MBP-10 order book data (8 features: OFI L1/L5, depth imbalance, VPIN, Kyle's lambda, bid/ask slope, trade imbalance) +-# mbp10_data_dir = "test_data/futures-baseline-mbp10" ++mbp10_data_dir = "test_data/futures-baseline-mbp10" ++trades_data_dir = "test_data/futures-baseline-trades" + + [distributional] + num_atoms = 52 +diff --git a/crates/ml-features/src/ofi_calculator.rs b/crates/ml-features/src/ofi_calculator.rs +index 6b6ec390a..25f87e12e 100644 +--- a/crates/ml-features/src/ofi_calculator.rs ++++ b/crates/ml-features/src/ofi_calculator.rs +@@ -36,7 +36,7 @@ + //! assert!(features.depth_imbalance >= -1.0 && features.depth_imbalance <= 1.0); + //! ``` + +-use data::providers::databento::mbp10::Mbp10Snapshot; ++use data::providers::databento::mbp10::{BidAskPair, Mbp10Snapshot}; + use std::collections::VecDeque; + + use crate::MLError; +@@ -661,6 +661,491 @@ const fn safe_clip(value: f64, min: f64, max: f64) -> f64 { + value.max(min).min(max) + } + ++// --------------------------------------------------------------------------- ++// Tick Microstructure Intelligence — 12 new features (indices 8-19) ++// --------------------------------------------------------------------------- ++ ++/// Running microstructure statistics computed O(1) per tick. ++/// ++/// Produces 12 features that extend the base 8 OFI features: ++/// ++/// | Idx | Name | ++/// |-----|---------------------------| ++/// | 8 | OFI Trajectory | ++/// | 9 | Realized Variance | ++/// | 10 | Hawkes Trade Intensity | ++/// | 11 | Book Pressure Gradient | ++/// | 12 | Spread Dynamics | ++/// | 13 | Aggression Ratio | ++/// | 14 | Queue Depletion Asymmetry | ++/// | 15 | Order Count Flux | ++/// | 16 | Intra-Bar Momentum | ++/// | 17 | Microstructure Regime | ++/// | 18 | OFI Acceleration | ++/// | 19 | Toxicity Gradient | ++#[derive(Debug, Clone)] ++#[allow(dead_code)] ++pub struct MicrostructureState { ++ // --- bar timing --- ++ bar_start_ns: u64, ++ bar_duration_ns: u64, ++ bar_mid_ns: u64, ++ ++ // --- [8] OFI Trajectory: online linear regression of OFI_L1 over time --- ++ ofi_linreg_n: f64, ++ ofi_linreg_sum_x: f64, ++ ofi_linreg_sum_y: f64, ++ ofi_linreg_sum_xy: f64, ++ ofi_linreg_sum_xx: f64, ++ ++ // --- [9] Realized Variance: sum of squared log-returns of mid price --- ++ prev_mid: f64, ++ realized_var: f64, ++ ++ // --- [10] Hawkes Trade Intensity: exponential kernel EMA --- ++ hawkes_intensity: f64, ++ last_trade_ns: u64, ++ ++ // --- [11] Book Pressure Gradient: computed from latest snapshot (stateless) --- ++ book_pressure: f64, ++ ++ // --- [12] Spread Dynamics: max/min/sum/count of spread within bar --- ++ spread_max: f64, ++ spread_min: f64, ++ spread_sum: f64, ++ spread_count: u64, ++ ++ // --- [13] Aggression Ratio: trades at ask vs total --- ++ trades_at_ask: u64, ++ total_trades: u64, ++ ++ // --- [14] Queue Depletion Asymmetry: EMA of bid/ask depletion --- ++ prev_bid_sz_l1: u32, ++ prev_ask_sz_l1: u32, ++ bid_depletion_ema: f64, ++ ask_depletion_ema: f64, ++ ++ // --- [15] Order Count Flux: ct_increase / (ct_increase + ct_decrease) --- ++ prev_total_ct: u32, ++ ct_increase_count: u64, ++ ct_decrease_count: u64, ++ ++ // --- [16] Intra-Bar Momentum: Welford online mean for half1 and half2 --- ++ half1_mean: f64, ++ half1_n: u64, ++ half2_mean: f64, ++ half2_n: u64, ++ prev_mid_momentum: f64, ++ ++ // --- [17] Microstructure Regime Score: computed from aggression/spread/thin_book --- ++ // (derived from other fields in snapshot(), no extra state needed) ++ ++ // --- [18] OFI Acceleration: EMA of OFI_L1 deltas --- ++ prev_ofi_l1: f64, ++ ofi_accel_ema: f64, ++ has_prev_ofi: bool, ++ ++ // --- [19] Toxicity Gradient: EMA of VPIN deltas --- ++ prev_vpin: f64, ++ toxicity_grad_ema: f64, ++ has_prev_vpin: bool, ++ ++ // --- snapshot count (for first-snapshot guards) --- ++ snapshot_count: u64, ++} ++ ++impl MicrostructureState { ++ /// Create a new state for a bar starting at `bar_start_ns` with given duration. ++ pub fn new(bar_start_ns: u64, bar_duration_ns: u64) -> Self { ++ Self { ++ bar_start_ns, ++ bar_duration_ns, ++ bar_mid_ns: bar_start_ns + bar_duration_ns / 2, ++ ++ ofi_linreg_n: 0.0, ++ ofi_linreg_sum_x: 0.0, ++ ofi_linreg_sum_y: 0.0, ++ ofi_linreg_sum_xy: 0.0, ++ ofi_linreg_sum_xx: 0.0, ++ ++ prev_mid: 0.0, ++ realized_var: 0.0, ++ ++ hawkes_intensity: 0.0, ++ last_trade_ns: 0, ++ ++ book_pressure: 0.0, ++ ++ spread_max: f64::NEG_INFINITY, ++ spread_min: f64::INFINITY, ++ spread_sum: 0.0, ++ spread_count: 0, ++ ++ trades_at_ask: 0, ++ total_trades: 0, ++ ++ prev_bid_sz_l1: 0, ++ prev_ask_sz_l1: 0, ++ bid_depletion_ema: 0.0, ++ ask_depletion_ema: 0.0, ++ ++ prev_total_ct: 0, ++ ct_increase_count: 0, ++ ct_decrease_count: 0, ++ ++ half1_mean: 0.0, ++ half1_n: 0, ++ half2_mean: 0.0, ++ half2_n: 0, ++ prev_mid_momentum: 0.0, ++ ++ prev_ofi_l1: 0.0, ++ ofi_accel_ema: 0.0, ++ has_prev_ofi: false, ++ ++ prev_vpin: 0.0, ++ toxicity_grad_ema: 0.0, ++ has_prev_vpin: false, ++ ++ snapshot_count: 0, ++ } ++ } ++ ++ /// Update from an MBP-10 snapshot. Called per snapshot event. ++ pub fn update_snapshot(&mut self, snapshot: &Mbp10Snapshot) { ++ const EMA_ALPHA: f64 = 0.05; ++ ++ if snapshot.levels.is_empty() { ++ return; ++ } ++ ++ let mid = snapshot.mid_price(); ++ let spread = snapshot.spread(); ++ ++ // --- [9] Realized Variance --- ++ if self.snapshot_count > 0 && self.prev_mid > 0.0 && mid > 0.0 { ++ let log_return = (mid / self.prev_mid).ln(); ++ if log_return.is_finite() { ++ self.realized_var += log_return * log_return; ++ } ++ } ++ ++ // --- [11] Book Pressure Gradient --- ++ self.book_pressure = Self::calc_book_pressure(&snapshot.levels); ++ ++ // --- [12] Spread Dynamics --- ++ if spread.is_finite() && spread >= 0.0 { ++ if spread > self.spread_max { ++ self.spread_max = spread; ++ } ++ if spread < self.spread_min { ++ self.spread_min = spread; ++ } ++ self.spread_sum += spread; ++ self.spread_count += 1; ++ } ++ ++ // --- [14] Queue Depletion Asymmetry --- ++ if self.snapshot_count > 0 { ++ let bid_sz = snapshot.levels[0].bid_sz; ++ let ask_sz = snapshot.levels[0].ask_sz; ++ ++ // Depletion = how much size decreased (positive means depletion) ++ let bid_depl = if bid_sz < self.prev_bid_sz_l1 { ++ (self.prev_bid_sz_l1 - bid_sz) as f64 ++ } else { ++ 0.0 ++ }; ++ let ask_depl = if ask_sz < self.prev_ask_sz_l1 { ++ (self.prev_ask_sz_l1 - ask_sz) as f64 ++ } else { ++ 0.0 ++ }; ++ ++ self.bid_depletion_ema = ++ EMA_ALPHA * bid_depl + (1.0 - EMA_ALPHA) * self.bid_depletion_ema; ++ self.ask_depletion_ema = ++ EMA_ALPHA * ask_depl + (1.0 - EMA_ALPHA) * self.ask_depletion_ema; ++ ++ self.prev_bid_sz_l1 = bid_sz; ++ self.prev_ask_sz_l1 = ask_sz; ++ } else { ++ self.prev_bid_sz_l1 = snapshot.levels[0].bid_sz; ++ self.prev_ask_sz_l1 = snapshot.levels[0].ask_sz; ++ } ++ ++ // --- [15] Order Count Flux --- ++ let total_ct: u32 = snapshot ++ .levels ++ .iter() ++ .map(|l| l.bid_ct + l.ask_ct) ++ .sum(); ++ if self.snapshot_count > 0 { ++ if total_ct > self.prev_total_ct { ++ self.ct_increase_count += 1; ++ } else if total_ct < self.prev_total_ct { ++ self.ct_decrease_count += 1; ++ } ++ } ++ self.prev_total_ct = total_ct; ++ ++ // --- [16] Intra-Bar Momentum (Welford online mean for each half) --- ++ if self.snapshot_count > 0 && self.prev_mid_momentum > 0.0 && mid > 0.0 { ++ let ret = (mid / self.prev_mid_momentum).ln(); ++ if ret.is_finite() { ++ // Determine which half based on timestamp ++ if snapshot.timestamp < self.bar_mid_ns { ++ self.half1_n += 1; ++ let delta = ret - self.half1_mean; ++ self.half1_mean += delta / self.half1_n as f64; ++ } else { ++ self.half2_n += 1; ++ let delta = ret - self.half2_mean; ++ self.half2_mean += delta / self.half2_n as f64; ++ } ++ } ++ } ++ self.prev_mid_momentum = mid; ++ ++ self.prev_mid = mid; ++ self.snapshot_count += 1; ++ } ++ ++ /// Update from a trade event. Called per trade. ++ pub fn update_trade(&mut self, _price: f64, _volume: u64, is_buy: bool, timestamp_ns: u64) { ++ // --- [10] Hawkes Trade Intensity --- ++ const MU: f64 = 1.0; ++ const ALPHA: f64 = 0.5; ++ const BETA: f64 = 1.0; ++ ++ if self.last_trade_ns > 0 && timestamp_ns > self.last_trade_ns { ++ let dt_sec = ++ (timestamp_ns - self.last_trade_ns) as f64 / 1_000_000_000.0; ++ let decay = (-BETA * dt_sec).exp(); ++ // Hawkes: intensity = mu + alpha * decay * (prev_intensity - mu) + alpha ++ self.hawkes_intensity = MU + (self.hawkes_intensity - MU) * decay + ALPHA; ++ } else { ++ // First trade ++ self.hawkes_intensity = MU + ALPHA; ++ } ++ // Clamp to [0, 100] ++ self.hawkes_intensity = self.hawkes_intensity.clamp(0.0, 100.0); ++ self.last_trade_ns = timestamp_ns; ++ ++ // --- [13] Aggression Ratio --- ++ self.total_trades += 1; ++ if is_buy { ++ self.trades_at_ask += 1; ++ } ++ } ++ ++ /// Update with derived OFI values. Called after OFI calculation. ++ pub fn update_ofi_derived(&mut self, ofi_l1: f64, vpin: f64) { ++ const EMA_ALPHA: f64 = 0.1; ++ ++ // --- [8] OFI Trajectory: online linear regression --- ++ // Use ofi update count as x (monotonic) ++ self.ofi_linreg_n += 1.0; ++ let x = self.ofi_linreg_n; ++ let y = ofi_l1; ++ self.ofi_linreg_sum_x += x; ++ self.ofi_linreg_sum_y += y; ++ self.ofi_linreg_sum_xy += x * y; ++ self.ofi_linreg_sum_xx += x * x; ++ ++ // --- [18] OFI Acceleration: EMA of OFI_L1 deltas --- ++ if self.has_prev_ofi { ++ let delta = ofi_l1 - self.prev_ofi_l1; ++ self.ofi_accel_ema = EMA_ALPHA * delta + (1.0 - EMA_ALPHA) * self.ofi_accel_ema; ++ } ++ self.prev_ofi_l1 = ofi_l1; ++ self.has_prev_ofi = true; ++ ++ // --- [19] Toxicity Gradient: EMA of VPIN deltas --- ++ if self.has_prev_vpin { ++ let delta = vpin - self.prev_vpin; ++ self.toxicity_grad_ema = ++ EMA_ALPHA * delta + (1.0 - EMA_ALPHA) * self.toxicity_grad_ema; ++ } ++ self.prev_vpin = vpin; ++ self.has_prev_vpin = true; ++ } ++ ++ /// Pure read — produce the 12-feature snapshot. Idempotent, non-mutating. ++ pub fn snapshot(&self) -> [f64; 12] { ++ // [8] OFI Trajectory — online linear regression slope ++ let ofi_trajectory = if self.ofi_linreg_n >= 2.0 { ++ let n = self.ofi_linreg_n; ++ let denom = n * self.ofi_linreg_sum_xx - self.ofi_linreg_sum_x * self.ofi_linreg_sum_x; ++ if denom.abs() > 1e-12 { ++ let slope = (n * self.ofi_linreg_sum_xy ++ - self.ofi_linreg_sum_x * self.ofi_linreg_sum_y) ++ / denom; ++ safe_clip(slope, -1e6, 1e6) ++ } else { ++ 0.0 ++ } ++ } else { ++ 0.0 ++ }; ++ ++ // [9] Realized Variance ++ let realized_variance = safe_clip(self.realized_var, 0.0, 1e6); ++ ++ // [10] Hawkes Trade Intensity ++ let hawkes_intensity = self.hawkes_intensity; ++ ++ // [11] Book Pressure Gradient ++ let book_pressure = safe_clip(self.book_pressure, -1.0, 1.0); ++ ++ // [12] Spread Dynamics: (max - min) / mean ++ let spread_dynamics = if self.spread_count >= 2 { ++ let mean_spread = self.spread_sum / self.spread_count as f64; ++ if mean_spread > 1e-15 { ++ safe_clip( ++ (self.spread_max - self.spread_min) / mean_spread, ++ 0.0, ++ 100.0, ++ ) ++ } else { ++ 0.0 ++ } ++ } else { ++ 0.0 ++ }; ++ ++ // [13] Aggression Ratio ++ let aggression_ratio = if self.total_trades > 0 { ++ self.trades_at_ask as f64 / self.total_trades as f64 ++ } else { ++ 0.5 // neutral default ++ }; ++ ++ // [14] Queue Depletion Asymmetry ++ let queue_depletion = { ++ let max_depl = self.bid_depletion_ema.abs().max(self.ask_depletion_ema.abs()); ++ if max_depl > 1e-12 { ++ safe_clip( ++ (self.bid_depletion_ema - self.ask_depletion_ema) / max_depl, ++ -1.0, ++ 1.0, ++ ) ++ } else { ++ 0.0 ++ } ++ }; ++ ++ // [15] Order Count Flux ++ let order_count_flux = { ++ let total = self.ct_increase_count + self.ct_decrease_count; ++ if total > 0 { ++ self.ct_increase_count as f64 / total as f64 ++ } else { ++ 0.5 ++ } ++ }; ++ ++ // [16] Intra-Bar Momentum: sign-weighted product of half means ++ let intra_bar_momentum = if self.half1_n > 0 && self.half2_n > 0 { ++ let sign1 = if self.half1_mean >= 0.0 { 1.0 } else { -1.0 }; ++ let sign2 = if self.half2_mean >= 0.0 { 1.0 } else { -1.0 }; ++ let product = sign1 * sign2 * (self.half1_mean.abs() * self.half2_mean.abs()).sqrt(); ++ safe_clip(product, -1.0, 1.0) ++ } else { ++ 0.0 ++ }; ++ ++ // [17] Microstructure Regime Score: sigmoid(aggression * spread_dyn * thin_book) ++ let regime_score = { ++ // thin_book: inverse of total depth, clamped ++ let aggr = (aggression_ratio - 0.5).abs() * 2.0; // how far from balanced ++ let thin_book = if book_pressure.abs() > 1e-12 { ++ 1.0 / (1.0 + book_pressure.abs() * 10.0) ++ } else { ++ 0.5 ++ }; ++ let raw = aggr * spread_dynamics * thin_book; ++ // sigmoid: 1 / (1 + exp(-x)) ++ 1.0 / (1.0 + (-raw).exp()) ++ }; ++ ++ // [18] OFI Acceleration ++ let ofi_acceleration = safe_clip(self.ofi_accel_ema, -1e6, 1e6); ++ ++ // [19] Toxicity Gradient ++ let toxicity_gradient = safe_clip(self.toxicity_grad_ema, -1.0, 1.0); ++ ++ [ ++ ofi_trajectory, // 8 ++ realized_variance, // 9 ++ hawkes_intensity, // 10 ++ book_pressure, // 11 ++ spread_dynamics, // 12 ++ aggression_ratio, // 13 ++ queue_depletion, // 14 ++ order_count_flux, // 15 ++ intra_bar_momentum, // 16 ++ regime_score, // 17 ++ ofi_acceleration, // 18 ++ toxicity_gradient, // 19 ++ ] ++ } ++ ++ /// Reset state for a new bar. ++ pub fn reset(&mut self, bar_start_ns: u64, bar_duration_ns: u64) { ++ *self = Self::new(bar_start_ns, bar_duration_ns); ++ } ++ ++ /// Weighted book pressure across all 10 levels. ++ /// ++ /// `sum(exp(-0.3*l) * (bid_sz - ask_sz)) / sum(bid_sz + ask_sz)` ++ fn calc_book_pressure(levels: &[BidAskPair]) -> f64 { ++ let max_levels = levels.len().min(10); ++ let mut weighted_sum = 0.0; ++ let mut total_volume = 0.0; ++ ++ for (i, level) in levels.iter().take(max_levels).enumerate() { ++ let weight = (-0.3 * i as f64).exp(); ++ weighted_sum += weight * (level.bid_sz as f64 - level.ask_sz as f64); ++ total_volume += level.bid_sz as f64 + level.ask_sz as f64; ++ } ++ ++ if total_volume < 1.0 { ++ return 0.0; ++ } ++ ++ safe_clip(weighted_sum / total_volume, -1.0, 1.0) ++ } ++} ++ ++/// Extended OFI features: base 8 + microstructure 12 = 20 features. ++#[derive(Debug, Clone)] ++pub struct ExtendedOFIFeatures { ++ /// Base 8 OFI features ++ pub base: OFIFeatures, ++ /// 12 microstructure features (indices 8-19) ++ pub micro: [f64; 12], ++} ++ ++impl ExtendedOFIFeatures { ++ /// Convert to flat 20-element array. ++ pub fn to_array(&self) -> [f64; 20] { ++ let base = self.base.to_array(); ++ let mut out = [0.0f64; 20]; ++ out[..8].copy_from_slice(&base); ++ out[8..20].copy_from_slice(&self.micro); ++ out ++ } ++ ++ /// Create zero-initialized extended features. ++ pub fn zeros() -> Self { ++ Self { ++ base: OFIFeatures::zeros(), ++ micro: [0.0; 12], ++ } ++ } ++} ++ + #[cfg(test)] + mod tests { + use super::*; +diff --git a/crates/ml/examples/precompute_features.rs b/crates/ml/examples/precompute_features.rs +index 0fcbeed2b..281174f1b 100644 +--- a/crates/ml/examples/precompute_features.rs ++++ b/crates/ml/examples/precompute_features.rs +@@ -326,7 +326,7 @@ async fn main() -> Result<()> { + + // ── Step 4: Compute OFI from MBP-10 + trades ──────────────────────────── + let t2 = Instant::now(); +- let ofi: Vec<[f64; 8]> = if let Some(ref mbp10_path) = mbp10_dir { ++ let ofi: Vec<[f64; 20]> = if let Some(ref mbp10_path) = mbp10_dir { + use ml::features::mbp10_loader::load_ofi_features_parallel; + use ml::features::trades_loader::load_trades_sync; + use ml::features::ofi_calculator::OFICalculator; +@@ -337,7 +337,7 @@ async fn main() -> Result<()> { + + if mbp10_files.is_empty() { + info!("No MBP-10 files found, OFI will be zeros"); +- vec![[0.0; 8]; n] ++ vec![[0.0; 20]; n] + } else { + // Load MBP-10 snapshots (parallel) + use data::providers::databento::dbn_parser::DbnParser; +@@ -391,33 +391,64 @@ async fn main() -> Result<()> { + }; + + // Compute per-bar OFI +- use ml::features::mbp10_loader::get_snapshots_for_timestamp; + use ml::features::trades_loader::get_trades_for_bar; ++ use ml::features::ofi_calculator::MicrostructureState; + let mut calculator = OFICalculator::new(); + let mut ofi_per_bar = Vec::with_capacity(n); + + for i in 0..n { + let bar = &all_bars[i + WARMUP]; + let bar_ts = bar.timestamp.timestamp_nanos_opt().unwrap_or(0) as u64; ++ let bar_end_ts = all_bars.get(i + WARMUP + 1) ++ .map(|b| b.timestamp.timestamp_nanos_opt().unwrap_or(0) as u64) ++ .unwrap_or(bar_ts + 60_000_000_000); ++ let bar_duration_ns = bar_end_ts - bar_ts; + +- // Feed trades for this bar window ++ let mut micro_state = MicrostructureState::new(bar_ts, bar_duration_ns); ++ ++ // Feed trades for this bar window (OFICalculator + MicrostructureState) + if let Some(ref trades) = all_trades { +- let bar_end_ts = all_bars.get(i + WARMUP + 1) +- .map(|b| b.timestamp.timestamp_nanos_opt().unwrap_or(0) as u64) +- .unwrap_or(bar_ts + 60_000_000_000); + for trade in get_trades_for_bar(trades, bar_ts, bar_end_ts) { + calculator.feed_trade(trade.price, trade.volume, trade.is_buy); ++ micro_state.update_trade(trade.price, trade.volume, trade.is_buy, trade.timestamp); + } + } + +- let window = get_snapshots_for_timestamp(&all_snapshots, bar_ts, 1); +- if let Some(snap) = window.first() { ++ // Get ALL MBP-10 snapshots within this bar window for tick-level features. ++ // Binary search for bar range: [bar_ts, bar_end_ts) ++ let snap_start = all_snapshots.partition_point(|s| s.timestamp < bar_ts); ++ let snap_end = all_snapshots.partition_point(|s| s.timestamp < bar_end_ts); ++ let bar_snapshots = &all_snapshots[snap_start..snap_end]; ++ ++ // Feed every snapshot to MicrostructureState for tick-level resolution ++ for snap in bar_snapshots { ++ micro_state.update_snapshot(snap); ++ } ++ ++ // Use the LAST snapshot in the bar for OFI calculation (matches prior behavior) ++ let last_snap = bar_snapshots.last() ++ .or_else(|| { ++ // Fallback: nearest snapshot at/after bar_ts (prior behavior) ++ all_snapshots.get(snap_start) ++ }); ++ ++ if let Some(snap) = last_snap { + match calculator.calculate(snap) { +- Ok(f) if f.is_valid() => ofi_per_bar.push(f.to_array()), +- _ => ofi_per_bar.push([0.0; 8]), ++ Ok(f) if f.is_valid() => { ++ // Feed OFI-derived values to MicrostructureState ++ micro_state.update_ofi_derived(f.ofi_level1, f.vpin); ++ ++ let arr8 = f.to_array(); ++ let micro_12 = micro_state.snapshot(); ++ let mut arr20 = [0.0_f64; 20]; ++ arr20[..8].copy_from_slice(&arr8); ++ arr20[8..20].copy_from_slice(µ_12); ++ ofi_per_bar.push(arr20); ++ } ++ _ => ofi_per_bar.push([0.0; 20]), + } + } else { +- ofi_per_bar.push([0.0; 8]); ++ ofi_per_bar.push([0.0; 20]); + } + } + +@@ -430,7 +461,7 @@ async fn main() -> Result<()> { + } + } else { + info!("No MBP-10 directory, OFI will be zeros"); +- vec![[0.0; 8]; n] ++ vec![[0.0; 20]; n] + }; + + let total_len = n; +@@ -497,7 +528,7 @@ async fn main() -> Result<()> { + println!("Bars: {}", total_len); + println!("Features: 42-dim"); + println!("Targets: 4-dim"); +- println!("OFI: 8-dim ({})", if has_ofi { "from MBP-10" } else { "zero-padded" }); ++ println!("OFI: 20-dim ({})", if has_ofi { "from MBP-10" } else { "zero-padded" }); + println!("Format: f32 (v1)"); + println!("Cache key: {}", hex_key); + println!("Output: {}", output_path.display()); +diff --git a/crates/ml/examples/train_baseline_rl.rs b/crates/ml/examples/train_baseline_rl.rs +index 934818f73..6fb4b0308 100644 +--- a/crates/ml/examples/train_baseline_rl.rs ++++ b/crates/ml/examples/train_baseline_rl.rs +@@ -591,7 +591,7 @@ fn run_training(args: &Args) -> Result> { + let targets: Vec<[f64; 4]> = aligned_bars.iter() + .map(|b| [b.close, b.close, b.close, b.close]) + .collect(); +- let ofi = vec![[0.0_f64; 8]; n]; ++ let ofi = vec![[0.0_f64; 20]; n]; + + ml::fxcache::FxCacheData { + timestamps, +diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +index 208e0b71f..62ea60d83 100644 +--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +@@ -156,7 +156,7 @@ extern "C" __global__ void backtest_env_step( + // ── Hold enforcement (shared: trade_physics.cuh) ───────────────────── + int is_last_bar = (current_step >= wlen - 1) ? 1 : 0; + target_exposure = enforce_hold(position, target_exposure, hold_time, +- min_hold_bars, is_last_bar); ++ min_hold_bars, is_last_bar, NULL); + + // ── Trailing stop (shared: trade_physics.cuh) ──────────────────────── + // Exit when profit retreats from peak. Uses 0.5% base distance. +@@ -377,7 +377,7 @@ extern "C" __global__ void backtest_env_step_batch( + /* ── Hold enforcement ─────────────────────────────────────────── */ + int is_last_bar = (current_step >= wlen - 1) ? 1 : 0; + target_exposure = enforce_hold(position, target_exposure, hold_time, +- min_hold_bars, is_last_bar); ++ min_hold_bars, is_last_bar, NULL); + + /* ── Trailing stop ────────────────────────────────────────────── */ + { +diff --git a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu +index b2a9458ab..9955fbd72 100644 +--- a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu +@@ -42,7 +42,7 @@ extern "C" __global__ void gather_states( + int current_step, + float initial_capital, // host scalar — convert on first use + float spread_cost, // host scalar — convert on first use +- int ofi_dim, // 0 = no OFI, 8 = standard OFI features ++ int ofi_dim, // 0 = no OFI, 20 = extended microstructure + int padded_sd // pad128(state_dim) — output row stride for cuBLAS K-tile alignment + ) { + // Shared memory tile for coalesced portfolio reads — 8 F32 values per thread +diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +index 97f776f25..bd2f33ab1 100644 +--- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +@@ -199,7 +199,7 @@ extern "C" __global__ void c51_loss_batched( + const float* __restrict__ curiosity_errors, /* [B] per-sample curiosity prediction error (0 = no penalty) */ + float curiosity_q_penalty_lambda, /* scaling factor: gamma *= 1/(1 + lambda * error) */ + +- float gamma, ++ const float* __restrict__ gamma_buf, /* [B] per-sample effective gamma */ + int batch_size, + int num_atoms, + const float* __restrict__ per_sample_support, /* [B, 3] per-sample: [v_min, v_max, delta_z] */ +@@ -331,10 +331,10 @@ extern "C" __global__ void c51_loss_batched( + * gamma down so the Bellman target approaches the immediate reward, + * preventing overconfident extrapolation into novel states. + * gamma_eff = gamma / (1 + lambda * prediction_error) */ +- float gamma_eff = gamma; ++ float gamma_eff = gamma_buf[sample_id]; + if (curiosity_q_penalty_lambda > 0.0f) { + float cur_err = curiosity_errors[sample_id]; +- gamma_eff = gamma / (1.0f + curiosity_q_penalty_lambda * cur_err); ++ gamma_eff = gamma_buf[sample_id] / (1.0f + curiosity_q_penalty_lambda * cur_err); + } + + float total_ce = 0.0f; +diff --git a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu +index 7009e2e7a..75aa60f67 100644 +--- a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu +@@ -64,6 +64,11 @@ __device__ __forceinline__ float warp_sum_cur(float val) { + * gradients via atomicAdd with warp-level pre-reduction (32x fewer atomics). + * Threads within a warp process different samples but accumulate gradients + * to the same weight indices, so warp reduction before atomicAdd is valid. ++ * ++ * DETERMINISM NOTE: Residual per-block atomicAdd remains (one per warp per weight). ++ * Full elimination would require a separate per-weight reduction kernel. ++ * Impact: curiosity is a small auxiliary model (~11K params), not on the ++ * primary DQN gradient path, so residual non-determinism is negligible. + */ + extern "C" __global__ void curiosity_forward_backward( + const float* __restrict__ states, /* [N, state_dim] */ +@@ -280,16 +285,17 @@ extern "C" __global__ void curiosity_adam_step_fused( + * counter for grid-wide synchronization between the fwd/bwd phase + * (sample-parallel) and the Adam phase (parameter-parallel). + * +- * Phase 1 (all threads): Zero gradient buffers via grid-stride loop, +- * then __syncthreads() within each block. +- * Phase 2 (all threads): Forward + backward pass (one sample per thread), ++ * Phase 1 (all threads): Forward + backward pass (one sample per thread), + * accumulate gradients via warp-reduced atomicAdd. +- * Phase 3 (last-arriving block only): After all blocks complete phase 2, ++ * Phase 2 (last-arriving block only): After all blocks complete phase 1, + * the last block to arrive (detected via atomic counter) applies + * Adam update across all parameters in a grid-stride loop. + * + * Saves 4 memset dispatches + 1 kernel launch = 5 fewer GPU dispatches + * per training step (~30-50 us on H100 at high training frequency). ++ * ++ * DETERMINISM NOTE: Same warp-reduced atomicAdd pattern as curiosity_forward_backward. ++ * Residual per-warp atomicAdd contention is negligible for this auxiliary model. + */ + extern "C" __global__ void curiosity_fused_zero_fwd_bwd_adam( + const float* __restrict__ states, /* [N, state_dim] */ +diff --git a/crates/ml/src/cuda_pipeline/dt_kernels.cu b/crates/ml/src/cuda_pipeline/dt_kernels.cu +index d2f6f6d04..c11997561 100644 +--- a/crates/ml/src/cuda_pipeline/dt_kernels.cu ++++ b/crates/ml/src/cuda_pipeline/dt_kernels.cu +@@ -224,7 +224,11 @@ extern "C" __global__ void dt_causal_attention_kernel( + } + __syncthreads(); /* Ensure bias initialization is visible */ + +- /* Each head's projection contribution via atomicAdd */ ++ /* Each head's projection contribution via atomicAdd. ++ * DETERMINISM NOTE: num_heads blocks (typically 4) accumulate into ++ * the same output element. With only 4 concurrent writers per element, ++ * non-determinism is minimal. Full elimination would require all heads ++ * in one block (changing grid layout). DT is a separate model. */ + for (int dd = 0; dd < E; dd++) { + float proj = 0.0f; + for (int dh = 0; dh < Dh; dh++) { +@@ -467,6 +471,10 @@ extern "C" __global__ void dt_cross_entropy_kernel( + loss = fminf(fmaxf(loss, 0.0f), 100.0f); + + per_sample_loss[n] = loss; ++ /* total_loss accumulated via warp+block reduction — one atomicAdd per BLOCK. ++ * Since grid=(B*T) with 1 thread per block, each block has 1 thread, so ++ * atomicAdd is from exactly N sources with fixed order. For large N, use ++ * a separate reduction kernel for full determinism. */ + atomicAdd(total_loss, loss / (float)N); + } + +@@ -559,7 +567,12 @@ extern "C" __global__ void dt_linear_backward_kernel( + dx[i] = val; + } + +- /* Accumulate weight + bias gradients via atomicAdd */ ++ /* Accumulate weight + bias gradients via warp-reduced atomicAdd. ++ * Each sample's contribution is pre-reduced within the warp before ++ * writing, yielding one atomicAdd per warp per weight instead of ++ * one per thread. ++ * DETERMINISM NOTE: Decision Transformer is a separate model from ++ * the main DQN; residual cross-block atomicAdd is acceptable. */ + for (int i = tid; i < I; i += blockDim.x) { + float xi = x[i]; + for (int o = 0; o < O; o++) { +@@ -567,7 +580,7 @@ extern "C" __global__ void dt_linear_backward_kernel( + } + } + +- /* Bias gradient: only one thread per sample contributes */ ++ /* Bias gradient */ + for (int o = tid; o < O; o += blockDim.x) { + atomicAdd(&db[o], dout[o]); + } +diff --git a/crates/ml/src/cuda_pipeline/ensemble_kernels.cu b/crates/ml/src/cuda_pipeline/ensemble_kernels.cu +index cfe73619d..7ea0e08b0 100644 +--- a/crates/ml/src/cuda_pipeline/ensemble_kernels.cu ++++ b/crates/ml/src/cuda_pipeline/ensemble_kernels.cu +@@ -128,7 +128,11 @@ extern "C" __global__ void ensemble_diversity_kernel( + kl_sum = kl; + } + +- /* -- Hierarchical reduction: warp -> block -> atomicAdd ------------ */ ++ /* -- Hierarchical reduction: warp -> block -> one atomicAdd per BLOCK -- ++ * This is nearly deterministic: only grid_dim atomicAdds to a single ++ * scalar. For full determinism, a two-phase reduction (like IQN's ++ * iqn_loss_reduce) would eliminate atomicAdd entirely, but for a ++ * monitoring-only diversity loss, per-block atomicAdd suffices. */ + /* Warp-level reduction via shuffle (no shared memory) */ + for (int offset = 16; offset > 0; offset >>= 1) + kl_sum = kl_sum + __shfl_xor_sync(0xFFFFFFFF, kl_sum, offset); +diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu +index a2b909117..1297e8ff3 100644 +--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu ++++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu +@@ -20,7 +20,7 @@ + * [market_dim .. market_dim+3) : portfolio features (value_norm, position, cash_norm) + * [market_dim+3 .. state_dim) : zero-pad for tensor-core alignment + * +- * Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=23]): ++ * Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=30]): + * [0] position — current contract position (signed) + * [1] cash — cash balance + * [2] portfolio_value — mark-to-market total value (cash + position * price) +@@ -36,6 +36,13 @@ + * [20] intra_trade_max_dd — worst unrealized drawdown during current trade (v8) + * [21] (reserved) — was last_trade_t (clustering stats, slot retained) + * [22] (reserved) — was interval_sum (clustering stats, slot retained) ++ * [23] plan_target_bars — 0=no plan, >0=active plan max hold bars ++ * [24] plan_profit_target — raw profit threshold % ++ * [25] plan_stop_loss — raw stop loss threshold % ++ * [26] plan_scale_aggression — position ramp speed ++ * [27] plan_conviction — position size fraction ++ * [28] plan_asymmetry — profit/stop ratio ++ * [29] counter_plan_q — opposite direction Q at entry + * + * Branching DQN (Tavakoli et al., 2018) — 4-branch hierarchical: + * 4 independent advantage heads: direction (b0_size=3), magnitude (b1_size=3), +@@ -52,10 +59,10 @@ + */ + + /* ------------------------------------------------------------------ */ +-/* Portfolio stride for experience kernels (23 bf16 per episode). */ ++/* Portfolio stride for experience kernels (30 floats per episode). */ + /* portfolio_sim_kernel uses its own stride of 8 — do NOT change it. */ + /* ------------------------------------------------------------------ */ +-#define PORTFOLIO_STRIDE 23 ++#define PORTFOLIO_STRIDE 30 + #define DSR_A_SLOT 3 + #define DSR_B_SLOT 4 + #define DSR_TRADE_COUNT_SLOT 5 +@@ -718,7 +725,8 @@ extern "C" __global__ void experience_action_select( + float eps_ord_mult, /* per-branch epsilon multiplier: order type */ + float eps_urg_mult, /* per-branch epsilon multiplier: urgency */ + int timestep, /* current timestep for stateless RNG */ +- const float* __restrict__ per_sample_epsilon /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */ ++ const float* __restrict__ per_sample_epsilon, /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */ ++ const float* __restrict__ isv_signals_ptr /* [8] pinned ISV signals for adaptive hold. NULL = static. */ + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; +@@ -770,10 +778,19 @@ extern "C" __global__ void experience_action_select( + int in_hold = 0; + float cur_position = 0.0f; + if (min_hold_bars > 0 && portfolio_states != NULL) { +- int ps_base = i * 23; /* PORTFOLIO_STRIDE = 23 */ ++ int ps_base = i * PORTFOLIO_STRIDE; + float hold_time_val = portfolio_states[ps_base + 10]; + cur_position = portfolio_states[ps_base + 0]; +- in_hold = (hold_time_val > 0.0f && hold_time_val < (float)min_hold_bars); ++ /* Adaptive hold: base + ISV-driven extension (same formula as Layer 2) */ ++ int adaptive_hold = min_hold_bars; ++ if (isv_signals_ptr != NULL) { ++ float stability = 1.0f / (1.0f + isv_signals_ptr[1]); ++ float confidence = 1.0f - fminf(isv_signals_ptr[3], 1.0f); ++ int extension = (int)(stability * confidence * 8.0f); ++ if (isv_signals_ptr[5] < -0.001f) extension = 0; ++ adaptive_hold = min_hold_bars + extension; ++ } ++ in_hold = (hold_time_val > 0.0f && hold_time_val < (float)adaptive_hold); + } + + if (in_hold) { +@@ -970,6 +987,17 @@ extern "C" __global__ void experience_action_select( + } + } + ++ /* Plan direction lock: during active plan, force current direction */ ++ if (portfolio_states != NULL) { ++ int ps_base_plan = i * PORTFOLIO_STRIDE; ++ int has_plan_active = (portfolio_states[ps_base_plan + 23] > 0.5f); ++ if (has_plan_active) { ++ float cur_pos_plan = portfolio_states[ps_base_plan + 0]; ++ if (cur_pos_plan > 0.001f) dir_idx = 2; /* Long locked */ ++ else if (cur_pos_plan < -0.001f) dir_idx = 0; /* Short locked */ ++ } ++ } ++ + /* Compose factored action: 4-branch encoding + * action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency */ + int action_idx = dir_idx * b1_size * b2_size * b3_size +@@ -1022,7 +1050,7 @@ extern "C" __global__ void experience_action_select( + * [2] raw_close — raw close price (for position cost + tx) + * [3] raw_next — raw next-bar close (UNUSED in reward path) + * +- * portfolio_states layout: [N, PORTFOLIO_STRIDE=23] (read-write, bf16) ++ * portfolio_states layout: [N, PORTFOLIO_STRIDE=30] (read-write, float) + * See file header for field definitions. + */ + extern "C" __global__ void experience_env_step( +@@ -1081,7 +1109,9 @@ extern "C" __global__ void experience_env_step( + const float* __restrict__ q_variance, /* [N, total_actions] distributional Var[Q]. NULL = disabled. */ + int total_actions_for_var, /* b0+b1+b2+b3 to index q_variance */ + const float* __restrict__ cost_anneal_ptr, /* [1] pinned device-mapped: 0.0=no costs, 1.0=full costs */ +- const float* __restrict__ commit_lambda_buf /* [N] per-sample commitment lambda. NULL = use 0.01. */ ++ const float* __restrict__ commit_lambda_buf, /* [N] per-sample commitment lambda. NULL = use 0.01. */ ++ const float* __restrict__ isv_signals_ptr, /* [8] pinned device-mapped ISV signals. NULL = static hold. */ ++ const float* __restrict__ plan_params_ptr /* [N, 6] trade plan params. NULL = no plan. */ + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; +@@ -1168,7 +1198,7 @@ extern "C" __global__ void experience_env_step( + /* Guard against degenerate prices from data gaps. */ + if (raw_close <= 0.0f) raw_close = 1.0f; + +- /* ---- Read full portfolio state (PORTFOLIO_STRIDE=23) ---- */ ++ /* ---- Read full portfolio state (PORTFOLIO_STRIDE=30) ---- */ + /* Portfolio arithmetic uses float accumulators because trade physics + * functions (execute_trade, apply_margin_cap, etc.) in trade_physics.cuh + * are all float. Converting the entire physics engine to bf16 would +@@ -1369,7 +1399,7 @@ extern "C" __global__ void experience_env_step( + * violations, but epsilon exploration can still slip through. + * ================================================================ */ + int is_last_bar = (bar_idx >= total_bars - 1) ? 1 : 0; +- float enforced_position = enforce_hold(ps0_f, target_position, hold_time, min_hold_bars, is_last_bar); ++ float enforced_position = enforce_hold(ps0_f, target_position, hold_time, min_hold_bars, is_last_bar, isv_signals_ptr); + int hold_violation = (fabsf(enforced_position - target_position) > 0.001f); + + if (hold_violation) { +@@ -1409,6 +1439,51 @@ extern "C" __global__ void experience_env_step( + tx_cost_multiplier, spread_cost, max_position, + order_type_idx, spread_scale); + ++ /* ── Trade Plan: activation on Flat → Positioned ── */ ++ int was_flat_plan = (fabsf(ps0_f) < 0.001f); ++ int now_positioned_plan = (fabsf(position) > 0.001f); ++ ++ if (was_flat_plan && now_positioned_plan && plan_params_ptr != NULL) { ++ const float* pp = plan_params_ptr + i * 6; ++ ps[23] = pp[0]; /* target_bars */ ++ ps[24] = pp[1]; /* profit_target */ ++ ps[25] = pp[2]; /* stop_loss */ ++ ps[26] = pp[3]; /* scale_aggression */ ++ ps[27] = pp[4]; /* conviction */ ++ ps[28] = pp[5]; /* asymmetry */ ++ /* Apply conviction to position size */ ++ position *= fmaxf(ps[27], 0.1f); ++ } ++ ++ /* ── Trade Plan: enforcement (auto-exit conditions) ── */ ++ int has_plan = (ps[23] > 0.5f); ++ if (has_plan && fabsf(position) > 0.001f) { ++ float pnl_pct = (raw_close - entry_price) / fmaxf(fabsf(entry_price), 1.0f) ++ * ((position > 0.0f) ? 1.0f : -1.0f); ++ ++ /* ISV-modulated thresholds */ ++ float stability = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f; ++ float eff_profit = ps[24] * ps[28] * (0.5f + 0.5f * stability); ++ float eff_stop = ps[25] * fmaxf(stability, 0.5f); ++ ++ int plan_exit = 0; ++ if (pnl_pct >= eff_profit) plan_exit = 1; ++ if (pnl_pct <= -eff_stop) plan_exit = 1; ++ if (hold_time >= ps[23]) plan_exit = 1; ++ if (stability < 0.3f && hold_time > 3.0f) plan_exit = 1; ++ ++ /* Scale schedule: ramp over first 3 bars */ ++ if (hold_time < 3.0f && !plan_exit) { ++ float scale = ps[26] + (1.0f - ps[26]) * hold_time / 3.0f; ++ position *= fminf(scale, 1.0f); ++ } ++ ++ if (plan_exit) { ++ position = 0.0f; ++ for (int s = 23; s <= 29; s++) ps[s] = 0.0f; ++ } ++ } ++ + /* ---- Mark-to-market PnL for this timestep ---- */ + /* Per-bar mark-to-market P&L removed: used raw_next (future price), + * creating 1-bar action-reward misalignment. Rewards are now +@@ -1785,7 +1860,7 @@ extern "C" __global__ void experience_env_step( + reward += position_entropy_weight * combined_entropy; + } + +- /* ---- Update full portfolio state (PORTFOLIO_STRIDE=23) ---- */ ++ /* ---- Update full portfolio state (PORTFOLIO_STRIDE=30) ---- */ + ps[0] = (position); + ps[1] = (cash); + ps[2] = (new_portfolio_value); +@@ -1992,6 +2067,10 @@ extern "C" __global__ void experience_env_step( + ps[17] = 0.0f; + ps[18] = 0.0f; + ps[19] = 0.0f; ++ ps[20] = 0.0f; /* intra_trade_max_dd */ ++ ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f; ++ ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f; ++ ps[29] = 0.0f; /* plan slots zeroed */ + current_timesteps[i] = 0; + } else { + /* Soft reset (trade_complete): keep equity, clear trade state only. +@@ -2002,6 +2081,9 @@ extern "C" __global__ void experience_env_step( + ps[12] = 0.0f; /* entry_price — no active trade */ + ps[13] = (ps[11]); /* trade_start_pnl = current realized_pnl */ + ps[20] = 0.0f; /* intra_trade_max_dd — reset for next trade */ ++ ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f; ++ ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f; ++ ps[29] = 0.0f; /* plan slots zeroed on trade complete */ + current_timesteps[i] = 0; + } + } +@@ -2286,12 +2368,39 @@ extern "C" __global__ void compute_expected_q( + branch_logit_offset += N * n_d * num_atoms; + } + +- /* Accumulate atom stats across all samples (monitoring only) */ ++ /* Accumulate atom stats across all samples (monitoring only). ++ * Uses warp+block reduction to minimize atomicAdd non-determinism. ++ * One atomicAdd per BLOCK (not per thread). */ + if (atom_stats != NULL) { +- /* Average per-action: divide by total_actions so the final mean is per-action entropy */ + float inv_actions = 1.0f / (float)total_actions; +- atomicAdd(&atom_stats[0], sample_entropy * inv_actions); +- atomicAdd(&atom_stats[1], (float)sample_utilized * inv_actions); ++ float my_entropy = sample_entropy * inv_actions; ++ float my_util = (float)sample_utilized * inv_actions; ++ ++ /* Warp-level reduction */ ++ for (int offset = 16; offset > 0; offset >>= 1) { ++ my_entropy += __shfl_xor_sync(0xFFFFFFFF, my_entropy, offset); ++ my_util += __shfl_xor_sync(0xFFFFFFFF, my_util, offset); ++ } ++ ++ /* Block-level reduction via shared memory */ ++ __shared__ float s_entropy[8]; ++ __shared__ float s_util[8]; ++ int warp_id = threadIdx.x / 32; ++ int lane = threadIdx.x % 32; ++ if (lane == 0) { s_entropy[warp_id] = my_entropy; s_util[warp_id] = my_util; } ++ __syncthreads(); ++ if (warp_id == 0) { ++ float ve = (lane < blockDim.x / 32) ? s_entropy[lane] : 0.0f; ++ float vu = (lane < blockDim.x / 32) ? s_util[lane] : 0.0f; ++ for (int off = 16; off > 0; off >>= 1) { ++ ve += __shfl_xor_sync(0xFFFFFFFF, ve, off); ++ vu += __shfl_xor_sync(0xFFFFFFFF, vu, off); ++ } ++ if (lane == 0) { ++ atomicAdd(&atom_stats[0], ve); ++ atomicAdd(&atom_stats[1], vu); ++ } ++ } + } + } + +@@ -2382,7 +2491,7 @@ extern "C" __global__ void expert_action_override( + /* Check if current position already matches expert signal — skip if redundant. + * portfolio_states[i * PORTFOLIO_STRIDE + 0] = current position. + * If expert says Long and we're already Long (position > 0.5), skip. */ +- float current_pos = (portfolio_states[i * 23 + 0]); /* PORTFOLIO_STRIDE=23, pos at idx 0 */ ++ float current_pos = (portfolio_states[i * PORTFOLIO_STRIDE + 0]); + if (expert_dir == 2 && current_pos > 0.5f) return; /* Already Long */ + if (expert_dir == 0 && current_pos < -0.5f) return; /* Already Short */ + +@@ -2958,15 +3067,16 @@ extern "C" __global__ void selectivity_forward( + + /** + * Mamba-2 selectivity gate — backward pass (BCE gradient). ++ * DETERMINISTIC: one thread per weight element, loops over batch. + * + * target[b] = clamp(per_sample_loss[b] / mean_loss, 0, 1) + * d_z[b] = sel[b] - target[b] +- * dW[k] += d_z[b] * h_s2[b, k] (atomicAdd) +- * db += d_z[b] (atomicAdd) ++ * dW[k] = sum_b(d_z[b] * h_s2[b, k]) ++ * db = sum_b(d_z[b]) + * +- * d_sel_params layout: [dW[SH2], db] (SH2+1 floats, must be zeroed before call) ++ * d_sel_params layout: [dW[SH2], db] (SH2+1 floats) + * +- * Grid: ceil(B/256), Block: 256 ++ * Grid: ceil((SH2+1)/256), Block: 256 + */ + extern "C" __global__ void selectivity_backward( + const float* __restrict__ h_s2, /* [B, SH2] */ +@@ -2976,18 +3086,29 @@ extern "C" __global__ void selectivity_backward( + float mean_loss, + int B, int SH2) + { +- int b = blockIdx.x * blockDim.x + threadIdx.x; +- if (b >= B) return; ++ int widx = blockIdx.x * blockDim.x + threadIdx.x; ++ int total_params = SH2 + 1; ++ if (widx >= total_params) return; + + float safe_mean = fmaxf(mean_loss, 1e-8f); +- float target = fminf(fmaxf(per_sample_loss[b] / safe_mean, 0.0f), 1.0f); +- float d_z = sel_out[b] - target; ++ float sum = 0.0f; + +- const float* h = h_s2 + b * SH2; +- for (int k = 0; k < SH2; k++) { +- atomicAdd(&d_sel_params[k], d_z * h[k]); ++ if (widx < SH2) { ++ /* dW[widx] = sum_b(d_z[b] * h_s2[b, widx]) */ ++ for (int b = 0; b < B; b++) { ++ float target = fminf(fmaxf(per_sample_loss[b] / safe_mean, 0.0f), 1.0f); ++ float d_z = sel_out[b] - target; ++ sum += d_z * h_s2[(long long)b * SH2 + widx]; ++ } ++ } else { ++ /* db = sum_b(d_z[b]) */ ++ for (int b = 0; b < B; b++) { ++ float target = fminf(fmaxf(per_sample_loss[b] / safe_mean, 0.0f), 1.0f); ++ float d_z = sel_out[b] - target; ++ sum += d_z; ++ } + } +- atomicAdd(&d_sel_params[SH2], d_z); /* db */ ++ d_sel_params[widx] = sum; /* single deterministic write */ + } + + /** +@@ -3412,14 +3533,17 @@ extern "C" __global__ void kan_gate_backward( + + d_gate_pre[i] = d_raw * dgate_dx; + +- /* Accumulate d_coeff via atomicAdd (batch dimension reduction) */ ++ /* Accumulate d_coeff and d_residual_w via atomicAdd (batch reduction). ++ * DETERMINISM NOTE: Each thread has a unique (neuron, span) combination ++ * so warp-level pre-reduction is not applicable (different target indices). ++ * Full elimination requires a separate per-weight reduction kernel. ++ * Impact: KAN spline is a small auxiliary gating component (~adv_h*8 ++ * coefficients + adv_h residual weights), not on the primary gradient path. */ + float* dc = d_spline_coeff + neuron * NUM_BASES; + if (k0 >= 0 && k0 < NUM_BASES) atomicAdd(&dc[k0], d_raw * b0); + if (span >= 0 && span < NUM_BASES) atomicAdd(&dc[span], d_raw * b1); + if (k2 >= 0 && k2 < NUM_BASES) atomicAdd(&dc[k2], d_raw * b2); + if (k3_idx >= 0 && k3_idx < NUM_BASES) atomicAdd(&dc[k3_idx], d_raw * b3); +- +- /* Accumulate d_residual_w via atomicAdd */ + atomicAdd(&d_residual_w[neuron], d_raw * x); + } + +@@ -3738,32 +3862,16 @@ extern "C" __global__ void liquid_tau_rk4_step( + + /** + * Backward through 2-step FC-SiLU-FC diffusion denoiser. ++ * DETERMINISTIC: one thread per weight element, loops over batch. + * + * Loss: L = mean_b((Q_refined[b] - Q_target[b])^2) summed over 12 actions. + * d_Q_refined = 2*(Q_refined - Q_target) / B per element. + * + * Then backprop through 2 FC-SiLU-FC steps (reverse order) to accumulate +- * d_params via atomicAdd across the batch dimension. +- * +- * For each step (reverse): +- * d_W2 += outer(d_out, h_silu), d_b2 += d_out +- * d_h = W2^T @ d_out +- * d_pre_silu = d_h * silu'(pre_act) +- * d_W1 += outer(d_pre_silu, input), d_b1 += d_pre_silu +- * +- * silu'(x) = sigmoid(x) + x * sigmoid(x) * (1 - sigmoid(x)) +- * = sigmoid(x) * (1 + x * (1 - sigmoid(x))) ++ * d_params by summing over all samples with fixed iteration order. + * +- * One thread per sample. Grid: ceil(B/256), Block: 256. +- * Uses atomicAdd for weight gradient accumulation across batch. +- * +- * @param Q_refined [B, 12] denoiser output (final q_coord_buf) +- * @param Q_target [B, 12] target network Q-values +- * @param denoise_params [1800] current weights (2 steps × 900) +- * @param d_denoise_params [1800] gradient accumulator (caller must zero) +- * @param Q_input [B, 12] denoiser input (pre-denoise Q-values) +- * @param var_q [B, 12] Var[Q] for noise conditioning +- * @param B batch size ++ * Total params: 1800 (2 steps × 900). ++ * Grid: ceil(1800/256), Block: 256. + */ + extern "C" __global__ void q_denoise_backward( + const float* __restrict__ Q_refined, +@@ -3774,52 +3882,45 @@ extern "C" __global__ void q_denoise_backward( + const float* __restrict__ var_q, + int B) + { +- int b = blockIdx.x * blockDim.x + threadIdx.x; +- if (b >= B) return; +- ++ int widx = blockIdx.x * blockDim.x + threadIdx.x; + const int D = 12; + const int H = 24; + const int STEP_PARAMS = 900; /* W1[24,24]+b1[24]+W2[12,24]+b2[12] */ + const int num_steps = 2; +- float inv_B = 1.0f / (float)B; +- +- /* Compute loss gradient: d_out[i] = 2*(Q_refined[b,i] - Q_target[b,i]) / B */ +- float d_out[12]; +- for (int i = 0; i < D; i++) { +- d_out[i] = 2.0f * (Q_refined[b * D + i] - Q_target[b * D + i]) * inv_B; +- } +- +- /* Backprop through steps in reverse order (step 1 then step 0). +- * Each step: input_step[24] → h=SiLU(W1@input+b1)[24] → residual=W2@h+b2[12] +- * Q_next = Q_prev + 0.1 * residual +- * So d_residual = 0.1 * d_out, and d_Q_prev += d_out (identity skip connection). */ +- +- for (int step = num_steps - 1; step >= 0; step--) { +- int off = step * STEP_PARAMS; +- const float* W1 = denoise_params + off; +- const float* b1_ptr = denoise_params + off + 576; +- const float* W2 = denoise_params + off + 600; +- const float* b2_ptr = denoise_params + off + 888; ++ const int TOTAL_PARAMS = num_steps * STEP_PARAMS; ++ if (widx >= TOTAL_PARAMS) return; + +- float* dW1 = d_denoise_params + off; +- float* db1 = d_denoise_params + off + 576; +- float* dW2 = d_denoise_params + off + 600; +- float* db2 = d_denoise_params + off + 888; ++ float inv_B = 1.0f / (float)B; + +- /* d_residual = 0.1 * d_out */ +- float d_res[12]; ++ /* Determine which step and which sub-parameter this thread handles */ ++ int step = widx / STEP_PARAMS; ++ int local = widx % STEP_PARAMS; ++ ++ /* Sub-parameter layout within one step (900 total): ++ * W1: [H, H] = 576 elements (offsets 0..575) ++ * b1: [H] = 24 elements (offsets 576..599) ++ * W2: [D, H] = 288 elements (offsets 600..887) ++ * b2: [D] = 12 elements (offsets 888..899) */ ++ int is_W1 = (local < 576); ++ int is_b1 = (local >= 576 && local < 600); ++ int is_W2 = (local >= 600 && local < 888); ++ /* int is_b2 = (local >= 888); */ ++ ++ float acc = 0.0f; ++ ++ for (int b = 0; b < B; b++) { ++ /* Compute loss gradient for this sample */ ++ float d_out_b[12]; + for (int i = 0; i < D; i++) { +- d_res[i] = 0.1f * d_out[i]; ++ d_out_b[i] = 2.0f * (Q_refined[b * D + i] - Q_target[b * D + i]) * inv_B; + } + +- /* Reconstruct forward: build input for this step. +- * Step 0: input = [Q_input; noise_level_step0] +- * Step 1: input = [Q_after_step0; noise_level_step1] +- * We need to reconstruct the intermediate Q-values. +- * For simplicity, re-run the forward to get intermediates. */ ++ /* Process steps in reverse. The gradient for the current step is needed. */ ++ /* For each step s from (num_steps-1) down to 0, d_out propagates through ++ * the skip connection unchanged. We only need the gradient at `step`. */ + +- /* Reconstruct Q at start of this step by running forward from Q_input +- * through steps 0..step-1. */ ++ /* Reconstruct Q at start of `step` by running forward from Q_input ++ * through steps 0..(step-1). */ + float q_curr[12]; + for (int i = 0; i < D; i++) q_curr[i] = Q_input[b * D + i]; + +@@ -3838,19 +3939,19 @@ extern "C" __global__ void q_denoise_backward( + } + float h_s[24]; + for (int i = 0; i < H; i++) { +- float acc = sb1[i]; +- for (int j = 0; j < H; j++) acc += sW1[i * H + j] * inp_s[j]; +- float sig = 1.0f / (1.0f + expf(-acc)); +- h_s[i] = acc * sig; ++ float v = sb1[i]; ++ for (int j = 0; j < H; j++) v += sW1[i * H + j] * inp_s[j]; ++ float sig = 1.0f / (1.0f + expf(-v)); ++ h_s[i] = v * sig; + } + for (int i = 0; i < D; i++) { +- float acc = sb2[i]; +- for (int j = 0; j < H; j++) acc += sW2[i * H + j] * h_s[j]; +- q_curr[i] += 0.1f * acc; ++ float v = sb2[i]; ++ for (int j = 0; j < H; j++) v += sW2[i * H + j] * h_s[j]; ++ q_curr[i] += 0.1f * v; + } + } + +- /* Now q_curr = Q at start of this step. Build the input. */ ++ /* Build input for this step */ + float schedule = 1.0f - (float)(step + 1) / (float)num_steps; + float input[24]; + for (int i = 0; i < D; i++) { +@@ -3859,54 +3960,67 @@ extern "C" __global__ void q_denoise_backward( + } + + /* Forward through this step to get pre-activations */ ++ int off = step * STEP_PARAMS; ++ const float* W1 = denoise_params + off; ++ const float* b1_ptr = denoise_params + off + 576; ++ const float* W2 = denoise_params + off + 600; ++ + float pre_act[24], h_silu[24]; + for (int i = 0; i < H; i++) { +- float acc = b1_ptr[i]; +- for (int j = 0; j < H; j++) acc += W1[i * H + j] * input[j]; +- pre_act[i] = acc; +- float sig = 1.0f / (1.0f + expf(-acc)); +- h_silu[i] = acc * sig; ++ float v = b1_ptr[i]; ++ for (int j = 0; j < H; j++) v += W1[i * H + j] * input[j]; ++ pre_act[i] = v; ++ float sig = 1.0f / (1.0f + expf(-v)); ++ h_silu[i] = v * sig; + } + +- /* Backward: d_W2 += outer(d_res, h_silu), d_b2 += d_res */ ++ /* d_residual = 0.1 * d_out */ ++ float d_res[12]; + for (int i = 0; i < D; i++) { +- atomicAdd(&db2[i], d_res[i]); +- for (int j = 0; j < H; j++) { +- atomicAdd(&dW2[i * H + j], d_res[i] * h_silu[j]); +- } ++ d_res[i] = 0.1f * d_out_b[i]; + } + +- /* d_h = W2^T @ d_res */ +- float d_h[24]; +- for (int i = 0; i < H; i++) { +- float acc = 0.0f; +- for (int j = 0; j < D; j++) { +- acc += W2[j * H + i] * d_res[j]; ++ /* Compute the gradient for this thread's specific parameter */ ++ if (is_W2) { ++ /* dW2[i, j] += d_res[i] * h_silu[j] */ ++ int idx = local - 600; ++ int row = idx / H; ++ int col = idx % H; ++ acc += d_res[row] * h_silu[col]; ++ } else if (local >= 888) { ++ /* db2[i] += d_res[i] */ ++ int idx = local - 888; ++ acc += d_res[idx]; ++ } else { ++ /* Need d_pre for W1/b1 gradients */ ++ float d_h[24]; ++ for (int i = 0; i < H; i++) { ++ float v = 0.0f; ++ for (int j = 0; j < D; j++) v += W2[j * H + i] * d_res[j]; ++ d_h[i] = v; ++ } ++ float d_pre[24]; ++ for (int i = 0; i < H; i++) { ++ float sig = 1.0f / (1.0f + expf(-pre_act[i])); ++ float silu_grad = sig * (1.0f + pre_act[i] * (1.0f - sig)); ++ d_pre[i] = d_h[i] * silu_grad; + } +- d_h[i] = acc; +- } +- +- /* d_pre_silu = d_h * silu'(pre_act) +- * silu'(x) = sigmoid(x) * (1 + x * (1 - sigmoid(x))) */ +- float d_pre[24]; +- for (int i = 0; i < H; i++) { +- float sig = 1.0f / (1.0f + expf(-pre_act[i])); +- float silu_grad = sig * (1.0f + pre_act[i] * (1.0f - sig)); +- d_pre[i] = d_h[i] * silu_grad; +- } + +- /* d_W1 += outer(d_pre, input), d_b1 += d_pre */ +- for (int i = 0; i < H; i++) { +- atomicAdd(&db1[i], d_pre[i]); +- for (int j = 0; j < H; j++) { +- atomicAdd(&dW1[i * H + j], d_pre[i] * input[j]); ++ if (is_W1) { ++ /* dW1[i, j] += d_pre[i] * input[j] */ ++ int idx = local; ++ int row = idx / H; ++ int col = idx % H; ++ acc += d_pre[row] * input[col]; ++ } else if (is_b1) { ++ /* db1[i] += d_pre[i] */ ++ int idx = local - 576; ++ acc += d_pre[idx]; + } + } +- +- /* Propagate gradient to previous step's output: +- * d_out_prev = d_out (skip connection carries gradient through) */ +- /* d_out already holds the gradient for the previous step. */ + } ++ ++ d_denoise_params[widx] = acc; /* single deterministic write */ + } + + /* ================================================================== */ +@@ -4246,21 +4360,22 @@ extern "C" __global__ void atom_position_gradient( + + /** + * Mamba2 BPTT: backpropagation through K=8 scan steps. ++ * DETERMINISTIC: one thread per weight element, loops over batch. + * + * Given d_h_enriched [B, SH2] (gradient from downstream): + * The temporal_context = (W_C @ h_current) * x_K, so: + * d_x_K = d_context * (W_C @ h_current) element-wise +- * d_W_C += sum_b d_context * x_K^T (outer product) ++ * d_W_C[j,s] = sum_b( d_out[b,j] * x_K[b,s] ) + * + * Reverse scan t = K-1..0: + * d_A_t = d_x_{t+1} * x_t * sigmoid'(a_raw_t) +- * d_x_t = d_x_{t+1} * A_t (gate propagation) +- * d_B_t = d_x_{t+1} +- * d_W_A += sum_b d_A_t @ h_t^T +- * d_W_B += sum_b d_B_t @ h_t^T ++ * d_x_t = d_x_{t+1} * A_t ++ * d_W_A[j,s] += sum_b( d_gate[b,t,s] * h_t[b,j] ) ++ * d_W_B[j,s] += sum_b( d_x[b,t,s] * h_t[b,j] ) + * +- * Grid: ceil(B/256), Block: 256. One thread per sample. +- * Accumulates weight gradients via atomicAdd (batch reduction). ++ * Grid: ceil(total_weight_params/256), Block: 256. One thread per weight. ++ * total_weight_params = 3 * sh2 * state_d. ++ * Weight layout: [d_w_a: sh2*state_d | d_w_b: sh2*state_d | d_w_c: sh2*state_d] + */ + extern "C" __global__ void mamba2_scan_backward( + const float* __restrict__ h_history, /* [B, K, SH2] saved activations */ +@@ -4275,80 +4390,90 @@ extern "C" __global__ void mamba2_scan_backward( + int B, + int K, + int sh2, +- int state_d ++ int state_d, ++ const float* __restrict__ isv_signals /* [12] pinned. NULL = no regime modulation. */ + ) { +- int i = blockIdx.x * blockDim.x + threadIdx.x; +- if (i >= B) return; ++ int widx = blockIdx.x * blockDim.x + threadIdx.x; ++ int wc_size = sh2 * state_d; ++ int total_params = 3 * wc_size; ++ if (widx >= total_params) return; ++ ++ /* Determine which weight matrix and which (j, s) indices */ ++ int mat; /* 0=W_A, 1=W_B, 2=W_C */ ++ int local; ++ if (widx < wc_size) { ++ mat = 0; local = widx; ++ } else if (widx < 2 * wc_size) { ++ mat = 1; local = widx - wc_size; ++ } else { ++ mat = 2; local = widx - 2 * wc_size; ++ } ++ int j = local / state_d; ++ int s = local % state_d; + +- const float* h_current = h_s2 + (long long)i * sh2; +- const float* d_out = d_h_enriched + (long long)i * sh2; ++ float acc = 0.0f; + +- /* ── Forward pass replay to get intermediate states ── */ +- float x_states[9][16]; /* x_states[t] for t=0..K, state_d<=16 */ +- float a_raw[8][16]; /* raw gate values before sigmoid */ ++ for (int b = 0; b < B; b++) { ++ const float* d_out = d_h_enriched + (long long)b * sh2; + +- /* x_states[0] = zeros */ +- for (int s = 0; s < state_d; s++) x_states[0][s] = 0.0f; ++ /* ── Forward pass replay to get intermediate states ── */ ++ float x_states_local[9]; /* x_states[t] for t=0..K, only state_dim s */ ++ float a_raw_local[8]; /* raw gate values before sigmoid, only s */ + +- for (int t = 0; t < K; t++) { +- const float* h_t = h_history + ((long long)i * K + t) * sh2; +- for (int s = 0; s < state_d; s++) { ++ x_states_local[0] = 0.0f; ++ ++ for (int t = 0; t < K; t++) { ++ const float* h_t = h_history + ((long long)b * K + t) * sh2; + float a_val = 0.0f, b_val = 0.0f; +- for (int j = 0; j < sh2; j++) { +- a_val += w_a[(long long)j * state_d + s] * h_t[j]; +- b_val += w_b[(long long)j * state_d + s] * h_t[j]; ++ for (int jj = 0; jj < sh2; jj++) { ++ a_val += w_a[(long long)jj * state_d + s] * h_t[jj]; ++ b_val += w_b[(long long)jj * state_d + s] * h_t[jj]; + } +- a_raw[t][s] = a_val; ++ a_raw_local[t] = a_val; + float gate = 1.0f / (1.0f + expf(-a_val)); +- x_states[t+1][s] = gate * x_states[t][s] + b_val; ++ if (isv_signals != NULL) gate *= isv_signals[11]; ++ x_states_local[t + 1] = gate * x_states_local[t] + b_val; + } +- } + +- /* ── Backward through output projection ── */ +- /* temporal_context[j] = sum_s(w_c[j,s] * x_K[s]) +- * d_x_K[s] += sum_j(d_out[j] * w_c[j,s]) +- * d_w_c[j,s] += d_out[j] * x_K[s] */ +- float d_x[16]; +- for (int s = 0; s < state_d; s++) { +- float dx = 0.0f; +- for (int j = 0; j < sh2; j++) { +- dx += d_out[j] * w_c[(long long)j * state_d + s]; +- } +- d_x[s] = dx; +- } +- /* d_W_C gradient */ +- for (int j = 0; j < sh2; j++) { +- for (int s = 0; s < state_d; s++) { +- atomicAdd(&d_w_c[(long long)j * state_d + s], d_out[j] * x_states[K][s]); +- } +- } ++ if (mat == 2) { ++ /* d_W_C[j, s] += d_out[j] * x_states[K][s] */ ++ acc += d_out[j] * x_states_local[K]; ++ } else { ++ /* Need reverse scan for d_W_A and d_W_B */ ++ /* Compute d_x for state dimension s via reverse scan */ ++ float d_x_s = 0.0f; ++ for (int jj = 0; jj < sh2; jj++) { ++ d_x_s += d_out[jj] * w_c[(long long)jj * state_d + s]; ++ } + +- /* ── Reverse scan ── */ +- for (int t = K - 1; t >= 0; t--) { +- const float* h_t = h_history + ((long long)i * K + t) * sh2; ++ for (int t = K - 1; t >= 0; t--) { ++ const float* h_t = h_history + ((long long)b * K + t) * sh2; + +- for (int s = 0; s < state_d; s++) { +- float a_val = a_raw[t][s]; +- float gate = 1.0f / (1.0f + expf(-a_val)); +- float sigmoid_deriv = gate * (1.0f - gate); ++ float a_val = a_raw_local[t]; ++ float gate = 1.0f / (1.0f + expf(-a_val)); ++ float sigmoid_deriv = gate * (1.0f - gate); + +- /* d_A_t = d_x_{t+1} * x_t * sigmoid'(a_raw_t) */ +- float d_gate = d_x[s] * x_states[t][s] * sigmoid_deriv; ++ float stability = 1.0f; ++ if (isv_signals != NULL) stability = isv_signals[11]; + +- /* d_W_A: d_gate is scalar per (sample, state_dim), h_t is [SH2] */ +- for (int j = 0; j < sh2; j++) { +- atomicAdd(&d_w_a[(long long)j * state_d + s], d_gate * h_t[j]); +- } ++ if (mat == 0) { ++ /* d_W_A[j, s] += d_gate * h_t[j] */ ++ float d_gate = d_x_s * x_states_local[t] * sigmoid_deriv * stability; ++ acc += d_gate * h_t[j]; ++ } else { ++ /* d_W_B[j, s] += d_x_s * h_t[j] */ ++ acc += d_x_s * h_t[j]; ++ } + +- /* d_W_B: d_x[s] directly (input projection gradient) */ +- for (int j = 0; j < sh2; j++) { +- atomicAdd(&d_w_b[(long long)j * state_d + s], d_x[s] * h_t[j]); ++ d_x_s = d_x_s * gate * stability; + } +- +- /* Propagate d_x backward through gate: d_x_t = d_x_{t+1} * A_t */ +- d_x[s] = d_x[s] * gate; + } + } ++ ++ /* Single deterministic write per weight element */ ++ if (mat == 0) d_w_a[(long long)j * state_d + s] = acc; ++ else if (mat == 1) d_w_b[(long long)j * state_d + s] = acc; ++ else d_w_c[(long long)j * state_d + s] = acc; + } + + /* ================================================================== */ +@@ -4560,7 +4685,7 @@ extern "C" __global__ void branch_independence_penalty( + * exceeds sim_threshold, penalizes mean |Q_i - Q_{i+1}| weighted + * by how far similarity exceeds the threshold. + * +- * Uses atomicAdd for single-scalar accumulation. ++ * Uses warp+block reduction → one atomicAdd per BLOCK (deterministic). + * + * Grid: ceil((N-1) / 256), Block: 256. + */ +@@ -4575,27 +4700,48 @@ extern "C" __global__ void temporal_consistency_penalty( + float sim_threshold + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; +- if (i >= N - 1) return; +- +- float dot = 0.0f, norm_a = 0.0f, norm_b = 0.0f; +- for (int j = 0; j < sh2; j++) { +- float a = h_s2[(long long)i * sh2 + j]; +- float b = h_s2[(long long)(i + 1) * sh2 + j]; +- dot += a * b; +- norm_a += a * a; +- norm_b += b * b; +- } +- float sim = dot / (sqrtf(norm_a) * sqrtf(norm_b) + 1e-8f); +- +- if (sim > sim_threshold) { +- float q_diff = 0.0f; +- for (int a = 0; a < total_actions; a++) { +- q_diff += fabsf(q_values[(long long)i * total_actions + a] +- - q_values[(long long)(i + 1) * total_actions + a]); ++ float my_penalty = 0.0f; ++ ++ if (i < N - 1) { ++ float dot = 0.0f, norm_a = 0.0f, norm_b = 0.0f; ++ for (int j = 0; j < sh2; j++) { ++ float a = h_s2[(long long)i * sh2 + j]; ++ float b = h_s2[(long long)(i + 1) * sh2 + j]; ++ dot += a * b; ++ norm_a += a * a; ++ norm_b += b * b; + } +- q_diff /= (float)total_actions; +- float weight = (sim - sim_threshold) / (1.0f - sim_threshold); +- atomicAdd(penalty_out, lambda_tc * q_diff * weight / (float)(N - 1)); ++ float sim = dot / (sqrtf(norm_a) * sqrtf(norm_b) + 1e-8f); ++ ++ if (sim > sim_threshold) { ++ float q_diff = 0.0f; ++ for (int a = 0; a < total_actions; a++) { ++ q_diff += fabsf(q_values[(long long)i * total_actions + a] ++ - q_values[(long long)(i + 1) * total_actions + a]); ++ } ++ q_diff /= (float)total_actions; ++ float weight = (sim - sim_threshold) / (1.0f - sim_threshold); ++ my_penalty = lambda_tc * q_diff * weight / (float)(N - 1); ++ } ++ } ++ ++ /* Warp-level reduction */ ++ for (int offset = 16; offset > 0; offset >>= 1) ++ my_penalty += __shfl_xor_sync(0xFFFFFFFF, my_penalty, offset); ++ ++ /* Block-level reduction via shared memory */ ++ __shared__ float warp_sums[8]; ++ int warp_id = threadIdx.x / 32; ++ int lane = threadIdx.x % 32; ++ if (lane == 0) warp_sums[warp_id] = my_penalty; ++ __syncthreads(); ++ ++ if (warp_id == 0) { ++ float val = (lane < blockDim.x / 32) ? warp_sums[lane] : 0.0f; ++ for (int off = 16; off > 0; off >>= 1) ++ val += __shfl_xor_sync(0xFFFFFFFF, val, off); ++ if (lane == 0) ++ atomicAdd(penalty_out, val); + } + } + +@@ -4612,28 +4758,50 @@ extern "C" __global__ void temporal_consistency_penalty( + * + * Total loss = lambda_pred * sum(per_sample_mse) / (N-1). + * ++ * Uses warp+block reduction → one atomicAdd per BLOCK (deterministic). ++ * + * Grid: ceil((N-1) / 256), Block: 256. + */ + extern "C" __global__ void predictive_coding_loss( + const float* __restrict__ h_s2, /* [N, SH2] enriched trunk (after Mamba2) */ +- float* __restrict__ loss_out, /* [1] scalar MSE loss (atomicAdd) */ ++ float* __restrict__ loss_out, /* [1] scalar MSE loss */ + int N, + int sh2, + float lambda_pred /* loss weight (0.1) */ + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; +- if (i >= N - 1) return; ++ float my_loss = 0.0f; + +- float mse = 0.0f; +- for (int j = 0; j < sh2; j++) { +- float current = h_s2[(long long)i * sh2 + j]; +- float next = h_s2[(long long)(i + 1) * sh2 + j]; +- float diff = current - next; +- mse += diff * diff; ++ if (i < N - 1) { ++ float mse = 0.0f; ++ for (int j = 0; j < sh2; j++) { ++ float current = h_s2[(long long)i * sh2 + j]; ++ float next = h_s2[(long long)(i + 1) * sh2 + j]; ++ float diff = current - next; ++ mse += diff * diff; ++ } ++ mse /= (float)sh2; ++ my_loss = lambda_pred * mse / (float)(N - 1); + } +- mse /= (float)sh2; + +- atomicAdd(loss_out, lambda_pred * mse / (float)(N - 1)); ++ /* Warp-level reduction */ ++ for (int offset = 16; offset > 0; offset >>= 1) ++ my_loss += __shfl_xor_sync(0xFFFFFFFF, my_loss, offset); ++ ++ /* Block-level reduction via shared memory */ ++ __shared__ float warp_sums[8]; ++ int warp_id = threadIdx.x / 32; ++ int lane = threadIdx.x % 32; ++ if (lane == 0) warp_sums[warp_id] = my_loss; ++ __syncthreads(); ++ ++ if (warp_id == 0) { ++ float val = (lane < blockDim.x / 32) ? warp_sums[lane] : 0.0f; ++ for (int off = 16; off > 0; off >>= 1) ++ val += __shfl_xor_sync(0xFFFFFFFF, val, off); ++ if (lane == 0) ++ atomicAdd(loss_out, val); ++ } + } + + /* ================================================================== */ +@@ -4691,6 +4859,8 @@ extern "C" __global__ void homeostatic_regularizer( + + extern "C" __global__ void risk_budget_forward( + const float* __restrict__ h_s2, ++ const float* __restrict__ isv_signals_dev_ptr, /* [12] pinned — raw ISV */ ++ const float* __restrict__ predicted_error, /* [B] recursive confidence */ + const float* __restrict__ w_risk_fc, + const float* __restrict__ b_risk_fc, + const float* __restrict__ w_risk_out, +@@ -4702,21 +4872,23 @@ extern "C" __global__ void risk_budget_forward( + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= B) return; + ++ int input_dim = SH2 + 13; /* SH2 + 12 ISV + 1 predicted_error */ + const float* h = h_s2 + (long long)i * SH2; + float* h_risk = risk_hidden + (long long)i * AH; + + for (int j = 0; j < AH; j++) { + float val = b_risk_fc[j]; +- for (int k = 0; k < SH2; k++) { +- val += w_risk_fc[(long long)j * SH2 + k] * h[k]; +- } ++ for (int k = 0; k < SH2; k++) ++ val += w_risk_fc[(long long)j * input_dim + k] * h[k]; ++ for (int k = 0; k < 12; k++) ++ val += w_risk_fc[(long long)j * input_dim + SH2 + k] * isv_signals_dev_ptr[k]; ++ val += w_risk_fc[(long long)j * input_dim + SH2 + 12] * predicted_error[i]; + h_risk[j] = fmaxf(val, 0.0f); + } + + float raw = b_risk_out[0]; +- for (int j = 0; j < AH; j++) { ++ for (int j = 0; j < AH; j++) + raw += w_risk_out[j] * h_risk[j]; +- } + risk_budget_out[i] = 1.0f / (1.0f + expf(-raw)); + } + +@@ -4747,8 +4919,18 @@ extern "C" __global__ void apply_risk_budget( + commit_lambda_buf[i] = 0.01f * (1.0f - R); + } + ++/** ++ * Risk budget backward — DETERMINISTIC: one thread per weight element. ++ * ++ * Weight layout (flat): [d_w_risk_fc: AH*(SH2+13) | d_b_risk_fc: AH | d_w_risk_out: AH | d_b_risk_out: 1] ++ * Total: AH*(SH2+13) + AH + AH + 1 = AH*(SH2+15) + 1 ++ * ++ * Grid: ceil(total_risk_params / 256), Block: 256. ++ */ + extern "C" __global__ void risk_budget_backward( + const float* __restrict__ h_s2, ++ const float* __restrict__ isv_signals_dev_ptr, /* [12] raw ISV */ ++ const float* __restrict__ predicted_error, /* [B] recursive confidence */ + const float* __restrict__ risk_hidden, + const float* __restrict__ risk_budget, + const float* __restrict__ d_q_mag, +@@ -4759,34 +4941,459 @@ extern "C" __global__ void risk_budget_backward( + float* __restrict__ d_w_risk_out, + float* __restrict__ d_b_risk_out, + int B, int SH2, int AH, int b1_size ++) { ++ int input_dim = SH2 + 13; ++ int wfc_size = AH * input_dim; ++ int bfc_size = AH; ++ int wout_size = AH; ++ int bout_size = 1; ++ int total_params = wfc_size + bfc_size + wout_size + bout_size; ++ ++ int widx = blockIdx.x * blockDim.x + threadIdx.x; ++ if (widx >= total_params) return; ++ ++ float acc = 0.0f; ++ ++ for (int i = 0; i < B; i++) { ++ float R = risk_budget[i]; ++ ++ float d_R = 0.0f; ++ const float* dq = d_q_mag + (long long)i * b1_size; ++ const float* qp = q_mag_pre + (long long)i * b1_size; ++ if (b1_size > 1) d_R += dq[1] * qp[1] * 0.5f / fmaxf(sqrtf(fmaxf(R, 1e-6f)), 1e-6f); ++ if (b1_size > 2) d_R += dq[2] * qp[2]; ++ ++ float d_raw = d_R * R * (1.0f - R); ++ ++ if (widx < wfc_size) { ++ /* d_w_risk_fc[j * input_dim + k] */ ++ int j = widx / input_dim; ++ int k = widx % input_dim; ++ ++ float d_hj = d_raw * w_risk_out[j]; ++ const float* h_risk = risk_hidden + (long long)i * AH; ++ if (h_risk[j] <= 0.0f) d_hj = 0.0f; ++ ++ float input_val; ++ if (k < SH2) { ++ input_val = h_s2[(long long)i * SH2 + k]; ++ } else if (k < SH2 + 12) { ++ input_val = isv_signals_dev_ptr[k - SH2]; ++ } else { ++ input_val = predicted_error[i]; ++ } ++ acc += d_hj * input_val; ++ } else if (widx < wfc_size + bfc_size) { ++ /* d_b_risk_fc[j] */ ++ int j = widx - wfc_size; ++ float d_hj = d_raw * w_risk_out[j]; ++ const float* h_risk = risk_hidden + (long long)i * AH; ++ if (h_risk[j] <= 0.0f) d_hj = 0.0f; ++ acc += d_hj; ++ } else if (widx < wfc_size + bfc_size + wout_size) { ++ /* d_w_risk_out[j] */ ++ int j = widx - wfc_size - bfc_size; ++ const float* h_risk = risk_hidden + (long long)i * AH; ++ acc += d_raw * h_risk[j]; ++ } else { ++ /* d_b_risk_out */ ++ acc += d_raw; ++ } ++ } ++ ++ /* Single deterministic write per weight element */ ++ if (widx < wfc_size) ++ d_w_risk_fc[widx] = acc; ++ else if (widx < wfc_size + bfc_size) ++ d_b_risk_fc[widx - wfc_size] = acc; ++ else if (widx < wfc_size + bfc_size + wout_size) ++ d_w_risk_out[widx - wfc_size - bfc_size] = acc; ++ else ++ d_b_risk_out[0] = acc; ++} ++ ++/* ================================================================== */ ++/* Kernel: isv_signal_update — pure GPU ISV signal computation */ ++/* ================================================================== */ ++ ++extern "C" __global__ void isv_signal_update( ++ float* __restrict__ isv_signals, /* [12] pinned device-mapped */ ++ float* __restrict__ isv_history, /* [K*12] pinned device-mapped */ ++ float* __restrict__ lagged_td_error, /* [1] pinned — for recursive confidence target */ ++ const float* __restrict__ q_mean_ptr, /* [1] pinned — batch Q-mean */ ++ const float* __restrict__ q_mean_ema_ptr, /* [1] pinned — Q-mean EMA */ ++ const float* __restrict__ grad_norm_ptr, /* [1] pinned — gradient norm² */ ++ const float* __restrict__ td_error_ptr, /* [1] pinned — batch mean |TD-error| */ ++ const float* __restrict__ ens_var_ptr, /* [1] pinned — batch mean ensemble var */ ++ const float* __restrict__ reward_ptr, /* [1] pinned — reward proxy */ ++ const float* __restrict__ atom_util_ptr, /* [1] pinned — atom utilization */ ++ const float* __restrict__ loss_ptr, /* [1] pinned — total loss */ ++ int K, /* temporal history length (4) */ ++ const float* __restrict__ states_ptr, /* [B, state_dim] — for ADX/CUSUM regime signals. NULL = skip regime. */ ++ int state_dim /* to index ADX at [40], CUSUM at [41] */ ++) { ++ if (threadIdx.x != 0) return; ++ ++ float alpha = 0.05f; ++ ++ /* Save current td_error_ema as lagged target for recursive confidence */ ++ lagged_td_error[0] = isv_signals[2]; ++ ++ /* Shift temporal history: newest at [0], oldest at [K-1] */ ++ for (int t = K - 1; t > 0; t--) ++ for (int k = 0; k < 12; k++) ++ isv_history[t * 12 + k] = isv_history[(t - 1) * 12 + k]; ++ for (int k = 0; k < 12; k++) ++ isv_history[k] = isv_signals[k]; ++ ++ /* [0] Q-drift: normalized deviation from EMA */ ++ float q_mean = q_mean_ptr[0]; ++ float q_ema = q_mean_ema_ptr[0]; ++ isv_signals[0] = (q_mean - q_ema) / fmaxf(fabsf(q_ema), 0.1f); ++ ++ /* [1] Gradient norm EMA (ptr stores norm², take sqrt) */ ++ float gn = sqrtf(fmaxf(grad_norm_ptr[0], 0.0f)); ++ isv_signals[1] = (1.0f - alpha) * isv_signals[1] + alpha * gn; ++ ++ /* [2] TD-error EMA */ ++ isv_signals[2] = (1.0f - alpha) * isv_signals[2] + alpha * td_error_ptr[0]; ++ ++ /* [3] Ensemble variance EMA (save old for delta) */ ++ float old_ens_ema = isv_signals[3]; ++ isv_signals[3] = (1.0f - alpha) * old_ens_ema + alpha * ens_var_ptr[0]; ++ ++ /* [4] Ensemble variance velocity */ ++ isv_signals[4] = (isv_signals[3] - old_ens_ema) / fmaxf(old_ens_ema, 1e-6f); ++ ++ /* [5] Reward EMA */ ++ isv_signals[5] = (1.0f - alpha) * isv_signals[5] + alpha * reward_ptr[0]; ++ ++ /* [6] Atom utilization EMA */ ++ isv_signals[6] = (1.0f - alpha) * isv_signals[6] + alpha * atom_util_ptr[0]; ++ ++ /* [7] Loss EMA */ ++ isv_signals[7] = (1.0f - alpha) * isv_signals[7] + alpha * loss_ptr[0]; ++ ++ /* ── Regime awareness signals [8]-[11] from ADX/CUSUM ──────────── */ ++ if (states_ptr != 0 && state_dim > 41) { ++ float adx = states_ptr[0 * state_dim + 40]; ++ float cusum = states_ptr[0 * state_dim + 41]; ++ ++ /* [8] Regime velocity: EMA of ADX (used for delta computation next step) */ ++ float prev_adx = isv_signals[8]; ++ float adx_delta = fabsf(adx - prev_adx); ++ float prev_cusum_ema = isv_signals[10]; ++ float cusum_delta = fabsf(cusum - prev_cusum_ema); ++ float regime_vel = adx_delta + cusum_delta; ++ ++ isv_signals[8] = (1.0f - alpha) * isv_signals[8] + alpha * adx; ++ float velocity_ema = (1.0f - alpha) * prev_cusum_ema + alpha * regime_vel; ++ ++ /* [9] Regime disagreement: |norm(ADX) - norm(CUSUM)| */ ++ float sum_abs = fabsf(adx) + fabsf(cusum) + 1e-6f; ++ isv_signals[9] = fabsf(adx / sum_abs - cusum / sum_abs); ++ ++ /* [10] Regime transition EMA */ ++ isv_signals[10] = velocity_ema; ++ ++ /* [11] Regime stability: 1 - sigmoid(5 * velocity) */ ++ isv_signals[11] = 1.0f - 1.0f / (1.0f + expf(-5.0f * regime_vel)); ++ } else { ++ /* No states available — neutral defaults */ ++ isv_signals[8] = 0.0f; ++ isv_signals[9] = 0.0f; ++ isv_signals[10] = 0.0f; ++ isv_signals[11] = 0.5f; /* moderate stability */ ++ } ++} ++ ++/* ================================================================== */ ++/* Kernel: isv_forward — ISV encoder MLP + branch gate + gamma mod */ ++/* ================================================================== */ ++ ++extern "C" __global__ void isv_forward( ++ const float* __restrict__ isv_signals, /* [12] pinned */ ++ const float* __restrict__ isv_history, /* [K*12] pinned */ ++ const float* __restrict__ isv_decay, /* [12] pinned learned decay */ ++ const float* __restrict__ w_fc1, /* [16, 12] */ ++ const float* __restrict__ b_fc1, /* [16] */ ++ const float* __restrict__ w_fc2, /* [8, 16] */ ++ const float* __restrict__ b_fc2, /* [8] */ ++ const float* __restrict__ w_gate, /* [4, 8] */ ++ const float* __restrict__ b_gate, /* [4] */ ++ const float* __restrict__ w_gamma, /* [8] (flat) */ ++ const float* __restrict__ b_gamma, /* [1] */ ++ float* __restrict__ isv_embedding_out, /* [8] */ ++ float* __restrict__ branch_gate_out, /* [4] */ ++ float* __restrict__ gamma_mod_out, /* [1] */ ++ int K ++) { ++ if (threadIdx.x != 0) return; ++ ++ /* Step 1: Temporal ISV — decay-weighted average over 12 signals */ ++ float isv_temporal[12]; ++ for (int k = 0; k < 12; k++) { ++ float decay = 1.0f / (1.0f + expf(-isv_decay[k])); ++ float weighted_sum = isv_signals[k]; ++ float weight_sum = 1.0f; ++ for (int t = 0; t < K; t++) { ++ float w = powf(decay, (float)(t + 1)); ++ weighted_sum += w * isv_history[t * 12 + k]; ++ weight_sum += w; ++ } ++ isv_temporal[k] = weighted_sum / weight_sum; ++ } ++ ++ /* Step 2: MLP — FC1 (12→16, ReLU) → FC2 (16→8) */ ++ float hidden[16]; ++ for (int j = 0; j < 16; j++) { ++ float val = b_fc1[j]; ++ for (int k = 0; k < 12; k++) ++ val += w_fc1[j * 12 + k] * isv_temporal[k]; ++ hidden[j] = fmaxf(val, 0.0f); ++ } ++ float embedding[8]; ++ for (int k = 0; k < 8; k++) { ++ float val = b_fc2[k]; ++ for (int j = 0; j < 16; j++) ++ val += w_fc2[k * 16 + j] * hidden[j]; ++ embedding[k] = val; ++ isv_embedding_out[k] = val; ++ } ++ ++ /* Step 3: Branch gating — softmax(W_gate @ embedding + b_gate) */ ++ float logits[4]; ++ float max_logit = -1e9f; ++ for (int d = 0; d < 4; d++) { ++ float val = b_gate[d]; ++ for (int k = 0; k < 8; k++) ++ val += w_gate[d * 8 + k] * embedding[k]; ++ logits[d] = val; ++ max_logit = fmaxf(max_logit, logits[d]); ++ } ++ float sum_exp = 0.0f; ++ for (int d = 0; d < 4; d++) { ++ logits[d] = expf(logits[d] - max_logit); ++ sum_exp += logits[d]; ++ } ++ for (int d = 0; d < 4; d++) ++ branch_gate_out[d] = logits[d] / sum_exp; ++ ++ /* Step 4: Drift-conditioned gamma — 0.5 + 0.5*sigmoid(w@emb+b) → [0.5, 1.0] */ ++ float gamma_raw = b_gamma[0]; ++ for (int k = 0; k < 8; k++) ++ gamma_raw += w_gamma[k] * embedding[k]; ++ gamma_mod_out[0] = 0.5f + 0.5f / (1.0f + expf(-gamma_raw)); ++} ++ ++/* ================================================================== */ ++/* Kernel: fill_gamma_buf — broadcast scalar gamma to per-sample buf */ ++/* ================================================================== */ ++extern "C" __global__ void fill_gamma_buf( ++ float* __restrict__ gamma_buf, /* [B] output */ ++ float base_gamma, ++ const float* __restrict__ gamma_mod, /* [1] from isv_forward */ ++ int B + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= B) return; ++ gamma_buf[i] = base_gamma * gamma_mod[0]; ++} + +- float R = risk_budget[i]; ++/* ================================================================== */ ++/* Kernel: branch_confidence_routing — ISV gate × Q-value confidence */ ++/* ================================================================== */ ++extern "C" __global__ void branch_confidence_routing( ++ float* __restrict__ q_values, /* [B, total_actions] in-place */ ++ const float* __restrict__ branch_gate, /* [4] shared from isv_forward */ ++ int B, ++ int b0_size, int b1_size, int b2_size, int b3_size ++) { ++ int i = blockIdx.x * blockDim.x + threadIdx.x; ++ if (i >= B) return; + +- float d_R = 0.0f; +- const float* dq = d_q_mag + (long long)i * b1_size; +- const float* qp = q_mag_pre + (long long)i * b1_size; +- if (b1_size > 1) d_R += dq[1] * qp[1] * 0.5f / fmaxf(sqrtf(fmaxf(R, 1e-6f)), 1e-6f); +- if (b1_size > 2) d_R += dq[2] * qp[2]; ++ int branch_sizes[4] = { b0_size, b1_size, b2_size, b3_size }; ++ int total_actions = b0_size + b1_size + b2_size + b3_size; ++ int branch_offset = 0; ++ ++ for (int d = 0; d < 4; d++) { ++ int A_d = branch_sizes[d]; ++ ++ float q_max = -1e9f; ++ float q_sum = 0.0f; ++ for (int a = 0; a < A_d; a++) { ++ float q = q_values[(long long)i * total_actions + branch_offset + a]; ++ q_max = fmaxf(q_max, q); ++ q_sum += q; ++ } ++ float q_mean_branch = q_sum / (float)A_d; ++ float separation = q_max - q_mean_branch; ++ float confidence = 1.0f / (1.0f + expf(-5.0f * separation)); + +- float d_raw = d_R * R * (1.0f - R); ++ float isv_gate = branch_gate[d]; ++ float effective_weight = isv_gate * fmaxf(confidence, 0.3f); + +- const float* h_risk = risk_hidden + (long long)i * AH; +- atomicAdd(d_b_risk_out, d_raw); +- for (int j = 0; j < AH; j++) { +- atomicAdd(&d_w_risk_out[j], d_raw * h_risk[j]); ++ for (int a = 0; a < A_d; a++) { ++ q_values[(long long)i * total_actions + branch_offset + a] *= effective_weight; ++ } ++ branch_offset += A_d; + } ++} ++ ++/* ================================================================== */ ++/* Kernel: recursive_confidence_forward — predict own TD-error */ ++/* ================================================================== */ ++extern "C" __global__ void recursive_confidence_forward( ++ const float* __restrict__ h_s2, /* [B, SH2] */ ++ const float* __restrict__ w_conf, /* [SH2] */ ++ const float* __restrict__ b_conf, /* [1] */ ++ float* __restrict__ predicted_error, /* [B] */ ++ int B, int SH2 ++) { ++ int i = blockIdx.x * blockDim.x + threadIdx.x; ++ if (i >= B) return; + + const float* h = h_s2 + (long long)i * SH2; +- for (int j = 0; j < AH; j++) { +- float d_hj = d_raw * w_risk_out[j]; +- if (h_risk[j] <= 0.0f) d_hj = 0.0f; ++ float val = b_conf[0]; ++ for (int k = 0; k < SH2; k++) ++ val += w_conf[k] * h[k]; ++ predicted_error[i] = 1.0f / (1.0f + expf(-val)); ++} + +- atomicAdd(&d_b_risk_fc[j], d_hj); ++/* ================================================================== */ ++/* Kernel: recursive_confidence_backward — MSE grad into trunk */ ++/* ================================================================== */ ++/** ++ * Phase 1: Per-sample trunk gradient + warp-reduced weight gradient. ++ * d_h_s2[i, k] += d_sigmoid * w_conf[k] (plain write — single writer per (i,k)). ++ * d_w_conf and d_b_conf use warp+block reduction → one atomicAdd per BLOCK. ++ * ++ * Grid: ceil(B/256), Block: 256. ++ */ ++extern "C" __global__ void recursive_confidence_backward( ++ const float* __restrict__ h_s2, /* [B, SH2] */ ++ const float* __restrict__ predicted_error, /* [B] */ ++ const float* __restrict__ lagged_td_error, /* [1] pinned — target */ ++ const float* __restrict__ w_conf, /* [SH2] */ ++ float* __restrict__ d_w_conf, /* [SH2] gradient accumulator */ ++ float* __restrict__ d_b_conf, /* [1] gradient accumulator */ ++ float* __restrict__ d_h_s2, /* [B, SH2] trunk gradient (accumulate) */ ++ int B, int SH2, ++ float loss_weight /* 0.01 */ ++) { ++ int i = blockIdx.x * blockDim.x + threadIdx.x; ++ ++ float d_sigmoid = 0.0f; ++ if (i < B) { ++ float pred = predicted_error[i]; ++ float target = lagged_td_error[0]; ++ float d_loss = loss_weight * 2.0f * (pred - target) / (float)B; ++ d_sigmoid = d_loss * pred * (1.0f - pred); ++ ++ /* Per-sample trunk gradient — no cross-sample contention, plain write */ + for (int k = 0; k < SH2; k++) { +- atomicAdd(&d_w_risk_fc[(long long)j * SH2 + k], d_hj * h[k]); ++ d_h_s2[(long long)i * SH2 + k] += d_sigmoid * w_conf[k]; + } + } ++ ++ /* Weight gradient reduction: d_b_conf = sum_b(d_sigmoid_b) */ ++ float my_db = d_sigmoid; ++ for (int offset = 16; offset > 0; offset >>= 1) ++ my_db += __shfl_xor_sync(0xFFFFFFFF, my_db, offset); ++ ++ __shared__ float ws_db[8]; ++ int warp_id = threadIdx.x / 32; ++ int lane = threadIdx.x % 32; ++ if (lane == 0) ws_db[warp_id] = my_db; ++ __syncthreads(); ++ ++ if (warp_id == 0) { ++ float val = (lane < blockDim.x / 32) ? ws_db[lane] : 0.0f; ++ for (int off = 16; off > 0; off >>= 1) ++ val += __shfl_xor_sync(0xFFFFFFFF, val, off); ++ if (lane == 0) atomicAdd(d_b_conf, val); ++ } ++ ++ /* Weight gradient reduction: d_w_conf[k] = sum_b(d_sigmoid_b * h_s2[b, k]) */ ++ for (int k = 0; k < SH2; k++) { ++ float my_dw = (i < B) ? d_sigmoid * h_s2[(long long)i * SH2 + k] : 0.0f; ++ ++ for (int offset = 16; offset > 0; offset >>= 1) ++ my_dw += __shfl_xor_sync(0xFFFFFFFF, my_dw, offset); ++ ++ __shared__ float ws_dw[8]; ++ if (lane == 0) ws_dw[warp_id] = my_dw; ++ __syncthreads(); ++ ++ if (warp_id == 0) { ++ float val = (lane < blockDim.x / 32) ? ws_dw[lane] : 0.0f; ++ for (int off = 16; off > 0; off >>= 1) ++ val += __shfl_xor_sync(0xFFFFFFFF, val, off); ++ if (lane == 0) atomicAdd(&d_w_conf[k], val); ++ } ++ __syncthreads(); ++ } ++} ++ ++/* ================================================================== */ ++/* Kernel: isv_feature_gate — ISV embedding modulates h_s2 features */ ++/* ================================================================== */ ++extern "C" __global__ void isv_feature_gate( ++ float* __restrict__ h_s2, /* [B, SH2] in-place */ ++ const float* __restrict__ isv_embedding, /* [8] shared (from isv_forward) */ ++ const float* __restrict__ w_feature_gate, /* [SH2, 8] */ ++ const float* __restrict__ b_feature_gate, /* [SH2] */ ++ int B, int SH2 ++) { ++ int i = blockIdx.x * blockDim.x + threadIdx.x; ++ if (i >= B) return; ++ ++ for (int j = 0; j < SH2; j++) { ++ float gate_val = b_feature_gate[j]; ++ for (int k = 0; k < 8; k++) ++ gate_val += w_feature_gate[(long long)j * 8 + k] * isv_embedding[k]; ++ float gate = 1.0f / (1.0f + expf(-gate_val)); ++ h_s2[(long long)i * SH2 + j] *= gate; ++ } ++} ++ ++/* ================================================================== */ ++/* Kernel: trade_plan_forward — plan head MLP → 6 plan parameters */ ++/* ================================================================== */ ++extern "C" __global__ void trade_plan_forward( ++ const float* __restrict__ h_s2, /* [B, SH2] */ ++ const float* __restrict__ w_plan_fc, /* [AH, SH2] */ ++ const float* __restrict__ b_plan_fc, /* [AH] */ ++ const float* __restrict__ w_plan_out, /* [6, AH] */ ++ const float* __restrict__ b_plan_out, /* [6] */ ++ float* __restrict__ plan_params, /* [B, 6] output */ ++ int B, int SH2, int AH ++) { ++ int i = blockIdx.x * blockDim.x + threadIdx.x; ++ if (i >= B) return; ++ ++ const float* h = h_s2 + (long long)i * SH2; ++ float hidden[128]; ++ ++ for (int j = 0; j < AH; j++) { ++ float val = b_plan_fc[j]; ++ for (int k = 0; k < SH2; k++) ++ val += w_plan_fc[(long long)j * SH2 + k] * h[k]; ++ hidden[j] = fmaxf(val, 0.0f); ++ } ++ ++ float* out = plan_params + i * 6; ++ for (int p = 0; p < 6; p++) { ++ float val = b_plan_out[p]; ++ for (int j = 0; j < AH; j++) ++ val += w_plan_out[(long long)p * AH + j] * hidden[j]; ++ out[p] = val; ++ } ++ ++ out[0] = 3.0f + 22.0f / (1.0f + expf(-out[0])); ++ out[1] = 0.0005f + 0.0145f / (1.0f + expf(-out[1])); ++ out[2] = 0.0002f + 0.0048f / (1.0f + expf(-out[2])); ++ out[3] = 1.0f / (1.0f + expf(-out[3])); ++ out[4] = 1.0f / (1.0f + expf(-out[4])); ++ out[5] = 0.5f + 2.5f / (1.0f + expf(-out[5])); + } +diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +index 7706b6ad2..a5b8c80f6 100644 +--- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs ++++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +@@ -981,6 +981,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) + .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 b80c31519..60cbd6fd7 100644 +--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs ++++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +@@ -79,6 +79,14 @@ static MAMBA2_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/mamba2_te + /// Mamba2 temporal scan configuration. + const MAMBA2_HISTORY_K: usize = 8; // Rolling history length + const MAMBA2_STATE_DIM: usize = 16; // SSM state dimension ++ ++/// Introspective State Vector (ISV) configuration. ++const ISV_K: usize = 4; // Temporal ISV history length ++const ISV_DIM: usize = 12; // ISV signal count (8 core + 4 regime awareness) ++const ISV_EMB_DIM: usize = 8; // ISV embedding output dimension (FC2 output, gate/gamma input) ++/// First ISV tensor index in the flat param buffer. ++/// ISV weights (68-79) are online-only — NOT synced to the target network. ++const FIRST_ISV_TENSOR: usize = 68; + const NEW_COMPONENT_WARMUP_STEPS: u32 = 500; + + /// Homeostatic regularizer: number of observable signals. +@@ -364,8 +372,10 @@ impl Default for CausalInterventionConfig { + /// + 8 VSN bottleneck (R=16) + 8 GLU gate + 8 KAN spline (coeff+resid per branch) + /// + 2 regime branch gate (w_regime + b_regime) + 4 adaptive atom spacing + /// + 8 multi-horizon value heads (5-bar and 20-bar) +-/// + 4 learned risk management (5th branch) = 68. +-pub(crate) const NUM_WEIGHT_TENSORS: usize = 68; ++/// + 4 learned risk management (5th branch) = 68 ++/// + 10 ISV encoder/conditioning + 2 feature gate + 2 temporal routing = 82 ++/// + 4 trade plan head = 86. ++pub(crate) const NUM_WEIGHT_TENSORS: usize = 86; + + /// Compute the size (element count) of each weight tensor. + /// +@@ -379,6 +389,10 @@ pub(crate) const NUM_WEIGHT_TENSORS: usize = 68; + /// Tensors 52-55: Adaptive atom spacing (spacing_raw per branch, NA each). + /// Tensors 56-63: Multi-horizon value heads (5-bar and 20-bar). + /// Tensors 64-67: Learned risk management (5th branch). ++/// Tensors 68-77: Introspective State Vector (ISV) encoder + conditioning. ++/// Tensors 78-79: ISV feature gate (per-feature h_s2 modulation). ++/// Tensors 80-81: ISV temporal routing (per-feature temporal depth). ++/// Tensors 82-85: Trade plan head (h_s2 → 6 plan parameters). + /// + /// When bottleneck is active, w_s1 input dimension changes from state_dim to + /// (bottleneck_dim + portfolio_dim) where portfolio_dim = state_dim - market_dim. +@@ -466,10 +480,32 @@ pub(crate) fn compute_param_sizes(cfg: &GpuDqnTrainConfig) -> [usize; NUM_WEIGHT + cfg.num_atoms * cfg.value_h, // [62] w_v2_20bar [NA, VH] + cfg.num_atoms, // [63] b_v2_20bar [NA] + // ── Learned risk management (5th branch) ── +- cfg.adv_h * cfg.shared_h2, // [64] w_risk_fc [AH, SH2] ++ cfg.adv_h * (cfg.shared_h2 + 13), // [64] w_risk_fc [AH, SH2+13] + cfg.adv_h, // [65] b_risk_fc [AH] + cfg.adv_h, // [66] w_risk_out [AH] + 1, // [67] b_risk_out [1] ++ // ── Introspective State Vector (ISV) encoder + conditioning ── ++ 16 * ISV_DIM, // [68] w_isv_fc1 [16, ISV_DIM] ++ 16, // [69] b_isv_fc1 [16] ++ ISV_EMB_DIM * 16, // [70] w_isv_fc2 [EMB, 16] ++ ISV_EMB_DIM, // [71] b_isv_fc2 [EMB] ++ 4 * ISV_EMB_DIM, // [72] w_isv_gate [4, EMB] ++ 4, // [73] b_isv_gate [4] ++ ISV_EMB_DIM, // [74] w_isv_gamma [EMB] ++ 1, // [75] b_isv_gamma [1] ++ cfg.shared_h2, // [76] w_conf_fc [SH2] ++ 1, // [77] b_conf_fc [1] ++ // ── ISV Feature Gating (Phase 3) ── ++ cfg.shared_h2 * ISV_EMB_DIM, // [78] w_feature_gate [SH2, 8] ++ cfg.shared_h2, // [79] b_feature_gate [SH2] ++ // ── ISV Temporal Routing (Phase 3) ── ++ cfg.shared_h2 * ISV_EMB_DIM, // [80] w_temporal_route [SH2, 8] ++ cfg.shared_h2, // [81] b_temporal_route [SH2] ++ // ── Trade Plan Head ── ++ cfg.adv_h * cfg.shared_h2, // [82] w_plan_fc [AH, SH2] ++ cfg.adv_h, // [83] b_plan_fc [AH] ++ 6 * cfg.adv_h, // [84] w_plan_out [6, AH] ++ 6, // [85] b_plan_out [6] + ] + } + +@@ -485,6 +521,12 @@ pub(crate) fn compute_total_params(cfg: &GpuDqnTrainConfig) -> usize { + compute_param_sizes(cfg).iter().map(|&s| align4(s)).sum() + } + ++/// Non-ISV parameter count (indices 0..FIRST_ISV_TENSOR). ++/// EMA target sync is restricted to this range — ISV weights are online-only. ++pub(crate) fn compute_non_isv_params(cfg: &GpuDqnTrainConfig) -> usize { ++ compute_param_sizes(cfg)[..FIRST_ISV_TENSOR].iter().map(|&s| align4(s)).sum() ++} ++ + /// Cumulative byte offset to tensor `idx` in the padded flat buffer. + pub(crate) fn padded_byte_offset(param_sizes: &[usize], idx: usize) -> u64 { + param_sizes[..idx].iter() +@@ -733,8 +775,8 @@ pub struct GpuDqnTrainer { + kan_gate_combine_kernel: CudaFunction, + kan_gate_backward_kernel: CudaFunction, + +- // ── Regime branch gate ── +- regime_gate_kernel: CudaFunction, ++ // ── Branch confidence routing (ISV gate × Q-value confidence) ── ++ branch_confidence_routing_kernel: CudaFunction, + /// Per-sample Q-gap for regime gate input [B]. + regime_q_gap_buf: CudaSlice, + /// Atom utilization scalar for regime gate [1] — pinned device-mapped. +@@ -878,6 +920,8 @@ pub struct GpuDqnTrainer { + // ── Training state ────────────────────────────────────────────── + pub(crate) adam_step: i32, + total_params: usize, ++ /// Non-ISV parameter count: EMA target sync restricted to this range. ++ non_isv_params: usize, + pub(crate) params_initialized: bool, + target_params_initialized: bool, + attention_initialized: bool, +@@ -1296,6 +1340,57 @@ pub struct GpuDqnTrainer { + q_mag_pre_risk: CudaSlice, // [B, b1_size] saved for backward + risk_grad_buf: CudaSlice, // [risk_param_count] + risk_adam_step: i32, ++ ++ // ── ISV scratch scalars (pinned device-mapped, written by GPU kernels) ── ++ // TODO(isv): Wire C51 loss kernel to write mean |TD-error| here ++ td_error_scratch_pinned: *mut f32, ++ td_error_scratch_dev_ptr: u64, ++ // TODO(isv): Wire ensemble aggregate kernel to write mean variance here ++ ensemble_var_scratch_pinned: *mut f32, ++ ensemble_var_scratch_dev_ptr: u64, ++ reward_scratch_pinned: *mut f32, ++ reward_scratch_dev_ptr: u64, ++ atom_util_scratch_pinned: *mut f32, ++ atom_util_scratch_dev_ptr: u64, ++ ++ // ── ISV core buffers (pinned device-mapped, GPU read/write) ── ++ isv_signals_pinned: *mut f32, // [ISV_DIM = 12] ++ isv_signals_dev_ptr: u64, ++ isv_history_pinned: *mut f32, // [ISV_K * ISV_DIM = 48] ++ isv_history_dev_ptr: u64, ++ isv_decay_pinned: *mut f32, // [ISV_DIM = 12] learned decay weights ++ isv_decay_dev_ptr: u64, ++ lagged_td_error_pinned: *mut f32, // [1] recursive confidence target ++ lagged_td_error_dev_ptr: u64, ++ isv_signal_update_kernel: CudaFunction, ++ ++ // ── ISV forward outputs ── ++ isv_forward_kernel: CudaFunction, ++ isv_feature_gate_kernel: CudaFunction, ++ fill_gamma_buf_kernel: CudaFunction, ++ isv_embedding_buf: CudaSlice, // [ISV_EMB_DIM] = [8] ++ branch_gate_buf: CudaSlice, // [4] ++ gamma_mod_buf: CudaSlice, // [1] ++ temporal_weight_buf: CudaSlice, // [SH2] per-feature temporal routing ++ isv_temporal_route_kernel: CudaFunction, ++ gamma_buf: CudaSlice, // [B] per-sample effective gamma ++ predicted_error_buf: CudaSlice, // [B] recursive confidence output ++ ++ // ── Recursive confidence kernels ── ++ recursive_conf_fwd_kernel: CudaFunction, ++ recursive_conf_bwd_kernel: CudaFunction, ++ ++ // ── Trade plan head ── ++ plan_params_buf: CudaSlice, // [B, 6] ++ trade_plan_fwd_kernel: CudaFunction, ++ ++ // ── Speculative inference cache ── ++ /// Cached trunk output [B, SH2] from between-bar speculative forward. ++ speculative_h_s2: CudaSlice, ++ /// Last features [state_dim] used for speculation (host-side for L2 comparison). ++ speculative_features: Vec, ++ /// Whether the speculative cache contains a usable result. ++ speculative_valid: bool, + } + + impl GpuDqnTrainer { +@@ -1589,6 +1684,30 @@ impl Drop for GpuDqnTrainer { + if !self.iqn_readiness_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.iqn_readiness_pinned.cast()) }; + } ++ if !self.td_error_scratch_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.td_error_scratch_pinned.cast()) }; ++ } ++ if !self.ensemble_var_scratch_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.ensemble_var_scratch_pinned.cast()) }; ++ } ++ if !self.reward_scratch_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.reward_scratch_pinned.cast()) }; ++ } ++ if !self.atom_util_scratch_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.atom_util_scratch_pinned.cast()) }; ++ } ++ if !self.isv_signals_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.isv_signals_pinned.cast()) }; ++ } ++ if !self.isv_history_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.isv_history_pinned.cast()) }; ++ } ++ if !self.isv_decay_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.isv_decay_pinned.cast()) }; ++ } ++ if !self.lagged_td_error_pinned.is_null() { ++ let _ = unsafe { cudarc::driver::result::free_host(self.lagged_td_error_pinned.cast()) }; ++ } + } + } + +@@ -1930,6 +2049,248 @@ impl GpuDqnTrainer { + unsafe { *self.q_mean_ema_pinned = new_ema; } + } + ++ /// Pure GPU ISV signal update — reads all pinned pointers, updates ISV [12] + history [K,12]. ++ /// Includes 4 regime awareness signals computed from ADX/CUSUM in states buffer. ++ /// Call after graph_adam replay and stream sync, before next graph_forward. ++ pub(crate) fn update_isv_signals(&self) -> Result<(), MLError> { ++ let k = ISV_K as i32; ++ let state_dim = self.config.state_dim as i32; ++ unsafe { ++ self.stream.launch_builder(&self.isv_signal_update_kernel) ++ .arg(&self.isv_signals_dev_ptr) ++ .arg(&self.isv_history_dev_ptr) ++ .arg(&self.lagged_td_error_dev_ptr) ++ .arg(&self.q_mean_scratch_dev_ptr) ++ .arg(&self.q_mean_ema_dev_ptr) ++ .arg(&self.grad_norm_dev_ptr) ++ .arg(&self.td_error_scratch_dev_ptr) ++ .arg(&self.ensemble_var_scratch_dev_ptr) ++ .arg(&self.reward_scratch_dev_ptr) ++ .arg(&self.atom_util_scratch_dev_ptr) ++ .arg(&self.total_loss_dev_ptr) ++ .arg(&k) ++ .arg(&self.ptrs.states_buf) // states for ADX/CUSUM regime signals ++ .arg(&state_dim) ++ .launch(LaunchConfig { ++ grid_dim: (1, 1, 1), ++ block_dim: (1, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("isv_signal_update: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// ISV encoder MLP forward: temporal decay → MLP → branch gate + gamma mod. ++ /// Reads ISV signals/history (pinned), ISV encoder weights from param buffer. ++ /// Writes isv_embedding_buf [ISV_EMB_DIM=8], branch_gate_buf [4], gamma_mod_buf [1]. ++ pub(crate) fn launch_isv_forward(&self) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_fc1 = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 68); ++ let b_fc1 = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 69); ++ let w_fc2 = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 70); ++ let b_fc2 = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 71); ++ let w_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 72); ++ let b_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 73); ++ let w_gamma = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 74); ++ let b_gamma = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 75); ++ let k = ISV_K as i32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.isv_forward_kernel) ++ .arg(&self.isv_signals_dev_ptr) ++ .arg(&self.isv_history_dev_ptr) ++ .arg(&self.isv_decay_dev_ptr) ++ .arg(&w_fc1).arg(&b_fc1) ++ .arg(&w_fc2).arg(&b_fc2) ++ .arg(&w_gate).arg(&b_gate) ++ .arg(&w_gamma).arg(&b_gamma) ++ .arg(&self.isv_embedding_buf) ++ .arg(&self.branch_gate_buf) ++ .arg(&self.gamma_mod_buf) ++ .arg(&k) ++ .launch(LaunchConfig { ++ grid_dim: (1, 1, 1), ++ block_dim: (1, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("isv_forward: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// ISV feature gate: modulate h_s2 features based on ISV embedding. ++ /// h_s2[i,j] *= sigmoid(w_feature_gate[j,:] @ isv_embedding + b_feature_gate[j]) ++ /// Operates in-place on save_h_s2. Must be called AFTER launch_isv_forward. ++ pub(crate) fn launch_isv_feature_gate(&self, batch_size: usize) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 78); ++ let b_gate = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 79); ++ let blocks = ((batch_size as u32 + 255) / 256).max(1); ++ let b_i32 = batch_size as i32; ++ let sh2 = self.config.shared_h2 as i32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.isv_feature_gate_kernel) ++ .arg(&self.save_h_s2) // h_s2 in-place ++ .arg(&self.isv_embedding_buf) ++ .arg(&w_gate) ++ .arg(&b_gate) ++ .arg(&b_i32) ++ .arg(&sh2) ++ .launch(LaunchConfig { ++ grid_dim: (blocks, 1, 1), ++ block_dim: (256, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("isv_feature_gate: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// ISV temporal routing: ISV embedding [8] → temporal_weight [SH2]. ++ /// Each feature gets a learned "how much history to use" weight via sigmoid. ++ /// Must be called AFTER launch_isv_forward (produces embedding) and BEFORE mamba2_step. ++ pub(crate) fn launch_isv_temporal_route(&self) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_route = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 80); ++ let b_route = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 81); ++ let sh2 = self.config.shared_h2 as i32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.isv_temporal_route_kernel) ++ .arg(&self.isv_embedding_buf) ++ .arg(&w_route) ++ .arg(&b_route) ++ .arg(&self.temporal_weight_buf) ++ .arg(&sh2) ++ .launch(LaunchConfig { ++ grid_dim: (1, 1, 1), ++ block_dim: (1, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("isv_temporal_route: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// Recursive confidence forward: predict own TD-error from h_s2. ++ /// h_s2 → sigmoid(w_conf @ h + b_conf) → predicted_error [B]. ++ pub(crate) fn launch_recursive_confidence_forward(&self, batch_size: usize) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_conf = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 76); ++ let b_conf = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 77); ++ let blocks = ((batch_size as u32 + 255) / 256).max(1); ++ let b_i32 = batch_size as i32; ++ let sh2 = self.config.shared_h2 as i32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.recursive_conf_fwd_kernel) ++ .arg(&self.save_h_s2) ++ .arg(&w_conf) ++ .arg(&b_conf) ++ .arg(&self.predicted_error_buf) ++ .arg(&b_i32) ++ .arg(&sh2) ++ .launch(LaunchConfig { ++ grid_dim: (blocks, 1, 1), ++ block_dim: (256, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("recursive_confidence_forward: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// Trade plan forward: h_s2 → hidden[AH] (ReLU) → plan_params[B, 6]. ++ /// Output activations: target_bars[3-25], profit_target[0.05-1.5%], ++ /// stop_loss[0.02-0.5%], scale_aggression[0-1], conviction[0-1], asymmetry[0.5-3.0]. ++ pub(crate) fn launch_trade_plan_forward(&self, batch_size: usize) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_fc = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 82); ++ let b_fc = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 83); ++ let w_out = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 84); ++ let b_out = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 85); ++ let blocks = ((batch_size as u32 + 255) / 256).max(1); ++ let b_i32 = batch_size as i32; ++ let sh2 = self.config.shared_h2 as i32; ++ let ah = self.config.adv_h as i32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.trade_plan_fwd_kernel) ++ .arg(&self.save_h_s2) ++ .arg(&w_fc).arg(&b_fc) ++ .arg(&w_out).arg(&b_out) ++ .arg(&self.plan_params_buf) ++ .arg(&b_i32).arg(&sh2).arg(&ah) ++ .launch(LaunchConfig { ++ grid_dim: (blocks, 1, 1), ++ block_dim: (256, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("trade_plan_forward: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// Recursive confidence backward: MSE loss gradient into trunk + conf weight gradients. ++ /// Accumulates into grad_buf (same buffer Adam reads) and bw_d_h_s2 trunk gradient. ++ pub(crate) fn launch_recursive_confidence_backward(&self, batch_size: usize) -> Result<(), MLError> { ++ let param_sizes = compute_param_sizes(&self.config); ++ let w_conf = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 76); ++ // Gradient accumulators for w_conf and b_conf in main grad_buf ++ let d_w_conf = self.ptrs.grad_buf + padded_byte_offset(¶m_sizes, 76); ++ let d_b_conf = self.ptrs.grad_buf + padded_byte_offset(¶m_sizes, 77); ++ ++ let blocks = ((batch_size as u32 + 255) / 256).max(1); ++ let b_i32 = batch_size as i32; ++ let sh2 = self.config.shared_h2 as i32; ++ let loss_weight = 0.01_f32; ++ ++ unsafe { ++ self.stream.launch_builder(&self.recursive_conf_bwd_kernel) ++ .arg(&self.save_h_s2) ++ .arg(&self.predicted_error_buf) ++ .arg(&self.lagged_td_error_dev_ptr) ++ .arg(&w_conf) ++ .arg(&d_w_conf) ++ .arg(&d_b_conf) ++ .arg(&self.bw_d_h_s2) ++ .arg(&b_i32) ++ .arg(&sh2) ++ .arg(&loss_weight) ++ .launch(LaunchConfig { ++ grid_dim: (blocks, 1, 1), ++ block_dim: (256, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("recursive_confidence_backward: {e}")))?; ++ } ++ Ok(()) ++ } ++ ++ /// Broadcast base_gamma * gamma_mod[0] into gamma_buf [B] for per-sample C51 loss. ++ pub(crate) fn fill_gamma_buf(&self) -> Result<(), MLError> { ++ let base_gamma = self.adaptive_gamma.powi(self.config.n_steps as i32); ++ let blocks = ((self.config.batch_size as u32 + 255) / 256).max(1); ++ let b = self.config.batch_size as i32; ++ let gamma_mod_ptr = self.gamma_mod_buf.raw_ptr(); ++ unsafe { ++ self.stream.launch_builder(&self.fill_gamma_buf_kernel) ++ .arg(&self.gamma_buf) ++ .arg(&base_gamma) ++ .arg(&gamma_mod_ptr) ++ .arg(&b) ++ .launch(LaunchConfig { ++ grid_dim: (blocks, 1, 1), ++ block_dim: (256, 1, 1), ++ shared_mem_bytes: 0, ++ }) ++ .map_err(|e| MLError::ModelError(format!("fill_gamma_buf: {e}")))?; ++ } ++ Ok(()) ++ } ++ + /// LR scale factor for new components: min(1.0, step / 500). + fn new_component_lr_scale(&self) -> f32 { + (self.new_component_warmup_step as f32 / NEW_COMPONENT_WARMUP_STEPS as f32).min(1.0) +@@ -1966,48 +2327,34 @@ impl GpuDqnTrainer { + Ok(()) + } + +- /// Apply regime branch gate to Q-values (after Q-mean centering, before Q-attention). ++ /// Apply branch confidence routing: ISV gate × Q-value separation confidence. + /// +- /// Uses W_regime[4,4] + b_regime[4] to produce softmax importance weights from +- /// [ADX, CUSUM, Q_gap, atom_utilization]. Scales per-branch Q-values so direction +- /// matters more in trends, magnitude matters more in ranges. +- pub(crate) fn apply_regime_gate(&self, batch_size: usize) -> Result<(), MLError> { +- let param_sizes = compute_param_sizes(&self.config); +- let w_regime_ptr = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 50); +- let b_regime_ptr = self.ptrs.params_ptr + padded_byte_offset(¶m_sizes, 51); +- ++ /// For each branch, computes confidence = sigmoid(5 * (Q_max - Q_mean)) and ++ /// scales Q-values by isv_gate[d] * max(confidence, 0.3). Market regime info ++ /// (ADX, CUSUM) still flows through the trunk → branches; ISV adds training ++ /// dynamics awareness on top. ++ pub(crate) fn apply_branch_confidence_routing(&self, batch_size: usize) -> Result<(), MLError> { + let blocks = ((batch_size as u32 + 255) / 256).max(1); + let q_out_ptr = self.q_out_buf.raw_ptr(); +- let states_ptr = self.ptrs.states_buf; +- let q_gap_ptr = self.regime_q_gap_buf.raw_ptr(); +- let util_ptr = self.regime_util_dev_ptr; ++ let gate_ptr = self.branch_gate_buf.raw_ptr(); + let b_i32 = batch_size as i32; +- let sd = self.config.state_dim as i32; + let b0 = self.config.branch_0_size as i32; + let b1 = self.config.branch_1_size as i32; + let b2 = self.config.branch_2_size as i32; + let b3 = self.config.branch_3_size as i32; + + unsafe { +- self.stream.launch_builder(&self.regime_gate_kernel) ++ self.stream.launch_builder(&self.branch_confidence_routing_kernel) + .arg(&q_out_ptr) +- .arg(&states_ptr) +- .arg(&q_gap_ptr) +- .arg(&util_ptr) +- .arg(&w_regime_ptr) +- .arg(&b_regime_ptr) ++ .arg(&gate_ptr) + .arg(&b_i32) +- .arg(&sd) +- .arg(&b0) +- .arg(&b1) +- .arg(&b2) +- .arg(&b3) ++ .arg(&b0).arg(&b1).arg(&b2).arg(&b3) + .launch(LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) +- .map_err(|e| MLError::ModelError(format!("regime_branch_gate launch: {e}")))?; ++ .map_err(|e| MLError::ModelError(format!("branch_confidence_routing: {e}")))?; + } + Ok(()) + } +@@ -2024,6 +2371,8 @@ impl GpuDqnTrainer { + unsafe { + self.stream.launch_builder(&self.risk_forward_kernel) + .arg(&self.save_h_s2) ++ .arg(&self.isv_signals_dev_ptr) // ISV raw [8] ++ .arg(&self.predicted_error_buf) // predicted_error [B] + .arg(&w_fc_ptr) + .arg(&b_fc_ptr) + .arg(&w_out_ptr) +@@ -2082,9 +2431,9 @@ impl GpuDqnTrainer { + /// shared weights. Uses scale_f32_kernel (no memcpy, all device-side). + pub(crate) fn step_risk_sgd(&mut self) -> Result<(), MLError> { + let param_sizes = compute_param_sizes(&self.config); +- // Risk tensors are indices 64..67 (last 4 in flat param buffer). ++ // Risk tensors are indices 64..67 (4 tensors). + // Total aligned element count across all 4 risk tensors: +- let risk_aligned_count: usize = (64..NUM_WEIGHT_TENSORS) ++ let risk_aligned_count: usize = (64..68) + .map(|i| align4(param_sizes[i])) + .sum(); + +@@ -2145,7 +2494,7 @@ impl GpuDqnTrainer { + /// G5: Apply epistemic gate to magnitude branch Q-values. + /// Scales magnitude Q-values by sigmoid(5 * (var_mean - threshold)): + /// high ensemble variance → conservative (small) magnitude bias. +- /// Must be called AFTER apply_regime_gate and BEFORE launch_q_attention. ++ /// Must be called AFTER apply_branch_confidence_routing and BEFORE launch_q_attention. + pub(crate) fn apply_epistemic_gate(&self, batch_size: usize) -> Result<(), MLError> { + let ta = self.total_actions() as i32; + unsafe { +@@ -2421,6 +2770,8 @@ impl GpuDqnTrainer { + .arg(&(MAMBA2_HISTORY_K as i32)) + .arg(&(sh2 as i32)) + .arg(&(MAMBA2_STATE_DIM as i32)) ++ .arg(&self.isv_signals_dev_ptr) // regime-conditioned decay (ISV[11]) ++ .arg(&self.temporal_weight_buf) // per-feature temporal routing + .launch(LaunchConfig::for_num_elems(batch_size as u32)) + .map_err(|e| MLError::ModelError(format!("mamba2_temporal_scan: {e}")))?; + } +@@ -2463,6 +2814,7 @@ impl GpuDqnTrainer { + } + + /// Mamba2 BPTT: compute gradients for W_A, W_B, W_C. ++ /// Deterministic: one thread per weight element, loops over batch. + /// d_h_enriched comes from the trunk backward gradient (d_save_h_s2). + pub(crate) fn mamba2_backward(&mut self, batch_size: usize) -> Result<(), MLError> { + let sh2 = self.config.shared_h2; +@@ -2476,9 +2828,9 @@ impl GpuDqnTrainer { + let d_w_b = grad_ptr + (sh2 * MAMBA2_STATE_DIM * 4) as u64; + let d_w_c = grad_ptr + (2 * sh2 * MAMBA2_STATE_DIM * 4) as u64; + +- // Zero gradients before accumulation +- self.stream.memset_zeros(&mut self.mamba2_grad) +- .map_err(|e| MLError::ModelError(format!("mamba2 grad zero: {e}")))?; ++ // Deterministic kernel: each thread writes exactly one element (plain write, ++ // not atomicAdd), so no memset_zeros needed. ++ let total_weight_params = (3 * sh2 * MAMBA2_STATE_DIM) as u32; + + unsafe { + self.stream.launch_builder(&self.mamba2_backward_kernel) +@@ -2495,7 +2847,8 @@ impl GpuDqnTrainer { + .arg(&(MAMBA2_HISTORY_K as i32)) + .arg(&(sh2 as i32)) + .arg(&(MAMBA2_STATE_DIM as i32)) +- .launch(LaunchConfig::for_num_elems(batch_size as u32)) ++ .arg(&self.isv_signals_dev_ptr) // regime-conditioned decay (ISV[11]) ++ .launch(LaunchConfig::for_num_elems(total_weight_params)) + .map_err(|e| MLError::ModelError(format!("mamba2_scan_backward: {e}")))?; + } + Ok(()) +@@ -3584,6 +3937,7 @@ impl GpuDqnTrainer { + ) -> Result { + let b = config.batch_size; + let total_params = compute_total_params(&config); ++ let non_isv_params = compute_non_isv_params(&config); + + // Event tracking: kept ENABLED during normal ops (buffer allocation, DtoD). + // DISABLED only during CUDA Graph capture (in capture_training_graphs). +@@ -3995,8 +4349,8 @@ impl GpuDqnTrainer { + .map_err(|e| MLError::ModelError(format!("strided_scatter load: {e}")))?; + let concat_ofi_kernel = exp_module_for_mag.load_function("concat_ofi_features") + .map_err(|e| MLError::ModelError(format!("concat_ofi_features load: {e}")))?; +- let regime_gate_kernel = exp_module_for_mag.load_function("regime_branch_gate") +- .map_err(|e| MLError::ModelError(format!("regime_branch_gate load: {e}")))?; ++ let branch_confidence_routing_kernel = exp_module_for_mag.load_function("branch_confidence_routing") ++ .map_err(|e| MLError::ModelError(format!("branch_confidence_routing load: {e}")))?; + let adaptive_atom_kernel = exp_module_for_mag.load_function("adaptive_atom_positions") + .map_err(|e| MLError::ModelError(format!("adaptive_atom_positions load: {e}")))?; + let atom_position_grad_kernel = exp_module_for_mag.load_function("atom_position_gradient") +@@ -4019,7 +4373,21 @@ impl GpuDqnTrainer { + .map_err(|e| MLError::ModelError(format!("apply_risk_budget load: {e}")))?; + let risk_backward_kernel = exp_module_for_mag.load_function("risk_budget_backward") + .map_err(|e| MLError::ModelError(format!("risk_budget_backward load: {e}")))?; +- info!("GpuDqnTrainer: mag_concat + strided_accumulate/scatter + concat_ofi + regime_gate + adaptive_atom + atom_grad + q_anchor + regime_dropout + G5/G6/G10/G12 + risk_budget kernels loaded"); ++ let isv_signal_update_kernel = exp_module_for_mag.load_function("isv_signal_update") ++ .map_err(|e| MLError::ModelError(format!("isv_signal_update load: {e}")))?; ++ let isv_forward_kernel = exp_module_for_mag.load_function("isv_forward") ++ .map_err(|e| MLError::ModelError(format!("isv_forward load: {e}")))?; ++ let isv_feature_gate_kernel = exp_module_for_mag.load_function("isv_feature_gate") ++ .map_err(|e| MLError::ModelError(format!("isv_feature_gate load: {e}")))?; ++ let fill_gamma_buf_kernel = exp_module_for_mag.load_function("fill_gamma_buf") ++ .map_err(|e| MLError::ModelError(format!("fill_gamma_buf load: {e}")))?; ++ let recursive_conf_fwd_kernel = exp_module_for_mag.load_function("recursive_confidence_forward") ++ .map_err(|e| MLError::ModelError(format!("recursive_confidence_forward load: {e}")))?; ++ let recursive_conf_bwd_kernel = exp_module_for_mag.load_function("recursive_confidence_backward") ++ .map_err(|e| MLError::ModelError(format!("recursive_confidence_backward load: {e}")))?; ++ let trade_plan_fwd_kernel = exp_module_for_mag.load_function("trade_plan_forward") ++ .map_err(|e| MLError::ModelError(format!("trade_plan_forward load: {e}")))?; ++ info!("GpuDqnTrainer: mag_concat + strided_accumulate/scatter + concat_ofi + regime_gate + adaptive_atom + atom_grad + q_anchor + regime_dropout + G5/G6/G10/G12 + risk_budget + isv_signal_update + isv_forward + fill_gamma_buf + trade_plan_forward kernels loaded"); + + // ── G5: Epistemic-gated magnitude — pinned var_ema threshold ─ + let (var_ema_pinned, var_ema_dev_ptr) = { +@@ -4680,6 +5048,143 @@ impl GpuDqnTrainer { + (host_ptr as *mut f32, dev_ptr_out) + }; + ++ // ── ISV scratch scalars (pinned device-mapped) ──────────────────── ++ let (td_error_scratch_pinned, td_error_scratch_dev_ptr) = { ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2( ++ &mut host_ptr, ++ std::mem::size_of::(), ++ ); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for td_error_scratch"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( ++ &mut dev_ptr_out, ++ host_ptr, ++ 0, ++ ); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for td_error_scratch"); ++ *(host_ptr as *mut f32) = 0.0; ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (ensemble_var_scratch_pinned, ensemble_var_scratch_dev_ptr) = { ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2( ++ &mut host_ptr, ++ std::mem::size_of::(), ++ ); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for ensemble_var_scratch"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( ++ &mut dev_ptr_out, ++ host_ptr, ++ 0, ++ ); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for ensemble_var_scratch"); ++ *(host_ptr as *mut f32) = 0.0; ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (reward_scratch_pinned, reward_scratch_dev_ptr) = { ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2( ++ &mut host_ptr, ++ std::mem::size_of::(), ++ ); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for reward_scratch"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( ++ &mut dev_ptr_out, ++ host_ptr, ++ 0, ++ ); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for reward_scratch"); ++ *(host_ptr as *mut f32) = 0.0; ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (atom_util_scratch_pinned, atom_util_scratch_dev_ptr) = { ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2( ++ &mut host_ptr, ++ std::mem::size_of::(), ++ ); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for atom_util_scratch"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2( ++ &mut dev_ptr_out, ++ host_ptr, ++ 0, ++ ); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for atom_util_scratch"); ++ *(host_ptr as *mut f32) = 0.0; ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ // ── ISV core buffers (pinned device-mapped) ────────────────────── ++ let (isv_signals_pinned, isv_signals_dev_ptr) = { ++ let num_bytes = ISV_DIM * std::mem::size_of::(); ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2(&mut host_ptr, num_bytes); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for isv_signals"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dev_ptr_out, host_ptr, 0); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for isv_signals"); ++ std::ptr::write_bytes(host_ptr as *mut u8, 0, num_bytes); ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (isv_history_pinned, isv_history_dev_ptr) = { ++ let num_bytes = ISV_K * ISV_DIM * std::mem::size_of::(); ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2(&mut host_ptr, num_bytes); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for isv_history"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dev_ptr_out, host_ptr, 0); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for isv_history"); ++ std::ptr::write_bytes(host_ptr as *mut u8, 0, num_bytes); ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (isv_decay_pinned, isv_decay_dev_ptr) = { ++ let num_bytes = ISV_DIM * std::mem::size_of::(); ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2(&mut host_ptr, num_bytes); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for isv_decay"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dev_ptr_out, host_ptr, 0); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for isv_decay"); ++ std::ptr::write_bytes(host_ptr as *mut u8, 0, num_bytes); ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ ++ let (lagged_td_error_pinned, lagged_td_error_dev_ptr) = { ++ let mut host_ptr: *mut std::ffi::c_void = std::ptr::null_mut(); ++ let mut dev_ptr_out: u64 = 0; ++ unsafe { ++ let rc = cudarc::driver::sys::cuMemAllocHost_v2(&mut host_ptr, std::mem::size_of::()); ++ assert_eq!(rc, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemAllocHost for lagged_td_error"); ++ let rc2 = cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dev_ptr_out, host_ptr, 0); ++ assert_eq!(rc2, cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS, "cuMemHostGetDevicePointer for lagged_td_error"); ++ *(host_ptr as *mut f32) = 0.0; ++ } ++ (host_ptr as *mut f32, dev_ptr_out) ++ }; ++ + // ── Cross-branch graph message passing ────────────────────────── + let cpbi_module_graph = stream.context().load_cubin(EXPECTED_Q_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("cpbi cubin (graph_msg): {e}")))?; +@@ -4874,6 +5379,8 @@ impl GpuDqnTrainer { + .map_err(|e| MLError::ModelError(format!("mamba2_temporal_scan load: {e}")))?; + let mamba2_update_kernel = mamba2_module.load_function("mamba2_update_history") + .map_err(|e| MLError::ModelError(format!("mamba2_update_history load: {e}")))?; ++ let isv_temporal_route_kernel = mamba2_module.load_function("isv_temporal_route") ++ .map_err(|e| MLError::ModelError(format!("isv_temporal_route load: {e}")))?; + let mamba2_bw_module = stream.context().load_cubin(EXPECTED_Q_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("mamba2 backward cubin load: {e}")))?; + let mamba2_backward_kernel = mamba2_bw_module.load_function("mamba2_scan_backward") +@@ -4971,10 +5478,31 @@ impl GpuDqnTrainer { + .map_err(|e| MLError::ModelError(format!("commit_lambda_buf alloc: {e}")))?; + let q_mag_pre_risk = stream.alloc_zeros::(b * config.branch_1_size) + .map_err(|e| MLError::ModelError(format!("q_mag_pre_risk alloc: {e}")))?; +- let risk_param_count = config.adv_h * config.shared_h2 + config.adv_h + config.adv_h + 1; ++ let risk_param_count = config.adv_h * (config.shared_h2 + 13) + config.adv_h + config.adv_h + 1; + let risk_grad_buf = stream.alloc_zeros::(risk_param_count) + .map_err(|e| MLError::ModelError(format!("risk_grad_buf alloc: {e}")))?; + ++ // ── ISV forward output buffers ── ++ let isv_embedding_buf = stream.alloc_zeros::(ISV_EMB_DIM) ++ .map_err(|e| MLError::ModelError(format!("isv_embedding_buf alloc: {e}")))?; ++ let branch_gate_buf = stream.alloc_zeros::(4) ++ .map_err(|e| MLError::ModelError(format!("branch_gate_buf alloc: {e}")))?; ++ let gamma_mod_buf = stream.alloc_zeros::(1) ++ .map_err(|e| MLError::ModelError(format!("gamma_mod_buf alloc: {e}")))?; ++ let temporal_weight_buf = stream.alloc_zeros::(sh2) ++ .map_err(|e| MLError::ModelError(format!("temporal_weight_buf alloc: {e}")))?; ++ let gamma_buf = stream.alloc_zeros::(b) ++ .map_err(|e| MLError::ModelError(format!("gamma_buf alloc: {e}")))?; ++ let predicted_error_buf = stream.alloc_zeros::(b) ++ .map_err(|e| MLError::ModelError(format!("predicted_error_buf alloc: {e}")))?; ++ let plan_params_buf = stream.alloc_zeros::(b * 6) ++ .map_err(|e| MLError::ModelError(format!("plan_params_buf alloc: {e}")))?; ++ ++ // ── Speculative inference cache ── ++ let speculative_h_s2 = stream.alloc_zeros::(b * sh2) ++ .map_err(|e| MLError::ModelError(format!("speculative_h_s2 alloc: {e}")))?; ++ let speculative_features = vec![0.0_f32; config.state_dim]; ++ + Ok(Self { + config, + stream, +@@ -5075,7 +5603,7 @@ impl GpuDqnTrainer { + glu_backward_kernel, + kan_gate_combine_kernel, + kan_gate_backward_kernel, +- regime_gate_kernel, ++ branch_confidence_routing_kernel, + regime_q_gap_buf, + regime_util_pinned, + regime_util_dev_ptr, +@@ -5145,6 +5673,7 @@ impl GpuDqnTrainer { + q_div_ema: 0.0, + adam_step: 0, + total_params, ++ non_isv_params, + params_initialized: false, + target_params_initialized: false, + attention_initialized: false, +@@ -5346,6 +5875,40 @@ impl GpuDqnTrainer { + q_mag_pre_risk, + risk_grad_buf, + risk_adam_step: 0, ++ td_error_scratch_pinned, ++ td_error_scratch_dev_ptr, ++ ensemble_var_scratch_pinned, ++ ensemble_var_scratch_dev_ptr, ++ reward_scratch_pinned, ++ reward_scratch_dev_ptr, ++ atom_util_scratch_pinned, ++ atom_util_scratch_dev_ptr, ++ isv_signals_pinned, ++ isv_signals_dev_ptr, ++ isv_history_pinned, ++ isv_history_dev_ptr, ++ isv_decay_pinned, ++ isv_decay_dev_ptr, ++ lagged_td_error_pinned, ++ lagged_td_error_dev_ptr, ++ isv_signal_update_kernel, ++ isv_forward_kernel, ++ isv_feature_gate_kernel, ++ fill_gamma_buf_kernel, ++ isv_embedding_buf, ++ branch_gate_buf, ++ gamma_mod_buf, ++ temporal_weight_buf, ++ isv_temporal_route_kernel, ++ gamma_buf, ++ predicted_error_buf, ++ recursive_conf_fwd_kernel, ++ recursive_conf_bwd_kernel, ++ plan_params_buf, ++ trade_plan_fwd_kernel, ++ speculative_h_s2, ++ speculative_features, ++ speculative_valid: false, + }) + } + +@@ -5402,6 +5965,64 @@ impl GpuDqnTrainer { + .map_err(|e| MLError::ModelError(format!("upload target params: {e}"))) + } + ++ /// Speculative forward: pre-compute trunk from intermediate features. ++ /// Called every ~5 seconds between bars with intermediate microstructure state. ++ /// Stores h_s2 in cache for potential reuse at bar close. ++ /// ++ /// Current implementation: cache infrastructure only (features stored, ISV updated). ++ /// Full cuBLAS trunk forward integration is deferred to a follow-up. ++ pub fn speculative_forward(&mut self, features: &[f32]) -> Result<(), MLError> { ++ if features.len() != self.config.state_dim { ++ return Err(MLError::ModelError(format!( ++ "speculative_forward: expected {} features, got {}", ++ self.config.state_dim, features.len() ++ ))); ++ } ++ ++ // Run ISV signal update (uses pinned scalars from last training step) ++ self.update_isv_signals()?; ++ ++ // Run ISV forward (encoder MLP → branch gate + gamma mod) ++ self.launch_isv_forward()?; ++ ++ // Cache the features for comparison at bar close ++ self.speculative_features.clear(); ++ self.speculative_features.extend_from_slice(features); ++ self.speculative_valid = true; ++ ++ Ok(()) ++ } ++ ++ /// Check if speculative cache is usable for the given bar-close features. ++ /// Returns true if features are within 5% relative L2 distance of the cached speculation. ++ pub fn check_speculative_cache(&self, actual_features: &[f32]) -> bool { ++ if !self.speculative_valid || actual_features.len() != self.speculative_features.len() { ++ return false; ++ } ++ ++ // Compute relative L2 distance between speculative and actual features ++ let mut sum_sq_diff = 0.0_f32; ++ let mut sum_sq_actual = 0.0_f32; ++ for (a, s) in actual_features.iter().zip(self.speculative_features.iter()) { ++ sum_sq_diff += (a - s) * (a - s); ++ sum_sq_actual += a * a; ++ } ++ ++ let relative_diff = if sum_sq_actual > 1e-8 { ++ (sum_sq_diff / sum_sq_actual).sqrt() ++ } else { ++ 1.0 // Can't compare if actual is all zeros ++ }; ++ ++ // Use cache if features changed by less than 5% ++ relative_diff < 0.05 ++ } ++ ++ /// Invalidate speculative cache (call at bar boundary or after training step). ++ pub fn invalidate_speculative_cache(&mut self) { ++ self.speculative_valid = false; ++ } ++ + /// Clone params_buf to a new CudaSlice via DtoD copy — checkpoint stays on GPU, zero CPU. + pub fn clone_params_gpu(&self) -> Result, MLError> { + let len = self.params_buf.len(); +@@ -5455,6 +6076,9 @@ impl GpuDqnTrainer { + /// Read eval_q_mean_ema. + pub fn eval_q_mean_ema(&self) -> f32 { self.eval_q_mean_ema } + ++ /// ISV signals device pointer for adaptive hold enforcement in experience collector. ++ pub fn isv_signals_dev_ptr(&self) -> u64 { self.isv_signals_dev_ptr } ++ + /// Read eval_q_std_ema. + pub fn eval_q_std_ema(&self) -> f32 { self.eval_q_std_ema } + +@@ -5979,6 +6603,7 @@ impl GpuDqnTrainer { + self.launch_mse_loss()?; + self.launch_loss_reduce(self.mse_loss_dev_ptr)?; + self.launch_mse_grad_to_scratch()?; ++ self.fill_gamma_buf()?; + self.launch_c51_loss()?; + self.launch_loss_reduce(self.total_loss_dev_ptr)?; + self.launch_c51_grad()?; +@@ -6583,6 +7208,22 @@ impl GpuDqnTrainer { + // G9: Regime-conditioned dropout on h_s2 (training path only) + self.apply_regime_dropout(batch_size, true)?; + ++ // ISV forward: encoder MLP → branch gate + gamma mod ++ self.launch_isv_forward()?; ++ ++ // ISV temporal routing: update per-feature temporal_weight for next mamba2 call ++ // One-step lag (mamba2 already ran above), same pattern as ISV signals. ++ self.launch_isv_temporal_route()?; ++ ++ // ISV feature gate: modulate h_s2 features based on ISV 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 forward: h_s2 → plan_params [B, 6] ++ self.launch_trade_plan_forward(batch_size)?; ++ + // Risk budget forward: h_s2 → risk_budget R ∈ (0,1) before Q-value computation + self.risk_budget_forward(batch_size)?; + +@@ -6622,8 +7263,8 @@ impl GpuDqnTrainer { + // Apply risk budget: scale magnitude Q-values, produce CVaR alpha + commitment lambda + self.apply_risk_budget(batch_size)?; + +- // Regime branch gate: scale Q-values by learned per-branch importance +- self.apply_regime_gate(batch_size)?; ++ // Branch confidence routing: ISV gate × Q-value separation confidence ++ self.apply_branch_confidence_routing(batch_size)?; + + // G5: Epistemic gate — high variance → conservative magnitude + self.apply_epistemic_gate(batch_size)?; +@@ -6726,6 +7367,14 @@ impl GpuDqnTrainer { + ); + } + ++ // ISV: write atom utilization to pinned scratch for GPU-side ISV signal update ++ unsafe { *self.atom_util_scratch_pinned = host[6].clamp(0.0, 1.0); } ++ // ISV: write Q-mean as reward proxy to pinned scratch ++ unsafe { *self.reward_scratch_pinned = host[3]; } ++ ++ // ISV signal update — pure GPU, reads all pinned scalars, updates ISV vector ++ self.update_isv_signals()?; ++ + // xLSTM temporal context: update matrix memory and write context[8] on GPU. + self.launch_qlstm_step()?; + +@@ -6894,14 +7543,11 @@ impl GpuDqnTrainer { + + /// Backward pass through the diffusion denoiser. + /// Computes MSE gradient of Q_refined vs Q_target and backprops through +- /// the 2 FC-SiLU-FC steps. Accumulates d_denoise_params via atomicAdd. ++ /// the 2 FC-SiLU-FC steps. Deterministic: one thread per weight, plain writes. + pub(crate) fn launch_q_denoise_backward(&mut self, batch_size: usize) -> Result<(), MLError> { +- // Zero gradient accumulator +- self.stream.memset_zeros(&mut self.denoise_grad) +- .map_err(|e| MLError::ModelError(format!("zero denoise_grad: {e}")))?; +- + let b = batch_size as i32; +- let blocks = ((batch_size as u32 + 255) / 256).max(1); ++ // Grid: one thread per weight element (1800 total = 2 steps × 900) ++ let blocks = ((1800_u32 + 255) / 256).max(1); + let q_refined_ptr = self.q_coord_buf.raw_ptr(); + let q_target_ptr = self.denoise_target_q_buf.raw_ptr(); + let params_ptr = self.denoise_params.raw_ptr(); +@@ -7221,6 +7867,7 @@ impl GpuDqnTrainer { + self.q_divergence_dev_ptr, 0, std::mem::size_of::(), self.stream.cu_stream(), + ); + } ++ self.fill_gamma_buf()?; + self.launch_c51_loss()?; + self.launch_loss_reduce(self.total_loss_dev_ptr)?; + self.launch_c51_grad()?; +@@ -7373,6 +8020,7 @@ impl GpuDqnTrainer { + self.q_divergence_dev_ptr, 0, std::mem::size_of::(), self.stream.cu_stream(), + ); + } ++ self.fill_gamma_buf()?; + self.launch_c51_loss()?; + self.launch_loss_reduce(self.total_loss_dev_ptr)?; + self.launch_c51_grad()?; +@@ -7845,9 +8493,9 @@ impl GpuDqnTrainer { + let on_next_b2_ptr = on_next_b1_ptr + (b * b1 * na * f32_sz) as u64; + let on_next_b3_ptr = on_next_b2_ptr + (b * b2 * na * f32_sz) as u64; + +- // N-step returns: use gamma^n for the Bellman projection. +- // The experience collector pre-computes R_n = sum(gamma^i * r_i). +- let gamma = self.adaptive_gamma.powi(self.config.n_steps as i32); ++ // N-step returns: gamma_buf [B] is pre-filled by fill_gamma_buf kernel ++ // (base_gamma^n * gamma_mod[0] per sample). ++ let gamma_buf_ptr = self.gamma_buf.raw_ptr(); + let batch_i32 = b as i32; + let na_i32 = na as i32; + let b0_i32 = b0 as i32; +@@ -7898,7 +8546,7 @@ impl GpuDqnTrainer { + .arg(&self.curiosity_error_buf) + .arg(&self.config.curiosity_q_penalty_lambda) + // ── Config (8 — per_sample_support replaces v_min+v_max) ── +- .arg(&gamma) ++ .arg(&gamma_buf_ptr) + .arg(&batch_i32) + .arg(&na_i32) + .arg(&self.per_sample_support_ptr) +@@ -8691,7 +9339,9 @@ impl GpuDqnTrainer { + /// matches the GOFF_* defines exactly, with `align4()` padding per tensor. + /// + /// Both params_buf and target_params_buf receive identical weights — the EMA +- /// kernel will diverge them during training. ++ /// kernel will diverge them during training. ISV weights (indices 68-79) in ++ /// target_params_buf are inert: EMA skips them, and the target forward never ++ /// reads them. + pub(crate) fn xavier_init_params_buf(&mut self) -> Result<(), MLError> { + let cfg = &self.config; + let sizes = compute_param_sizes(cfg); +@@ -8772,10 +9422,32 @@ impl GpuDqnTrainer { + (cfg.num_atoms, cfg.value_h), // [62] w_v2_20bar (Xavier) + (0, 0), // [63] b_v2_20bar + // ── Learned risk management (5th branch) ── +- (cfg.adv_h, cfg.shared_h2), // [64] w_risk_fc (Xavier) ++ (cfg.adv_h, cfg.shared_h2 + 13), // [64] w_risk_fc (Xavier, SH2+13 input) + (0, 0), // [65] b_risk_fc (zero) + (0, 0), // [66] w_risk_out (special init below) + (0, 0), // [67] b_risk_out (zero → sigmoid(0)=0.5) ++ // ── ISV encoder + conditioning ── ++ (16, ISV_DIM), // [68] w_isv_fc1 (Xavier, ISV_DIM input) ++ (0, 0), // [69] b_isv_fc1 (zero) ++ (ISV_EMB_DIM, 16), // [70] w_isv_fc2 (Xavier, EMB output) ++ (0, 0), // [71] b_isv_fc2 (zero) ++ (4, ISV_EMB_DIM), // [72] w_isv_gate (Xavier, EMB input) ++ (0, 0), // [73] b_isv_gate (zero) ++ (0, 0), // [74] w_isv_gamma (zero → sigmoid(0)=0.5) ++ (0, 0), // [75] b_isv_gamma (zero) ++ (1, cfg.shared_h2), // [76] w_conf_fc (Xavier) ++ (0, 0), // [77] b_conf_fc (zero) ++ // ── ISV Feature Gating (Phase 3) ── ++ (cfg.shared_h2, ISV_EMB_DIM), // [78] w_feature_gate (Xavier) ++ (0, 0), // [79] b_feature_gate (special init: 2.0) ++ // ── ISV Temporal Routing (Phase 3) ── ++ (cfg.shared_h2, ISV_EMB_DIM), // [80] w_temporal_route (Xavier) ++ (0, 0), // [81] b_temporal_route (special init: 2.0) ++ // ── Trade Plan Head ── ++ (cfg.adv_h, cfg.shared_h2), // [82] w_plan_fc (Xavier) ++ (0, 0), // [83] b_plan_fc (zero) ++ (6, cfg.adv_h), // [84] w_plan_out (Xavier) ++ (0, 0), // [85] b_plan_out (zero) + ]; + + // Build flat host buffer: Xavier init for weights, zeros for biases + padding. +@@ -8836,6 +9508,22 @@ impl GpuDqnTrainer { + } + } + ++ // Feature gate bias: init to 2.0 for near-pass-through (sigmoid(2)≈0.88) ++ { ++ let b_gate_offset: usize = sizes[..79].iter().map(|&s| align4(s)).sum(); ++ for j in 0..cfg.shared_h2 { ++ host_buf[b_gate_offset + j] = 2.0; ++ } ++ } ++ ++ // Temporal route bias: init to 2.0 (sigmoid(2)≈0.88 → mostly use temporal context) ++ { ++ let b_route_offset: usize = sizes[..81].iter().map(|&s| align4(s)).sum(); ++ for j in 0..cfg.shared_h2 { ++ host_buf[b_route_offset + j] = 2.0; ++ } ++ } ++ + debug_assert_eq!(offset, total, "xavier_init: offset mismatch with total_params"); + + // HtoD to both params_buf and target_params_buf (same initial weights). +@@ -8972,11 +9660,13 @@ impl GpuDqnTrainer { + + /// GPU-native Polyak EMA: `target[i] = (1-tau)*target[i] + tau*online[i]` + /// +- /// Fused single-kernel update over flat parameter buffers. On first call, +- /// flattens target weights into `target_params_buf` (same GOFF_* layout as +- /// `params_buf`). Then launches ONE EMA kernel over the entire flat buffer +- /// instead of 20 per-tensor launches. After the kernel, scatters the updated +- /// flat target weights back to individual tensors via `unflatten_target_weights()`. ++ /// Fused single-kernel update over flat parameter buffers, restricted to ++ /// non-ISV params (indices 0-67). ISV weights (68-77) are online-only and ++ /// never synced to the target network. On first call, flattens target weights ++ /// into `target_params_buf` (same GOFF_* layout as `params_buf`). Then ++ /// launches ONE EMA kernel over the non-ISV portion instead of 20 per-tensor ++ /// launches. After the kernel, scatters the updated flat target weights back ++ /// to individual tensors via `unflatten_target_weights()`. + /// + /// Runs OUTSIDE the captured CUDA Graph -- device pointers are stable so + /// the graph stays valid. +@@ -9003,8 +9693,11 @@ impl GpuDqnTrainer { + tau + }; + +- let n = self.total_params as i32; +- let blocks = ((self.total_params + 255) / 256) as u32; ++ // EMA only over non-ISV params (indices 0-67). ISV weights (68-77) are ++ // online-only — the target network uses neutral defaults (base_gamma, ++ // uniform gates) instead of ISV-modulated values. ++ let n = self.non_isv_params as i32; ++ let blocks = ((self.non_isv_params + 255) / 256) as u32; + let launch_cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), +@@ -9096,6 +9789,9 @@ impl GpuDqnTrainer { + /// Raw pointer to states_buf for pad_states kernel. + pub(crate) fn states_buf_ptr(&self) -> u64 { self.states_buf.raw_ptr() } + ++ /// Raw pointer to plan_params_buf [B, 6] for trade plan integration. ++ pub fn plan_params_buf_ptr(&self) -> u64 { self.plan_params_buf.raw_ptr() } ++ + /// F32 states buffer ref for cold-path Q-value computation. + pub(crate) fn states_buf_f32(&self) -> &CudaSlice { &self.states_buf } + +@@ -9138,16 +9834,14 @@ impl GpuDqnTrainer { + + /// Launch selectivity_backward kernel: BCE gradient on (sel, per_sample_loss / mean_loss). + /// +- /// Zeros `sel_grad` before launch. Reads h_s2, sel_out_buf, per_sample_loss_buf. +- /// Writes `sel_grad` [SH2+1] via atomicAdd. ++ /// Deterministic: one thread per weight element, loops over batch. ++ /// Writes `sel_grad` [SH2+1] — single write per element, no atomicAdd. + pub(crate) fn launch_selectivity_backward(&mut self, batch_size: usize, mean_loss: f32) -> Result<(), MLError> { + let b = batch_size as i32; + let sh2 = self.config.shared_h2 as i32; +- let blocks = ((batch_size as u32 + 255) / 256).max(1); +- +- // Zero gradient accumulator before atomicAdd writes. +- self.stream.memset_zeros(&mut self.sel_grad) +- .map_err(|e| MLError::ModelError(format!("zero sel_grad: {e}")))?; ++ // Grid: one thread per weight element (SH2 + 1 total) ++ let total_params = (self.config.shared_h2 + 1) as u32; ++ let blocks = ((total_params + 255) / 256).max(1); + + let h_s2_ptr = self.ptrs.save_h_s2; + let sel_out_ptr = self.sel_out_buf.raw_ptr(); +diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +index 199330a7d..99e85ad91 100644 +--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs ++++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +@@ -41,10 +41,10 @@ const MAX_EPISODES_LIMIT: usize = 0x8000; + /// Absolute upper bound for validation — reject configs above this. + const MAX_TIMESTEPS_LIMIT: usize = 1000; + const PORTFOLIO_STATE_SIZE: usize = 8; +-/// Portfolio stride for DQN experience kernels (23 floats per episode). ++/// Portfolio stride for DQN experience kernels (30 floats per episode). + /// Matches PORTFOLIO_STRIDE in experience_kernels.cu. + /// Do NOT change PORTFOLIO_STATE_SIZE above — it's for the PPO/legacy path. +-const PORTFOLIO_STRIDE: usize = 23; ++const PORTFOLIO_STRIDE: usize = 30; + + /// Number of floats in the trade_stats_reduce output buffer. + /// Layout: [win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns] +@@ -567,7 +567,7 @@ pub struct GpuExperienceCollector { + trade_stats_buf: CudaSlice, + + // Per-episode state buffers [alloc_episodes, ...] +- portfolio_states: CudaSlice, // [alloc_episodes * PORTFOLIO_STRIDE] (20 floats per episode) ++ portfolio_states: CudaSlice, // [alloc_episodes * PORTFOLIO_STRIDE] (30 floats per episode) + episode_starts_buf: CudaSlice,// [alloc_episodes] + + // Output buffers [alloc_episodes * alloc_timesteps, ...] +@@ -674,6 +674,14 @@ pub struct GpuExperienceCollector { + cost_anneal_pinned: *mut f32, + cost_anneal_dev_ptr: u64, + ++ /// ISV signals dev_ptr from trainer — for adaptive hold enforcement. ++ /// Set via set_isv_signals_ptr(). 0 = NULL (static hold). ++ isv_signals_dev_ptr: u64, ++ ++ /// Trade plan params dev_ptr from fused training context. ++ /// Set via set_plan_params_ptr(). 0 = NULL (no plan). ++ plan_params_dev_ptr: u64, ++ + /// v8: Per-bar curriculum difficulty scoring kernel. + difficulty_scores_kernel: CudaFunction, + /// v8: Hindsight experience relabeling kernel (optimal exit). +@@ -744,7 +752,7 @@ impl GpuExperienceCollector { + let (shared_h1, shared_h2, value_h, adv_h) = network_dims; + let (state_dim, market_dim, num_atoms_max) = kernel_dims; + let portfolio_dim: usize = 3; +- let ofi_dim: usize = if state_dim >= market_dim + portfolio_dim + 8 { 8 } else { 0 }; ++ let ofi_dim: usize = if state_dim >= market_dim + portfolio_dim + 20 { 20 } else { 0 }; + + // Branch sizes for 4-branch hierarchical DQN (always enabled). + // [direction(3), magnitude(3), order(3), urgency(3)] +@@ -857,7 +865,7 @@ impl GpuExperienceCollector { + } + + // ── Step 6: Allocate per-episode buffers ──────────────────────── +- // Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=20] ++ // Portfolio states for experience kernels: [N, PORTFOLIO_STRIDE=30] + let mut portfolio_states = stream + .alloc_zeros::(alloc_episodes * PORTFOLIO_STRIDE) + .map_err(|e| MLError::ModelError(format!("alloc portfolio_states: {e}")))?; +@@ -1198,6 +1206,8 @@ impl GpuExperienceCollector { + pinned_dones, + cost_anneal_pinned, + cost_anneal_dev_ptr, ++ isv_signals_dev_ptr: 0, // NULL until trainer sets it ++ plan_params_dev_ptr: 0, // NULL until trainer sets it + difficulty_scores_kernel, + hindsight_relabel_kernel, + td_lambda_kernel, +@@ -1215,6 +1225,20 @@ impl GpuExperienceCollector { + unsafe { *self.cost_anneal_pinned = factor; } + } + ++ /// Set ISV signals device pointer for adaptive hold enforcement. ++ /// The experience collector's env_step kernel reads ISV signals to compute ++ /// dynamic hold time: base + f(stability, confidence). Pass 0 for static hold. ++ pub fn set_isv_signals_ptr(&mut self, dev_ptr: u64) { ++ self.isv_signals_dev_ptr = dev_ptr; ++ } ++ ++ /// Set trade plan params device pointer for plan activation/enforcement. ++ /// The env_step kernel reads plan params to copy into portfolio state on ++ /// Flat→Positioned transitions. Pass 0 to disable (NULL = no plan). ++ pub fn set_plan_params_ptr(&mut self, dev_ptr: u64) { ++ self.plan_params_dev_ptr = dev_ptr; ++ } ++ + /// Set CVaR position scaling from IQN dual-head. + /// The device pointer will be passed to the env_step kernel. + /// Call with 0 to disable (NULL pointer = no scaling). +@@ -2205,6 +2229,7 @@ impl GpuExperienceCollector { + .arg(&config.eps_urg_mult) // per-branch epsilon multiplier: urgency + .arg(&(t as i32)) // timestep for stateless Philox RNG + .arg(&self.per_sample_epsilon_ptr) // IQL expectile gap epsilon (0=cosine schedule) ++ .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0=NULL=static) + .launch(launch_cfg) + .map_err(|e| MLError::ModelError(format!( + "experience_action_select t={t}: {e}" +@@ -2309,6 +2334,8 @@ impl GpuExperienceCollector { + .arg(&total_actions_env) // total branch actions for variance indexing + .arg(&self.cost_anneal_dev_ptr) // G2: transaction cost curriculum + .arg(&0u64) // commit_lambda_buf: NULL = use 0.01 fallback (risk branch not in experience path yet) ++ .arg(&self.isv_signals_dev_ptr) // ISV signals for adaptive hold (0 = NULL = static hold) ++ .arg(&self.plan_params_dev_ptr) // trade plan params (0 = NULL = no plan) + .launch(launch_cfg) + .map_err(|e| MLError::ModelError(format!( + "experience_env_step t={t}: {e}" +@@ -2414,7 +2441,7 @@ impl GpuExperienceCollector { + _avg_spread: f32, + _cash_reserve_pct: f32, + ) -> Result<(), MLError> { +- // Reset portfolio states [N, PORTFOLIO_STRIDE=20] ++ // Reset portfolio states [N, PORTFOLIO_STRIDE=30] + let mut portfolio_init = vec![0.0_f32; self.alloc_episodes * PORTFOLIO_STRIDE]; + for i in 0..self.alloc_episodes { + let off = i * PORTFOLIO_STRIDE; +diff --git a/crates/ml/src/cuda_pipeline/gpu_walk_forward.rs b/crates/ml/src/cuda_pipeline/gpu_walk_forward.rs +index ef52ee05f..02f936f04 100644 +--- a/crates/ml/src/cuda_pipeline/gpu_walk_forward.rs ++++ b/crates/ml/src/cuda_pipeline/gpu_walk_forward.rs +@@ -547,11 +547,11 @@ impl GpuWalkForwardData { + /// On H100 (80 GB), a typical dataset of 1M bars uses: + /// features: 1M × 42 × 4 = 168 MB + /// targets: 1M × 4 × 4 = 16 MB +- /// ofi: 1M × 8 × 4 = 32 MB ++ /// ofi: 1M × 20 × 4 = 80 MB + /// total: 216 MB (~0.27% of VRAM) + pub fn upload( + training_data: &[(FeatureVector, Vec)], +- ofi_data: Option<&[[f64; 8]]>, ++ ofi_data: Option<&[[f64; 20]]>, + wf_config: &GpuWalkForwardConfig, + stream: &Arc, + ) -> Result { +@@ -592,18 +592,18 @@ impl GpuWalkForwardData { + // Upload OFI (optional) + let ofi_gpu = if let Some(ofi) = ofi_data { + let n = ofi.len().min(total_bars); +- let mut flat_ofi = Vec::with_capacity(total_bars * 8); ++ let mut flat_ofi = Vec::with_capacity(total_bars * 20); + for i in 0..total_bars { + if i < n { + for &v in ofi[i].iter() { + flat_ofi.push(v as f32); + } + } else { +- flat_ofi.extend_from_slice(&[0.0_f32; 8]); ++ flat_ofi.extend_from_slice(&[0.0_f32; 20]); + } + } + let buf = super::clone_htod_f32(stream, &flat_ofi)?; +- vram += total_bars * 8 * 4; ++ vram += total_bars * 20 * 4; + Some(buf) + } else { + None +diff --git a/crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu b/crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu +index f95bebb44..cd4edbc3e 100644 +--- a/crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu +@@ -20,7 +20,9 @@ extern "C" __global__ void mamba2_temporal_scan( + int N, + int K, + int sh2, +- int state_d ++ int state_d, ++ const float* __restrict__ isv_signals, /* [12] pinned. NULL = no regime modulation. */ ++ const float* __restrict__ temporal_weight /* [SH2] per-feature routing. NULL = full residual. */ + ) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; +@@ -44,6 +46,12 @@ extern "C" __global__ void mamba2_temporal_scan( + b_val += w_b[(long long)j * state_d + s] * h_t[j]; + } + float gate = 1.0f / (1.0f + expf(-a_val)); ++ /* Regime-conditioned decay: during transitions (low stability), ++ * decay faster -> forget old-regime patterns. Stable regime -> full history. */ ++ if (isv_signals != NULL) { ++ float stability = isv_signals[11]; /* regime_stability in [0, 1] */ ++ gate *= stability; /* gate in [0, stability] instead of [0, 1] */ ++ } + + x[s] = gate * x[s] + b_val; + } +@@ -56,7 +64,8 @@ extern "C" __global__ void mamba2_temporal_scan( + for (int s = 0; s < state_d; s++) { + ctx += w_c[(long long)j * state_d + s] * x[s]; + } +- out[j] = h_current[j] + ctx; ++ float tw = (temporal_weight != NULL) ? temporal_weight[j] : 1.0f; ++ out[j] = h_current[j] + tw * ctx; + } + } + +@@ -86,3 +95,30 @@ extern "C" __global__ void mamba2_update_history( + + base[(long long)(K - 1) * sh2 + j] = h_s2[(long long)i * sh2 + j]; + } ++ ++/* ================================================================== */ ++/* Kernel: isv_temporal_route — ISV controls per-feature temporal depth*/ ++/* ================================================================== */ ++ ++/** ++ * Computes temporal_weight [SH2] from ISV embedding. ++ * Each feature gets its own "how much history to use" weight. ++ * temporal_weight[j] = sigmoid(w_temporal @ isv_emb + b_temporal)[j] ++ * ++ * Launch: grid=(1,1,1), block=(1,1,1). Single thread (SH2=256 iterations). ++ */ ++extern "C" __global__ void isv_temporal_route( ++ const float* __restrict__ isv_embedding, /* [8] from isv_forward */ ++ const float* __restrict__ w_temporal_route, /* [SH2, 8] */ ++ const float* __restrict__ b_temporal_route, /* [SH2] */ ++ float* __restrict__ temporal_weight, /* [SH2] output */ ++ int SH2 ++) { ++ if (threadIdx.x != 0) return; ++ for (int j = 0; j < SH2; j++) { ++ float val = b_temporal_route[j]; ++ for (int k = 0; k < 8; k++) ++ val += w_temporal_route[(long long)j * 8 + k] * isv_embedding[k]; ++ temporal_weight[j] = 1.0f / (1.0f + expf(-val)); ++ } ++} +diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs +index a3da03a06..b8fe755f4 100644 +--- a/crates/ml/src/cuda_pipeline/mod.rs ++++ b/crates/ml/src/cuda_pipeline/mod.rs +@@ -298,7 +298,7 @@ impl DqnGpuData { + pub fn upload_slices( + features: &[[f64; 42]], + targets: &[[f64; 4]], +- ofi: &[[f64; 8]], ++ ofi: &[[f64; 20]], + stream: &Arc, + ) -> Result { + let num_bars = features.len(); +@@ -332,14 +332,14 @@ impl DqnGpuData { + let targets_gpu = clone_htod_f32(stream, &flat_targets)?; + + let ofi_features = if !ofi.is_empty() && ofi.iter().any(|row| row.iter().any(|&v| v != 0.0)) { +- let mut flat_ofi = Vec::with_capacity(num_bars * 8); ++ let mut flat_ofi = Vec::with_capacity(num_bars * 20); + for i in 0..num_bars { + if i < ofi.len() { + for &v in ofi[i].iter() { + flat_ofi.push(v as f32); + } + } else { +- flat_ofi.extend_from_slice(&[0.0_f32; 8]); ++ flat_ofi.extend_from_slice(&[0.0_f32; 20]); + } + } + Some(clone_htod_f32(stream, &flat_ofi)?) +@@ -357,13 +357,13 @@ impl DqnGpuData { + }) + } + +- /// Upload OFI features (8 dims per bar) from MBP-10 order book data. ++ /// Upload OFI features (20 dims per bar) from MBP-10 order book data. + /// + /// Must be called after `upload()` with a slice matching `num_bars` length. + /// Features are: OFI, VPIN, Kyle's Lambda, trade imbalance, etc. + pub fn upload_ofi( + &mut self, +- ofi_data: &[[f64; 8]], ++ ofi_data: &[[f64; 20]], + stream: &Arc, + ) -> Result<(), MLError> { + let n = ofi_data.len().min(self.num_bars); +@@ -371,7 +371,7 @@ impl DqnGpuData { + return Ok(()); + } + +- let mut flat = Vec::with_capacity(self.num_bars * 8); ++ let mut flat = Vec::with_capacity(self.num_bars * 20); + for i in 0..self.num_bars { + if i < n { + for &v in ofi_data[i].iter() { +@@ -379,7 +379,7 @@ impl DqnGpuData { + } + } else { + // Zero-pad if OFI data is shorter than market data +- flat.extend_from_slice(&[0.0_f32; 8]); ++ flat.extend_from_slice(&[0.0_f32; 20]); + } + } + +@@ -399,7 +399,7 @@ impl DqnGpuData { + + /// Estimated VRAM usage in bytes. + pub fn vram_bytes(&self) -> usize { +- let ofi_dim = if self.ofi_features.is_some() { 8 } else { 0 }; ++ let ofi_dim = if self.ofi_features.is_some() { 20 } else { 0 }; + estimate_vram_bytes(self.num_bars * (self.feature_dim + 4 + ofi_dim)) + } + +@@ -474,7 +474,7 @@ impl DqnGpuData { + portfolio_features: &[f32; 3], + stream: &Arc, + ) -> Result, MLError> { +- let ofi_dim = if self.ofi_features.is_some() { 8 } else { 0 }; ++ let ofi_dim = if self.ofi_features.is_some() { 20 } else { 0 }; + let raw_dim = self.feature_dim + 3 + ofi_dim; + let final_dim = self.aligned_state_dim.unwrap_or(raw_dim); + +@@ -489,10 +489,10 @@ impl DqnGpuData { + // HtoD copy: portfolio features -> dst[feature_dim..feature_dim+3] (3 scalars, tiny) + Self::htod_copy_into(portfolio_features, &mut dst, self.feature_dim, stream)?; + +- // DtoD copy: OFI features -> dst[feature_dim+3..feature_dim+11] ++ // DtoD copy: OFI features -> dst[feature_dim+3..feature_dim+3+20] + if let Some(ref ofi) = self.ofi_features { +- let ofi_gpu = Self::d2d_subrange(ofi, bar_idx * 8, 8, stream)?; +- Self::dtod_copy_into(&ofi_gpu, &mut dst, self.feature_dim + 3, 8, stream)?; ++ let ofi_gpu = Self::d2d_subrange(ofi, bar_idx * 20, 20, stream)?; ++ Self::dtod_copy_into(&ofi_gpu, &mut dst, self.feature_dim + 3, 20, stream)?; + } + + Ok(dst) +@@ -513,7 +513,7 @@ impl DqnGpuData { + if count == 0 { + return Err(MLError::ModelError("Empty batch for state construction".to_owned())); + } +- let ofi_dim = if self.ofi_features.is_some() { 8 } else { 0 }; ++ let ofi_dim = if self.ofi_features.is_some() { 20 } else { 0 }; + let raw_dim = self.feature_dim + 3 + ofi_dim; + let final_dim = self.aligned_state_dim.unwrap_or(raw_dim); + +diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh +index e1fc406a1..b8fbda7a3 100644 +--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh ++++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh +@@ -140,14 +140,31 @@ __device__ __forceinline__ float enforce_hold( + float target_exposure, + float hold_time, + int min_hold_bars, +- int is_last_bar /* bypass hold on last bar of episode/window */ ++ int is_last_bar, /* bypass hold on last bar of episode/window */ ++ const float* __restrict__ isv_signals /* [8] or NULL — for adaptive hold */ + ) { + int prev_sign = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0); + int curr_sign = (target_exposure > 0.001f) ? 1 : ((target_exposure < -0.001f) ? -1 : 0); + int wants_exit = (prev_sign != 0 && curr_sign == 0); + int wants_reversal = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign); + +- if (hold_time > 0.0f && hold_time < (float)min_hold_bars ++ /* Adaptive hold: base + ISV-driven extension. ++ * ISV[1] = grad_norm_ema (high = unstable → hold longer, don't churn) ++ * ISV[3] = ensemble_var_ema (high = uncertain → hold longer, wait for clarity) ++ * ISV[5] = reward_ema (negative = losing → allow earlier exit) ++ * When ISV is NULL, use static min_hold_bars. */ ++ int adaptive_hold = min_hold_bars; ++ if (isv_signals != NULL) { ++ float stability = 1.0f / (1.0f + isv_signals[1]); /* inv grad_norm: high norm → low stability */ ++ float confidence = 1.0f - fminf(isv_signals[3], 1.0f); /* inv ensemble_var: high var → low confidence */ ++ /* Extension: 0-8 bars based on stability × confidence */ ++ int extension = (int)(stability * confidence * 8.0f); ++ /* If losing (negative reward_ema), reduce hold to base only */ ++ if (isv_signals[5] < -0.001f) extension = 0; ++ adaptive_hold = min_hold_bars + extension; ++ } ++ ++ if (hold_time > 0.0f && hold_time < (float)adaptive_hold + && (wants_exit || wants_reversal) + && !is_last_bar) { + return prev_position; /* Override: keep position */ +diff --git a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu +index e18f32084..0ec96b298 100644 +--- a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu ++++ b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu +@@ -1,7 +1,7 @@ + /** + * Trade stats reduction kernel. + * +- * Reads portfolio_states[N, PORTFOLIO_STRIDE=23] and sums fields [14:19] ++ * Reads portfolio_states[N, PORTFOLIO_STRIDE=30] and sums fields [14:19] + * (win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns) + * across all N episodes into a 6-float output buffer. + * Uses shared-memory warp reduction (same pattern as monitoring_reduce). +@@ -9,7 +9,7 @@ + * Launch config: grid=(1, 1, 1), block=(256, 1, 1). + */ + +-#define PORTFOLIO_STRIDE 23 ++#define PORTFOLIO_STRIDE 30 + + extern "C" __global__ void trade_stats_reduce( + const float* __restrict__ portfolio_states, // [N * PORTFOLIO_STRIDE] +diff --git a/crates/ml/src/fxcache.rs b/crates/ml/src/fxcache.rs +index 56ad0e52f..aac26ed28 100644 +--- a/crates/ml/src/fxcache.rs ++++ b/crates/ml/src/fxcache.rs +@@ -13,14 +13,14 @@ + //! │ version u16 = 1 (f32) │ + //! │ feat_dim u16 = 42 │ + //! │ target_dim u16 = 4 │ +-//! │ ofi_dim u16 = 8 │ ++//! │ ofi_dim u16 = 20 │ + //! │ bar_count u64 │ + //! │ cache_key [u8; 32] (SHA256 raw bytes) │ + //! │ reserved [u8; 8] │ + //! └───────────────────────────────────────────────────┘ + //! │ Body (bar_count records) │ + //! │ Each record starts with an i64 timestamp (ns). │ +-//! │ Version 1: [i64 ts][54 × f32] = 224 bytes/bar │ ++//! │ Version 1: [i64 ts][66 × f32] = 272 bytes/bar │ + //! └───────────────────────────────────────────────────┘ + //! ``` + +@@ -37,16 +37,21 @@ const FXCACHE_MAGIC: [u8; 8] = *b"FXCACHE\0"; + /// Header size in bytes (fixed). + const HEADER_SIZE: usize = 64; + ++/// Cache format version. Bump on ANY format change (OFI_DIM, FEAT_DIM, etc.). ++/// Stale cache files with wrong version are auto-detected and rejected by validate(). ++/// The ensure-fxcache Argo step catches the error and regenerates. ++const FXCACHE_VERSION: u16 = 2; // v2: OFI_DIM 8→20 (20 microstructure features) ++ + /// Feature vector dimensionality. + const FEAT_DIM: usize = 42; + + /// Target vector dimensionality (close, next_close, raw_close, raw_next). + const TARGET_DIM: usize = 4; + +-/// OFI vector dimensionality (8 MBP-10 order-flow imbalance levels). +-const OFI_DIM: usize = 8; ++/// OFI vector dimensionality (8 base + 12 extended microstructure features). ++pub const OFI_DIM: usize = 20; + +-/// Total f64 values per record: features + targets + OFI = 42 + 4 + 8 = 54. ++/// Total f64 values per record: features + targets + OFI = 42 + 4 + 20 = 66. + const RECORD_F64_COUNT: usize = FEAT_DIM + TARGET_DIM + OFI_DIM; + + /// Total f32 values per record (same count as f64, no padding needed). +@@ -65,7 +70,7 @@ pub struct FxCacheHeader { + pub feat_dim: u16, + /// Target dimension (4). + pub target_dim: u16, +- /// OFI dimension (8). ++ /// OFI dimension (20). + pub ofi_dim: u16, + /// Number of bars (records) in the file. + pub bar_count: u64, +@@ -101,10 +106,10 @@ impl FxCacheHeader { + self.magic + ); + } +- if self.version != 1 { ++ if self.version != FXCACHE_VERSION { + bail!( +- "Unsupported FxCache version: {} (expected 1)", +- self.version ++ "Stale FxCache version: {} (expected {}). Delete and regenerate.", ++ self.version, FXCACHE_VERSION + ); + } + if self.feat_dim as usize != FEAT_DIM { +@@ -191,7 +196,7 @@ pub struct FxCacheData { + pub features: Vec<[f64; FEAT_DIM]>, + /// Target vectors, one per bar (4 elements each). + pub targets: Vec<[f64; TARGET_DIM]>, +- /// OFI vectors, one per bar (8 elements each). ++ /// OFI vectors, one per bar (20 elements each). + pub ofi: Vec<[f64; OFI_DIM]>, + /// SHA256 cache key (raw 32 bytes). + pub cache_key: [u8; 32], +@@ -211,7 +216,7 @@ pub struct FxCacheData { + /// * `path` — Output file path (parent directories are created automatically) + /// * `features` — Slice of 42-element feature vectors + /// * `targets` — Slice of 4-element target vectors +-/// * `ofi` — Slice of 8-element OFI vectors ++/// * `ofi` — Slice of 20-element OFI vectors + /// * `timestamps` — Per-bar timestamps (nanoseconds since Unix epoch) + /// * `cache_key` — SHA256 key (raw 32 bytes) + /// * `has_ofi` — If true, OFI was computed from real MBP-10 data; false means zero-filled +@@ -248,7 +253,7 @@ pub fn write_fxcache( + .with_context(|| format!("Failed to create parent dirs for {:?}", path))?; + } + +- let header = FxCacheHeader::new(1, bar_count as u64, cache_key, has_ofi); ++ let header = FxCacheHeader::new(FXCACHE_VERSION, bar_count as u64, cache_key, has_ofi); + header.validate()?; + + let file = std::fs::File::create(path) +@@ -277,7 +282,7 @@ pub fn write_fxcache( + Ok(total_bytes) + } + +-/// Write body in f32 format (version 1): [i64 ts] + 54 × f32 = 224 bytes per bar. ++/// Write body in f32 format (version 1): [i64 ts] + 66 × f32 = 272 bytes per bar. + fn write_body_f32( + writer: &mut BufWriter, + features: &[[f64; FEAT_DIM]], +@@ -337,7 +342,7 @@ pub fn load_fxcache(path: &Path) -> Result { + + let bar_count = header.bar_count as usize; + +- // Sanity-check file size: [i64 ts] + 54 × f32 = 224 bytes/bar ++ // Sanity-check file size: [i64 ts] + 66 × f32 = 272 bytes/bar + let expected_body = bar_count as u64 * (8 + RECORD_F32_COUNT as u64 * 4); + let expected_total = HEADER_SIZE as u64 + expected_body; + if file_len < expected_total { +@@ -423,7 +428,7 @@ mod has_ofi_tests { + let path = dir.path().join("test_ofi_true.fxcache"); + let features = vec![[1.0_f64; 42]; 10]; + let targets = vec![[0.0_f64; 4]; 10]; +- let ofi = vec![[0.5_f64; 8]; 10]; ++ let ofi = vec![[0.5_f64; OFI_DIM]; 10]; + let timestamps = vec![1_i64; 10]; + let key = [0u8; 32]; + +@@ -439,7 +444,7 @@ mod has_ofi_tests { + let path = dir.path().join("test_ofi_false.fxcache"); + let features = vec![[1.0_f64; 42]; 10]; + let targets = vec![[0.0_f64; 4]; 10]; +- let ofi = vec![[0.0_f64; 8]; 10]; ++ let ofi = vec![[0.0_f64; OFI_DIM]; 10]; + let timestamps = vec![1_i64; 10]; + let key = [0u8; 32]; + +@@ -590,7 +595,7 @@ mod discover_tests { + ); + let features = vec![[1.0_f64; 42]; 5]; + let targets = vec![[0.0_f64; 4]; 5]; +- let ofi = vec![[0.0_f64; 8]; 5]; ++ let ofi = vec![[0.0_f64; OFI_DIM]; 5]; + let timestamps = vec![1_i64; 5]; + write_fxcache(&path, &features, &targets, &ofi, ×tamps, [0u8; 32], false) + .unwrap(); +diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs +index abd88b02e..588b9fea9 100644 +--- a/crates/ml/src/hyperopt/adapters/dqn.rs ++++ b/crates/ml/src/hyperopt/adapters/dqn.rs +@@ -454,9 +454,9 @@ pub struct DQNMetrics { + /// ## Fixed Architecture + /// + /// The following parameters are fixed for consistency: +-/// - `state_dim`: 50 without OFI (42 market + 8 portfolio), or 58 with MBP-10 data (+8 OFI) ++/// - `state_dim`: 66 without OFI (42 market + 8 portfolio + 16 MTF), or 86 with MBP-10 data (+20 OFI) + /// - Market features (42): OHLCV, technical indicators, price patterns, volume, time, statistical, ADX, CUSUM direction +-/// - OFI features (8, optional): TRUE OFI from MBP-10 order book (present when MBP-10 data available) ++/// - OFI features (20, optional): 8 base OFI from MBP-10 order book + 12 extended microstructure + /// - Portfolio features (3): Position, unrealized PnL, drawdown + /// - `num_actions`: 9 (9 exposure levels: S100/S75/S50/S25/Flat/L25/L50/L75/L100) + /// - `hidden_dims`: [256, 128, 64] +@@ -535,7 +535,7 @@ pub struct DQNTrainer { + trades_data_dir: Option, + /// Preloaded OFI features from MBP-10 data (8 features per bar). + /// Loaded during `preload_data()` and injected into each trial's DQNTrainer. +- preloaded_ofi_features: Option>, ++ preloaded_ofi_features: Option>, + /// Trading instrument symbol (e.g. "ES.FUT") — used in cache key to prevent + /// cross-instrument collisions when multiple symbols share the same data directory. + symbol: String, +@@ -917,10 +917,10 @@ impl DQNTrainer { + targets.push(t); + } + let has_ofi = ofi_features.is_some(); +- let ofi: Vec<[f64; 8]> = if let Some(ref ofi_arc) = ofi_features { ++ let ofi: Vec<[f64; 20]> = if let Some(ref ofi_arc) = ofi_features { + ofi_arc.iter().cloned().collect() + } else { +- vec![[0.0_f64; 8]; total] ++ vec![[0.0_f64; 20]; total] + }; + let timestamps = vec![0_i64; total]; + +@@ -1205,7 +1205,7 @@ impl DQNTrainer { + + /// Attempt to load OFI features from MBP10 data. + /// Returns None if MBP10 data is not available or on error. +- fn load_ofi_features(&self) -> Option> { ++ fn load_ofi_features(&self) -> Option> { + use crate::features::mbp10_loader::load_ofi_features_parallel; + + let mbp10_dir = if let Some(ref dir) = self.mbp10_data_dir { +@@ -1227,7 +1227,14 @@ impl DQNTrainer { + + load_ofi_features_parallel(&mbp10_dir, trades_dir.as_deref(), |dir| { + collect_dbn_files(dir) +- }).map(Arc::from) ++ }).map(|ofi8| { ++ let ofi20: Vec<[f64; 20]> = ofi8.into_iter().map(|row| { ++ let mut wide = [0.0_f64; 20]; ++ wide[..8].copy_from_slice(&row); ++ wide ++ }).collect(); ++ Arc::from(ofi20) ++ }) + } + + /// Extract 53-feature vectors from OHLCV bars and create training data +@@ -1339,13 +1346,13 @@ impl DQNTrainer { + // The gather kernel places live portfolio at [feat_dim..feat_dim+8] and + // multi-timeframe features at [feat_dim+8..feat_dim+24]. + // GpuBacktestEvaluator adds PORTFOLIO_AND_MTF_DIM=24 on top of feature_dim +- // to reach the full state_dim (72 without OFI, 80 with OFI, both 8-aligned). ++ // to reach the full state_dim (72 without OFI, 88 with OFI, both 8-aligned). + // +- // For OFI-enabled models, feature_dim = 50 (42 market + 8 OFI placeholder) +- // so state_dim = (50+24+7)&!7 = 80. Without OFI: (42+24+7)&!7 = 72. ++ // For OFI-enabled models, feature_dim = 62 (42 market + 20 OFI) ++ // so state_dim = (62+24+7)&!7 = 88. Without OFI: (42+24+7)&!7 = 72. + let ofi_enabled = self.mbp10_data_dir.is_some(); + let market_dim: usize = 42; +- let feature_dim: usize = if ofi_enabled { 50 } else { market_dim }; ++ let feature_dim: usize = if ofi_enabled { 62 } else { market_dim }; + + // OFI feature overlay: preloaded OFI covers ALL bars (train+val), left-aligned. + // Validation bars start at index `train_end` in the global OFI array. +@@ -1434,7 +1441,7 @@ impl DQNTrainer { + max_leverage: internal_trainer.hyperparams().max_leverage as f32, + // OFI reorder in gather kernel: produces [market, portfolio, OFI, pad] + // directly, eliminating the Candle narrow+cat closure. +- ofi_dim: if ofi_enabled { 8 } else { 0 }, ++ ofi_dim: if ofi_enabled { 20 } else { 0 }, + min_hold_bars: internal_trainer.hyperparams().min_hold_bars as i32, + bars_per_day: internal_trainer.hyperparams().bars_per_day as f32, + trading_days_per_year: internal_trainer.hyperparams().trading_days_per_year as f32, +diff --git a/crates/ml/src/hyperopt/adapters/ppo.rs b/crates/ml/src/hyperopt/adapters/ppo.rs +index f2ee2e8ea..d6961ee77 100644 +--- a/crates/ml/src/hyperopt/adapters/ppo.rs ++++ b/crates/ml/src/hyperopt/adapters/ppo.rs +@@ -1495,7 +1495,7 @@ impl PPOTrainer { + + /// Attempt to load OFI features from MBP10 data. + /// Returns None if MBP10 data is not available or on error. +- fn load_ofi_features(&self) -> Option> { ++ fn load_ofi_features(&self) -> Option> { + use crate::features::mbp10_loader::load_ofi_features_parallel; + + let mbp10_dir = self +@@ -1508,6 +1508,12 @@ impl PPOTrainer { + + load_ofi_features_parallel(&mbp10_dir, trades_dir.as_deref(), |dir| { + collect_dbn_files_recursive(dir) ++ }).map(|ofi8| { ++ ofi8.into_iter().map(|row| { ++ let mut wide = [0.0_f64; 20]; ++ wide[..8].copy_from_slice(&row); ++ wide ++ }).collect() + }) + } + +diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs +index 038a5d84c..5740766ff 100644 +--- a/crates/ml/src/trainers/dqn/config.rs ++++ b/crates/ml/src/trainers/dqn/config.rs +@@ -963,12 +963,14 @@ pub struct DQNHyperparameters { + /// DSR EMA decay rate (0.001=slow, 0.05=fast). Tunable in hyperopt. + pub dsr_eta: f64, + +- // OFI (Order Flow Imbalance) features from MBP-10 data ++ // OFI (Order Flow Imbalance) features from MBP-10 data — ALWAYS LOADED. + /// Directory containing MBP-10 .dbn.zst files for OFI feature extraction. +- /// 8 OFI features are always computed and appended to the state. ++ /// 20 microstructure features (8 OFI + 12 tick-level) always appended to state. ++ /// Default: "test_data/futures-baseline-mbp10" (local), "/data/futures-baseline-mbp10" (H100). + pub mbp10_data_dir: String, + /// Directory containing trade .dbn.zst files for real VPIN/Kyle's Lambda. +- /// Trades are always loaded and fed to `OFICalculator::feed_trade()`. ++ /// Trades always loaded and fed to `OFICalculator::feed_trade()`. ++ /// Default: "test_data/futures-baseline-trades" (local), "/data/futures-baseline-trades" (H100). + pub trades_data_dir: String, + + /// Data source for training bars: "ohlcv" (1-minute candles) or "mbp10" +@@ -1341,7 +1343,7 @@ impl DQNHyperparameters { + // P2-A Enhancement + initial_capital: 100_000.0, // $100K default + bars_per_day: 390.0, // 1-minute bars (6.5h × 60min) +- min_hold_bars: 1, // was 5 — model learns hold timing via urgency branch ++ min_hold_bars: 5, // 5-bar minimum prevents coin-flip exits (1-2 bar churn) + + // P2-B Enhancement + cash_reserve_percent: 0.0, // Default: no reserve (backward compatible) +diff --git a/crates/ml/src/trainers/dqn/data_loading.rs b/crates/ml/src/trainers/dqn/data_loading.rs +index 7063e7a71..cf47afbd6 100644 +--- a/crates/ml/src/trainers/dqn/data_loading.rs ++++ b/crates/ml/src/trainers/dqn/data_loading.rs +@@ -397,15 +397,18 @@ impl DQNTrainer { + match ofi_calculator.calculate(snap) { + Ok(features) => { + if features.is_valid() { +- ofi_per_bar.push(features.to_array()); ++ let arr8 = features.to_array(); ++ let mut arr20 = [0.0_f64; 20]; ++ arr20[..8].copy_from_slice(&arr8); ++ ofi_per_bar.push(arr20); + } else { +- ofi_per_bar.push([0.0; 8]); ++ ofi_per_bar.push([0.0; 20]); + } + } +- Err(_) => ofi_per_bar.push([0.0; 8]), ++ Err(_) => ofi_per_bar.push([0.0; 20]), + } + } else { +- ofi_per_bar.push([0.0; 8]); ++ ofi_per_bar.push([0.0; 20]); + } + } + +diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs +index 4bc1f5fc0..9875c4fa4 100644 +--- a/crates/ml/src/trainers/dqn/fused_training.rs ++++ b/crates/ml/src/trainers/dqn/fused_training.rs +@@ -1373,6 +1373,11 @@ impl FusedTrainingCtx { + } + } + ++ // Recursive confidence backward: MSE grad into trunk + conf weight gradients. ++ // Must run before Adam (which reads grad_buf for the parameter update). ++ self.trainer.launch_recursive_confidence_backward(self.batch_size) ++ .map_err(|e| anyhow::anyhow!("Recursive confidence backward: {e}"))?; ++ + // Regime-adaptive PER scaling. + self.trainer.regime_scale_td_errors() + .map_err(|e| anyhow::anyhow!("Regime PER scaling: {e}"))?; +@@ -1885,6 +1890,12 @@ impl FusedTrainingCtx { + self.gpu_iqn.as_ref().map(|iqn| iqn.iqr_buf_ptr()) + } + ++ /// Device pointer to ISV signals [8] pinned buffer for adaptive hold enforcement. ++ /// Returns the dev_ptr (u64) that the experience collector passes to env_step. ++ pub(crate) fn isv_signals_dev_ptr(&self) -> u64 { ++ self.trainer.isv_signals_dev_ptr() ++ } ++ + /// Raw device pointer to the ensemble Q-value variance buffer and its length. + /// The buffer contains per-atom variance across K ensemble heads: [B * num_atoms] f32. + /// Populated after each training step by the aggregate kernel. +@@ -2230,6 +2241,10 @@ impl FusedTrainingCtx { + pub(crate) fn num_atoms(&self) -> usize { self.trainer.config().num_atoms } + pub(crate) fn update_adaptive_clip(&mut self, grad_norm: f32) { self.trainer.update_adaptive_clip(grad_norm); } + pub(crate) fn adaptive_clip_value(&self) -> f32 { self.trainer.adaptive_clip_value() } ++ /// Raw pointer to plan_params_buf [B, 6] for trade plan integration. ++ pub(crate) fn plan_params_buf_ptr(&self) -> u64 { ++ self.trainer.plan_params_buf_ptr() ++ } + + /// Raw device pointer to the flat f32 target params buffer (for DtoD sync). + pub(crate) fn target_params_flat_ptr(&self) -> u64 { +diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs +index 7ee62899a..2d21515c1 100644 +--- a/crates/ml/src/trainers/dqn/trainer/constructor.rs ++++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs +@@ -49,8 +49,8 @@ impl DQNTrainer { + // Align input_dim to 8 so the log matches the actual model dimensions. + // (device not yet created, so use the formula directly -- CUDA always aligns) + let ofi_pre = !hyperparams.mbp10_data_dir.is_empty(); +- // 42 market + 8 portfolio + 16 multi-timeframe = 66 base, +8 OFI = 74 +- let input_dim: usize = if ofi_pre { 80 } else { 72 }; // (74+7)&!7=80, (66+7)&!7=72 ++ // 42 market + 8 portfolio + 16 multi-timeframe = 66 base, +20 OFI = 86 ++ let input_dim: usize = if ofi_pre { 88 } else { 72 }; // (86+7)&!7=88, (66+7)&!7=72 + let output_dim: usize = 5; + let hidden_dims: Vec = match hyperparams.hidden_dim_base { + Some(base) => { +@@ -242,18 +242,18 @@ impl DQNTrainer { + // [0..42) 42 market features (OHLCV, technical, patterns, volume, time, statistical) + // [42..50) 8 portfolio features (position, P&L, drawdown, etc.) + // [50..66) 16 multi-timeframe features (4 windows x 4 features) +- // [66..74) 8 OFI features (from MBP-10 order book data, when mbp10_data_dir set) +- // Raw state_dim: 66 without OFI, 74 with OFI. ++ // [66..86) 20 OFI features (from MBP-10 order book data, when mbp10_data_dir set) ++ // Raw state_dim: 66 without OFI, 86 with OFI. + // market_dim: always 42 (bottleneck compresses only base market features). + // OFI features bypass bottleneck via portfolio_dim (fed directly to shared trunk). + // + // GpuTensor core alignment: state_dim is rounded up to the next multiple of 8 +- // (66->72, 74->80) so that cuBLAS dispatches BF16 HMMA instructions instead ++ // (66->72, 86->88) so that cuBLAS dispatches BF16 HMMA instructions instead + // of falling back to scalar FMA. The extra columns are zero-padded at the + // data pipeline boundaries (GpuPreloadedData and train_batch CPU path). + let ofi_enabled = !hyperparams.mbp10_data_dir.is_empty(); +- let raw_state_dim: usize = if ofi_enabled { 74 } else { 66 }; // 42 market + 8 portfolio + 16 MTF + (8 OFI) +- let full_state_dim = (raw_state_dim + 7) & !7; // aligned: 80 with OFI, 72 without ++ let raw_state_dim: usize = if ofi_enabled { 86 } else { 66 }; // 42 market + 8 portfolio + 16 MTF + (20 OFI) ++ let full_state_dim = (raw_state_dim + 7) & !7; // aligned: 88 with OFI, 72 without + + // state_dim is always the FULL padded dimension. The GpuDqnTrainer + // handles bottleneck reduction internally via compute_param_sizes(). +diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs +index 1496323c1..0c7e56116 100644 +--- a/crates/ml/src/trainers/dqn/trainer/metrics.rs ++++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs +@@ -70,6 +70,8 @@ impl DQNTrainer { + metrics.add_metric("best_sharpe", self.best_sharpe); + metrics.add_metric("best_val_loss", self.best_val_loss); + metrics.add_metric("best_epoch", self.best_epoch as f64); ++ metrics.add_metric("epoch_sharpe", self.best_sharpe); ++ metrics.add_metric("val_loss", self.best_val_loss); + + // Per-epoch history for overfitting detection (last fold only — history resets per fold) + let first_epoch_loss = self.loss_history.first().copied().unwrap_or(f64::NAN); +@@ -442,7 +444,7 @@ impl DQNTrainer { + if self.gpu_evaluator.is_none() { + let ofi_enabled = self.ofi_features.is_some(); + let market_dim: usize = 42; +- let feature_dim: usize = if ofi_enabled { 50 } else { market_dim }; ++ let feature_dim: usize = if ofi_enabled { 62 } else { market_dim }; + + // Build a single window from all val_data + let mut prices: Vec<[f32; 4]> = Vec::with_capacity(self.val_data.len()); +@@ -456,7 +458,7 @@ impl DQNTrainer { + let close = if target.len() >= 2 { target[0] as f32 } else { fv[3] as f32 }; + prices.push([close, close, close, close]); + +- // Features: 42 market features, plus 8 OFI if enabled ++ // Features: 42 market features, plus 20 OFI if enabled + let fv_slice = &fv[..market_dim.min(fv.len())]; + let mut fv_f32: Vec = fv_slice.iter().map(|&v| v as f32).collect(); + if ofi_enabled { +@@ -485,7 +487,7 @@ impl DQNTrainer { + contract_multiplier: hp.contract_multiplier as f32, + margin_pct: hp.margin_pct as f32, + max_leverage: hp.max_leverage as f32, +- ofi_dim: if ofi_enabled { 8 } else { 0 }, ++ ofi_dim: if ofi_enabled { 20 } else { 0 }, + min_hold_bars: hp.min_hold_bars as i32, + bars_per_day: hp.bars_per_day as f32, + trading_days_per_year: hp.trading_days_per_year as f32, +diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs +index 7a65dbce9..91ae0fcf8 100644 +--- a/crates/ml/src/trainers/dqn/trainer/mod.rs ++++ b/crates/ml/src/trainers/dqn/trainer/mod.rs +@@ -407,8 +407,8 @@ pub struct DQNTrainer { + /// Pre-computed OFI features per bar (indexed by global bar position). + /// Populated during data loading when MBP-10 order book data is available. + /// Passed as `regime_features` in `TradingState::from_normalized()`. +- /// Arc-shared to avoid 2.68 GB copy per hyperopt trial (41.9M × 8 × 8 bytes). +- pub ofi_features: Option>, ++ /// Arc-shared to avoid copy per hyperopt trial. ++ pub ofi_features: Option>, + /// Number of training bars (OFI offset for validation data). + /// val_data[i] corresponds to ofi_features[ofi_val_offset + i]. + pub(crate) ofi_val_offset: usize, +@@ -736,7 +736,7 @@ impl DQNTrainer { + }; + + // Upload FULL dataset to GPU FIRST so we can use GPU prefix sums for fold generation +- let ofi_slice: Vec<[f64; 8]> = self.ofi_features.as_ref() ++ let ofi_slice: Vec<[f64; 20]> = self.ofi_features.as_ref() + .map(|ofi| ofi.to_vec()) + .unwrap_or_default(); + self.init_from_fxcache(&features, &targets, &ofi_slice).await?; +@@ -1212,7 +1212,7 @@ impl DQNTrainer { + &mut self, + features: &[[f64; 42]], + targets: &[[f64; 4]], +- ofi: &[[f64; 8]], ++ ofi: &[[f64; 20]], + ) -> Result<()> { + let stream = self.cuda_stream.as_ref() + .ok_or_else(|| anyhow::anyhow!("CUDA stream required for init_from_fxcache"))?; +@@ -1223,14 +1223,14 @@ impl DQNTrainer { + ).map_err(|e| anyhow::anyhow!("DqnGpuData::upload_slices failed: {e}"))?; + + let ofi_enabled = gpu_data.ofi_features.is_some(); +- let raw_dim = if ofi_enabled { 53 } else { 45 }; ++ let raw_dim = if ofi_enabled { 65 } else { 45 }; + let aligned_dim = (raw_dim + 7) & !7; + gpu_data.set_aligned_state_dim(aligned_dim); + + tracing::info!( + "init_from_fxcache: {} bars uploaded to GPU ({:.1} MB, OFI={})", + features.len(), +- (features.len() * (42 + 4 + 8) * 4) as f64 / 1_048_576.0, ++ (features.len() * (42 + 4 + 20) * 4) as f64 / 1_048_576.0, + ofi_enabled, + ); + self.gpu_data = Some(gpu_data); +diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs +index e073722b1..c9ae3ac07 100644 +--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs ++++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs +@@ -518,6 +518,22 @@ impl DQNTrainer { + } + } + ++ // ISV signals for adaptive hold enforcement ++ if let Some(ref fused) = self.fused_ctx { ++ let isv_ptr = fused.isv_signals_dev_ptr(); ++ if let Some(ref mut collector) = self.gpu_experience_collector { ++ collector.set_isv_signals_ptr(isv_ptr); ++ } ++ } ++ ++ // Trade plan params for plan activation/enforcement ++ if let Some(ref fused) = self.fused_ctx { ++ let plan_ptr = fused.plan_params_buf_ptr(); ++ if let Some(ref mut collector) = self.gpu_experience_collector { ++ collector.set_plan_params_ptr(plan_ptr); ++ } ++ } ++ + // IQR exploration bonus from IQN quantile spread + if let Some(ref fused) = self.fused_ctx { + let iqr_ptr = fused.iqn_iqr_ptr().unwrap_or(0); +diff --git a/crates/ml/tests/fxcache_roundtrip_test.rs b/crates/ml/tests/fxcache_roundtrip_test.rs +index d14573203..98b6b2710 100644 +--- a/crates/ml/tests/fxcache_roundtrip_test.rs ++++ b/crates/ml/tests/fxcache_roundtrip_test.rs +@@ -35,12 +35,12 @@ fn make_targets(n: usize) -> Vec<[f64; 4]> { + } + + /// Deterministic OFI vector for bar `i`. +-fn make_ofi(n: usize) -> Vec<[f64; 8]> { ++fn make_ofi(n: usize) -> Vec<[f64; 20]> { + (0..n) + .map(|i| { +- let mut row = [0.0_f64; 8]; +- for j in 0..8 { +- row[j] = (i * 8 + j) as f64 * 0.001; ++ let mut row = [0.0_f64; 20]; ++ for j in 0..20 { ++ row[j] = (i * 20 + j) as f64 * 0.001; + } + row + }) +@@ -119,7 +119,7 @@ fn test_fxcache_f32_roundtrip() { + data.targets[i][j] + ); + } +- for j in 0..8 { ++ for j in 0..20 { + let diff = (data.ofi[i][j] - ofi[i][j]).abs(); + assert!( + diff < 1e-4, +@@ -187,7 +187,7 @@ fn test_fxcache_empty() { + // Writer rejects 0 bars. + let features: Vec<[f64; 42]> = vec![]; + let targets: Vec<[f64; 4]> = vec![]; +- let ofi: Vec<[f64; 8]> = vec![]; ++ let ofi: Vec<[f64; 20]> = vec![]; + let timestamps: Vec = vec![]; + let key = test_cache_key(); + +@@ -205,7 +205,7 @@ fn test_fxcache_empty() { + header[8..10].copy_from_slice(&1u16.to_le_bytes()); // version=1 + header[10..12].copy_from_slice(&42u16.to_le_bytes()); // feat_dim + header[12..14].copy_from_slice(&4u16.to_le_bytes()); // target_dim +- header[14..16].copy_from_slice(&8u16.to_le_bytes()); // ofi_dim ++ header[14..16].copy_from_slice(&20u16.to_le_bytes()); // ofi_dim + header[16..24].copy_from_slice(&0u64.to_le_bytes()); // bar_count=0 + // cache_key + reserved stay zeroed. + +diff --git a/crates/ml/tests/smoke_test_real_data.rs b/crates/ml/tests/smoke_test_real_data.rs +index 047d41b6b..f291c66d9 100644 +--- a/crates/ml/tests/smoke_test_real_data.rs ++++ b/crates/ml/tests/smoke_test_real_data.rs +@@ -441,7 +441,7 @@ mod gpu_smoke { + assert!(n > 200, "Need at least 200 bars for GPU test, got {n}"); + + // 3. Optionally load OFI features from MBP-10 +- let ofi_features: Option> = mbp10_dir.and_then(|dir| { ++ let ofi_features: Option> = mbp10_dir.and_then(|dir| { + let files: Vec<_> = std::fs::read_dir(dir) + .ok()? + .filter_map(|e| e.ok()) +@@ -453,7 +453,13 @@ mod gpu_smoke { + .collect(); + let file = files.first()?; + info!(file = %file.display(), "Loading OFI"); +- ml::features::mbp10_loader::compute_ofi_from_file(file).ok() ++ ml::features::mbp10_loader::compute_ofi_from_file(file).ok().map(|ofi8| { ++ ofi8.into_iter().map(|row| { ++ let mut wide = [0.0_f64; 20]; ++ wide[..8].copy_from_slice(&row); ++ wide ++ }).collect::>() ++ }) + }); + + let ofi_count = ofi_features.as_ref().map_or(0, |v| v.len()); +diff --git a/infra/k8s/argo/train-template.yaml b/infra/k8s/argo/train-template.yaml +index 818fd64b3..316f2630b 100644 +--- a/infra/k8s/argo/train-template.yaml ++++ b/infra/k8s/argo/train-template.yaml +@@ -334,7 +334,19 @@ spec: + exit 1 + fi + ++ # Auto-invalidate stale fxcache (version mismatch → delete and regenerate) ++ echo "=== Checking existing fxcache ===" ++ for f in /feature-cache/*.fxcache; do ++ if [ -f "$f" ]; then ++ # Try to validate — precompute_features will fail on version mismatch ++ echo "Found existing: $f" ++ fi ++ done ++ + echo "=== Running precompute_features (SHA: $SHA) ===" ++ # --yes auto-confirms. If fxcache exists with wrong version, ++ # the binary detects the mismatch and regenerates. ++ rm -f /feature-cache/*.fxcache + $BINARY \ + --data-dir /data/futures-baseline \ + --mbp10-data-dir /data/futures-baseline-mbp10 \ +diff --git a/scripts/argo-precompute.sh b/scripts/argo-precompute.sh +new file mode 100755 +index 000000000..d81397007 +--- /dev/null ++++ b/scripts/argo-precompute.sh +@@ -0,0 +1,52 @@ ++#!/usr/bin/env bash ++# Regenerate fxcache on PVC — uses the train workflow's ensure-binary + ensure-fxcache steps. ++# Submits a lightweight workflow that compiles the binary and regenerates the feature cache. ++# No GPU needed — runs entirely on ci-compile-cpu. ++# ++# Usage: ./scripts/argo-precompute.sh [--branch main] [--watch] ++set -euo pipefail ++ ++BRANCH="main" ++WATCH=false ++ ++while [[ $# -gt 0 ]]; do ++ case "$1" in ++ --branch) BRANCH="$2"; shift 2;; ++ --watch) WATCH=true; shift;; ++ -h|--help) ++ echo "Usage: $0 [--branch ] [--watch]" ++ echo "" ++ echo "Regenerates fxcache on PVC with current code." ++ echo "Uses ensure-binary + ensure-fxcache from the train workflow." ++ echo "Runs on ci-compile-cpu — no GPU required." ++ echo "" ++ echo "The fxcache auto-invalidates on version mismatch" ++ echo "(FXCACHE_VERSION in fxcache.rs). Old cache is deleted" ++ echo "and regenerated automatically." ++ exit 0;; ++ *) echo "Unknown option: $1"; exit 1;; ++ esac ++done ++ ++echo "=== Argo Precompute fxcache ===" ++echo " branch: $BRANCH" ++echo "" ++ ++# Submit a training workflow with 0 epochs — it will compile + ensure-fxcache, then stop. ++# hyperopt-trials=0 + train-epochs=0 means only ensure-binary + ensure-fxcache run. ++WORKFLOW_NAME=$(argo submit -n foxhunt --from=wftmpl/train \ ++ -p git-branch="$BRANCH" \ ++ -p hyperopt-trials=0 \ ++ -p train-epochs=0 \ ++ -o name 2>&1) ++ ++echo "Submitted: $WORKFLOW_NAME" ++echo " ensure-binary → ensure-fxcache (compile + regenerate)" ++echo "" ++echo "Monitor: kubectl -n foxhunt logs -f $WORKFLOW_NAME --tail=50" ++ ++if $WATCH; then ++ echo "" ++ echo "Watching..." ++ argo watch -n foxhunt "$WORKFLOW_NAME" ++fi +diff --git a/services/trading_agent_service/Cargo.toml b/services/trading_agent_service/Cargo.toml +index 0d9fe65de..a7416a22e 100644 +--- a/services/trading_agent_service/Cargo.toml ++++ b/services/trading_agent_service/Cargo.toml +@@ -55,6 +55,8 @@ common = { workspace = true, features = ["database"] } + config = { workspace = true, features = ["postgres"] } + risk = { path = "../../crates/risk" } + ml = { path = "../../crates/ml", default-features = false, features = ["minimal-inference"] } ++ml-features.workspace = true ++data.workspace = true + + # Utilities + thiserror.workspace = true +diff --git a/services/trading_agent_service/src/lib.rs b/services/trading_agent_service/src/lib.rs +index c55aef827..fe23225ab 100644 +--- a/services/trading_agent_service/src/lib.rs ++++ b/services/trading_agent_service/src/lib.rs +@@ -27,4 +27,5 @@ pub mod orders; + pub mod regime; + pub mod service; + pub mod strategies; ++pub mod tick_processor; + pub mod universe; +diff --git a/services/trading_agent_service/src/tick_processor.rs b/services/trading_agent_service/src/tick_processor.rs +new file mode 100644 +index 000000000..d0e67227c +--- /dev/null ++++ b/services/trading_agent_service/src/tick_processor.rs +@@ -0,0 +1,201 @@ ++//! Real-time tick processor for Databento MBP-10 -> microstructure features. ++//! ++//! Feeds each MBP-10 event and trade into OFICalculator + MicrostructureState. ++//! At bar close, snapshots 20 features for the DQN forward pass. ++//! Between bars, provides intermediate snapshots for speculative compute. ++ ++use data::providers::databento::mbp10::Mbp10Snapshot; ++use ml_features::ofi_calculator::{MicrostructureState, OFICalculator, OFIFeatures}; ++use tokio::sync::watch; ++use tracing::info; ++ ++/// Real-time tick processor for live trading. ++/// ++/// Wraps `OFICalculator` (8 base features) and `MicrostructureState` (12 extended ++/// features) to produce a 20-feature vector per bar from streaming MBP-10 ticks. ++pub struct TickProcessor { ++ /// OFI calculator (base 8 features) ++ ofi: OFICalculator, ++ /// Microstructure state (12 extended features) ++ micro: MicrostructureState, ++ /// Last computed OFI features (cached for mid-bar snapshots) ++ last_ofi: OFIFeatures, ++ /// Current bar start timestamp (nanoseconds) ++ bar_start_ns: u64, ++ /// Bar duration (nanoseconds, typically 60_000_000_000 for 60s) ++ bar_duration_ns: u64, ++ /// Channel to broadcast feature snapshots at bar close ++ feature_tx: watch::Sender<[f64; 20]>, ++ /// Tick count in current bar ++ tick_count: u64, ++} ++ ++impl TickProcessor { ++ /// Create a new `TickProcessor`. ++ /// ++ /// Returns the processor and a `watch::Receiver` that receives the 20-feature ++ /// snapshot every time a bar closes. ++ pub fn new(bar_duration_ns: u64) -> (Self, watch::Receiver<[f64; 20]>) { ++ let (tx, rx) = watch::channel([0.0; 20]); ++ let now_ns = now_unix_ns(); ++ ( ++ Self { ++ ofi: OFICalculator::new(), ++ micro: MicrostructureState::new(now_ns, bar_duration_ns), ++ last_ofi: OFIFeatures::zeros(), ++ bar_start_ns: now_ns, ++ bar_duration_ns, ++ feature_tx: tx, ++ tick_count: 0, ++ }, ++ rx, ++ ) ++ } ++ ++ /// Process an MBP-10 snapshot event. ++ /// ++ /// Updates both the OFI calculator and microstructure state. The OFI ++ /// features are cached so `current_features()` can be called without ++ /// requiring the snapshot again. ++ pub fn on_snapshot(&mut self, snapshot: &Mbp10Snapshot) { ++ // Update microstructure state (spread dynamics, queue depletion, etc.) ++ self.micro.update_snapshot(snapshot); ++ self.tick_count += 1; ++ ++ // Compute OFI features from snapshot and cache the result ++ if let Ok(ofi_features) = self.ofi.calculate(snapshot) { ++ self.micro ++ .update_ofi_derived(ofi_features.ofi_level1, ofi_features.vpin); ++ self.last_ofi = ofi_features; ++ } ++ } ++ ++ /// Process a trade event. ++ /// ++ /// Feeds trade data to both the OFI calculator (VPIN, Kyle's Lambda, ++ /// trade imbalance) and microstructure state (Hawkes intensity, aggression). ++ pub fn on_trade(&mut self, price: f64, volume: u64, is_buy: bool, timestamp_ns: u64) { ++ self.ofi.feed_trade(price, volume, is_buy); ++ self.micro.update_trade(price, volume, is_buy, timestamp_ns); ++ } ++ ++ /// Snapshot current features (non-mutating, safe to call mid-bar for speculative compute). ++ /// ++ /// Returns 20 features: `[0..8]` = base OFI, `[8..20]` = microstructure state. ++ pub fn current_features(&self) -> [f64; 20] { ++ let base_8 = self.last_ofi.to_array(); ++ let micro_12 = self.micro.snapshot(); ++ let mut out = [0.0; 20]; ++ out[..8].copy_from_slice(&base_8); ++ out[8..20].copy_from_slice(µ_12); ++ out ++ } ++ ++ /// Called at bar close -- snapshot final features, broadcast, and reset for next bar. ++ /// ++ /// Returns the 20-feature vector that was broadcast. ++ pub fn bar_close(&mut self) -> [f64; 20] { ++ let features = self.current_features(); ++ let _ = self.feature_tx.send(features); ++ ++ info!( ++ tick_count = self.tick_count, ++ "Bar close: {} ticks processed, features broadcast", ++ self.tick_count ++ ); ++ ++ // Reset for next bar ++ let now_ns = now_unix_ns(); ++ self.bar_start_ns = now_ns; ++ self.micro.reset(now_ns, self.bar_duration_ns); ++ self.ofi = OFICalculator::new(); ++ self.last_ofi = OFIFeatures::zeros(); ++ self.tick_count = 0; ++ ++ features ++ } ++ ++ /// Returns the tick count in the current bar. ++ pub const fn tick_count(&self) -> u64 { ++ self.tick_count ++ } ++ ++ /// Returns the current bar start timestamp in nanoseconds. ++ pub const fn bar_start_ns(&self) -> u64 { ++ self.bar_start_ns ++ } ++} ++ ++/// Helper: current wall-clock time as nanoseconds since Unix epoch. ++fn now_unix_ns() -> u64 { ++ std::time::SystemTime::now() ++ .duration_since(std::time::UNIX_EPOCH) ++ .map_or(0, |d| d.as_nanos() as u64) ++} ++ ++#[cfg(test)] ++mod tests { ++ use super::*; ++ use data::providers::databento::mbp10::BidAskPair; ++ ++ fn make_snapshot(bid_px: i64, ask_px: i64, bid_sz: u32, ask_sz: u32) -> Mbp10Snapshot { ++ let level = BidAskPair { ++ bid_px, ++ bid_sz, ++ bid_ct: 10, ++ ask_px, ++ ask_sz, ++ ask_ct: 10, ++ }; ++ Mbp10Snapshot { ++ symbol: "ES.FUT".to_string(), ++ timestamp: now_unix_ns(), ++ levels: vec![level; 10], ++ sequence: 1, ++ trade_count: 0, ++ } ++ } ++ ++ #[test] ++ fn test_new_returns_zero_features() { ++ let (proc, rx) = TickProcessor::new(60_000_000_000); ++ assert_eq!(*rx.borrow(), [0.0; 20]); ++ assert_eq!(proc.tick_count(), 0); ++ } ++ ++ #[test] ++ fn test_on_snapshot_increments_tick_count() { ++ let (mut proc, _rx) = TickProcessor::new(60_000_000_000); ++ let snap = make_snapshot(5000_000_000_000, 5001_000_000_000, 100, 100); ++ proc.on_snapshot(&snap); ++ assert_eq!(proc.tick_count(), 1); ++ } ++ ++ #[test] ++ fn test_bar_close_resets_state() { ++ let (mut proc, rx) = TickProcessor::new(60_000_000_000); ++ let snap = make_snapshot(5000_000_000_000, 5001_000_000_000, 100, 100); ++ proc.on_snapshot(&snap); ++ proc.on_snapshot(&snap); ++ assert_eq!(proc.tick_count(), 2); ++ ++ let features = proc.bar_close(); ++ assert_eq!(proc.tick_count(), 0); ++ assert_eq!(*rx.borrow(), features); ++ } ++ ++ #[test] ++ fn test_current_features_is_20_dim() { ++ let (proc, _rx) = TickProcessor::new(60_000_000_000); ++ let features = proc.current_features(); ++ assert_eq!(features.len(), 20); ++ } ++ ++ #[test] ++ fn test_on_trade_feeds_through() { ++ let (mut proc, _rx) = TickProcessor::new(60_000_000_000); ++ // Should not panic ++ proc.on_trade(5000.25, 10, true, now_unix_ns()); ++ proc.on_trade(5000.50, 5, false, now_unix_ns()); ++ } ++}