fix(aux-head): stop-gradient on aux's h_s2 input — fixes aux-loss-rises-during-training pathology

Root cause from 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)

Aux head's backward gradient was flowing back to shared trunk activation
h_s2 via dh_s2_out write at aux_heads_kernel.cu:599-613. 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.

Fix: stop-gradient. Aux reads h_s2 via forward, trains its own w1/b1/w2/b2
from CE loss, but does NOT propagate to h_s2. Q-loss is 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).

This was the SEVENTH fix in today's chain (after 6 SP14 EGF cadence/
gate/saturation fixes). The EGF was a scaffold over a broken aux head;
fixing aux first is the architectural prerequisite for EGF to route
useful signal.

Verification: cargo check clean; sp14_oracle_tests 7/7 pass.
Validation: aux next_bar_mse should now DECREASE during training in
the next L40S smoke (vs the rising-from-0.35-to-0.72 pattern in v8ztm).

Deferred follow-up: aux_regime_backward has the same architecture
(propagates dh_s2 to trunk). Same fix is a candidate once next_bar
result validates the approach.

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

crates/ml/src/cuda_pipeline/aux_heads_kernel.cu

@@ -596,19 +596,43 @@ extern "C" __global__ void aux_next_bar_backward(
         }
     }
 
-    /* Step 3: dh_s2[b, j] = sum_k sh_dh_pre[k] * w1[k, j].
-     * Per-sample, per-feature serial reduction over H lanes — H is small
-     * (32) and SH2 is moderate; stride loop over j. Masked rows propagate
-     * zero d_h_pre through this sum, so dh_s2 row is zero for skipped
-     * samples (the trunk SAXPY adds zero — no spurious gradient). */
+    /* Step 3: SP14 stop-gradient (2026-05-07).
+     *
+     * Diagnostic from L40S `train-v8ztm` 9-epoch HEALTH_DIAG aux next_bar_mse:
+     *   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)
+     *
+     * Root cause: the aux 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. 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.
+     *
+     * Fix: stop-gradient on aux's input. We zero `dh_s2_out` rather than
+     * computing the sum-of-w1·dh_pre. The aux head still trains its own
+     * params (w1/b1/w2/b2) from the CE loss via the dW1/db1/dW2/db2
+     * partials written above — those are unaffected. But h_s2 is no longer
+     * pulled by aux's gradient. Q-loss is the sole shaping force on h_s2;
+     * aux must adapt to whatever h_s2 happens to be.
+     *
+     * If aux still cannot learn after this fix, the deeper diagnosis is that
+     * h_s2 carries no direction signal at all — separate-trunk Option 2
+     * (deferred). The trunk SAXPY adds zero either way (same arithmetic the
+     * masked-row branch already produced for skipped samples), so no other
+     * consumer is affected — `dh_s2_out` is read only by the trunk
+     * 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. */
     {
         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;
         }
     }
 }