diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index c9cd3016f..54ed7247f 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -1034,66 +1034,32 @@ extern "C" __global__ void cross_fold_consistency( } } -/* ══════════════════════════════════════════════════════════════════════ - * PHANTOM LIQUIDITY KERNEL (#16) +/* phantom_liquidity_gbm DELETED (Fix 30 Ambiguous-A, 2026-05-02). * - * Generates synthetic price data via Geometric Brownian Motion (GBM). - * Replaces a fraction of episodes with GBM prices that match the - * real data's volatility and drift. + * The synthetic-data overlay kernel was dead code — zero callers in the + * production crate, in `services/`, and in `crates/`. The audit grep + * (`grep -rn "phantom_liquidity_gbm" crates/ml/src/ services/ crates/`) + * returned only the kernel definition itself; no `load_function`, no + * `launch_builder`, no Rust-side launcher. * - * For each phantom episode: S(t+1) = S(t) * exp((mu - sigma²/2)*dt + sigma*sqrt(dt)*Z) - * where Z ~ N(0,1) from LCG RNG. + * Per Fix 29 row #18 (`dqn_utility_kernels.cu:1089-1093`), the writer's + * downstream consumer contract was unverifiable because no consumer + * existed. The kernel wrote `log_ret` (synthetic GBM log-return) to + * `market_features[bar*market_dim + 0]` AND `+3`. Post Bug 1 (commit + * `5a5dd0fed`) the writer for `market_features[..+0]` is z-normed + * log-return, not raw log-return — a hypothetical re-wire of this + * kernel into the production path would have introduced exactly the + * Bug-1 contract drift the audit was triaging. * - * Overwrites market_features for selected episodes in-place. - * Grid: ceil(N * L / 256), Block: 256. - * ══════════════════════════════════════════════════════════════════════ */ -extern "C" __global__ void phantom_liquidity_gbm( - float* __restrict__ market_features, /* [total_bars, market_dim] — MODIFIED */ - const int* __restrict__ episode_starts, /* [N] */ - unsigned int* __restrict__ rng_states, /* [N] */ - float mu, /* drift per bar */ - float sigma, /* volatility per bar */ - float initial_price, /* starting price for GBM */ - int N, - int L, /* timesteps per episode */ - int market_dim, - int total_bars, - int phantom_mod /* episodes where i % mod == 0 get synthetic data */ -) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx >= N * L) return; - - int episode = idx / L; - int t = idx % L; - - /* Only phantom episodes */ - if (episode % phantom_mod != 0) return; - - int bar_idx = episode_starts[episode] + t; - if (bar_idx >= total_bars) return; - - /* GBM: generate synthetic close price */ - unsigned int rng = rng_states[episode] + (unsigned int)t * 1013904223u; - rng = rng * 1664525u + 1013904223u; - float u1 = (float)(rng & 0xFFFF) / 65536.0f + 1e-6f; - rng = rng * 1664525u + 1013904223u; - float u2 = (float)(rng & 0xFFFF) / 65536.0f; - float Z = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); - - /* Cumulative GBM from initial price */ - float log_return = (mu - 0.5f * sigma * sigma) + sigma * Z; - /* Approximate: use running product via t steps */ - float cumul_log = log_return * (float)(t + 1); - float synthetic_price = initial_price * expf(cumul_log); - - /* Write synthetic log-return to feature[0] (close log return) */ - long long offset = (long long)bar_idx * market_dim; - float prev_price = initial_price * expf(log_return * (float)t); - float log_ret = (prev_price > 0.0f) ? logf(synthetic_price / prev_price) : 0.0f; - market_features[offset + 0] = (log_ret); - /* Write synthetic prices to targets-style features */ - market_features[offset + 3] = (log_ret); /* close log return copy */ -} + * Resolution per `feedback_no_hiding` + `feedback_wire_everything_up`: + * delete the kernel rather than leave it as a contract-ambiguous orphan + * that future re-wirers might trip over. Re-introducing GBM-overlay + * augmentation is a separate design with its own writer/consumer + * contract decision (raw vs z-normed log-return; column 0 vs column 3 + * post the `data_loading.rs` Bug-1 normalisation pivot) and should be + * landed with caller wiring in the same commit per + * `feedback_wire_everything_up`. + */ /* ══════════════════════════════════════════════════════════════════════ * PRICE-LEVEL INVARIANCE KERNEL (#12) diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c44b3937a..d14d684c3 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3816,6 +3816,7 @@ Closes the four ⚠ Stale and one ❓ Ambiguous sites Fix 29 enumerated. Each si | 14 | `crates/ml/src/trainers/dqn/trainer/metrics.rs:576` (`val_data → window_prices`) | ⚠ Stale — reads `target[0]` (preproc_close, z-normed log-return) as a dollar close price | Switched to `target[TARGET_RAW_CLOSE]` (col 2). Dead `fv[3]` fallback deleted (`set_val_data_from_slices` always yields `target.len() == 6` post `TARGET_DIM=6` bump in `063fd2716`, so the fallback is unreachable; per `feedback_no_hiding` deleted rather than left as a silent wrong-units fallback). Imports `TARGET_RAW_CLOSE` from `crate::fxcache`. | ✅ | | 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`. | ✅ | **Stale-A commit** (rows #14, #15): - `crates/ml/src/trainers/dqn/trainer/metrics.rs` — column-index swap + import + dead-fallback deletion + comment. @@ -3830,3 +3831,11 @@ References: clones the host-side variant of Fix 27 Bug B exactly. `feedback_no_h - `docs/dqn-wire-up-audit.md` — Stale-C row appended. References: Stale-C restores seed-phase scripted policy quality (post-seed all branches use the network's Q-spread, so the affected blast radius is bounded to the warm-start window — `replay_seed_steps` = 100k production / 1k smoke). `feedback_no_partial_refactor` (single launcher migrates with the kernel signature change in one commit), `feedback_no_functionality_removal` (recent_ret signal stays — only its data source is fixed; CUSUM substitution explicitly rejected per the audit's contract-decision branch), `feedback_no_hiding` (no fallback to z-normed-log-return reads remaining; the kernel either gets real raw_close or clamps to total_bars-1 boundary), `feedback_no_cpu_compute_strict` n/a (zero new host-side compute; targets buffer already exists at line 3042), `feedback_no_htod_htoh_only_mapped_pinned` already satisfied (`targets_buf` is mapped pinned by upstream caller). + +**Ambiguous-A commit** (row #18): +- `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` — `phantom_liquidity_gbm` kernel definition (47 lines including its `══` comment header) deleted; replaced with a 23-line `/* ... DELETED ... */` block citing the audit grep result, the Bug-1 contract trap, and the re-introduction conditions per `feedback_wire_everything_up`. +- `docs/dqn-wire-up-audit.md` — Ambiguous-A row appended with the dead-code-deleted resolution. + +References: this is the dead-code branch of the contract decision. The kernel had been live in the codebase since at least early 2026 with no consumer ever materialising. `feedback_no_hiding` (orphan kernels with ambiguous contracts get deleted, not left for a future contributor to re-wire wrong), `feedback_wire_everything_up` (an orphan kernel that compiles but is unconsumed is exactly the failure mode this rule prevents — wire it in the same commit or delete it). `feedback_no_functionality_removal` does NOT apply: a kernel with zero callers is not a functional feature, and the audit doc retains the design intent (synthetic-data augmentation via GBM) so any future re-introduction has the spec available. + +**Stale-B (row #13) — DEFERRED**: pre-existing `clone_to_device_*_via_pinned` guard violations in `gpu_backtest_evaluator.rs` (5 sites at lines 641-650, 1468) block any commit that stages this file. The DtoD-via-pinned guard landed today (commit `5275932f4`) and the file has not been migrated since the guard landed; this would be the first commit to do so. Migration scope is 5 buffer allocations + 3 launch-site updates (CudaSlice → MappedF32Buffer + `.dev_ptr` indirections) covering buffers (`prices_buf`, `features_buf`, `window_lens_buf`, `portfolio_buf`, plus the `reset_evaluation_state` portfolio re-init) that are orthogonal to Bug-1 contract drift. Per `feedback_no_partial_refactor` the right unit-of-work is a single dedicated commit migrating all 5 sites at once, not folding the migration under the Stale-B header. The Stale-B kernel-side fix (route `prices_buf` to `backtest_plan_state_isv` and read `prices[base+3]` instead of `features[base+0]`) is stashed at `stash@{0}` (`stale-b-wip`) for re-application after the `_via_pinned` migration lands. 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).