feat(phase-e-4-a-T16): vol-regime detection + cost-aware training

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 <noreply@anthropic.com>

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)

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<ExecutionEnv> = (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::<f32>(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,
                 )?;
             }

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,
             )?;

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<usize> = indices.to_vec();
         sorted.sort();

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(&regime_vol_ema_idx)
+        .arg(&regime_vol_ref_idx)
+        .arg(&regime_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,
                 )

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 <cuda_runtime.h>
+
+// 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;
+}

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

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

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 <hi>` (combined with `--train-cost <lo>`): 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.

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