025554afcd
Cluster smoke 88dbp (post-Fix-S1.19 parameterized price range) still produces 97 zero + 85 i32::MAX sentinel entry_px values despite source data being fully filtered. Sentinels originate in kernel arithmetic paths post-sanitization, not from book input. Adds 6 per-backtest u32 counters incremented at NaN-producing sites: - nan_avg_px: avg_px = total_cost/filled_lots -> NaN/Inf - nan_realised: (avg_px - vwap_entry) * dir * unwind -> NaN/Inf - nan_realized_pnl: pos.realized_pnl becomes non-finite after += or -= - zero_vwap_at_open: pnl_track open branch saw vwap_entry == 0 - saturated_vwap_at_open: pnl_track open saw |vwap_entry| > 21M or NaN/Inf - defensive_exit_clamp: pnl_track close defensive clamp fired Each counter printed at end of smoke via nan_counters: log line. apply_fill_to_pos (resting_orders.cu) gains 3 counter args; NaN-producing paths return early WITHOUT propagating into pos state. pnl_track_step gains 3 counter args. All call sites threaded through sim/mod.rs launches. NanCounters struct + read_nan_counters() accessor added to LobSimCuda. Unit test nan_counters_initialize_to_zero confirms zero-init and accessor compile. 18/18 stop_controller, 5/5 decision_floor_coldstart pass. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
558 lines
25 KiB
Plaintext
558 lines
25 KiB
Plaintext
// resting_orders.cu — limit + stop + OCO machinery.
|
||
//
|
||
// Runs on every event (after book_update, before pnl_track). Per-block,
|
||
// single-writer (thread 0). Processes:
|
||
//
|
||
// 1. In-flight → resting promotion when arrival_ts_ns reached.
|
||
// Marketable on arrival? Cross immediately at the just-arrived book.
|
||
// 2. Resting limits: queue-decay against trade_signed_vol at level,
|
||
// check marketability against the new book (price moved through us).
|
||
// 3. Stop arming/triggering: best ask crossing trigger up → buy stop,
|
||
// best bid crossing trigger down → sell stop. On trigger:
|
||
// - StopMarket: walk book immediately, fill, emit OrderEvent
|
||
// - StopLimit: convert to a resting limit (allocate a LimitSlot)
|
||
// 4. OCO mutual exclusion: when one leg fills or stops triggers, free
|
||
// the paired slot via oco_pair index (0..31 = limit, 32..47 = stop).
|
||
//
|
||
// Per feedback_no_atomicadd.md: no atomics — single writer per block.
|
||
|
||
#include "lob_state.cuh"
|
||
|
||
#define KIND_LIMIT 1
|
||
#define KIND_IOC 2
|
||
#define KIND_FOK 3
|
||
#define KIND_STOP_MARKET 4
|
||
#define KIND_STOP_LIMIT 5
|
||
|
||
// Book walking helpers — same shape as order_match.cu but inlined here
|
||
// to keep this kernel's symbol set self-contained.
|
||
__device__ static float walk_ask_for_buy(const Book& book, float size, float& filled) {
|
||
float cost = 0.0f; filled = 0.0f;
|
||
#pragma unroll
|
||
for (int k = 0; k < 10 && size > 0.0f; ++k) {
|
||
const float lvl_sz = book.ask_sz[k];
|
||
const float lvl_px = book.ask_px[k];
|
||
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue;
|
||
const float take = (size < lvl_sz) ? size : lvl_sz;
|
||
cost += take * lvl_px; filled += take; size -= take;
|
||
}
|
||
return cost;
|
||
}
|
||
__device__ static float walk_bid_for_sell(const Book& book, float size, float& filled) {
|
||
float cost = 0.0f; filled = 0.0f;
|
||
#pragma unroll
|
||
for (int k = 0; k < 10 && size > 0.0f; ++k) {
|
||
const float lvl_sz = book.bid_sz[k];
|
||
const float lvl_px = book.bid_px[k];
|
||
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue;
|
||
const float take = (size < lvl_sz) ? size : lvl_sz;
|
||
cost += take * lvl_px; filled += take; size -= take;
|
||
}
|
||
return cost;
|
||
}
|
||
|
||
// Mirror of decision_policy.cu's emit_event — duplicated to avoid
|
||
// a cross-kernel header dependency. Layout matches OrderEvent (24 B)
|
||
// in crates/ml-backtesting/src/order.rs.
|
||
__device__ static void emit_audit(
|
||
unsigned char* audit_base, unsigned int* head_d, int b, unsigned int ring_cap,
|
||
unsigned long long ts_ns, unsigned int slot_tag, unsigned char kind,
|
||
unsigned char side, unsigned short size, int px_ticks, int trigger_px_ticks
|
||
) {
|
||
const unsigned int idx = head_d[b] % ring_cap;
|
||
head_d[b] += 1;
|
||
unsigned char* rec = audit_base + ((size_t)b * ring_cap + idx) * 24;
|
||
*reinterpret_cast<unsigned long long*>(rec + 0) = ts_ns;
|
||
*reinterpret_cast<unsigned int*>(rec + 8) = slot_tag;
|
||
rec[12] = kind;
|
||
rec[13] = side;
|
||
*reinterpret_cast<unsigned short*>(rec + 14) = size;
|
||
*reinterpret_cast<int*>(rec + 16) = px_ticks;
|
||
*reinterpret_cast<int*>(rec + 20) = trigger_px_ticks;
|
||
}
|
||
|
||
// SlotTag pack: kind | (idx << 8) | (gen << 16). See src/order.rs.
|
||
__device__ static unsigned int pack_slot_tag(unsigned char kind, unsigned char idx, unsigned short gen) {
|
||
return (unsigned int)kind | ((unsigned int)idx << 8) | ((unsigned int)gen << 16);
|
||
}
|
||
|
||
// Free the slot referenced by an oco_pair byte: 0..31 = limits, 32..47 = stops.
|
||
__device__ static void cancel_oco_pair(Orders& orders, unsigned char oco_pair) {
|
||
if (oco_pair == 0xFF) return;
|
||
if (oco_pair < MAX_LIMITS) {
|
||
orders.limits[oco_pair].active = 0;
|
||
} else if (oco_pair - MAX_LIMITS < MAX_STOPS) {
|
||
orders.stops[oco_pair - MAX_LIMITS].active = 0;
|
||
}
|
||
}
|
||
|
||
// Apply a fill to per-backtest position state. Same VWAP-scale-in /
|
||
// counter-direction realised-P&L rule as order_match.cu's market-order
|
||
// path; isolated here as a __device__ helper.
|
||
//
|
||
// P4: per-fill cost integration. After the close-leg realized_pnl math
|
||
// runs (so the unwind P&L is computed from gross prices), deduct the
|
||
// fill cost from realized_pnl AND accumulate into total_fees_per_b[b].
|
||
// Net-of-cost semantics: isv_kelly_update_on_close reads realized_pnl
|
||
// delta which is now net of cost — Kelly state learns realistic returns.
|
||
//
|
||
// Plain `+=` for total_fees_per_b is safe under the single-writer-per-block
|
||
// convention enforced by `if (threadIdx.x != 0) return;` at the top of
|
||
// every block-per-backtest kernel.
|
||
//
|
||
// S1.20: NaN instrumentation counters.
|
||
// nan_avg_px — incremented when avg_px = total_cost/filled_lots is NaN/Inf.
|
||
// nan_realised — incremented when realised P&L in counter-direction branch is NaN/Inf.
|
||
// nan_realized_pnl — incremented when pos.realized_pnl becomes non-finite after += or -=.
|
||
// On NaN-producing events the function returns early WITHOUT poisoning pos state.
|
||
__device__ static void apply_fill_to_pos(
|
||
Pos& pos,
|
||
float filled_lots,
|
||
float total_cost,
|
||
unsigned char side,
|
||
int b,
|
||
const float* __restrict__ cost_per_lot_per_side_per_b,
|
||
float* __restrict__ total_fees_per_b,
|
||
unsigned int* __restrict__ nan_avg_px,
|
||
unsigned int* __restrict__ nan_realised,
|
||
unsigned int* __restrict__ nan_realized_pnl
|
||
) {
|
||
if (filled_lots <= 0.0f) return;
|
||
const float avg_px = total_cost / filled_lots;
|
||
if (!isfinite(avg_px)) {
|
||
nan_avg_px[b] += 1u;
|
||
return; // do NOT propagate NaN into pos state
|
||
}
|
||
const float sgn = (side == 0) ? 1.0f : -1.0f;
|
||
const float prev_pos_f = (float)pos.position_lots;
|
||
const float new_pos_f = prev_pos_f + sgn * filled_lots;
|
||
|
||
if (pos.position_lots == 0) {
|
||
pos.vwap_entry = avg_px;
|
||
} else if ((prev_pos_f > 0.0f) == (sgn > 0.0f)) {
|
||
const float prev_notional = prev_pos_f * pos.vwap_entry;
|
||
const float new_notional = prev_notional + sgn * total_cost;
|
||
const float new_abs = (new_pos_f > 0.0f) ? new_pos_f : -new_pos_f;
|
||
pos.vwap_entry = (new_abs > 0.0f) ? (new_notional / (sgn * new_abs)) : 0.0f;
|
||
} else {
|
||
const float prev_abs = (prev_pos_f > 0.0f) ? prev_pos_f : -prev_pos_f;
|
||
const float unwind = (filled_lots < prev_abs) ? filled_lots : prev_abs;
|
||
const float dir_was_long = (prev_pos_f > 0.0f) ? 1.0f : -1.0f;
|
||
const float realised = (avg_px - pos.vwap_entry) * dir_was_long * unwind;
|
||
if (!isfinite(realised)) {
|
||
nan_realised[b] += 1u;
|
||
return; // do not poison realized_pnl
|
||
}
|
||
pos.realized_pnl += realised;
|
||
if (!isfinite(pos.realized_pnl)) {
|
||
nan_realized_pnl[b] += 1u;
|
||
}
|
||
if (filled_lots > prev_abs) { pos.vwap_entry = avg_px; }
|
||
}
|
||
pos.position_lots = (int)new_pos_f;
|
||
|
||
// P4: per-fill cost deduction (single-writer-per-block; plain += safe).
|
||
const float fill_cost = filled_lots * cost_per_lot_per_side_per_b[b];
|
||
pos.realized_pnl -= fill_cost;
|
||
if (!isfinite(pos.realized_pnl)) {
|
||
nan_realized_pnl[b] += 1u;
|
||
}
|
||
total_fees_per_b[b] += fill_cost;
|
||
}
|
||
|
||
// Free-slot search. Returns slot index 0..31 or -1 if all in use.
|
||
__device__ static int find_free_limit_slot(Orders& orders) {
|
||
#pragma unroll
|
||
for (int i = 0; i < MAX_LIMITS; ++i) {
|
||
if (orders.limits[i].active == 0) return i;
|
||
}
|
||
return -1;
|
||
}
|
||
|
||
// Main resting-orders processing kernel.
|
||
//
|
||
// `tick_size` is the price-to-tick conversion for audit-event recording.
|
||
extern "C" __global__ void resting_orders_step(
|
||
const Book* __restrict__ books, // [n_backtests]
|
||
unsigned char* orders_base, // [n_backtests * ORDERS_BYTES]
|
||
unsigned char* pos_base,
|
||
unsigned char* audit_base,
|
||
unsigned int* audit_head,
|
||
unsigned long long current_ts_ns,
|
||
float trade_signed_vol,
|
||
int orders_bytes,
|
||
int pos_bytes,
|
||
int ring_cap_events,
|
||
float tick_size,
|
||
// P4: per-fill cost integration (per-backtest arrays).
|
||
const float* __restrict__ cost_per_lot_per_side_per_b,
|
||
float* __restrict__ total_fees_per_b,
|
||
// Fix 3: session-gap force-close. Tracks last event ts per backtest.
|
||
unsigned long long* __restrict__ last_event_ts, // [n_backtests]
|
||
// S1.20: NaN instrumentation counters (per-backtest).
|
||
unsigned int* __restrict__ nan_avg_px,
|
||
unsigned int* __restrict__ nan_realised,
|
||
unsigned int* __restrict__ nan_realized_pnl,
|
||
int n_backtests
|
||
) {
|
||
int b = blockIdx.x;
|
||
if (b >= n_backtests || threadIdx.x != 0) return;
|
||
|
||
const Book& book = books[b];
|
||
Orders& orders = *reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
||
Pos& pos = *reinterpret_cast<Pos*>(pos_base + (size_t)b * pos_bytes);
|
||
|
||
// Session-gap force-close: when event timestamps jump > 1 hour (weekend
|
||
// halt, session boundary, gap-fill), max_hold can't fire because no
|
||
// intervening events advance current_ts. Force-close all open positions
|
||
// before processing the new event's fills.
|
||
const unsigned long long SESSION_GAP_NS = 3600000000000ull; // 1 hour
|
||
const unsigned long long last_ts = last_event_ts[b];
|
||
if (last_ts > 0ull && current_ts_ns > last_ts
|
||
&& (current_ts_ns - last_ts) >= SESSION_GAP_NS
|
||
&& pos.position_lots != 0)
|
||
{
|
||
// Force-flat: zero position_lots directly. pnl_track_step (called by
|
||
// step_resting_orders after this kernel) will detect the prev!=0 && now==0
|
||
// transition via open_trade_state scratch and emit a close TradeRecord.
|
||
// No synthetic P&L is added — records show realised_pnl from whatever
|
||
// was accumulated before the gap (honest: cannot fill across a halt).
|
||
pos.position_lots = 0;
|
||
}
|
||
last_event_ts[b] = current_ts_ns;
|
||
|
||
// (1) Promote in-flight limits to resting; (2) queue-decay; (3) marketability check.
|
||
const float trade_abs = (trade_signed_vol < 0.0f) ? -trade_signed_vol : trade_signed_vol;
|
||
for (int i = 0; i < MAX_LIMITS; ++i) {
|
||
LimitSlot& s = orders.limits[i];
|
||
if (s.active == 2 && current_ts_ns >= s.arrival_ts_ns) {
|
||
s.active = 1;
|
||
// Initialise queue position pessimistically: full level depth
|
||
// ahead of us at the just-arrived book (per spec §9 queue-decay
|
||
// risk row).
|
||
int level_idx = -1;
|
||
if (s.side == 0) {
|
||
#pragma unroll
|
||
for (int k = 0; k < 10; ++k) if (book.bid_px[k] == s.price) { level_idx = k; break; }
|
||
s.queue_position = (level_idx >= 0) ? book.bid_sz[level_idx] : 0.0f;
|
||
} else {
|
||
#pragma unroll
|
||
for (int k = 0; k < 10; ++k) if (book.ask_px[k] == s.price) { level_idx = k; break; }
|
||
s.queue_position = (level_idx >= 0) ? book.ask_sz[level_idx] : 0.0f;
|
||
}
|
||
}
|
||
if (s.active != 1) continue;
|
||
|
||
// Queue decay: trade flow on our side at our price level eats ahead-of-us first,
|
||
// then us. Only count trades on the SAME side as our resting order.
|
||
if (trade_abs > 0.0f) {
|
||
const bool same_side_trade = (s.side == 0 && trade_signed_vol < 0.0f) // resting bid + seller hits us
|
||
|| (s.side == 1 && trade_signed_vol > 0.0f); // resting ask + buyer lifts
|
||
if (same_side_trade) {
|
||
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);
|
||
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;
|
||
}
|
||
}
|
||
}
|
||
}
|
||
}
|
||
|
||
// Marketability check: book moved through our price → cross at the
|
||
// arrival book (slippage-aware fill). Triggers on:
|
||
// buy limit @ p ≥ ask_px[0] — we'd be taking liquidity at our price or better
|
||
// sell limit @ p ≤ bid_px[0] — same on opposite side
|
||
const bool marketable = (s.side == 0 && s.price >= book.ask_px[0])
|
||
|| (s.side == 1 && s.price <= book.bid_px[0]);
|
||
if (marketable) {
|
||
float filled = 0.0f;
|
||
const float cost = (s.side == 0)
|
||
? walk_ask_for_buy (book, s.size_remaining, filled)
|
||
: walk_bid_for_sell(book, s.size_remaining, filled);
|
||
if (filled > 0.0f) {
|
||
apply_fill_to_pos(pos, filled, cost, s.side,
|
||
b, cost_per_lot_per_side_per_b, total_fees_per_b,
|
||
nan_avg_px, nan_realised, nan_realized_pnl);
|
||
s.size_remaining -= filled;
|
||
emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events,
|
||
current_ts_ns,
|
||
pack_slot_tag(0, (unsigned char)i, s.gen_counter),
|
||
s.kind /*Limit|Ioc|Fok*/, s.side, (unsigned short)filled,
|
||
(int)(s.price / tick_size + 0.5f), 0);
|
||
}
|
||
// IOC: any unfilled remainder cancelled immediately.
|
||
// FOK: if not fully filled, cancel everything (no partial leaves resting).
|
||
// Limit: leave residual resting at our price.
|
||
if (s.kind == KIND_IOC || (s.kind == KIND_FOK && s.size_remaining > 0.0f)) {
|
||
if (s.kind == KIND_FOK && filled > 0.0f) {
|
||
// FOK partial fill is forbidden — should not happen; defensive guard
|
||
// documents intent but the actual prevention belongs in submission code.
|
||
}
|
||
cancel_oco_pair(orders, s.oco_pair);
|
||
s.active = 0;
|
||
continue;
|
||
}
|
||
if (s.size_remaining <= 0.0f) {
|
||
cancel_oco_pair(orders, s.oco_pair);
|
||
s.active = 0;
|
||
continue;
|
||
}
|
||
}
|
||
}
|
||
|
||
// (4) Stop trigger detection.
|
||
for (int j = 0; j < MAX_STOPS; ++j) {
|
||
StopSlot& st = orders.stops[j];
|
||
if (st.active == 2 && current_ts_ns >= st.arrival_ts_ns) st.active = 1;
|
||
if (st.active != 1) continue;
|
||
|
||
const bool triggered = (st.side == 0 && book.ask_px[0] >= st.trigger_price)
|
||
|| (st.side == 1 && book.bid_px[0] <= st.trigger_price);
|
||
if (!triggered) continue;
|
||
st.active = 0;
|
||
|
||
if (st.kind == KIND_STOP_MARKET) {
|
||
float filled = 0.0f;
|
||
const float cost = (st.side == 0)
|
||
? walk_ask_for_buy (book, st.size, filled)
|
||
: walk_bid_for_sell(book, st.size, filled);
|
||
if (filled > 0.0f) {
|
||
apply_fill_to_pos(pos, filled, cost, st.side,
|
||
b, cost_per_lot_per_side_per_b, total_fees_per_b,
|
||
nan_avg_px, nan_realised, nan_realized_pnl);
|
||
emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events,
|
||
current_ts_ns,
|
||
pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter),
|
||
4 /*SubmitStopMarket*/, st.side, (unsigned short)filled,
|
||
0, (int)(st.trigger_price / tick_size + 0.5f));
|
||
}
|
||
cancel_oco_pair(orders, st.oco_pair);
|
||
} else if (st.kind == KIND_STOP_LIMIT) {
|
||
// Convert to a resting limit at st.limit_price (may not fill immediately
|
||
// if the limit is not marketable at current book).
|
||
const int free_idx = find_free_limit_slot(orders);
|
||
if (free_idx < 0) {
|
||
pos.submission_overflow += 1u;
|
||
cancel_oco_pair(orders, st.oco_pair);
|
||
continue;
|
||
}
|
||
LimitSlot& ns = orders.limits[free_idx];
|
||
ns.active = 1;
|
||
ns.side = st.side;
|
||
ns.kind = KIND_LIMIT;
|
||
ns.oco_pair = st.oco_pair;
|
||
ns.price = st.limit_price;
|
||
ns.size_remaining = st.size;
|
||
ns.queue_position = 0.0f; // marketable check next iter will fire if applicable
|
||
ns.arrival_ts_ns = current_ts_ns;
|
||
ns.gen_counter = (unsigned short)(ns.gen_counter + 1);
|
||
emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events,
|
||
current_ts_ns,
|
||
pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter),
|
||
5 /*SubmitStopLimit*/, st.side, (unsigned short)st.size,
|
||
(int)(st.limit_price / tick_size + 0.5f),
|
||
(int)(st.trigger_price / tick_size + 0.5f));
|
||
}
|
||
}
|
||
}
|
||
|
||
// Host-callable seed helpers — write a single slot. Single-block,
|
||
// single-thread. Used for fixture testing + as the kernel side of
|
||
// LobSimCuda::submit_limit. SlotKind::Limit only (Stop seed below).
|
||
|
||
extern "C" __global__ void seed_limit_slot(
|
||
unsigned char* orders_base,
|
||
int orders_bytes,
|
||
int b,
|
||
int slot_idx,
|
||
unsigned char side,
|
||
unsigned char kind,
|
||
unsigned char active_state, // 1=resting immediately, 2=in-flight
|
||
unsigned char oco_pair,
|
||
float price,
|
||
float size,
|
||
float queue_position_init,
|
||
unsigned long long arrival_ts_ns,
|
||
unsigned int* overflow_flag // [1] — set to 1 if slot_idx already in use
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
Orders* orders = reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
||
if (slot_idx < 0 || slot_idx >= MAX_LIMITS) { *overflow_flag = 1u; return; }
|
||
LimitSlot& s = orders->limits[slot_idx];
|
||
if (s.active != 0) { *overflow_flag = 1u; return; }
|
||
s.active = active_state;
|
||
s.side = side;
|
||
s.kind = kind;
|
||
s.oco_pair = oco_pair;
|
||
s.price = price;
|
||
s.size_remaining = size;
|
||
s.queue_position = queue_position_init;
|
||
s.arrival_ts_ns = arrival_ts_ns;
|
||
s.gen_counter = (unsigned short)(s.gen_counter + 1);
|
||
}
|
||
|
||
extern "C" __global__ void seed_stop_slot(
|
||
unsigned char* orders_base,
|
||
int orders_bytes,
|
||
int b,
|
||
int slot_idx,
|
||
unsigned char side,
|
||
unsigned char kind,
|
||
unsigned char active_state,
|
||
unsigned char oco_pair,
|
||
float trigger_price,
|
||
float limit_price,
|
||
float size,
|
||
unsigned long long arrival_ts_ns,
|
||
unsigned int* overflow_flag
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
Orders* orders = reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
||
if (slot_idx < 0 || slot_idx >= MAX_STOPS) { *overflow_flag = 1u; return; }
|
||
StopSlot& st = orders->stops[slot_idx];
|
||
if (st.active != 0) { *overflow_flag = 1u; return; }
|
||
st.active = active_state;
|
||
st.side = side;
|
||
st.kind = kind;
|
||
st.oco_pair = oco_pair;
|
||
st.trigger_price = trigger_price;
|
||
st.limit_price = limit_price;
|
||
st.size = size;
|
||
st.arrival_ts_ns = arrival_ts_ns;
|
||
st.gen_counter = (unsigned short)(st.gen_counter + 1);
|
||
}
|
||
|
||
// Try to allocate a limit slot from device-side submission code (e.g. C7
|
||
// decision kernel later) — returns slot index or -1 if all in use.
|
||
// Bumps pos.submission_overflow on failure. Single-block, single-thread.
|
||
extern "C" __global__ void submit_limit_alloc(
|
||
unsigned char* orders_base,
|
||
unsigned char* pos_base,
|
||
int orders_bytes,
|
||
int pos_bytes,
|
||
int b,
|
||
unsigned char side,
|
||
unsigned char kind,
|
||
unsigned char oco_pair,
|
||
float price,
|
||
float size,
|
||
unsigned long long arrival_ts_ns,
|
||
int* allocated_slot_out // [1] — written: slot idx or -1
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
Orders* orders = reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
||
Pos* pos = reinterpret_cast<Pos*>(pos_base + (size_t)b * pos_bytes);
|
||
const int idx = find_free_limit_slot(*orders);
|
||
if (idx < 0) {
|
||
pos->submission_overflow += 1u;
|
||
*allocated_slot_out = -1;
|
||
return;
|
||
}
|
||
LimitSlot& s = orders->limits[idx];
|
||
s.active = 2; // in-flight
|
||
s.side = side;
|
||
s.kind = kind;
|
||
s.oco_pair = oco_pair;
|
||
s.price = price;
|
||
s.size_remaining = size;
|
||
s.queue_position = 0.0f;
|
||
s.arrival_ts_ns = arrival_ts_ns;
|
||
s.gen_counter = (unsigned short)(s.gen_counter + 1);
|
||
*allocated_slot_out = idx;
|
||
}
|
||
|
||
// P2: GPU-side seeder for in-flight LimitSlots driven by market_targets[b].
|
||
// Replaces the host-roundtrip loop in sim.rs::dispatch_latent_market_orders.
|
||
// Per-backtest single-writer (threadIdx.x == 0).
|
||
//
|
||
// §7 truth table:
|
||
// side=0, abs_sz=N → target signed position = +N (long)
|
||
// side=1, abs_sz=N → target signed position = -N (short)
|
||
// side=2 → no-op: preserve current position
|
||
// side=3 → force flat: target signed position = 0
|
||
//
|
||
// effective_position = pos.position_lots + Σ unfilled signed slot sizes
|
||
// where unfilled = active∈{1,2} (resting or in-flight).
|
||
// delta = target_signed − effective_position.
|
||
// If delta == 0: nothing to do (idempotent re-submission under latency).
|
||
// If delta != 0: seed an in-flight order for |delta| lots on the
|
||
// appropriate side. If all 32 slots are occupied, increments
|
||
// pos.submission_overflow.
|
||
extern "C" __global__ void seed_inflight_limits_batched(
|
||
unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes]
|
||
int orders_bytes,
|
||
Pos* __restrict__ positions, // [n_backtests]
|
||
const int* __restrict__ market_targets, // [n_backtests * 2]
|
||
const unsigned int* __restrict__ latency_ns_per_b, // [n_backtests]
|
||
unsigned long long current_ts_ns,
|
||
int n_backtests
|
||
) {
|
||
int b = blockIdx.x;
|
||
if (b >= n_backtests || threadIdx.x != 0) return;
|
||
|
||
const int side = market_targets[b * 2 + 0];
|
||
const int abs_sz = market_targets[b * 2 + 1];
|
||
|
||
if (side == 2) return; // §7: no-op = preserve position
|
||
|
||
int target_signed;
|
||
if (side == 0) target_signed = +abs_sz;
|
||
else if (side == 1) target_signed = -abs_sz;
|
||
else target_signed = 0; // side == 3: force flat
|
||
|
||
Orders* orders = reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
||
|
||
// Effective position = filled position + all unfilled signed lots.
|
||
int effective = positions[b].position_lots;
|
||
#pragma unroll
|
||
for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) {
|
||
const LimitSlot& s = orders->limits[s_idx];
|
||
if (s.active == 0) continue;
|
||
const int slot_lots = (int)lroundf(s.size_remaining);
|
||
effective += (s.side == 0) ? slot_lots : -slot_lots;
|
||
}
|
||
|
||
const int delta = target_signed - effective;
|
||
if (delta == 0) return;
|
||
|
||
const int order_lots = (delta > 0) ? delta : -delta;
|
||
const int order_side = (delta > 0) ? 0 : 1;
|
||
|
||
// Find a free slot and seed an in-flight order.
|
||
for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) {
|
||
LimitSlot& s = orders->limits[s_idx];
|
||
if (s.active == 0) {
|
||
s.active = 2; // in-flight
|
||
s.side = (unsigned char)order_side;
|
||
s.kind = 1; // Limit
|
||
s.oco_pair = 0xFF;
|
||
s.price = (order_side == 0) ? 1.0e9f : 0.0f; // always crosses
|
||
s.size_remaining = (float)order_lots;
|
||
s.queue_position = 0.0f;
|
||
s.arrival_ts_ns = current_ts_ns + (unsigned long long)latency_ns_per_b[b];
|
||
s.gen_counter = (unsigned short)(s.gen_counter + 1);
|
||
return;
|
||
}
|
||
}
|
||
positions[b].submission_overflow += 1u;
|
||
}
|