feat(ml-backtesting): bytecode VM dispatcher for Strategy compositions (C13)

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<Strategy>) — 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 <yaml> 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 <noreply@anthropic.com>

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)

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<const IsvKellyState*>(
         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<<h)). The
+// aggregators pop n values + n masks, push the aggregated value with
+// OR'd mask. The terminal WriteOrder pops one value+mask, converts to
+// {side, abs_size} in market_targets[b] and records the attribution
+// mask in open_horizon_masks[b] for entry-time crediting.
+//
+// Layout MUST stay in sync with src/policy/mod.rs::{OpCode, Instruction}.
+
+#define OP_NOOP                  0
+#define OP_EVAL_REGIME           1
+#define OP_EMIT_PER_HORIZON_SIZE 2
+#define OP_AGG_WEIGHTED_SHARPE   3
+#define OP_AGG_MEAN              4
+#define OP_AGG_MAX_CONFIDENCE    5
+#define OP_APPLY_CONFLICT        6
+#define OP_WRITE_ORDER           7
+#define OP_PUSH_SCALAR           8
+#define OP_BRANCH_IF_REGIME      9
+#define STACK_CAP                32
+
+struct Instruction {
+    unsigned char  op;
+    unsigned char  arg0;
+    unsigned short arg1;
+    unsigned int   arg2;
+};
+
+extern "C" __global__ void decision_policy_program(
+    const float*       __restrict__ alpha_probs,
+    const unsigned char* __restrict__ isv_kelly_base,
+    const Pos*         __restrict__ positions,
+    const Instruction* __restrict__ programs,           // [n_backtests * max_instructions]
+    const unsigned int* __restrict__ program_lens,      // [n_backtests]
+    const unsigned int* __restrict__ regimes,           // [n_backtests] — current regime id
+    int*               __restrict__ market_targets,     // [n_backtests * 2]
+    unsigned int*      __restrict__ open_horizon_masks,
+    float                          target_annual_vol_units,
+    float                          annualisation_factor,
+    int                            max_instructions,
+    int                            max_lots,
+    int                            n_backtests
+) {
+    int b = blockIdx.x;
+    if (b >= 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<const IsvKellyState*>(
+        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

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<crate::policy::Strategy>,
     /// 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<Self> {
+        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 {

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)]

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<u8>,                 // [n_backtests * RING_CAP_EVENTS * ORDER_EVENT_BYTES]
     audit_head_d: CudaSlice<u32>,           // [n_backtests]
+
+    // Bytecode program upload (C13 — per-backtest custom compositions).
+    program_table_d: CudaSlice<u8>,         // [n_backtests * MAX_PROGRAM_INSTRUCTIONS * INSTRUCTION_BYTES]
+    program_lens_d:  CudaSlice<u32>,        // [n_backtests]
+    regimes_d:       CudaSlice<u32>,        // [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::<u32>(n_backtests)
             .context("alloc audit_head_d")?;
 
+        let program_table_d = stream
+            .alloc_zeros::<u8>(n_backtests * crate::lob::MAX_PROGRAM_INSTRUCTIONS * crate::lob::INSTRUCTION_BYTES)
+            .context("alloc program_table_d")?;
+        let program_lens_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc program_lens_d")?;
+        let regimes_d = stream
+            .alloc_zeros::<u32>(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)

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
+    }
+  ]
+}

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();
+}