From d49003de6d1014653b0a26e11b754007f00a5611 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 16 May 2026 09:22:20 +0200 Subject: [PATCH] feat(regime): vol_ref floor controller + cross-invocation disk persist MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes the TrainingPersist loop for the regime defense per the KELLY_F_SMOOTH precedent (pearl_kelly_cap_signal_driven_floors). Three layers: (1) Per-step: alpha_regime_vol_update_kernel tracks min(vol_obs > 1e-12) in new slot 553 (REGIME_VOL_OBS_MIN_INDEX). Filtered against artifact- zero observations from stationary snapshots where curr == prev. (2) Per-cell: stacker_threshold_controller_update gains a fifth branch that reads vol_ref (slot 550) at cell-end, computes `target = floor_target_ratio × vol_ref`, slow-EMAs the floor anchor (slot 552) toward the target with rate `floor_update_rate` (0.1 per cell). Subfloor 1e-12 inside the kernel guards against the anchor collapsing to zero. (3) Per-invocation: alpha_baseline reads `config/ml/alpha_baseline_state.json` at startup and seeds slot 552 from the `regime_vol_ref_floor` field. At end of main(), the learned floor is written back via tmp+rename atomic write so concurrent walk-forward invocations see a consistent file. Matches the cross-fold-persistent shape of KELLY_F_SMOOTH. Block extended to 15 slots (539..=553). Smoke + kernel unit test pass -1/-1 for the new floor-controller indices (backward compat). Walk-forward CV verdict (Q1 fxcache, 3 sequential folds): iteration fold-A fold-B fold-C mean ± SD pre-defense (no regime) +91.52 -21.44 +46.74 +38.94 ± 56.88 hardcoded 1e-9 floor -19.77 +65.04 +6.45 +17.24 ± 43.42 learned floor (0.5 × cell_min) +74.78 -12.72 +8.80 +23.62 ± 45.59 learned floor (0.1 × vol_ref) -19.53 -26.51 +15.46 -10.20 ± 22.49 Controller infrastructure is structurally correct (loop closes, floor persists across invocations, kernel + disk + ISV all roundtrip). The TUNING is data-dependent — single-quarter CV doesn't have enough regime diversity to anchor the floor against. Multi-quarter fxcache validation is the next step (built cluster-side on the 9-quarter 2024-Q1..2026-Q1 ES futures dataset, downloaded as a single artifact for local CV). Co-Authored-By: Claude Opus 4.7 --- crates/ml/examples/alpha_baseline.rs | 78 +++++++++++++++++-- crates/ml/examples/alpha_dqn_h600_smoke.rs | 9 ++- .../ml/src/cuda_pipeline/alpha_isv_slots.rs | 21 ++++- crates/ml/src/cuda_pipeline/alpha_kernels.rs | 23 ++++++ .../cuda_pipeline/alpha_regime_vol_update.cu | 13 ++++ .../stacker_threshold_controller.cu | 45 +++++++++++ docs/isv-slots.md | 45 +++++++++++ 7 files changed, 220 insertions(+), 14 deletions(-) diff --git a/crates/ml/examples/alpha_baseline.rs b/crates/ml/examples/alpha_baseline.rs index 327e83081..ef91c4129 100644 --- a/crates/ml/examples/alpha_baseline.rs +++ b/crates/ml/examples/alpha_baseline.rs @@ -471,17 +471,32 @@ 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 = vec![0.0; 553]; + let mut isv_host: Vec = vec![0.0; 554]; 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; + // T16 vol_ref floor seed (TrainingPersist anchor). Read from disk if + // a previous training run wrote a learned value; otherwise seed with + // a sensible default based on ES MBP-10 squared-log-return range + // (typical vol_obs ~1e-10..1e-6). Hardcoded sub-floor 1e-12 inside + // the regime kernel still catches the slot collapsing to zero. + let floor_state_path = std::path::PathBuf::from("config/ml/alpha_baseline_state.json"); + let seeded_floor: f32 = std::fs::read_to_string(&floor_state_path) + .ok() + .and_then(|s| serde_json::from_str::(&s).ok()) + .and_then(|v| v.get("regime_vol_ref_floor").and_then(|x| x.as_f64())) + .map(|x| x as f32) + .unwrap_or(1.0e-9); + isv_host[ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_FLOOR_INDEX] = seeded_floor; + info!( + "Regime floor seed: {:.3e} ({})", + seeded_floor, + if floor_state_path.exists() { "from disk" } else { "default" }, + ); + // T16 vol_obs running-min slot: reset to sentinel so the first cell + // accumulates min from any positive vol_obs observation. + isv_host[ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_INDEX] = + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_SENTINEL; let mut isv_dev = stream.clone_htod(&isv_host).context("upload isv")?; let mut ctl_wiener_dev = stream.alloc_zeros::(3).context("alloc ctl wiener")?; @@ -877,6 +892,18 @@ fn main() -> Result<()> { mids_curr_pinned.write(&zeros); mids_prev_pinned.write(&zeros); } + // T16: reset the vol_obs min slot to its sentinel at cell start + // so each cell's min accumulates independently. The floor + // controller (fired at cell-end inside the stacker controller) + // will consume this slot's final value to update the floor + // anchor. + { + let sentinel = vec![ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_SENTINEL]; + let start = ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_INDEX; + let mut slot_view = isv_dev.slice_mut(start..start + 1); + stream.memcpy_htod(&sentinel, &mut slot_view) + .context("reset vol_obs min slot")?; + } // Lockstep step loop. for _step in 0..cli.horizon { // Gather current states (CPU; <100μs for N=500). @@ -912,6 +939,7 @@ fn main() -> Result<()> { ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_INDEX as i32, 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, + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_INDEX as i32, 100, // bootstrap samples: running mean over first 100 obs 0.005, // β: 200-step horizon for vol_ref tracking after bootstrap )?; @@ -1015,6 +1043,17 @@ fn main() -> Result<()> { 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.25, + // T16 floor controller: refine REGIME_VOL_REF_FLOOR + // toward 0.1 × vol_ref (slot 550). Floor sits at + // 10% of the trained-regime baseline — small + // enough that the spike detector still fires on + // realistic vol_ema excursions, large enough to + // prevent the deadband deadlock. + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_FLOOR_INDEX as i32, + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_INDEX as i32, + 0.1, // floor sits at 10% of vol_ref + 0.1, // slow EMA toward target (10% per cell) + ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_OBS_MIN_SENTINEL, isv_ptr, wv_ptr, )?; } @@ -1119,6 +1158,29 @@ fn main() -> Result<()> { write!(f, "{}", serde_json::to_string_pretty(&json)?)?; info!("Wrote table to {}", cli.out_path.display()); + // T16: persist the learned vol_ref floor across invocations so the + // KELLY_F_SMOOTH-equivalent cross-fold-persistent pattern works for + // walk-forward CV. Read the final floor value off the device, write + // to disk; next invocation reloads via the seed-from-disk path at + // startup. Atomic write: tmp file + rename so concurrent reads see + // either the old or the new file, never a partial one. + { + stream.synchronize().context("sync before floor persist")?; + let isv_final = stream.clone_dtoh(&isv_dev).context("dtoh isv for floor persist")?; + let learned_floor = isv_final + [ml::cuda_pipeline::alpha_isv_slots::REGIME_VOL_REF_FLOOR_INDEX]; + info!( + "Regime floor learned: {:.3e} (was {:.3e})", + learned_floor, seeded_floor, + ); + let json = serde_json::json!({ "regime_vol_ref_floor": learned_floor as f64 }); + let tmp_path = floor_state_path.with_extension("json.tmp"); + std::fs::write(&tmp_path, serde_json::to_string_pretty(&json)?) + .context("write floor state tmp")?; + std::fs::rename(&tmp_path, &floor_state_path) + .context("rename floor state into place")?; + } + // Touch unused ISV slot constants so they're imported for future expansion. let _ = (RANDOM_BASELINE_MEAN_INDEX, RANDOM_BASELINE_STD_INDEX); let _: Option = None; diff --git a/crates/ml/examples/alpha_dqn_h600_smoke.rs b/crates/ml/examples/alpha_dqn_h600_smoke.rs index 7cd3b3a5b..d9b090891 100644 --- a/crates/ml/examples/alpha_dqn_h600_smoke.rs +++ b/crates/ml/examples/alpha_dqn_h600_smoke.rs @@ -1479,11 +1479,16 @@ fn main() -> Result<()> { 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. + // scale here (backward compat). -1, -1, 0.25, + // T16 floor controller — also disabled in smoke. + -1, + -1, + 0.5, + 0.1, + 1.0e10, 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 38a52b518..96a2e9715 100644 --- a/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/alpha_isv_slots.rs @@ -1,6 +1,6 @@ //! Alpha trading system ISV slot block. //! -//! Reserves indices 539..=552 (14 slots, contiguous) for durable +//! Reserves indices 539..=553 (15 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 @@ -27,7 +27,8 @@ //! | 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 | +//! | 552 | REGIME_VOL_REF_FLOOR_INDEX | Anchor | TrainingPersist | host seed + stacker_threshold_controller | +//! | 553 | REGIME_VOL_OBS_MIN_INDEX | Diagnostic | per-inference | alpha_regime_vol_update.cu | //! //! Reservation discipline: this block had two spares at first-commit //! (549, 550) for absorbing growth. T16 consumed those (vol_ema + @@ -36,7 +37,7 @@ //! HERE first, not by extending past 551 ad-hoc. pub const ALPHA_ISV_BLOCK_LO: usize = 539; -pub const ALPHA_ISV_BLOCK_HI: usize = 552; // inclusive upper +pub const ALPHA_ISV_BLOCK_HI: usize = 553; // inclusive upper // Diagnostics (read-only for circuit breakers; not controllers) pub const Q_SPREAD_EMA_INDEX: usize = 539; @@ -92,6 +93,16 @@ pub const REGIME_VOL_REF_SAMPLES_INDEX: usize = 551; // 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; +// Running min of vol_obs observed within the current eval cell. +// Updated by `alpha_regime_vol_update_kernel` per inference step, but +// filtered to skip vol_obs <= 1e-12 (artifact-zero observations from +// stationary snapshots where curr == prev). Reset by host to a large +// sentinel at the start of each cell so the next cell's min +// accumulates cleanly. Consumed by the floor-update branch of the +// stacker controller at cell-end. Sentinel "large value" is 1e10 +// (anything above any realistic vol_obs). +pub const REGIME_VOL_OBS_MIN_INDEX: usize = 553; +pub const REGIME_VOL_OBS_MIN_SENTINEL: f32 = 1.0e10; // Sentinels: zero by convention. Fallbacks use the first-observation // bootstrap pattern (see pearl_first_observation_bootstrap.md). @@ -106,7 +117,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, 14, "alpha-system reserved 14 contiguous slots"); + assert_eq!(block_size, 15, "alpha-system reserved 15 contiguous slots"); } #[test] @@ -126,6 +137,7 @@ mod tests { REGIME_VOL_REF_INDEX, REGIME_VOL_REF_SAMPLES_INDEX, REGIME_VOL_REF_FLOOR_INDEX, + REGIME_VOL_OBS_MIN_INDEX, ]; for &i in &indices { assert!(i >= ALPHA_ISV_BLOCK_LO && i <= ALPHA_ISV_BLOCK_HI, @@ -150,6 +162,7 @@ mod tests { REGIME_VOL_REF_INDEX, REGIME_VOL_REF_SAMPLES_INDEX, REGIME_VOL_REF_FLOOR_INDEX, + REGIME_VOL_OBS_MIN_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 6962efe72..282ea2ef1 100644 --- a/crates/ml/src/cuda_pipeline/alpha_kernels.rs +++ b/crates/ml/src/cuda_pipeline/alpha_kernels.rs @@ -267,6 +267,12 @@ pub unsafe fn launch_stacker_threshold_controller( regime_vol_ema_idx: i32, regime_vol_ref_idx: i32, regime_scale_floor: f32, + // T16 floor-controller args. Pass -1 / -1 to disable the floor update. + regime_vol_ref_floor_idx: i32, + regime_vol_obs_min_idx: i32, + floor_target_ratio: f32, + floor_update_rate: f32, + vol_obs_min_sentinel: f32, isv_dev: u64, wiener_state_dev: u64, ) -> Result<(), MLError> { @@ -303,6 +309,11 @@ pub unsafe fn launch_stacker_threshold_controller( .arg(®ime_vol_ema_idx) .arg(®ime_vol_ref_idx) .arg(®ime_scale_floor) + .arg(®ime_vol_ref_floor_idx) + .arg(®ime_vol_obs_min_idx) + .arg(&floor_target_ratio) + .arg(&floor_update_rate) + .arg(&vol_obs_min_sentinel) .arg(&isv_dev) .arg(&wiener_state_dev) .launch(cfg) @@ -482,6 +493,7 @@ pub unsafe fn launch_alpha_regime_vol_update( vol_ref_index: i32, vol_ref_samples_index: i32, vol_ref_floor_index: i32, + vol_obs_min_index: i32, bootstrap_samples: i32, ref_track_rate: f32, ) -> Result<(), MLError> { @@ -503,6 +515,7 @@ pub unsafe fn launch_alpha_regime_vol_update( .arg(&vol_ref_index) .arg(&vol_ref_samples_index) .arg(&vol_ref_floor_index) + .arg(&vol_obs_min_index) .arg(&bootstrap_samples) .arg(&ref_track_rate) .launch(cfg) @@ -1850,6 +1863,11 @@ mod tests { -1, // regime_vol_ema_idx (disabled in test) -1, // regime_vol_ref_idx (disabled in test) 0.25, // regime_scale_floor (unused when disabled) + -1, // regime_vol_ref_floor_idx (disabled in test) + -1, // regime_vol_obs_min_idx (disabled in test) + 0.5, // floor_target_ratio (unused when disabled) + 0.1, // floor_update_rate (unused when disabled) + 1.0e10, // vol_obs_min_sentinel (unused when disabled) isv_ptr, w_ptr, ) @@ -1897,6 +1915,11 @@ mod tests { -1, // regime_vol_ema_idx (disabled in test) -1, // regime_vol_ref_idx (disabled in test) 0.25, // regime_scale_floor (unused when disabled) + -1, // regime_vol_ref_floor_idx (disabled in test) + -1, // regime_vol_obs_min_idx (disabled in test) + 0.5, // floor_target_ratio (unused when disabled) + 0.1, // floor_update_rate (unused when disabled) + 1.0e10, // vol_obs_min_sentinel (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 index e8075bee6..405b01f50 100644 --- a/crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu +++ b/crates/ml/src/cuda_pipeline/alpha_regime_vol_update.cu @@ -65,6 +65,7 @@ extern "C" __global__ void alpha_regime_vol_update_kernel( int vol_ref_index, // 550 int vol_ref_samples_index, // 551 (bootstrap counter) int vol_ref_floor_index, // 552 (TrainingPersist floor anchor) + int vol_obs_min_index, // 553 (running min — input to floor controller) int bootstrap_samples, // N — typ. 100 float ref_track_rate // β (per-step), e.g. 0.005 ) { @@ -171,4 +172,16 @@ extern "C" __global__ void alpha_regime_vol_update_kernel( 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; + + // Per-cell running min of vol_obs, FILTERED for vol_obs > 1e-12. At + // MBP-10 cadence many snapshots see no mid-price change (curr == + // prev → vol_obs = 0), so an unfiltered min would pin at 0 on the + // first such bar and the floor controller would converge to 0. The + // filter skips those artifact-zero observations and tracks the min + // over "actually moved" steps — that's the signal the floor + // controller needs. + if (vol_obs > 1.0e-12f) { + const float min_prev = isv[vol_obs_min_index]; + isv[vol_obs_min_index] = fminf(min_prev, vol_obs); + } } diff --git a/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu index 8f931093c..c185e7891 100644 --- a/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu +++ b/crates/ml/src/cuda_pipeline/stacker_threshold_controller.cu @@ -63,6 +63,14 @@ extern "C" __global__ void stacker_threshold_controller_update( int regime_vol_ema_idx, int regime_vol_ref_idx, float regime_scale_floor, // typ. 0.25 — bottoms the multiplicative defense + /* T16 vol_ref floor controller (set indices < 0 to disable). At + cell-end, observe the filtered cell-min vol_obs and slow-EMA + the floor anchor toward `floor_target_ratio × cell_min`. */ + int regime_vol_ref_floor_idx, + int regime_vol_obs_min_idx, + float floor_target_ratio, // typ. 0.5 — floor sits half the cell-min + float floor_update_rate, // typ. 0.1 — slow EMA toward target + float vol_obs_min_sentinel, // typ. 1e10 — large value reset by host /* Device buffers. */ float* __restrict__ isv, float* __restrict__ wiener_state /* [sample_var, diff_var, x_lag] for slot 545 */ @@ -176,5 +184,42 @@ extern "C" __global__ void stacker_threshold_controller_update( const float final_atten = fmaxf(ATTEN_LO, reactive_atten * regime_scale); isv[stacker_kelly_atten_idx] = final_atten; + // ---- (5) T16 vol_ref floor controller update ---- + // + // The vol_ref floor anchor (slot 552) prevents the deadband + // deadlock; it has to be SIZED to the instrument's typical + // vol_obs range. Initial seed is a host hyperparam; this + // controller refines it from observed signal so the same binary + // works across instruments / time regimes without re-tuning. + // + // Signal: the trained-regime baseline vol_ref itself (slot 550). + // Target: `floor_target_ratio × vol_ref` — floor sits at a small + // fraction of the regime baseline (typically 0.1 = 10%). This is + // architecturally coherent: floor is a fraction of the same signal + // the spike detector compares against (vol_ema vs vol_ref), so the + // floor scales naturally with regime shifts as vol_ref tracks the + // current regime. + // Update: slow EMA `floor_new = floor_old + α × (target − old)`. + // + // The per-step min(vol_obs) slot 553 is now an observability-only + // diagnostic — written by the regime kernel for inspection but no + // longer consumed by the controller (the 0.5 × cell_min target + // overshot in walk-forward and effectively disabled the defense). + // + // Disabled when either floor or vol_ref index < 0 + // (backward-compat for smoke + kernel unit test). + (void)regime_vol_obs_min_idx; // unused signal channel — see above + (void)vol_obs_min_sentinel; + if (regime_vol_ref_floor_idx >= 0 && regime_vol_ref_idx >= 0) { + const float vol_ref_current = isv[regime_vol_ref_idx]; + if (vol_ref_current > 1.0e-12f) { + const float floor_prev = isv[regime_vol_ref_floor_idx]; + const float target = floor_target_ratio * vol_ref_current; + const float floor_new = + floor_prev + floor_update_rate * (target - floor_prev); + isv[regime_vol_ref_floor_idx] = fmaxf(floor_new, 1.0e-12f); + } + } + __threadfence_system(); } diff --git a/docs/isv-slots.md b/docs/isv-slots.md index e713f06a4..ec62d2177 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -964,3 +964,48 @@ and one int hyperparam (`bootstrap_samples`). The previous f32 - 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. + +## 2026-05-16 — Slot 553 + floor controller wired into stacker_threshold_controller + +The hardcoded 1e-9 floor was the obvious anti-pattern; this round +makes the floor signal-driven per `pearl_controller_anchors_isv_driven`. + +**Slot 553 — REGIME_VOL_OBS_MIN_INDEX (diagnostic)** +- Producer: `alpha_regime_vol_update.cu` per inference step. Tracks + `min(vol_obs)` filtered for `vol_obs > 1e-12` (skips artifact-zero + observations from stationary snapshots where curr == prev). +- Reset to `REGIME_VOL_OBS_MIN_SENTINEL = 1e10` by host at cell start. +- Originally intended as the input to the floor controller, but the + empirical tuning (target = 0.5 × cell_min) overshot — the floor + ended at 7.787e-3 across 3 folds, an order of magnitude above + typical vol_ema, which effectively disabled the regime defense. + Retained as observability-only diagnostic. + +**Floor controller — branch (5) of `stacker_threshold_controller.cu`** +- Reads slot 550 (REGIME_VOL_REF) at cell-end, computes + `target = floor_target_ratio × vol_ref`, slow-EMAs slot 552 toward + the target with rate `floor_update_rate` (typ. 0.1 per cell). +- Subfloor 1e-12 inside the kernel guards against the anchor + collapsing to zero. +- Disabled when either `regime_vol_ref_floor_idx` or + `regime_vol_ref_idx` is -1 (backward-compat for smoke + kernel test). + +**Disk persistence** +- `alpha_baseline` reads `config/ml/alpha_baseline_state.json` at + startup (key `regime_vol_ref_floor`) and seeds slot 552. At end of + main(), the learned floor is written back via tmp+rename atomic + write so concurrent walk-forward invocations see a consistent file. +- Matches the `KELLY_F_SMOOTH` cross-fold-persistent pattern, applied + to the regime defense. + +**Tuning observation (open)** +- Walk-forward CV with target = 0.1 × vol_ref settled the floor at + 1.792e-6 (across 3 folds). Defense fired aggressively → cross-fold + std-dev was the lowest measured (22.49 at cost=0) but mean Sharpe + the lowest measured (-10.20 at cost=0). The fraction-of-vol_ref + target makes the "floor" track vol_ref rather than acting as a + fixed lower bound — wrong architectural shape. +- The right tune is data-dependent: needs multi-quarter validation + to anchor the floor against a diverse vol regime. Local Q1+Q2 + fxcache rebuild OOMs on RTX 3050 Ti; cluster-side build is the + next step.