diff --git a/crates/ml-backtesting/src/lib.rs b/crates/ml-backtesting/src/lib.rs
index 8e6d3db83..9bec3ff3e 100644
--- a/crates/ml-backtesting/src/lib.rs
+++ b/crates/ml-backtesting/src/lib.rs
@@ -18,6 +18,7 @@ pub use ml_core::MLError;
 pub mod action_loader;
 pub mod barrier_backtest;
 pub mod order;
+pub mod policy;
 pub mod report;
 
 pub use action_loader::{load_actions_from_csv, DQNActionRecord};
diff --git a/crates/ml-backtesting/src/policy/composition.rs b/crates/ml-backtesting/src/policy/composition.rs
new file mode 100644
index 000000000..76209dd3b
--- /dev/null
+++ b/crates/ml-backtesting/src/policy/composition.rs
@@ -0,0 +1,29 @@
+//! Aggregator + conflict + regime enums for the multi-strategy
+//! composition tree. See spec §6.
+
+use serde::{Deserialize, Serialize};
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
+pub enum EnsembleAggregator {
+    Mean,
+    MaxConfidence,
+    WeightedByValAuc,
+    /// ADAPTIVE — per-horizon recent_sharpe weights. Auto-shifts capital
+    /// toward empirically winning horizons. See spec §5 + §6.
+    WeightedByRealizedSharpe,
+}
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
+pub enum ConflictPolicy {
+    NetOut,
+    FirstSubmitWins,
+    BlockOpposing,
+}
+
+#[derive(Clone, Copy, Debug, Hash, PartialEq, Eq, Serialize, Deserialize)]
+pub enum Regime {
+    SpreadQ1,
+    SpreadQ2,
+    SpreadQ3,
+    SpreadQ4,
+}
diff --git a/crates/ml-backtesting/src/policy/latency.rs b/crates/ml-backtesting/src/policy/latency.rs
new file mode 100644
index 000000000..779091348
--- /dev/null
+++ b/crates/ml-backtesting/src/policy/latency.rs
@@ -0,0 +1,88 @@
+//! Latency model. Default: Constant(100ms) reflects IBKR + Scaleway.
+//! See spec §4.
+
+use serde::{Deserialize, Serialize};
+
+#[derive(Clone, Debug, Serialize, Deserialize)]
+pub enum LatencyConfig {
+    Constant {
+        total_ns: u32,
+    },
+    /// Per-decision latency drawn from a host-uploaded empirical
+    /// distribution. Indices wrap modulo `samples.len()`.
+    Empirical {
+        samples: Vec<f32>,
+    },
+    Decomposed {
+        inference_ns: u32,
+        decision_ns: u32,
+        wire_ns: u32,
+        match_ns: u32,
+    },
+}
+
+impl Default for LatencyConfig {
+    fn default() -> Self {
+        LatencyConfig::Constant { total_ns: 100_000_000 }
+    }
+}
+
+impl LatencyConfig {
+    /// Sample a latency in ns for the given decision-event index.
+    pub fn sample_ns(&self, decision_idx: u64) -> u32 {
+        match self {
+            LatencyConfig::Constant { total_ns } => *total_ns,
+            LatencyConfig::Empirical { samples } => {
+                if samples.is_empty() {
+                    return 0;
+                }
+                let i = (decision_idx as usize) % samples.len();
+                samples[i].max(0.0) as u32
+            }
+            LatencyConfig::Decomposed {
+                inference_ns,
+                decision_ns,
+                wire_ns,
+                match_ns,
+            } => inference_ns + decision_ns + wire_ns + match_ns,
+        }
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn default_is_100ms() {
+        assert!(matches!(
+            LatencyConfig::default(),
+            LatencyConfig::Constant { total_ns: 100_000_000 }
+        ));
+    }
+
+    #[test]
+    fn empirical_wraps() {
+        let cfg = LatencyConfig::Empirical { samples: vec![10.0, 20.0, 30.0] };
+        assert_eq!(cfg.sample_ns(0), 10);
+        assert_eq!(cfg.sample_ns(3), 10);
+        assert_eq!(cfg.sample_ns(7), 20);
+    }
+
+    #[test]
+    fn empirical_empty_returns_zero() {
+        let cfg = LatencyConfig::Empirical { samples: vec![] };
+        assert_eq!(cfg.sample_ns(0), 0);
+    }
+
+    #[test]
+    fn decomposed_sums() {
+        let cfg = LatencyConfig::Decomposed {
+            inference_ns: 1_000_000,
+            decision_ns: 100_000,
+            wire_ns: 50_000_000,
+            match_ns: 5_000_000,
+        };
+        assert_eq!(cfg.sample_ns(0), 56_100_000);
+    }
+}
diff --git a/crates/ml-backtesting/src/policy/mod.rs b/crates/ml-backtesting/src/policy/mod.rs
new file mode 100644
index 000000000..4b43ec611
--- /dev/null
+++ b/crates/ml-backtesting/src/policy/mod.rs
@@ -0,0 +1,287 @@
+//! Multi-strategy composition tree + flat bytecode program for the
+//! decision-policy kernel.
+//!
+//! See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md §6.
+
+pub mod composition;
+pub mod latency;
+pub mod sizing;
+pub mod stop_rules;
+
+pub use composition::{ConflictPolicy, EnsembleAggregator, Regime};
+pub use latency::LatencyConfig;
+pub use sizing::{IsvKellyStateHost, SizingPolicyId};
+pub use stop_rules::StopRules;
+
+use serde::{Deserialize, Serialize};
+use std::collections::HashMap;
+
+/// Mirrors `crates/ml-alpha/src/heads.rs::N_HORIZONS`.
+pub const N_HORIZONS: usize = 5;
+
+/// One horizon-anchored leaf strategy. See spec §5 (per-horizon ISV-Kelly) + §6.
+#[derive(Clone, Debug, Serialize, Deserialize)]
+pub struct StrategyConfig {
+    pub horizon_idx: u8, // 0..N_HORIZONS
+    pub sizing_policy: SizingPolicyId,
+    pub sl_tp_rules: StopRules,
+    pub max_concurrent_lots: u16,
+}
+
+/// Recursive composition tree. Default constructor `default_for(max_lots)`
+/// emits the adaptive 5-horizon Ensemble with WeightedByRealizedSharpe.
+#[derive(Clone, Debug, Serialize, Deserialize)]
+pub enum Strategy {
+    Leaf(StrategyConfig),
+    Ensemble {
+        children: Vec<Strategy>,
+        aggregator: EnsembleAggregator,
+    },
+    Portfolio {
+        children: Vec<Strategy>,
+        capital_allocation: Vec<f32>,
+        conflict_policy: ConflictPolicy,
+    },
+    RegimeSwitch {
+        regime_to_child: HashMap<Regime, Strategy>,
+    },
+}
+
+impl Strategy {
+    /// Default v1 policy: per-horizon adaptive Ensemble. No static
+    /// horizon mask; WeightedByRealizedSharpe auto-shifts capital from
+    /// in-run evidence. See spec §6 `default_strategy(...)`.
+    pub fn default_for(max_lots: u16) -> Self {
+        let children = (0..N_HORIZONS as u8)
+            .map(|h| {
+                Strategy::Leaf(StrategyConfig {
+                    horizon_idx: h,
+                    sizing_policy: SizingPolicyId::IsvKelly,
+                    sl_tp_rules: StopRules::default(),
+                    max_concurrent_lots: max_lots,
+                })
+            })
+            .collect();
+        Strategy::Ensemble {
+            children,
+            aggregator: EnsembleAggregator::WeightedByRealizedSharpe,
+        }
+    }
+
+    /// Walk the tree and emit a flat bytecode Program. See spec §6
+    /// host-side flattening; layout must stay in sync with
+    /// cuda/decision_policy.cu (added in C7).
+    pub fn flatten(&self) -> Program {
+        let mut prog = Program::default();
+        flatten_into(self, &mut prog);
+        prog.finalize();
+        prog
+    }
+}
+
+fn flatten_into(s: &Strategy, prog: &mut Program) {
+    match s {
+        Strategy::Leaf(cfg) => {
+            prog.instructions.push(Instruction {
+                op: OpCode::EmitPerHorizonSize as u8,
+                arg0: cfg.horizon_idx,
+                arg1: cfg.max_concurrent_lots,
+                arg2: 0,
+            });
+        }
+        Strategy::Ensemble { children, aggregator } => {
+            for c in children {
+                flatten_into(c, prog);
+            }
+            let agg_op = match aggregator {
+                EnsembleAggregator::Mean => OpCode::AggMean,
+                EnsembleAggregator::MaxConfidence => OpCode::AggMaxConfidence,
+                // v1: treat ValAuc as Mean (no static-weights upload yet).
+                EnsembleAggregator::WeightedByValAuc => OpCode::AggMean,
+                EnsembleAggregator::WeightedByRealizedSharpe => OpCode::AggWeightedSharpe,
+            };
+            prog.instructions.push(Instruction {
+                op: agg_op as u8,
+                arg0: children.len() as u8,
+                arg1: 0,
+                arg2: 0,
+            });
+        }
+        Strategy::Portfolio { children, capital_allocation: _, conflict_policy } => {
+            for c in children {
+                flatten_into(c, prog);
+            }
+            prog.instructions.push(Instruction {
+                op: OpCode::AggMean as u8,
+                arg0: children.len() as u8,
+                arg1: 0,
+                arg2: 0,
+            });
+            prog.instructions.push(Instruction {
+                op: OpCode::ApplyConflict as u8,
+                arg0: *conflict_policy as u8,
+                arg1: 0,
+                arg2: 0,
+            });
+        }
+        Strategy::RegimeSwitch { regime_to_child } => {
+            let mut branch_indices: Vec<usize> = Vec::with_capacity(regime_to_child.len());
+            for (regime, child) in regime_to_child {
+                let branch_idx = prog.instructions.len();
+                prog.instructions.push(Instruction {
+                    op: OpCode::BranchIfRegime as u8,
+                    arg0: *regime as u8,
+                    arg1: 0,
+                    arg2: 0, // patched below
+                });
+                flatten_into(child, prog);
+                branch_indices.push(branch_idx);
+            }
+            let end = prog.instructions.len() as u32;
+            for branch_idx in branch_indices {
+                prog.instructions[branch_idx].arg2 = end;
+            }
+        }
+    }
+}
+
+/// Bytecode opcodes. The decision-policy kernel (C7) interprets these
+/// against current alpha probs + per-horizon ISV-Kelly state to mutate
+/// shared-mem orders[]. Layout MUST stay in sync with
+/// `cuda/decision_policy.cu`. See spec §6.
+#[repr(u8)]
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub enum OpCode {
+    NoOp = 0,
+    EvalRegime = 1,
+    EmitPerHorizonSize = 2,
+    AggWeightedSharpe = 3,
+    AggMean = 4,
+    AggMaxConfidence = 5,
+    ApplyConflict = 6,
+    WriteOrder = 7,
+    PushScalar = 8,
+    BranchIfRegime = 9,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug, bytemuck::Pod, bytemuck::Zeroable)]
+pub struct Instruction {
+    pub op: u8,
+    pub arg0: u8,
+    pub arg1: u16,
+    pub arg2: u32,
+}
+
+/// Flat device-uploadable program. One per backtest cell; indexed by
+/// blockIdx.x in device-global memory (NOT __constant__, see spec §6
+/// "Why not __constant__").
+#[derive(Clone, Debug, Default)]
+pub struct Program {
+    pub instructions: Vec<Instruction>,
+}
+
+impl Program {
+    /// Append the top-level WriteOrder. Called once by Strategy::flatten().
+    pub(crate) fn finalize(&mut self) {
+        self.instructions.push(Instruction {
+            op: OpCode::WriteOrder as u8,
+            arg0: 0,
+            arg1: 0,
+            arg2: 0,
+        });
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn default_strategy_has_5_horizon_leaves() {
+        let s = Strategy::default_for(3);
+        match s {
+            Strategy::Ensemble { children, aggregator } => {
+                assert_eq!(children.len(), 5);
+                assert!(matches!(
+                    aggregator,
+                    EnsembleAggregator::WeightedByRealizedSharpe
+                ));
+                for (i, c) in children.iter().enumerate() {
+                    match c {
+                        Strategy::Leaf(cfg) => {
+                            assert_eq!(cfg.horizon_idx as usize, i);
+                            assert_eq!(cfg.max_concurrent_lots, 3);
+                        }
+                        _ => panic!("expected Leaf at index {i}"),
+                    }
+                }
+            }
+            _ => panic!("expected top-level Ensemble"),
+        }
+    }
+
+    #[test]
+    fn default_strategy_flattens_to_5_emits_plus_agg_plus_write() {
+        let s = Strategy::default_for(3);
+        let p = s.flatten();
+        assert_eq!(p.instructions.len(), 7);
+        for i in 0..5 {
+            assert_eq!(p.instructions[i].op, OpCode::EmitPerHorizonSize as u8);
+            assert_eq!(p.instructions[i].arg0, i as u8);
+            assert_eq!(p.instructions[i].arg1, 3);
+        }
+        assert_eq!(p.instructions[5].op, OpCode::AggWeightedSharpe as u8);
+        assert_eq!(p.instructions[5].arg0, 5);
+        assert_eq!(p.instructions[6].op, OpCode::WriteOrder as u8);
+    }
+
+    #[test]
+    fn single_leaf_flattens_to_emit_plus_write() {
+        let leaf = Strategy::Leaf(StrategyConfig {
+            horizon_idx: 4,
+            sizing_policy: SizingPolicyId::IsvKelly,
+            sl_tp_rules: StopRules::default(),
+            max_concurrent_lots: 1,
+        });
+        let p = leaf.flatten();
+        assert_eq!(p.instructions.len(), 2);
+        assert_eq!(p.instructions[0].op, OpCode::EmitPerHorizonSize as u8);
+        assert_eq!(p.instructions[0].arg0, 4);
+        assert_eq!(p.instructions[1].op, OpCode::WriteOrder as u8);
+    }
+
+    #[test]
+    fn portfolio_flattens_with_conflict_op() {
+        let p = Strategy::Portfolio {
+            children: vec![
+                Strategy::Leaf(StrategyConfig {
+                    horizon_idx: 0,
+                    sizing_policy: SizingPolicyId::IsvKelly,
+                    sl_tp_rules: StopRules::default(),
+                    max_concurrent_lots: 1,
+                }),
+                Strategy::Leaf(StrategyConfig {
+                    horizon_idx: 1,
+                    sizing_policy: SizingPolicyId::IsvKelly,
+                    sl_tp_rules: StopRules::default(),
+                    max_concurrent_lots: 1,
+                }),
+            ],
+            capital_allocation: vec![0.6, 0.4],
+            conflict_policy: ConflictPolicy::NetOut,
+        }
+        .flatten();
+        // 2 × EmitPerHorizonSize + AggMean + ApplyConflict + WriteOrder
+        assert_eq!(p.instructions.len(), 5);
+        assert_eq!(p.instructions[2].op, OpCode::AggMean as u8);
+        assert_eq!(p.instructions[3].op, OpCode::ApplyConflict as u8);
+        assert_eq!(p.instructions[3].arg0, ConflictPolicy::NetOut as u8);
+        assert_eq!(p.instructions[4].op, OpCode::WriteOrder as u8);
+    }
+
+    #[test]
+    fn instruction_is_8_bytes_pod() {
+        assert_eq!(std::mem::size_of::<Instruction>(), 8);
+    }
+}
diff --git a/crates/ml-backtesting/src/policy/sizing.rs b/crates/ml-backtesting/src/policy/sizing.rs
new file mode 100644
index 000000000..da92197a6
--- /dev/null
+++ b/crates/ml-backtesting/src/policy/sizing.rs
@@ -0,0 +1,27 @@
+//! Per-leaf sizing policy id. See spec §5.
+
+use serde::{Deserialize, Serialize};
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
+pub enum SizingPolicyId {
+    /// Per-horizon adaptive Kelly with ISV-driven cap (see spec §5).
+    IsvKelly,
+    /// Constant lot count regardless of signal magnitude. Useful for
+    /// fixture testing and as a baseline sanity check vs IsvKelly.
+    FixedLots(u16),
+}
+
+/// Host-side mirror of cuda/lob_state.cuh IsvKellyState.
+/// Used by the harness for warm-start file I/O + tests.
+#[repr(C)]
+#[derive(
+    Clone, Copy, Debug, Default, bytemuck::Pod, bytemuck::Zeroable, Serialize, Deserialize,
+)]
+pub struct IsvKellyStateHost {
+    pub pnl_ema_win: f32,
+    pub pnl_ema_loss: f32,
+    pub win_rate_ema: f32,
+    pub n_trades_seen: u32,
+    pub realised_return_var: f32,
+    pub recent_sharpe: f32,
+}
diff --git a/crates/ml-backtesting/src/policy/stop_rules.rs b/crates/ml-backtesting/src/policy/stop_rules.rs
new file mode 100644
index 000000000..ae0bae464
--- /dev/null
+++ b/crates/ml-backtesting/src/policy/stop_rules.rs
@@ -0,0 +1,14 @@
+//! Per-leaf stop / take-profit / max-hold config. See spec §6 StrategyConfig.
+
+use serde::{Deserialize, Serialize};
+
+#[derive(Clone, Copy, Debug, PartialEq, Default, Serialize, Deserialize)]
+pub struct StopRules {
+    /// Hard stop-loss distance in price-unit ticks (i32). `0` = disabled.
+    pub stop_loss_ticks: i32,
+    /// Take-profit distance in ticks. `0` = disabled.
+    pub take_profit_ticks: i32,
+    /// Max holding time in nanoseconds. `0` = disabled (hold until
+    /// counter-signal closes).
+    pub max_hold_ns: u64,
+}