diff --git a/crates/ml/src/cuda_pipeline/decision_transformer.rs b/crates/ml/src/cuda_pipeline/decision_transformer.rs index e9711ebac..9229db795 100644 --- a/crates/ml/src/cuda_pipeline/decision_transformer.rs +++ b/crates/ml/src/cuda_pipeline/decision_transformer.rs @@ -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 /// diff --git a/crates/ml/src/cuda_pipeline/dt_kernels.cu b/crates/ml/src/cuda_pipeline/dt_kernels.cu index 0aa6cb71a..cd8f6cf9e 100644 --- a/crates/ml/src/cuda_pipeline/dt_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dt_kernels.cu @@ -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 { diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index df864cd3c..085a5043c 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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; } diff --git a/crates/ml/src/fxcache.rs b/crates/ml/src/fxcache.rs index bd71e3a14..eb4ebb8ab 100644 --- a/crates/ml/src/fxcache.rs +++ b/crates/ml/src/fxcache.rs @@ -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::*; diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index a1e0bf883..0072888b4 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 027fffd14..8c636e7ce 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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` 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).