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

fix(sp21): T2.2 Phase 8.7 — default ISV buffer in evaluator (atomic)

Browse Source 
v8 smoke (train-96wfk, commit 5694eb4df) eval pod crashed at the
same point as v7 with CUDA_ERROR_ILLEGAL_ADDRESS despite Phase 8.5's
set_branch_sizes fix. 15-minute runtime confirmed env_step decoder
no longer crashes (8.5 worked); a later kernel still OOBed.

Localization via local compute-sanitizer (RTX 3050 Ti):

  Updated gpu_backtest_validation.rs to mirror production eval-baseline
  (FEATURE_DIM=42, portfolio_dim=24, set_branch_sizes), reproducing the
  v8 crash in 1.66s locally.

  compute-sanitizer --tool=memcheck pinpointed:
    Invalid __global__ read of size 4 bytes
      at cost_net_sharpe_kernel+0x90
      by thread (32,0,0) in block (0,0,0)
      Access at 0x65c is out of bounds

  0x65c = 1628 bytes = float index 407 = OFI_IMPACT_LAMBDA_INDEX.
  Kernel reads isv[407]; isv pointer was 0 (null) because
  isv_signals_ptr defaults to 0 in constructor and set_isv_signals_ptr
  is only called by production training. Closure-path callers
  (eval-baseline) dereferenced null + slot×4 bytes.

Fix:

  Allocate zero-filled default_isv_buf of size ISV_TOTAL_DIM=536 f32
  in constructor. Wire isv_signals_ptr to its dev_ptr by default.
  Production training still overrides via set_isv_signals_ptr.

  Zero-init semantics:
    - isv[407]=0 → ofi_lambda=0 → c_ofi=0 in cost-net sharpe
      (degraded but valid; matches LobBar.ofi=0.0 placeholder)
    - Other slots default to 0 — Kelly health, controller anchors,
      etc. all see degraded-but-valid defaults

Test updates (consumer migration):
  - FEATURE_DIM 10 → 42 (production value; FEATURE_DIM < 32 makes
    gather kernel's `market_dim = feat_dim - SL_OFI_DIM (32)` negative
    → OOB; previous tests were already broken even before our changes)
  - portfolio_dim 3 → 24 (Phase 8.3+9 contract)
  - set_branch_sizes call added (Phase 8.5 contract)

Verification:
  cargo test -p ml --test gpu_backtest_validation \
    gpu_tests::test_always_long_on_uptrend --features cuda --release
    # PASSES

  compute-sanitizer --tool=memcheck <test_bin>
    # 0 CUDA errors across all 6 tests
    # (2 PnL-assertion test failures are pre-existing data-expectation
    #  issues with new FEATURE_DIM=42, not OOB bugs)

Pearls honoured:
  - feedback_no_hiding: null pointer surfaced via compute-sanitizer
    instead of silently crashing in production
  - feedback_no_partial_refactor: closure-path callers now have
    self-contained ISV setup matching training path; consumer
    migration (test FEATURE_DIM/portfolio_dim/set_branch_sizes)
    atomic with the producer-side default ISV alloc
  - pearl_no_deferrals_for_complementary_fixes: same SP cycle as
    8.3+9, 8.5 (the layer-by-layer eval rot peeling)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-12 16:37:28 +02:00
parent cb7ace0063
commit 9f2d0fffb5
3 changed files with 202 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

37
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -682,9 +682,23 @@ pub struct GpuBacktestEvaluator {
config: GpuBacktestConfig,
/// Precomputed annualization factor: sqrt(bars_per_day * 252).
annualization_factor: f32,
/// ISV signals device pointer for adaptive hold. 0 = NULL (static hold).
/// ISV signals device pointer for adaptive hold. Initialised to point at
/// `default_isv_buf` (zero-filled). Production training overrides via
/// `set_isv_signals_ptr` to share the master training-ISV bus; closure-
/// path callers (eval-baseline) inherit the safe zero-default. Pre-Phase
/// 8.7 this was `0` (null) — every kernel that read `isv[slot]` OOBed in
/// the closure path (CUDA_ERROR_ILLEGAL_ADDRESS at
/// `cost_net_sharpe_kernel+0x90` reading slot 407 → byte 1628 = 0x65c).
isv_signals_ptr: u64,
/// Default zero-filled ISV buffer owned by this evaluator. Backs
/// `isv_signals_ptr` when no external ISV is wired (closure-path eval).
/// Sized at `ISV_TOTAL_DIM=536` f32 to cover all current ISV slot
/// indices. Zero-init means `ofi_lambda=0 → c_ofi=0` in cost-net sharpe
/// (degraded but valid — matches the "no MBP-10 / no OFI in eval"
/// semantics anyway).
default_isv_buf: CudaSlice<f32>,
/// Actions buffer for forward kernel output (reused across steps).
forward_actions_buf: CudaSlice<i32>,
@@ -1154,6 +1168,22 @@ impl GpuBacktestEvaluator {
.alloc_zeros::<i32>(n_windows)
.map_err(|e| MLError::ModelError(format!("forward_actions alloc: {e}")))?;
// Phase 8.7 (2026-05-12) — default zero-filled ISV bus for the
// closure-path callers (eval-baseline). Pre-fix, `isv_signals_ptr=0`
// (null) and every kernel reading any `isv[slot]` OOBed at
// `slot×4` bytes. compute-sanitizer caught it first at
// `cost_net_sharpe_kernel+0x90` reading slot 407 → 0x65c. Sized at
// ISV_TOTAL_DIM (536 f32s) to cover all current ISV slot indices.
// Production training overrides via `set_isv_signals_ptr` to share
// the master training ISV bus.
let default_isv_buf = stream
.alloc_zeros::<f32>(super::gpu_dqn_trainer::ISV_TOTAL_DIM)
.map_err(|e| MLError::ModelError(format!("default_isv_buf alloc: {e}")))?;
let default_isv_ptr = {
let (ptr, _g) = default_isv_buf.device_ptr(&stream);
ptr
};
// ── SP15 Wave 3b — val-cost-streams buffer alloc + LobBar SoA upload ──
//
// Three input streams sourced from `window_lob_bars` (zero-padded to
@@ -1331,7 +1361,10 @@ impl GpuBacktestEvaluator {
let effective_bars_per_day = config.bars_per_day / stride;
(effective_bars_per_day * config.trading_days_per_year).sqrt()
},
isv_signals_ptr: 0,
// Phase 8.7 (2026-05-12) — point at the zero-filled default ISV
// buffer. Production training overrides via `set_isv_signals_ptr`.
isv_signals_ptr: default_isv_ptr,
default_isv_buf,
config,
forward_actions_buf,
action_select_kernel: None,

76
crates/ml/tests/gpu_backtest_validation.rs
View File

@@ -118,7 +118,7 @@ mod gpu_tests {
use super::*;
use cudarc::driver::{CudaContext, CudaSlice, CudaStream};
use ml::cuda_pipeline::gpu_backtest_evaluator::{
GpuBacktestConfig, GpuBacktestEvaluator,
DqnBacktestConfig, GpuBacktestConfig, GpuBacktestEvaluator,
};
use ml::cuda_pipeline::lob_bar::LobBar;
use ml::MLError;
@@ -180,7 +180,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 500;
let prices = generate_prices(N_BARS, 42, 0.001); // positive drift
@@ -205,10 +209,15 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Without this, `evaluate()` bails on the
// zero-b-size defensive guard. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
// Action 4 = Long100
let model = constant_action_model(4, &stream);
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
assert_eq!(metrics.len(), 1, "expected exactly one window result");
@@ -242,7 +251,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 500;
let prices = generate_prices(N_BARS, 77, -0.001); // negative drift
@@ -267,9 +280,13 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
let model = constant_action_model(4, &stream); // Always Long100
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
assert_eq!(metrics.len(), 1);
@@ -297,7 +314,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 200;
let prices = generate_prices(N_BARS, 99, 0.0);
@@ -316,10 +337,15 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Without this, `evaluate()` bails on the
// zero-b-size defensive guard. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
// Action 2 = Flat -- never enters a position
let model = constant_action_model(2, &stream);
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
assert_eq!(metrics.len(), 1);
@@ -343,7 +369,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 300;
let prices_up = generate_prices(N_BARS, 42, 0.001);
@@ -371,9 +401,13 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
let model = constant_action_model(4, &stream); // Always Long100
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
assert_eq!(metrics.len(), 2, "expected exactly 2 window results");
@@ -402,7 +436,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 500;
let prices = generate_prices(N_BARS, 42, 0.001);
@@ -420,9 +458,13 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
let model = constant_action_model(4, &stream);
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
let m = metrics
@@ -453,7 +495,11 @@ mod gpu_tests {
None => return,
};
const FEATURE_DIM: usize = 10;
// FEATURE_DIM=42 mirrors production eval-baseline.rs (Market 42 from
// extract_ml_features). The gather kernel computes `market_dim =
// feat_dim - SL_OFI_DIM (32)` so anything < 32 produces a negative
// market_dim → kernel OOB. 42 is the canonical production value.
const FEATURE_DIM: usize = 42;
const N_BARS: usize = 300;
let prices = generate_prices(N_BARS, 55, 0.001);
@@ -471,9 +517,13 @@ mod gpu_tests {
)
.expect("evaluator creation should succeed");
// Phase 8.5 (2026-05-12) — wire factored-action branch sizes for
// env_step decoder. Match production eval-baseline.rs.
evaluator.set_branch_sizes(&DqnBacktestConfig::from_network_dims((256, 256, 128, 128)));
let model = constant_action_model(4, &stream); // Always Long100
let metrics = evaluator
.evaluate(&model, 3)
.evaluate(&model, 24)
.expect("evaluation should succeed");
let m = metrics

104
docs/dqn-wire-up-audit.md
View File

@@ -15934,3 +15934,107 @@ range with std ~0.005):
- weights roughly `[0.14, 7.4]` pre-normalisation
- post-normalisation: max/min ratio ~50× — strong differential signal
- `curric_conc` (Phase 8.4 CV formula): should reach ~0.3-0.6
## 2026-05-12 — SP21 T2.2 Phase 8.7: default ISV buffer in evaluator (atomic)
### Scope (atomic single commit)
v8 smoke (train-96wfk, commit 5694eb4df) eval pod crashed at the
same point as v7 with `CUDA_ERROR_ILLEGAL_ADDRESS` despite the
Phase 8.5 `set_branch_sizes` fix. The 15-minute runtime confirmed
the env_step decoder no longer crashed (so 8.5 worked), but a later
kernel still OOBed.
### Localization via compute-sanitizer (local repro)
Updated `crates/ml/tests/gpu_backtest_validation.rs` to mirror
production eval-baseline (FEATURE_DIM=42, portfolio_dim=24,
`set_branch_sizes(&DqnBacktestConfig::from_network_dims(...))`),
reproducing the v8 crash in **1.66s locally** on RTX 3050 Ti.
`compute-sanitizer --tool=memcheck` pinpointed:
```
Invalid __global__ read of size 4 bytes
at cost_net_sharpe_kernel+0x90
by thread (32,0,0) in block (0,0,0)
Access at 0x65c is out of bounds
```
`0x65c` = 1628 bytes = float index 407 = `OFI_IMPACT_LAMBDA_INDEX`.
The kernel reads `isv[407]`; `isv` pointer was 0 (null) because the
evaluator's `isv_signals_ptr` field defaults to `0` in the
constructor and `pub fn set_isv_signals_ptr` is only called by the
production training path — closure-path callers (eval-baseline)
hit the null pointer.
### Root cause
```rust
// gpu_backtest_evaluator.rs line 1348:
isv_signals_ptr: 0, // ← null device pointer
```
Every kernel that reads `isv[slot]` then dereferences a null pointer
plus the slot offset. `cost_net_sharpe_kernel` was the first to do
so (slot 407 → byte 1628 → 0x65c).
### Fix
Allocate a zero-filled default ISV buffer in the constructor, sized
`ISV_TOTAL_DIM=536` f32s (covers all current ISV slot indices). Wire
`isv_signals_ptr` to point at it by default. Production training
still overrides via `set_isv_signals_ptr` to share the master
training-ISV bus. Closure-path callers inherit the safe zero-default.
Zero-init semantics:
- `isv[407]=0` → `ofi_lambda=0` → `c_ofi=0` in cost-net sharpe
(degraded but valid — matches the "no MBP-10 / no OFI in eval"
semantics already in place via `LobBar.ofi=0.0` placeholder).
- Other slots default to 0 — Kelly health, controller anchors,
etc. all see degraded-but-valid defaults.
### Files changed
- `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`:
added `default_isv_buf: CudaSlice<f32>` field, allocated in
constructor, `isv_signals_ptr` now defaults to its dev_ptr.
- `crates/ml/tests/gpu_backtest_validation.rs`: updated all 6
tests to mirror production eval-baseline (FEATURE_DIM 10→42,
portfolio_dim 3→24, added `set_branch_sizes` call). 4/6 pass
cleanly under compute-sanitizer; 2 failures are pre-existing
PnL-expectation assertions on synthetic data, NOT CUDA errors.
- `docs/dqn-wire-up-audit.md`: this entry.
### Verification
```
cargo test -p ml --test gpu_backtest_validation \
gpu_tests::test_always_long_on_uptrend --features cuda --release # PASSES
compute-sanitizer --tool=memcheck <test_bin> \
gpu_tests::test_always_long_on_uptrend # 0 errors
compute-sanitizer --tool=memcheck <test_bin> # 4 pass, 0 CUDA errors
# (2 PnL-assertion failures, not OOB)
```
### Pearls + invariants honoured
- **feedback_no_hiding**: null pointer surfaced; eval no longer
silently crashes
- **feedback_no_partial_refactor**: closure-path callers now have
the same self-contained ISV setup as the training path; the
production override path is preserved
- **pearl_no_deferrals_for_complementary_fixes**: same SP cycle as
8.3+9, 8.5 (the layer-by-layer eval rot)
### Expected v9 smoke
eval phase completes for all 3 folds with:
- `[DQN GPU] Fold N - Sharpe=... TotalPnL=... WinRate=...` lines
- `Report saved to:` confirmation
- `DQN - avg Sharpe=... avg WR=...` aggregate at the end
If v9 still fails, compute-sanitizer next time on the larger
production data — but Phase 8.7's local repro caught the canonical
bug and the test now exercises the same code path.