feat(ml-backtesting): trail-TP trigger with HWM ratchet in stop_check_isv

trail_distance = max(pnl_ema_win, atr_mid_ema). HWM ratchets up via
fmaxf each event while position open; trail fires when HWM clears
the arming threshold AND unrealized drops by trail_distance from
HWM. Same force-flat (3, 0) write as SL.

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

crates/ml-backtesting/cuda/decision_policy.cu

@@ -59,18 +59,24 @@ __device__ static int stop_check_isv(
         isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState)
     );
     float ema_loss = 0.0f;
+    float ema_win  = 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;
+            ema_win  += isv[h].pnl_ema_win;
             n_open   += 1;
         }
     }
-    if (n_open > 0) ema_loss /= (float)n_open;
+    if (n_open > 0) {
+        ema_loss /= (float)n_open;
+        ema_win  /= (float)n_open;
+    }
 
     // pearl_blend_formulas_must_have_permanent_floor — max, not blend.
-    const float sl_distance = fmaxf(ema_loss, atr);
+    const float sl_distance    = fmaxf(ema_loss, atr);
+    const float trail_distance = fmaxf(ema_win,  atr);
 
     const bool   is_long  = pos.position_lots > 0;
     const Book&  bk       = books[b];
@@ -81,10 +87,13 @@ __device__ static int stop_check_isv(
 
     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;
+    const float new_hwm = fmaxf(trail_hwm[b], unrealized_pl_per_lot);
+    trail_hwm[b] = new_hwm;
+    const bool trail_fired =
+           (new_hwm > trail_distance)
+        && (unrealized_pl_per_lot <= new_hwm - trail_distance);
 
-    if (sl_fired) {
+    if (sl_fired || trail_fired) {
         // §7 force-flat encoding.
         market_targets[b * 2 + 0] = 3;
         market_targets[b * 2 + 1] = 0;

crates/ml-backtesting/tests/stop_controller.rs

@@ -151,3 +151,71 @@ fn sl_fires_when_unrealized_breaks_distance() -> Result<()> {
     assert_eq!(size, 0, "force-flat encoding has abs_sz=0; got size={size}");
     Ok(())
 }
+
+#[test]
+#[ignore = "requires CUDA"]
+fn trail_arms_then_fires() -> 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)?;
+
+    // Cold-start ISV (forces kelly_frac_floor + max_lots cap) so position opens.
+    // pnl_ema_loss large enough that SL won't fire during the +5pt walk.
+    let cold_start: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost {
+        pnl_ema_win:  0.0,
+        pnl_ema_loss: 10.0,    // big SL so trail fires first
+        win_rate_ema: 0.0,
+        n_trades_seen: 0,
+        realised_return_var: 0.0,
+        recent_sharpe: 0.0,
+    });
+    sim.write_isv_kelly(0, &cold_start)?;
+
+    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 long.
+    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; }
+    assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opened");
+
+    // After opening, switch ISV to one with pnl_ema_win=1.5 (trail_distance target).
+    // The trail check reads this fresh state on subsequent decision steps.
+    let trail_state: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost {
+        pnl_ema_win:  1.5,
+        pnl_ema_loss: 10.0,
+        win_rate_ema: 0.5,
+        n_trades_seen: 20,
+        realised_return_var: 0.5,
+        recent_sharpe: 0.5,
+    });
+    sim.write_isv_kelly(0, &trail_state)?;
+
+    // Walk mid up to ratchet trail HWM past trail_distance=1.5.
+    for &m in &[5501.0_f32, 5502.0, 5503.0, 5504.0, 5505.0] {
+        let (b, bs, a, asz) = level_book(m, 0.25);
+        sim.apply_snapshot(&b, &bs, &a, &asz)?;
+        sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?;
+        sim.step_decision_with_latency(ts, &cfg_default(1))?;
+        ts += 1_000_000;
+    }
+    let hwm_armed = sim.read_trail_hwm(0)?;
+    assert!(hwm_armed > 1.5, "HWM should ratchet past trail_distance=1.5, got {hwm_armed}");
+
+    // Drop mid by trail_distance+ε from peak to fire trail.
+    let (b, bs, a, asz) = level_book(5502.0, 0.25);  // drop from peak 5505 to 5502 = 3.0
+    sim.apply_snapshot(&b, &bs, &a, &asz)?;
+    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, "trail must fire force-flat; got side={side}");
+    assert_eq!(size, 0);
+    Ok(())
+}