fix(dqn): SP3 Mech 6 v2 — tighten clip multiplier 100x -> 5x; revert Mech 7

Smoke smoke-test-ftdjz (commit d9a4d98a3, Mech 6 + Mech 7) regressed
both F0 (44 -> 38.82, per-element cap clipped legitimate outliers) and
F1 (grad-collapse at step 2040 vs 3720 with Mech 6 alone). Slots 36-42
STILL fired — Mech 7's per-element clip didn't prevent Adam EMA
saturation, just slowed training to grad-collapse.

Re-analysis: Mech 6's 100x multiplier on the upper-bound formula was
mismatched with the slot 36 threshold ratio. Per-element gradient max
<= adaptive_clip = 100 x slow_ema x isv ~= 200. Adam m_X steady-state
reaches 200, exceeding slot 36 threshold of 100*isv = 100. Mech 6 was
doing its job but the bound was wider than the diagnostic threshold.

Two coordinated changes (per feedback_no_partial_refactor):

1. REVERT Mech 7 (per-element clip in dqn_adam_update_kernel). The
   per-element approach was misdiagnosis — clipping post-global-clip
   gradients tighter than legitimate per-element variance harms F0
   training without addressing Adam saturation root cause. Kernel
   returns to its post-Mech-6 state (blob 546feee48).

2. TIGHTEN Mech 6's upper-bound multiplier from 100x to 5x. Standard
   DL practice (5-10x steady-state grad norm). Per-element gradient
   max becomes <= 5 x slow_ema x isv ~= 10, well below slot 36
   threshold of 100. Adam m_X EMA stays bounded <= 10 — slots 36-42
   should not fire.

Why 5x and not 10x: slot 36 threshold is 100 x isv. 5x slow_ema
(~=10 absolute) leaves 10x headroom against the threshold, providing
robust margin. 10x slow_ema would be 20 absolute, only 5x margin —
risk of fluctuations triggering slot 36.

Why not tighter (e.g., 2x): per-step gradient norms can legitimately
spike to 5x slow_ema in normal training (e.g., gradient resumption
after warmup); tighter bounds would over-clip.

F0 risk: low — F0 typical adaptive_clip values are bounded by Mech 6's
upper anchor only when the EMA-driven clip exceeds 5x slow_ema, which
is rare in steady F0 training. Cap should be a no-op for F0.

F1 risk: prevents the saturation pathway diagnosed by smoke
smoke-test-ftdjz. Validates by running the next smoke at this commit.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu

@@ -136,27 +136,6 @@ extern "C" __global__ void dqn_adam_update_kernel(
     float clip_scale_f = (norm_f > max_grad_norm) ? (max_grad_norm / norm_f) : 1.0f;
     float clipped_g_f = g_f * clip_scale_f;
 
-    /* SP3 Mech 7: per-element gradient clip — prevents Adam EMA saturation
-     * at single-element extreme gradients that pass the global L2 norm clip.
-     * Mech 6 (anchored upper-bound on adaptive_clip) bounds the AGGREGATE
-     * gradient magnitude, but a gradient with one large element + rest small
-     * can satisfy the L2 budget while poisoning the Adam m EMA at that
-     * element's index. Smoke smoke-test-fxvkk (commit 48c25d999) F1-NaN'd at
-     * step 3720 with slots 36-38, 40-42 firing despite Mech 6 — confirming
-     * per-element saturation is the residual pathology.
-     *
-     * per_element_cap = 10 × sqrt(adaptive_clip / total_params)
-     *   - sqrt(adaptive_clip / N) is the AVERAGE per-element L2 budget
-     *     (if all elements equal, each contributes sqrt(clip²/N) = clip/sqrt(N))
-     *   - 10× allows legitimate per-element deviations up to 10× average
-     *     (some elements legitimately have larger gradients than others)
-     *   - When the global clip is doing its job (most steps), per_element_cap
-     *     is tight enough to prevent saturation
-     *   - When global clip is loose (e.g., post-fold-warmup), per_element_cap
-     *     scales with it — never tighter than the global clip's intent */
-    const float per_element_cap = 10.0f * sqrtf(max_grad_norm / (float)total_params);
-    clipped_g_f = copysignf(fminf(fabsf(clipped_g_f), per_element_cap), clipped_g_f);
-
     /* Adam bias correction: native float — no bf16 rounding issues. */
     float beta1_t = fmaxf(1.0f - powf(beta1, (float)t), 1e-8f);
     float beta2_t = fmaxf(1.0f - powf(beta2, (float)t), 1e-8f);

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -20317,7 +20317,7 @@ impl GpuDqnTrainer {
         }
         let new_clip = (self.grad_norm_ema * CLIP_MULTIPLIER).max(MIN_CLIP);
 
-        // SP3 Mech 6 (2026-04-29): anchored upper bound on `new_clip`.
+        // SP3 Mech 6 v2 (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
@@ -20327,14 +20327,18 @@ impl GpuDqnTrainer {
         // 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)
+        // Bound: 5 × grad_norm_slow_ema × ISV[Q_ABS_REF=16].max(1.0)
         // - `grad_norm_slow_ema` (existing α=0.001 slow EMA, updated below
         //   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).
-        // - 100×: headroom for legitimate per-step deviations; single
-        //   steps can have norm 10–100× the slow average without being
-        //   pathological.
+        // - 5× multiplier (revised from 100× — see "v2" rationale below):
+        //   matches standard DL practice (5–10× steady-state grad norm).
+        //   At 5×, per-element gradient max ≤ adaptive_clip = 5 × slow_ema
+        //   × isv ≈ 10 (typical slow_ema ≈ 2, isv-floor = 1), well below
+        //   slot 36 firing threshold of 100 × isv = 100. Adam m_X EMA
+        //   steady-state with worst-case constant g = 10 reaches m_X = 10
+        //   — slots 36–42 should not fire.
         // - 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.
@@ -20344,9 +20348,26 @@ impl GpuDqnTrainer {
         //   `grad_norm_slow_ema` ≈ 0 must not pin upper_bound at 0 (which
         //   would force `new_clip` down to MIN_CLIP every step until the
         //   slow EMA warms up).
+        //
+        // v2 rationale (multiplier 100 → 5):
+        // - Smoke smoke-test-fxvkk (Mech 6 v1, 100×) F1-NaN'd at step 3720
+        //   with slots 36-42 still firing. Per-element gradient max under
+        //   v1 reached ≤ 100 × slow_ema × isv ≈ 200, and Adam m_X EMA
+        //   steady-state matched, exceeding slot 36 threshold of 100. The
+        //   100× bound was wider than the diagnostic threshold — Mech 6
+        //   was active but the cap was the wrong width.
+        // - Why 5× and not tighter (e.g., 2×): per-step gradient norms can
+        //   legitimately spike to ~5× slow_ema in normal training (e.g.,
+        //   gradient resumption after warmup, occasional curvature
+        //   transients); tighter bounds would over-clip and harm fit.
+        // - Why not wider (e.g., 10× or 20×): >10× restores the slot 36
+        //   threshold-overshoot pathology with diminishing margin. 5×
+        //   leaves robust 10× headroom against the 100× threshold.
+        // - F0 risk: low — F0 typical adaptive_clip stays under 5×
+        //   slow_ema in steady training; the cap should be a no-op for F0.
         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);
+        let upper_bound = (prev_slow_ema * 5.0_f32 * abs_mult).max(MIN_CLIP);
         let new_clip = new_clip.min(upper_bound);
 
         // Write directly to pinned memory — GPU sees it on next kernel read

docs/dqn-wire-up-audit.md

@@ -2285,6 +2285,8 @@ SP3 Task B6 — fused kernel extended for slots 36-47 threshold checks (2026-04-
 
 SP3 Task B7 — name-table entries for slots 36-47 readback log (2026-04-30): replaced the SP1 Phase B placeholder strings (`"rsv36"`-`"rsv47"`) in both `training_loop.rs` `let names = [...]` tables (the `halt_nan` block ~L1992 and the `halt_grad_collapse` block ~L2089) with the SP3 Mech 5 diagnostic slot names per the B1/B6 accessor allocation: 36 `trunk_adam_m_max`, 37 `value_adam_m_max`, 38 `branch_adam_m_max`, 39 `iqn_adam_m_max`, 40 `trunk_adam_v_max`, 41 `value_adam_v_max`, 42 `branch_adam_v_max`, 43 `iqn_adam_v_max`, 44 `trunk_weight_max`, 45 `heads_weight_max`, 46 `target_q_post_clip`, 47 `atom_span_max`. Both name tables receive byte-identical replacement content (modulo the indentation difference between the two enclosing blocks) per `feedback_no_partial_refactor` — the post-SP2 stale-doc cleanup commit `387335e2b` already established that the two tables must remain identical, and the same shared-contract migration principle applies here. When the fused kernel (B6) sets a slot bit, the readback log line names the buffer that exceeded its ISV-derived threshold (e.g. `flagged=[42=branch_adam_v_max, 46=target_q_post_clip]`) instead of the opaque `rsv*` placeholder, providing direct observability for SP3 mechanism effectiveness. Pure name-string replacement — no logic changes, no kernel changes, no buffer changes, no ISV changes. `cargo check -p ml --lib` clean.
 
+SP3 Mech 6 v2 + Mech 7 revert — tighten anchored clip multiplier 100× → 5× (2026-04-29): coordinated change across `dqn_utility_kernels.cu` (revert) + `gpu_dqn_trainer.rs` (multiplier tighten) per `feedback_no_partial_refactor`. **Smoke evidence**: smoke `smoke-test-ftdjz` (commit `d9a4d98a3`, Mech 6 v1 + Mech 7 combined) regressed BOTH F0 and F1: F0 dropped from ~44 to 38.82 (per-element cap clipped legitimate gradient outliers, harming fit on the cleanest fold) and F1 grad-collapse moved earlier (step 2040 vs 3720 with Mech 6 v1 alone). Slots 36-42 STILL fired under the combined fix — Mech 7 didn't prevent Adam EMA saturation, just slowed training to grad-collapse on F1 while harming F0. **Re-analysis** of the saturation pathway: Mech 6 v1's 100× multiplier on `upper_bound = prev_slow_ema × 100 × isv.max(1)` was mismatched with the slot-36 firing threshold ratio (slot 36 fires at `100 × isv.max(1)`). With slow_ema ≈ 2 and isv = 1, upper_bound ≈ 200, so per-element gradient max ≤ adaptive_clip = 200 (worst case all energy in one element). Adam m_X EMA steady-state with constant g = 200 reaches m_X = 200 — exceeding slot 36 threshold of 100. Mech 6 was active and BOUNDED the clip, but the bound was wider than the diagnostic threshold by 2×. **Two coordinated changes**: (1) REVERT Mech 7's per-element clip block from `dqn_adam_update_kernel` — the per-element approach was misdiagnosis; it caps post-global-clip gradients tighter than legitimate per-element variance, harming F0 without addressing the Adam saturation root cause. Kernel returns to its post-Mech-6 state for that section (`clipped_g_f = g_f * clip_scale_f` directly followed by Adam moment update). (2) TIGHTEN Mech 6 v1's `upper_bound` multiplier from `100.0_f32` to `5.0_f32` in `update_adaptive_clip`. Standard DL practice for clip thresholds is 5–10× steady-state grad norm (vs Mech 6 v1's 100× which was 10–20× standard). Per-element gradient max becomes ≤ 5 × slow_ema × isv ≈ 10 (worst case), well below slot 36 threshold of 100; Adam m_X steady-state caps at ≈ 10. **Why 5× and not 10×**: slot 36 threshold is 100 × isv. 5× slow_ema (≈ 10 absolute) leaves robust 10× headroom against the threshold; 10× slow_ema would be 20 absolute, only 5× margin — risk of fluctuations triggering slot 36. **Why not tighter (e.g., 2×)**: per-step gradient norms can legitimately spike to ~5× slow_ema in normal training (gradient resumption after warmup, occasional curvature transients); tighter would over-clip and harm fit. **F0 risk**: low — F0 typical adaptive_clip values are bounded by Mech 6's anchored upper bound only when the EMA-driven clip exceeds 5× slow_ema, which is rare in steady F0 training (typical clip = `grad_norm_ema × 2` stays under 5× slow_ema). Cap should be a no-op for F0; F1 v1's 38.82 regression should reverse to the 44 baseline once Mech 7's per-element cap is removed. **F1 expected behaviour**: per-element gradient max ≤ 10, Adam m_X ≤ 10, slots 36-42 should not fire — validates by next smoke at this commit. Composes with: Mech 1 (target_q clipping at single-point source, B2), Mech 2 (atom-position growth bounds, B3), Mech 3 (CQL aux loss clamping, deferred), Mech 4 (Adam EMA reset comprehensive, B5), Mech 5 (fused threshold-check kernel for slots 36-47, B6+B7). Mech 6 v2 is now the single Adam-saturation defence (Mech 7 retired). Zero new HtoD/DtoD/HtoH copies; zero new ISV slots (reuses `Q_ABS_REF_INDEX = 16` and `grad_norm_slow_ema_pinned`); zero new buffers; zero new kernels. `cargo check -p ml --lib` clean.
+
 SP3 Mech 7 — per-element gradient clip in Adam kernel (2026-04-29): added a per-element gradient cap inside `dqn_adam_update_kernel` (`crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`) immediately after the existing global L2 norm clip (`clipped_g_f = g_f * clip_scale_f`) and before the Adam m/v moment updates. **Why Mech 6 wasn't enough**: smoke `smoke-test-fxvkk` (commit `48c25d999`, Mech 6 anchored upper-bound clip) F1-NaN'd at step 3720 (vs step 3060 in prior smoke; Mech 6 helped marginally, ~22% later) with slots 36-38, 40-42 still firing — DQN main Adam m + v saturated despite Mech 6 capping the AGGREGATE clip threshold. The L2 norm is an aggregate quantity but Adam EMAs are PER-ELEMENT: a gradient with one large element (e.g. `element_X = 100`, rest small) has `||g||₂ ≈ 100` and passes a clip threshold of 1000 untouched, but the Adam m_X EMA accumulates the large element. With β1=0.9 steady-state, `m_X ≈ 100 / (1 - 0.9) = 1000`, exceeding slot 36's threshold (`100 × ISV[Q_ABS_REF].max(1.0)`). **Cap formula**: `per_element_cap = 10 × sqrt(max_grad_norm / total_params)` where `max_grad_norm` is the existing kernel arg sourced from `max_grad_norm_buf` (= the Mech-6-bounded `adaptive_clip` mapped-pinned scalar), and `total_params` is the existing kernel arg counting all DQN main params. **Rationale**: `sqrt(adaptive_clip / N)` is the AVERAGE per-element contribution to the L2 norm budget — if all elements equal in magnitude, each contributes `sqrt(clip² / N) = clip / sqrt(N)`, and the L2 norm equals `clip` exactly. The 10× multiplier allows legitimate per-element deviations up to 10× average (some elements legitimately have larger gradients than others — e.g. branch heads vs trunk parameters); together this means the per-element cap kicks in only for the genuinely outlying elements that would saturate Adam m EMAs. **Coupling with Mech 6**: the per-element cap scales with `adaptive_clip`, so when Mech 6's anchored bound keeps `adaptive_clip` reasonable (the typical case), `per_element_cap` is correspondingly tight (e.g. `adaptive_clip = 10`, `total_params ≈ 200k` → `per_element_cap = 10 × sqrt(10 / 200000) ≈ 0.0707` — a tight bound on per-element gradient contribution). When Mech 6's bound is loose (e.g. immediately post-fold-warmup before the slow-EMA stabilises), `per_element_cap` scales up with it — never tighter than the global clip's intent. **Implementation**: pure kernel modification using existing args; no new kernel arg, no new buffer, no new ISV slot, no new launch site, no graph recapture. The cap is applied via `clipped_g_f = copysignf(fminf(fabsf(clipped_g_f), per_element_cap), clipped_g_f)` — preserves sign, bounds magnitude. **What stays unchanged**: the global L2 norm clip itself (computed identically), Adam bias correction, m/v EMA β1/β2 coefficients, AdamW decoupled weight decay, L1 proximal step on `w_s1`, the NaN/Inf gradient skip guard, every kernel launch site and graph layout. Composes with: Mech 1 (target_q clipping at single-point source, B2), Mech 2 (atom-position growth bounds, B3), Mech 3 (CQL aux loss clamping, deferred), Mech 4 (Adam EMA reset comprehensive, B5), Mech 5 (fused threshold-check kernel for diagnostic slots 36-47, B6+B7), Mech 6 (anchored upper bound on aggregate adaptive grad clip). Mechs 6+7 together prevent Adam saturation at BOTH the aggregate (L2 norm) and per-element levels — closes the residual pathology after Mech 6's partial 3060→3720 improvement. F0 risk is low: F0-typical `adaptive_clip` is on the order of 5-20 with `total_params ≈ 200k`, giving `per_element_cap ≈ 10 × sqrt(10/200000) = 0.0707`, well above F0-typical per-element gradients (~1e-3 to 1e-2 for a converged trunk on stable input), so Mech 7 is invisible in normal F0 operation. F1+F2 benefit by preventing single-element Adam EMA poisoning. Zero new HtoD/DtoD/HtoH copies; zero new ISV slots; zero new buffers; zero new kernels. `cargo check -p ml --lib` clean.
 
 SP3 Mech 6 — anchored upper bound on adaptive grad clip (2026-04-29): added an upper bound to `GpuDqnTrainer::update_adaptive_clip`'s `new_clip` formula in `gpu_dqn_trainer.rs`. The existing winsorizer (Plan C T11 follow-up N) caps a SINGLE input sample at `K=100 × prev_clip` before the EMA absorbs it, 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 are poisoned, and they saturate at the SP3 Mech 5 slot 36-43 thresholds. This is the F1 NaN root cause from `smoke-test-5rqzs` (commit `b9edccfc1`) at step 3060: Mech 5 diagnostic flags `[36=trunk_adam_m_max, 37=value_adam_m_max, 38=branch_adam_m_max, 40=trunk_adam_v_max, 41=value_adam_v_max, 42=branch_adam_v_max]` fired with target_q + atoms bounded (Mechs 1+2 working) and weights still finite — narrowing the divergence to the Adam state itself. **Bound formula**: `upper_bound = (grad_norm_slow_ema × 100 × ISV[Q_ABS_REF=16].max(1.0)).max(MIN_CLIP=1.0)`; final `new_clip = (grad_norm_ema × CLIP_MULTIPLIER).max(MIN_CLIP).min(upper_bound)`. **Anchor**: `grad_norm_slow_ema` is the existing α=0.001 slow-EMA scalar (mapped-pinned, updated later in this same function via `*self.grad_norm_slow_ema_pinned`) — read from pinned memory BEFORE the slow-EMA update on this step, so it reflects the previous step's slow EMA (the legitimate steady-state grad norm at the time of clip computation). **Headroom 100×**: legitimate per-step deviations can be 10-100× the slow average without being pathological, so `100 ×` keeps Mech 6 invisible in normal training and only kicks in when the EMA-driven clip ratchets past plausible-deviation bounds. **ISV-adaptive multiplier**: `ISV[Q_ABS_REF_INDEX = 16].max(1.0_f32)` scales the cap with the Q-magnitude regime per `feedback_isv_for_adaptive_bounds` — same ISV slot used by SP3 Mechs 1, 2, 3, and 5 (no new slot). ε on the multiplier (`.max(1.0)`) per the SP1 ε-floor pearl — cold-start `ISV[16] ≈ 0` would otherwise collapse `upper_bound` toward zero. **ε-floor on the bound itself** (`.max(MIN_CLIP=1.0)`): cold-start `grad_norm_slow_ema ≈ 0` (first ~200 steps before the slow EMA warms up) would otherwise pin `upper_bound` at 0, which combined with `new_clip.min(upper_bound)` would force `new_clip = MIN_CLIP=1.0` every step until the slow EMA established a meaningful baseline — exactly the cold-start ratcheting the upper bound is meant to PREVENT. The `MIN_CLIP` floor on the bound aligns with the existing `MIN_CLIP` floor on `new_clip` itself, so the bound is at minimum a no-op (matching the floor below) until the slow EMA warms up. **What stays unchanged**: the existing winsorizer (single-sample input cap before EMA update), the `EMA_BETA = 0.95` adaptive_clip EMA, the `CLIP_MULTIPLIER = 2.0` and `MIN_CLIP = 1.0` constants, the `grad_norm_fast_ema` / `grad_norm_slow_ema` updates further down in the same function (still receive the RAW `observed_grad_norm` per follow-up K's "fast/slow EMAs are stability signal that should respond to outliers" rationale), the `fold_warmup_factor_update` kernel that consumes the fast/slow EMAs, and every ISV slot. Pure formula change in one function — no new buffer, no new kernel, no new ISV slot, no new launch site, no graph recapture. F0 risk is low: F0-typical `grad_norm_slow_ema` is on the order of 1-10 → `upper_bound = 100-1000 × ISV[16].max(1)`; F0-typical `new_clip` (= `grad_norm_ema × 2`) is single-digits to low tens, well below the cap, so Mech 6 is invisible in normal F0 operation. F1+F2 benefit by preventing the EMA-ratchet pathology. Composes with: Mech 1 (target_q clipping at single-point source, B2), Mech 2 (atom-position growth bounds, B3), Mech 3 (CQL aux loss clamping, deferred), Mech 4 (Adam EMA reset comprehensive, B5), Mech 5 (fused threshold-check kernel for diagnostic slots 36-47, B6+B7). Zero new HtoD/DtoD/HtoH copies; zero new ISV slots; zero new buffers; zero new kernels. `cargo check -p ml --lib` clean.