gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(dqn): Plan C K — Adam shrink-and-perturb + adaptive fold-warmup ISV

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-29 20:24:36 +02:00
parent ef243e771d
commit 4ef1d8ebb7
9 changed files with 611 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
crates/ml/build.rs
View File

@@ -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`

88
crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu Normal file
View File

@@ -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 */
}

324
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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::<f32>(), 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::<f32>(), 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) {

20
crates/ml/src/trainers/dqn/fused_training.rs
View File

@@ -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()

27
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -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 }
}

11
crates/ml/src/trainers/dqn/trainer/constructor.rs
View File

@@ -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

13
crates/ml/src/trainers/dqn/trainer/mod.rs
View File

@@ -503,6 +503,19 @@ pub struct DQNTrainer {
/// distributions, the same rationale as `prev_epoch_q_mean`'s reset.
pub(crate) q_mean_window: VecDeque<f64>,
/// 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,

100
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -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.

24
docs/dqn-wire-up-audit.md
View File

@@ -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.