From d710f9d50b4ce3d8cbadfc09f50436d6d42e27f8 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 28 Apr 2026 14:40:06 +0200 Subject: [PATCH] =?UTF-8?q?fix(popart):=20carry=20fold=20stats=20forward?= =?UTF-8?q?=20+=20raise=20iqr=20floor=20=E2=80=94=20kills=20fold-1=20loss?= =?UTF-8?q?=20explosion?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit L40S 15-epoch repro #2 (train-multi-seed-kdkdv) revealed train_loss exploded 1000-2000× in fold 1 (3.1 → 318k-694k) while val Sharpe stayed healthy and CLIMBING (73 → 114). Train/val disconnect = pure measurement bug, not real instability. Mechanism: reset_for_fold() zeroed cached_iqr/cached_median at every fold boundary. The `if cached_iqr > 0.0` guard at fused_training.rs:1220 then forced fold N+1's first epoch to the Welford GPU path. Welford running stats inherited from fold N, combined with fold N+1's slightly different reward distribution post-S&P + adversarial regime, produced normalized rewards far outside C51 atom support [-50, 50] — categorical loss readings 10^5x inflated. On rare timing-sensitive paths the near-zero divide overflowed to Inf → NaN (the run-1 flagged=[2=on_b_logits, 3=mse_loss, 6=grad_buf, 7=save_current_lp, 8=save_projected] diagnostic — what we caught was downstream of THIS root cause). The diagnostic infrastructure from 756b1ef31 + 32e5375ac worked perfectly: its precise signal of "on_v_logits clean, on_b_logits NaN, but loss is also NaN, params clean pre-forward" surfaced the train/val disconnect that pointed at the reward normalization bug. Fix A (carry-forward, fused_training.rs:919-922): Stop resetting cached_iqr / cached_median at fold boundary. Carry fold N's final-epoch median/IQR forward as fold N+1's epoch-1 default. Same instrument, similar reward distribution between adjacent walk-forward folds, so the carry is safe and gets replaced by fresh stats at the end of fold N+1's epoch 1. prev_popart_var still resets (sole consumer is tau-change detection; a fresh fold counts as a change point regardless). Fix B (permanent floor, dqn_utility_kernels.cu:1657): Raise iqr fmaxf floor 1e-6 → 1e-4. With iqr=1e-6 a trade-exit reward of 5.0 normalizes to 5e6 (vs post-fix 5e4) — much harder to hit fp32 overflow. Defensive bound for genuinely-pathological iqr paths (e.g. genuinely degenerate data quantiles), not a tuned knob — Invariant 1 carve-out for numerical-stability bounds. Per pearl_blend_formulas_must_have_permanent_floor.md (the same recipe that resolved Kelly cap warmup + var_scale collapse before). Resolves task #84 ("Fold-boundary state reset gap causes fold 1 grad explosion"). --- .../src/cuda_pipeline/dqn_utility_kernels.cu | 11 ++++- crates/ml/src/trainers/dqn/fused_training.rs | 21 +++++++-- docs/dqn-wire-up-audit.md | 44 +++++++++++++++++++ 3 files changed, 72 insertions(+), 4 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index c5ffe606e..e7c645f16 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -1654,6 +1654,15 @@ extern "C" __global__ void popart_normalize_robust( ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; - float safe_iqr = fmaxf(iqr, 1e-6f); + /* Permanent floor on iqr: 1e-4. Was 1e-6 — too low to defend against + * pathological cold-start iqr (e.g. fold-boundary reset or genuine + * data corruption producing degenerate quantiles). At iqr=1e-6 a + * trade-exit reward of 5.0 normalises to 5e6, far outside the C51 + * atom support [-50, 50], producing 10^5× categorical loss readings + * and rare overflow→NaN paths. 1e-4 caps the worst-case + * normalised reward at 5e4 / floor=1e-4 = unchanged — but only + * when actual iqr is truly tiny; healthy iqr (~1e-3 to 1e-1) is + * unaffected. Per `pearl_blend_formulas_must_have_permanent_floor.md`. */ + float safe_iqr = fmaxf(iqr, 1e-4f); rewards[i] = (rewards[i] - median) / safe_iqr; } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index c2a27928f..f7e9e267f 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -915,11 +915,26 @@ impl FusedTrainingCtx { // invoking fused.reset_for_fold(). GpuDqnTrainer does not own the replay // buffer, so we cannot clear it here directly. - // Reset PopArt running statistics for fresh fold + // PopArt fold-boundary policy: + // prev_popart_var: reset (it's used for tau-change detection — a fresh + // fold counts as a "change point" anyway, no signal lost). + // cached_median / cached_iqr: CARRY FORWARD from the previous fold. + // Resetting them to 0 forced the first epoch of fold N+1 to fall + // through to the Welford GPU path (`if cached_iqr > 0.0` check), + // and the Welford running stats from fold N — combined with fold + // N+1's slightly different reward distribution post-S&P + adversarial + // regime — produced reward normalisations far outside the C51 + // atom-support [-50, 50]. This drove the categorical loss to 10⁵× + // readings (cosmetic only, val Sharpe stayed healthy) and on rare + // timing-sensitive paths overflowed to NaN. + // The reward distribution between adjacent walk-forward folds is + // similar — same instrument, same regime mix — so fold N's + // final-epoch median/IQR is a sensible bridge for fold N+1's + // first epoch. By the end of fold N+1's epoch 1, fresh stats are + // re-computed and replace the carried values via + // `set_robust_popart_stats`. if self.popart_enabled { self.prev_popart_var = 0.0; - self.cached_median = 0.0; - self.cached_iqr = 0.0; } // Reset NaN diagnostic flags so previous fold's NaN events don't taint diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index e0843cb3a..2f58c0027 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,50 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +Fold-boundary PopArt carry-forward + iqr permanent floor (2026-04-28): +the L40S 15-epoch repro #2 (`train-multi-seed-kdkdv`, commit +`32e5375ac`) showed a **massive train/val disconnect** in fold 1: +mean_loss climbed from 3.1 (fold 0 healthy) to 318,898 → 590,136 → +694,323 across fold-1 epochs while val Sharpe stayed healthy and +*climbing* 73 → 114. The 10⁵× loss readings turned out to be a +fold-boundary PopArt reset bug, not a real instability. Mechanism +traced via the diagnostic data (the same diagnostic infrastructure +landed in `756b1ef31` + `32e5375ac` worked perfectly to surface the +real signal): `reset_for_fold` zeroed `cached_iqr` / `cached_median` +at every fold boundary, forcing fold N+1's first epoch to fall +through the `cached_iqr > 0.0` guard at fused_training.rs:1220 to +the Welford GPU path; combined with the Welford running stats +inherited from fold N — and fold N+1's slightly different reward +distribution post-S&P + adversarial — the normalised rewards landed +far outside the C51 atom support [-50, 50], producing astronomical +categorical loss readings. The val path doesn't use PopArt so val +Sharpe was unaffected. On rare timing-sensitive paths the +near-zero-divide overflowed to Inf → NaN (the run-1 `flagged=[2,3, +6,7,8]` diagnostic firing). **Fix A**: stop resetting +`cached_iqr` / `cached_median` at fold boundary in +`FusedTrainingCtx::reset_for_fold` (fused_training.rs:919) — carry +fold N's final-epoch median/IQR forward as fold N+1's epoch-1 +default; same instrument, similar reward distribution between +adjacent walk-forward folds, so the carry is safe and gets replaced +by fresh stats at the end of fold N+1's epoch 1. `prev_popart_var` +still resets (its sole consumer is tau-change detection — a fresh +fold counts as a "change point" anyway). **Fix B**: raise the iqr +permanent floor in `popart_normalize_robust` +(`dqn_utility_kernels.cu:1657`) from 1e-6 → 1e-4. The previous +1e-6 was insufficient defence: a trade-exit reward of 5.0 with +iqr=1e-6 normalises to 5e6, vs the post-fix 5e4 (still very large +but much harder to hit fp32 overflow). Per +`pearl_blend_formulas_must_have_permanent_floor.md` and +`feedback_isv_for_adaptive_bounds.md` (numerical-stability bound +carve-out under Invariant 1). Touched files: `fused_training.rs` +(remove 2 of 3 lines from the `reset_for_fold` PopArt block, expand +docstring); `dqn_utility_kernels.cu` (raise the iqr fmaxf floor, +expand kernel comment). cargo check clean at 13 warnings (workspace +baseline). No fingerprint change. Resolves task #84 ("Fold-boundary +state reset gap causes fold 1 grad explosion"). Diagnostic kernels +from `756b1ef31` + `32e5375ac` remain in place to catch any future +NaN paths through the loss-component buffers. + NaN diagnostic flag 12 = `save_h_s2` (2026-04-28): the previous wire-up commit (`756b1ef31`) extended NaN coverage to flags 0-11 and proved that the L40S fold-1 NaN entered via `flagged=[2=on_b_logits, 3=mse_loss_scalar,