From ed19985c9546b4f04e5b2f37a18e24bfe928c38e Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 18 May 2026 09:26:22 +0200 Subject: [PATCH] feat(ml-backtesting): bytecode VM dispatcher for Strategy compositions (C13) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The hardcoded WeightedByRealizedSharpe path from C7 is now joined by a stack-based bytecode interpreter that consumes Strategy::flatten() output, unlocking RegimeSwitch / Portfolio / non-default Ensemble / single-Leaf compositions specified via policy-grid YAML. cuda/decision_policy.cu — new kernel `decision_policy_program`: Stack-based VM with parallel (value, attribution_mask) stacks. Opcodes mirror src/policy/mod.rs::OpCode exactly: NoOp / PushScalar / EvalRegime / BranchIfRegime EmitPerHorizonSize (computes sized intent from alpha[h] + IsvKellyState[h] using same Kelly + ISV-cap formula as the hardcoded default) AggMean / AggWeightedSharpe / AggMaxConfidence (pop n values, push aggregated, OR attribution masks) ApplyConflict (v1 no-op, reserved for Portfolio) WriteOrder (terminal — converts top-of-stack to market_target + open_horizon_masks attribution if currently flat) AggWeightedSharpe recovers the source horizon from a single-bit attribution mask to look up recent_sharpe; multi-bit masks (nested aggregators that collapsed horizons) fall back to uniform weight. decision_policy_default extended with a program_lens param: skips any backtest whose plen > 0 (the program kernel handled it). The two kernels run sequentially in step_decision_with_latency with mutual exclusivity on each backtest slot. LobSimCuda gains: upload_program(b, &Program) — uploads a Strategy::flatten() output to backtest b's slot in program_table_d, updates program_lens_d. set_regime(b, regime_id) — writes regimes_d for OP_EVAL_REGIME / OP_BRANCH_IF_REGIME consumption. BacktestHarnessConfig.strategies (Vec) — empty means every cell uses the hardcoded default; non-empty len must equal n_parallel and each strategy is flattened + uploaded at construction. bin/fxt-backtest --policy-grid path now actually plumbs through to the kernel (was parsed-but-ignored in C9). Empty grid keeps the default behaviour. New fixture decision_program_h4_only: uploads Strategy::Leaf(h4_only) flattened to (EmitPerHorizonSize, WriteOrder) — 2 instructions, 16 bytes — and verifies the bytecode kernel produces equivalent end-state to the hardcoded default (3 lot buy at vwap=5500). Proves the VM dispatcher works end-to-end. 12/12 GPU fixtures green. 33 lib unit tests + 3 Ring 2 fuzz tests still green. Co-Authored-By: Claude Opus 4.7 --- bin/fxt-backtest/src/main.rs | 15 +- crates/ml-backtesting/cuda/decision_policy.cu | 214 ++++++++++++++++++ crates/ml-backtesting/src/harness.rs | 22 +- crates/ml-backtesting/src/lob/mod.rs | 7 + crates/ml-backtesting/src/sim.rs | 99 +++++++- .../fixtures/decision_program_h4_only.json | 40 ++++ .../ml-backtesting/tests/lob_sim_fixtures.rs | 21 ++ 7 files changed, 411 insertions(+), 7 deletions(-) create mode 100644 crates/ml-backtesting/tests/fixtures/decision_program_h4_only.json diff --git a/bin/fxt-backtest/src/main.rs b/bin/fxt-backtest/src/main.rs index c52d2cf2b..3b882494c 100644 --- a/bin/fxt-backtest/src/main.rs +++ b/bin/fxt-backtest/src/main.rs @@ -129,11 +129,15 @@ fn run(args: RunArgs) -> Result<()> { .map(|_| Strategy::default_for(args.max_lots)) .collect() }; - // Strategies parsed but not yet plumbed into the v1 LobSimCuda which - // hardcodes the WeightedByRealizedSharpe default policy (see C7 - // commit message). Keep the parse step so the YAML path is validated - // end-to-end and ready when the v2 bytecode VM lands. - let _ = strategies; + // If the user provided a policy grid YAML, strategies will be flattened + // and uploaded to the per-backtest bytecode-program slots in the sim; + // otherwise each cell uses the hardcoded default policy + // (WeightedByRealizedSharpe across all 5 horizons). + let strategy_grid = if args.policy_grid.is_some() { + strategies + } else { + Vec::new() + }; let dev = MlDevice::cuda(0).map_err(|e| anyhow::anyhow!("MlDevice::cuda(0): {e}"))?; let trunk = CfcTrunk::new_random(&dev, &CfcConfig::default(), args.seed) @@ -149,6 +153,7 @@ fn run(args: RunArgs) -> Result<()> { max_lots: args.max_lots, max_events: args.max_events, latency_ns: args.latency_ns, + strategies: strategy_grid, }; std::fs::create_dir_all(&args.out) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 268d6c1fb..f5e61e296 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -21,6 +21,7 @@ extern "C" __global__ void decision_policy_default( const float* __restrict__ alpha_probs, // [N_HORIZONS] broadcast const unsigned char* __restrict__ isv_kelly_base, // [n_backtests * N_HORIZONS * sizeof(IsvKellyState)] const Pos* __restrict__ positions, // [n_backtests] + const unsigned int* __restrict__ program_lens, // [n_backtests] — non-zero ⇒ skip (program kernel handles it) int* __restrict__ market_targets, // [n_backtests * 2] — written unsigned int* __restrict__ open_horizon_masks, // [n_backtests] — written (entry attribution) float target_annual_vol_units, @@ -30,6 +31,7 @@ extern "C" __global__ void decision_policy_default( ) { int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; + if (program_lens[b] != 0u) return; // backtest b uses a custom bytecode program; skip default const IsvKellyState* isv = reinterpret_cast( isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) @@ -101,6 +103,218 @@ extern "C" __global__ void decision_policy_default( market_targets[b * 2 + 1] = abs_sz; } +// Bytecode VM — interprets the flat Program emitted by Strategy::flatten() +// (Rust src/policy/mod.rs) against current alpha probs + per-horizon +// IsvKellyState. Replaces the hardcoded default policy when the host +// has uploaded a non-empty program for backtest b. +// +// Stack layout: parallel float-value + u32-attribution-mask stacks. +// Each EmitPerHorizonSize pushes its (signed_size, mask=(1<= n_backtests || threadIdx.x != 0) return; + + const unsigned int plen = program_lens[b]; + if (plen == 0u) { + // Skip — this backtest uses the hardcoded default kernel separately. + return; + } + const Instruction* prog = programs + (size_t)b * max_instructions; + const IsvKellyState* isv = reinterpret_cast( + isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) + ); + const unsigned int my_regime = regimes[b]; + + float stack_val[STACK_CAP]; + unsigned int stack_mask[STACK_CAP]; + int sp = 0; + + unsigned int pc = 0; + while (pc < plen && pc < (unsigned int)max_instructions) { + const Instruction ins = prog[pc]; + pc += 1; + switch (ins.op) { + case OP_NOOP: break; + case OP_PUSH_SCALAR: { + if (sp >= STACK_CAP) break; + const float v = __int_as_float((int)ins.arg2); + stack_val[sp] = v; + stack_mask[sp] = 0u; + sp += 1; + break; + } + case OP_EVAL_REGIME: { + if (sp >= STACK_CAP) break; + stack_val[sp] = (my_regime == (unsigned int)ins.arg0) ? 1.0f : 0.0f; + stack_mask[sp] = 0u; + sp += 1; + break; + } + case OP_BRANCH_IF_REGIME: { + if (my_regime != (unsigned int)ins.arg0) { + pc = ins.arg2; + } + break; + } + case OP_EMIT_PER_HORIZON_SIZE: { + if (sp >= STACK_CAP) break; + const int h = (int)ins.arg0; + const int leaf_max_lots = (int)ins.arg1; + 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 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; + } + stack_val[sp] = ss; + stack_mask[sp] = 1u << h; + sp += 1; + break; + } + case OP_AGG_MEAN: { + const int n = (int)ins.arg0; + if (n <= 0 || sp < n) break; + float sum = 0.0f; + unsigned int mask_or = 0u; + for (int i = sp - n; i < sp; ++i) { + sum += stack_val[i]; + mask_or |= stack_mask[i]; + } + sp -= n; + stack_val[sp] = sum / (float)n; + stack_mask[sp] = mask_or; + sp += 1; + break; + } + case OP_AGG_WEIGHTED_SHARPE: { + const int n = (int)ins.arg0; + if (n <= 0 || sp < n) break; + float w_sum = 0.0f, wx_sum = 0.0f; + unsigned int mask_or = 0u; + for (int i = sp - n; i < sp; ++i) { + const unsigned int m = stack_mask[i]; + // Look up recent_sharpe for the single horizon this slot + // attributes to. If the mask has multiple bits set (a + // nested aggregator already merged horizons) we fall back + // to uniform weight 1.0 because per-horizon attribution + // collapsed. + float w; + if (m != 0u && (m & (m - 1u)) == 0u) { + // Single-bit mask — recover horizon index. + 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); + } else { + w = 1.0f; + } + w_sum += w; + wx_sum += w * stack_val[i]; + mask_or |= m; + } + sp -= n; + stack_val[sp] = (w_sum > 1e-9f) ? (wx_sum / w_sum) : 0.0f; + stack_mask[sp] = mask_or; + sp += 1; + break; + } + case OP_AGG_MAX_CONFIDENCE: { + const int n = (int)ins.arg0; + if (n <= 0 || sp < n) break; + int best = sp - n; + for (int i = sp - n + 1; i < sp; ++i) { + if (fabsf(stack_val[i]) > fabsf(stack_val[best])) best = i; + } + const float bv = stack_val[best]; + const unsigned int bm = stack_mask[best]; + sp -= n; + stack_val[sp] = bv; + stack_mask[sp] = bm; + sp += 1; + break; + } + case OP_APPLY_CONFLICT: { + // v1: no-op (mean of children already aggregated; conflict + // resolution only matters when multiple cells share state, + // which v1 doesn't). Preserved for v2 Portfolio mode. + break; + } + case OP_WRITE_ORDER: { + if (sp <= 0) break; + sp -= 1; + const float final_size = stack_val[sp]; + const unsigned int attribution = stack_mask[sp]; + int lots = (int)truncf(final_size + (final_size >= 0.0f ? 0.5f : -0.5f)); + if (lots == 0) { + market_targets[b * 2 + 0] = 2; // no-op + market_targets[b * 2 + 1] = 0; + } else { + if (positions[b].position_lots == 0) { + open_horizon_masks[b] = attribution; + } + const int side = (lots > 0) ? 0 : 1; + const int abs_sz = (lots > 0) ? lots : -lots; + market_targets[b * 2 + 0] = side; + market_targets[b * 2 + 1] = abs_sz; + } + break; + } + default: break; + } + } +} + // Updates per-horizon IsvKellyState[h] for every horizon flagged in // `closed_horizon_mask[b]`. Called by the host after pnl_track_step // detects close transitions. `realised_return[b]` is the closed-trade diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index c3162940c..434236a4a 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -30,6 +30,11 @@ pub struct BacktestHarnessConfig { pub max_lots: u16, /// Max number of inference inputs to consume. 0 = exhaust loader. pub max_events: u64, + /// Per-backtest Strategy compositions. Empty = every cell uses the + /// hardcoded default policy (WeightedByRealizedSharpe across all 5 + /// horizons). Non-empty: len MUST equal n_parallel; each entry is + /// flattened to bytecode and uploaded to its backtest slot. + pub strategies: Vec, /// Total submission→fill latency in nanoseconds. 0 = immediate /// (legacy path). 100_000_000 = 100ms IBKR + Scaleway baseline. /// When non-zero, the orchestrator submits aggressive IOC limits @@ -59,6 +64,14 @@ impl BacktestHarness { dev: &MlDevice, mut trunk: CfcTrunk, ) -> Result { + if !cfg.strategies.is_empty() { + anyhow::ensure!( + cfg.strategies.len() == cfg.n_parallel, + "strategies len {} ≠ n_parallel {}", + cfg.strategies.len(), + cfg.n_parallel + ); + } let files = discover_mbp10_files_sorted(&cfg.data_root)?; let loader_cfg = MultiHorizonLoaderConfig { files, @@ -77,7 +90,14 @@ impl BacktestHarness { let first = loader.peek_first().context("peek_first")?; trunk.capture_graph_a(&first).context("capture_graph_a")?; - let sim = LobSimCuda::new(cfg.n_parallel, dev)?; + let mut sim = LobSimCuda::new(cfg.n_parallel, dev)?; + + // Upload per-cell bytecode programs if the caller provided any. + // Empty Vec means every cell uses the hardcoded default policy. + for (b, strat) in cfg.strategies.iter().enumerate() { + let prog = strat.flatten(); + sim.upload_program(b, &prog)?; + } let pnl_curves = (0..cfg.n_parallel).map(|_| Vec::with_capacity(1024)).collect(); Ok(Self { diff --git a/crates/ml-backtesting/src/lob/mod.rs b/crates/ml-backtesting/src/lob/mod.rs index e2b6d55d3..b142f885f 100644 --- a/crates/ml-backtesting/src/lob/mod.rs +++ b/crates/ml-backtesting/src/lob/mod.rs @@ -22,6 +22,13 @@ pub const OPEN_TRADE_STATE_BYTES: usize = 24; pub const ORDER_EVENT_BYTES: usize = 24; /// Max OrderEvent records buffered per backtest in the audit ring. pub const RING_CAP_EVENTS: usize = 256; +/// Max bytecode Instructions per backtest program. The Strategy::flatten() +/// emission size is bounded by ~(N_HORIZONS + 2) for the default policy +/// and scales linearly with composition tree depth; 64 covers a 3-deep +/// RegimeSwitch×Portfolio×Ensemble nesting comfortably. +pub const MAX_PROGRAM_INSTRUCTIONS: usize = 64; +/// Bytes per Instruction — matches src/policy/mod.rs (u8 + u8 + u16 + u32 = 8). +pub const INSTRUCTION_BYTES: usize = 8; /// Flat host-side mirror of `cuda/lob_state.cuh::Book`. #[derive(Clone, Copy, Debug, Default, PartialEq)] diff --git a/crates/ml-backtesting/src/sim.rs b/crates/ml-backtesting/src/sim.rs index d5a1922f1..55c8941f6 100644 --- a/crates/ml-backtesting/src/sim.rs +++ b/crates/ml-backtesting/src/sim.rs @@ -37,6 +37,7 @@ pub struct LobSimCuda { submit_market_fn: cudarc::driver::CudaFunction, pnl_track_fn: cudarc::driver::CudaFunction, decision_fn: cudarc::driver::CudaFunction, + decision_program_fn: cudarc::driver::CudaFunction, isv_kelly_update_fn: cudarc::driver::CudaFunction, resting_orders_fn: cudarc::driver::CudaFunction, seed_limit_fn: cudarc::driver::CudaFunction, @@ -68,6 +69,11 @@ pub struct LobSimCuda { // Audit ring for OrderEvent emission (deferred from C2; needed by C11). audit_d: CudaSlice, // [n_backtests * RING_CAP_EVENTS * ORDER_EVENT_BYTES] audit_head_d: CudaSlice, // [n_backtests] + + // Bytecode program upload (C13 — per-backtest custom compositions). + program_table_d: CudaSlice, // [n_backtests * MAX_PROGRAM_INSTRUCTIONS * INSTRUCTION_BYTES] + program_lens_d: CudaSlice, // [n_backtests] + regimes_d: CudaSlice, // [n_backtests] — host-set current regime id } impl LobSimCuda { @@ -100,6 +106,9 @@ impl LobSimCuda { let decision_fn = decision_module .load_function("decision_policy_default") .context("load decision_policy_default")?; + let decision_program_fn = decision_module + .load_function("decision_policy_program") + .context("load decision_policy_program")?; let isv_kelly_update_fn = decision_module .load_function("isv_kelly_update_on_close") .context("load isv_kelly_update_on_close")?; @@ -169,6 +178,16 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc audit_head_d")?; + let program_table_d = stream + .alloc_zeros::(n_backtests * crate::lob::MAX_PROGRAM_INSTRUCTIONS * crate::lob::INSTRUCTION_BYTES) + .context("alloc program_table_d")?; + let program_lens_d = stream + .alloc_zeros::(n_backtests) + .context("alloc program_lens_d")?; + let regimes_d = stream + .alloc_zeros::(n_backtests) + .context("alloc regimes_d")?; + Ok(Self { n_backtests, stream, @@ -176,6 +195,7 @@ impl LobSimCuda { submit_market_fn, pnl_track_fn, decision_fn, + decision_program_fn, isv_kelly_update_fn, resting_orders_fn, seed_limit_fn, @@ -199,9 +219,61 @@ impl LobSimCuda { overflow_flag_d, audit_d, audit_head_d, + program_table_d, + program_lens_d, + regimes_d, }) } + /// Upload a flattened bytecode Program for one backtest. The kernel + /// `decision_policy_program` interprets it if program_lens[b] > 0; + /// otherwise it returns immediately and the host orchestrator falls + /// back to `decision_policy_default`. + pub fn upload_program( + &mut self, + backtest_idx: usize, + program: &crate::policy::Program, + ) -> Result<()> { + anyhow::ensure!(backtest_idx < self.n_backtests); + anyhow::ensure!( + program.instructions.len() <= crate::lob::MAX_PROGRAM_INSTRUCTIONS, + "program length {} exceeds MAX_PROGRAM_INSTRUCTIONS={}", + program.instructions.len(), + crate::lob::MAX_PROGRAM_INSTRUCTIONS + ); + + // Read back current program_table, mutate slot, push back. + let stride = crate::lob::MAX_PROGRAM_INSTRUCTIONS * crate::lob::INSTRUCTION_BYTES; + let mut raw = vec![0u8; self.n_backtests * stride]; + self.stream.memcpy_dtoh(&self.program_table_d, raw.as_mut_slice())?; + let off = backtest_idx * stride; + let bytes: &[u8] = bytemuck::cast_slice(program.instructions.as_slice()); + raw[off..off + bytes.len()].copy_from_slice(bytes); + // Zero out trailing slots in this backtest's row beyond program length. + for b in &mut raw[off + bytes.len()..off + stride] { + *b = 0; + } + self.stream.memcpy_htod(raw.as_slice(), &mut self.program_table_d)?; + + // Update program_lens[b]. + let mut lens = vec![0u32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.program_lens_d, lens.as_mut_slice())?; + lens[backtest_idx] = program.instructions.len() as u32; + self.stream.memcpy_htod(lens.as_slice(), &mut self.program_lens_d)?; + Ok(()) + } + + /// Set the current regime id for a backtest. Read by + /// `decision_policy_program` for OP_EVAL_REGIME / OP_BRANCH_IF_REGIME. + pub fn set_regime(&mut self, backtest_idx: usize, regime_id: u32) -> Result<()> { + anyhow::ensure!(backtest_idx < self.n_backtests); + let mut regimes = vec![0u32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.regimes_d, regimes.as_mut_slice())?; + regimes[backtest_idx] = regime_id; + self.stream.memcpy_htod(regimes.as_slice(), &mut self.regimes_d)?; + Ok(()) + } + /// Read the audit-ring head count + records for one backtest. /// Returns the raw bytes; tests parse via bytemuck::pod_read_unaligned. pub fn read_audit_records( @@ -449,7 +521,9 @@ impl LobSimCuda { max_lots: u16, latency_ns: u32, ) -> Result<()> { - // Step 1: decision kernel writes market_targets + open_horizon_mask. + // Step 1: decision kernels write market_targets + open_horizon_mask. + // 1a — bytecode program kernel handles backtests with plen > 0. + // 1b — default hardcoded kernel handles backtests with plen == 0. let cfg = LaunchConfig { grid_dim: (self.n_backtests as u32, 1, 1), block_dim: (32, 1, 1), @@ -457,12 +531,35 @@ impl LobSimCuda { }; let n = self.n_backtests as i32; let max_lots_i = max_lots as i32; + let max_instrs = crate::lob::MAX_PROGRAM_INSTRUCTIONS as i32; + + // 1a: bytecode VM kernel (skips backtests with plen==0 internally). + let mut launch = self.stream.launch_builder(&self.decision_program_fn); + unsafe { + launch + .arg(&self.alpha_probs_d) + .arg(&self.isv_kelly_d) + .arg(&self.pos_d) + .arg(&self.program_table_d) + .arg(&self.program_lens_d) + .arg(&self.regimes_d) + .arg(&mut self.market_targets_d) + .arg(&mut self.open_horizon_mask_d) + .arg(&target_annual_vol_units) + .arg(&annualisation_factor) + .arg(&max_instrs) + .arg(&max_lots_i) + .arg(&n) + .launch(cfg)?; + } + // 1b: default kernel (skips backtests with plen>0 internally). let mut launch = self.stream.launch_builder(&self.decision_fn); unsafe { launch .arg(&self.alpha_probs_d) .arg(&self.isv_kelly_d) .arg(&self.pos_d) + .arg(&self.program_lens_d) .arg(&mut self.market_targets_d) .arg(&mut self.open_horizon_mask_d) .arg(&target_annual_vol_units) diff --git a/crates/ml-backtesting/tests/fixtures/decision_program_h4_only.json b/crates/ml-backtesting/tests/fixtures/decision_program_h4_only.json new file mode 100644 index 000000000..ba00e1fe0 --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/decision_program_h4_only.json @@ -0,0 +1,40 @@ +{ + "name": "decision_program_h4_only", + "n_backtests": 1, + "comment": "Uploads a custom Strategy bytecode program (single Leaf at horizon 4) that bypasses the hardcoded WeightedByRealizedSharpe default. Should produce equivalent end-state when h4 is the only weighted horizon.", + "warm_start_isv_kelly": [ + [ + { "pnl_ema_win": 0.0, "pnl_ema_loss": 0.0, "win_rate_ema": 0.0, "n_trades_seen": 0, "realised_return_var": 0.0, "recent_sharpe": 0.0 }, + { "pnl_ema_win": 0.0, "pnl_ema_loss": 0.0, "win_rate_ema": 0.0, "n_trades_seen": 0, "realised_return_var": 0.0, "recent_sharpe": 0.0 }, + { "pnl_ema_win": 0.0, "pnl_ema_loss": 0.0, "win_rate_ema": 0.0, "n_trades_seen": 0, "realised_return_var": 0.0, "recent_sharpe": 0.0 }, + { "pnl_ema_win": 0.0, "pnl_ema_loss": 0.0, "win_rate_ema": 0.0, "n_trades_seen": 0, "realised_return_var": 0.0, "recent_sharpe": 0.0 }, + { "pnl_ema_win": 2.0, "pnl_ema_loss": 0.5, "win_rate_ema": 0.8, "n_trades_seen": 50, "realised_return_var": 0.25, "recent_sharpe": 1.0 } + ] + ], + "upload_h4_leaf_program": true, + "events": [ + { + "type": "snapshot", + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "decision", + "alpha_probs": [0.5, 0.5, 0.5, 0.5, 0.9], + "target_annual_vol_units": 50.0, + "annualisation_factor": 1.0, + "max_lots": 5, + "ts_ns": 1000000000 + } + ], + "expected_pos": [ + { + "position_lots": 3, + "vwap_entry": 5500.00, + "vwap_tol": 0.01, + "realized_pnl_tol": 0.001 + } + ] +} diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs index fdbe10ee2..456543eb4 100644 --- a/crates/ml-backtesting/tests/lob_sim_fixtures.rs +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -165,6 +165,8 @@ struct Fixture { fill_all_limit_slots: bool, #[serde(default)] expect_33rd_overflow: bool, + #[serde(default)] + upload_h4_leaf_program: bool, } fn parse_kind_limit(s: &str) -> u8 { @@ -200,6 +202,19 @@ fn run_book_fixture(path: &Path) -> Result<()> { .map_err(|e| anyhow::anyhow!("cuda device: {e}"))?; let mut sim = LobSimCuda::new(fx.n_backtests, &dev)?; + // Upload a single-leaf h4 Strategy program for backtest 0 if requested. + if fx.upload_h4_leaf_program { + use ml_backtesting::policy::{SizingPolicyId, StopRules, Strategy, StrategyConfig}; + let leaf = Strategy::Leaf(StrategyConfig { + horizon_idx: 4, + sizing_policy: SizingPolicyId::IsvKelly, + sl_tp_rules: StopRules::default(), + max_concurrent_lots: 5, + }); + let prog = leaf.flatten(); + sim.upload_program(0, &prog)?; + } + // Apply ISV-Kelly warm-start if provided (one row per backtest). for (b, states) in fx.warm_start_isv_kelly.iter().enumerate() { let host_states: [IsvKellyStateHost; 5] = [ @@ -469,3 +484,9 @@ fn fix_submission_overflow() { fn fix_latency_in_flight_miss() { run_book_fixture(Path::new("tests/fixtures/latency_in_flight_miss.json")).unwrap(); } + +#[test] +#[ignore = "requires CUDA"] +fn fix_decision_program_h4_only() { + run_book_fixture(Path::new("tests/fixtures/decision_program_h4_only.json")).unwrap(); +}