fix(data): Fix 30 Stale-B — backtest_plan_kernel raw_close from prices_buf (post-MappedF32Buffer migration)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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;
|
||||
}
|
||||
|
||||
|
||||
@@ -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 {
|
||||
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user