feat(regime): vol_ref bootstrap window + ISV-driven permanent floor

Two coupled fixes to the vol-EMA regime detector exposed by walk-forward
CV after all features were promoted to always-on:

(1) Bootstrap window for vol_ref (slot 551 = REGIME_VOL_REF_SAMPLES)
- Replace the Pearl-A "first observation replaces directly" bootstrap
  with a running mean over the first N=100 vol_ema observations, then
  switch to β-tracking. For IID observations the running-mean estimator
  has variance σ²/N — a 100-sample mean is 10× less noisy than the
  single-shot replace.

(2) Permanent floor on vol_ref (slot 552 = REGIME_VOL_REF_FLOOR)
- The bootstrap alone exposed the asymmetric deadband-deadlock: if
  vol_ref converged to a tiny value during a calm initial stretch,
  vol_ref / vol_ema fired the moment any realistic vol resumed and
  Kelly stayed trapped at regime_scale_floor=0.25 forever. Floor
  lives in ISV slot (TrainingPersist) with a hardcoded 1e-12 sub-floor
  inside the kernel as numerical-underflow guard.
- Host seeds slot 552 with 1e-9. Future controller kernel will refine
  this from observed cell-level vol minima with cross-fold persistence.

Walk-forward CV on Q1 fxcache (3 folds, window=700K, train_frac=0.6):
  cost   fold-A   fold-B          fold-C    mean ± SD
  0.00   -19.77   +65.04 (100%)   +6.45     +17.24 ± 43.42

Fold B turnaround is the headline: -47.64 (bootstrap-only) → +65.04
(bootstrap + floor) confirms the floor is the load-bearing fix.
Cross-fold std-dev compressed 24% at cost=0; mean dropped from +38.94
(pre-defense) to +17.24 (with defense). Classic mean/variance trade.

Block extended to 14 slots (539..=552). Kernel sig: vol_ref_floor
moved from f32 scalar to vol_ref_floor_index i32, so the anchor is
named/addressable in ISV.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

crates/ml/examples/alpha_baseline.rs

@@ -471,10 +471,17 @@ fn main() -> Result<()> {
         .context("train_window_store kernel load")?;
     let n_atoms_i = c51_n_atoms as i32;
     // ISV buffer + Wiener state for the eval-time controller path.
-    let mut isv_host: Vec<f32> = vec![0.0; 552];
+    let mut isv_host: Vec<f32> = vec![0.0; 553];
     isv_host[ml::cuda_pipeline::alpha_isv_slots::RANDOM_BASELINE_MEAN_INDEX] = -5191.53;
     isv_host[ml::cuda_pipeline::alpha_isv_slots::RANDOM_BASELINE_STD_INDEX] = 4963.62;
     isv_host[ml::cuda_pipeline::alpha_isv_slots::TRADE_RATE_TARGET_INDEX] = 0.08;
+    // T16 vol_ref floor seed (TrainingPersist anchor). Initial guess based
+    // on ES MBP-10 squared-log-return range (typical vol_obs ~1e-10..1e-6,
+    // so a floor of 1e-9 sits an order of magnitude below typical realized
+    // vol). A future controller kernel can refine this from observed
+    // bootstrap statistics; the hardcoded sub-floor 1e-12 inside the
+    // regime kernel still catches the slot collapsing to zero.
+    isv_host[ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_FLOOR_INDEX] = 1.0e-9;
     let mut isv_dev = stream.clone_htod(&isv_host).context("upload isv")?;
     let mut ctl_wiener_dev = stream.alloc_zeros::<f32>(3).context("alloc ctl wiener")?;
 
@@ -903,7 +910,10 @@ fn main() -> Result<()> {
                         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
+                        ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_SAMPLES_INDEX as i32,
+                        ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_FLOOR_INDEX as i32,
+                        100,    // bootstrap samples: running mean over first 100 obs
+                        0.005,  // β: 200-step horizon for vol_ref tracking after bootstrap
                     )?;
                 }
                 // Swap prev ← curr for the next step. Mapped-pinned write

crates/ml/src/cuda_pipeline/alpha_isv_slots.rs

@@ -1,6 +1,6 @@
 //! Alpha trading system ISV slot block.
 //!
-//! Reserves indices 539..=550 (12 slots, contiguous) for durable
+//! Reserves indices 539..=552 (14 slots, contiguous) for durable
 //! infrastructure of the alpha trading system: stacker/execution-policy
 //! diagnostics, controller anchors, and population baselines. Unlike the
 //! SP4..SP22 slot files (which were per-iteration milestone-scoped), this
@@ -24,16 +24,19 @@
 //! | 546   | STACKER_KELLY_ATTENUATION_IX | Controller   | per-trade-close | stacker_threshold_controller.cu          |
 //! | 547   | RANDOM_BASELINE_MEAN_INDEX   | Anchor       | per-fold        | host computes once per fold              |
 //! | 548   | RANDOM_BASELINE_STD_INDEX    | Anchor       | per-fold        | host computes once per fold              |
-//! | 549   | (reserved spare)             | —            | —               | —                                        |
-//! | 550   | (reserved spare)             | —            | —               | —                                        |
+//! | 549   | REGIME_VOL_EMA_INDEX         | Diagnostic   | per-inference   | alpha_regime_vol_update.cu             |
+//! | 550   | REGIME_VOL_REF_INDEX         | Controller   | per-inference   | alpha_regime_vol_update.cu             |
+//! | 551   | REGIME_VOL_REF_SAMPLES_INDEX | Counter      | per-inference   | alpha_regime_vol_update.cu             |
+//! | 552   | REGIME_VOL_REF_FLOOR_INDEX   | Anchor       | TrainingPersist | host seed + future controller-kernel   |
 //!
-//! Reservation discipline (per the alpha-system ISV research memo): slots 549..550
-//! are intentionally unused at first-commit to absorb growth without bumping
-//! `ISV_TOTAL_DIM` again. Add new alpha-system slots HERE first, not by extending
-//! past 550 ad-hoc.
+//! Reservation discipline: this block had two spares at first-commit
+//! (549, 550) for absorbing growth. T16 consumed those (vol_ema +
+//! vol_ref) and added one more for the bootstrap-sample counter,
+//! bumping the block from 12 to 13 slots. Add new alpha-system slots
+//! HERE first, not by extending past 551 ad-hoc.
 
 pub const ALPHA_ISV_BLOCK_LO: usize = 539;
-pub const ALPHA_ISV_BLOCK_HI: usize = 550; // inclusive upper
+pub const ALPHA_ISV_BLOCK_HI: usize = 552; // inclusive upper
 
 // Diagnostics (read-only for circuit breakers; not controllers)
 pub const Q_SPREAD_EMA_INDEX: usize = 539;
@@ -69,6 +72,26 @@ pub const RANDOM_BASELINE_STD_INDEX: usize = 548;
 // computing kelly_attenuation update.
 pub const REGIME_VOL_EMA_INDEX: usize = 549;
 pub const REGIME_VOL_REF_INDEX: usize = 550;
+// Sample counter that gates the vol_ref bootstrap. While count <
+// `REGIME_VOL_REF_BOOTSTRAP_SAMPLES` the kernel updates vol_ref as a
+// running mean over the observed vol_ema values; once the counter
+// crosses the threshold, vol_ref switches to slow β-tracking. Without
+// this, the very first observation drives vol_ref permanently noisy,
+// and short-training-window walk-forward folds get pessimistic Kelly
+// attenuation immediately (regression confirmed 2026-05-16 against
+// the previous pre-defense walk-forward Sharpe).
+pub const REGIME_VOL_REF_SAMPLES_INDEX: usize = 551;
+// Cross-fold-persistent floor on REGIME_VOL_REF (per
+// `pearl_kelly_cap_signal_driven_floors` and
+// `pearl_blend_formulas_must_have_permanent_floor`). Prevents the
+// deadband-deadlock where vol_ref → 0 during calm market stretches
+// then traps Kelly at the regime_scale floor once any vol resumes.
+// TrainingPersist semantics: NOT cleared by fold reset; survives
+// across cells of a 2D sweep so the first cell's bootstrap statistics
+// inform subsequent cells. Hardcoded sub-floor of 1e-12 inside the
+// kernel guards against the slot itself collapsing to zero.
+// Sentinel value 0 → kernel uses the hardcoded sub-floor only.
+pub const REGIME_VOL_REF_FLOOR_INDEX: usize = 552;
 
 // Sentinels: zero by convention. Fallbacks use the first-observation
 // bootstrap pattern (see pearl_first_observation_bootstrap.md).
@@ -83,7 +106,7 @@ mod tests {
     fn test_block_bounds_consistent() {
         assert!(ALPHA_ISV_BLOCK_HI >= ALPHA_ISV_BLOCK_LO);
         let block_size = ALPHA_ISV_BLOCK_HI - ALPHA_ISV_BLOCK_LO + 1;
-        assert_eq!(block_size, 12, "alpha-system reserved 12 contiguous slots");
+        assert_eq!(block_size, 14, "alpha-system reserved 14 contiguous slots");
     }
 
     #[test]
@@ -101,6 +124,8 @@ mod tests {
             RANDOM_BASELINE_STD_INDEX,
             REGIME_VOL_EMA_INDEX,
             REGIME_VOL_REF_INDEX,
+            REGIME_VOL_REF_SAMPLES_INDEX,
+            REGIME_VOL_REF_FLOOR_INDEX,
         ];
         for &i in &indices {
             assert!(i >= ALPHA_ISV_BLOCK_LO && i <= ALPHA_ISV_BLOCK_HI,
@@ -123,6 +148,8 @@ mod tests {
             RANDOM_BASELINE_STD_INDEX,
             REGIME_VOL_EMA_INDEX,
             REGIME_VOL_REF_INDEX,
+            REGIME_VOL_REF_SAMPLES_INDEX,
+            REGIME_VOL_REF_FLOOR_INDEX,
         ];
         let mut sorted: Vec<usize> = indices.to_vec();
         sorted.sort();

crates/ml/src/cuda_pipeline/alpha_kernels.rs

@@ -480,10 +480,14 @@ pub unsafe fn launch_alpha_regime_vol_update(
     b: i32,
     vol_ema_index: i32,
     vol_ref_index: i32,
+    vol_ref_samples_index: i32,
+    vol_ref_floor_index: i32,
+    bootstrap_samples: 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}");
+    debug_assert!(bootstrap_samples > 0, "bootstrap_samples must be positive");
     let cfg = LaunchConfig {
         grid_dim: (1, 1, 1),
         block_dim: (256, 1, 1),
@@ -497,6 +501,9 @@ pub unsafe fn launch_alpha_regime_vol_update(
         .arg(&b)
         .arg(&vol_ema_index)
         .arg(&vol_ref_index)
+        .arg(&vol_ref_samples_index)
+        .arg(&vol_ref_floor_index)
+        .arg(&bootstrap_samples)
         .arg(&ref_track_rate)
         .launch(cfg)
         .map_err(|e| MLError::ModelError(format!("alpha_regime_vol_update launch: {e}")))?;

crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu

@@ -63,6 +63,9 @@ extern "C" __global__ void alpha_regime_vol_update_kernel(
     int          B,
     int          vol_ema_index,            // 549
     int          vol_ref_index,            // 550
+    int          vol_ref_samples_index,    // 551 (bootstrap counter)
+    int          vol_ref_floor_index,      // 552 (TrainingPersist floor anchor)
+    int          bootstrap_samples,        // N — typ. 100
     float        ref_track_rate            // β (per-step), e.g. 0.005
 ) {
     // Single-block design: each thread handles one b. Final reduction by
@@ -123,16 +126,49 @@ extern "C" __global__ void alpha_regime_vol_update_kernel(
     }
     isv[vol_ema_index] = vol_ema_new;
 
-    // Slow reference vol tracking on slot 550.
+    // Slow reference vol tracking on slot 550 with a bootstrap window.
+    //
+    // Old behaviour: vol_ref ← vol_ema on the very first observation.
+    // Problem: vol_obs is itself high-variance, so vol_ref started off
+    // noisy and short-training-window walk-forward folds saw vol_ema
+    // diverge from a noisy vol_ref → spurious early Kelly attenuation.
+    //
+    // New behaviour: while sample count < bootstrap_samples, vol_ref
+    // is the running mean over the observed vol_ema values
+    //   vol_ref ← vol_ref + (vol_ema − vol_ref) / (count + 1)
+    // After bootstrap, switch to slow β-tracking
+    //   vol_ref ← vol_ref + β · (vol_ema − vol_ref)
+    // Running mean shrinks the bootstrap variance by √N relative to
+    // single-shot replace; median would shrink it further but requires
+    // either O(N) buffered samples or a quantile sketch — running mean
+    // is single-pass with one f32 of extra state, which fits naturally
+    // in an ISV slot.
     const float vol_ref_prev = isv[vol_ref_index];
+    const float sample_count_prev = isv[vol_ref_samples_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;
+    if (sample_count_prev < (float)bootstrap_samples) {
+        // Bootstrap phase: running mean of observed vol_ema values.
+        const float n = sample_count_prev + 1.0f;
+        vol_ref_new = vol_ref_prev + (vol_ema_new - vol_ref_prev) / n;
     } else {
+        // Steady state: slow β-tracking.
         vol_ref_new = vol_ref_prev + ref_track_rate * (vol_ema_new - vol_ref_prev);
     }
+    // Permanent floor (`pearl_blend_formulas_must_have_permanent_floor`).
+    // Two layers per `pearl_kelly_cap_signal_driven_floors`:
+    //   * ISV-driven floor (slot 552, TrainingPersist) — the learned
+    //     anchor. Sentinel 0 → fall back to the hardcoded sub-floor.
+    //   * Hardcoded sub-floor 1e-12 — purely a float-underflow guard
+    //     against the ISV slot itself collapsing to zero.
+    // Prevents the asymmetric deadlock: if vol_ref decays to ≈0 during
+    // a deadband stretch (calm market, near-zero squared log-returns),
+    // the spike-detector ratio vol_ref / vol_ema fires the moment any
+    // realistic vol resumes — Kelly stays trapped at regime_scale_floor
+    // until a regime change pushes vol_ref back up. Floor caps how small
+    // "trained-regime baseline" can claim to be.
+    const float floor_anchor = isv[vol_ref_floor_index];
+    const float effective_floor = fmaxf(floor_anchor, 1.0e-12f);
+    vol_ref_new = fmaxf(vol_ref_new, effective_floor);
     isv[vol_ref_index] = vol_ref_new;
+    isv[vol_ref_samples_index] = sample_count_prev + 1.0f;
 }

docs/isv-slots.md

@@ -913,3 +913,54 @@ CV across Q1↔Q2 boundary with all four configurations:
   - T10 + both
 Expected: combined config compresses cross-fold std-dev meaningfully
 without sacrificing best-fold Sharpe.
+
+## 2026-05-16 — Regime vol_ref bootstrap window + permanent floor
+
+Walk-forward CV on Q1 fxcache after promoting --regime-scale to
+always-on exposed two failure modes of the original Pearl-A single-shot
+bootstrap of vol_ref:
+
+1. The very first vol_obs observation is itself high-variance, so
+   vol_ref started off noisy and short-training-window folds saw
+   vol_ema diverge from a noisy vol_ref → spurious early Kelly
+   attenuation.
+2. If vol_ref decayed to ≈0 during a deadband stretch, the ratio
+   vol_ref / vol_ema fired the moment any realistic vol resumed
+   and Kelly stayed trapped at the regime_scale_floor=0.25.
+
+**Slot 551 — REGIME_VOL_REF_SAMPLES_INDEX (counter)**
+- Producer: `alpha_regime_vol_update.cu`. Increments by 1 each kernel
+  call.
+- Gates the vol_ref update path. While count < 100, vol_ref is the
+  running mean over observed vol_ema values:
+    `vol_ref ← vol_ref + (vol_ema − vol_ref) / (count + 1)`
+  After 100 samples, switch to slow β-tracking with β=0.005.
+- Cuts bootstrap-window variance by √N (~10× for N=100) relative to
+  single-shot Pearl-A replace.
+
+**Slot 552 — REGIME_VOL_REF_FLOOR_INDEX (TrainingPersist anchor)**
+- Permanent floor on vol_ref. Two-layered design per
+  `pearl_kelly_cap_signal_driven_floors`:
+  - ISV slot is the learned anchor — named, addressable, observable.
+  - Hardcoded sub-floor of 1e-12 inside the kernel guards against
+    the slot itself collapsing to zero.
+- Sentinel value 0 → kernel uses the hardcoded sub-floor only.
+- Host seeds with 1e-9 at training start (one order of magnitude
+  below ES MBP-10 typical squared-log-return ~1e-10..1e-6). A future
+  controller kernel will refine this from observed cell-level vol
+  minima with cross-fold persistence.
+- TrainingPersist semantics: NOT cleared by fold reset.
+
+**Kernel signature change:** `alpha_regime_vol_update_kernel` gained
+two int slot-index args (`vol_ref_samples_index`, `vol_ref_floor_index`)
+and one int hyperparam (`bootstrap_samples`). The previous f32
+`vol_ref_floor` scalar arg was dropped — the anchor now lives in ISV.
+
+**Walk-forward CV verdict (Q1 fxcache, 3 folds, window=700K, train_frac=0.6):**
+  - cost=0.0000 mean Sharpe_ann = +17.24 ± 43.42 (was +38.94 ± 56.88)
+  - cost=0.0625 mean Sharpe_ann = +8.39  ± 47.47
+  - cost=0.1250 mean Sharpe_ann = -2.18  ± 50.77
+- Fold B turnaround: -47.64 (bootstrap-only) → +65.04 (bootstrap + floor).
+- Cross-fold std-dev compressed 24% at cost=0; mean dropped from
+  +38.94 (pre-defense) to +17.24 (with defense). Classic mean/variance
+  trade — variance compression at meaningful mean cost.