feat(ml-backtesting): hard-SL trigger in stop_check_isv

sl_distance = max(pnl_ema_loss, atr_mid_ema) per
pearl_blend_formulas_must_have_permanent_floor. Fires force-flat
(3, 0) when unrealized_pl_per_lot drops below -sl_distance.
Multi-horizon mask averaging + trail trigger land in Tasks 5-6.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

crates/ml-backtesting/cuda/decision_policy.cu

@@ -48,13 +48,48 @@ __device__ static int stop_check_isv(
     const Book*          __restrict__         books,
     int*                 __restrict__         market_targets
 ) {
-    if (positions[b].position_lots == 0) {
-        return 0;  // flat — entry path handles this
+    const Pos& pos = positions[b];
+    if (pos.position_lots == 0) return 0;
+
+    // Spec §5 controller math — SL only in this task; trail in Task 5.
+    const float atr = atr_mid_ema[b];
+    const unsigned int mask = open_horizon_masks[b];
+
+    const IsvKellyState* isv = reinterpret_cast<const IsvKellyState*>(
+        isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState)
+    );
+    float ema_loss = 0.0f;
+    int   n_open   = 0;
+    #pragma unroll
+    for (int h = 0; h < N_HORIZONS; ++h) {
+        if (mask & (1u << h)) {
+            ema_loss += isv[h].pnl_ema_loss;
+            n_open   += 1;
+        }
+    }
+    if (n_open > 0) ema_loss /= (float)n_open;
+
+    // pearl_blend_formulas_must_have_permanent_floor — max, not blend.
+    const float sl_distance = fmaxf(ema_loss, atr);
+
+    const bool   is_long  = pos.position_lots > 0;
+    const Book&  bk       = books[b];
+    const float  close_px = is_long ? bk.bid_px[0] : bk.ask_px[0];
+    const float  entry_px = pos.vwap_entry;
+    const float  unrealized_pl_per_lot =
+        is_long ? (close_px - entry_px) : (entry_px - close_px);
+
+    const bool sl_fired = unrealized_pl_per_lot <= -sl_distance;
+
+    // trail_hwm is unused in this task — Task 5 will wire it.
+    (void)trail_hwm;
+
+    if (sl_fired) {
+        // §7 force-flat encoding.
+        market_targets[b * 2 + 0] = 3;
+        market_targets[b * 2 + 1] = 0;
+        return 1;
     }
-    // Stop trigger logic added in Tasks 4-6.
-    // Mark args as referenced so -Wunused-parameter stays quiet across cubins:
-    (void)isv_kelly_base; (void)open_horizon_masks; (void)atr_mid_ema;
-    (void)trail_hwm; (void)books; (void)market_targets;
     return 0;
 }
 

crates/ml-backtesting/tests/stop_controller.rs

@@ -96,3 +96,58 @@ fn stop_check_skipped_when_flat() -> Result<()> {
     assert!(size >= 1, "entry should fire under p=0.8 + flat (size={size})");
     Ok(())
 }
+
+use ml_backtesting::policy::IsvKellyStateHost;
+
+#[test]
+#[ignore = "requires CUDA"]
+fn sl_fires_when_unrealized_breaks_distance() -> Result<()> {
+    let dev = match MlDevice::cuda(0) {
+        Ok(d) => d,
+        Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); }
+    };
+    let mut sim = LobSimCuda::new(1, &dev)?;
+
+    // Seed isv: 5 horizons with pnl_ema_loss=2.0 → sl_distance=2.0.
+    // pnl_ema_win=0.0 forces cold-start branch in decision kernel (kelly_frac →
+    // floor=0.20, cap_lots → max_lots=5). With alpha=0.95 this gives
+    // lots=round(0.9*0.20*5)=round(0.9)=1 — position opens.
+    let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost {
+        pnl_ema_win:  0.0,
+        pnl_ema_loss: 2.0,
+        win_rate_ema: 0.5,
+        n_trades_seen: 0,
+        realised_return_var: 0.0,
+        recent_sharpe: 0.5,
+    });
+    sim.write_isv_kelly(0, &seeded)?;
+
+    // Two snapshots so atr_mid_ema becomes nonzero (small).
+    let (bp, bs, ap, az) = level_book(5500.0, 0.25);
+    sim.apply_snapshot(&bp, &bs, &ap, &az)?;
+    let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25);
+    sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?;
+
+    // Open a long position via strong-alpha + step_decision + manual fill.
+    sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?;
+    sim.step_decision_with_latency(0, &cfg_default(1))?;
+    let mut ts: u64 = 1_000_000;
+    for _ in 0..3 {
+        sim.step_resting_orders(ts, 0.0)?;
+        ts += 1_000_000;
+    }
+    let pos_open = sim.read_pos(0)?;
+    assert!(pos_open.position_lots > 0,
+            "test setup: long position should open, got {}", pos_open.position_lots);
+
+    // Drive mid down by 3.0 (>= sl_distance=2.0). Stop must fire force-flat.
+    let (bp3, bs3, ap3, az3) = level_book(5497.0, 0.25);
+    sim.apply_snapshot(&bp3, &bs3, &ap3, &az3)?;
+    sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?;
+    sim.step_decision_with_latency(ts, &cfg_default(1))?;
+
+    let (side, size) = sim.read_market_target(0)?;
+    assert_eq!(side, 3, "SL must fire with force-flat (3, 0); got side={side}");
+    assert_eq!(size, 0, "force-flat encoding has abs_sz=0; got size={size}");
+    Ok(())
+}