fix(ml-backtesting): realize PnL on event-rate max_hold force-close

S2.2 moved max_hold to event-rate at step_resting_orders (cap is now
enforced — mean hold 432s → 58.43s, max 173893s → 96s). But it mirrored
the session-gap pattern `pos.position_lots = 0;` which intentionally
skips PnL ("cannot fill across halt"). Max_hold differs: the market is
live so the close SHOULD realize PnL via the current top-of-book.

Smoke t9msj (b92bd72c7) showed the consequence: 84% of trades record
$0 realised_pnl, win_rate dropped to 3.3%, total_pnl looks artificially
better (-$225k) only because losses aren't recorded.

Fix: route the force-close through apply_fill_to_pos by synthesizing a
closing fill at book.bid_px[0] (long) or book.ask_px[0] (short). The
existing counter-direction branch (realised = (avg_px − vwap) × dir ×
unwind) correctly realizes PnL on the unwound position. Top-of-book is
already validated by book_update_apply_snapshot's skip, so close_px is
guaranteed in range.

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

crates/ml-backtesting/cuda/resting_orders.cu

@@ -321,7 +321,30 @@ extern "C" __global__ void resting_orders_step(
                 open_trade_state + (size_t)b * 24);  // OPEN_TRADE_STATE_BYTES = 24
         if (entry_ts > 0ull && current_ts_ns >= entry_ts
             && (current_ts_ns - entry_ts) >= max_hold) {
-            pos.position_lots = 0;   // force-flat; pnl_track close fires next step
+            // S2.3: realize PnL on the force-close. apply_fill_to_pos's
+            // counter-direction branch handles the unwind:
+            //   realised = (avg_px − vwap) × dir × unwind
+            // Synthesize the closing fill at top-of-book on the closing side:
+            // long → sell at bid[0], short → buy at ask[0].
+            // book[b]'s top-of-book is already validated by book_update's skip
+            // (book_update_apply_snapshot rejects bad-top snapshots), so
+            // close_px is guaranteed in [min_px, max_px] and > 0.
+            const int  pos_lots   = pos.position_lots;
+            const bool is_long    = pos_lots > 0;
+            const float close_px  = is_long ? book.bid_px[0] : book.ask_px[0];
+            const float lots_abs  = is_long ? (float)pos_lots : -(float)pos_lots;
+            const float total_cost = lots_abs * close_px;
+            const unsigned char close_side = is_long ? 1 : 0;  // sell to close long, buy to close short
+            apply_fill_to_pos(
+                pos, lots_abs, total_cost, close_side,
+                b, cost_per_lot_per_side_per_b, total_fees_per_b,
+                nan_avg_px, nan_realised, nan_realized_pnl,
+                vwap_zero_open_flat, vwap_huge_open_flat,
+                vwap_zero_scale_in,  vwap_huge_scale_in,
+                vwap_zero_flip,      vwap_huge_flip,
+                last_bad_vwap,       last_bad_path
+            );
+            // After apply_fill_to_pos with unwind == prev_abs, position_lots = 0.
         }
     }