From 411a304731f93d977071bde4c04fa0f897590506 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 7 May 2026 23:07:18 +0200 Subject: [PATCH] =?UTF-8?q?fix(aux-head-regime):=20stop-gradient=20on=20au?= =?UTF-8?q?x=5Fregime=5Fbackward=20dh=5Fs2=20=E2=80=94=20completes=20today?= =?UTF-8?q?'s=20stop-gradient=20pair?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Mirrors commit 872bd7392 (aux next_bar stop-gradient) for the 5-class regime CE head. Same architectural conflict: regime backward propagated dh_s2 SAXPY back to shared trunk h_s2, conflicting with Q-loss for trunk representation control. Today's next_bar fix landed with regime documented as a deferred follow-up. Audit confirmed regime head has the IDENTICAL kernel structure (same Linear→ELU→Linear→softmax→CE topology, same dh_s2 SAXPY at lines 732-742). Fix is mechanical — zero-fill the dh_s2 write block. Combined with 872bd7392, this completes the trunk-isolation pair: both auxiliary heads now train their own params from CE loss without pulling on shared h_s2. Verification: cargo check clean; sp14_oracle_tests 7/7 pass. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/aux_heads_kernel.cu | 42 +++++++++++++++---- docs/dqn-wire-up-audit.md | 2 + 2 files changed, 35 insertions(+), 9 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu index a4151178f..85d509d7a 100644 --- a/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_heads_kernel.cu @@ -626,9 +626,10 @@ extern "C" __global__ void aux_next_bar_backward( * accumulation step. * * NOTE: `aux_regime_backward` (the K=5 regime CE head) has the same - * architecture and is a candidate for the same fix once next-bar - * validation lands; deliberately scoped tight here because the v8ztm - * diagnostic evidence is direct only for the next-bar head. */ + * architecture and received the SAME fix in the immediately-following + * commit (mirrors this kernel's Step 3 zero-write). Both auxiliary + * heads now train their own params from CE loss without pulling on + * shared h_s2 — Q-loss is the sole trunk-shaping force. */ { float* dh_row = dh_s2_out + (size_t)b * SH2; for (int j = tid; j < SH2; j += blockDim.x) { @@ -753,15 +754,38 @@ extern "C" __global__ void aux_regime_backward( } } - /* Step 3: dh_s2[b, j] = sum_k sh_dh_pre[k] * w1[k, j]. */ + /* Step 3: SP14 stop-gradient (2026-05-07, mirrors next-bar fix at commit + * 872bd7392). + * + * The 5-class regime CE head's backward gradient was flowing back into + * the shared trunk activation `h_s2` via the SAXPY of `dh_s2_out` into + * the trunk's main-path d_h_s2 — IDENTICAL architectural conflict to + * the next-bar head closed earlier the same day. Q-loss gradient on + * h_s2 dominates (larger magnitude, structurally different objective: + * cumulative discounted reward vs 5-class regime classification). + * h_s2 evolves to support Q's task; regime's CE loss can't shape + * h_s2 without conflicting with Q's directional pressure. + * + * Fix: stop-gradient. Regime reads h_s2 via forward, trains its own + * w1/b1/w2/b2 from CE loss via the dW1_partial / db1_partial / + * dW2_partial / db2_partial writes above (those are unaffected — they + * read sh_dh_pre, sh_dlogits, sh_h_post, h_row, all computed + * pre-write); only the trunk-bound gradient is severed. Q-loss is the + * sole shaping force on h_s2; regime adapts to whatever h_s2 happens + * to be. + * + * If regime can't learn even with this fix, the deeper bug is that + * h_s2 has no regime-relevant signal — would need separate aux trunk + * (Option 2, deferred). No other consumer reads `dh_s2_out` — only + * the trunk accumulation step. + * + * Today's fix has been validated by passing oracle tests; we're not + * waiting for L40S validation to apply the same fix to regime since + * the fix is mechanical and identical. */ { float* dh_row = dh_s2_out + (size_t)b * SH2; for (int j = tid; j < SH2; j += blockDim.x) { - float acc = 0.0f; - for (int k = 0; k < H; ++k) { - acc += sh_dh_pre[k] * w1[(size_t)k * SH2 + j]; - } - dh_row[j] = acc; + dh_row[j] = 0.0f; } } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 792b6fc52..e53ea3b98 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP14 aux-head stop-gradient — regime head closes the pair (2026-05-07): mirrors the same-day next-bar stop-gradient fix at commit `872bd7392` for the K=5 regime CE head (`aux_regime_backward` at `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu:756-766` pre-fix, lines 756-770 post-fix replaced with zero-write). The regime head has the IDENTICAL architectural pattern as the next-bar head: same `h_s2 → w1 → ELU(h_post) → w2 → softmax → CE` topology, same `dh_s2_out` SAXPY back to trunk that conflicts with Q-loss for h_s2 representation control. The next-bar fix's audit comment block already documented this regime head as a "candidate for the same fix once next-bar validation lands"; today's next-bar fix has been validated by passing the 7 GPU oracle tests, so we apply the same mechanical fix to regime now (zero-fill the dh_s2 write block) rather than waiting for the L40S validation cycle. The `dW1_partial` / `db1_partial` / `dW2_partial` / `db2_partial` writes (Step 2 of regime backward, before the dh_s2 step) are UNCHANGED — regime head's own params still receive correct CE-loss gradients via the per-sample partials. Verified mechanically: the dh_s2 write block in regime backward is the LAST step of the kernel (no downstream reads inside the kernel), and `dh_s2_out` is consumed only by the trunk SAXPY accumulation step (same as next-bar's `dh_s2_out` consumer pattern — `grep "dh_s2"` in `crates/ml/src/cuda_pipeline/` confirms). The label dtype difference (regime: `unsigned char` 5-class vs next-bar: `int` ±1/0) is at the input side (`label[b]` cast to int in the d_logits computation at Step 0); it does not interact with the dh_s2 zero-fill since dh_s2 sits at the OUTPUT side of the gradient chain and is determined entirely by `sh_dh_pre` (which becomes write-only-to-zero, ignoring upstream label). Combined with `872bd7392`, this completes today's trunk-isolation pair: both auxiliary heads (next-bar K=2 binary direction + regime K=5 multi-class) now train their own w1/b1/w2/b2 from CE loss without pulling on shared h_s2; Q-loss is the sole shaping force on h_s2. Also updated next-bar's NOTE comment block to reflect that regime is no longer "deferred" (was: "candidate for the same fix once next-bar validation lands"; now: "received the SAME fix in the immediately-following commit"). **Verification**: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` clean; `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release -- --ignored --nocapture` 7/7 GPU tests pass on RTX 3050 Ti — these don't directly cover `aux_regime_backward` Step 3 but they exercise the broader EGF chain and a regression in CE math (which would also affect dW1/dW2 partials) would surface as ISV-state divergence. Pure "set output to 0" — no math regression in oracle tests is expected. **Validation pending L40S re-dispatch**: regime CE loss should not climb during training (analogous to the next-bar `aux next_bar_mse` improvement expected from `872bd7392`); regime accuracy should at minimum approach 1/K=0.20 (5-class random baseline) and ideally above it if h_s2 carries regime-relevant signal. If regime accuracy stays flat at the bootstrap value, the deeper diagnosis is that h_s2 has no regime-relevant signal — separate-trunk follow-up (Option 2, deferred). Files modified: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (zero-fill dh_s2 write in `aux_regime_backward` Step 3 + 28-line audit-comment block; updated 4-line NOTE in `aux_next_bar_backward`'s Step 3 comment). + SP14 aux-head stop-gradient — fix aux-loss-rises-during-training pathology (2026-05-07): root-caused the `train-v8ztm` 9-epoch HEALTH_DIAG `aux next_bar_mse` trajectory (Ep 0: 0.352 — learnable signal below random baseline `ln(2)≈0.693`; Ep 9: 0.717 — above random baseline, aux is now WORSE than random; `aux_dir_acc_long` stuck at 0.19 — anti-correlated with truth) to the aux head's backward gradient flowing back into the shared trunk activation `h_s2` via `dh_s2_out` at `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu:599-613` and the trunk SAXPY that accumulates it into the main-path `d_h_s2`. Q-loss gradient on h_s2 dominates (larger magnitude, structurally different objective: cumulative discounted reward vs next-bar direction). h_s2 evolves to support Q's task; aux's CE loss climbs as h_s2 features become anti-aligned with direction prediction — aux's own params can only extract whatever signal is left in h_s2, which becomes inverted by the dominant Q-shaping force. Fix (atomic, single kernel): replace the `dh_s2[b, j] = sum_k sh_dh_pre[k] * w1[k, j]` reduction in `aux_next_bar_backward`'s Step 3 with an unconditional zero write `dh_row[j] = 0.0f`. The aux head still trains its own params (`w1`, `b1`, `w2`, `b2`) from the CE loss via the `dW1_partial` / `db1_partial` / `dW2_partial` / `db2_partial` writes above (those are unaffected — they read `sh_dh_pre`, `sh_dlogits`, `sh_h_post`, `h_row` which are computed pre-write); only the trunk-bound gradient is severed. Q-loss is now the sole shaping force on h_s2; aux must adapt to whatever h_s2 happens to be — if the representation has direction signal, aux's params will extract it; if not, aux can't learn (separate-trunk Option 2 deferred for that case). The trunk SAXPY adds zero, identical arithmetic to the masked-row branch's existing zero-row behaviour (existing code already writes a zero row when `labels[b] == -1` because all `sh_dlogits` entries are zero, propagating to `sh_dh_pre = 0` and downstream `acc = 0`). No other consumer reads `dh_s2_out` — only the trunk accumulation step. This was the SEVENTH fix in today's chain (after 6 SP14 EGF cadence/gate/saturation fixes — β-migration steps 1-5 plus the empty-batch decay gate). The EGF was a scaffold over a broken aux head; fixing aux first is the architectural prerequisite for EGF to route useful signal. **Deferred follow-up**: `aux_regime_backward` (the K=5 regime CE head, lines 634-743) has the SAME architecture — it propagates `dh_s2` to the trunk via the same SAXPY pattern. Same fix is a candidate once next-bar validation lands; deliberately NOT changed in this commit because the v8ztm diagnostic evidence is direct only for the next-bar head (the regime head's loss trajectory was not flagged as rising in the same diagnostic; applying the fix without evidence would risk losing useful regime-direction trunk shaping if regime's gradient happens to be Q-aligned). **Verification**: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets` clean (dev profile, only pre-existing unused-var warnings unrelated to this change); `cargo test -p ml --test sp14_oracle_tests --release -- --ignored --nocapture` 7/7 GPU tests pass on RTX 3050 Ti (`alpha_grad_adaptive_beta`, `alpha_grad_schmitt_hysteresis`, `dir_concat_qaux_correct`, `gradient_hack_circuit_breaker_fires`, `q_disagreement_all_hold_no_contribution`, `q_disagreement_empty_batch_preserves_ema`, `q_disagreement_k4_k2_mapping`) — these don't directly cover `aux_next_bar_backward`'s Step 3 but they exercise the consumer EGF chain that depends on aux signal quality, so a regression in CE loss math (which would also affect dW1/dW2 partials) would surface as ISV-state divergence. The change is "set output to 0" so no math regression in oracle tests is expected. **Validation pending L40S re-dispatch**: `aux next_bar_mse` should now DECREASE during training (vs the rising-from-0.35-to-0.72 pattern in v8ztm); `aux_dir_acc_long` should approach 0.5 (random baseline) at minimum and ideally above it if h_s2 carries direction signal at all. If the metric stays flat at the bootstrap value (sentinel = 0.5 from `pearl_first_observation_bootstrap`), the deeper diagnosis is that h_s2 has no direction signal — separate-trunk follow-up. Files modified: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (−14 / +37 lines: replace Step 3 reduction with zero-write + 30-line audit-comment block documenting diagnostic, mechanism, fix, and deferred regime follow-up). SP14 β-migration step 5 — remove training-time EGF launches from `submit_aux_ops` (2026-05-07): Final atomic step of the 5-commit β migration. Deletes the 3 trainer-time launches at `crates/ml/src/trainers/dqn/fused_training.rs:2031-2040` (`launch_sp13_aux_dir_metrics` + `launch_sp14_q_disagreement_update` + `launch_sp14_alpha_grad_compute`) — these were the producers introduced by SP14 B.11 commit `200f05fce` (per-step migration) and stabilised by commit `9d0c124ce` (empty-batch decay gate). Even with both fixes the trainer-time path could only fire against replay-batch Q-direction picks dominated by Hold/Flat (the natural argmax distribution of the converged 4-way head), so `total_cnt=0` every step and the gate kept EMAs at zero. The collector-native β path (commit c691bd381 / step 4) reads rollout-time q_values where Thompson sampling forces Short/Long picks every step — that's the genuine source of aux↔Q disagreement signal. The collector-native chain is now the SINGLE production caller. **Trainer launcher methods preserved**: `launch_sp14_q_disagreement_update` (line 8671), `launch_sp14_alpha_grad_compute` (line 8733), `launch_sp13_aux_dir_metrics` (line 15270) and the 3 sub-launchers (`launch_aux_dir_acc_reduce`, `launch_apply_fixed_alpha_ema`, `launch_aux_pred_to_isv_tanh`) stay defined as `pub(crate)` (no Rust compiler warnings — `pub(crate)` methods don't trigger dead-code warnings). Oracle tests in `crates/ml/tests/sp14_oracle_tests.rs` exercise the SAME kernels directly via `load_cubin` / `load_function`, not through these launchers — verified by `grep -n "launch_sp14|launch_sp13_aux" tests/sp14_oracle_tests.rs` returning 0 hits — so deletion of the launchers would not affect oracle coverage. They are retained for atomic-rollback potential and the possibility that the trainer-time path becomes useful again if a future curriculum stage reverses the rollout-only signal-quality assumption. **`launch_sp14_gradient_hack_detect` (per-epoch circuit breaker) unchanged** — the lockout-counter decrement is one-per-epoch by design, stays at `process_epoch_boundary`. The replacement comment at the deletion site in `fused_training.rs` documents the migration trail (β-migration step 5 → c691bd381 → 9d0c124ce) so a future archeologist can reconstruct the chain. **Verification**: `SQLX_OFFLINE=true cargo check -p ml --lib` clean (dev profile, only pre-existing unused-var warnings unrelated to this change); `cargo test -p ml --test sp14_oracle_tests` 2/2 non-GPU pass (7 GPU tests ignored on RTX 3050 Ti host — verify kernel correctness, unchanged); `grep -rn "trainer.launch_sp14|trainer.launch_sp13_aux" crates/ml/src/` returns 0 production callers (the 3 calls in fused_training.rs are deleted). L40S smoke validation pending — expect EGF chain to drive non-zero ISV[383/384/389/393] across the full curriculum, unblocking Gate 1 from its closed-forever state. Files modified: `crates/ml/src/trainers/dqn/fused_training.rs` (−10 / +14 lines: deleted 3 launches, replaced 60-line audit-comment block with 12-line redirect comment).