gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(ml-backtesting): per-level price-range validation in walk_ask_for_buy/walk_bid_for_sell — eliminates sized-but-bad-priced sentinel propagation

Browse Source 
v2 NaN instrumentation (S1.21) localized the bug to apply_fill_to_pos's
open-from-flat branch writing pos.vwap_entry = avg_px where avg_px was 0
(194 cases) or > 21M finite (216 cases). The arithmetic was clean — root
cause is walk_ask_for_buy/walk_bid_for_sell consuming size from deep
levels (k=1..9) that have lvl_sz > 0 but lvl_px = 0 (or huge sentinel).
MBP-10 fills empty depth slots beyond available levels with these
sentinels.

Existing per-level filter checked lvl_sz <= 0, !isfinite(lvl_sz),
!isfinite(lvl_px) — but allowed lvl_px = 0 and lvl_px > max range.

Fix: thread per-backtest min_reasonable_px / max_reasonable_px (already
uploaded for the top-of-book skip in book_update_apply_snapshot via
S1.19) through to walk_*, and reject any level whose price falls
outside [min_px, max_px]. Same fix shape as Bug C-b (top-of-book skip),
now applied at all 10 depths.

Changes: resting_orders.cu (walk_* signatures + kernel param + 2 call
sites), order_match.cu (walk_* signatures + kernel param + 1 call site),
sim/mod.rs (submit_market + step_resting_orders launch args). 19 CUDA
tests pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 13:49:28 +02:00
parent fc41440dc9
commit 76ec68c1c9
3 changed files with 57 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files

33
crates/ml-backtesting/cuda/order_match.cu
View File

@@ -12,8 +12,14 @@
// Walk ask side for a buy of `size`. Returns total notional cost in
// price-units. Writes filled_lots out-param.
//
// min_px / max_px: per-backtest price-range bounds (from upload_price_range /
// S1.19). Levels with lvl_px outside [min_px, max_px] are MBP-10 sentinels
// for empty depth slots — consuming them writes avg_px=0 or avg_px>21M into
// pos.vwap_entry (S1.22 root-cause fix).
__device__ static float walk_ask_for_buy(
const Book& book, float size, float& filled_lots
const Book& book, float size, float& filled_lots,
float min_px, float max_px
) {
float cost = 0.0f;
filled_lots = 0.0f;
@@ -22,7 +28,8 @@ __device__ static float walk_ask_for_buy(
if (size <= 0.0f) break;
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;
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)
|| lvl_px < min_px || lvl_px > max_px) continue;
const float take = (size < lvl_sz) ? size : lvl_sz;
cost += take * lvl_px;
filled_lots += take;
@@ -32,7 +39,8 @@ __device__ static float walk_ask_for_buy(
}
__device__ static float walk_bid_for_sell(
const Book& book, float size, float& filled_lots
const Book& book, float size, float& filled_lots,
float min_px, float max_px
) {
float cost = 0.0f;
filled_lots = 0.0f;
@@ -41,7 +49,8 @@ __device__ static float walk_bid_for_sell(
if (size <= 0.0f) break;
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;
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)
|| lvl_px < min_px || lvl_px > max_px) continue;
const float take = (size < lvl_sz) ? size : lvl_sz;
cost += take * lvl_px;
filled_lots += take;
@@ -55,10 +64,15 @@ __device__ static float walk_bid_for_sell(
//
// `targets` is a [n_backtests, 2]-shaped i32 array: { side, size }.
// side: 0=buy, 1=sell, 2=no-op (skip this backtest this step).
//
// S1.22: min_reasonable_px / max_reasonable_px added so walk_* can reject
// sentinel-priced depth levels (zero-priced or huge-priced empty MBP-10 slots).
extern "C" __global__ void submit_market_immediate(
const Book* __restrict__ books,
const int* __restrict__ targets, // [n_backtests, 2]
Pos* __restrict__ positions, // [n_backtests]
const int* __restrict__ targets, // [n_backtests, 2]
Pos* __restrict__ positions, // [n_backtests]
const float* __restrict__ min_reasonable_px, // [n_backtests]
const float* __restrict__ max_reasonable_px, // [n_backtests]
int n_backtests
) {
int b = blockIdx.x;
@@ -71,10 +85,13 @@ extern "C" __global__ void submit_market_immediate(
const Book& book = books[b];
Pos& pos = positions[b];
const float min_px = min_reasonable_px[b];
const float max_px = max_reasonable_px[b];
float filled = 0.0f;
const float cost = (side == 0)
? walk_ask_for_buy (book, (float)sz_in, filled)
: walk_bid_for_sell(book, (float)sz_in, filled);
? walk_ask_for_buy (book, (float)sz_in, filled, min_px, max_px)
: walk_bid_for_sell(book, (float)sz_in, filled, min_px, max_px);
if (filled <= 0.0f) return;
const float avg_px = cost / filled;
const float sgn = (side == 0) ? 1.0f : -1.0f;

36
crates/ml-backtesting/cuda/resting_orders.cu
View File

@@ -26,25 +26,38 @@
// 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) {
//
// min_px / max_px: per-backtest price-range bounds (from upload_price_range /
// S1.19). Levels with lvl_px outside [min_px, max_px] are MBP-10 sentinels
// for empty depth slots — consuming them writes avg_px=0 or avg_px>21M into
// pos.vwap_entry (194 zero + 216 huge sentinel writes, S1.22 root-cause).
__device__ static float walk_ask_for_buy(
const Book& book, float size, float& filled,
float min_px, float max_px
) {
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;
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)
|| lvl_px < min_px || lvl_px > max_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) {
__device__ static float walk_bid_for_sell(
const Book& book, float size, float& filled,
float min_px, float max_px
) {
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;
if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)
|| lvl_px < min_px || lvl_px > max_px) continue;
const float take = (size < lvl_sz) ? size : lvl_sz;
cost += take * lvl_px; filled += take; size -= take;
}
@@ -249,6 +262,9 @@ extern "C" __global__ void resting_orders_step(
unsigned int* __restrict__ vwap_session_gap_was_bad,
float* __restrict__ last_bad_vwap,
unsigned int* __restrict__ last_bad_path,
// S1.22: per-level price-range validation in walk_ask_for_buy/walk_bid_for_sell.
const float* __restrict__ min_reasonable_px, // [n_backtests]
const float* __restrict__ max_reasonable_px, // [n_backtests]
int n_backtests
) {
int b = blockIdx.x;
@@ -258,6 +274,10 @@ extern "C" __global__ void resting_orders_step(
Orders& orders = *reinterpret_cast<Orders*>(orders_base + (size_t)b * orders_bytes);
Pos& pos = *reinterpret_cast<Pos*>(pos_base + (size_t)b * pos_bytes);
// S1.22: per-backtest price bounds for walk_* level filtering.
const float min_px = min_reasonable_px[b];
const float max_px = max_reasonable_px[b];
// 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
@@ -354,8 +374,8 @@ extern "C" __global__ void resting_orders_step(
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);
? walk_ask_for_buy (book, s.size_remaining, filled, min_px, max_px)
: walk_bid_for_sell(book, s.size_remaining, filled, min_px, max_px);
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,
@@ -405,8 +425,8 @@ extern "C" __global__ void resting_orders_step(
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);
? walk_ask_for_buy (book, st.size, filled, min_px, max_px)
: walk_bid_for_sell(book, st.size, filled, min_px, max_px);
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,

4
crates/ml-backtesting/src/sim/mod.rs
View File

@@ -810,6 +810,8 @@ impl LobSimCuda {
.arg(&self.books_d)
.arg(&self.market_targets_d)
.arg(&mut self.pos_d)
.arg(&self.min_reasonable_px_d)
.arg(&self.max_reasonable_px_d)
.arg(&n)
.launch(cfg)?;
}
@@ -1390,6 +1392,8 @@ impl LobSimCuda {
.arg(&mut self.vwap_session_gap_was_bad_d)
.arg(&mut self.last_bad_vwap_d)
.arg(&mut self.last_bad_path_d)
.arg(&self.min_reasonable_px_d)
.arg(&self.max_reasonable_px_d)
.arg(&n)
.launch(cfg)?;
}