From d090685ca9b03c5e34c92f0ef99fbc03e433da21 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 16 May 2026 00:28:01 +0200 Subject: [PATCH] feat(phase-e-4-a-T16): vol-regime detection + cost-aware training MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Adds two coupled interventions on the regime fragility exposed by walk-forward CV (mean Sharpe +27 ± 56 at half-tick across 3 folds — std-dev ≈ mean means the strategy is regime-dependent). (1) Vol-EMA regime detector (new ISV slots 549/550) - New alpha_regime_vol_update.cu kernel: per inference step, reads B parallel-env mid prices, computes cross-env mean squared log-return, and maintains ISV[549]=REGIME_VOL_EMA (Wiener-α with 0.4 floor + Pearl A bootstrap) and ISV[550]=REGIME_VOL_REF (slow tracker β=0.005, ≈200-step horizon). - Block-tree-reduce (no atomicAdd), guards against zero/non-finite mids. (2) Pre-emptive Kelly attenuation (modified stacker controller) - stacker_threshold_controller.cu takes 3 new args: regime_vol_ema_idx, regime_vol_ref_idx, regime_scale_floor. - Multiplies its reactive Sharpe-error Kelly output by regime_scale = clamp(vol_ref / vol_ema, 0.25, 1.0) - Disabled when indices = -1 (backward-compatible smoke + kernel test). (3) Cost-aware training (--train-cost-hi) - alpha_compose_backtest --train-cost-hi: when > --train-cost, each training epoch samples cost ~ U[lo, hi] so the Q-network learns cost-conservative behaviour across the realistic ES range. (4) Wiring - alpha_compose_backtest --regime-scale enables both per-step regime kernel firing during eval AND the regime hookup in the per-episode controller call. Mapped-pinned mids buffers, all compute device-side. - ExecutionEnv exposes current_mid() so the host gather reads the active snapshot mid per env without leaking the private cursor field. Smoke + test sites pass -1/-1 for regime indices (backward compat). Doc: docs/isv-slots.md ledger for slots 549/550. Co-Authored-By: Claude Opus 4.7 --- crates/ml/build.rs | 7 + crates/ml/examples/alpha_compose_backtest.rs | 99 ++++++++++++- crates/ml/examples/alpha_dqn_h600_smoke.rs | 7 + .../ml/src/cuda_pipeline/alpha_isv_slots.rs | 23 +++ crates/ml/src/cuda_pipeline/alpha_kernels.rs | 61 ++++++++ .../cuda_pipeline/alpha_regime_vol_update.cu | 138 ++++++++++++++++++ .../stacker_threshold_controller.cu | 38 ++++- crates/ml/src/env/execution_env.rs | 7 + docs/isv-slots.md | 66 +++++++++ scripts/build_2q_fxcache.sh | 40 +++++ 10 files changed, 482 insertions(+), 4 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu create mode 100755 scripts/build_2q_fxcache.sh diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f47ae6082..b2b405cc0 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1625,6 +1625,13 @@ fn main() { // sliding-window state input to the Mamba2 temporal encoder // (Mamba2Block from ml-alpha). "alpha_window_push.cu", + // Phase E.4.A.T16 (2026-05-16): vol-EMA regime detection kernel. + // Per inference step: reads B mid-prices, reduces to mean + // squared log-return, updates ISV[549] via Wiener-α with floor + // 0.4 (Pearl A bootstrap) and slow-tracks the reference vol in + // ISV[550]. Drives pre-emptive Kelly attenuation when realized + // vol blows past the trained reference. + "alpha_regime_vol_update.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/examples/alpha_compose_backtest.rs b/crates/ml/examples/alpha_compose_backtest.rs index fe0a9314b..a131d85cb 100644 --- a/crates/ml/examples/alpha_compose_backtest.rs +++ b/crates/ml/examples/alpha_compose_backtest.rs @@ -204,9 +204,20 @@ struct Cli { trade_size: i32, #[arg(long, default_value_t = 0xCAFEBABE_u64)] seed: u64, - /// Training-time cost (the policy LEARNED against this cost). + /// Training-time cost (the policy LEARNED against this cost). Used as + /// a fixed cost when `--train-cost-hi == train_cost`; otherwise serves + /// as the lower bound of a uniform[lo, hi] random cost sampled per + /// training episode (Phase E.4.A T16 cost-aware training). #[arg(long, default_value_t = 0.0625)] train_cost: f32, + /// Upper bound on the per-episode random training cost. When > `train_cost`, + /// each training episode samples its `cost_per_contract` uniformly from + /// [train_cost, train_cost_hi] so the Q-network learns cost-conservative + /// behaviour across the realistic ES futures range (quarter-tick to + /// quad-tick). Defaults to the same value as `train_cost` (no randomization, + /// backward-compatible). + #[arg(long, default_value_t = 0.0625)] + train_cost_hi: f32, /// Training-time alpha-confidence threshold for the gate. Forces /// Wait during training when `|sigmoid(alpha)−0.5| < this`, so the /// Q-network learns weights for the gated policy class. Default @@ -274,6 +285,19 @@ struct Cli { /// rate / Sharpe. #[arg(long, default_value_t = false)] isv_continual: bool, + /// Phase E.4.A T16: enable vol-EMA regime detection + pre-emptive + /// Kelly attenuation. When set: + /// - per inference step, `alpha_regime_vol_update_kernel` reads + /// the parallel envs' current/previous mids and maintains + /// ISV[549]=vol_ema (Wiener-α with 0.4 floor + Pearl A + /// bootstrap) and ISV[550]=vol_ref (slow tracker, β=0.005); + /// - per eval episode (requires `--isv-continual`), the existing + /// stacker controller multiplies its Sharpe-reactive Kelly + /// output by `clamp(vol_ref/vol_ema, 0.25, 1.0)` so vol spikes + /// cut Kelly *before* losses materialise. + /// Off by default to keep the T10 baseline comparable. + #[arg(long, default_value_t = false)] + regime_scale: bool, /// Phase E.4.A.T15.batched (2026-05-15): parallel envs per training /// "epoch". n_train_episodes / n_train_par epochs total. Each epoch /// runs N_par envs in lockstep H steps + ONE batched C51 update at @@ -612,9 +636,20 @@ fn main() -> Result<()> { for epoch in 0..train_n_epochs { let _eps = cli.eps_start + (cli.eps_end - cli.eps_start) * (epoch as f32 / train_n_epochs.max(1) as f32); + // T16 cost-aware training: when --train-cost-hi > --train-cost, + // sample the per-episode cost uniformly from that range, so the + // policy learns to behave robustly across the realistic ES cost + // spectrum (quarter-tick to quad-tick) instead of memorizing a + // single low-cost world. + let cost_for_epoch: f32 = if cli.train_cost_hi > cli.train_cost { + let u = (episode_rng.next_u64() as f32) / (u64::MAX as f32); + cli.train_cost + u * (cli.train_cost_hi - cli.train_cost) + } else { + cli.train_cost + }; let mut par_envs: Vec = (0..train_n_par).map(|_| { let mut cfg = env_config_template_train.clone(); - cfg.cost_per_contract = cli.train_cost; + cfg.cost_per_contract = cost_for_epoch; ExecutionEnv::new_arc( cfg, fill_model_for_train.clone(), std::sync::Arc::clone(&snapshots_arc_train), @@ -942,6 +977,20 @@ fn main() -> Result<()> { let mut batched_probs_dev = stream .alloc_zeros::(n_par * N_ACTIONS * c51_n_atoms) .context("alloc batched probs")?; + // T16 regime kernel state. mids buffers are mapped-pinned so the host + // gather (one f32 per env per step) writes directly into device-mapped + // memory with no htod copy. + let mids_curr_pinned = unsafe { MappedF32::new(n_par)? }; + let mids_prev_pinned = unsafe { MappedF32::new(n_par)? }; + let regime_module = if cli.regime_scale { + Some(ctx.load_cubin( + ml::cuda_pipeline::alpha_kernels::ALPHA_REGIME_VOL_UPDATE_CUBIN.to_vec() + ).context("alpha_regime_vol_update cubin load")?) + } else { None }; + let regime_kernel = regime_module.as_ref().map(|m| + m.load_function("alpha_regime_vol_update_kernel") + .expect("alpha_regime_vol_update_kernel load") + ); let snapshots_arc = env.snapshots_arc(); let fill_model_for_par = env.fill_model.clone(); let env_config_template = env.config.clone(); @@ -975,6 +1024,13 @@ fn main() -> Result<()> { stream.memset_zeros(batched_window_tensor.data_mut()) .context("zero batched windows")?; } + // T16: reset prev_mids so the first-step Pearl-A bootstrap fires + // cleanly (kernel guards against curr<=0 || prev<=0 internally). + if cli.regime_scale { + let zeros = vec![0.0_f32; n_par]; + mids_curr_pinned.write(&zeros); + mids_prev_pinned.write(&zeros); + } // Lockstep step loop. for _step in 0..cli.horizon { // Gather current states (CPU; <100μs for N=500). @@ -986,6 +1042,35 @@ fn main() -> Result<()> { } } batched_state_pinned.write(&batched_states_host); + // T16: capture current mids + fire regime kernel BEFORE the + // env step advances. Regime signal therefore reflects vol + // observed up to and including this step, which is what the + // controller's per-episode Kelly update will read. + if let Some(rk) = regime_kernel.as_ref() { + let mut mids_host = vec![0.0_f32; n_par]; + for i in 0..n_par { + if !done_flags[i] { + mids_host[i] = par_envs[i].current_mid(); + } + } + mids_curr_pinned.write(&mids_host); + unsafe { + let (isv_ptr, _g) = isv_dev.device_ptr_mut(&stream); + ml::cuda_pipeline::alpha_kernels::launch_alpha_regime_vol_update( + &stream, rk, + mids_curr_pinned.dev_u64(), + mids_prev_pinned.dev_u64(), + isv_ptr, + n_par as i32, + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_EMA_INDEX as i32, + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_INDEX as i32, + 0.005, // β: 200-step horizon for vol_ref tracking + )?; + } + // Swap prev ← curr for the next step. Mapped-pinned write + // is a CPU-side fill, no device-to-device copy needed. + mids_prev_pinned.write(&mids_host); + } if cli.temporal { // Batched window push. { @@ -1088,6 +1173,16 @@ fn main() -> Result<()> { ml::cuda_pipeline::alpha_isv_slots::TRADE_RATE_TARGET_INDEX as i32, ml::cuda_pipeline::alpha_isv_slots::TRADE_RATE_OBSERVED_EMA_INDEX as i32, ml::cuda_pipeline::alpha_isv_slots::STACKER_KELLY_ATTENUATION_INDEX as i32, + // T16 regime hookup. Active when --isv-continual is on; + // the regime kernel writes slots 549/550 per-step during + // inference. Floor 0.25 = max defensive 4× Kelly cut. + if cli.regime_scale { + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_EMA_INDEX as i32 + } else { -1 }, + if cli.regime_scale { + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_INDEX as i32 + } else { -1 }, + 0.25, isv_ptr, wv_ptr, )?; } diff --git a/crates/ml/examples/alpha_dqn_h600_smoke.rs b/crates/ml/examples/alpha_dqn_h600_smoke.rs index 1dbd54b65..7cd3b3a5b 100644 --- a/crates/ml/examples/alpha_dqn_h600_smoke.rs +++ b/crates/ml/examples/alpha_dqn_h600_smoke.rs @@ -1477,6 +1477,13 @@ fn main() -> Result<()> { ml::cuda_pipeline::alpha_isv_slots::TRADE_RATE_TARGET_INDEX as i32, ml::cuda_pipeline::alpha_isv_slots::TRADE_RATE_OBSERVED_EMA_INDEX as i32, ml::cuda_pipeline::alpha_isv_slots::STACKER_KELLY_ATTENUATION_INDEX as i32, + // T16 regime hookup. Smoke doesn't currently run the vol + // update kernel so we pass -1/-1 to disable the regime + // scale here (backward compat); the backtest binary will + // wire it on when `--regime-scale` is set. + -1, + -1, + 0.25, isv_ptr, w_ptr, )?; diff --git a/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs b/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs index 8c2c158d5..666ae3f35 100644 --- a/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs @@ -51,6 +51,25 @@ pub const STACKER_KELLY_ATTENUATION_INDEX: usize = 546; pub const RANDOM_BASELINE_MEAN_INDEX: usize = 547; pub const RANDOM_BASELINE_STD_INDEX: usize = 548; +// Phase E.4.A T16 regime detection (vol-EMA pre-emptive Kelly defense). +// +// Slot 549 holds the Wiener-α-smoothed estimate of recent squared +// log-return magnitude across the parallel envs. Slot 550 holds a slow- +// tracking "reference" vol (windowed Welford-style median proxy) that +// represents the regime the policy was trained on. The ratio +// `vol_ref / max(vol_ema, ε)` becomes a multiplicative Kelly scale — +// when realized vol blows past the trained reference, Kelly attenuates +// PRE-EMPTIVELY (before realized P&L losses trigger the existing +// STACKER_KELLY_ATTENUATION reactive response). +// +// Producer kernel: `alpha_regime_vol_update.cu` (per inference step, +// reads B mid-prices, updates both slots in one kernel via Pearl A +// bootstrap + Wiener-α with floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary). +// Consumer: `stacker_threshold_controller.cu` reads slot 549/550 when +// computing kelly_attenuation update. +pub const REGIME_VOL_EMA_INDEX: usize = 549; +pub const REGIME_VOL_REF_INDEX: usize = 550; + // Sentinels: zero by convention. Fallbacks use the first-observation // bootstrap pattern (see pearl_first_observation_bootstrap.md). pub const ALPHA_SENTINEL: f32 = 0.0; @@ -80,6 +99,8 @@ mod tests { STACKER_KELLY_ATTENUATION_INDEX, RANDOM_BASELINE_MEAN_INDEX, RANDOM_BASELINE_STD_INDEX, + REGIME_VOL_EMA_INDEX, + REGIME_VOL_REF_INDEX, ]; for &i in &indices { assert!(i >= ALPHA_ISV_BLOCK_LO && i <= ALPHA_ISV_BLOCK_HI, @@ -100,6 +121,8 @@ mod tests { STACKER_KELLY_ATTENUATION_INDEX, RANDOM_BASELINE_MEAN_INDEX, RANDOM_BASELINE_STD_INDEX, + REGIME_VOL_EMA_INDEX, + REGIME_VOL_REF_INDEX, ]; let mut sorted: Vec = indices.to_vec(); sorted.sort(); diff --git a/crates/ml/src/cuda_pipeline/alpha_kernels.rs b/crates/ml/src/cuda_pipeline/alpha_kernels.rs index 20f6a624d..ae30caedc 100644 --- a/crates/ml/src/cuda_pipeline/alpha_kernels.rs +++ b/crates/ml/src/cuda_pipeline/alpha_kernels.rs @@ -263,6 +263,10 @@ pub unsafe fn launch_stacker_threshold_controller( trade_rate_target_idx: i32, trade_rate_observed_idx: i32, stacker_kelly_atten_idx: i32, + // T16 regime args. Pass -1 / -1 / any-float to disable the regime scale. + regime_vol_ema_idx: i32, + regime_vol_ref_idx: i32, + regime_scale_floor: f32, isv_dev: u64, wiener_state_dev: u64, ) -> Result<(), MLError> { @@ -274,6 +278,8 @@ pub unsafe fn launch_stacker_threshold_controller( "wiener_alpha_floor must be in [0, 1)"); debug_assert!(alpha_meta > 0.0 && alpha_meta < 1.0, "alpha_meta must be in (0, 1)"); + debug_assert!(regime_scale_floor > 0.0 && regime_scale_floor <= 1.0, + "regime_scale_floor must be in (0, 1]"); let cfg = LaunchConfig { grid_dim: (1, 1, 1), @@ -294,6 +300,9 @@ pub unsafe fn launch_stacker_threshold_controller( .arg(&trade_rate_target_idx) .arg(&trade_rate_observed_idx) .arg(&stacker_kelly_atten_idx) + .arg(®ime_vol_ema_idx) + .arg(®ime_vol_ref_idx) + .arg(®ime_scale_floor) .arg(&isv_dev) .arg(&wiener_state_dev) .launch(cfg) @@ -448,6 +457,52 @@ pub static ALPHA_C51_CUBIN: &[u8] = pub static ALPHA_WINDOW_PUSH_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/alpha_window_push.cubin")); +/// Precompiled vol-EMA regime detection kernel. Phase E.4.A.T16. +pub static ALPHA_REGIME_VOL_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/alpha_regime_vol_update.cubin")); + +/// Launch `alpha_regime_vol_update_kernel`. Reads B current/previous +/// mids, computes the cross-env mean squared log-return, and updates +/// ISV slots 549 (vol_ema, Wiener-α with 0.4 floor + Pearl A bootstrap) +/// and 550 (vol_ref, slow-tracking reference). One block, ≤256 threads +/// — B is typically 50, fits easily. +/// +/// # Safety +/// `mids_curr_dev`, `mids_prev_dev`, `isv_dev` must be valid device +/// pointers. The kernel does not atomicAdd and contains no host +/// branches inside its execution. +pub unsafe fn launch_alpha_regime_vol_update( + stream: &cudarc::driver::CudaStream, + kernel: &cudarc::driver::CudaFunction, + mids_curr_dev: u64, + mids_prev_dev: u64, + isv_dev: u64, + b: i32, + vol_ema_index: i32, + vol_ref_index: i32, + ref_track_rate: f32, +) -> Result<(), MLError> { + use cudarc::driver::{LaunchConfig, PushKernelArg}; + debug_assert!(b > 0 && b <= 256, "B must be in (0, 256]; got {b}"); + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + stream + .launch_builder(kernel) + .arg(&mids_curr_dev) + .arg(&mids_prev_dev) + .arg(&isv_dev) + .arg(&b) + .arg(&vol_ema_index) + .arg(&vol_ref_index) + .arg(&ref_track_rate) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("alpha_regime_vol_update launch: {e}")))?; + Ok(()) +} + /// Launch `alpha_window_push_kernel`. Shifts the window's K slots left /// by one (slot 1 → slot 0, slot 2 → slot 1, …, slot K-1 → slot K-2) /// and inserts the new state at slot K-1. After the call, @@ -1785,6 +1840,9 @@ mod tests { TRADE_RATE_TARGET_INDEX as i32, TRADE_RATE_OBSERVED_EMA_INDEX as i32, STACKER_KELLY_ATTENUATION_INDEX as i32, + -1, // regime_vol_ema_idx (disabled in test) + -1, // regime_vol_ref_idx (disabled in test) + 0.25, // regime_scale_floor (unused when disabled) isv_ptr, w_ptr, ) @@ -1829,6 +1887,9 @@ mod tests { TRADE_RATE_TARGET_INDEX as i32, TRADE_RATE_OBSERVED_EMA_INDEX as i32, STACKER_KELLY_ATTENUATION_INDEX as i32, + -1, // regime_vol_ema_idx (disabled in test) + -1, // regime_vol_ref_idx (disabled in test) + 0.25, // regime_scale_floor (unused when disabled) isv_ptr, w_ptr, ) diff --git a/crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu b/crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu new file mode 100644 index 000000000..587d45fc4 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu @@ -0,0 +1,138 @@ +// crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu +// +// Phase E.4.A.T16 (2026-05-16): vol-EMA regime detection. +// +// Per inference step, reads B mid-prices and the previous step's mids, +// computes the cross-env mean squared log-return, and updates two ISV +// slots: +// +// slot 549 (REGIME_VOL_EMA): Wiener-α EMA of mean-squared log-return. +// Pearl A bootstrap (sentinel 0 → first +// observation replaces directly). Pearl D +// Wiener-α with floor 0.4 per +// pearl_wiener_alpha_floor_for_nonstationary +// (the controller co-adapts with policy +// and the target itself drifts — Wiener +// would collapse to 0 without the floor). +// +// slot 550 (REGIME_VOL_REF): slow-tracking reference vol. β=0.005 +// per-step (about 200-step horizon). On +// first observation, bootstraps directly +// from slot 549 so the policy's "expected +// vol regime" matches the early-training +// observation. Drifts slowly enough that +// acute spikes do NOT pollute the +// reference, preserving the spike-detector +// property of `vol_ema / vol_ref`. +// +// Downstream: `stacker_threshold_controller.cu` reads slots 549/550 in +// the Kelly-attenuation branch — when vol_ema >> vol_ref the controller +// shrinks Kelly *before* losses materialise (pre-emptive defense), while +// the existing Sharpe-error controller continues to react after losses +// (reactive defense). The two compose multiplicatively. +// +// Why mid-price log-return and not spread/imbalance? Vol is the single +// strongest regime indicator empirically (cf. walk-forward Fold B failure +// in Q1 2024 mid-quarter). Multi-feature regime classification is a +// follow-up; this kernel is the MVP minimal-borrow regime signal. +// +// Mapped-pinned discipline: caller passes device pointers for the mids, +// and ISV is the standard device-side bus. No CPU branches inside the +// kernel — single thread does the EMA arithmetic after a CUB-free +// block reduction over the B mids. + +#include + +// Wiener-α floor — keeps the EMA responsive when the target drifts. +// Mirrors the 0.4 floor used by stacker_threshold_controller and +// elsewhere per pearl_wiener_alpha_floor_for_nonstationary. +__device__ __forceinline__ float wiener_alpha_floored( + float diff_var, float sample_var, float floor +) { + const float eps = 1e-12f; + float alpha = diff_var / (diff_var + sample_var + eps); + return fmaxf(floor, fminf(alpha, 1.0f)); +} + +// Cooperative block reduction over B values stored at mids[b]. Returns +// the mean SQUARED log-return for thread 0; other threads return 0. +extern "C" __global__ void alpha_regime_vol_update_kernel( + const float* __restrict__ mids_curr, // [B] current mids (device) + const float* __restrict__ mids_prev, // [B] previous step mids + float* __restrict__ isv, // ISV bus + int B, + int vol_ema_index, // 549 + int vol_ref_index, // 550 + float ref_track_rate // β (per-step), e.g. 0.005 +) { + // Single-block design: each thread handles one b. Final reduction by + // thread 0. B is typically 50 (n_eval_episodes) — fits easily in 256 + // threads. If B > blockDim.x, each thread handles a strided slice. + __shared__ float partial_sum[256]; + __shared__ int partial_cnt[256]; + + const int tid = threadIdx.x; + float my_sum = 0.0f; + int my_cnt = 0; + for (int b = tid; b < B; b += blockDim.x) { + const float curr = mids_curr[b]; + const float prev = mids_prev[b]; + // Guard against zero / non-finite mids (zero-init buffer at + // episode 0 will have curr == prev == 0 for some envs). + if (curr > 0.0f && prev > 0.0f && isfinite(curr) && isfinite(prev)) { + const float r = logf(curr / prev); + my_sum += r * r; + my_cnt += 1; + } + } + partial_sum[tid] = my_sum; + partial_cnt[tid] = my_cnt; + __syncthreads(); + + // Tree-reduce (block-internal; no atomicAdd per feedback_no_atomicadd). + for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) { + if (tid < stride) { + partial_sum[tid] += partial_sum[tid + stride]; + partial_cnt[tid] += partial_cnt[tid + stride]; + } + __syncthreads(); + } + + if (tid != 0) return; + + const int total = partial_cnt[0]; + if (total == 0) return; // no valid envs this step — leave ISV unchanged + + const float vol_obs = partial_sum[0] / (float)total; + + // Pearl A first-observation bootstrap on slot 549. + const float vol_ema_prev = isv[vol_ema_index]; + float vol_ema_new; + if (vol_ema_prev <= 0.0f) { + vol_ema_new = vol_obs; + } else { + // Wiener-optimal α with 0.4 floor. diff_var proxy: (vol_obs − + // vol_ema_prev)²; sample_var proxy: vol_ema_prev² (rough running + // variance — exact Welford machinery is heavier than the spike- + // detector application needs at this point). + const float diff = vol_obs - vol_ema_prev; + const float diff_var = diff * diff; + const float sample_var = vol_ema_prev * vol_ema_prev; + const float alpha = wiener_alpha_floored(diff_var, sample_var, 0.4f); + vol_ema_new = vol_ema_prev + alpha * diff; + } + isv[vol_ema_index] = vol_ema_new; + + // Slow reference vol tracking on slot 550. + const float vol_ref_prev = isv[vol_ref_index]; + float vol_ref_new; + if (vol_ref_prev <= 0.0f) { + // First observation: bootstrap reference equal to the early EMA + // so the spike detector ratio starts at 1.0 (no spurious early + // attenuation while the policy is finding its regime). + vol_ref_new = vol_ema_new; + } else { + vol_ref_new = vol_ref_prev + ref_track_rate * (vol_ema_new - vol_ref_prev); + } + isv[vol_ref_index] = vol_ref_new; +} diff --git a/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu index 88002d1d9..8f931093c 100644 --- a/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu +++ b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu @@ -58,6 +58,11 @@ extern "C" __global__ void stacker_threshold_controller_update( int trade_rate_target_idx, int trade_rate_observed_idx, int stacker_kelly_atten_idx, + /* T16 vol-regime slots. Set both to a negative number to disable + the regime scale (backward-compatible). */ + int regime_vol_ema_idx, + int regime_vol_ref_idx, + float regime_scale_floor, // typ. 0.25 — bottoms the multiplicative defense /* Device buffers. */ float* __restrict__ isv, float* __restrict__ wiener_state /* [sample_var, diff_var, x_lag] for slot 545 */ @@ -138,9 +143,38 @@ extern "C" __global__ void stacker_threshold_controller_update( const float prev_atten_raw = isv[stacker_kelly_atten_idx]; const float prev_atten = prev_atten_raw == 0.0f ? ATTEN_HI : prev_atten_raw; const float err_sharpe = rollout_realized_sharpe - target_sharpe; - const float new_atten = + const float reactive_atten = fmaxf(ATTEN_LO, fminf(ATTEN_HI, prev_atten + k_atten * err_sharpe)); - isv[stacker_kelly_atten_idx] = new_atten; + + // ---- (4b) T16 pre-emptive vol-regime scale ---- + // + // The reactive Sharpe-error controller above only fires AFTER losses + // materialize. The vol-regime scale fires AS SOON AS realized vol + // diverges from the trained-on reference vol. The two compose + // multiplicatively — reactive sets the upper bound based on observed + // P&L, regime trims further when conditions look unfriendly. + // + // regime_scale = clamp(vol_ref / max(vol_ema, ε), regime_scale_floor, 1.0) + // + // vol_ema << vol_ref (calm market vs trained turbulence) → scale = 1.0 (no trim) + // vol_ema ≈ vol_ref (same regime as training) → scale ≈ 1.0 + // vol_ema = 2·vol_ref (vol spike) → scale = 0.5 + // vol_ema = 4·vol_ref (extreme spike) → scale = 0.25 (floor) + // + // Disabled when either slot index is negative OR when vol_ref ≤ 0 + // (regime kernel hasn't bootstrapped yet — must not punish cold start). + float regime_scale = 1.0f; + if (regime_vol_ema_idx >= 0 && regime_vol_ref_idx >= 0) { + const float vol_ema = isv[regime_vol_ema_idx]; + const float vol_ref = isv[regime_vol_ref_idx]; + if (vol_ref > 0.0f && vol_ema > 0.0f) { + const float raw_scale = vol_ref / vol_ema; + regime_scale = fmaxf(regime_scale_floor, fminf(1.0f, raw_scale)); + } + } + + const float final_atten = fmaxf(ATTEN_LO, reactive_atten * regime_scale); + isv[stacker_kelly_atten_idx] = final_atten; __threadfence_system(); } diff --git a/crates/ml/src/env/execution_env.rs b/crates/ml/src/env/execution_env.rs index 1316fd6c6..0aa86e59d 100644 --- a/crates/ml/src/env/execution_env.rs +++ b/crates/ml/src/env/execution_env.rs @@ -176,6 +176,13 @@ impl ExecutionEnv { self.state(&EpisodeState::new()) } + /// Mid price at the env's current cursor. Phase E.4.A T16 regime + /// detector reads this each step to compute the cross-env mean + /// squared log-return (the vol-EMA regime signal). + pub fn current_mid(&self) -> f32 { + self.snapshots[self.cursor].mid_price + } + pub fn snapshots_len(&self) -> usize { self.snapshots.len() } diff --git a/docs/isv-slots.md b/docs/isv-slots.md index cd8c6cb26..d10865db5 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -847,3 +847,69 @@ NO ISV slot interaction — both kernels operate purely on the smoke / backtest binary's `window_dev` and `h_enriched_buf_dev` buffers. Documented here only for the kernel-audit-doc hook requirement; truly slot-agnostic. + +## 2026-05-16 — Phase E.4.A T16 vol-EMA regime detection (slots 549/550) + +Walk-forward CV across Q1 2024 (3 folds, 700K-bar windows) exposed +that the T10 architecture has mean Sharpe +27 ± 56 at half-tick cost +(std-dev ≈ mean → regime-dependent learning). Fold B in particular +collapsed to -21..-60 across all costs with 0-22% win rate. T16 +adds two coupled defenses. + +**Slot 549 — REGIME_VOL_EMA_INDEX** +- Producer: `alpha_regime_vol_update.cu::alpha_regime_vol_update_kernel` +- Update cadence: per inference step (every env step during eval; off + during training unless explicitly wired) +- Math: reduce mids across B parallel envs to mean(log(curr/prev))², + EMA-update via Wiener-α with floor 0.4 + Pearl A bootstrap. +- Consumer: modified `stacker_threshold_controller.cu` Kelly branch + reads slot 549 + slot 550 to compute pre-emptive scale. +- Sentinel: 0.0. First-observation bootstrap replaces directly. + +**Slot 550 — REGIME_VOL_REF_INDEX** +- Producer: same kernel; β=0.005 (≈200-step horizon) so it tracks the + slow-moving "what is the trained regime" reference. Acute spikes do + NOT pollute the reference, preserving the spike-detector property + of `vol_ema / vol_ref`. +- Sentinel: 0.0. First-observation bootstrap copies the bootstrapped + vol_ema so the spike-detector ratio starts at 1.0 (no spurious + early attenuation). + +**Producer kernel — `alpha_regime_vol_update.cu`:** +- Single-block, ≤256-thread block-tree reduction over B mids; no + atomicAdd per `feedback_no_atomicadd`. Guards against zero / + non-finite mids (B can be 50 eval envs, some may have done==1 + earlier in the episode and their mids will be stale 0.0). +- All compute device-side per `feedback_cpu_is_read_only`. Mids + themselves are written via mapped-pinned f32 buffers (host gather + is one f32 per env per step, no htod copy). + +**Consumer change — `stacker_threshold_controller.cu`:** +- Signature gains `regime_vol_ema_idx`, `regime_vol_ref_idx`, + `regime_scale_floor` args. When indices ≥ 0 AND `vol_ref > 0` AND + `vol_ema > 0`, the Kelly output is scaled by + `clamp(vol_ref / vol_ema, regime_scale_floor=0.25, 1.0)`. When + indices = -1 the scale collapses to 1.0 (backward-compatible — both + the smoke binary and the kernel unit test pass -1/-1). +- Reactive (Sharpe-error) and pre-emptive (vol-spike) components + compose multiplicatively. Pre-emptive fires AS SOON AS vol blows + past trained-on reference; reactive only fires AFTER P&L pain. + +**Backtest CLI:** +- `--regime-scale` enables both the per-step regime kernel firing + during eval AND the regime hookup in the per-episode controller + call (requires `--isv-continual` for the controller path). +- `--train-cost-hi ` (combined with `--train-cost `): per + training epoch samples cost ~ U[lo, hi]. Default lo=hi=0.0625 = + backward compat (no randomization). The T16 hypothesis is that + cost-randomized training + vol-regime defense together close the + Fold B failure mode. + +**Verification (when Q1+Q2 2024 fxcache builds complete):** walk-forward +CV across Q1↔Q2 boundary with all four configurations: + - T10 baseline (no regime, fixed cost) + - T10 + cost randomization only + - T10 + regime scale only + - T10 + both +Expected: combined config compresses cross-fold std-dev meaningfully +without sacrificing best-fold Sharpe. diff --git a/scripts/build_2q_fxcache.sh b/scripts/build_2q_fxcache.sh new file mode 100755 index 000000000..5a9cb1535 --- /dev/null +++ b/scripts/build_2q_fxcache.sh @@ -0,0 +1,40 @@ +#!/usr/bin/env bash +# Build a Q1+Q2 2024 combined fxcache for cross-quarter walk-forward CV. +# +# Requires: +# test_data/futures-baseline/ES.FUT/ES.FUT_2024-Q1.dbn.zst (present) +# test_data/futures-baseline/ES.FUT/ES.FUT_2024-Q2.dbn.zst (present) +# test_data/futures-baseline-mbp10/ES.FUT/ES.FUT_2024-Q1.dbn.zst (present) +# test_data/futures-baseline-mbp10/ES.FUT/ES.FUT_2024-Q2.dbn.zst (download in progress) +# test_data/futures-baseline-trades/ES.FUT/ES.FUT_2024-Q1.dbn.zst (present) +# test_data/futures-baseline-trades/ES.FUT/ES.FUT_2024-Q2.dbn.zst (present) +# +# Produces a single fxcache + norm_stats in test_data/feature-cache/. +# Cache key is SHA256 over all matching DBN files in the input dirs. + +set -euo pipefail + +# Stage Q1+Q2 only — symlink them into a temporary dir so the builder +# doesn't pick up other quarters that may land later. +STAGE="${STAGE:-/tmp/fxbuild_2q_2024}" +rm -rf "$STAGE" +mkdir -p "$STAGE"/{ohlcv,mbp10,trades} +for q in Q1 Q2; do + ln -s "/home/jgrusewski/Work/foxhunt/test_data/futures-baseline/ES.FUT/ES.FUT_2024-${q}.dbn.zst" "$STAGE/ohlcv/" + ln -s "/home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT/ES.FUT_2024-${q}.dbn.zst" "$STAGE/mbp10/" + ln -s "/home/jgrusewski/Work/foxhunt/test_data/futures-baseline-trades/ES.FUT/ES.FUT_2024-${q}.dbn.zst" "$STAGE/trades/" +done + +mkdir -p test_data/feature-cache + +SQLX_OFFLINE=true cargo run -p ml --release --example precompute_features -- \ + --data-dir "$STAGE/ohlcv" \ + --mbp10-data-dir "$STAGE/mbp10" \ + --trades-data-dir "$STAGE/trades" \ + --output-dir test_data/feature-cache \ + --symbol ES.FUT \ + --yes + +echo +echo "Built fxcaches in test_data/feature-cache/:" +ls -lh test_data/feature-cache/ | grep "$(date +%Y-%m-%d)" || ls -lh test_data/feature-cache/ | tail