cd82f9a4a0
Adds the two sweep axes that the spec's deployability grid needs but were missing from the kernels: Threshold gate (decision_policy.cu, both kernels): - New per-backtest `threshold_per_b` array kernel arg. - Pre-Kelly prelude: if max_h |alpha[h] - 0.5| * 2 < threshold[b], emit noop and return. Kept deterministic from alpha alone so the threshold pre-registration step (p60-p95 absolute calibration on a validation window, future P6) reflects exactly what gets gated in deployment. Per-fill cost integration (resting_orders.cu / apply_fill_to_pos): - apply_fill_to_pos signature grows three args: b, cost_per_lot_per_side_per_b, total_fees_per_b. Single insertion point at line 90. - After the close-leg realized_pnl math runs (so the gross unwind P&L is preserved), deduct fill_cost = filled_lots * cost_per_lot_per_side[b] from pos.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 from realistic return distribution. - All 3 apply_fill_to_pos call sites in step_resting_orders updated. order_match.cu's submit_market_immediate path is dead code in the post-P1 flow (everything routes through seed_inflight_limits_batched → step_resting_orders → apply_fill_to_pos) so not touched here. BatchedSimConfig + UniformSimParams + BacktestHarnessConfig gain threshold + cost_per_lot_per_side fields. All UniformSimParams constructors in tests and main.rs updated with defaults (0.0, 0.0 = gate disabled, frictionless). Regression: - threshold_gate_skips_low_conviction (p=0.51 + threshold=0.10 → noop) - threshold_gate_allows_high_conviction (p=0.8 + threshold=0.10 → buy 1+) - threshold_zero_is_passthrough (sanity) - All P1+P2+P3 tests continue to pass via the new ABI. cost_deducted_at_each_fill + kelly_state_sees_net_return end-to-end tests deferred — they require a full submit_market → fill → close sequence, which the production smoke exercises. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
476 lines
21 KiB
Plaintext
476 lines
21 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];
|
|
if (lvl_sz <= 0.0f) continue;
|
|
const float take = (size < lvl_sz) ? size : lvl_sz;
|
|
cost += take * book.ask_px[k]; 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];
|
|
if (lvl_sz <= 0.0f) continue;
|
|
const float take = (size < lvl_sz) ? size : lvl_sz;
|
|
cost += take * book.bid_px[k]; 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.
|
|
__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
|
|
) {
|
|
if (filled_lots <= 0.0f) return;
|
|
const float avg_px = total_cost / filled_lots;
|
|
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;
|
|
pos.realized_pnl += realised;
|
|
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;
|
|
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,
|
|
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);
|
|
|
|
// (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);
|
|
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);
|
|
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);
|
|
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). Scans this backtest's own
|
|
// MAX_LIMITS slot range for an `active == 0` slot; if found, writes the
|
|
// in-flight Limit with arrival_ts = current_ts + latency_ns_per_b[b]. If
|
|
// all 32 slots are in use, increments pos.submission_overflow.
|
|
//
|
|
// noop targets (side == 2) and zero-size targets are skipped — the prior
|
|
// host loop had the same behaviour.
|
|
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 size = market_targets[b * 2 + 1];
|
|
if (side == 2 || size <= 0) return; // noop / no order to seed
|
|
|
|
Orders* orders = reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
|
|
|
|
// Scan per-backtest slots for a free one.
|
|
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)side;
|
|
s.kind = 1; // Limit
|
|
s.oco_pair = 0xFF;
|
|
s.price = (side == 0) ? 1.0e9f : 0.0f; // always crosses
|
|
s.size_remaining = (float)size;
|
|
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;
|
|
}
|
|
}
|
|
// No free slot — record overflow.
|
|
positions[b].submission_overflow += 1u;
|
|
}
|