From 7a3d88646221758c6971020d85b9509629e49898 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 26 Apr 2026 22:27:49 +0200 Subject: [PATCH] fix(dqn): three concerns from kelly-fix-multifold val run MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CONCERN 1 — train Return display (financials.rs) Replace `mean_per_bar × bars_per_year` ("annualized arithmetic return") with the actual cumulative compounded return: `prod(1+r) − 1` via log-space sum for f64 stability. Step floor at 1+r >= 1e-10 so a -100% bar gives -23 log contribution rather than -inf; result bounded above -100%. The old formula multiplied per-bar mean by ~98K, producing values like Return=-51234% that LOOKED like portfolio collapse but were actually just label artefact: -52 bps mean × 98,280 bars/year = -51,234%, while the actual compounded return over the rollout was bounded and finite. "Total Return" now means what a trader expects. CONCERN 2 — val_Sharpe annualization source-of-truth (training_loop.rs + gpu_backtest_evaluator.rs) val_Sharpe was using `m.sharpe` from the kernel (annualised by `sqrt(bars_per_day × trading_days_per_year)` at evaluator init), but val_Sharpe_raw recomputed the divisor from `self.hyperparams.bars_per_day * 252.0` in Rust. Empirically the two diverged by 5.4× (val_Sharpe=143.38 / val_Sharpe_raw=0.0841 = 1704, expected sqrt(98280) = 313.5). Root cause unclear without runtime instrumentation but the structural fix is to read the EXACT factor the kernel applied: new `GpuBacktestEvaluator::annualization_factor()` returns the f32 stored at init time. Rust now divides by that, so the two numbers can never drift regardless of how `bars_per_day` flows through the config layers. Per `feedback_no_partial_refactor.md` — every consumer of a shared contract must use the same source. CONCERN 3 — C51 expected-Q Hold/Flat bias (experience_kernels.cu) Direction-branch Boltzmann tau floor: replace static `0.01` with ISV-adaptive `max(ISV[21], 0.01)` where ISV[21] is q_dir_abs_ref (EMA of mean |Q| across direction bins, same signal that drives conviction). The C51 expected-Q biases Flat above directional actions when the policy has no measurable edge — Flat returns are exactly 0 (no position change), collapsing C51's distribution to a delta at 0; directional returns concentrate slightly below 0 from tx_cost without compensating edge, giving E[Q_directional] < E[Q_flat] = 0. With Q-spread on the order of 0.01-0.1 and the static 0.01 tau floor, exp(Q/tau) ratios approached deterministic argmax, locking the policy into Hold/Flat within ~1 epoch and preventing the exploration needed to discover real edge. Validated empirically by the kelly-fix-multifold run (train-multi-seed-bs9m5): val_dir_dist epoch 1 [S=.23 H=.17 L=.42 F=.18] (active 65%) val_dir_dist epoch 2 [S=.14 H=.35 L=.15 F=.37] (active 29%) That 35→72% Hold+Flat shift in one epoch is the C51 attractor in action. The ISV-adaptive floor preserves relative spread for sampling: when Q-magnitudes grow during training, the floor scales with them, so the policy never collapses to deterministic argmax until q_range exceeds the network's typical Q magnitude — i.e., until spread represents real, scale-significant edge. Coherent with the existing conviction formula (same ISV reference). No tuned constants per `feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`. Cold-start fallback retains 0.01 minimum so kernel doesn't divide by zero pre-first-update. Once ISV producer fires (every epoch), adaptive floor takes over. 7 financials unit tests pass. Workspace cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 56 ++++++++++++++----- .../cuda_pipeline/gpu_backtest_evaluator.rs | 13 +++++ crates/ml/src/trainers/dqn/financials.rs | 40 ++++++++++--- .../src/trainers/dqn/trainer/training_loop.rs | 14 ++++- docs/dqn-wire-up-audit.md | 1 + 5 files changed, 102 insertions(+), 22 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 60b1ca9d7..cc944ade1 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -841,20 +841,41 @@ extern "C" __global__ void experience_action_select( (void)min_hold_bars; (void)portfolio_states; - /* Branch 0: direction — Boltzmann softmax with adaptive tau (q_range, - * floor 0.01) for both training and eval. + /* Branch 0: direction — Boltzmann softmax with ISV-adaptive tau floor. * - * State-adaptive without tuned constants: large per-sample q_range - * (confident state) gives large tau → softmax picks the best direction - * near-deterministically; small q_range (ambiguous state) sits at the - * 0.01 floor → softmax samples uniformly across directions. The Philox - * stream is seeded by (i, timestep) so eval is bit-reproducible across - * runs at the same checkpoint. + * tau = max(q_range, q_dir_abs_ref) where q_dir_abs_ref is ISV[21], the + * EMA of mean |Q| across direction bins. This makes the tau floor scale + * with the network's actual Q magnitude rather than relying on a tuned + * 0.01 constant. * - * C51's expected Q structurally favors Flat (zero PnL variance → tightest - * distribution → highest expected Q). Boltzmann samples PROPORTIONALLY to - * exp(Q/tau) so Short/Long retain probability mass even when their Q is - * mildly worse, preserving the diversity training relies on. */ + * Why the old `max(q_range, 0.01)` was dangerous: C51's expected Q + * structurally biases Flat above directional actions when the policy has + * no measurable edge — Flat returns are exactly 0 (no position change), + * collapsing C51's distribution to a delta at 0; directional returns + * concentrate slightly below 0 from tx_cost without compensating edge, + * giving E[Q_directional] < E[Q_flat] = 0. With Q-spread on the order + * of 0.01-0.1 (typical at fresh init or post-collapse), the static 0.01 + * floor produced exp(Q/tau) ratios approaching deterministic argmax, + * locking the policy into Hold/Flat within ~1 epoch and preventing the + * exploration needed to discover real edge. Validated empirically: the + * Kelly val-Flat-collapse fix (commit 0c9d1ee39) exposed val_dir_dist + * shifting from S=.23/H=.17/L=.42/F=.18 (epoch 1) to S=.14/H=.35/L=.15/ + * F=.37 (epoch 2) — a 35→72% Hold+Flat shift in one epoch. + * + * With the ISV-driven floor: when Q-magnitudes grow during training, + * the floor scales with them, preserving the SAME relative spread for + * sampling. The policy never collapses to deterministic argmax until + * q_range exceeds the network's typical Q magnitude — i.e., until the + * spread represents real, scale-significant edge over the baseline. + * Both conviction (computed below) and tau use the SAME ISV reference, + * keeping the two adaptive mechanisms coherent. + * + * Cold-start fallback: when ISV[21] is < 1e-6 (pre-first-update), use + * the legacy 0.01 floor so the kernel doesn't divide by zero. Once the + * ISV producer fires (every epoch), the adaptive floor takes over. + * + * Philox stream is seeded by (i, timestep) so eval is bit-reproducible + * across runs at the same checkpoint. */ if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_dir) { int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b0_size); dir_idx = (r >= b0_size) ? b0_size - 1 : r; @@ -868,7 +889,16 @@ extern "C" __global__ void experience_action_select( q_min_d = fminf(q_min_d, qv); } { - float tau_d = fmaxf(q_max_d - q_min_d, 0.01f); + /* ISV-adaptive tau floor: same q_dir_abs_ref signal that drives + * the conviction-based Kelly warmup floor (kelly_position_cap). + * Scales the exploration temperature with the network's actual + * Q magnitude so softmax never becomes deterministic for tiny + * absolute Q-spreads — preventing the C51 Hold/Flat attractor. */ + float q_dir_abs_ref = (isv_signals_ptr != NULL) + ? isv_signals_ptr[21] + : 0.0f; + float tau_floor = fmaxf(q_dir_abs_ref, 0.01f); + float tau_d = fmaxf(q_max_d - q_min_d, tau_floor); float sum_e = 0.0f; float exps_d[4]; /* b0_size=4: Short/Hold/Long/Flat */ for (int a = 0; a < b0_size; a++) { diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index a11e08ac4..f4c94c2d3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -749,6 +749,19 @@ impl GpuBacktestEvaluator { self.n_windows } + /// Sharpe annualization factor used by `compute_backtest_metrics` kernel. + /// + /// Returns the EXACT factor applied to `(mean/std)` inside the kernel + /// (`sqrt(bars_per_day × trading_days_per_year)` — see line 667 init). + /// Callers that need to un-annualize the kernel's reported Sharpe MUST + /// use this getter rather than recomputing from their own copy of + /// `bars_per_day`, otherwise the two numbers can drift if the evaluator's + /// config was overridden between construction and read (e.g. via + /// hyperopt). Source-of-truth alignment per `feedback_no_partial_refactor.md`. + pub fn annualization_factor(&self) -> f32 { + self.annualization_factor + } + /// C.3 Plan 3 Task 7: device pointer + sample count for the val-state /// batch the most recent `evaluate_dqn_graphed` call left in /// `chunked_states_buf`. Layout: `[chunk_len × n_windows, state_dim_padded]` diff --git a/crates/ml/src/trainers/dqn/financials.rs b/crates/ml/src/trainers/dqn/financials.rs index eeb037ead..5f98906fc 100644 --- a/crates/ml/src/trainers/dqn/financials.rs +++ b/crates/ml/src/trainers/dqn/financials.rs @@ -58,19 +58,43 @@ pub(crate) fn compute_epoch_financials( 0.0 }; - // Annualized return from mean per-bar return (not compounded). - // Compounding even 10K bars produces extreme numbers when per-bar returns - // are biased. Annualized mean is more informative for monitoring: - // annual_return = mean_per_bar × bars_per_year - let bars_per_year = bars_per_day * 252.0; + // Total return: actual cumulative compounded portfolio return over the + // rollout — the standard meaning of "Return". Uses log-space sum to avoid + // f64 overflow on long rollouts and clamps step returns at -100% (loss + // of all capital) so a single extreme bar can't push (1+r) to ≤ 0. + // + // Previously this was `mean_per_bar × bars_per_year` ("annualized mean + // arithmetic return") which produced misleading numbers like -51,234% + // when per-bar mean was -52 bps × 98,280 bars/year. Multiplying any + // moderately negative per-bar mean by ~100K creates a value that LOOKS + // like portfolio collapse but is actually a label artefact — the formula + // assumed compounding when the metric used arithmetic annualization. + // + // The compounded form below cannot exceed -100% (clamped at log-space + // step floor) and matches what a trader would actually realise from + // following the policy over the rollout. Annualization is left to + // downstream callers if needed (multiply by bars_per_year/n_returns). + let bars_per_year = bars_per_day * 252.0; // retained for trade-annualization below let returns_slice = &trade_stats.step_returns; let n_returns = returns_slice.len(); - let mean_return = if n_returns > 0 { - returns_slice.iter().sum::() / n_returns as f64 + let total_return = if n_returns > 0 { + // Stable log-space cumulative: log(prod(1+r)) = sum(log(1+r)). + // Clamp 1+r to >= 1e-10 so a -100% bar gives a finite -23 log + // contribution rather than -inf; result then bounded above -100%. + let log_growth: f64 = returns_slice.iter() + .map(|&r| { + let g = (1.0_f64 + r).max(1e-10_f64); + if g.is_finite() { g.ln() } else { 0.0 } + }) + .sum(); + if log_growth.is_finite() { + log_growth.exp() - 1.0 + } else { + 0.0 + } } else { 0.0 }; - let total_return = mean_return * bars_per_year; // annualized fractional return let avg_return = if total_trades > 0 { total_return / total_trades as f64 } else { diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 98f6dc738..fc7dd5422 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -3735,7 +3735,19 @@ impl DQNTrainer { "Epoch {}/{}: val_Sharpe={:.2} (deterministic backtest on fixed validation window)", epoch + 1, self.hyperparams.epochs, val_sharpe, ); - let val_sharpe_raw = val_sharpe / f64::from(self.hyperparams.bars_per_day * 252.0).sqrt(); + // Source-of-truth alignment: read the EXACT annualization factor the + // kernel applied (`sqrt(bars_per_day × trading_days_per_year)`) rather + // than recomputing from `self.hyperparams.bars_per_day` here, which + // could drift if hyperopt or downstream callers override the + // evaluator's config between construction and metric emit. The + // evaluator's getter returns the same f32 stored at init time. + let val_sharpe_raw = match self.gpu_evaluator.as_ref() { + Some(ev) => { + let ann = f64::from(ev.annualization_factor()); + if ann > 1e-10 { val_sharpe / ann } else { 0.0 } + } + None => 0.0, + }; info!(" val_Sharpe_raw={:.6} (un-annualized per-bar)", val_sharpe_raw); // C.3 Plan 3 Task 7: state-distribution KL between train sample and val diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6d73a1127..a4d2fac06 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -184,6 +184,7 @@ P5T5 Phase E (2026-04-26): per-fold reset for `MetricBandsRegistry` regression-d | C.3 B.1/B.2 KL-amp consumers (`experience_kernels.cu` Flat opp_cost + entering_trade B.2 bonus) | `experience_env_step` kernel: B.1 site multiplies `reward = -shaping_scale × holding_cost_rate × conviction_core × vol_proxy_flat × q_abs_ref × kl_amp_b1`; B.2 site multiplies `bonus_scaled = shaping_scale × bonus × kl_amp_b2`. Both consumers use `fmaxf(1.0, isv_signals_ptr[ISV_STATE_KL_AMP_IDX])` to no-op against cold-start. `ISV_STATE_KL_AMP_IDX=79` macro defined in `state_layout.cuh`. | Wired | Plan 3 Task 7 C.3 — bounded amp ∈ [1, 2] stacks safely with B.1's existing `q_abs_ref` unbounded multiplicand per `pearl_one_unbounded_signal_per_reward.md`. | — | | `trainers/dqn/monitors/state_kl_monitor.rs` | Read-only observer for `state_kl_moment_match` kernel output (ISV slots 78 + 79); consumers: HEALTH_DIAG `state_kl.train_val_ema` / `state_kl.amp` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 7 C.3 | — | | `cuda_pipeline/scripted_policy_kernel.cu` | `GpuExperienceCollector::launch_timestep_loop` (per-timestep dispatch when `seed_phase_active_cache=true`); 4 policies (UNIFORM 40% / MOMENTUM 20% / MEAN_REV 20% / VWAP_DEV 20%) deterministically mixed by `i % 5`. Per-thread one sample. Reads `batch_states[i, MARKET_START]` for current close + `portfolio_states[i, PS_PREV_CLOSE]` for prev close; writes `batch_actions[i]` (factored `dir*27 + mag*9 + ord*3 + urg`) and `conviction_buf[i] ∈ [0, 1]` — same outputs the network's `experience_action_select` would have written. Direction-flip thresholds (`±0.0001f` momentum/reversal, `±5bp` VWAP band) are noise-floor cutoffs, not scaling coefficients. | Wired | Plan 3 Task 8 B.3 — GPU-only seeded warm-start. CPU only orchestrates per-epoch dispatch (cold-path read of ISV[SEED_STEPS_DONE/TARGET]); the action computation itself is 100% GPU. No CPU physics mirror — existing `experience_env_step` runs unchanged. | — | +| C51 bias C.1 ISV-adaptive direction-Boltzmann tau floor (`experience_kernels.cu` `experience_action_select` direction branch) | `experience_action_select` Branch 0 direction softmax: `tau_d = max(q_range, max(ISV[Q_DIR_ABS_REF_INDEX=21], 0.01))`. Replaces static `0.01` floor (tuned constant, dangerous when Q-magnitudes grow during training: spread that's small in absolute terms becomes deterministic argmax even though it represents no real edge relative to scale). Now scales with the network's actual Q magnitude EMA, preserving relative spread for sampling and preventing the C51 expected-Q Hold/Flat attractor that the val-Flat-collapse Kelly fix exposed (val_dir_dist S=.23/H=.17/L=.42/F=.18 epoch 1 → S=.14/H=.35/L=.15/F=.37 epoch 2 = 35→72% Hold+Flat shift in one epoch). Cold-start fallback retains 0.01 minimum so kernel doesn't divide by zero pre-first-update. Same ISV[21] reference used by conviction below — coherent adaptive mechanism. | Wired | val-collapse follow-up — C51 expected-Q bias fix. ISV-driven, no tuned constants per `feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`. | — | | `cuda_pipeline/seed_step_counter_update_kernel.cu` | `GpuExperienceCollector::launch_seed_step_counter_update_inplace` → `training_loop.rs` (called alongside the other Plan 3 EMA producers each epoch); single-block single-thread cold-path. Increments ISV[SEED_STEPS_DONE_INDEX=83] by `n_samples`, capped at ISV[SEED_STEPS_TARGET_INDEX=82], and EMAs the derived `max(0, 1 - DONE/TARGET)` into ISV[SEED_FRAC_EMA_INDEX=84]. Adaptive α=α_base × (1+0.5×\|clamp(sharpe,−2,2)\|), α_base=0.05. | Wired | Plan 3 Task 8 B.3 — cold-path (per-epoch). No atomicAdd. SEED_FRAC_EMA cold-start 1.0 ensures Task 9's CQL ramp sees `target=0` until the seed phase actually decays. | — | | `cuda_pipeline/cql_alpha_seed_update_kernel.cu` | `GpuExperienceCollector::launch_cql_alpha_seed_update_inplace` → `training_loop.rs` (called alongside `launch_seed_step_counter_update_inplace` each epoch); single-block single-thread cold-path. EMAs ISV[CQL_ALPHA_INDEX=48] toward `cql_alpha_final × max(0, 1 - ISV[SEED_FRAC_EMA_INDEX=84])` where `cql_alpha_final = config.cql_alpha`. Adaptive α matches Task 8 convention. | Wired | Plan 3 Task 9 C.5 — producer upgrade for ISV[CQL_ALPHA_INDEX=48]: SchemaContract → FoldReset. CQL gradient kernel consumer (already reading slot 48 per Plan 1 Task 12) unchanged. | — | | `trainers/dqn/monitors/seed_monitor.rs` | Read-only observer for `seed_step_counter_update` kernel output (ISV slots 82 / 83 / 84); consumers: HEALTH_DIAG `seed.steps_target` / `seed.steps_done` / `seed.frac_ema` + `controller_activity` smoke fire-rate tracking | Wired | Plan 3 Task 8 B.3 | — |