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