diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index dcc39fccf..94c2bee9f 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -336,29 +336,41 @@ extern "C" __global__ void resting_orders_step( if (s.queue_position >= trade_abs) { s.queue_position -= trade_abs; } else { - const float over = trade_abs - s.queue_position; - s.queue_position = 0.0f; - const float take = (over < s.size_remaining) ? over : s.size_remaining; - if (take > 0.0f) { - // Filled `take` lots at our limit price (queue position reached us). - const float cost = take * s.price; - apply_fill_to_pos(pos, take, cost, s.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); - s.size_remaining -= take; - emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, - current_ts_ns, - pack_slot_tag(0 /*SlotKind::Limit*/, (unsigned char)i, s.gen_counter), - 1 /*SubmitLimit*/, s.side, (unsigned short)take, - (int)(s.price / tick_size + 0.5f), 0); - if (s.size_remaining <= 0.0f) { - cancel_oco_pair(orders, s.oco_pair); - s.active = 0; - continue; + // Sentinel-priced (market-like) orders use s.price = 1e9 (buy) + // or 0 (sell) so the marketability check below always crosses. + // The queue-decay arm MUST NOT use s.price for cost because it + // would inject the sentinel into avg_px (S1.23 root cause: + // zero_flat=194, huge_flat=216). Absorb the queue depletion + // only; the marketability check fires this same iteration via + // walk_* with book-derived prices. + const bool is_sentinel_price = (s.price < min_px) || (s.price > max_px); + if (is_sentinel_price) { + s.queue_position = 0.0f; + } else { + const float over = trade_abs - s.queue_position; + s.queue_position = 0.0f; + const float take = (over < s.size_remaining) ? over : s.size_remaining; + if (take > 0.0f) { + // Filled `take` lots at our limit price (queue position reached us). + const float cost = take * s.price; + apply_fill_to_pos(pos, take, cost, s.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); + s.size_remaining -= take; + emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + pack_slot_tag(0 /*SlotKind::Limit*/, (unsigned char)i, s.gen_counter), + 1 /*SubmitLimit*/, s.side, (unsigned short)take, + (int)(s.price / tick_size + 0.5f), 0); + if (s.size_remaining <= 0.0f) { + cancel_oco_pair(orders, s.oco_pair); + s.active = 0; + continue; + } } } }