diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 7ad40bbf3..4ee4608e6 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -158,6 +158,71 @@ __device__ static int stop_check_isv( return 0; } +// CRT-A2: Wiener-α adaptive EMA on max-conviction-across-horizons. +// Reads alpha_probs[h] (broadcast across backtests), computes the current +// max |alpha-0.5|*2, updates per-backtest second-order EMAs of sample +// variance + difference variance, derives adaptive α with floor 0.4 per +// pearl_wiener_alpha_floor_for_nonstationary, and blends. +// First-observation bootstrap (prev_ema == 0): replace directly per +// pearl_first_observation_bootstrap. +// +// Writes new EMA to conviction_ema[b], updates the two second-order +// variance EMAs, and returns the smoothed max-conviction. Direction is +// recovered separately at each per-horizon sizing call (raw alpha sign); +// only the magnitude is smoothed so genuine reversals respond at event +// rate while micro-jitter is suppressed. +// +// Single-thread-per-backtest invariant: caller must already have +// gated on threadIdx.x == 0. +__device__ static float update_conviction_ema( + int b, + const float* __restrict__ alpha_probs, + float* __restrict__ conviction_ema, + float* __restrict__ conviction_diff_var_ema, + float* __restrict__ conviction_sample_var_ema +) { + float max_conv = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float c = fabsf(alpha_probs[h] - 0.5f) * 2.0f; + if (c > max_conv) max_conv = c; + } + // Defensive clamp to [0, 1] — alpha_probs should be in [0, 1] by + // contract but bound it before feeding into a variance EMA. + if (max_conv > 1.0f) max_conv = 1.0f; + if (max_conv < 0.0f) max_conv = 0.0f; + + // Update second-order EMAs (diff_var, sample_var) with fixed α=0.1. + // These track the variance of the signal and the variance of changes, + // and are themselves bootstrapped with the first observation. + const float prev_ema = conviction_ema[b]; + const float diff = max_conv - prev_ema; + const float diff_var_prev = conviction_diff_var_ema[b]; + const float new_diff_var = (diff_var_prev == 0.0f) + ? (diff * diff) // first-observation bootstrap + : (0.9f * diff_var_prev + 0.1f * diff * diff); + conviction_diff_var_ema[b] = new_diff_var; + + const float sample_var_prev = conviction_sample_var_ema[b]; + const float new_sample_var = (sample_var_prev == 0.0f) + ? (max_conv * max_conv) // first-observation bootstrap + : (0.9f * sample_var_prev + 0.1f * max_conv * max_conv); + conviction_sample_var_ema[b] = new_sample_var; + + // Wiener α with floor at 0.4 (pearl_wiener_alpha_floor_for_nonstationary). + const float eps_v = 1e-6f; + const float alpha_raw = new_diff_var / (new_diff_var + new_sample_var + eps_v); + const float alpha_active = (alpha_raw > 0.4f) ? alpha_raw : 0.4f; + + // First-observation bootstrap: prev_ema==0 → replace directly. Otherwise + // blend with Wiener-α. + const float new_ema = (prev_ema == 0.0f) + ? max_conv + : (alpha_active * max_conv + (1.0f - alpha_active) * prev_ema); + conviction_ema[b] = new_ema; + return new_ema; +} + extern "C" __global__ void decision_policy_default( const float* __restrict__ alpha_probs, // [N_HORIZONS] broadcast const unsigned char* __restrict__ isv_kelly_base, // [n_backtests * N_HORIZONS * sizeof(IsvKellyState)] @@ -187,6 +252,10 @@ extern "C" __global__ void decision_policy_default( const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] (via lob_state.cuh) // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] + // CRT-A2: Wiener-α conviction-EMA state (per-backtest). + float* __restrict__ conviction_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conviction_diff_var_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conviction_sample_var_ema_per_b, // [n_backtests] — read + write int n_backtests ) { int b = blockIdx.x; @@ -202,23 +271,33 @@ extern "C" __global__ void decision_policy_default( return; } + // Compute raw max conviction once (used by threshold gate AND by the + // CRT-A2 magnitude-scale below). + float raw_max_conv = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + raw_max_conv = fmaxf(raw_max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); + } + + // CRT-A2: Wiener-α adaptive EMA on max-conviction. Runs ahead of the + // threshold gate so the EMA tracks all events (gating off an unsmoothed + // tick would let the EMA grow stale on long gated stretches). + const float conv_ema = update_conviction_ema( + b, alpha_probs, + conviction_ema_per_b, + conviction_diff_var_ema_per_b, + conviction_sample_var_ema_per_b); + // P4: threshold gate. Skip entirely if max conviction below threshold. // Pre-Kelly placement keeps the gate deterministic from alpha alone, // so the threshold pre-registration step (compute p60-p95 absolute // values from observed convictions) reflects exactly what gets gated // in deployment. - { - float max_conv = 0.0f; - #pragma unroll - for (int h = 0; h < N_HORIZONS; ++h) { - max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); - } - if (max_conv < threshold_per_b[b]) { - market_targets[b * 2 + 0] = 2; - market_targets[b * 2 + 1] = 0; - open_horizon_masks[b] = 0u; - return; - } + if (raw_max_conv < threshold_per_b[b]) { + market_targets[b * 2 + 0] = 2; + market_targets[b * 2 + 1] = 0; + open_horizon_masks[b] = 0u; + return; } // Per-backtest scalar reads (host-uploaded arrays). @@ -295,6 +374,26 @@ extern "C" __global__ void decision_policy_default( } } + // CRT-A2: rescale the aggregate magnitude by (conv_ema / raw_max_conv). + // The smoothed conviction acts as a unified scalar gain over the raw + // per-horizon contributions — per-horizon sizing logic (Kelly, cap_units, + // sharpe weights, attribution mask) is unchanged. Spec §4.2. + // + // The scale factor is in [0, 1] when raw_max_conv > 0 (since the EMA + // is a convex blend of historical raw_max_conv values, bounded by their + // running max). At the first observation the EMA == raw_max_conv exactly + // (bootstrap), so the scale is 1.0 and behaviour is bit-identical to the + // pre-A2 kernel for the very first decision. + // + // Guard: when raw_max_conv ≈ 0 the threshold gate above (when configured + // to non-zero) would already have returned. Under threshold=0.0 a raw_max_conv + // of 0 means every alpha is exactly 0.5 → final_size is also 0 (any dir + // contributions cancel in expectation under uniform Kelly/cap); the scale + // doesn't change that. We still divide-guard with eps to be safe. + if (raw_max_conv > 1e-9f) { + final_size *= (conv_ema / raw_max_conv); + } + // Round-to-nearest (truncation alone bites here: 0.8(f32) * 0.75 * 5.0 // = 2.99999976 → (int)2, off by one. Floor truncation for sub-1 // suppression is preserved via the explicit |lots| < 1 check. @@ -375,6 +474,10 @@ extern "C" __global__ void decision_policy_program( const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] + // CRT-A2: Wiener-α conviction-EMA state (per-backtest). + float* __restrict__ conviction_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conviction_diff_var_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conviction_sample_var_ema_per_b, // [n_backtests] — read + write int n_backtests ) { int b = blockIdx.x; @@ -395,19 +498,30 @@ extern "C" __global__ void decision_policy_program( return; } + // Compute raw max conviction once (used by threshold gate AND by the + // CRT-A2 final-aggregate rescale in OP_WRITE_ORDER). + float raw_max_conv = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + raw_max_conv = fmaxf(raw_max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); + } + + // CRT-A2: Wiener-α adaptive EMA on max-conviction. See decision_policy_default + // for rationale. The smoothed magnitude rescales final_size at the + // OP_WRITE_ORDER terminal — per-horizon attribution + Kelly remain + // bit-identical with raw |alpha-0.5|*2. + const float conv_ema = update_conviction_ema( + b, alpha_probs, + conviction_ema_per_b, + conviction_diff_var_ema_per_b, + conviction_sample_var_ema_per_b); + // P4: threshold gate. - { - float max_conv = 0.0f; - #pragma unroll - for (int h = 0; h < N_HORIZONS; ++h) { - max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); - } - if (max_conv < threshold_per_b[b]) { - market_targets[b * 2 + 0] = 2; - market_targets[b * 2 + 1] = 0; - open_horizon_masks[b] = 0u; - return; - } + if (raw_max_conv < threshold_per_b[b]) { + market_targets[b * 2 + 0] = 2; + market_targets[b * 2 + 1] = 0; + open_horizon_masks[b] = 0u; + return; } // Per-backtest scalar reads (host-uploaded arrays). @@ -564,8 +678,13 @@ extern "C" __global__ void decision_policy_program( case OP_WRITE_ORDER: { if (sp <= 0) break; sp -= 1; - const float final_size = stack_val[sp]; + float final_size = stack_val[sp]; const unsigned int attribution = stack_mask[sp]; + // CRT-A2: rescale the aggregate magnitude by (conv_ema / raw_max_conv). + // See decision_policy_default for rationale. + if (raw_max_conv > 1e-9f) { + final_size *= (conv_ema / raw_max_conv); + } int lots = (int)truncf(final_size + (final_size >= 0.0f ? 0.5f : -0.5f)); if (lots == 0) { market_targets[b * 2 + 0] = 2; // no-op diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 74b5a8d2b..54271c311 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -244,6 +244,16 @@ pub struct LobSimCuda { // S2.1/S2.2: max_hold diagnostic counters (2 remain after S2.2 prune). mh_kernel_calls_d: CudaSlice, // [n_backtests] mh_force_flat_seen_by_seed_d: CudaSlice, // [n_backtests] + + // CRT-A2: Wiener-α adaptive EMA state on max-conviction. Smooths + // event-rate alpha jitter so sizing magnitude doesn't oscillate + // bar-by-bar. All three slots start at zero (sentinel = "no + // observation yet") so the first decision-kernel call replaces them + // directly per pearl_first_observation_bootstrap. Updated in-place + // by both decision_policy_default and decision_policy_program. + conviction_ema_d: CudaSlice, // [n_backtests] — smoothed conviction + conviction_diff_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of (new−prev)² + conviction_sample_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of x² } impl LobSimCuda { @@ -488,6 +498,19 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc mh_force_flat_seen_by_seed_d")?; + // CRT-A2: Wiener-α conviction-EMA state. Zero-init means + // "first observation pending" — the decision kernel branches on + // prev_ema == 0 to bootstrap directly (no warmup zero-bias). + let conviction_ema_d = stream + .alloc_zeros::(n_backtests) + .context("alloc conviction_ema_d")?; + let conviction_diff_var_ema_d = stream + .alloc_zeros::(n_backtests) + .context("alloc conviction_diff_var_ema_d")?; + let conviction_sample_var_ema_d = stream + .alloc_zeros::(n_backtests) + .context("alloc conviction_sample_var_ema_d")?; + // CRT Phase A0.5 corrective: on-device conviction history buffer. // Capacity = 5M decisions × 4B/f32 = 20MB. Matches the prior // host-side `Vec::with_capacity(3_000_000)` plus headroom for @@ -578,6 +601,9 @@ impl LobSimCuda { last_bad_path_d, mh_kernel_calls_d, mh_force_flat_seen_by_seed_d, + conviction_ema_d, + conviction_diff_var_ema_d, + conviction_sample_var_ema_d, }) } @@ -713,6 +739,20 @@ impl LobSimCuda { Ok(buf[backtest_idx]) } + /// CRT-A2: read `conviction_ema_d[backtest_idx]`. Test-only accessor — + /// production paths inspect smoothed conviction on-device only. One DtoH. + pub fn read_conviction_ema(&self, backtest_idx: usize) -> Result { + anyhow::ensure!( + backtest_idx < self.n_backtests, + "backtest_idx {} >= n_backtests {}", + backtest_idx, + self.n_backtests + ); + let mut buf = vec![0.0f32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.conviction_ema_d, buf.as_mut_slice())?; + Ok(buf[backtest_idx]) + } + /// Read `snapshots_skipped_d`. Returns per-backtest count of snapshots skipped /// due to corrupt top-of-book (NaN/Inf price or zero/negative size at level 0). /// Used by tests to verify the skip logic fires; production code can sum for @@ -1170,6 +1210,10 @@ impl LobSimCuda { .arg(&self.books_d) // S2.1: mh_kernel_calls diagnostic counter. .arg(&mut self.mh_kernel_calls_d) + // CRT-A2: Wiener-α conviction-EMA state. + .arg(&mut self.conviction_ema_d) + .arg(&mut self.conviction_diff_var_ema_d) + .arg(&mut self.conviction_sample_var_ema_d) .arg(&n) .launch(cfg)?; } @@ -1195,6 +1239,10 @@ impl LobSimCuda { .arg(&self.books_d) // S2.1: mh_kernel_calls diagnostic counter. .arg(&mut self.mh_kernel_calls_d) + // CRT-A2: Wiener-α conviction-EMA state. + .arg(&mut self.conviction_ema_d) + .arg(&mut self.conviction_diff_var_ema_d) + .arg(&mut self.conviction_sample_var_ema_d) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index c138acee5..e8718e92b 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1234,3 +1234,161 @@ fn max_hold_counters_initialize_to_zero() -> Result<()> { assert_eq!(c.mh_force_flat_seen_by_seed[0], 0, "mh_force_flat_seen_by_seed must initialize to 0"); Ok(()) } + +/// CRT-A2: Wiener-α adaptive EMA on max-conviction-across-horizons. +/// +/// After A1 the controller fires every event. Without smoothing, sizing +/// magnitude would oscillate with the per-event alpha jitter (spread-bleed). +/// A2 adds a per-backtest Wiener-α EMA with floor 0.4 (non-stationary +/// control loop per pearl_wiener_alpha_floor_for_nonstationary) on the +/// max |alpha-0.5|*2 across horizons. Direction continues to come from +/// the raw alpha sign — only magnitude is smoothed. +/// +/// This test drives alternating high/low all-bullish alpha probs and +/// asserts: +/// (a) the on-device EMA is in fact populated (sentinel 0 → nonzero +/// after first observation); +/// (b) the EMA value sits strictly between the high and low input +/// convictions (smoothing produces a bounded mix, never extrapolates); +/// (c) under high enough kelly-floor to clear the integer-rounding bar, +/// direction (side) stays stable across alternation — no sign flips +/// on co-directional inputs. +fn cfg_high_kelly_floor(n: usize) -> BatchedSimConfig { + // kelly_frac_floor=1.0 keeps |signed_size| ≥ ~sig_mag * max_lots = 0.1*5 = 0.5 + // for the low-conv alpha (sig_mag ~ EMA), rounding to ±1 lots — so the + // *direction* of the target is stable across alternation regardless of + // sizing oscillation. + BatchedSimConfig::from_uniform(n, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 5, + latency_ns: 0, + kelly_frac_floor: 1.0, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, + max_hold_ns: 0, + min_reasonable_px: 0.0, + max_reasonable_px: f32::INFINITY, + }) +} + +#[test] +#[ignore = "requires CUDA"] +fn conviction_ema_smooths_micro_oscillations() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + let alphas_high: [f32; N_HORIZONS] = [0.8; N_HORIZONS]; // strong bullish, conv=0.6 + let alphas_low: [f32; N_HORIZONS] = [0.55; N_HORIZONS]; // weak bullish, conv=0.1 + let conv_high = 0.6_f32; + let conv_low = 0.1_f32; + + let cfg = cfg_high_kelly_floor(1); + + // Pre-condition: EMA slot starts at sentinel 0 (no observation yet). + assert_eq!(sim.read_conviction_ema(0)?, 0.0, + "conviction_ema must initialize to 0 (first-observation sentinel)"); + + // First decision: alphas_high. Bootstraps the EMA directly per + // pearl_first_observation_bootstrap (prev_ema == 0 → replace). + sim.broadcast_alpha(&alphas_high)?; + sim.step_decision_with_latency(1_000_000_000u64, &cfg)?; + let ema_after_first = sim.read_conviction_ema(0)?; + assert!((ema_after_first - conv_high).abs() < 1e-5, + "first observation must replace directly (bootstrap); got EMA={ema_after_first}, expected {conv_high}"); + + // Drive 10 alternations. Capture each EMA + target sign. + let mut prev_target_signed: Option = None; + let mut sign_flips = 0; + let mut ema_min = ema_after_first; + let mut ema_max = ema_after_first; + + for i in 0..10 { + let probs = if i % 2 == 0 { alphas_low } else { alphas_high }; + sim.broadcast_alpha(&probs)?; + let ts = 1_000_000_000u64 * ((i as u64) + 2); + sim.step_decision_with_latency(ts, &cfg)?; + let (side, size) = sim.read_market_target(0)?; + // side ∈ {0=buy, 1=sell, 2=noop, 3=force-flat}. For all-bullish + // alphas the raw direction is +; smoothed magnitude shouldn't flip + // it under the high-kelly-floor config. + let target_signed: i32 = match side { + 0 => size, + 1 => -size, + _ => 0, + }; + if let Some(p) = prev_target_signed { + if target_signed != 0 && p != 0 && (target_signed > 0) != (p > 0) { + sign_flips += 1; + } + } + prev_target_signed = Some(target_signed); + + let ema = sim.read_conviction_ema(0)?; + if ema < ema_min { ema_min = ema; } + if ema > ema_max { ema_max = ema; } + } + + // (b): EMA stays bounded between the two input convictions (post-bootstrap, + // the blend formula α·new + (1-α)·old is a strict convex combination of + // observed values, so the running EMA is bounded by the observed range). + assert!(ema_min > conv_low - 1e-4 && ema_max < conv_high + 1e-4, + "EMA must stay between conv_low={conv_low} and conv_high={conv_high}; got min={ema_min}, max={ema_max}"); + // The EMA should also have moved away from the bootstrap value (proves + // the second-and-onward update path executed, not just bootstrap). + assert!(ema_min < conv_high - 1e-4, + "EMA must drop below initial bootstrap conv_high={conv_high} after seeing alphas_low; got min={ema_min}"); + + // (c): direction stable under co-directional alternation. The raw alpha + // sign is positive in both cases (>0.5); smoothed magnitude shouldn't + // produce a sign flip. + assert_eq!(sign_flips, 0, + "direction must stay stable under EMA smoothing of co-directional alphas; flips={sign_flips}"); + + Ok(()) +} + +/// CRT-A2 corollary: when a true reversal arrives (alpha sign flips from +/// >0.5 to <0.5), the target direction must respond at the next event — +/// the smoothing is on magnitude only, not on sign. This is the property +/// that distinguishes Wiener-α smoothing on max-conviction from smoothing +/// the alpha probabilities themselves (which would lag reversals). +#[test] +#[ignore = "requires CUDA"] +fn conviction_ema_does_not_lag_reversals() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + let cfg = cfg_high_kelly_floor(1); + + // Bootstrap with strong bullish. + sim.broadcast_alpha(&[0.8; N_HORIZONS])?; + sim.step_decision_with_latency(1_000_000_000, &cfg)?; + let (side_buy, size_buy) = sim.read_market_target(0)?; + assert_eq!(side_buy, 0, "bootstrap bullish → buy; got side={side_buy} size={size_buy}"); + assert!(size_buy >= 1, "bootstrap bullish → size>=1; got {size_buy}"); + + // Hard reversal: strong bearish. Direction must flip on next event. + sim.broadcast_alpha(&[0.2; N_HORIZONS])?; + sim.step_decision_with_latency(2_000_000_000, &cfg)?; + let (side_sell, size_sell) = sim.read_market_target(0)?; + // The position may now be long (from the prior buy fill), but the + // *target* side written by the decision kernel reflects raw alpha sign. + // Stop checks only fire on open positions with SL/trail breach; with + // identical book + zero ATR they won't fire here, so the alpha-driven + // target dictates. + assert_eq!(side_sell, 1, + "post-reversal target must flip to sell; got side={side_sell} size={size_sell}"); + Ok(()) +}