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()); + } +}