From 2fb7d7f57c574bb84af544708c84af254bfcc57a Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 2 May 2026 22:56:00 +0200 Subject: [PATCH] =?UTF-8?q?fix(data):=20Fix=2030=20Stale-B=20=E2=80=94=20b?= =?UTF-8?q?acktest=5Fplan=5Fkernel=20raw=5Fclose=20from=20prices=5Fbuf=20(?= =?UTF-8?q?post-MappedF32Buffer=20migration)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes Fix 29 audit row #13 — the last ⚠ Stale row from the Bug-1 contract drift triage. Pre-fix, `backtest_plan_state_isv` extracted raw_close as `features[bar*feat_dim + 0]`. Post Bug-1 (commit `5a5dd0fed`) `features[..+0]` is z-normed log-return, NOT raw_close. The resulting `equity = cash + position*raw_close` and `unrealized = position*(raw_close - entry_price)` formulas mixed z-normed-log-return as a dollar price, corrupting val plan_isv slots [PNL_VS_TARGET] (slot 1) and [PNL_VS_STOP] (slot 2). Other plan_isv slots (progress, conviction, drift, regime, remaining) don't depend on raw_close and were correct pre-fix. Resolution: route raw_close from the upload-once `prices` buffer (layout `[n*max_len*4]` raw OHLC). Close is at column index 3 — the same index `backtest_env_kernel.cu` already reads from for portfolio mark-to-market (line 64 of that kernel; OHLC layout is canonical across the val backtest path). Reading from `prices` aligns the val plan_isv path with env_step's source-of-truth, eliminating the mixed-units pathology end-to-end. Sites fixed: - crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu:74-100 Kernel signature: `const float* features` and `int feat_dim` parameters dropped, replaced by `const float* prices` (the [n*max_len*4] raw OHLC buffer). Raw_close read becomes `prices[(w*max_len + current_step)*4 + 3]`. Multi-line comment block documents the Bug-1 origin and the env_step parity reference. The `bool have_close` / `prices != nullptr` guard semantics preserved — callers without OHLC data fall back to `unrealized = 0` exactly as pre-fix. - crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs:~1956 Single launcher (`evaluate_dqn_graphed` chunk loop, the only invocation site of `plan_state_isv_kernel`) updated to pass `&self.prices_buf.dev_ptr` instead of `&self.features_buf.dev_ptr` and to drop the now-unused `feat_dim_i32` local + `.arg()` call. Inline comment explains the Bug-1 origin. - docs/dqn-wire-up-audit.md Stale-B row appended to Fix 30's table. The standalone Stale-B DEFERRED paragraph at the bottom replaced by the commit summary + the Fix 30 closure note (all 4 ⚠ Stale and 1 ❓ Ambiguous rows from Fix 29's deferred follow-ups now resolved). Migration scope per `feedback_no_partial_refactor`: kernel signature changed → every consumer migrates in the same commit. Single launcher; verified via `grep -rn backtest_plan_state_isv` (only the kernel definition + the gpu_backtest_evaluator launcher + load site appear). Verification: - SQLX_OFFLINE=true cargo check -p ml --offline (43.68s) clean. - SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda (1m 30s) clean; cubin recompiled via nvcc. - Pre-commit DtoD-via-pinned guard passes (the prereq commit `4d966e62f` migrated `gpu_backtest_evaluator.rs`'s buffers to MappedF32Buffer, eliminating the 5 `_via_pinned` callers that had blocked any prior staging of this file). Refs Fix 29 row #13. `feedback_no_partial_refactor` (single launcher migrated alongside kernel signature change in one commit), `feedback_no_functionality_removal` (PNL_VS_TARGET / PNL_VS_STOP slots preserved — only their data source corrected; the audit's "drop the slots" alternative explicitly rejected), `feedback_no_hiding` (no fallback to z-normed reads remaining; kernel either gets real raw_close or falls through to `have_close=false` with `unrealized=0`, identical to pre-fix smoke-test semantics where prices==NULL), `feedback_no_cpu_compute_strict` n/a (zero new host-side compute), `feedback_no_htod_htoh_only_mapped_pinned` already satisfied (`prices_buf` is `MappedF32Buffer` post the prereq migration), `feedback_trust_code_not_docs` (the kernel comment said "raw_close from features buffer" for months — accurate-when-written pre-Bug-1, stale-after-Bug-1; verify-against-code disambiguates). Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/backtest_plan_kernel.cu | 23 +++++++++++++------ .../cuda_pipeline/gpu_backtest_evaluator.rs | 13 ++++++++--- docs/dqn-wire-up-audit.md | 10 +++++++- 3 files changed, 35 insertions(+), 11 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu index 70c7ee323..d29bfdb78 100644 --- a/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu @@ -74,10 +74,9 @@ extern "C" __global__ void backtest_plan_state_isv( float* __restrict__ plan_state, /* [N, 7] persistent plan state (in-place) */ const float* __restrict__ portfolio, /* [N, 8] val portfolio state */ const float* __restrict__ isv_signals, /* [ISV_TOTAL_DIM] for slot [11] */ - const float* __restrict__ features, /* [N * max_len * feat_dim] for raw_close extraction. NULL → skip P&L ratios */ + const float* __restrict__ prices, /* [N * max_len * 4] raw OHLC; close at col 3. NULL → skip P&L ratios */ int current_step, /* current bar index within window */ int max_len, /* bars per window */ - int feat_dim, /* feature stride (raw_close at offset 0) */ float* __restrict__ plan_isv_out, /* [N, SL_PORTFOLIO_PLAN_DIM=7] plan_isv output (consumed by next state gather) */ int n_windows ) { @@ -89,13 +88,23 @@ extern "C" __global__ void backtest_plan_state_isv( float entry_price = portfolio[w * 8 + 3]; float hold_time = portfolio[w * 8 + 5]; - /* Extract raw_close from features buffer: layout is - * [window, bar, feat] with raw_close at feat index 0. */ + /* Fix 30 Stale-B (2026-05-02): pre-fix this kernel read raw_close from + * `features[bar*feat_dim + 0]`, but post Bug-1 (commit 5a5dd0fed) + * `features[..+0]` is z-normed log-return, NOT raw_close. The + * resulting `equity = cash + position*raw_close` and + * `unrealized = position*(raw_close - entry_price)` formulas mixed + * z-normed-log-return as a dollar price, corrupting val plan_isv + * slots [PNL_VS_TARGET] and [PNL_VS_STOP]. Post-fix raw_close comes + * from the upload-once `prices` buffer (layout [n*max_len*4] OHLC, + * close column = index 3) — the same source the env_step kernel + * reads from for portfolio mark-to-market (`backtest_env_kernel.cu:64`). + * The other plan_isv slots (progress, conviction, drift, regime, + * remaining) don't depend on raw_close and were correct pre-fix. */ float raw_close = 0.0f; bool have_close = false; - if (features != nullptr && current_step >= 0 && current_step < max_len) { - long long feat_off = ((long long)w * max_len + current_step) * feat_dim; - raw_close = features[feat_off]; + if (prices != nullptr && current_step >= 0 && current_step < max_len) { + long long price_off = ((long long)w * max_len + current_step) * 4 + 3; + raw_close = prices[price_off]; have_close = true; } diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 057048c7c..b0c3d508e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -1948,11 +1948,19 @@ impl GpuBacktestEvaluator { // (b) Launch plan_state_isv kernel — updates plan_state in-place // and writes plan_isv_buf[N, 7] consumed by the next chunk's // first launch_gather() call. + // + // Fix 30 Stale-B (2026-05-02): the kernel previously took + // `features` + `feat_dim` to extract raw_close as + // `features[bar*feat_dim + 0]`, but post-Bug-1 that slot is + // z-normed log-return (NOT raw_close). Now passes the + // upload-once `prices` buffer (layout [n*max_len*4] OHLC, + // close at col 3) — the same source backtest_env_kernel.cu + // already uses for portfolio mark-to-market. `feat_dim` arg + // dropped because the kernel no longer needs feature stride. let plan_blocks = ((n as u32) + 255) / 256; let n_i32 = n as i32; let last_step_i32 = last_step as i32; let max_len_i32 = self.max_len as i32; - let feat_dim_i32 = self.feature_dim as i32; unsafe { self.stream .launch_builder(plan_kernel) @@ -1960,10 +1968,9 @@ impl GpuBacktestEvaluator { .arg(&self.plan_state_buf) .arg(&self.portfolio_buf.dev_ptr) .arg(&self.isv_signals_ptr) - .arg(&self.features_buf.dev_ptr) + .arg(&self.prices_buf.dev_ptr) .arg(&last_step_i32) .arg(&max_len_i32) - .arg(&feat_dim_i32) .arg(&self.plan_isv_buf) .arg(&n_i32) .launch(LaunchConfig { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c5ff18f1e..8d5e482f6 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3817,6 +3817,7 @@ Closes the four ⚠ Stale and one ❓ Ambiguous sites Fix 29 enumerated. Each si | 15 | `crates/ml/src/hyperopt/adapters/dqn.rs:2492` (`val_close_prices`) | ⚠ Stale — same `target[0]` / `fv[3]` pattern as #14 in the HPO val backtest path | Same fix shape as #14: `target[TARGET_RAW_CLOSE]`, dead fallback deleted, named constant import. Affects HPO val Sharpe + window aggregation. | ✅ | | 12 | `crates/ml/src/cuda_pipeline/scripted_policy_kernel.cu:59-65` (seed-phase MOMENTUM / MEAN_REV / VWAP_DEV `recent_ret`) | ⚠ Stale — reads `state[MARKET_START]` (z-normed log-return post Bug 1) as `close_now`, then computes `(close_now − prev_close) / prev_close` where `prev_close` is raw_close. Mixed units made `recent_ret` meaningless; seed-phase MOMENTUM/MEAN_REV/VWAP_DEV branches degenerated to noise-floor below the ±0.0001f cutoffs | Kernel signature extended with `const float* targets`, `const int* episode_starts`, `int t`, `int total_bars`. Per-thread `bar_idx = episode_starts[i] + t`, then `close_now = targets[bar_idx*6 + TARGET_RAW_CLOSE]` matching env_step's source (`experience_kernels.cu:1769`). Bounds-clamp to `[0, total_bars-1]` mirrors env_step's out-of-bounds early-return. The previous `bar_idx` parameter (per-step `t` mixed into the LCG seed) renamed `t`; the LCG seed now mixes per-thread `bar_idx = episode_starts[i] + t` instead — stronger entropy across episodes, no behavioural regression (UNIFORM policy still produces a 4-direction uniform distribution). Single launcher (`gpu_experience_collector.rs:3735`) migrated in the same commit; `total_bars` already in scope from line 3301. Recent_ret signal preserved per `feedback_no_functionality_removal`; the audit's CUSUM-substitution alternative path explicitly rejected (the recent_ret signal is the seed-policy contract, not the implementation accident). NO new `PS_*` slot or PORTFOLIO_STRIDE bump required — the targets buffer is the canonical raw_close source and routing it directly avoids cascading through ~12 PS_STRIDE consumers. | ✅ | | 18 | `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu:1050-1096` (`phantom_liquidity_gbm` synthetic-feature overlay) | ❓ Ambiguous — writes `log_ret` to `market_features[bar*market_dim + 0]` and `+3` for synthetic GBM episodes; downstream consumer contract unverifiable | Investigation: `grep -rn "phantom_liquidity_gbm" crates/ml/src/ services/ crates/ bin/` returns ZERO callers — no `load_function`, no `launch_builder`, no Rust-side launcher. Kernel is dead code. Per `feedback_no_hiding` + `feedback_wire_everything_up`: deleted the kernel definition; replaced with a 23-line explanatory comment block documenting why (no consumer existed, contract unverifiable, hypothetical re-wire would have introduced exactly the Bug-1 contract drift the audit was triaging). Re-introducing GBM-overlay augmentation must land caller wiring in the same commit per `feedback_wire_everything_up`. | ✅ | +| 13 | `crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu:77-100` (val plan_isv) | ⚠ Stale — reads `features[bar*feat_dim + 0]` (post-Bug-1 z-normed log-return) as raw_close → `equity = cash + position*raw_close` and `unrealized = position*(raw_close - entry_price)` mixed wrong units → val plan_isv slots [PNL_VS_TARGET], [PNL_VS_STOP] corrupted | Kernel signature: `const float* features` + `int feat_dim` parameters replaced by `const float* prices` (the [n*max_len*4] raw OHLC buffer the env_step kernel already reads from). Raw_close source switched to `prices[(w*max_len + current_step)*4 + 3]` — close column of the OHLC layout, identical to the index `backtest_env_kernel.cu` uses for portfolio mark-to-market. Single launcher (`gpu_backtest_evaluator.rs::evaluate_dqn_graphed` chunk loop) updated in the same commit per `feedback_no_partial_refactor` to pass `&self.prices_buf.dev_ptr` instead of `&self.features_buf.dev_ptr` and to drop the `feat_dim_i32` arg. PNL_VS_TARGET / PNL_VS_STOP slots now compute on dollar-denominated unrealized P&L. Other plan_isv slots (progress, conviction, drift, regime, remaining) didn't depend on raw_close and were correct pre-fix. Prereq commit `4d966e62f` migrated the file's CudaSlice buffers to MappedF32Buffer (DtoD-via-pinned guard had blocked any prior staging). | ✅ | **Stale-A commit** (rows #14, #15): - `crates/ml/src/trainers/dqn/trainer/metrics.rs` — column-index swap + import + dead-fallback deletion + comment. @@ -3847,4 +3848,11 @@ References: this is the dead-code branch of the contract decision. The kernel ha This is a pure structural migration — orthogonal to Bug-1 contract drift, but a hard prerequisite for the Stale-B kernel-side fix (which has to stage `gpu_backtest_evaluator.rs` to thread `prices_buf` into the `backtest_plan_state_isv` launcher). Per `feedback_no_partial_refactor` every consumer of the field-type change migrates in this commit; per `feedback_no_htod_htoh_only_mapped_pinned` the migration eliminates the last 5 `_via_pinned` callers in this file. -**Stale-B (row #13) — DEFERRED**: kernel-side fix follows in the next commit. The migration above unblocks staging `gpu_backtest_evaluator.rs`; the actual Stale-B fix routes `prices_buf` into `backtest_plan_state_isv` and reads `prices[bar*4 + 3]` (raw_close column) instead of `features[bar*feat_dim + 0]` (z-normed log-return). Affects val plan_isv slots [PNL_VS_TARGET], [PNL_VS_STOP] only — gradient flow on the main training graph is unaffected (Fix 29 `Summary` line confirms). +**Stale-B commit** (row #13): +- `crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu` — kernel signature: `const float* features` + `int feat_dim` parameters dropped, replaced by `const float* prices` (the [n*max_len*4] raw OHLC buffer). Raw_close read becomes `prices[(w*max_len + current_step)*4 + 3]` — same source `backtest_env_kernel.cu` reads from for portfolio mark-to-market. Multi-line comment block documents the Bug-1 origin and the env_step parity reference. +- `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` — single launcher (`evaluate_dqn_graphed` chunk loop, ~line 1956) updated to pass `&self.prices_buf.dev_ptr` instead of `&self.features_buf.dev_ptr` and to drop the `feat_dim_i32` local + arg. Inline comment explains the Bug-1 origin. +- `docs/dqn-wire-up-audit.md` — Stale-B row appended in Fix 30's table; Stale-B DEFERRED paragraph replaced with this commit summary. + +References: Stale-B closes the last ⚠ Stale row from Fix 29's audit. `feedback_no_partial_refactor` (single launcher migrated alongside kernel signature change in one commit), `feedback_no_functionality_removal` (PNL_VS_TARGET / PNL_VS_STOP slots preserved — only their data source corrected; the audit's "drop the slots" alternative explicitly rejected), `feedback_no_hiding` (no fallback to z-normed reads remaining; kernel either gets real raw_close from `prices` or falls back to `unrealized=0` when `prices==NULL`, matching the existing `have_close=false` semantics for callers without OHLC data — same as pre-fix behaviour for the smoke-test paths), `feedback_no_cpu_compute_strict` n/a (zero new host-side compute), `feedback_no_htod_htoh_only_mapped_pinned` already satisfied (`prices_buf` is `MappedF32Buffer` post the prereq commit `4d966e62f`). + +**Fix 30 closure**: all four ⚠ Stale rows (#12 Stale-C, #13 Stale-B, #14/#15 Stale-A) and the one ❓ Ambiguous row (#18 Ambiguous-A) from Fix 29's deferred follow-ups are now resolved. Bug-1 contract drift triage is complete; the production hot path was already clean post-Fix-29's host-side `vol_normalizer` deletion, and the val/HPO/seed-phase derived metrics now read raw_close from the correct source on every path.