From da21feb1b1ab98751f627319def54387380fc406 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 19 May 2026 13:41:17 +0200 Subject: [PATCH] fix(ml-backtesting): cold-start Kelly + Sharpe-weight floors in decision kernel MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Production smoke completed end-to-end but produced n_trades=0 across 99,969 decisions — `decision_policy_default` and `decision_policy_program` both applied a sentinel-skip pattern: if `isv_kelly_d` had not been seeded (pnl_ema_win == 0), each horizon's signed-size stayed zero, AND each horizon's aggregation weight (= recent_sharpe) also stayed zero. The cross-horizon w_sum was therefore 0, final_size was 0, every market_target was noop. State only updates on trade close → no trade ever fires → infinite cold-start. Per pearl_blend_formulas_must_have_permanent_floor (`max(real, floor)`, not blend) and pearl_kelly_cap_signal_driven_floors, replace the sentinel- skip with a two-layer floor on each kernel: 1. Kelly fraction: `max(kelly_frac_floor, computed_kelly)` — when state is sentinel, falls back to the floor directly. Cap_lots falls back to `max_lots` when realised_return_var is sentinel. 2. Aggregation weight: `max(sharpe_weight_floor, recent_sharpe)` — lets cross-horizon sum produce a non-zero size before recent_sharpe is populated. Once a horizon shows positive sharpe it dominates. Plumbed through `step_decision_with_latency` / `step_decision` as two new f32 args (atomic contract change, every caller migrated). Defaults 0.20 / 0.10 chosen so a strong-conviction signal (sig_mag ≥ 0.5) fires 1 lot at cold-start under max_lots=5 while weaker signals stay flat (see `default_kelly_frac_floor` comment for the arithmetic). Exposed as CLI flags + sweep-grid base/cell overrides. Regression test `decision_floor_coldstart` proves: - default floors (0.20/0.10) fire a 1-lot buy with p_h=0.8 and zero state - zero floors reproduce the original noop bug Also moves `aggregate` step to the GPU pool because fxt-backtest is dynamically linked against libcuda.so.1 (the ci-compile-cpu hosts don't expose CUDA driver libs). Verified locally on RTX 3050 Ti — workspace cargo check passes, both regression tests pass, trunk save/load roundtrip still passes. Co-Authored-By: Claude Opus 4.7 (1M context) --- bin/fxt-backtest/src/main.rs | 30 ++++++ crates/ml-backtesting/cuda/decision_policy.cu | 95 +++++++++++++------ crates/ml-backtesting/src/harness.rs | 13 +++ crates/ml-backtesting/src/sim.rs | 27 +++++- .../tests/decision_floor_coldstart.rs | 71 ++++++++++++++ .../ml-backtesting/tests/lob_sim_fixtures.rs | 9 +- .../tests/lob_sim_integrated_fuzz.rs | 2 + .../k8s/argo/lob-backtest-sweep-template.yaml | 11 ++- 8 files changed, 225 insertions(+), 33 deletions(-) create mode 100644 crates/ml-backtesting/tests/decision_floor_coldstart.rs diff --git a/bin/fxt-backtest/src/main.rs b/bin/fxt-backtest/src/main.rs index f147ce222..3a9ebccf5 100644 --- a/bin/fxt-backtest/src/main.rs +++ b/bin/fxt-backtest/src/main.rs @@ -90,6 +90,17 @@ struct RunArgs { /// Cap on events processed; 0 = exhaust the input stream. #[arg(long, default_value_t = 0)] max_events: u64, + /// Cold-start Kelly fraction floor (see BacktestHarnessConfig + + /// pearl_blend_formulas_must_have_permanent_floor). When isv_kelly + /// has no closed-trade history, kelly_frac defaults to this floor so + /// the first trade can fire and bootstrap state. + #[arg(long, default_value_t = 0.20)] + kelly_frac_floor: f32, + /// Cold-start Sharpe weight floor — applied to recent_sharpe in the + /// cross-horizon aggregation. Lets the aggregate produce a non-zero + /// signed size before recent_sharpe is populated. + #[arg(long, default_value_t = 0.10)] + sharpe_weight_floor: f32, /// Random seed for trunk initialisation when --checkpoint is not given. #[arg(long, default_value_t = 0xC0FFEE)] seed: u64, @@ -175,6 +186,8 @@ struct SweepBase { #[serde(default)] max_events: u64, #[serde(default = "default_seed")] seed: u64, #[serde(default)] checkpoint: Option, + #[serde(default = "default_kelly_frac_floor")] kelly_frac_floor: f32, + #[serde(default = "default_sharpe_weight_floor")] sharpe_weight_floor: f32, } fn default_n_parallel() -> usize { 1 } @@ -184,6 +197,15 @@ fn default_target_vol() -> f32 { 50.0 } fn default_ann_factor() -> f32 { 825.0 } fn default_max_lots() -> u16 { 5 } fn default_seed() -> u64 { 0xC0FFEE } +// Sized so a strong-conviction signal (|p-0.5| ≥ 0.2 ⇒ sig_mag ≥ 0.4) +// produces ≥1 lot at cold-start for max_lots=5: +// ss = sig_mag * floor * cap_lots → 0.4 * 0.20 * 5 = 0.4 → rounds to 0 +// ss = sig_mag * floor * cap_lots → 0.5 * 0.20 * 5 = 0.5 → rounds to 1 +// Weaker signals stay at 0 lots cold-start. Once kelly state bootstraps +// (first close), `max(floor, computed_kelly)` keeps the floor as the +// minimum trading intensity. +fn default_kelly_frac_floor() -> f32 { 0.20 } +fn default_sharpe_weight_floor() -> f32 { 0.10 } #[derive(Debug, serde::Deserialize, Clone)] struct SweepCell { @@ -196,6 +218,8 @@ struct SweepCell { #[serde(default)] max_events: Option, #[serde(default)] seed: Option, #[serde(default)] checkpoint: Option, + #[serde(default)] kelly_frac_floor: Option, + #[serde(default)] sharpe_weight_floor: Option, } fn main() -> Result<()> { @@ -279,6 +303,10 @@ fn sweep(args: SweepArgs) -> Result<()> { max_events: cell.max_events.unwrap_or(grid.base.max_events), seed: cell.seed.unwrap_or(grid.base.seed), checkpoint: cell.checkpoint.clone().or_else(|| grid.base.checkpoint.clone()), + kelly_frac_floor: + cell.kelly_frac_floor.unwrap_or(grid.base.kelly_frac_floor), + sharpe_weight_floor: + cell.sharpe_weight_floor.unwrap_or(grid.base.sharpe_weight_floor), out: cell_out.clone(), }; run(run_args).with_context(|| format!("sweep cell {}", cell.name))?; @@ -363,6 +391,8 @@ fn run(args: RunArgs) -> Result<()> { max_lots: args.max_lots, max_events: args.max_events, latency_ns: args.latency_ns, + kelly_frac_floor: args.kelly_frac_floor, + sharpe_weight_floor: args.sharpe_weight_floor, strategies: strategy_grid, }; diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index f5e61e296..c7677a7bf 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -27,6 +27,15 @@ extern "C" __global__ void decision_policy_default( float target_annual_vol_units, float annualisation_factor, int max_lots, + // Cold-start floors per pearl_blend_formulas_must_have_permanent_floor: + // max(floor, real), not blend; pearl_kelly_cap_signal_driven_floors. + // When isv_kelly state is at sentinel (no closed trades), kelly_frac + // falls back to `kelly_frac_floor`; the aggregate weight floor lets + // cross-horizon sum produce a non-zero size before recent_sharpe is + // populated. Floors are non-zero by contract — passing 0.0 reverts + // to the old sentinel-skip behaviour. + float kelly_frac_floor, + float sharpe_weight_floor, int n_backtests ) { int b = blockIdx.x; @@ -47,25 +56,41 @@ extern "C" __global__ void decision_policy_default( const float p_h = alpha_probs[h]; const IsvKellyState& s = isv[h]; - // Per-spec §5: no floor, cap derives from ISV. - // If pnl_ema_win is sentinel (0), no kelly_frac estimate available. - float ss = 0.0f; - if (s.pnl_ema_win > 1e-9f && s.realised_return_var > 1e-9f) { - const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; // ∈ [0, 1] - const float dir = (p_h > 0.5f) ? 1.0f : -1.0f; - // Kelly fraction; clamp [0, 1]. - float kelly_frac = (s.win_rate_ema * s.pnl_ema_win - - (1.0f - s.win_rate_ema) * s.pnl_ema_loss) - / s.pnl_ema_win; - kelly_frac = fmaxf(0.0f, fminf(1.0f, kelly_frac)); - // Cap units from realized variance + target vol budget. + const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; // ∈ [0, 1] — always derivable + const float dir = (p_h > 0.5f) ? 1.0f : -1.0f; + + // Kelly fraction with floor. When pnl_ema_win is sentinel, the + // ratio is undefined → floor directly. Otherwise compute the + // realised Kelly and take max(floor, kelly) so the floor never + // drives sizing down once we have signal. + float kelly_frac; + if (s.pnl_ema_win > 1e-9f) { + float kf = (s.win_rate_ema * s.pnl_ema_win + - (1.0f - s.win_rate_ema) * s.pnl_ema_loss) + / s.pnl_ema_win; + kf = fmaxf(0.0f, fminf(1.0f, kf)); + kelly_frac = fmaxf(kelly_frac_floor, kf); + } else { + kelly_frac = kelly_frac_floor; + } + + // Cap units: use realised variance when available, otherwise trust + // the host-supplied `max_lots` cap (cold-start cannot derive a + // variance-based cap; the same `max_lots` already bounds it). + float cap_lots; + if (s.realised_return_var > 1e-9f) { const float cap_units = target_annual_vol_units / sqrtf(s.realised_return_var * annualisation_factor); - const float cap_lots = fminf(cap_units, (float)max_lots); - ss = dir * sig_mag * kelly_frac * cap_lots; + cap_lots = fminf(cap_units, (float)max_lots); + } else { + cap_lots = (float)max_lots; } - signed_sizes[h] = ss; - const float w = fmaxf(0.0f, s.recent_sharpe); + + signed_sizes[h] = dir * sig_mag * kelly_frac * cap_lots; + // Recent-Sharpe weight with floor: lets cold-start aggregate fire + // (uniform weights = sharpe_weight_floor everywhere) until one + // horizon's recent_sharpe rises above the floor and dominates. + const float w = fmaxf(sharpe_weight_floor, s.recent_sharpe); weights[h] = w; w_sum += w; } @@ -149,6 +174,10 @@ extern "C" __global__ void decision_policy_program( float annualisation_factor, int max_instructions, int max_lots, + // Cold-start floors — see decision_policy_default header for the + // contract. Both kernels share the same Kelly + weight floor pattern. + float kelly_frac_floor, + float sharpe_weight_floor, int n_backtests ) { int b = blockIdx.x; @@ -203,20 +232,28 @@ extern "C" __global__ void decision_policy_program( if (h < 0 || h >= N_HORIZONS) { stack_val[sp]=0.0f; stack_mask[sp]=0u; sp+=1; break; } const float p_h = alpha_probs[h]; const IsvKellyState& s = isv[h]; - float ss = 0.0f; - if (s.pnl_ema_win > 1e-9f && s.realised_return_var > 1e-9f) { - const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; - const float dir = (p_h > 0.5f) ? 1.0f : -1.0f; - float kelly_frac = (s.win_rate_ema * s.pnl_ema_win - - (1.0f - s.win_rate_ema) * s.pnl_ema_loss) - / s.pnl_ema_win; - kelly_frac = fmaxf(0.0f, fminf(1.0f, kelly_frac)); + const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; + const float dir = (p_h > 0.5f) ? 1.0f : -1.0f; + float kelly_frac; + if (s.pnl_ema_win > 1e-9f) { + float kf = (s.win_rate_ema * s.pnl_ema_win + - (1.0f - s.win_rate_ema) * s.pnl_ema_loss) + / s.pnl_ema_win; + kf = fmaxf(0.0f, fminf(1.0f, kf)); + kelly_frac = fmaxf(kelly_frac_floor, kf); + } else { + kelly_frac = kelly_frac_floor; + } + const float effective_cap = fminf((float)leaf_max_lots, (float)max_lots); + float cap_lots; + if (s.realised_return_var > 1e-9f) { const float cap_units = target_annual_vol_units / sqrtf(s.realised_return_var * annualisation_factor); - const float effective_cap = fminf((float)leaf_max_lots, (float)max_lots); - const float cap_lots = fminf(cap_units, effective_cap); - ss = dir * sig_mag * kelly_frac * cap_lots; + cap_lots = fminf(cap_units, effective_cap); + } else { + cap_lots = effective_cap; } + const float ss = dir * sig_mag * kelly_frac * cap_lots; stack_val[sp] = ss; stack_mask[sp] = 1u << h; sp += 1; @@ -255,7 +292,9 @@ extern "C" __global__ void decision_policy_program( int h = 0; unsigned int mm = m; while ((mm & 1u) == 0u && h < N_HORIZONS) { mm >>= 1; h += 1; } - w = fmaxf(0.0f, isv[h].recent_sharpe); + // Floor lets cold-start aggregate; matches the + // decision_policy_default weight pattern. + w = fmaxf(sharpe_weight_floor, isv[h].recent_sharpe); } else { w = 1.0f; } diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index 6113c42ac..bcf7ba486 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -45,6 +45,17 @@ pub struct BacktestHarnessConfig { /// with `active=2`/`arrival_ts_ns = current + latency_ns` so the /// book can move under them during the in-flight window. pub latency_ns: u32, + /// Cold-start Kelly fraction floor — used when isv_kelly is at + /// sentinel (no closed trades yet). Per + /// pearl_blend_formulas_must_have_permanent_floor: max(floor, real), + /// not blend. Default 0.10 = trade conservatively until Kelly state + /// bootstraps. + pub kelly_frac_floor: f32, + /// Cold-start Sharpe weight floor — applied to recent_sharpe when + /// aggregating signed-sizes across horizons. Lets the cross-horizon + /// sum produce a non-zero size before recent_sharpe is populated. + /// Default 0.10 = uniform 1/5 weight per horizon at cold-start. + pub sharpe_weight_floor: f32, } pub struct BacktestHarness { @@ -175,6 +186,8 @@ impl BacktestHarness { self.cfg.annualisation_factor, self.cfg.max_lots, self.cfg.latency_ns, + self.cfg.kelly_frac_floor, + self.cfg.sharpe_weight_floor, )?; self.decision_count += 1; total_decisions += 1; diff --git a/crates/ml-backtesting/src/sim.rs b/crates/ml-backtesting/src/sim.rs index 55c8941f6..3722c1419 100644 --- a/crates/ml-backtesting/src/sim.rs +++ b/crates/ml-backtesting/src/sim.rs @@ -470,6 +470,21 @@ impl LobSimCuda { Ok(pos) } + /// Read back the market_target written by `step_decision*` for one + /// backtest. Returns `(side, abs_size)` where `side ∈ {0=buy, 1=sell, 2=noop}`. + /// Used by integration tests that need to inspect decision-kernel output + /// without going through the full submit/fill chain. + pub fn read_market_target(&self, backtest_idx: usize) -> Result<(i32, i32)> { + anyhow::ensure!( + backtest_idx < self.n_backtests, + "backtest_idx {} >= n_backtests {}", + backtest_idx, self.n_backtests + ); + let mut raw = vec![0i32; self.n_backtests * 2]; + self.stream.memcpy_dtoh(&self.market_targets_d, raw.as_mut_slice())?; + Ok((raw[backtest_idx * 2], raw[backtest_idx * 2 + 1])) + } + /// Broadcast alpha probs (per-horizon) to the decision kernel. /// Single source of truth for v1: all N backtests see the same probs. pub fn broadcast_alpha(&mut self, probs: &[f32; 5]) -> Result<()> { @@ -478,13 +493,15 @@ impl LobSimCuda { } /// Backward-compat wrapper for the no-latency immediate-match path. - /// Equivalent to `step_decision_with_latency(.., 0)`. + /// Equivalent to `step_decision_with_latency(.., 0, .., ..)`. pub fn step_decision( &mut self, current_ts_ns: u64, target_annual_vol_units: f32, annualisation_factor: f32, max_lots: u16, + kelly_frac_floor: f32, + sharpe_weight_floor: f32, ) -> Result<()> { self.step_decision_with_latency( current_ts_ns, @@ -492,6 +509,8 @@ impl LobSimCuda { annualisation_factor, max_lots, 0, + kelly_frac_floor, + sharpe_weight_floor, ) } @@ -520,6 +539,8 @@ impl LobSimCuda { annualisation_factor: f32, max_lots: u16, latency_ns: u32, + kelly_frac_floor: f32, + sharpe_weight_floor: f32, ) -> Result<()> { // Step 1: decision kernels write market_targets + open_horizon_mask. // 1a — bytecode program kernel handles backtests with plen > 0. @@ -549,6 +570,8 @@ impl LobSimCuda { .arg(&annualisation_factor) .arg(&max_instrs) .arg(&max_lots_i) + .arg(&kelly_frac_floor) + .arg(&sharpe_weight_floor) .arg(&n) .launch(cfg)?; } @@ -565,6 +588,8 @@ impl LobSimCuda { .arg(&target_annual_vol_units) .arg(&annualisation_factor) .arg(&max_lots_i) + .arg(&kelly_frac_floor) + .arg(&sharpe_weight_floor) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/decision_floor_coldstart.rs b/crates/ml-backtesting/tests/decision_floor_coldstart.rs new file mode 100644 index 000000000..0d2b56179 --- /dev/null +++ b/crates/ml-backtesting/tests/decision_floor_coldstart.rs @@ -0,0 +1,71 @@ +//! Regression test for the cold-start sentinel-skip bug surfaced during +//! the trunk-grows smoke (2026-05-19). Before the kernel floor was +//! added, `step_decision*` would observe sentinel `isv_kelly_d` (all +//! zeros from `alloc_zeros`) and skip every horizon → `market_target = +//! (noop, 0)` forever, which prevented any trade from ever firing. +//! +//! Per `pearl_blend_formulas_must_have_permanent_floor.md` and +//! `pearl_kelly_cap_signal_driven_floors.md`: the decision policy uses +//! `max(floor, computed)` on Kelly fraction AND on the recent-Sharpe +//! aggregation weight so cold-start produces a non-zero target. +//! +//! Asserts: with sentinel isv_kelly_d, strong directional alpha +//! (p_h=0.8 ⇒ long) + floors > 0 ⇒ market_target side=0 (buy), +//! size >= 1 lot. + +use anyhow::Result; +use ml_backtesting::sim::LobSimCuda; +use ml_core::device::MlDevice; + +#[test] +#[ignore = "requires CUDA"] +fn cold_start_sentinel_state_still_fires_a_trade() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("skipping: cuda device unavailable ({e})"); + return Ok(()); + } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + // Do NOT seed isv_kelly — leave at zeros (alloc_zeros' sentinel). + // Strong directional alpha across all horizons → conviction-driven + // sig_mag = 0.6 for every horizon, dir = +1. + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency( + 0, // current_ts_ns + 50.0, // target_annual_vol_units + 825.0, // annualisation_factor + 5, // max_lots + 0, // latency_ns (immediate path; not relevant — we only read targets) + 0.20, // kelly_frac_floor — the fix under test (default) + 0.10, // sharpe_weight_floor + )?; + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 0, "cold-start with p_h=0.8 must produce a long; got side={side}"); + assert!(size >= 1, "cold-start size {size} < 1 — the kernel floor isn't firing"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn cold_start_with_zero_floor_reproduces_old_bug() -> Result<()> { + // Mirror of the test above but with floors = 0 — the kernel must + // then behave like the old sentinel-skip pre-fix code: no trade ever + // fires. Lets us prove the fix actually changes behaviour (and not + // some other unrelated code path). + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("skipping: cuda device unavailable ({e})"); + return Ok(()); + } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency(0, 50.0, 825.0, 5, 0, 0.0, 0.0)?; + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 2, "with zero floors, sentinel state must still skip (side=noop)"); + assert_eq!(size, 0); + Ok(()) +} diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs index 456543eb4..e836338b4 100644 --- a/crates/ml-backtesting/tests/lob_sim_fixtures.rs +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -248,7 +248,14 @@ fn run_book_fixture(path: &Path) -> Result<()> { ts_ns, } => { sim.broadcast_alpha(alpha_probs)?; - sim.step_decision(*ts_ns, *target_annual_vol_units, *annualisation_factor, *max_lots)?; + sim.step_decision( + *ts_ns, + *target_annual_vol_units, + *annualisation_factor, + *max_lots, + 0.10, // kelly_frac_floor — cold-start prior + 0.10, // sharpe_weight_floor + )?; } FixtureEvent::SeedLimit { backtest_idx, slot_idx, side, kind, active_state, diff --git a/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs b/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs index fe2b732ea..985569305 100644 --- a/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs +++ b/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs @@ -126,6 +126,8 @@ fn run_integrated_fuzz(n_backtests: usize, n_events: usize, seed: u64) -> Result rng.gen_range(500.0..1200.0), // annualisation_factor rng.gen_range(2..6), // max_lots latency, + 0.10, // kelly_frac_floor + 0.10, // sharpe_weight_floor )?; decision_idx += 1; } diff --git a/infra/k8s/argo/lob-backtest-sweep-template.yaml b/infra/k8s/argo/lob-backtest-sweep-template.yaml index 6da6d2832..f4b673051 100644 --- a/infra/k8s/argo/lob-backtest-sweep-template.yaml +++ b/infra/k8s/argo/lob-backtest-sweep-template.yaml @@ -307,16 +307,21 @@ spec: echo "=== cell $CELL done ===" ls -lh "$OUT" - # ── aggregate: final CPU pod runs fxt-backtest aggregate ─────── + # ── aggregate: runs fxt-backtest aggregate on the GPU pool. ─── + # The binary is dynamically linked against libcuda.so.1, which the + # ci-compile-cpu pool's host doesn't expose; the aggregate workload + # itself is CPU-only (~seconds) but the binary needs the driver libs. + # Reusing the GPU pool is cheap (single fast pod) and avoids needing + # a separate CPU-only binary. - name: aggregate inputs: parameters: - name: sha nodeSelector: - k8s.scaleway.com/pool-name: ci-compile-cpu + k8s.scaleway.com/pool-name: "{{workflow.parameters.gpu-pool}}" topology.kubernetes.io/zone: fr-par-2 tolerations: - - key: node.cilium.io/agent-not-ready + - key: nvidia.com/gpu operator: Exists effect: NoSchedule container: