fix(data): target stride/column drift — host + GPU consumer cleanup
Two latent bugs converged in the cancelled 50-epoch run (train-multi-seed-p5qzw at96769d171, label_scale=808 vs smoke baseline 22-28): Bug A (latent since063fd2716, 2026-04-19): TARGET_DIM was bumped 4→6 in fxcache (added raw_open at col 4 + mid_price_open at col 5), but N consumer kernels in experience_kernels.cu + dt_kernels.cu hardcoded stride 4 when indexing targets[bar*4+col]. Reading at the old stride against the stride-6 buffer slid every lookup into the wrong bar's data. Smoke harness paths (multi_fold_convergence) didn't exercise expert_action_override / compute_difficulty_scores / hindsight_relabel heavily; production 368k-bar run surfaced the drift via the compounded label_scale = 30× expected magnitude. Bug B (introduced today at5a5dd0fed, 2026-05-02): the Bug 1 fix changed targets[0] from raw_close to log-return-normalized preproc_close. training_loop.rs::epoch_vol_normalizer's Welford pass still read targets[0] expecting raw_close → output was ln(log_return/log_return) ~ stddev 0.6 → triggered sanity-band warning + default fallback (5e-4); but downstream label_scale consumers fed the corrupt value through. Sites fixed: HOST (1 site): - training_loop.rs:570-573, :603 — w[i].1[0] → w[i].1[TARGET_RAW_CLOSE] + warning message updated GPU (5 sites in experience_kernels.cu): - line 3768 — kernel param doc OHLCV → fxcache TARGET_DIM=6 - lines 3806-3808 — bar*4+3 → bar*6+2 (raw_close column) - lines 3974-3975 — i*4+2 → i*6+2 (stride-only) - lines 4015, 4020 — bar*4+2 → bar*6+2 (stride-only) - lines 3400-3404 — t_offset = global_idx*4 → *6 (stride-only) - lines 1553-1561 — kernel header doc updated to 6-column layout GPU (1 site in dt_kernels.cu): - line 793 — kernel param doc OHLC → fxcache TARGET_DIM=6 - lines 803, 807 — i*4+3 → i*6+2 (raw_close col 2) HOST docstring (decision_transformer.rs:1049) — [num_bars, 4] → TARGET_DIM=6 with reference to fxcache constant module. Structural prevention: 6 named column constants (TARGET_PREPROC_CLOSE / TARGET_PREPROC_NEXT / TARGET_RAW_CLOSE / TARGET_RAW_NEXT / TARGET_RAW_OPEN / TARGET_MID_OPEN) added to crates/ml/src/fxcache.rs co-located with the existing TARGET_DIM (promoted from private to pub). Co-locating in the contract-owner module means future renames or column adds force every consumer to update at the same call site. Compile-time density test asserts columns are dense and exhaustive. Host-side consumer training_loop.rs imports TARGET_RAW_CLOSE; GPU kernels reference the constant module in comments + use literal stride 6 with a header block documenting the contract (kernels can't import Rust constants). PPO consumer NOT in scope: ppo_experience_kernel.cu reads at stride 4 but PPO has its own set_raw_market_data path uploading 4 columns from Vec<f64>. PPO write/read pair is internally consistent at stride 4, unaffected by fxcache stride bump. Production train_baseline_rl.rs flow doesn't invoke set_raw_market_data, so PPO GPU collector is currently orphaned. Consolidating PPO onto fxcache target buffer is a separate refactor. Verification: - SQLX_OFFLINE=true cargo check -p ml --offline clean - cargo build -p ml --release --offline --features cuda clean (cubin compiles via nvcc; release build under 2 min) - cargo test -p ml --lib -- target_layout 1/1 pass (target_columns_dense_and_exhaustive) - audit grep "targets[var * 4 +]" returns ZERO hits in production - re-run of train-multi-seed-p5qzw will show label_scale ~22-30 (not 808) — gating production validation Refs: cancelled train-multi-seed-p5qzw at96769d171, 50-epoch validation blocker. Bug origins:063fd2716(target_dim 4→6 bump) and5a5dd0fed(Bug 1 column-0 semantic fix). feedback_no_partial_refactor (every consumer of shared buffer contract migrates in same commit), feedback_trust_code_not_docs (kernel doc OHLCV/OHLC strings stale post-bump), feedback_no_quickfixes (proper structural fix not one-line patch), feedback_wire_everything_up (column constants land co-located with fxcache::TARGET_DIM so writer + caller share single contract module). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -1046,7 +1046,8 @@ impl DecisionTransformer {
|
||||
/// * `features_gpu` - mapped-pinned buffer `[num_bars, feat_dim]` (42-dim
|
||||
/// feature vectors). Kernel reads via `dev_ptr`; host writes happen at
|
||||
/// fxcache load.
|
||||
/// * `targets_gpu` - mapped-pinned buffer `[num_bars, 4]` (OHLC per bar).
|
||||
/// * `targets_gpu` - mapped-pinned buffer `[num_bars, TARGET_DIM=6]`
|
||||
/// (fxcache target layout; see `crate::fxcache::TARGET_*` constants).
|
||||
/// * `num_bars` - Number of bars in the dataset
|
||||
/// * `gamma` - Discount factor for return-to-go
|
||||
///
|
||||
|
||||
@@ -790,7 +790,7 @@ extern "C" __global__ void dt_build_trajectories_kernel(
|
||||
* Block: (256, 1, 1)
|
||||
* ══════════════════════════════════════════════════════════════════════ */
|
||||
extern "C" __global__ void dt_compute_rewards_actions_kernel(
|
||||
const float* __restrict__ targets, /* [num_bars, 4] — O,H,L,C per bar */
|
||||
const float* __restrict__ targets, /* [num_bars, TARGET_DIM=6] — fxcache targets */
|
||||
float* __restrict__ rewards, /* [num_bars] */
|
||||
int* __restrict__ actions, /* [num_bars] */
|
||||
int num_bars,
|
||||
@@ -799,12 +799,20 @@ extern "C" __global__ void dt_compute_rewards_actions_kernel(
|
||||
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i >= num_bars) return;
|
||||
|
||||
/* Close price is at offset 3 (O=0, H=1, L=2, C=3) — F32 for division */
|
||||
float close_t = (float)targets[i * 4 + 3];
|
||||
/* Targets stride per `crates/ml/src/fxcache.rs`
|
||||
* (TARGET_DIM=6): col 0=preproc_close, 1=preproc_next, 2=raw_close,
|
||||
* 3=raw_next (FUTURE), 4=raw_open, 5=mid_open. Read raw_close at col 2.
|
||||
*
|
||||
* Pre-`063fd2716` the stride was 4 and this kernel read `i * 4 + 3`
|
||||
* (raw_next column) — the stale "OHLC" docstring was incorrect: the
|
||||
* fxcache never had an OHLC column layout, only the close-row variants.
|
||||
* Reading at the old stride against the new stride-6 buffer slid every
|
||||
* lookup into the wrong bar's data. */
|
||||
float close_t = (float)targets[i * 6 + 2];
|
||||
|
||||
float reward;
|
||||
if (i < num_bars - 1) {
|
||||
float close_t1 = (float)targets[(i + 1) * 4 + 3];
|
||||
float close_t1 = (float)targets[(i + 1) * 6 + 2];
|
||||
/* Avoid division by zero */
|
||||
reward = (close_t > 1e-8f) ? (close_t1 / close_t - 1.0f) : 0.0f;
|
||||
} else {
|
||||
|
||||
@@ -1550,11 +1550,15 @@ extern "C" __global__ void experience_action_select(
|
||||
*
|
||||
* Grid: ceil(N / 256), Block: 256. One thread per episode.
|
||||
*
|
||||
* targets layout: [total_bars, 4]
|
||||
* targets layout: [total_bars, TARGET_DIM=6]
|
||||
* See `crates/ml/src/fxcache.rs` for the canonical
|
||||
* contract; mirror below for human readers (kernels can't import Rust):
|
||||
* [0] preproc_close — log-return-normalised close (network only)
|
||||
* [1] preproc_next — log-return-normalised next close (unused here)
|
||||
* [2] raw_close — raw close price (for position cost + tx)
|
||||
* [3] raw_next — raw next-bar close (UNUSED in reward path)
|
||||
* [3] raw_next — raw next-bar close (UNUSED — FUTURE info)
|
||||
* [4] raw_open — raw open price
|
||||
* [5] mid_price_open — MBP-10 midpoint at bar open (fallback raw_open)
|
||||
*
|
||||
* portfolio_states layout: [N, PORTFOLIO_STRIDE=40] (read-write, float)
|
||||
* See file header for field definitions.
|
||||
@@ -3397,11 +3401,16 @@ extern "C" __global__ void portfolio_sim_kernel(
|
||||
int global_idx = batch_start + b;
|
||||
int action_idx = actions[b];
|
||||
|
||||
int t_offset = global_idx * 4;
|
||||
/* Targets stride per `crates/ml/src/fxcache.rs`
|
||||
* (TARGET_DIM=6): col 0=preproc_close, 1=preproc_next, 2=raw_close,
|
||||
* 3=raw_next (FUTURE), 4=raw_open, 5=mid_open. Pre-`063fd2716` the
|
||||
* stride was 4 with the same column semantics for cols 0-3; only
|
||||
* the stride bumps 4 → 6 here. */
|
||||
int t_offset = global_idx * 6;
|
||||
float current_close = (targets[t_offset + 0]);
|
||||
float next_close = (targets[t_offset + 1]);
|
||||
float current_close_raw = (targets[t_offset + 2]);
|
||||
(void)(targets[t_offset + 3]); /* next_close_raw: FUTURE — must NOT be used */
|
||||
(void)(targets[t_offset + 3]); /* raw_next: FUTURE — must NOT be used */
|
||||
|
||||
float price = (current_close_raw != 0.0f) ? current_close_raw : current_close;
|
||||
if (price <= 0.0f) price = 1.0f;
|
||||
@@ -3765,7 +3774,7 @@ extern "C" __global__ void atom_stats_finalize(
|
||||
*/
|
||||
extern "C" __global__ void expert_action_override(
|
||||
int* out_actions, /* [N] Q-network actions (overwritten) */
|
||||
const float* __restrict__ targets, /* [total_bars, 4] OHLCV price data */
|
||||
const float* __restrict__ targets, /* [total_bars, TARGET_DIM=6] fxcache targets */
|
||||
const float* __restrict__ features, /* [total_bars, MARKET_DIM] market features */
|
||||
const float* __restrict__ portfolio_states, /* [N, PORTFOLIO_STRIDE] per-episode state */
|
||||
const int* __restrict__ episode_starts, /* [N] bar offset per episode */
|
||||
@@ -3802,10 +3811,17 @@ extern "C" __global__ void expert_action_override(
|
||||
float adx = (adx_bf);
|
||||
float cusum = (cusum_bf);
|
||||
|
||||
/* 5-bar return (fast proxy) */
|
||||
float close_now = (targets[bar * 4 + 3]);
|
||||
float close_5 = (targets[(bar - 5) * 4 + 3]);
|
||||
float close_20 = (targets[(bar - 20) * 4 + 3]);
|
||||
/* 5-bar return (fast proxy).
|
||||
* Targets layout per `crates/ml/src/fxcache.rs`
|
||||
* (TARGET_DIM=6): col 0=preproc_close, 1=preproc_next, 2=raw_close,
|
||||
* 3=raw_next (FUTURE), 4=raw_open, 5=mid_open. Read raw_close (col 2)
|
||||
* for momentum — pre-`063fd2716` this was `bar*4+3` (= col 3 = raw_next
|
||||
* in stride-4 = same value as raw_close[bar+1]); the stride bump made
|
||||
* that read slide into the wrong bar. Switch to raw_close at col 2,
|
||||
* stride 6. */
|
||||
float close_now = (targets[bar * 6 + 2]);
|
||||
float close_5 = (targets[(bar - 5) * 6 + 2]);
|
||||
float close_20 = (targets[(bar - 20) * 6 + 2]);
|
||||
float ret_5 = (close_now - close_5) / (close_5 + 1e-8f);
|
||||
float ret_20 = (close_now - close_20) / (close_20 + 1e-8f);
|
||||
|
||||
@@ -3971,8 +3987,13 @@ extern "C" __global__ void compute_difficulty_scores(
|
||||
if (i >= total_bars) return;
|
||||
if (i >= total_bars - 5) { scores[i] = 1e6f; return; }
|
||||
|
||||
float close_0 = (targets[i * 4 + 2]);
|
||||
float close_5 = (targets[(i + 5) * 4 + 2]);
|
||||
/* Per `crates/ml/src/fxcache.rs` (TARGET_DIM=6):
|
||||
* col 2 = raw_close. Pre-`063fd2716` stride was 4 with raw_close still
|
||||
* at col 2 — the `+ 2` column index is unchanged; only the stride bumps
|
||||
* 4 → 6. Reading at the old stride against the new layout slid every
|
||||
* lookup into the wrong bar's data. */
|
||||
float close_0 = (targets[i * 6 + 2]);
|
||||
float close_5 = (targets[(i + 5) * 6 + 2]);
|
||||
float ret_5 = fabsf(close_5 - close_0) / fmaxf(close_0, 1.0f);
|
||||
scores[i] = 1.0f / (ret_5 + 0.001f);
|
||||
}
|
||||
@@ -4012,12 +4033,14 @@ extern "C" __global__ void hindsight_relabel_kernel(
|
||||
float pos = compute_target_position_4branch(dir_idx_hr, mag_idx_hr, max_position);
|
||||
if (fabsf(pos) < 0.001f) return;
|
||||
|
||||
float entry_price = (targets[bar * 4 + 2]);
|
||||
/* Per `crates/ml/src/fxcache.rs` (TARGET_DIM=6):
|
||||
* col 2 = raw_close. Stride 4 → 6 fix; column index unchanged. */
|
||||
float entry_price = (targets[bar * 6 + 2]);
|
||||
if (entry_price < 1.0f) return;
|
||||
|
||||
float best_pnl = 0.0f;
|
||||
for (int k = 1; k <= lookahead && (bar + k) < total_bars; k++) {
|
||||
float future_price = (targets[(bar + k) * 4 + 2]);
|
||||
float future_price = (targets[(bar + k) * 6 + 2]);
|
||||
float pnl = pos * (future_price - entry_price);
|
||||
if (pnl > best_pnl) best_pnl = pnl;
|
||||
}
|
||||
|
||||
@@ -96,8 +96,37 @@ pub const FEATURE_SCHEMA_HASH: u64 = {
|
||||
/// Feature vector dimensionality.
|
||||
const FEAT_DIM: usize = 42;
|
||||
|
||||
/// Target vector dimensionality (close, next_close, raw_close, raw_next, raw_open, mid_price_open).
|
||||
const TARGET_DIM: usize = 6;
|
||||
/// Target vector dimensionality. Mirrors the named column constants below.
|
||||
///
|
||||
/// Public so consumer kernels and host-side readers can index by name rather
|
||||
/// than literal stride. Two latent bugs converged in the cancelled 50-epoch
|
||||
/// run `train-multi-seed-p5qzw` because consumers hardcoded stride 4 (left
|
||||
/// over from pre-`063fd2716`) and column 0 (pre-`5a5dd0fed` raw_close, post-
|
||||
/// fix preproc_close). Centralizing the layout here so future renames force
|
||||
/// every consumer to update at the same call site.
|
||||
pub const TARGET_DIM: usize = 6;
|
||||
|
||||
/// Column 0: `preproc_close` — log-return-normalized close (network input).
|
||||
pub const TARGET_PREPROC_CLOSE: usize = 0;
|
||||
|
||||
/// Column 1: `preproc_next` — log-return-normalized next-bar close.
|
||||
pub const TARGET_PREPROC_NEXT: usize = 1;
|
||||
|
||||
/// Column 2: `raw_close` — raw close price for portfolio simulation P&L + tx.
|
||||
pub const TARGET_RAW_CLOSE: usize = 2;
|
||||
|
||||
/// Column 3: `raw_next` — raw next-bar close. **FUTURE information** — must
|
||||
/// not be used for any reward, P&L, equity, or portfolio computation. Kept
|
||||
/// in the layout for legacy back-compat; consumers should `(void)` this slot
|
||||
/// to suppress unused-element warnings.
|
||||
pub const TARGET_RAW_NEXT: usize = 3;
|
||||
|
||||
/// Column 4: `raw_open` — raw open price (OHLCV).
|
||||
pub const TARGET_RAW_OPEN: usize = 4;
|
||||
|
||||
/// Column 5: `mid_price_open` — MBP-10 midpoint at bar open; falls back to
|
||||
/// `raw_open` when no MBP-10 snapshot is available at the bar boundary.
|
||||
pub const TARGET_MID_OPEN: usize = 5;
|
||||
|
||||
/// OFI vector dimensionality — mirrors `ml_core::state_layout::OFI_DIM`
|
||||
/// (20 legacy slots + 12 new microstructure slots).
|
||||
@@ -499,6 +528,29 @@ fn read_body_f32(
|
||||
|
||||
// ── Finder ───────────────────────────────────────────────────────────────────
|
||||
|
||||
#[cfg(test)]
|
||||
mod target_layout_tests {
|
||||
use super::*;
|
||||
|
||||
/// Compile-time guard: the named target columns are dense and exhaust
|
||||
/// the layout. If a future patch reorders or removes a slot, this fires.
|
||||
#[test]
|
||||
fn target_columns_dense_and_exhaustive() {
|
||||
let cols = [
|
||||
TARGET_PREPROC_CLOSE,
|
||||
TARGET_PREPROC_NEXT,
|
||||
TARGET_RAW_CLOSE,
|
||||
TARGET_RAW_NEXT,
|
||||
TARGET_RAW_OPEN,
|
||||
TARGET_MID_OPEN,
|
||||
];
|
||||
for (i, &c) in cols.iter().enumerate() {
|
||||
assert_eq!(c, i, "Column constant out of order at index {i}");
|
||||
}
|
||||
assert_eq!(cols.len(), TARGET_DIM);
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod has_ofi_tests {
|
||||
use super::*;
|
||||
|
||||
@@ -22,6 +22,7 @@ use common::metrics::{questdb_sink, training_metrics};
|
||||
use tracing::{debug, info, warn};
|
||||
|
||||
use crate::cuda_pipeline::mapped_pinned::MappedI32Buffer;
|
||||
use crate::fxcache::TARGET_RAW_CLOSE;
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::{
|
||||
EPOCH_IDX_INDEX, EPSILON_EFF_INDEX, GAMMA_DIR_EFF_INDEX, Q_ABS_REF_INDEX,
|
||||
Q_DIR_ABS_REF_INDEX, TLOB_REGIME_FOCUS_EMA_INDEX,
|
||||
@@ -567,8 +568,12 @@ impl DQNTrainer {
|
||||
let mut mean: f64 = 0.0;
|
||||
let mut m2: f64 = 0.0;
|
||||
for w in training_data.windows(2) {
|
||||
let prev_close = w[0].1[0];
|
||||
let curr_close = w[1].1[0];
|
||||
// Bug B fix (2026-05-02): post-`5a5dd0fed` `targets[0]` is
|
||||
// log-return-normalized `preproc_close`, not raw_close.
|
||||
// Welford needs raw price ratios — read `TARGET_RAW_CLOSE`
|
||||
// (column 2 per `crate::fxcache::TARGET_*` constants).
|
||||
let prev_close = w[0].1[TARGET_RAW_CLOSE];
|
||||
let curr_close = w[1].1[TARGET_RAW_CLOSE];
|
||||
if prev_close > 0.0 {
|
||||
let r = (curr_close / prev_close).ln();
|
||||
if r.is_finite() {
|
||||
@@ -600,7 +605,7 @@ impl DQNTrainer {
|
||||
tracing::warn!(
|
||||
"epoch_vol_normalizer raw value {:.3e} outside sanity band [{:.0e}, {:.0e}] \
|
||||
(n_returns={}, mean={:.3e}); using default {:.3e}. \
|
||||
Likely cause: targets[0] not raw_close on this dataset.",
|
||||
Likely cause: targets[TARGET_RAW_CLOSE=2] not raw_close on this dataset.",
|
||||
raw_vol, VOL_LOW, VOL_HIGH, count, mean, VOL_DEFAULT
|
||||
);
|
||||
VOL_DEFAULT
|
||||
|
||||
@@ -3661,3 +3661,75 @@ Each of those independently injects per-branch variability into the magnitude he
|
||||
**Follow-up trigger:** if a future smoke / training run shows `intent_dist` re-freezing at exactly `0.25/0.50/0.25` for ≥3 consecutive epochs, re-open the investigation as a separate spec (`2026-XX-XX-action-select-binary-decomposition.md` per the SP5 plan §C4 escalation path) and search for whichever shared-infra site SP5 missed.
|
||||
|
||||
Refs: SP5 plan §C Task C4 (line 1256); `smoke-test-ks2wf` (`5845e4403`), `smoke-test-7pv9v` (`f42b5fff8`), `smoke-test-w9nsw` (`2e9e276a0`); pearls landed in SP5 §C5 (`pearl_per_branch_c51_atom_span`, `pearl_per_branch_loss_budget`, `pearl_per_branch_noisy_sigma`, `pearl_per_group_adam_hyperparams`, `pearl_per_branch_iqn_tau_schedule`, `pearl_kelly_cap_signal_driven_floors`, `pearl_trail_stop_signal_driven`); new memory pearl `pearl_intent_dist_freeze_resolved.md` records the implicit lesson. Closes plan task #296.
|
||||
|
||||
### Fix 27: target stride/column drift cleanup (host + GPU consumers, 2026-05-02)
|
||||
|
||||
Two latent bugs converged in the cancelled 50-epoch run `train-multi-seed-p5qzw` at HEAD `96769d171` (label_scale=808 vs smoke baseline 22-28, ~30× over expected magnitude). Root-cause and atomic fix:
|
||||
|
||||
**Bug A — fxcache target stride 4 → 6 not propagated to consumer kernels** (latent since `063fd2716`, 2026-04-19). Commit `063fd2716` bumped `TARGET_DIM` from 4 to 6 and added two columns (`raw_open` at col 4, `mid_price_open` at col 5) for the SP5 MFT mid-price mark-to-market signal. The fxcache writer + the production hot-path consumer (`experience_env_step` via `tgt = targets + bar_idx * 6`) were updated in the same commit, but five other consumer kernels still hardcoded `targets[bar * 4 + col]`. Reading at the old stride against the stride-6 buffer slid every lookup into the wrong bar's data — for the curriculum difficulty kernel (368k bars), reading `targets[i*4+2]` retrieved bytes from somewhere between bar `i*4/6` and `i*4/6 + 1`, producing nonsense `ret_5` values.
|
||||
|
||||
**Bug B — `targets[0]` semantic flip not propagated to host consumer** (introduced today at `5a5dd0fed`, 2026-05-02). The Bug 1 fix changed `target[0]` from raw_close to log-return-normalized `preproc_close`. The kernel comment at `experience_kernels.cu:1556` and the host-side docstring at `cuda_pipeline/mod.rs:508` were updated. The writers (`precompute_features.rs`, `data_loading.rs`) were updated. But one host-side consumer in `training_loop.rs::epoch_vol_normalizer` (Welford streaming variance over per-bar log-returns) still read `w[0].1[0]` and `w[1].1[0]` expecting raw_close. Post-fix that returns log-returns ≈ ±0.001; the Welford then takes `ln(curr/prev)` which produces `ln(small)/ln(small) → 0.6 ≈ stddev` — way outside the sanity band `[1e-5, 1e-1]`. The sanity-band code triggered `VOL_DEFAULT = 5e-4`, which is the only reason this didn't NaN out — but downstream consumers fed off this "default" without checking, and the compounded label_scale=808 surfaced in the 50-epoch run.
|
||||
|
||||
**Why the smokes missed it.** The local `multi_fold_convergence::test_multi_fold_convergence` smoke harness exercises only a 5-epoch path with `expert_action_override` rate gated to 0 (no expert override) and `compute_difficulty_scores` only invoked at curriculum init (which is itself gated off in smoke configs). The `hindsight_relabel_kernel` is only invoked when `hindsight_fraction > 0` — also gated off by default in smoke. The `portfolio_sim_kernel` (legacy `GpuPortfolioSimulator`) has no live caller. The Welford normalizer is invoked once per epoch over the full training_data, so it ran at smoke scale (1024 bars) but the sanity-band fallback hid the bug — and the final value used (`5e-4` default) coincidentally matched the equity-index 1-min vol baseline. Production `train-multi-seed-p5qzw` with 368k bars across 50 epochs amplified the compounded effect; once vol_normalizer fed into `epoch_label_scale = mean_abs_label / vol_normalizer` it inverted to ~30× the expected scale, and the cluster killed the run on the Q-drift kill criterion at epoch 7.
|
||||
|
||||
**Sites fixed**:
|
||||
- HOST (1 site, training_loop.rs):
|
||||
- Lines 570-573 (`epoch_vol_normalizer` Welford): `w[0].1[0]` / `w[1].1[0]` → `w[0].1[TARGET_RAW_CLOSE]` / `w[1].1[TARGET_RAW_CLOSE]` (column 2 per the contract). Imports `TARGET_RAW_CLOSE` from `crate::fxcache` (the contract-owner module).
|
||||
- Line 603 (warning message): "Likely cause: targets[0] not raw_close" → "targets[TARGET_RAW_CLOSE=2] not raw_close" so future operators read the correct column when this fires.
|
||||
- GPU (5 sites in experience_kernels.cu):
|
||||
- Line 3768 (`expert_action_override` parameter doc): `[total_bars, 4] OHLCV` → `[total_bars, TARGET_DIM=6] fxcache targets`.
|
||||
- Lines 3806-3808 (`expert_action_override` 5-bar momentum): `targets[bar*4+3]` → `targets[bar*6+2]` (raw_close, col 2). Old stride-4 layout had `[3] = raw_next ≈ close[bar+1]` which served as a momentum proxy; the stride bump broke that proxy and reading raw_close directly is the cleaner semantic anyway.
|
||||
- Lines 3974-3975 (`compute_difficulty_scores` directional clarity): `targets[i*4+2]` / `targets[(i+5)*4+2]` → `targets[i*6+2]` / `targets[(i+5)*6+2]`. Column index unchanged (raw_close at col 2 in both old + new layouts); only stride bumps.
|
||||
- Line 4015 (`hindsight_relabel_kernel` entry price): `targets[bar*4+2]` → `targets[bar*6+2]`. Stride only.
|
||||
- Line 4020 (`hindsight_relabel_kernel` future price): `targets[(bar+k)*4+2]` → `targets[(bar+k)*6+2]`. Stride only.
|
||||
- Lines 3400-3404 (`portfolio_sim_kernel`): `t_offset = global_idx * 4` → `t_offset = global_idx * 6`. Column indices unchanged for cols 0-3 (preproc_close, preproc_next, raw_close, raw_next); kernel is currently orphaned (no live `GpuPortfolioSimulator` caller in production path) but kept consistent per `feedback_no_partial_refactor` — every consumer of the same buffer contract migrates in the same commit.
|
||||
- GPU (1 site in dt_kernels.cu):
|
||||
- Line 793 (`dt_compute_rewards_actions_kernel` parameter doc): `[num_bars, 4] — O,H,L,C per bar` → `[num_bars, TARGET_DIM=6] — fxcache targets`. The "OHLC" docstring was historically incorrect — fxcache never had OHLC layout, only the close-row variants.
|
||||
- Lines 803, 807: `targets[i*4+3]` / `targets[(i+1)*4+3]` → `targets[i*6+2]` / `targets[(i+1)*6+2]`. Pre-bump this kernel read `raw_next` (col 3 in stride-4); post-bump it now reads `raw_close` (col 2 in stride-6) which is the correct semantic for "close at bar t".
|
||||
- HOST docstring (1 site, decision_transformer.rs:1049): `[num_bars, 4]` → `[num_bars, TARGET_DIM=6]` with reference to the new `target_layout` module.
|
||||
- HOST docstring (1 site, gpu_walk_forward.rs:12): VRAM-layout ASCII art `targets [total_bars, 4]` → `targets [total_bars, 6]`.
|
||||
|
||||
**PPO consumer NOT fixed** (intentional, deferred): `ppo_experience_kernel.cu:570-574, 805-806` and `gpu_ppo_collector.rs:488` use stride 4 — but PPO has its own `set_raw_market_data` path (`ppo.rs:382-417`) that uploads only 4 columns from `Vec<f64>` targets. The PPO write/read pair is internally consistent at stride 4 and is unaffected by the fxcache stride bump. Production `train_baseline_rl.rs::train_from_slices` flow does NOT invoke `set_raw_market_data`, so the PPO GPU collector is currently orphaned (returns error at `ppo.rs:592` if reached). Consolidating PPO onto the same fxcache target buffer is a separate refactor.
|
||||
|
||||
**Structural prevention** — named column constants in `crates/ml/src/fxcache.rs` (the contract-owner module that already houses `TARGET_DIM`):
|
||||
- Defines `TARGET_PREPROC_CLOSE=0`, `TARGET_PREPROC_NEXT=1`, `TARGET_RAW_CLOSE=2`, `TARGET_RAW_NEXT=3`, `TARGET_RAW_OPEN=4`, `TARGET_MID_OPEN=5`. `TARGET_DIM=6` was already there but was promoted from private to `pub` so consumers can import it.
|
||||
- Consumers import via `use crate::fxcache::TARGET_RAW_CLOSE;` — co-locating the constants with the writer means future renames or column adds force every consumer to update at the same call site (the contract owner is the writer; if you change the layout you must change the writer; if you change the writer you change the constant; if you change the constant the call site moves with it).
|
||||
- Compile-time test `fxcache::target_layout_tests::target_columns_dense_and_exhaustive` asserts dense layout — if a future patch reorders or removes a column, the test fails.
|
||||
- GPU kernels reference the constant module path in comments (kernels can't import Rust constants); future stride/column changes are visible in every call site because the comment block and the literal `* 6` / `+ 2` sit adjacent.
|
||||
|
||||
**Exhaustive audit results** (full grep sweep before committing):
|
||||
|
||||
```bash
|
||||
# Hardcoded stride-4 in targets reads (production, post-fix):
|
||||
$ grep -rnE 'targets\[[a-zA-Z_][a-zA-Z_0-9]*[[:space:]]*\*[[:space:]]*4|targets\[\([^)]+\)[[:space:]]*\*[[:space:]]*4' \
|
||||
crates/ml/src/cuda_pipeline/ crates/ml/src/trainers/ | grep -v test
|
||||
# (zero matches in production code; only doc-comment hits in audit doc describe the bug)
|
||||
|
||||
# Host-side .1[0/1] tuple reads against (features, targets) pairs:
|
||||
$ grep -rnE '\.1\[[0-1]\]' crates/ml/src/ crates/ml/examples/
|
||||
# (zero matches post-fix)
|
||||
|
||||
# Stride-4 multiplications in target context (caught the t_offset / t_off variants):
|
||||
$ grep -rnE 'global_idx\s*\*\s*4|global_bar\s*\*\s*4|t_offset\s*=.*\*\s*4|t_off\s*=.*\*\s*4' \
|
||||
crates/ml/src/cuda_pipeline/experience_kernels.cu crates/ml/src/cuda_pipeline/dt_kernels.cu | grep -v "//"
|
||||
# (zero matches post-fix in DQN/DT kernels; PPO `t_off = global_bar * 4` retained as it's
|
||||
# internally consistent with PPO's own 4-col upload — out-of-scope per above)
|
||||
```
|
||||
|
||||
**Verification gates**:
|
||||
- `SQLX_OFFLINE=true cargo check -p ml --offline` clean.
|
||||
- `SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda` clean (cubin compiles via nvcc; release build under 2 min).
|
||||
- `cargo test -p ml --lib --offline -- target_layout` 1/1 pass (the new contiguity test).
|
||||
- Existing smokes (`smoke-test-w9nsw` baseline) re-run will pick up the fix; cancelled `train-multi-seed-p5qzw` re-dispatch is the gating production validation.
|
||||
|
||||
**What this change touches**:
|
||||
- `crates/ml/src/fxcache.rs` — adds 6 named column constants (`TARGET_PREPROC_CLOSE` etc.) co-located with the existing `TARGET_DIM`; promotes `TARGET_DIM` from private to `pub`; new compile-time density test.
|
||||
- `crates/ml/src/cuda_pipeline/mod.rs` — wires the new module in.
|
||||
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — Bug B host fix + warning-message update.
|
||||
- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 5 GPU sites + kernel header doc.
|
||||
- `crates/ml/src/cuda_pipeline/dt_kernels.cu` — 1 GPU site (2 reads) + kernel header doc.
|
||||
- `crates/ml/src/cuda_pipeline/decision_transformer.rs` — host docstring contract update.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_walk_forward.rs` — host docstring (ASCII layout art) update.
|
||||
- `docs/dqn-wire-up-audit.md` — this entry.
|
||||
|
||||
Refs: cancelled `train-multi-seed-p5qzw` at HEAD `96769d171`; bug origins `063fd2716` (target_dim 4→6 bump) and `5a5dd0fed` (Bug 1 column-0 semantic fix). `feedback_no_partial_refactor` (every consumer of the shared buffer contract migrates in the same commit), `feedback_trust_code_not_docs` (kernel doc comments saying `[..., 4] OHLCV` were stale post-bump and misled future readers), `feedback_no_quickfixes` (proper fix with shared-constants prevention, not just a one-line patch), `feedback_wire_everything_up` (the column constants land co-located with the existing `fxcache::TARGET_DIM` so callers and writer share a single contract module).
|
||||
|
||||
Reference in New Issue
Block a user