fix(ml-backtesting): cold-start Kelly + Sharpe-weight floors in decision kernel

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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-19 13:41:17 +02:00
parent d87b3a5491
commit da21feb1b1
8 changed files with 225 additions and 33 deletions

View File

@@ -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<PathBuf>,
#[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<u64>,
#[serde(default)] seed: Option<u64>,
#[serde(default)] checkpoint: Option<PathBuf>,
#[serde(default)] kelly_frac_floor: Option<f32>,
#[serde(default)] sharpe_weight_floor: Option<f32>,
}
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,
};

View File

@@ -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;
}

View File

@@ -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;

View File

@@ -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)?;
}

View File

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

View File

@@ -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,

View File

@@ -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;
}

View File

@@ -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: