fix(sp14): A.1 — C51 inv_a_std floor lift (1e-6 → 1e-3)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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:
|
||||
|
||||
Reference in New Issue
Block a user