diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 44399fc6d..fa51d0650 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -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, + /// Actions buffer for forward kernel output (reused across steps). forward_actions_buf: CudaSlice, @@ -1154,6 +1168,22 @@ impl GpuBacktestEvaluator { .alloc_zeros::(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::(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, diff --git a/crates/ml/tests/gpu_backtest_validation.rs b/crates/ml/tests/gpu_backtest_validation.rs index 157853372..ebdef2932 100644 --- a/crates/ml/tests/gpu_backtest_validation.rs +++ b/crates/ml/tests/gpu_backtest_validation.rs @@ -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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index cebe64383..fdb983c23 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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` 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 \ + gpu_tests::test_always_long_on_uptrend # 0 errors + +compute-sanitizer --tool=memcheck # 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.