fix(dqn): SP3 close-out — Mech 9 post-Adam weight clamp + Mech 8 revert
Coordinated close-out of SP3 Q-learning numerical-stability per
feedback_no_partial_refactor.
REVERT Mech 8 (slow_ema fold-boundary reset). smoke-test-rxhjh on b8a7ac6f7
showed F1 NaN at step 2300 — WORSE than the 3720 ceiling with Mech 6 alone.
Persistent slow_ema across folds was providing unintentional F1 protection
(anchored at F0's smaller grad scale → tighter Mech 6 upper_bound during F1
ramp-up). Resetting loosened the clip and accelerated Adam saturation.
Removed: reset_grad_norm_slow_ema() method + reset_for_fold call site.
Kept: grad_norm_slow_ema_pinned field (consumed by Mech 6).
Kept: Mech 6 multiplier at 100× (already restored from the 5× experiment
in the Mech 8 commit; correct steady-state value).
ADD Mech 9 (post-Adam ISV-driven weight clamp). Root cause being targeted:
cuBLAS sgemm f32 accumulator overflow at slots 26 (iqn_trunk_m) + 32
(bn_d_concat_buf) at F1 step ~3540 with INPUTS clean (slots 27, 35
unflagged). The matmul output saturates because the WEIGHT matrices have
drifted to extreme-but-finite magnitudes via Adam updates over thousands
of steps. Mech 1 (gradient clamp) and Mech 6 (grad-norm clip) bound
gradients, not parameters — neither prevents this drift.
Mech 9 clamps |p_val| ≤ 100 × Q_ABS_REF.max(1.0) inside dqn_adam_update_kernel
after the L1 proximal step. ISV-driven (slot 16, ε on multiplier per SP1
pearl). 1 OOM below slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so
the regression sentinel retains 10× firing headroom — symmetric with
Mech 1's clamp/diagnostic ratio.
Implementation:
- dqn_adam_update_kernel: new trailing arg `weight_clamp_max_abs`.
Clamp `p_val = fminf(fmaxf(p_val, -bound), bound)` after L1 step.
- All Adam launch sites in gpu_dqn_trainer.rs: compute bound from
read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0) × 100.0 host-side, pass
as final arg.
- No new ISV slots. No new kernels. No graph topology change.
Validation: deferred to one L40S smoke at the new HEAD. Expected:
- F1 trains past step 3720 (Mech-6-only ceiling) — Mech 9 closes SP3.
- F1 still NaNs at similar steps — close SP3 honestly with documented
residual pointing at SP4/SP5.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -977,6 +977,13 @@ impl DecisionTransformer {
|
||||
let wd_mask_ptr = self.wd_mask.dev_ptr;
|
||||
let l1_end_dt: i32 = 0; // no L1 for DT
|
||||
let l1_lambda_dt: f32 = 0.0; // disabled
|
||||
// SP3 Mech 9: post-Adam |p_val| clamp is keyed on the main DQN
|
||||
// trainer's ISV[Q_ABS_REF_INDEX=16]; DecisionTransformer is a
|
||||
// separate offline-RL trainer with no ISV bus access and is
|
||||
// outside the SP3 cuBLAS-overflow scope. Pass 0.0 to disable
|
||||
// the clamp here — the kernel branch is `if (bound > 0.0f)`,
|
||||
// so Adam behaves exactly as before for DT.
|
||||
let weight_clamp_max_abs_dt: f32 = 0.0_f32;
|
||||
unsafe {
|
||||
stream.launch_builder(&adam.adam_update)
|
||||
.arg(¶ms_ptr_mut)
|
||||
@@ -991,6 +998,7 @@ impl DecisionTransformer {
|
||||
.arg(&wd_mask_ptr)
|
||||
.arg(&l1_end_dt)
|
||||
.arg(&l1_lambda_dt)
|
||||
.arg(&weight_clamp_max_abs_dt) // SP3 Mech 9: disabled (no ISV in DT scope)
|
||||
.launch(lc)
|
||||
.map_err(|e| MLError::ModelError(format!("DT adam: {e}")))?;
|
||||
}
|
||||
|
||||
@@ -96,6 +96,19 @@ extern "C" __global__ void dqn_grad_norm_finalize(
|
||||
* applies gradient clipping, then AdamW parameter update.
|
||||
* Trivially parallel: 1 thread per parameter element.
|
||||
*
|
||||
* SP3 Mech 9 (2026-04-30): post-Adam ISV-driven weight clamp. After the
|
||||
* Adam + AdamW + L1-proximal step writes p_val, |p_val| is clamped to
|
||||
* `weight_clamp_max_abs` (computed host-side from
|
||||
* ISV[Q_ABS_REF_INDEX=16].max(1.0) × 100 — see `update_adaptive_clip` and
|
||||
* the Adam launch sites in `gpu_dqn_trainer.rs`). Targets cuBLAS sgemm
|
||||
* f32-accumulator overflow at `iqn_trunk_m` / `bn_d_concat_buf` GEMMs
|
||||
* with finite inputs but extreme-magnitude weights (root cause of the
|
||||
* F1 step-3540 NaN observed at slots 26+32 with slots 27+35 unflagged).
|
||||
* Bound is one OOM below the slot 44-45 diagnostic threshold
|
||||
* (1e3 × Q_ABS_REF) so the regression sentinel keeps 10× firing
|
||||
* headroom. Disabled when `weight_clamp_max_abs == 0.0` (debug
|
||||
* fast-path; production never sets 0).
|
||||
*
|
||||
* Launch config: grid=(ceil(TOTAL_PARAMS/256), 1, 1), block=(256, 1, 1).
|
||||
* ══════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
@@ -115,7 +128,8 @@ extern "C" __global__ void dqn_adam_update_kernel(
|
||||
int total_params,
|
||||
const float* __restrict__ weight_decay_mask, /* [TOTAL_PARAMS] 1.0=decay, 0.0=no decay */
|
||||
int l1_end, /* param index where L1 stops (end of w_s1) */
|
||||
float l1_lambda /* L1 penalty (1e-3). 0=disabled */
|
||||
float l1_lambda, /* L1 penalty (1e-3). 0=disabled */
|
||||
float weight_clamp_max_abs /* SP3 Mech 9: ISV-driven |p_val| bound (100 * Q_ABS_REF.max(1.0)). 0=disabled. */
|
||||
) {
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (idx >= total_params) return;
|
||||
@@ -160,6 +174,16 @@ extern "C" __global__ void dqn_adam_update_kernel(
|
||||
p_val = sign_val * fmaxf(fabsf(p_val) - lr * l1_lambda, 0.0f);
|
||||
}
|
||||
|
||||
/* SP3 Mech 9: post-Adam weight clamp. ISV-driven bound
|
||||
* (100 × Q_ABS_REF.max(1.0)) computed host-side and passed in. Targets
|
||||
* cuBLAS sgemm f32 accumulator overflow at slots 26+32 — bounds weights
|
||||
* 1 OOM below slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so the
|
||||
* diagnostic retains regression-sentinel headroom. Disabled when
|
||||
* bound==0 (debug fast-path; production never sets 0). */
|
||||
if (weight_clamp_max_abs > 0.0f) {
|
||||
p_val = fminf(fmaxf(p_val, -weight_clamp_max_abs), weight_clamp_max_abs);
|
||||
}
|
||||
|
||||
params[idx] = p_val;
|
||||
}
|
||||
|
||||
|
||||
@@ -3855,25 +3855,6 @@ impl GpuDqnTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP3 Mech 8: hard-reset grad_norm_slow_ema at fold boundary. Without
|
||||
/// this reset, the α=0.001 EMA persists across folds with the prior
|
||||
/// fold's value (half-life ~693 steps). Mech 6's upper_bound formula
|
||||
/// (`100 × slow_ema × isv`) is then anchored to the wrong scale during
|
||||
/// fold-N+1 ramp-up — either over-clipping (if slow_ema is too small
|
||||
/// relative to current grads) or under-clipping (after slow_ema catches
|
||||
/// up to the larger F1 grad scale, allowing Adam EMA saturation).
|
||||
///
|
||||
/// Reset philosophy matches Mech 4 (Adam state reset) and Mech 3 (hard
|
||||
/// target sync): fold boundaries are distribution-shift events; ALL
|
||||
/// running EMAs of training dynamics should reinitialize, not just
|
||||
/// some of them. Without Mech 8, Mech 6's anchor was the lone outlier
|
||||
/// — every other distribution-tracking signal already resets per
|
||||
/// existing infrastructure (`reset_for_fold` in fused_training.rs).
|
||||
pub fn reset_grad_norm_slow_ema(&mut self) -> Result<(), MLError> {
|
||||
unsafe { *self.grad_norm_slow_ema_pinned = 0.0_f32; }
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Reset Adam momentum (m/v) for branch weights only (indices 8-49).
|
||||
/// Trunk momentum (indices 0-7) preserved — carries valuable convergence history.
|
||||
/// Lighter than full rewind — tries to break Adam fixed point without changing weights.
|
||||
@@ -5749,6 +5730,9 @@ impl GpuDqnTrainer {
|
||||
let wd_mask_ptr = self.weight_decay_mask.raw_ptr();
|
||||
let l1_end_aux: i32 = 0;
|
||||
let l1_lambda_aux: f32 = 0.0;
|
||||
// SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`).
|
||||
let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
|
||||
let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff;
|
||||
let blocks = ((OFI_EMBED_TOTAL_PARAMS as u32 + 255) / 256).max(1);
|
||||
unsafe {
|
||||
self.stream
|
||||
@@ -5769,6 +5753,7 @@ impl GpuDqnTrainer {
|
||||
.arg(&wd_mask_ptr)
|
||||
.arg(&l1_end_aux)
|
||||
.arg(&l1_lambda_aux)
|
||||
.arg(&weight_clamp_max_abs) // SP3 Mech 9: post-Adam |p_val| clamp
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (blocks, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
@@ -17001,6 +16986,9 @@ impl GpuDqnTrainer {
|
||||
let wd_mask_ptr = self.weight_decay_mask.raw_ptr();
|
||||
let l1_end_aux: i32 = 0;
|
||||
let l1_lambda_aux: f32 = 0.0;
|
||||
// SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`).
|
||||
let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
|
||||
let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff;
|
||||
let blocks = ((DENOISE_TOTAL_PARAMS as u32 + 255) / 256).max(1);
|
||||
unsafe {
|
||||
self.stream
|
||||
@@ -17021,6 +17009,7 @@ impl GpuDqnTrainer {
|
||||
.arg(&wd_mask_ptr)
|
||||
.arg(&l1_end_aux)
|
||||
.arg(&l1_lambda_aux)
|
||||
.arg(&weight_clamp_max_abs) // SP3 Mech 9: post-Adam |p_val| clamp
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (blocks, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
@@ -19454,7 +19443,7 @@ impl GpuDqnTrainer {
|
||||
/// Argument order matches `dqn_adam_update_kernel` in `dqn_utility_kernels.cu`:
|
||||
/// params, grads, m, v, grad_norm_sq, lr, beta1, beta2, epsilon,
|
||||
/// weight_decay, max_grad_norm, t_ptr, total_params,
|
||||
/// weight_decay_mask, l1_end, l1_lambda = 16 args.
|
||||
/// weight_decay_mask, l1_end, l1_lambda, weight_clamp_max_abs = 17 args.
|
||||
///
|
||||
/// Must be launched AFTER `launch_grad_norm` so that `grad_norm_buf`
|
||||
/// contains the completed sum of squares (no race condition).
|
||||
@@ -19483,6 +19472,15 @@ impl GpuDqnTrainer {
|
||||
let l1_end: i32 = align4(param_sizes[0]) as i32; // end of w_s1
|
||||
let l1_lambda: f32 = 1e-3; // L1 penalty strength on w_s1
|
||||
|
||||
// SP3 Mech 9: ISV-driven post-Adam weight clamp.
|
||||
// Bound = 100 × Q_ABS_REF.max(1.0). ε on the multiplier per the
|
||||
// SP1 pearl (`isv.max(1.0)` — never floor the bound itself, that
|
||||
// re-introduces the cold-start clamp pathology). One OOM below the
|
||||
// slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so the
|
||||
// regression sentinel retains 10× firing headroom.
|
||||
let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
|
||||
let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff;
|
||||
|
||||
let count = self.total_params as i32;
|
||||
let blocks = ((self.total_params + 255) / 256) as u32;
|
||||
let cfg = LaunchConfig {
|
||||
@@ -19516,6 +19514,7 @@ impl GpuDqnTrainer {
|
||||
.arg(&wd_mask_base) // G1: per-param weight decay mask
|
||||
.arg(&l1_end) // G8: L1 boundary (end of w_s1)
|
||||
.arg(&l1_lambda) // G8: L1 penalty strength
|
||||
.arg(&weight_clamp_max_abs) // SP3 Mech 9: post-Adam |p_val| clamp
|
||||
.launch(cfg)
|
||||
.map_err(|e| {
|
||||
MLError::ModelError(format!("dqn_adam_update_kernel launch: {e}"))
|
||||
@@ -20336,15 +20335,15 @@ impl GpuDqnTrainer {
|
||||
}
|
||||
let new_clip = (self.grad_norm_ema * CLIP_MULTIPLIER).max(MIN_CLIP);
|
||||
|
||||
// SP3 Mech 6 (anchored) + Mech 8 (anchor reset, 2026-04-29):
|
||||
// anchored upper bound on `new_clip`. The winsorizer above caps a
|
||||
// SINGLE sample at K × prev_clip but does NOT prevent CONSECUTIVE
|
||||
// elevated samples from compounding the EMA upward without bound.
|
||||
// Over hundreds of steps the clip threshold ratchets to thousands
|
||||
// while actual grad_norm tracks it from below → clipping becomes
|
||||
// a no-op against in-distribution drift → Adam m/v EMAs poisoned
|
||||
// → saturate. This is the F1 NaN root cause from smoke-test-5rqzs
|
||||
// (commit b9edccfc1) at step 3060 (Mech 5 diag slots 36–38, 40–42
|
||||
// SP3 Mech 6 (anchored): anchored upper bound on `new_clip`. The
|
||||
// winsorizer above caps a SINGLE sample at K × prev_clip but does
|
||||
// NOT prevent CONSECUTIVE elevated samples from compounding the
|
||||
// EMA upward without bound. Over hundreds of steps the clip
|
||||
// threshold ratchets to thousands while actual grad_norm tracks
|
||||
// it from below → clipping becomes a no-op against
|
||||
// in-distribution drift → Adam m/v EMAs poisoned → saturate.
|
||||
// Original F1 NaN root cause from smoke-test-5rqzs (commit
|
||||
// b9edccfc1) at step 3060 (Mech 5 diag slots 36–38, 40–42
|
||||
// firing).
|
||||
//
|
||||
// Bound: 100 × grad_norm_slow_ema × ISV[Q_ABS_REF=16].max(1.0)
|
||||
@@ -20352,43 +20351,40 @@ impl GpuDqnTrainer {
|
||||
// in this same function) anchors against the legitimate
|
||||
// steady-state grad norm; here we read the PREVIOUS step's slow
|
||||
// EMA from pinned memory (it gets overwritten further down).
|
||||
// Mech 8 (`reset_grad_norm_slow_ema`) hard-resets this scalar at
|
||||
// every fold boundary so the anchor tracks the CURRENT fold's
|
||||
// grad-norm regime from step 1 instead of lagging the prior
|
||||
// fold's value via the α=0.001 EMA half-life (~693 steps). The
|
||||
// pair Mech 6 + Mech 8 is what makes the 100× multiplier safe;
|
||||
// Mech 6 alone with a stale anchor was the F1-saturation
|
||||
// pathology from smoke-test-fxvkk / -ftdjz.
|
||||
// - 100× multiplier: matches the principled per-step headroom
|
||||
// over CURRENT-fold steady-state grad norm. Earlier v2 of this
|
||||
// bound used 5× to compensate for the stale anchor — but that
|
||||
// over-clipped F0 ramp-up gradients (where slow_ema lags
|
||||
// grad_norm_ema during cold-start) and still let F1 saturate
|
||||
// once the EMA caught up. With Mech 8 keeping the anchor fresh,
|
||||
// 100× provides legitimate 10× headroom over the slot-36
|
||||
// diagnostic threshold of 100 × isv without enabling the Adam-EMA
|
||||
// saturation pathology.
|
||||
// The slow EMA INTENTIONALLY persists across fold boundaries —
|
||||
// anchoring at F0's grad scale provides tighter clipping during
|
||||
// F1 ramp-up. Resetting the anchor (Mech 8 experiment, reverted
|
||||
// 2026-04-30 per smoke-test-rxhjh F1 NaN @ 2300) loosened the
|
||||
// bound during F1 transients and accelerated Adam saturation —
|
||||
// worse than the 3720-step ceiling that Mech 6 alone produced.
|
||||
// Cross-fold Adam saturation is now addressed by Mech 9 (post-
|
||||
// Adam weight clamp), not by anchor manipulation.
|
||||
// - 100× multiplier: principled per-step headroom over the
|
||||
// steady-state grad norm. Earlier v2 of this bound used 5× —
|
||||
// that over-clipped F0 ramp-up gradients while still letting
|
||||
// F1 saturate. The 100× value provides legitimate 10× headroom
|
||||
// over the slot-36 diagnostic threshold of 100 × isv.
|
||||
// - ISV[Q_ABS_REF=16] multiplier: scales the cap with the
|
||||
// Q-magnitude regime per `feedback_isv_for_adaptive_bounds.md`.
|
||||
// Same ISV slot used by SP3 Mechs 1+2+5 — no new slot.
|
||||
// Same ISV slot used by SP3 Mechs 1+2+5+9 — no new slot.
|
||||
// - `.max(1.0_f32)` ε on the multiplier per the SP1 ε-floor pearl:
|
||||
// cold-start ISV[16] ≈ 0 must not collapse the bound.
|
||||
// - `.max(MIN_CLIP)` ε-floor on the bound itself: cold-start
|
||||
// `grad_norm_slow_ema` ≈ 0 (true on F0 step 1 AND on every
|
||||
// post-Mech-8 fold-boundary first step) must not pin upper_bound
|
||||
// at 0; MIN_CLIP=1.0 floor stays active for the brief transient
|
||||
// while the slow EMA warms back up over the new fold.
|
||||
// `grad_norm_slow_ema` ≈ 0 (true on F0 step 1) must not pin
|
||||
// upper_bound at 0; MIN_CLIP=1.0 floor stays active for the
|
||||
// brief F0 cold-start transient.
|
||||
//
|
||||
// F1+ history that motivated Mech 8 (anchor reset):
|
||||
// F1+ smoke history (multiplier search converged on 100×):
|
||||
// - smoke-test-fxvkk (Mech 6 only, 100×): F0=44, F1 NaN @ 3720
|
||||
// - smoke-test-ftdjz (Mech 6 + Mech 7, 100×): F0=38, F1 collapse
|
||||
// @ 2040
|
||||
// @ 2040 — Mech 7 reverted (per-element clip misdiagnosis)
|
||||
// - smoke-test-d25vq (Mech 6 v2, 5×): F0=21, F1 collapse @ 2820
|
||||
// The multiplier was searching the wrong dimension — the anchor
|
||||
// itself was stale across fold boundaries, while every other
|
||||
// distribution-tracking signal (Adam m/v, target nets, atom
|
||||
// positions) reset per existing infrastructure. Mech 8 closes the
|
||||
// partial-refactor gap; Mech 6 stays at the principled 100×.
|
||||
// - smoke-test-rxhjh (Mech 6 100× + Mech 8 anchor reset): F1 NaN
|
||||
// @ 2300 (worse than Mech 6 alone) — Mech 8 reverted; the
|
||||
// persistent slow_ema across folds was unintentional protection.
|
||||
// SP3 closes the cross-fold Adam pathology with Mech 9 (post-Adam
|
||||
// weight clamp at 100 × Q_ABS_REF.max(1.0)) rather than further
|
||||
// tuning the clip-anchor side of the chain.
|
||||
let abs_mult = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
|
||||
let prev_slow_ema = unsafe { *self.grad_norm_slow_ema_pinned };
|
||||
let upper_bound = (prev_slow_ema * 100.0_f32 * abs_mult).max(MIN_CLIP);
|
||||
@@ -20836,6 +20832,9 @@ impl GpuDqnTrainer {
|
||||
let wd_mask_ptr = self.weight_decay_mask.raw_ptr();
|
||||
let l1_end_aux: i32 = 0;
|
||||
let l1_lambda_aux: f32 = 0.0;
|
||||
// SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`).
|
||||
let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
|
||||
let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff;
|
||||
let blocks = ((sel_dim as u32 + 255) / 256).max(1);
|
||||
unsafe {
|
||||
self.stream
|
||||
@@ -20856,6 +20855,7 @@ impl GpuDqnTrainer {
|
||||
.arg(&wd_mask_ptr)
|
||||
.arg(&l1_end_aux)
|
||||
.arg(&l1_lambda_aux)
|
||||
.arg(&weight_clamp_max_abs) // SP3 Mech 9: post-Adam |p_val| clamp
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (blocks, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
|
||||
@@ -1064,16 +1064,18 @@ impl FusedTrainingCtx {
|
||||
}
|
||||
}
|
||||
|
||||
// SP3 Mech 8: reset grad_norm_slow_ema at fold boundary. Aligns
|
||||
// Mech 6's anchor with the new fold's grad scale from step 1
|
||||
// (instead of lagging the prior fold's value via α=0.001 EMA).
|
||||
// Without this, Mech 6's upper_bound was either over-clipping
|
||||
// (slow_ema too small) or under-clipping (slow_ema caught up but
|
||||
// F1 grad scale already too large for slot 36 threshold).
|
||||
// Resetting here lets Mech 6 re-anchor to the new fold's scale
|
||||
// properly while preserving the 100× multiplier.
|
||||
self.trainer.reset_grad_norm_slow_ema()?;
|
||||
|
||||
// SP3 Mech 8 (slow_ema fold-boundary reset, 2026-04-29) was
|
||||
// tested here and reverted 2026-04-30 after smoke-test-rxhjh
|
||||
// showed F1 NaN @ 2300 — WORSE than the 3720-step ceiling Mech 6
|
||||
// alone produced. Persisting `grad_norm_slow_ema` across folds
|
||||
// (anchored at F0's smaller grad scale) provides unintentional
|
||||
// F1 protection by tightening Mech 6's upper_bound during F1
|
||||
// ramp-up. Resetting it loosened the bound and accelerated Adam
|
||||
// saturation. SP3 Mech 9 (post-Adam ISV-driven weight clamp in
|
||||
// `dqn_adam_update_kernel`) addresses cross-fold Adam pathology
|
||||
// by bounding parameter magnitudes directly rather than tuning
|
||||
// the clip anchor.
|
||||
//
|
||||
// FoldReset state (eval_q_mean_ema, eval_q_std_ema, isv_v_range_slots,
|
||||
// isv_learning_health, isv_sharpe_ema, isv_q_means) is now handled by
|
||||
// StateResetRegistry-driven dispatch in DQNTrainer::reset_for_fold
|
||||
|
||||
@@ -815,11 +815,15 @@ If Gate 1 fails, halt before SP3 starts and re-investigate F0 cause (alignment e
|
||||
|
||||
| # | Mechanism | Commit | File(s) |
|
||||
|---|---|---|---|
|
||||
| Mech 1 | Target-Q clipping at single-point source (`compute_denoise_target_q` end) | `27ed7daa6` | gpu_dqn_trainer.rs |
|
||||
| Mech 2 | C51 atom-position growth bounds (3 sites: atoms_update, iql_value, experience) | `96b77043e` | atoms_update_kernel.cu, iql_value_kernel.cu, experience_kernels.cu |
|
||||
| Mech 3 | Hard target sync at fold boundary (DQN main + IQN target) | pre-existing | fused_training.rs::reset_for_fold |
|
||||
| Mech 4 | Comprehensive Adam EMA reset at fold transitions (incl. NEW: GpuAttention) | `ef429c25d` | gpu_attention.rs, fused_training.rs |
|
||||
| Mech 5 | Adam-centric diagnostic in slots 36-47 via fused-kernel extension | `45e077188` + `5f7db7d8d` | dqn_utility_kernels.cu, gpu_dqn_trainer.rs, training_loop.rs |
|
||||
| Mech 1 | Target-Q clipping at single-point source | `27ed7daa6` | gpu_dqn_trainer.rs |
|
||||
| Mech 2 | C51 atom-position growth bounds (3 sites) | `96b77043e` | atoms_update_kernel.cu, iql_value_kernel.cu, experience_kernels.cu |
|
||||
| Mech 3 | Hard target sync at fold boundary | pre-existing | fused_training.rs::reset_for_fold |
|
||||
| Mech 4 | Comprehensive Adam EMA reset (+GpuAttention) | `ef429c25d` | gpu_attention.rs, fused_training.rs |
|
||||
| Mech 5 | Adam-centric diagnostic in slots 36-47 | `45e077188` + `5f7db7d8d` | dqn_utility_kernels.cu, gpu_dqn_trainer.rs, training_loop.rs |
|
||||
| Mech 6 | Anchored upper_bound on adaptive grad clip (100× multiplier) | `48c25d999` | gpu_dqn_trainer.rs |
|
||||
| ~~Mech 7~~ | ~~Per-element gradient clip in Adam~~ | reverted in `f67ede94f` | (n/a) |
|
||||
| ~~Mech 8~~ | ~~slow_ema fold-boundary reset~~ | reverted in <this commit> | (n/a) |
|
||||
| Mech 9 | Post-Adam ISV-driven weight clamp | <this commit> | dqn_utility_kernels.cu, gpu_dqn_trainer.rs |
|
||||
|
||||
**Supporting:** Task B1 (`aee11f321`) — 24 diagnostic accessors (Adam m/v + weight + target_q + atom_positions).
|
||||
|
||||
@@ -866,3 +870,59 @@ Per slot 36-47 firing pattern:
|
||||
- Slot 47 (atom span) fires often: Mech 2 doing real work; no action
|
||||
- Slots 24-35 fire post-fix: SP3 missed pathology class; new investigation thread
|
||||
|
||||
### Mech 8 reverted + Mech 9 added — post-Adam weight clamp
|
||||
|
||||
**Mech 8 outcome (reverted 2026-04-30).** Mech 8 added
|
||||
`reset_grad_norm_slow_ema()` zeroing the persistent α=0.001 slow-EMA
|
||||
on every fold boundary, intending to keep Mech 6's `100 × slow_ema ×
|
||||
isv` upper_bound anchored to the CURRENT fold's grad-norm regime.
|
||||
Smoke-test-rxhjh (HEAD `b8a7ac6f7` = Mech 8 active) recorded F1 NaN at
|
||||
step 2300 — WORSE than the 3720-step ceiling that Mech 6 alone produced
|
||||
on smoke-test-fxvkk. Diagnosis: persisting `slow_ema` across folds
|
||||
(anchored at F0's smaller grad scale) was providing UNINTENTIONAL F1
|
||||
protection. Mech 6's upper_bound stayed tight during F1 ramp-up because
|
||||
the cross-fold lag of the EMA half-life (~693 steps) kept the anchor
|
||||
small while F1 grad scale grew. Resetting the anchor loosened the
|
||||
bound during the very transient when F1 needed it tightest, and Adam
|
||||
EMAs saturated faster. Mech 8 reverted in this commit; the `100×`
|
||||
multiplier in Mech 6 stays (already restored from the 5× experiment in
|
||||
the Mech 8 commit — correct steady-state value).
|
||||
|
||||
**Mech 9 mechanism.** Mech 1 (gradient clamp at target-Q source) and
|
||||
Mech 6 (gradient-norm clip via adaptive_clip) bound GRADIENTS, not
|
||||
PARAMETERS. The slot 26 (`iqn_trunk_m`) + slot 32 (`bn_d_concat_buf`)
|
||||
F1-step-3540 NaN signature has the cuBLAS sgemm INPUTS clean (slots 27,
|
||||
35 stay at 0) — the GEMM output saturates because the WEIGHT tensors
|
||||
have drifted to extreme-but-finite magnitudes via Adam updates over
|
||||
thousands of steps. Once `|w| > O(1e3 × Q_ABS_REF)` and `K ≈ 256` inner
|
||||
dim, the f32 accumulator under TF32 path overflows even with bounded
|
||||
inputs. Mech 9 closes the gap by clamping `|p_val|` directly inside
|
||||
`dqn_adam_update_kernel` after the L1 proximal step, before the writeback.
|
||||
|
||||
**ISV-driven design.** Bound = `100 × Q_ABS_REF.max(1.0)`. ε floor on the
|
||||
multiplier per the SP1 pearl (`isv.max(1.0)` not `bound.max(constant)`
|
||||
— flooring the bound itself re-introduces the cold-start clamp pathology
|
||||
SP1 closed). `Q_ABS_REF_INDEX = 16` of the ISV signal bus, same slot
|
||||
already consumed by Mechs 1+2+5+6 — no new slot. Bound is one OOM below
|
||||
the slot 44-45 diagnostic threshold (`1e3 × Q_ABS_REF`) so the
|
||||
regression sentinel retains a 10× firing headroom — symmetric with
|
||||
Mech 1's clamp/diagnostic ratio.
|
||||
|
||||
**Implementation.** New trailing arg `weight_clamp_max_abs` on
|
||||
`dqn_adam_update_kernel`. CPU computes the bound host-side once per Adam
|
||||
launch from `read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0) × 100` and
|
||||
passes it as a scalar — keeps the kernel signature clean and avoids
|
||||
per-thread re-reading of the ISV buffer. Wired at all four Adam launch
|
||||
sites: main `launch_adam_update` (DQN params), `step_selectivity_adam`
|
||||
(selectivity gate), denoise-head Adam, OFI-embed Adam. The Decision
|
||||
Transformer Adam launch in `decision_transformer.rs` passes 0.0 to
|
||||
disable the clamp (DT is a separate offline-RL trainer outside SP3
|
||||
scope, with no ISV bus access). Kernel branch `if (bound > 0.0f)` makes
|
||||
the disabled path zero-cost for DT.
|
||||
|
||||
**Residual scope.** If Mech 9 still leaves F1 NaN at similar steps, the
|
||||
pathology is not weight-magnitude growth and the next investigation
|
||||
should attack the GEMM accumulator path itself (TF32 disable, FP32
|
||||
enforcement, or per-layer gain rescaling). Close SP3 honestly with the
|
||||
documented residual rather than tuning further.
|
||||
|
||||
|
||||
@@ -2,6 +2,8 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
SP3 Mech 9 — post-Adam ISV-driven weight clamp + Mech 8 revert (2026-04-30): coordinated close-out of SP3 Q-learning numerical-stability per `feedback_no_partial_refactor.md`. Mech 8 (slow_ema fold-boundary reset, commit `b8a7ac6f7`) reverted after smoke-test-rxhjh F1 NaN @ step 2300 — WORSE than the 3720-step ceiling Mech 6 alone produced. Persisting the α=0.001 `grad_norm_slow_ema` across folds (anchored at F0's smaller grad scale) was providing UNINTENTIONAL F1 protection by tightening Mech 6's `100 × slow_ema × isv` upper_bound during F1 ramp-up; resetting it loosened the bound during the very transient when F1 needed it tightest, accelerating Adam saturation. Removed `GpuDqnTrainer::reset_grad_norm_slow_ema` method and the `fused_training.rs::reset_for_fold` call site; kept the `grad_norm_slow_ema_pinned` field (still consumed by Mech 6) and Mech 6 multiplier at 100× (correct steady-state value, already restored from the 5× experiment in the Mech 8 commit). Mech 9 is the real fix for the cross-fold Adam pathology — root cause being targeted is cuBLAS sgemm f32-accumulator overflow at slots 26 (`iqn_trunk_m`) + 32 (`bn_d_concat_buf`) at F1 step ~3540 with INPUTS clean (slots 27, 35 unflagged). The matmul output saturates because WEIGHT tensors have drifted to extreme-but-finite magnitudes via Adam updates over thousands of steps. Mech 1 (gradient clamp) and Mech 6 (grad-norm clip) bound gradients, not parameters; neither prevents this drift. Mech 9 clamps `|p_val| ≤ 100 × Q_ABS_REF.max(1.0)` inside `dqn_adam_update_kernel` after the L1 proximal step (new trailing arg `weight_clamp_max_abs` on the kernel signature). ISV-driven: bound reads `Q_ABS_REF_INDEX = 16` host-side (same slot consumed by Mechs 1+2+5+6 — no new slot), ε on the multiplier per the SP1 pearl (`isv.max(1.0)`, never floor the bound itself or re-introduce the cold-start clamp pathology), 1 OOM below the slot 44-45 diagnostic threshold (`1e3 × Q_ABS_REF`) so the regression sentinel retains 10× firing headroom — symmetric with Mech 1's clamp/diagnostic ratio. All four Adam launch sites in `gpu_dqn_trainer.rs` wired with the bound (`launch_adam_update` at ~line 19470, `step_selectivity_adam` at ~line 20860, denoise-head Adam at ~line 17014, OFI-embed Adam at ~line 5758); the Decision Transformer Adam launch in `decision_transformer.rs` passes 0.0 to disable the clamp (DT is a separate offline-RL trainer outside SP3 scope, no ISV bus access; kernel branch `if (bound > 0.0f)` makes the disabled path zero-cost). No new ISV slots, no new kernels, no graph topology change. Touched: `cuda_pipeline/dqn_utility_kernels.cu` (kernel signature + clamp), `cuda_pipeline/gpu_dqn_trainer.rs` (4 launch sites + Mech 6 comment update + `reset_grad_norm_slow_ema` method removed), `cuda_pipeline/decision_transformer.rs` (DT launch site, disabled), `trainers/dqn/fused_training.rs` (Mech 8 call removed), `docs/dqn-backward-nan-audit.md` (mechanism table + Mech 8/9 subsection). cargo check clean. No fingerprint change — kernel signature change but param-tensor layout unchanged.
|
||||
|
||||
SP3 Mech 8 — `grad_norm_slow_ema` fold-boundary reset (2026-04-29): closed the partial-refactor gap where the α=0.001 grad-norm slow EMA (driving Mech 6's anchored upper bound on `adaptive_clip`) persisted across fold boundaries while every other distribution-tracking signal — Adam m/v (`reset_adam_state`), main DQN target params (`sync_target_from_online`), IQN target + Adam (`iqn.sync_target_from_online` + `iqn.reset_adam_state`), TLOB / IQL / IQL-low / 4-head attention Adam (`reset_adam_state` on each), C51 atom-position EMAs, replay buffer — already reset per existing `reset_for_fold` infrastructure. The α=0.001 half-life (~693 steps) meant `grad_norm_slow_ema` lagged the new fold's grad-norm regime by hundreds of steps; Mech 6's `100 × slow_ema × isv` upper bound was anchored to the WRONG fold's scale during F1 ramp-up, producing the non-monotonic multiplier-tuning dance observed across smokes (smoke-test-fxvkk 100× F1-NaN @ 3720 / smoke-test-ftdjz 100×+Mech7 F1-collapse @ 2040 / smoke-test-d25vq 5× F1-collapse @ 2820 — the multiplier was searching the wrong dimension). Three coordinated changes per `feedback_no_partial_refactor.md`: (1) Mech 6 multiplier restored 5× → 100× in `gpu_dqn_trainer.rs::update_adaptive_clip` (the original principled setting; v2's 5× was over-clipping legitimate F0-ramp gradients while still allowing F1-anchor saturation), (2) new `GpuDqnTrainer::reset_grad_norm_slow_ema` method zeroes the existing mapped-pinned scalar (no new buffer, no new ISV slot — reuses Plan C Phase 2 follow-up K's grad_norm_slow_ema_pinned at `gpu_dqn_trainer.rs:3337`), (3) wired into `fused_training.rs::reset_for_fold` immediately after the SP3 Mech 4 attention Adam reset block (line ~1066) — same fold-boundary contract consumer chain as Mech 3+4. F0 unaffected (cold-starts at slow_ema=0 already, so Mech 8 is a no-op on F0); F1+F2 first step transient sees `upper_bound = 100 × 0 × isv → MIN_CLIP=1.0` floor for ~10-50 steps as the EMA warms, then converges to legitimate 100× headroom over CURRENT fold's grad norm. Mech 7 stays reverted (per-element clip was misdiagnosis). Touched: `cuda_pipeline/gpu_dqn_trainer.rs` (Mech 6 comment + multiplier 5.0→100.0 line ~20370; new `reset_grad_norm_slow_ema` method line ~3858), `trainers/dqn/fused_training.rs` (call wired line ~1067). cargo check clean. No fingerprint change — touches an already-allocated mapped-pinned scalar, not param-tensor layout.
|
||||
|
||||
HEALTH_DIAG GPU port — Phase 1 (2026-04-28, follow-up to commit `333ea7184`'s Phase 0 inventory): added `HealthDiagSnapshot` `#[repr(C)]` POD struct and `MappedHealthDiagSnapshot` mapped-pinned wrapper in new `crates/ml/src/cuda_pipeline/health_diag.rs`. The struct holds 147 fields (every numeric value emitted across the 7 HEALTH_DIAG log sites) as `f32` or `u32` so CPU and GPU agree byte-for-byte; controller fire booleans (`fire_lr`/`fire_tau`/etc.) promoted from `u8` → `u32` per the design review to avoid `#[repr(C)]` padding mismatch when followed by `f32` fields. `MappedHealthDiagSnapshot::new()` wraps `cuMemHostAlloc(DEVICEMAP|PORTABLE)` of `sizeof(HealthDiagSnapshot)` + `cuMemHostGetDevicePointer_v2` — the same pattern as `MappedF32Buffer` but typed for the snapshot. Re-exported from `cuda_pipeline/mod.rs` for downstream consumers. **No producer kernel and no consumer wiring in this commit** — Phase 2 lands the kernel family (`health_diag_per_sample_reduce`, `health_diag_q_mag_reduce`, `health_diag_eval_histogram`, `health_diag_isv_mirror`, `health_diag_finalise`) and the `gpu_health_diag.rs` orchestrator; Phase 3 wires the launch into the captured graph alongside `launch_h_s2_rms_ema`; Phase 4 rewrites the CPU emit path and deletes the 12 CPU-bound feeder methods identified in the Phase 0 inventory. Three unit tests (`snapshot_size_is_stable`, `default_is_zeroed`, `alignment_is_4_bytes`) pin the layout so any future field add/reorder requires an explicit test update — guards the kernel-side offset table from silent drift. Touched: `cuda_pipeline/health_diag.rs` (+339 LOC, new), `cuda_pipeline/mod.rs` (+8 LOC re-exports). cargo check clean at 12 warnings (workspace baseline). No fingerprint change — separate mapped-pinned alloc, not part of param-tensor layout the fingerprint guards.
|
||||
|
||||
Reference in New Issue
Block a user