From 617eb61aabf406f2208ea1a0921004c1db9383e6 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 13 May 2026 09:00:33 +0200 Subject: [PATCH] diagnostic(sp22): zero W prior + beta prior to isolate NaN source First smoke at ff98edc77 produced NaN in: - q_by_action (dir Qs from q_out_buf) - a_mag (mag SGEMM forward output) - v_share_traj + grad_decomp_pinned (all-branch backward) - WR = 43.40% (below ~46% baseline) Key clue: a_dir = -0.0029 finite but q_by_action NaN. Raw dir logits fine, expected Q (with atom-shift) NaN. So atom-shift produces NaN from finite inputs. Either W or state_121 contains NaN. This diagnostic zeroes both: - aux_w_prior_init: W [-0.5, 0, +0.5, 0] -> [0, 0, 0, 0] - beta scale prior: 0.5 -> 0.0 With W=0 + beta=0, atom-shift and beta are runtime no-ops. All wiring remains intact (kernels loaded, scratch populated, dW/Adam launched). Diagnostic outcomes: - Clean smoke -> bug was W magnitude x downstream interaction. Restore W at smaller magnitude. - Still NaN -> bug is in Step 8/11 backward logic. Investigate c51_aux_dw_kernel + adam_w_aux interaction. Cargo check clean. Ready for diagnostic smoke. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../cuda_pipeline/aux_w_prior_init_kernel.cu | 26 ++++++++++++++----- .../src/trainers/dqn/trainer/training_loop.rs | 14 +++++++--- docs/dqn-wire-up-audit.md | 23 ++++++++++++++++ 3 files changed, 53 insertions(+), 10 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/aux_w_prior_init_kernel.cu b/crates/ml/src/cuda_pipeline/aux_w_prior_init_kernel.cu index c8936358a..ed9494677 100644 --- a/crates/ml/src/cuda_pipeline/aux_w_prior_init_kernel.cu +++ b/crates/ml/src/cuda_pipeline/aux_w_prior_init_kernel.cu @@ -51,12 +51,26 @@ extern "C" __global__ void aux_w_prior_init( int a = threadIdx.x; if (a >= 4) return; - /* Per-action structural prior: - * a == 0 (Short) → -0.5 - * a == 1 (Hold) → 0.0 - * a == 2 (Long) → +0.5 - * a == 3 (Flat) → 0.0 + /* SP22 H6 Phase 3 α DIAGNOSTIC (2026-05-13): W zeroed to isolate + * the NaN source. First smoke at commit ff98edc77 with W=[-0.5,0,+0.5,0] + * showed NaN in q_by_action (dir Qs), a_mag (mag SGEMM output), and + * grad_decomp_pinned (all branches' backward gradients). Forward + * a_dir, v, a_ord, a_urg remained finite, ruling out atom-shift in + * compute_expected_q as the source. Hypothesis: either atom-shift + * magnitude × downstream interaction destabilises training, or + * Step 8/11 (c51_aux_dw + adam_w_aux) backward path corrupts W. + * + * With W = [0, 0, 0, 0]: + * - aux_atom_shift = W[a] * state_121 = 0 for all (a, state_121) + * - Equivalent to atom-shift inactive throughout + * - Step 8/11 still run, dW computed from SP (which depends only on + * state_121 + projection, not W), Adam updates W with this dW + * - If smoke is clean: bug was W magnitude × downstream interaction + * - If smoke NaNs: bug is in Step 8/11 logic (independent of W init) + * + * Original structural prior [-0.5, 0, +0.5, 0] will be restored once + * the NaN source is identified. */ - float prior_values[4] = { -0.5f, 0.0f, +0.5f, 0.0f }; + float prior_values[4] = { 0.0f, 0.0f, 0.0f, 0.0f }; w_aux_to_q_dir[a] = prior_values[a]; } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 173e0d80d..e6aa78dd4 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -10589,10 +10589,16 @@ impl DQNTrainer { "sp22_aux_align_scale" => { if let Some(ref fused) = self.fused_ctx { use crate::cuda_pipeline::sp22_isv_slots::SP22_AUX_ALIGN_SCALE_INDEX; - // Structural prior 0.5 — activates β with fixed magnitude - // from epoch 0. SP11 controller's Phase B6 extension will - // later overwrite with adaptive emit anchored on slot 536. - fused.trainer().write_isv_signal_at(SP22_AUX_ALIGN_SCALE_INDEX, 0.5); + // SP22 H6 Phase 3 α DIAGNOSTIC (2026-05-13): β scale + // ZEROED to deactivate β reward shaping. Earlier 0.5 + // structural prior caused NaN in smoke at ff98edc77 + // — to isolate the NaN source, deactivate β entirely + // alongside W=[0,0,0,0]. If diagnostic smoke is clean, + // bug was in α/β magnitude × downstream interaction. + // If still NaN, bug is in Step 8/11 (c51_aux_dw + + // adam_w_aux) backward logic. Original 0.5 prior to + // be restored once NaN source is identified. + fused.trainer().write_isv_signal_at(SP22_AUX_ALIGN_SCALE_INDEX, 0.0); } } _ => { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 16de58d2f..ad6b0fb2d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -17333,3 +17333,26 @@ Estimated ~1-2 hr engineering (single-kernel edit + launcher update if N_COMPONE **Once full B6 lands** (SP11 controller emits adaptive scale_β per bounds [0.05, 2.0]), the 0.5 prior is overwritten on first emit per `pearl_first_observation_bootstrap`. **Status update**: the previous "α ACTIVE β DEAD" scope clarification is superseded. After this commit, both α and β are active when the smoke re-runs at this commit. The current train-qsltr smoke (at parent commit 5106e3b117) still has β-dead since it predates this fix. + +#### Diagnostic — W and β zeroed to isolate NaN source (2026-05-13) + +**First smoke `train-th8pj` @ commit `ff98edc77` produced NaN**: +- `q_by_action [hold/long/short/flat = NaN]` (dir Qs from q_out_buf NaN) +- `a_mag = NaN` (mag SGEMM output NaN) but `a_dir = -0.0029` finite +- `v_share_traj` + `grad_decomp_pinned` all-branch NaN (backward path) +- WR = 43.40% (below ~46% baseline) +- Forward `a_ord = 0.0042`, `a_urg = 0.0029` remained finite + +**Key clue**: a_dir is finite but q_by_action is NaN. Raw dir advantage logits are fine, but expected Q (which adds atom-shift) is NaN. The atom-shift `W[a] * state_121` must be NaN — either W contains NaN (Adam writing NaN) or state_121 contains NaN (corrupted state from prior NaN propagation). + +**This commit (diagnostic)**: Zero out W structural prior AND β scale prior: +- `aux_w_prior_init_kernel.cu`: prior `[-0.5, 0, +0.5, 0]` → `[0, 0, 0, 0]`. atom-shift = W[a] * state_121 = 0 regardless of state_121. +- `training_loop.rs` reset arm: β scale prior `0.5` → `0.0`. β reward bonus is 0 regardless of alignment. + +With W=0 + β=0, atom-shift and β are runtime no-ops. The full Phase 3 wiring (kernels loaded, scratch buffers populated, c51_aux_dw + adam_w_aux launched, dispatch arms present) remains intact for testing. + +**Diagnostic outcomes**: +1. **Smoke clean (WR ≥ 46%, no NaN)** → bug was W magnitude × downstream interaction. Action: restore W prior at smaller magnitude (e.g., `[-0.05, 0, +0.05, 0]` instead of `[-0.5, 0, +0.5, 0]`). +2. **Smoke still NaN** → bug is in Step 8/11 (c51_aux_dw + adam_w_aux) backward logic, NOT atom-shift forward. Action: investigate the dW + Adam path for the specific NaN source (likely in c51_aux_dw_kernel reads of scratch buffers, or adam_w_aux interaction). + +Cargo check: 0 errors, 21 pre-existing warnings.