From f75786fc5afc004e83e7502fb375e398a5f1ba3f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 5 May 2026 18:09:21 +0200 Subject: [PATCH] =?UTF-8?q?fix(sp14):=20A.1=20=E2=80=94=20C51=20inv=5Fa=5F?= =?UTF-8?q?std=20floor=20lift=20(1e-6=20=E2=86=92=201e-3)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit c51_grad_kernel.cu line 275: lift floor from 1e-6 to 1e-3 in \`inv_a_std = 1.0f / (a_std + 1e-3f)\`, capping the magnitude-branch gradient amplifier at 1000 instead of ~1e6 in the degenerate case. Why: Smoke A produced 1109 GRAD_CLIP_OUTLIER events with C51 grad reaching 9.5e6 — the SP7 budget controller saturated at the EPS_DIV floor instead of rebalancing proportionally. Phase-0 verification against the actual kernel found the amplifier is NOT the spec's claimed −log(p)/p divide (which does not exist; the kernel uses the CE-stable expf(lp) - proj form at line 81). The actual amplifier is inv_a_std = 1/(a_std + 1e-6) at line 275, gated by \`if (d == 1) grad_val *= inv_a_std\` at line 282. When magnitude advantage logits collapse near-uniform (Smoke A: var_q=9e-10), a_std → ~1e-9, so inv_a_std → ~1e6. Per feedback_isv_for_adaptive_bounds, this is a numerical-stability anchor (Invariant 1: prevent division-by-near-zero amplification), not a behavioural bound. ISV-driven bounds govern behavior; the existing 1e-12f floor on a_std at line 274 is also a structural anchor — same class of fix. Validation gate: Smoke A2-A GRAD_CLIP_OUTLIER count <100 in fold 2 (was 1109 pre-fix). The 3-order-of-magnitude reduction in worst- case amplification should bring C51 grad spikes back under SP7 budget controller authority. Audit doc: Fix 40 added (parallel to Fix 39 for A.2 and Fix 41 for A.3); the stale "A.1 deferred" note was removed in the A.3 commit. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/c51_grad_kernel.cu | 10 ++++- docs/dqn-wire-up-audit.md | 45 +++++++++++++++++++ 2 files changed, 54 insertions(+), 1 deletion(-) diff --git a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu index 1130cc335..1b4c502c0 100644 --- a/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_grad_kernel.cu @@ -272,7 +272,15 @@ extern "C" __global__ void c51_grad_kernel( sq_sum += v * v; } float a_std = sqrtf(sq_sum / (float)A1 + 1e-12f); - inv_a_std = 1.0f / (a_std + 1e-6f); + /* SP14 A.1: lift floor 1e-6 → 1e-3 to cap inv_a_std at 1000 + * (was ~1e6 when magnitude advantage logits collapse near- + * uniform — Smoke A captured var_q≈9e-10 producing 1109 + * GRAD_CLIP_OUTLIER events with C51 grad reaching 9.5e6). + * Numerical-stability anchor (Invariant 1: prevent division- + * by-near-zero amplification), not an ISV-driven behavioural + * bound; same class of fix as the existing 1e-12f a_std floor + * above. */ + inv_a_std = 1.0f / (a_std + 1e-3f); } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3229d5a06..6656726ce 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -6164,6 +6164,51 @@ Layer A is a small pre-Smoke-A2-A sweep addressing aux_w controller pathologies - `SQLX_OFFLINE=true cargo check --workspace` — clean - `SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests set_aux_weight_clamp_range` — passes +### Fix 40 — SP14 A.1: C51 `inv_a_std` floor lift (2026-05-05) + +**Problem.** Smoke A surfaced 1109 `GRAD_CLIP_OUTLIER` events in fold 2 with C51 grad reaching 9.5e6 — far above the SP7 budget controller's effective range. The plan §A.1 (writeup at lines 113–308 of `docs/superpowers/plans/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md`) initially attributed this to a `−log(p)/p · ∂p/∂z` divide; Phase-0 verification against `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` at HEAD `037c24116` proved that mechanism does not exist. Line 81 of the kernel uses the CE-stable form `d_combined = inv_batch * isw * (expf(lp) - proj)` — no `−log(p)/p` divide present. + +The actual amplifier is at lines 274–282 in the magnitude-branch d==1 code path: + +```cuda +float a_std = sqrtf(sq_sum / (float)A1 + 1e-12f); +inv_a_std = 1.0f / (a_std + 1e-6f); // pre-fix: floor 1e-6 → cap 1e6 +... +if (d == 1) grad_val *= inv_a_std; // applies amplifier in mag branch +``` + +When the magnitude advantage logits collapse near-uniform (Smoke A captured `var_q ≈ 9e-10`, `q_full → q_half → q_quarter` flat to 0.005), `a_std → ~1e-9` and `inv_a_std → ~1e6`. That 1e6 multiplicative factor is what produces the 9.5e6 grad spikes in the magnitude-branch d==1 path. + +This is a misdiagnosis-corrected fix: the writeup-claimed `−log(p)/p` cause was never in the kernel; the cause is the magnitude-branch standardization-derivative chain, surfacing only when forward-pass advantage collapse drives `a_std` to numerical-floor territory. + +**Smoke A trace.** 1109 `GRAD_CLIP_OUTLIER` events in fold 2; C51 grad max 9.5e6; SP7 budget controller saturated at the `EPS_DIV` floor instead of rebalancing proportionally (it could not — the per-loss budget mechanism cannot soak a 6-order-of-magnitude spike originating inside a single per-thread kernel write). + +**Fix.** One-line floor lift at line 275 of `c51_grad_kernel.cu`: + +```cuda +inv_a_std = 1.0f / (a_std + 1e-3f); // post-fix: floor 1e-3 → cap 1000 +``` + +Caps `inv_a_std` at 1000 (vs ~1e6 pre-fix) — a 3-order-of-magnitude reduction in the worst-case amplification. Justification per `feedback_isv_for_adaptive_bounds`: + +- The `1e-3` is a numerical-stability anchor (Invariant 1: prevent division-by-near-zero amplification), not a tuned behavioural bound. ISV-driven bounds govern *behavior*; structural numerical safety floors are constants. +- The existing `1e-12f` floor on `a_std` itself (line 274) is in the same class — kernel-internal numerical safety constant, no ISV slot. +- A behavioural alternative (`max(1e-7, atom_pos_p99 * 1e-3)` proposed in the original spec) would add a new ISV slot read in a hot per-batch kernel without addressing a behavioural problem; the issue is purely numerical safety. + +**Wiring contract.** Kernel-only change (no ISV slots, no host code, no graph-capture invalidation, no contract change for callers). The `inv_a_std` is computed and consumed within the same thread inside the same kernel; callers see no change. Cubin recompiles automatically via `build.rs`. + +**Files changed.** +- `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` (single-line edit + comment) — `inv_a_std = 1.0f / (a_std + 1e-3f)` at line 275. +- `docs/dqn-wire-up-audit.md` — this section. + +**Build + test verification.** +- `SQLX_OFFLINE=true cargo check --workspace` — clean (the .cu change rebuilds the cubin via `build.rs`). +- No host-side oracle test added: a CPU-side test of "after the fix, the kernel caps the amplifier at 1000" would be tautological — it would just re-implement the kernel's clamp on the host. The Smoke A2-A `GRAD_CLIP_OUTLIER` count is the validation gate. + +**Validation gate (Smoke A2-A).** +- `GRAD_CLIP_OUTLIER` count drops from 1109 (Smoke A fold 2) to <100 in fold 2 of Smoke A2-A. The 3-order-of-magnitude reduction in worst-case amplification should bring C51 grad spikes back under SP7 budget controller authority. +- If `GRAD_CLIP_OUTLIER` count remains high but max C51 grad is now in the 1e3–1e4 range (vs 1e6 pre-fix), the fix is working but the magnitude-branch advantage collapse itself is a separate signal-quality issue (likely Layer B's aux→Q wire or the EGF pearl that motivated SP14 Layer B). That diagnosis routes to SP14 Layer B, not back to A.1. + ### Fix 41 — SP14 A.3: stagnation warmup gate at fold boundary (2026-05-05) **Problem.** `compute_aux_w_p0b` in `training_loop.rs` (the SP13 P0b deficit + stagnation controller) was firing the stagnation decay term inappropriately at every fold reset. The mechanism is interaction between two correct subsystems producing an incorrect outcome: