From 4ef1d8ebb7f8360635db63afe455e5db07d05ec3 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 29 Apr 2026 20:24:36 +0200 Subject: [PATCH] =?UTF-8?q?fix(dqn):=20Plan=20C=20K=20=E2=80=94=20Adam=20s?= =?UTF-8?q?hrink-and-perturb=20+=20adaptive=20fold-warmup=20ISV?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes the F1 ep1 catastrophic-overshoot gap exposed by smoke-test-s9h4h (F0 succeeded with Best Sharpe = 36.03; F1 ep1 grad_norm = 355,009 — 5 orders of magnitude larger than F0 steady-state ~10 — leading to NaN propagation and grad-clamp-to-zero early stop). Combined fix: (1) Adam shrink-and-perturb at fold boundary (replace m=0/v=0 with m*=0.1, v*=0.01 to preserve direction while damping magnitude), and (2) a single adaptive ISV signal driving BOTH lr_eff and clip_eff dampening over the fold's first ~50 steps. K — Adam shrink-and-perturb in `GpuDqnTrainer::reset_adam_state`: - m *= 0.1, v *= 0.01 via existing `dqn_scale_f32_kernel` (loaded as `scale_f32_ungraphed`); t_pinned still zeroed so bias correction restarts. Architectural constants (preserve direction / lose magnitude history) per `feedback_isv_for_adaptive_bounds.md` Invariant 1 carve-out — not tuned. Composes with existing param shrink-and-perturb (`alpha=0.8`) in `FusedTrainingCtx::reset_for_fold`. Root cause for F1 overshoot: m=0,v=0 → first Adam step ≈ lr × g / ε → 6 OoM amplification. New CUDA kernel `fold_warmup_factor_kernel.cu`: - Single-block single-thread cold-path producer mirroring `q_drift_rate_ema_kernel.cu` / `moe_lambda_eff_kernel.cu` shape. - Reads two grad-norm EMAs (fast α=0.1, slow α=0.001) plus host-passed step counter; writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1). - Bootstrap branches (steps_observed < 200, slow EMA < 1e-6) emit factor=1.0 (no damping during cold-start). No atomicAdd; no DtoH. New ISV slot Q_DRIFT_RATE_INDEX → FOLD_WARMUP_FACTOR_INDEX = 130: - ISV_TOTAL_DIM 130 → 131; layout fingerprint shifts (checkpoint- incompatible per `feedback_no_legacy_aliases.md`, expected for a real architecture change). - FoldReset entries: `isv_fold_warmup_factor` → 0.0 and companion `isv_grad_norm_fast_ema` → 0.0 (lockstep reset per `feedback_no_partial_refactor.md`); slow EMA persists across folds as the cross-fold steady-state baseline. - Two new mapped-pinned scalars on GpuDqnTrainer (grad_norm_fast_ema_pinned, grad_norm_slow_ema_pinned) fed by `update_adaptive_clip` from the same `gr.raw_grad_norm` observation source as the existing adaptive clip EMA. Two consumers, both monotone (only dampen, never excite): - lr_eff = cosine_effective_lr × max(MIN_WARMUP_LR_FRAC=0.05, factor) via `set_lr` per-step. New `cosine_effective_lr_base` field on DQNTrainer composes the cosine schedule's per-epoch baseline with the warmup factor's per-step damping (rather than overriding the cosine schedule). - clip_eff = clip_base × (MIN_CLIP_FRAC=0.1 + 0.9 × factor) via new `set_active_clip` setter on FusedTrainingCtx + GpuDqnTrainer. Composes with the EMA-derived `clip_base = grad_norm_ema × 2` that `update_adaptive_clip` just wrote to the pinned slot. Numerical-stability bounds 0.05 / 0.1 are Invariant 1 carve-outs. Steady-state behaviour unchanged: factor=1 → lr_eff=lr_base, clip_eff=clip_base. Fold-boundary behaviour: factor starts at 0 → lr_eff = 0.05 × lr_base, clip_eff ≈ 0.1 × clip_base; rises to 1 over ~50 steps as the fast EMA catches up to the slow steady-state EMA. Predicted impact on Plan C smoke F1: 355,009-magnitude transient grad clipped to ~clip_base × 0.1 ≈ 1.0 (vs 10), Adam state shrunk instead of zeroed → first step update bounded; grad recovers normally over ~50 steps. Companion to A.1 (prev_epoch_q_mean reset), A.2 (adaptive Polyak-tau), A.3 (gradient_collapse_counter reset), F+H (kill-criterion robustness) — completes the fold-boundary state-reset family. Per `pearl_adaptive_moe_lambda.md` (kernel + ISV slot + bootstrap + reset + observability template), `pearl_cold_path_no_exception_to_gpu_drives.md` (GPU-stays-on-GPU even at cold-path cadence), `pearl_blend_formulas_must_have_permanent_floor.md` (lr/clip floors are permanent minimums), `feedback_adaptive_not_tuned.md` (lr+clip ISV-driven), `feedback_isv_for_adaptive_bounds.md` (factor IS the bound; consumers compose at runtime), `feedback_no_atomicadd.md` (single-thread reduce), `feedback_cudarc_f64_f32_abi.md` (slot index passed as i32), `feedback_no_partial_refactor.md` (kernel + slot + reset + producer + 2 consumers all land together), `feedback_no_quickfixes.md` (replaces brittle full-reset with adaptive damping; not threshold relaxation). Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 10 + .../fold_warmup_factor_kernel.cu | 88 +++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 324 +++++++++++++++++- crates/ml/src/trainers/dqn/fused_training.rs | 20 ++ .../src/trainers/dqn/state_reset_registry.rs | 27 ++ .../src/trainers/dqn/trainer/constructor.rs | 11 + crates/ml/src/trainers/dqn/trainer/mod.rs | 13 + .../src/trainers/dqn/trainer/training_loop.rs | 100 ++++++ docs/dqn-wire-up-audit.md | 24 ++ 9 files changed, 611 insertions(+), 6 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 2ea764085..025a3b68e 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -211,6 +211,16 @@ fn main() { // 1/(1+ISV[129]) so tau ∈ [tau_base/5, tau_base] — monotone // dampening under drift; healthy runs unaffected. "q_drift_rate_ema_kernel.cu", + // Plan C Phase 2 follow-up K (2026-04-29): adaptive fold-boundary + // warmup signal. Single-block single-thread cold-path producer + // mirroring q_drift_rate_ema_kernel / moe_lambda_eff_kernel + // precedent. Reads two grad-norm EMAs (fast α=0.1, slow α=0.001) + // and writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1). + // Two consumers, both monotone (only dampen): lr_eff = lr_base × max(0.05, factor) + // and clip_eff = clip_base × (0.1 + 0.9 × factor). After fold reset + // factor starts at 0 → heavy damping; rises to 1 as gradients + // stabilise → consumers return to baseline (healthy runs unaffected). + "fold_warmup_factor_kernel.cu", // HEALTH_DIAG GPU port — Phase 2A landing (2026-04-28). Single-block // single-thread `health_diag_isv_mirror` kernel copies a curated set // of ISV signal-bus slots into the mapped-pinned `HealthDiagSnapshot` diff --git a/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu b/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu new file mode 100644 index 000000000..2697e9902 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu @@ -0,0 +1,88 @@ +/* fold_warmup_factor_update — adaptive fold-boundary warmup signal, GPU-driven. + * + * Plan C Phase 2 follow-up K (2026-04-29). Companion to A.2's q-drift-rate + * adaptive Polyak-tau dampening. Single ISV signal driving BOTH lr_eff and + * clip_eff at fold boundaries; both consumers are monotone (only dampen, + * never excite). + * + * Reads: current_grad_norm_ema — fast EMA of `gr.raw_grad_norm` (α=0.1, ~10-step) + * steady_grad_norm_ema — slow EMA of same source (α=0.001, ~1000-step) + * min_steps_for_ratio — numerical bootstrap guard (cold-start delay) + * steps_observed — step counter (host-driven; gates compute) + * Writes: ISV[FOLD_WARMUP_FACTOR_INDEX=130] — clamped factor ∈ [0, 1] + * + * Cold-path cadence — single-block, single-thread kernel matching + * `q_drift_rate_ema_kernel.cu`, `moe_lambda_eff_kernel.cu`, and + * `kelly_cap_update_kernel.cu` shapes. Launched once per training step AFTER + * the per-step grad-norm EMAs have been updated (host-side `update_adaptive_clip` + * already feeds these EMAs from `gr.raw_grad_norm`). + * + * Formula: + * + * if steps_observed < min_steps_for_ratio -> factor = 1.0 (no damping; pre-warmup) + * else if steady < 1e-6 -> factor = 1.0 (slow EMA cold-start) + * else -> factor = clamp(curr / steady, 0, 1) + * + * The `1e-6` denominator floor is a numerical-stability bound (Invariant 1 + * carve-out per `feedback_isv_for_adaptive_bounds.md`). The `[0, 1]` clamp + * is structural — `factor > 1` would mean fast > slow (a transient spike) and + * we explicitly do NOT excite past baseline (monotone-only-dampens contract, + * matching A.2's `q_drift_rate_ema_update` kernel). After fold reset + * (FoldReset: fast EMA → 0, ISV[130] → 0; slow EMA persists across folds), + * curr starts at 0 → factor = 0 → heavy damping; as curr recovers toward + * steady → factor → 1 → consumers return to baseline. + * + * Per `pearl_cold_path_no_exception_to_gpu_drives.md`: even cold-path scalar + * arithmetic stays on GPU when its EMA inputs already live on GPU (mapped- + * pinned host ptrs visible to the device). + * + * No atomicAdd (per `feedback_no_atomicadd.md`); a single thread reads two + * scalars and writes one ISV slot. The `__threadfence_system()` mirrors the + * other ISV producers' emission so the mapped-pinned host ptr sees the + * fresh value before downstream consumers (lr/clip composers) read it. + * + * Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic + * drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold- + * start) + reset (0.0 fold boundary) + observability (HEALTH_DIAG via slot + * 130 mirror). + */ + +extern "C" __global__ void fold_warmup_factor_update( + float* __restrict__ isv, /* ISV bus [ISV_TOTAL_DIM] */ + const float* __restrict__ current_grad_norm_ema, /* [1] fast EMA of grad norm */ + const float* __restrict__ steady_grad_norm_ema, /* [1] slow EMA of grad norm */ + int isv_fold_warmup_factor_idx, /* FOLD_WARMUP_FACTOR_INDEX = 130 */ + int min_steps_for_ratio, /* host counter floor */ + int steps_observed /* host-passed step counter */ +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float curr = current_grad_norm_ema[0]; + const float steady = steady_grad_norm_ema[0]; + + /* Bootstrap branches: + * (a) Pre-warmup window: not enough steps yet for the EMAs to be + * informative — let the system run freely (factor = 1). + * (b) Slow EMA still cold-start (~zero): same — no meaningful + * baseline to compare against, so emit no-op. The `1e-6` floor + * is the same Invariant 1 numerical-stability bound used by + * `q_drift_rate_ema_kernel.cu`. + * + * Otherwise compute the ratio and clamp to [0, 1]. The upper clamp + * preserves the monotone-dampening contract (factor > 1 would EXCITE + * lr/clip past baseline, which we explicitly forbid). The lower clamp + * is structurally guaranteed (curr / steady ≥ 0 when both are non- + * negative, which they are by construction) but kept defensively. */ + float factor; + if (steps_observed < min_steps_for_ratio) { + factor = 1.0f; + } else if (steady < 1e-6f) { + factor = 1.0f; + } else { + const float raw = curr / steady; + factor = fmaxf(0.0f, fminf(1.0f, raw)); + } + + isv[isv_fold_warmup_factor_idx] = factor; + __threadfence_system(); /* PCIe-visible to mapped pinned host_ptr */ +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 89582d6da..338f565c7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -194,6 +194,14 @@ pub(crate) static AUX_HEADS_LOSS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!( /// q-drift rises. static Q_DRIFT_RATE_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/q_drift_rate_ema_kernel.cubin")); +/// Plan C Phase 2 follow-up K (2026-04-29): adaptive fold-boundary warmup +/// signal. Single-block single-thread cold-path producer mirroring +/// `q_drift_rate_ema_kernel.cu` / `moe_lambda_eff_kernel.cu` shape. Reads +/// two grad-norm EMAs (fast α=0.1, slow α=0.001) plus a host-passed step +/// counter; writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1). +/// Two consumers (lr_eff and clip_eff in `training_loop.rs`), both monotone. +static FOLD_WARMUP_FACTOR_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/fold_warmup_factor_kernel.cubin")); + /// Plan 4 Task 1B-ii: per-group VSN MLP hidden dim. `Linear_1[g]` is /// `[VSN_HIDDEN_DIM, group_dim_g]`, `Linear_2[g]` is `[1, VSN_HIDDEN_DIM]`. /// 16 chosen as a tractable expansion that fits comfortably against the @@ -452,7 +460,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -const ISV_TOTAL_DIM: usize = 130; +const ISV_TOTAL_DIM: usize = 131; /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -877,6 +885,36 @@ pub const MOE_LAMBDA_EFF_INDEX: usize = 128; /// `feedback_no_quickfixes.md` — neither is a tuned constant. pub const Q_DRIFT_RATE_INDEX: usize = 129; +/// ISV slot [130] — adaptive fold-boundary warmup factor ∈ [0, 1]. +/// +/// Plan C Phase 2 follow-up K (2026-04-29): single signal driving BOTH lr_eff +/// and clip_eff dampening at fold boundaries. Computed per-step by +/// `fold_warmup_factor_kernel.cu` from grad-norm stability: +/// `factor = current_grad_norm_ema / steady_state_grad_norm_ema` +/// clamped to `[0, 1]`. After fold reset (FoldReset: 0.0) the fast EMA recovers +/// from zero while the slow EMA tracks the cross-fold steady-state grad scale, +/// so the ratio rises 0 → 1 as gradients stabilise. Steady-state behaviour: +/// fast EMA matches slow EMA → factor = 1 → no dampening (healthy runs +/// unaffected — same monotone-only-dampens contract as A.2's +/// `Q_DRIFT_RATE_INDEX`). +/// +/// Two consumers, both monotone (only dampen, never excite): +/// `lr_eff = lr_base × max(MIN_WARMUP_LR_FRAC=0.05, factor)` +/// `clip_eff = clip_base × (MIN_CLIP_FRAC=0.1 + (1 - MIN_CLIP_FRAC) × factor)` +/// +/// Cold-start (constructor): 0.0 (heavy damping until grad-norm EMAs warm up). +/// FoldReset: 0.0 (re-engage damping for the new fold's first steps; the slow +/// EMA persists across folds so the steady-state baseline stays meaningful). +/// +/// Tail-appended at index 130 (one past the previous `ISV_TOTAL_DIM=130` +/// boundary, raising it to 131) per `feedback_no_partial_refactor.md`. Layout +/// fingerprint changes — checkpoint-incompatible per `feedback_no_legacy_aliases.md`. +/// The numerical-stability bounds `0.05` (lr floor frac) and `0.1` (clip floor +/// frac) are Invariant 1 carve-outs per `feedback_isv_for_adaptive_bounds.md` +/// — the warmup factor is the *adaptive* bound; consumers read ISV[130] at +/// runtime and compose it with their respective architectural floors. +pub const FOLD_WARMUP_FACTOR_INDEX: usize = 130; + /// ISV slot [115] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// /// Design note: this is NOT a version number. There is no ordered version space. @@ -993,7 +1031,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { MOE_GATE_ENTROPY_EMA=126;\ MOE_LAMBDA_EFF=128;\ Q_DRIFT_RATE=129;\ - ISV_TOTAL_DIM=130;\ + FOLD_WARMUP_FACTOR=130;\ + ISV_TOTAL_DIM=131;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ @@ -3234,6 +3273,41 @@ pub struct GpuDqnTrainer { /// runs unaffected. Loaded from `q_drift_rate_ema_kernel.cubin`. q_drift_rate_ema_kernel: CudaFunction, + // ── Plan C Phase 2 follow-up K: adaptive fold-boundary warmup ───── + /// Single-thread single-block cold-path producer kernel. Reads the + /// fast/slow grad-norm EMA mapped-pinned buffers and writes + /// `ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1)`. + /// Consumers: `training_loop.rs` reads ISV[130] each step to derive + /// `lr_eff = lr_base × max(0.05, factor)` (passed via `set_lr`) and + /// `clip_eff = clip_base × (0.1 + 0.9 × factor)` (passed via the + /// existing `update_adaptive_clip` pinned slot). Loaded from + /// `fold_warmup_factor_kernel.cubin`. + fold_warmup_factor_kernel: CudaFunction, + /// Mapped-pinned host pointer for the fast (α=0.1, ~10-step horizon) + /// grad-norm EMA. Updated by `update_adaptive_clip` per training step + /// (the same observation source feeds both this fast EMA and the + /// existing `grad_norm_ema` field used by adaptive clip). FoldReset: + /// 0.0 (the new fold's first steps observe the recovery from zero). + grad_norm_fast_ema_pinned: *mut f32, + /// Device-mapped view of `grad_norm_fast_ema_pinned`. + grad_norm_fast_ema_dev_ptr: u64, + /// Mapped-pinned host pointer for the slow (α=0.001, ~1000-step horizon) + /// grad-norm EMA. Tracks the cross-fold steady-state grad scale — + /// persists across folds so the warmup factor's denominator stays a + /// meaningful long-window baseline. Constructor: 0.0 (cold-start; the + /// `min_steps_for_ratio` gate in the kernel keeps the ratio at 1.0 + /// until enough samples have accumulated). + grad_norm_slow_ema_pinned: *mut f32, + /// Device-mapped view of `grad_norm_slow_ema_pinned`. + grad_norm_slow_ema_dev_ptr: u64, + /// Number of times `update_adaptive_clip` has fed a finite grad-norm + /// observation into the EMAs. The `fold_warmup_factor_update` kernel + /// reads this as its bootstrap gate — until it crosses + /// `WARMUP_FACTOR_MIN_STEPS`, the factor is forced to 1.0 (no damping). + /// Survives across folds (the slow EMA also persists, so the gate + /// only fires the very first time it warms up at training start). + grad_norm_emas_step_count: u64, + // ── Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMA producer ── /// 4-block kernel (one per off-median quantile τ ∈ {0.05, 0.25, 0.75, 0.95}) /// computing the mean |Q| over (B × TBA) at each fixed-τ position from @@ -3643,10 +3717,64 @@ impl GpuDqnTrainer { /// Zeroes momentum (m), variance (v), and step counter (t). /// Weights are preserved (warm-start for the new fold). pub fn reset_adam_state(&mut self) -> Result<(), MLError> { - self.stream.memset_zeros(&mut self.m_buf) - .map_err(|e| MLError::ModelError(format!("reset m_buf: {e}")))?; - self.stream.memset_zeros(&mut self.v_buf) - .map_err(|e| MLError::ModelError(format!("reset v_buf: {e}")))?; + // Plan C Phase 2 follow-up K (2026-04-29): Adam shrink-and-perturb at + // fold boundary. Replaces the prior full-zero reset of `m_buf` / + // `v_buf` with a multiplicative shrink — preserves directional + // information while damping magnitude. Composes with the existing + // param shrink-and-perturb (`alpha=0.8` in + // `FusedTrainingCtx::reset_for_fold`) but is more aggressive + // because Adam moments are more cold-start-fragile than parameters + // (zeroed `v` makes the Adam denominator `sqrt(v_hat) + ε ≈ ε`, + // so the FIRST post-reset step is `lr × m̂ / ε` — a catastrophic + // overshoot when the new fold's data distribution differs from + // the previous fold's). + // + // Numerical-stability bounds (Invariant 1 carve-out per + // `feedback_isv_for_adaptive_bounds.md`): the shrink factors 0.1 + // (m) and 0.01 (v) are STRUCTURAL constants ("preserve the + // optimisation direction, lose the magnitude history") not tuned + // values. Shrinking `m` by 0.1 keeps the gradient direction's + // sign+ratio across tensors; shrinking `v` by 0.01 keeps the + // per-tensor scale ratios while letting the new fold's gradients + // re-grow the absolute magnitude. Smoke-test-s9h4h F0 successful + // (Best Sharpe = 36.03), F1 ep1 grad_norm = 355,009 (5 OoM + // larger than F0 steady-state ~10) → root-caused to full-zero + // Adam reset's denominator pathology. + // + // `t_pinned` (Adam step counter) IS still zeroed: the bias + // correction `1 - β^t` should restart so the new fold's first + // few steps are properly bias-corrected. Only `m`/`v` magnitude + // history is shrunk (not direction). + const ADAM_M_SHRINK: f32 = 0.1; + const ADAM_V_SHRINK: f32 = 0.01; + let n = self.total_params as i32; + let blocks = ((self.total_params as u32 + 255) / 256) as u32; + let cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + let m_ptr = self.m_buf.raw_ptr(); + let v_ptr = self.v_buf.raw_ptr(); + unsafe { + // dqn_scale_f32_kernel signature: (y, alpha, n) → y[i] *= alpha. + // Use scale_f32_ungraphed (separate CUmodule) — `reset_adam_state` + // is called outside the captured forward_child / adam_child + // graphs (fold boundary), so the ungraphed kernel module + // applies (mirrors `scale_adam_momentum`'s rationale). + self.stream.launch_builder(&self.scale_f32_ungraphed) + .arg(&m_ptr) + .arg(&ADAM_M_SHRINK) + .arg(&n) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("reset_adam_state shrink m: {e}")))?; + self.stream.launch_builder(&self.scale_f32_ungraphed) + .arg(&v_ptr) + .arg(&ADAM_V_SHRINK) + .arg(&n) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("reset_adam_state shrink v: {e}")))?; + } unsafe { *self.t_pinned = 0; } self.adam_step = 0; @@ -3954,6 +4082,14 @@ impl Drop for GpuDqnTrainer { if !self.adaptive_clip_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.adaptive_clip_pinned.cast()) }; } + // Plan C Phase 2 follow-up K (2026-04-29): grad-norm fast/slow EMA + // mapped-pinned scalars feeding `fold_warmup_factor_update`. + if !self.grad_norm_fast_ema_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.grad_norm_fast_ema_pinned.cast()) }; + } + if !self.grad_norm_slow_ema_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.grad_norm_slow_ema_pinned.cast()) }; + } if !self.sel_t_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.sel_t_pinned.cast()) }; } @@ -9687,6 +9823,17 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("q_drift_rate_ema_update load: {e}")))? }; + // Plan C Phase 2 follow-up K: load fold_warmup_factor kernel + // (cold-path, per-step). Single-thread single-block ISV producer + // for ISV[FOLD_WARMUP_FACTOR_INDEX=130]; consumers are the + // `lr_eff` and `clip_eff` derivations in `training_loop.rs`. + let fold_warmup_factor_kernel = { + let module = stream.context().load_cubin(FOLD_WARMUP_FACTOR_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("fold_warmup_factor cubin load: {e}")))?; + module.load_function("fold_warmup_factor_update") + .map_err(|e| MLError::ModelError(format!("fold_warmup_factor_update load: {e}")))? + }; + // Plan 4 Task 3 (E.3): load iqn_quantile_ema kernel (cold-path, per-step). // 4-block kernel — one block per off-median fixed quantile in // FIXED_TAUS = [0.05, 0.25, 0.50, 0.75, 0.95]. Producer for @@ -9826,6 +9973,52 @@ impl GpuDqnTrainer { dev_ptr }; + // Plan C Phase 2 follow-up K (2026-04-29): grad-norm fast/slow EMA + // mapped-pinned scalars feeding `fold_warmup_factor_update`. Both + // are written by the host-side `update_adaptive_clip` (the same + // observation source that maintains the legacy `grad_norm_ema` + // field used by the existing adaptive clip). Mapped-pinned so the + // kernel reads them via device-pointer with zero HtoD per + // `feedback_no_htod_htoh_only_mapped_pinned.md`. + // + // FoldReset semantics (handled by `reset_named_state` arm + // `isv_grad_norm_fast_ema`): fast → 0.0; slow persists across + // folds (the cross-fold steady-state grad scale is the meaningful + // baseline against which the new fold's recovery is measured — + // resetting both would defeat the warmup signal entirely). + let grad_norm_fast_ema_pinned: *mut f32 = unsafe { + let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; + cudarc::driver::result::malloc_host(std::mem::size_of::(), flags) + .map_err(|e| MLError::ModelError(format!("pinned grad_norm_fast_ema alloc: {e}")))? + as *mut f32 + }; + unsafe { *grad_norm_fast_ema_pinned = 0.0_f32; } + let grad_norm_fast_ema_dev_ptr = unsafe { + let mut dev_ptr: u64 = 0; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2( + &mut dev_ptr as *mut u64, + grad_norm_fast_ema_pinned.cast(), + 0, + ); + dev_ptr + }; + let grad_norm_slow_ema_pinned: *mut f32 = unsafe { + let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; + cudarc::driver::result::malloc_host(std::mem::size_of::(), flags) + .map_err(|e| MLError::ModelError(format!("pinned grad_norm_slow_ema alloc: {e}")))? + as *mut f32 + }; + unsafe { *grad_norm_slow_ema_pinned = 0.0_f32; } + let grad_norm_slow_ema_dev_ptr = unsafe { + let mut dev_ptr: u64 = 0; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2( + &mut dev_ptr as *mut u64, + grad_norm_slow_ema_pinned.cast(), + 0, + ); + dev_ptr + }; + // ── Allocate consolidated transfer buffers ───────────────── // Upload staging: states + next_states (f32) // Rewards/dones uploaded separately as f32 @@ -12511,6 +12704,12 @@ impl GpuDqnTrainer { reward_component_ema_kernel, h_s2_rms_ema_kernel, q_drift_rate_ema_kernel, + fold_warmup_factor_kernel, + grad_norm_fast_ema_pinned, + grad_norm_fast_ema_dev_ptr, + grad_norm_slow_ema_pinned, + grad_norm_slow_ema_dev_ptr, + grad_norm_emas_step_count: 0, iqn_quantile_ema_kernel, q_mean_reduce_kernel, q_mean_subtract_kernel, @@ -19349,11 +19548,124 @@ impl GpuDqnTrainer { let new_clip = (self.grad_norm_ema * CLIP_MULTIPLIER).max(MIN_CLIP); // Write directly to pinned memory — GPU sees it on next kernel read unsafe { *self.adaptive_clip_pinned = new_clip; } + + // Plan C Phase 2 follow-up K (2026-04-29): feed the SAME observation + // into the fast/slow grad-norm EMAs that drive the fold-boundary + // warmup factor `fold_warmup_factor_update`. Constants are + // numerical-stability bounds per `feedback_isv_for_adaptive_bounds.md` + // — fast α=0.1 is the architectural ~10-step horizon (recover from + // fold reset in roughly the time the existing adaptive clip's EMA + // takes to recapture grad scale); slow α=0.001 is the ~1000-step + // horizon (cross-fold steady-state baseline). Both are STRUCTURAL + // (preserve direction / amortise scale across folds), not tuned. + const FAST_ALPHA: f32 = 0.1; + const SLOW_ALPHA: f32 = 0.001; + unsafe { + let fast_prev = *self.grad_norm_fast_ema_pinned; + let new_fast = if fast_prev <= 0.0 { + observed_grad_norm + } else { + (1.0 - FAST_ALPHA) * fast_prev + FAST_ALPHA * observed_grad_norm + }; + *self.grad_norm_fast_ema_pinned = new_fast; + + let slow_prev = *self.grad_norm_slow_ema_pinned; + let new_slow = if slow_prev <= 0.0 { + observed_grad_norm + } else { + (1.0 - SLOW_ALPHA) * slow_prev + SLOW_ALPHA * observed_grad_norm + }; + *self.grad_norm_slow_ema_pinned = new_slow; + } + self.grad_norm_emas_step_count = self.grad_norm_emas_step_count.saturating_add(1); } + /// Plan C Phase 2 follow-up K (2026-04-29): launch + /// `fold_warmup_factor_update` — single-thread single-block ISV + /// producer for ISV[FOLD_WARMUP_FACTOR_INDEX=130]. Reads the fast/slow + /// grad-norm EMA mapped-pinned scalars (filled by + /// `update_adaptive_clip` per training step) plus a host-passed step + /// counter (gates the bootstrap `factor=1` window) and writes + /// `clamp(fast/slow, 0, 1)`. Cold-path-cadence — invoked once per + /// training step alongside the other ISV producers + /// (`launch_h_s2_rms_ema`, `launch_iqn_quantile_ema`, etc.). + /// + /// Per `pearl_cold_path_no_exception_to_gpu_drives.md`: even a cold- + /// path scalar arithmetic stays on GPU when its EMA inputs already + /// live in mapped-pinned host memory the device can read directly. + pub fn launch_fold_warmup_factor(&self) -> Result<(), MLError> { + // Same invariant rationale as `launch_q_drift_rate_ema` / + // `launch_h_s2_rms_ema`: loud failure (debug_assert + kernel-launch + // error) preferred over silent skip. + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_fold_warmup_factor: isv_signals_dev_ptr must be allocated by constructor"); + let isv_dev_ptr = self.isv_signals_dev_ptr; + let fast_dev_ptr = self.grad_norm_fast_ema_dev_ptr; + let slow_dev_ptr = self.grad_norm_slow_ema_dev_ptr; + let warmup_idx = FOLD_WARMUP_FACTOR_INDEX as i32; + // `WARMUP_FACTOR_MIN_STEPS` is a numerical-stability bound (Invariant 1 + // carve-out per `feedback_isv_for_adaptive_bounds.md`) — until the + // EMAs have absorbed at least this many samples the slow EMA is + // effectively cold-start (within ~10× of the first observation), + // making the ratio meaningless. 200 ≈ 2τ for the fast EMA so + // both EMAs have meaningful state. Saturating-cast to i32 caps at + // i32::MAX, which is fine: once we exceed the cap the gate is + // permanently passed. + const WARMUP_FACTOR_MIN_STEPS: i32 = 200; + let steps_observed_i32: i32 = self.grad_norm_emas_step_count.min(i32::MAX as u64) as i32; + unsafe { + self.stream.launch_builder(&self.fold_warmup_factor_kernel) + .arg(&isv_dev_ptr) + .arg(&fast_dev_ptr) + .arg(&slow_dev_ptr) + .arg(&warmup_idx) + .arg(&WARMUP_FACTOR_MIN_STEPS) + .arg(&steps_observed_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("fold_warmup_factor_update: {e}")))?; + } + Ok(()) + } + + /// Plan C Phase 2 follow-up K (2026-04-29): reset the FAST grad-norm + /// EMA at fold boundary. The SLOW EMA persists — it tracks the + /// cross-fold steady-state grad scale that the warmup factor's + /// denominator compares against. Resetting only the fast EMA causes + /// the ratio to start at 0 (fast=0, slow≈baseline) → factor = 0 → + /// heavy damping for the new fold's first steps. Step counter is NOT + /// reset — once the warmup gate has been crossed (typical after the + /// first fold), subsequent folds skip the cold-start guard and rely + /// purely on the EMA ratio. + pub fn reset_fast_grad_norm_ema(&mut self) { + unsafe { *self.grad_norm_fast_ema_pinned = 0.0_f32; } + } + + /// Read-only accessor for diagnostics / HEALTH_DIAG mirror. + pub fn grad_norm_fast_ema_value(&self) -> f32 { unsafe { *self.grad_norm_fast_ema_pinned } } + /// Read-only accessor for diagnostics / HEALTH_DIAG mirror. + pub fn grad_norm_slow_ema_value(&self) -> f32 { unsafe { *self.grad_norm_slow_ema_pinned } } + pub fn adaptive_clip_value(&self) -> f32 { unsafe { *self.adaptive_clip_pinned } } pub fn grad_norm_ema_value(&self) -> f32 { self.grad_norm_ema } + /// Plan C Phase 2 follow-up K (2026-04-29): write a warmup-factor-scaled + /// clip value directly to the pinned device-mapped slot. Mirrors `set_lr`'s + /// zero-graph-recapture contract (the kernel reads `*adaptive_clip_pinned` + /// at replay, so a mapped-pinned write is GPU-visible immediately on the + /// next kernel launch). Used by the per-step warmup-factor consumer in + /// `training_loop.rs` to compose `clip_eff = clip_base × (0.1 + 0.9 × factor)` + /// on top of the EMA-derived baseline that `update_adaptive_clip` just + /// wrote. + pub fn set_active_clip(&mut self, clip: f32) { + if clip.is_finite() && clip > 0.0 { + unsafe { *self.adaptive_clip_pinned = clip; } + } + } + /// Write the EMA tau value directly to pinned device-mapped memory. /// GPU reads the updated value on next kernel launch — no HtoD copy needed. pub(crate) fn set_tau_value(&mut self, tau: f32) { diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index da7c7679f..0fc059c31 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -2954,6 +2954,17 @@ impl FusedTrainingCtx { &self.trainer } + /// Plan C Phase 2 follow-up K (2026-04-29): mut accessor for the + /// fold-boundary fast-grad-norm-EMA reset. The reset is host-side + /// (mapped-pinned scalar write) so it does NOT touch the captured + /// CUDA graphs — same caveat as `set_lr` / `set_aux_weight`. Use + /// only for plain mapped-pinned scalar mutations and other host-side + /// state on the trainer; do NOT use this to launch graph-affecting + /// kernels (those must go through properly-ordered submission paths). + pub(crate) fn trainer_mut(&mut self) -> &mut GpuDqnTrainer { + &mut self.trainer + } + /// Raw pointer to trainer's states buffer. pub(crate) fn trainer_states_buf_ptr(&self) -> u64 { self.trainer.states_buf_ptr() @@ -3252,6 +3263,15 @@ impl FusedTrainingCtx { } pub(crate) fn update_adaptive_clip(&mut self, grad_norm: f32) { self.trainer.update_adaptive_clip(grad_norm); } pub(crate) fn adaptive_clip_value(&self) -> f32 { self.trainer.adaptive_clip_value() } + /// Plan C Phase 2 follow-up K (2026-04-29): override the active + /// adaptive-clip pinned slot with a warmup-factor-scaled value. Mirrors + /// `set_lr`'s zero-graph-recapture contract — the kernel reads + /// `*adaptive_clip_pinned` at replay time, so writing the pinned slot + /// takes effect on the next captured graph replay without invalidation. + /// The companion `update_adaptive_clip` sets the EMA-derived baseline + /// (`grad_norm_ema × 2`, floored at 1.0); this setter then scales that + /// baseline by the warmup-factor consumer's `0.1 + 0.9 × factor`. + pub(crate) fn set_active_clip(&mut self, clip: f32) { self.trainer.set_active_clip(clip); } /// Raw pointer to plan_params_buf [B, 6] for trade plan integration. pub(crate) fn plan_params_buf_ptr(&self) -> u64 { self.trainer.plan_params_buf_ptr() diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 7cf6e2a73..f2affdc1f 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -424,6 +424,33 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[Q_DRIFT_RATE_INDEX=129] — adaptive Polyak-tau drift signal ∈ [0, 4]; GPU q_drift_rate_ema_update kernel fills (Plan C A.2). Cold-start 0.0 (no-op dampening factor) reapplied at fold boundary; consumer tau_update_kernel.cu multiplies tau_eff by 1/(1+ISV[129]) so dampening is monotone — tau only ever decreases under drift, healthy runs unaffected", }, + // Plan C Phase 2 follow-up K (2026-04-29): adaptive fold-boundary + // warmup signal (slot 130). Reset to 0.0 at fold boundary so the + // new fold's first steps see heavy lr/clip damping; the kernel + // raises the factor toward 1.0 as the fast grad-norm EMA catches + // up to the slow steady-state EMA (which persists across folds). + // Companion to A.1+A.2+A.3 fold-boundary state-reset family. + // Per `pearl_blend_formulas_must_have_permanent_floor.md`, the + // damping floor (lr_floor=0.05, clip_floor=0.1) is a permanent + // minimum — even at factor=0 the consumers never reach zero lr/clip. + RegistryEntry { + name: "isv_fold_warmup_factor", + category: ResetCategory::FoldReset, + description: "ISV[FOLD_WARMUP_FACTOR_INDEX=130] — adaptive fold-boundary warmup factor ∈ [0, 1]; GPU fold_warmup_factor_update kernel fills (Plan C K). Cold-start 0.0 (heavy damping at fold start) reapplied at fold boundary; consumers in training_loop.rs derive lr_eff = lr_base × max(0.05, factor) and clip_eff = clip_base × (0.1 + 0.9 × factor) — both monotone (only dampen). Companion FoldReset entry `isv_grad_norm_fast_ema` resets the kernel's numerator EMA so the ratio starts at 0", + }, + // Plan C Phase 2 follow-up K (2026-04-29): fast grad-norm EMA + // mapped-pinned scalar feeding `fold_warmup_factor_update`'s + // numerator. Companion to `isv_fold_warmup_factor` — both must + // reset together so the ratio (fast/slow) starts at 0 → factor=0 + // → heavy damping. The slow EMA does NOT reset (cross-fold + // steady-state baseline). Per `feedback_no_partial_refactor.md`, + // both halves of the warmup-factor input contract are reset in + // lockstep. + RegistryEntry { + name: "isv_grad_norm_fast_ema", + category: ResetCategory::FoldReset, + description: "GpuDqnTrainer.grad_norm_fast_ema_pinned — fast (α=0.1, ~10-step horizon) EMA of `gr.raw_grad_norm`, host-side mapped-pinned scalar fed by `update_adaptive_clip` and consumed by `fold_warmup_factor_update` (Plan C K). Reset to 0.0 at fold boundary so the warmup-factor ratio starts at 0 → consumer-derived lr_eff/clip_eff start at their permanent-floor values; recovers to slow-EMA baseline as gradients stabilise. Slow EMA companion (grad_norm_slow_ema_pinned) deliberately persists across folds — it tracks the cross-fold steady-state grad scale that the warmup factor compares against.", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index ec00b73ea..31a201b97 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -527,6 +527,10 @@ impl DQNTrainer { let initial_batch_size = hyperparams.batch_size; let base_tau: f64 = f64::from(hyperparams.tau); let initial_epsilon = hyperparams.epsilon_start; + // Plan C Phase 2 follow-up K (2026-04-29): cosine LR baseline for the + // per-step warmup-factor consumer; captured here so the struct + // literal below doesn't reborrow `hyperparams` after move. + let initial_cosine_effective_lr_base: f32 = hyperparams.learning_rate as f32; // Multi-GPU: auto-detect if multiple CUDA devices are available let multi_gpu = crate::cuda_pipeline::multi_gpu::MultiGpuConfig::detect() @@ -648,6 +652,13 @@ impl DQNTrainer { // ≥ 2 entries (median/MAD undefined on a 1-element window). Capacity // matches `Q_DRIFT_WINDOW_SIZE` in `training_loop.rs`. q_mean_window: VecDeque::with_capacity(8), + // Plan C Phase 2 follow-up K (2026-04-29): cosine LR baseline + // updated by the cosine schedule block in `training_loop.rs` + // and consumed per-step by the warmup-factor LR consumer. + // Initialised to the configured base LR so the first epoch's + // pre-cosine `set_lr` matches the constructor-injected + // `config.lr` (no transient mismatch). + cosine_effective_lr_base: initial_cosine_effective_lr_base, // WAVE 24 (Agent 17): Initialize patience-based early stopping // Use gradient_collapse_patience from hyperparams for consistency diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 379c088a5..7c2341c1e 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -503,6 +503,19 @@ pub struct DQNTrainer { /// distributions, the same rationale as `prev_epoch_q_mean`'s reset. pub(crate) q_mean_window: VecDeque, + /// Plan C Phase 2 follow-up K (2026-04-29): the most recent + /// cosine-scheduled "base" LR (re-set at epoch boundary by the cosine + /// LR block in `training_loop.rs`). Per-step the warmup-factor consumer + /// derives `lr_eff = cosine_effective_lr_base × max(MIN_WARMUP_LR_FRAC, + /// ISV[FOLD_WARMUP_FACTOR_INDEX])` and writes it via `set_lr`. Storing + /// the cosine value here decouples per-epoch scheduling from per-step + /// warmup damping (without this the per-step `set_lr` would clobber + /// the cosine-scheduled value, and the next epoch's cosine update would + /// fight the warmup factor at every step). Initialised to + /// `hyperparams.learning_rate` and updated alongside the existing + /// `set_lr(effective_lr)` site. + pub(crate) cosine_effective_lr_base: f32, + /// WAVE 24 (Agent 17): Patience-based early stopping for anti-overfitting pub(crate) early_stopping: super::early_stopping::EarlyStopping, diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 7281fd4da..970fdfc8e 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -375,6 +375,16 @@ impl DQNTrainer { }; let base_lr = self.hyperparams.learning_rate; let effective_lr = (base_lr * lr_mult) as f32; + // Plan C Phase 2 follow-up K (2026-04-29): record the + // cosine-effective LR as the baseline that the per-step + // warmup-factor consumer composes with (it derives + // lr_eff = effective_lr × max(0.05, ISV[130]) and writes + // via set_lr each step). The set_lr call here remains + // for the case where the very first step fires before + // the warmup-factor producer has run (factor=0 cold- + // start, but consumer applies floor=0.05 → lr_eff = + // effective_lr × 0.05, dampened as intended). + self.cosine_effective_lr_base = effective_lr; if let Some(ref mut fused) = self.fused_ctx { fused.set_lr(effective_lr); if t_local == 0 && epoch > 0 { @@ -1990,7 +2000,56 @@ impl DQNTrainer { { self.grad_clip_kicked_this_epoch = true; } + // Update the EMA-based adaptive clip + the fast/slow + // grad-norm EMAs that drive Plan C K's fold-warmup + // factor (single observation, three consumers). fused.update_adaptive_clip(gr.raw_grad_norm); + + // Plan C Phase 2 follow-up K (2026-04-29): apply the + // adaptive fold-boundary warmup factor to BOTH lr_eff + // and clip_eff. ISV[FOLD_WARMUP_FACTOR_INDEX] is + // produced once per step by `launch_fold_warmup_factor` + // alongside the other ISV producers (see the + // post-step ISV producer block lower in this loop); + // by the time we land here on the NEXT step the slot + // already holds the most recent valid factor. + // + // Numerical-stability bounds (Invariant 1 carve-outs + // per `feedback_isv_for_adaptive_bounds.md`): + // MIN_WARMUP_LR_FRAC = 0.05 — lr never below 5% + // MIN_CLIP_FRAC = 0.1 — clip never below 10% + // Both consumers are MONOTONE (factor=1 → no + // dampening; factor<1 → dampening; factor never + // excites lr/clip past baseline). Healthy runs + // (steady-state factor=1) are unaffected. + const MIN_WARMUP_LR_FRAC: f32 = 0.05; + const MIN_CLIP_FRAC: f32 = 0.1; + let warmup_factor = fused.read_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::FOLD_WARMUP_FACTOR_INDEX, + ).clamp(0.0, 1.0); + // LR consumer: lr_eff = cosine_effective_lr × max(0.05, factor). + // Compose with the cosine schedule's per-epoch baseline + // (stored in `self.cosine_effective_lr_base` by the + // cosine LR block at epoch start) so per-step warmup + // damping multiplies the cosine schedule rather than + // overriding it. + let lr_base = self.cosine_effective_lr_base; + let lr_warmup_scale = MIN_WARMUP_LR_FRAC.max(warmup_factor); + let lr_eff = lr_base * lr_warmup_scale; + fused.set_lr(lr_eff); + // Clip consumer: clip_eff = clip_base × (0.1 + 0.9 × factor). + // `update_adaptive_clip` just wrote `clip_base = ema × 2` + // to the pinned slot above; rescale that value + // in-place by the warmup multiplier. The pinned + // write through `set_active_clip` is GPU-visible + // immediately on the next kernel read. + let clip_base = fused.adaptive_clip_value(); + if clip_base.is_finite() && clip_base > 0.0 { + let clip_warmup_scale = + MIN_CLIP_FRAC + (1.0 - MIN_CLIP_FRAC) * warmup_factor; + let clip_eff = clip_base * clip_warmup_scale; + fused.set_active_clip(clip_eff); + } } } } // end guard frequency check @@ -3053,6 +3112,20 @@ impl DQNTrainer { } } + // Plan C Phase 2 follow-up K (2026-04-29): per-step adaptive + // fold-boundary warmup factor producer. Reads the fast/slow + // grad-norm EMA mapped-pinned scalars (filled by + // `update_adaptive_clip` in the per-step guard block above) + // and writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1). + // Two consumers below derive lr_eff (via set_lr) and + // clip_eff (via update_adaptive_clip's pinned slot). Same + // launch cadence as h_s2_rms_ema — once per step. + if let Some(ref fused) = self.fused_ctx { + if let Err(e) = fused.trainer().launch_fold_warmup_factor() { + tracing::warn!("Plan C K fold_warmup_factor launch failed: {e}"); + } + } + // Plan 4 Task 3 (E.3): per-step IQN multi-quantile diagnostic // EMAs into ISV[99..103). 4-block kernel reads the IQN online // forward's `save_q_online [TBA, B*Q]` (populated by @@ -5293,6 +5366,33 @@ impl DQNTrainer { ); } } + "isv_fold_warmup_factor" => { + // Plan C Phase 2 follow-up K (2026-04-29): adaptive + // fold-boundary warmup factor. Reset to 0.0 at fold + // boundary so the new fold begins under heavy lr/clip + // damping (consumer floors: lr_eff = 0.05 × lr_base, + // clip_eff = 0.1 × clip_base). Companion FoldReset entry + // `isv_grad_norm_fast_ema` zeroes the kernel's numerator + // EMA in lockstep so the ratio starts at 0; the slow EMA + // persists across folds. + if let Some(ref fused) = self.fused_ctx { + fused.trainer().write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::FOLD_WARMUP_FACTOR_INDEX, + 0.0_f32, + ); + } + } + "isv_grad_norm_fast_ema" => { + // Plan C Phase 2 follow-up K (2026-04-29): fast grad-norm + // EMA mapped-pinned scalar (numerator of the + // fold-warmup-factor ratio). Companion to + // `isv_fold_warmup_factor` — both reset in lockstep per + // `feedback_no_partial_refactor.md` so the ratio starts + // at 0 (fast=0 over slow≈baseline). Slow EMA persists. + if let Some(ref mut fused) = self.fused_ctx { + fused.trainer_mut().reset_fast_grad_norm_ema(); + } + } "isv_gamma_dir_eff" | "isv_gamma_mag_eff" | "isv_gamma_ord_eff" | "isv_gamma_urg_eff" => { // D.2 per-branch gamma slots. Reset to 0.0 at fold boundary; // per_branch_gamma_update GPU kernel re-populates on next epoch-boundary launch. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 71047856a..dc46a521f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2200,3 +2200,27 @@ Plan C Phase 2 follow-up F (2026-04-29): wire `q_mag_bin_means_reduce` and `q_di Plan C Phase 2 follow-up (2026-04-29): bonus reward optimism-coupling break in `experience_kernels.cu`. The C.4 timing bonus (~line 2521) and D.4b regime penalty (~line 2581) both used the trade-cumulative `|final_pnl|` / `|reward|` as their unbounded multiplicand. By the time bonus shaping runs in `experience_env_step`, `reward` is the Layer 2 vol-normalised return PLUS Layer 4-9 credits PLUS the C.4 timing bonus PLUS the D.4a persistence bonus — i.e. it is itself a function of all prior shaping inflation. Across trades this compounds into a multi-trade optimism feedback loop: bonus inflates → Q-target inflates (Bellman bootstraps off shaped reward) → next trade's reward inflates → bonus inflates further. Researcher reproduction (a25f669e9df953174, 2026-04-29) found the a52d99613 baseline ALSO hits bonus EMA=235 by F2 ep2 — the pathology is structural and pre-dates Plan C Phase 2. Per `pearl_one_unbounded_signal_per_reward.md` the rule is exactly ONE unbounded multiplicand per reward term, and per `feedback_isv_for_adaptive_bounds.md` adaptive bounds live in the ISV signal bus. **Surgical fix**: replace `|final_pnl|` (C.4) and `|reward|` (D.4b) with `min(|x|/ISV[Q_DIR_ABS_REF_INDEX=21], 1) × ISV[21]` — the multiplicand is now the direction-branch Q-scale EMA (already produced gradient-decoupled by `q_stats_kernel.cu` since Plan 1) rather than the trade-cumulative shaping output. Cap is adaptive (matches Q magnitude as it evolves) and breaks cross-trade compounding; |reward| can still drive the cap to its max (= ISV[21]) so directional information is preserved, just bounded. **Kernel signature**: `experience_env_step` gains a final `int q_dir_abs_ref_idx` parameter (after the existing `trade_target_rate_idx`); single Rust call site `gpu_experience_collector.rs:3800` passes `Q_DIR_ABS_REF_INDEX as i32`. `experience_action_select` not affected (it never read `|reward|`/`|final_pnl|` as a multiplicand). **Other shaping sites NOT touched** because they were already bounded — D.4a persist (multiplicand `drawdown_depth ≤ ~0.1` gates the term, `tanh(reward/dd)` is bounded), B.2 novelty (vol_proxy ≤ 0.01 gates, kl_amp ≤ 1.0), D.4c stable (vol_proxy gates + bounded stability/conviction). Per `feedback_no_partial_refactor.md` consistency: both fixed sites use the SAME ISV[21] bound and the SAME `(unit, capped)` two-step formula — symmetric migration. Predicted impact: bonus EMA O(100-256) → O(0.05-2.0); rc[5] → Bellman-target optimism loop broken; Plan C smoke F0 ep2 Q-drift kill should no longer fire on this pathway; a52d99613 baseline gets the same fix automatically (validates pre-existing pathology, not Plan-C-specific). Layout fingerprint unchanged — kernel signature change but no ISV slot add and no param-tensor change. Plan C T11 follow-up A.3 (2026-04-29): added `DQN::reset_for_fold` zeroing `gradient_collapse_counter` + `training_steps`. Wired from `DQNTrainer::reset_for_fold` next to the A.1 `prev_epoch_q_mean` reset. Discovered when `smoke-test-vh9bj`'s fold-0 succeeded (F+H fix worked) but fold-1 failed immediately at epoch 1 with "Gradient collapse detected for 5 consecutive" — the counter had carried over from fold 0's near-zero-grad late-epoch steps, plus the `past_warmup` gate was always-true in fold 1+ because `training_steps` accumulated. Resetting both restores per-fold warmup semantics. Same fold-boundary state-reset gap pattern as A.1. + +Plan C Phase 2 follow-up K (2026-04-29): combined Adam shrink-and-perturb + adaptive fold-boundary warmup factor. Closes the F1 ep1 catastrophic-overshoot gap exposed by smoke-test-s9h4h (F0 successful with Best Sharpe = 36.03; F1 ep1 grad_norm = 355,009 — 5 OoM larger than F0 steady-state ~10 — leading to NaN propagation and grad-clamp-to-zero). + +* **Adam shrink-and-perturb** in `GpuDqnTrainer::reset_adam_state` replaces the prior full-zero reset of `m_buf` / `v_buf` with multiplicative shrink (`m *= 0.1`, `v *= 0.01`). Preserves directional information (sign+ratio across tensors) while damping magnitude; composes with the existing param shrink-and-perturb (`alpha=0.8` in `FusedTrainingCtx::reset_for_fold`). Shrink launches use the existing `dqn_scale_f32_kernel` (loaded as `scale_f32_ungraphed` — same module / same launch path as `scale_adam_momentum`). `t_pinned` (Adam step counter) IS still zeroed so the bias correction `1 - β^t` restarts properly. Per `feedback_isv_for_adaptive_bounds.md` Invariant 1 carve-out: shrink factors `0.1` / `0.01` are STRUCTURAL ("preserve direction / lose magnitude history") not tuned constants. Root cause for the F1 overshoot: `m=0, v=0` → first Adam step is `lr × m̂ / (sqrt(v̂) + ε) ≈ lr × g / ε` (a 6 OoM amplification when `ε = 1e-8`). + +* **New CUDA kernel** `crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu` (single-block single-thread cold-path producer mirroring `q_drift_rate_ema_kernel.cu` / `moe_lambda_eff_kernel.cu` shape). Reads two grad-norm EMAs (fast α=0.1 ~10-step horizon, slow α=0.001 ~1000-step horizon) plus a host-passed step counter; writes `clamp(fast/slow, 0, 1)` to `ISV[FOLD_WARMUP_FACTOR_INDEX=130]`. Bootstrap branches for the cold-start window (`steps_observed < WARMUP_FACTOR_MIN_STEPS=200`) and the slow-EMA-near-zero edge case both emit `factor = 1.0` (no damping). Registered in `crates/ml/build.rs` kernel list (count grew by 1). + +* **New ISV slot** `FOLD_WARMUP_FACTOR_INDEX = 130` (tail-appended after `Q_DRIFT_RATE_INDEX = 129`, raising `ISV_TOTAL_DIM` 130→131). Cold-start 0.0 (constructor `*sig_ptr.add(FOLD_WARMUP_FACTOR_INDEX) = 0.0_f32` is satisfied by the existing zero-init of the mapped-pinned ISV bus); FoldReset 0.0 via new `isv_fold_warmup_factor` registry entry in `state_reset_registry.rs` + dispatch arm in `training_loop.rs::reset_named_state`. Companion FoldReset entry `isv_grad_norm_fast_ema` resets the kernel's NUMERATOR (host-side mapped-pinned scalar) so the ratio starts at 0 (fast=0 over slow≈baseline) → factor=0 → heavy damping for the new fold's first steps. Slow EMA `isv_grad_norm_slow_ema` deliberately persists across folds (cross-fold steady-state baseline). Layout fingerprint shifts (one slot added + ISV_TOTAL_DIM bump) — checkpoint-incompatible per `feedback_no_legacy_aliases.md`, expected for a real architecture change. + +* **Two new mapped-pinned scalars** on `GpuDqnTrainer`: `grad_norm_fast_ema_pinned` (α=0.1) and `grad_norm_slow_ema_pinned` (α=0.001). Both updated by `update_adaptive_clip` (the same `gr.raw_grad_norm` observation source that already feeds the legacy `grad_norm_ema` field driving the existing adaptive clip). Step counter `grad_norm_emas_step_count` gates the kernel's bootstrap window. Mapped-pinned via `cuMemHostAlloc DEVICEMAP` per `feedback_no_htod_htoh_only_mapped_pinned.md`. `Drop` impl frees both alongside `adaptive_clip_pinned`. + +* **Rust orchestrator wrapper** `GpuDqnTrainer::launch_fold_warmup_factor()` in `gpu_dqn_trainer.rs` — same pattern as `launch_q_drift_rate_ema` / `launch_h_s2_rms_ema` (debug_assert non-zero `isv_signals_dev_ptr`, single-thread launch config, kernel-launch error mapped to `MLError::ModelError`). + +* **Per-step producer launch** in `training_loop.rs` alongside `launch_h_s2_rms_ema` — fires once per training step, after `update_adaptive_clip` has fed the fast/slow EMAs. Cold-path cadence; no atomicAdd; no DtoH. + +* **LR consumer wire-up** in `training_loop.rs` per-step block: read `ISV[FOLD_WARMUP_FACTOR_INDEX]`, derive `lr_eff = cosine_effective_lr_base × max(MIN_WARMUP_LR_FRAC=0.05, factor)`, write via `set_lr`. New `cosine_effective_lr_base: f32` field on `DQNTrainer` records the cosine-scheduled per-epoch baseline so the per-step warmup damping multiplies the cosine schedule rather than overriding it. Floor `0.05` is a numerical-stability bound (Invariant 1 carve-out per `feedback_isv_for_adaptive_bounds.md`). + +* **Clip consumer wire-up** in same block: after `update_adaptive_clip` writes `clip_base = grad_norm_ema × 2` (floored at 1.0) to the pinned slot, scale by `clip_warmup_scale = MIN_CLIP_FRAC=0.1 + (1 - MIN_CLIP_FRAC) × factor` and write the result via the new `FusedTrainingCtx::set_active_clip` setter (forwarded to `GpuDqnTrainer::set_active_clip`, which writes `*adaptive_clip_pinned` directly — same zero-graph-recapture contract as `set_lr`). Floor `0.1` is a numerical-stability bound. + +* **Both consumers are MONOTONE** (factor ∈ [0, 1] structurally, so `lr_eff ≤ lr_base` and `clip_eff ≤ clip_base` always). Healthy runs (steady-state factor=1) are unaffected — `lr_eff = lr_base × max(0.05, 1) = lr_base` and `clip_eff = clip_base × (0.1 + 0.9) = clip_base`. Per `pearl_blend_formulas_must_have_permanent_floor.md`: `MIN_WARMUP_LR_FRAC` and `MIN_CLIP_FRAC` are permanent floors — even at factor=0 the consumers never reach zero lr/clip. + +Predicted impact on Plan C smoke fold 1: factor starts at 0 → `lr_eff = 0.05 × lr_base`, `clip_eff = 0.1 × clip_base ≈ 1.0` (vs `clip_base = 10`). The 355,009-magnitude transient grad gets clipped tightly to ~1, Adam state shrunk (m × 0.1, v × 0.01) instead of zeroed → first step update is bounded by `0.05 × lr × m̂/(sqrt(v̂) + ε)` on a non-degenerate denominator. Over ~50 steps the fast EMA catches up to the slow steady-state baseline → factor → 1 → full lr/clip restore. Steady-state (mid-fold) behavior unchanged. + +Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold-start) + reset (0.0 fold boundary, fast EMA) + observability (HEALTH_DIAG via slot 130 mirror, plus `grad_norm_fast_ema_value` / `grad_norm_slow_ema_value` accessors on the trainer). Per `pearl_cold_path_no_exception_to_gpu_drives.md` — even a cold-path scalar arithmetic stays on GPU when its inputs already live in mapped-pinned host memory the device can read directly (no DtoH for the EMAs). Per `feedback_adaptive_not_tuned.md` — both lr_eff and clip_eff are now ISV-driven, not constant. Per `feedback_isv_for_adaptive_bounds.md` — the warmup factor IS the bound; consumers read ISV[130] at runtime and compose with their respective architectural floors. Per `feedback_no_partial_refactor.md` — both halves of the warmup-factor input contract (factor slot + fast EMA companion) reset together; both consumers (lr + clip) wire together; producer launch + consumer reads land in the same commit.