gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
Files
74bdc58536b73ecaa3aa0bdeb766811d41602b4c
foxhunt/pal_precommit.changeset
jgrusewski 11b1a1ca9f fix: zen precommit — 2 criticals + 3 lows resolved 
CRITICAL: ISV+plan heads wired into submit_forward_ops_main (training).
CRITICAL: backtest single-step 3-branch→4-branch action encoding.
LOW: Remove dead bar_start_ns/bar_duration_ns from MicrostructureState.
LOW: Pre-allocated eps_host_buf replaces Vec allocation in hot path.
LOW: training_sharpe + val_loss added to metrics additional_metrics.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
2026-04-17 02:05:20 +02:00

4743 lines
214 KiB
Plaintext
Raw Blame History

This file contains ambiguous Unicode characters
This file contains Unicode characters that might be confused with other characters. If you think that this is intentional, you can safely ignore this warning. Use the Escape button to reveal them.
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(&micro_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<Vec<RlTrainingResult>> {
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<f32>,
/// 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<f32>, // [B, b1_size] saved for backward
risk_grad_buf: CudaSlice<f32>, // [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<f32>, // [ISV_EMB_DIM] = [8]
+ branch_gate_buf: CudaSlice<f32>, // [4]
+ gamma_mod_buf: CudaSlice<f32>, // [1]
+ temporal_weight_buf: CudaSlice<f32>, // [SH2] per-feature temporal routing
+ isv_temporal_route_kernel: CudaFunction,
+ gamma_buf: CudaSlice<f32>, // [B] per-sample effective gamma
+ predicted_error_buf: CudaSlice<f32>, // [B] recursive confidence output
+
+ // ── Recursive confidence kernels ──
+ recursive_conf_fwd_kernel: CudaFunction,
+ recursive_conf_bwd_kernel: CudaFunction,
+
+ // ── Trade plan head ──
+ plan_params_buf: CudaSlice<f32>, // [B, 6]
+ trade_plan_fwd_kernel: CudaFunction,
+
+ // ── Speculative inference cache ──
+ /// Cached trunk output [B, SH2] from between-bar speculative forward.
+ speculative_h_s2: CudaSlice<f32>,
+ /// Last features [state_dim] used for speculation (host-side for L2 comparison).
+ speculative_features: Vec<f32>,
+ /// 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(&param_sizes, 68);
+ let b_fc1 = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 69);
+ let w_fc2 = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 70);
+ let b_fc2 = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 71);
+ let w_gate = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 72);
+ let b_gate = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 73);
+ let w_gamma = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 74);
+ let b_gamma = self.ptrs.params_ptr + padded_byte_offset(&param_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(&param_sizes, 78);
+ let b_gate = self.ptrs.params_ptr + padded_byte_offset(&param_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(&param_sizes, 80);
+ let b_route = self.ptrs.params_ptr + padded_byte_offset(&param_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(&param_sizes, 76);
+ let b_conf = self.ptrs.params_ptr + padded_byte_offset(&param_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(&param_sizes, 82);
+ let b_fc = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 83);
+ let w_out = self.ptrs.params_ptr + padded_byte_offset(&param_sizes, 84);
+ let b_out = self.ptrs.params_ptr + padded_byte_offset(&param_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(&param_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(&param_sizes, 76);
+ let d_b_conf = self.ptrs.grad_buf + padded_byte_offset(&param_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(&param_sizes, 50);
- let b_regime_ptr = self.ptrs.params_ptr + padded_byte_offset(&param_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<Self, MLError> {
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::<f32>(),
+ );
+ 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::<f32>(),
+ );
+ 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::<f32>(),
+ );
+ 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::<f32>(),
+ );
+ 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::<f32>();
+ 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::<f32>();
+ 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::<f32>();
+ 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::<f32>());
+ 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::<f32>(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::<f32>(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::<f32>(ISV_EMB_DIM)
+ .map_err(|e| MLError::ModelError(format!("isv_embedding_buf alloc: {e}")))?;
+ let branch_gate_buf = stream.alloc_zeros::<f32>(4)
+ .map_err(|e| MLError::ModelError(format!("branch_gate_buf alloc: {e}")))?;
+ let gamma_mod_buf = stream.alloc_zeros::<f32>(1)
+ .map_err(|e| MLError::ModelError(format!("gamma_mod_buf alloc: {e}")))?;
+ let temporal_weight_buf = stream.alloc_zeros::<f32>(sh2)
+ .map_err(|e| MLError::ModelError(format!("temporal_weight_buf alloc: {e}")))?;
+ let gamma_buf = stream.alloc_zeros::<f32>(b)
+ .map_err(|e| MLError::ModelError(format!("gamma_buf alloc: {e}")))?;
+ let predicted_error_buf = stream.alloc_zeros::<f32>(b)
+ .map_err(|e| MLError::ModelError(format!("predicted_error_buf alloc: {e}")))?;
+ let plan_params_buf = stream.alloc_zeros::<f32>(b * 6)
+ .map_err(|e| MLError::ModelError(format!("plan_params_buf alloc: {e}")))?;
+
+ // ── Speculative inference cache ──
+ let speculative_h_s2 = stream.alloc_zeros::<f32>(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<CudaSlice<f32>, 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::<f32>(), 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::<f32>(), 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<f32> { &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<f32>,
// Per-episode state buffers [alloc_episodes, ...]
- portfolio_states: CudaSlice<f32>, // [alloc_episodes * PORTFOLIO_STRIDE] (20 floats per episode)
+ portfolio_states: CudaSlice<f32>, // [alloc_episodes * PORTFOLIO_STRIDE] (30 floats per episode)
episode_starts_buf: CudaSlice<i32>,// [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::<f32>(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<f64>)],
- ofi_data: Option<&[[f64; 8]]>,
+ ofi_data: Option<&[[f64; 20]]>,
wf_config: &GpuWalkForwardConfig,
stream: &Arc<CudaStream>,
) -> Result<Self, MLError> {
@@ -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<CudaStream>,
) -> Result<Self, MLError> {
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<CudaStream>,
) -> 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<CudaStream>,
) -> Result<CudaSlice<f32>, 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<std::fs::File>,
features: &[[f64; FEAT_DIM]],
@@ -337,7 +342,7 @@ pub fn load_fxcache(path: &Path) -> Result<FxCacheData> {
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, &timestamps, [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<String>,
/// 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<Arc<[[f64; 8]]>>,
+ preloaded_ofi_features: Option<Arc<[[f64; 20]]>>,
/// 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<Arc<[[f64; 8]]>> {
+ fn load_ofi_features(&self) -> Option<Arc<[[f64; 20]]>> {
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<Vec<[f64; 8]>> {
+ fn load_ofi_features(&self) -> Option<Vec<[f64; 20]>> {
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<usize> = 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<f32> = 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<[[f64; 8]]>>,
+ /// Arc-shared to avoid copy per hyperopt trial.
+ pub ofi_features: Option<Arc<[[f64; 20]]>>,
/// 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<i64> = 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<Vec<[f64; 8]>> = mbp10_dir.and_then(|dir| {
+ let ofi_features: Option<Vec<[f64; 20]>> = 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::<Vec<_>>()
+ })
});
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 <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(&micro_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());
+ }
+}
Reference in New Issue View Git Blame Copy Permalink