diff --git a/crates/ml-backtesting/cuda/order_match.cu b/crates/ml-backtesting/cuda/order_match.cu
index 1814fed1c..55a192840 100644
--- a/crates/ml-backtesting/cuda/order_match.cu
+++ b/crates/ml-backtesting/cuda/order_match.cu
@@ -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;
diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu
index b99539a50..dcc39fccf 100644
--- a/crates/ml-backtesting/cuda/resting_orders.cu
+++ b/crates/ml-backtesting/cuda/resting_orders.cu
@@ -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,
diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs
index a432c4a80..f4ca95e4b 100644
--- a/crates/ml-backtesting/src/sim/mod.rs
+++ b/crates/ml-backtesting/src/sim/mod.rs
@@ -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)?;
         }