From 5f7db7d8d8d29f5d09ee7fc42ae91119253dd85b Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 30 Apr 2026 10:29:28 +0200 Subject: [PATCH] =?UTF-8?q?feat(dqn):=20SP3=20=E2=80=94=20name=20table=20e?= =?UTF-8?q?ntries=20for=20slots=2036-47=20diagnostic?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replaces rsv36-rsv47 placeholders with the SP3 Mech 5 diagnostic slot names. Both name table sites (halt_nan + halt_grad_collapse) updated identically per shared-contract migration principle (feedback_no_partial_refactor). Slot names mirror the audit doc table: - 36-39: Adam m max-abs (trunk/value/branch/IQN) - 40-43: Adam v max-abs (trunk/value/branch/IQN) - 44-45: Weight max-abs (trunk/heads) - 46: target_q post-clip - 47: atom_span_max When the fused kernel sets a slot bit, the readback log line names the buffer that exceeded its ISV-derived threshold — providing direct observability for SP3 mechanism effectiveness. --- .../src/trainers/dqn/trainer/training_loop.rs | 32 +++++++++++++++---- docs/dqn-wire-up-audit.md | 2 ++ 2 files changed, 28 insertions(+), 6 deletions(-) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 492d72faa..2f897ccbc 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -2022,9 +2022,19 @@ impl DQNTrainer { "bw_d_h_s2_post_main", // 33 bw_d_h_s2 after main backward chain "bw_d_h_s2_post_aux", // 34 bw_d_h_s2 after aux head SAXPY "bw_d_h_s2_post_iqn", // 35 bw_d_h_s2 after IQN DtoD - "rsv36", "rsv37", "rsv38", "rsv39", // 36-39 SP2/SP3 headroom - "rsv40", "rsv41", "rsv42", "rsv43", // 40-43 SP2/SP3 headroom - "rsv44", "rsv45", "rsv46", "rsv47", // 44-47 SP2/SP3 headroom + // SP3 Mech 5 (slots 36-47): Adam-centric diagnostic threshold checks. + "trunk_adam_m_max", // 36 + "value_adam_m_max", // 37 + "branch_adam_m_max", // 38 + "iqn_adam_m_max", // 39 + "trunk_adam_v_max", // 40 + "value_adam_v_max", // 41 + "branch_adam_v_max", // 42 + "iqn_adam_v_max", // 43 + "trunk_weight_max", // 44 + "heads_weight_max", // 45 + "target_q_post_clip", // 46 + "atom_span_max", // 47 ]; let flagged: Vec = flags.iter().enumerate() .filter(|(_, &f)| f != 0) @@ -2109,9 +2119,19 @@ impl DQNTrainer { "bw_d_h_s2_post_main", // 33 bw_d_h_s2 after main backward chain "bw_d_h_s2_post_aux", // 34 bw_d_h_s2 after aux head SAXPY "bw_d_h_s2_post_iqn", // 35 bw_d_h_s2 after IQN DtoD - "rsv36", "rsv37", "rsv38", "rsv39", // 36-39 SP2/SP3 headroom - "rsv40", "rsv41", "rsv42", "rsv43", // 40-43 SP2/SP3 headroom - "rsv44", "rsv45", "rsv46", "rsv47", // 44-47 SP2/SP3 headroom + // SP3 Mech 5 (slots 36-47): Adam-centric diagnostic threshold checks. + "trunk_adam_m_max", // 36 + "value_adam_m_max", // 37 + "branch_adam_m_max", // 38 + "iqn_adam_m_max", // 39 + "trunk_adam_v_max", // 40 + "value_adam_v_max", // 41 + "branch_adam_v_max", // 42 + "iqn_adam_v_max", // 43 + "trunk_weight_max", // 44 + "heads_weight_max", // 45 + "target_q_post_clip", // 46 + "atom_span_max", // 47 ]; let flagged: Vec = flags.iter().enumerate() .filter(|(_, &f)| f != 0) diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index e82cd06ac..7882e9361 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2282,3 +2282,5 @@ SP3 Task B1 — slot 36-47 diagnostic accessors (2026-04-29): added 24 accessor SP3 Task B5 — comprehensive Adam EMA reset (2026-04-29): audited every Adam optimizer state buffer in `crates/ml/src/` (grep over `adam_m`/`adam_v`/`m_buf`/`v_buf`/`reset_adam_state` across the crate) against the existing 5 `reset_adam_state` calls in `fused_training.rs::reset_for_fold` (DQN main at line 971 → `GpuDqnTrainer::reset_adam_state` which covers `m_buf`/`v_buf`, `iqn_trunk_*`, `q_attn_adam_*`, `sel_adam_*`, `denoise_adam_*`, `mamba2_adam_*`, `ofi_embed_adam_*`, plus PopArt + Q-div EMA; IQN head at line 1011; TLOB at line 1026; IQL-high at line 1031; IQL-low at line 1036). VSN params and aux next_bar/regime heads (slots 119-126) live inside the unified `m_buf`/`v_buf`, so they are covered by the main DQN reset already. **One missing optimizer found**: `GpuAttention` (4-head feature attention on `h_s2`, "always active" per the constructor comment at `fused_training.rs:676`) owns its own `(attn_m, attn_v, attn_adam_step)` tuple — called every training step via `attn.adam_step` at `fused_training.rs:1962/1995` but never zeroed at fold boundary. Same pathology as IQN/TLOB/IQL: fold N momentum oversizes fold N+1's first-epoch SDP + output projection updates, compounding through the trunk gradient. **Fix landed**: added `GpuAttention::reset_adam_state` mirroring `gpu_iqn_head` / `gpu_tlob` pattern (memset_zeros m/v + zero step counter + write 0 through pinned `t_pinned`); wired under `if let Some(ref mut attn) = self.gpu_attention` immediately after the IQL-low reset with the same warn-on-failure / info-on-success log pattern. Fold-boundary Adam reset coverage is now comprehensive across every optimizer in `FusedTrainingCtx` per `feedback_no_partial_refactor`. Out-of-scope optimizers (documented, not changed): `DecisionTransformer` (scoped within DT pretrain block, dropped at end of pretrain — no fold leak); `GpuCuriosityTrainer` (owned by `GpuExperienceCollector`, external to `FusedTrainingCtx` — reset belongs at the collector layer if needed); `MetaQNetwork` (host-side observability MLP for collapse prediction, does not feed `LearningHealth` or trunk gradients); `GpuMoeHead` (forward-only, no Adam state of its own — uses params from main DQN). Zero new ISV slots; zero new buffers; zero new kernels. `cargo check -p ml --lib` clean. SP3 Task B6 — fused kernel extended for slots 36-47 threshold checks (2026-04-30): extended `dqn_nan_check_fused_f32_kernel` (SP2 Task A4) from 12-slot NaN-only coverage (slots 24-35) to 24-slot NaN + threshold coverage (slots 24-47). Kernel signature gained one f32 by-value arg `q_abs_ref_eff` (passed as `f32` to match the kernel's `float` param per `feedback_cudarc_f64_f32_abi`); per-slot threshold computed inline via the relative `slot = blockIdx.x` index. Slots 0-11 (= absolute 24-35) keep `thr = INFINITY` so the magnitude check is trivially false (NaN-only behaviour preserved bit-exactly). Slots 12-15 (= absolute 36-39) check Adam m max-abs ≥ `100 × q_abs_ref_eff`; 16-19 (= 40-43) check Adam v max-abs ≥ `1e6 × q_abs_ref_eff²`; 20-21 (= 44-45) check weight max-abs ≥ `1e3 × q_abs_ref_eff`; 22 (= 46) checks `target_q` post-clip ≥ `95 × q_abs_ref_eff` (= `9.5 × max_abs_target_q` where `max_abs_target_q = 10 × q_abs_ref_eff` per B2); 23 (= 47) checks atom-positions span ≥ `190 × q_abs_ref_eff` (= `9.5 × max_atom_abs × 2` where `max_atom_abs = 10 × q_abs_ref_eff` per B3). Threshold logic is inline `if (slot >= 12 && slot < 16) thr = 100*q...` — no per-slot threshold buffer, no per-step HtoD. `q_abs_ref_eff = max(ISV[Q_ABS_REF=16], 1.0)` computed host-side at launch via `read_isv_signal_at` and floored at 1.0 per the SP1 ε-floor pearl (cold-start ISV ~0 yields q_abs_ref_eff = 1, so all thresholds floor at their ε-floor multiplier × 1). **Buffer + populate + launch wiring**: `nan_check_buf_ptrs` / `nan_check_buf_lens` resized 12 → 24 entries (mapped-pinned host write at construction; no HtoD copy per `feedback_no_htod_htoh_only_mapped_pinned`); `populate_nan_check_meta` extended with 4 new args `iqn_adam_m_ptr / iqn_adam_m_len / iqn_adam_v_ptr / iqn_adam_v_len: Option/Option` (None when IQN inactive — entry becomes (0, 0), kernel's null-pointer guard skips the block; symmetric with slots 27/28); 12 new entries appended for slots 36-47 calling the SP3 Task B1 accessors (`trunk_adam_{m,v}_{ptr,len}`, `value_adam_{m,v}_{ptr,len}`, `branch_adam_{m,v}_{ptr,len}`, `trunk_params_{ptr,len}`, `heads_params_{ptr,len}`, `target_q_{ptr,len}`, `atom_positions_{ptr,len}`); caller `fused_training.rs::FusedTrainingCtx::new` updated to pass the IQN Adam ptrs via `gpu_iqn.as_ref().map(|h| h.adam_m_ptr())` etc. **Launch wrapper** `launch_nan_check_fused_f32` reads `q_abs_ref` via `read_isv_signal_at(Q_ABS_REF_INDEX)`, computes `q_abs_ref_eff: f32 = q_abs_ref.max(1.0)`, and grows the grid from 12 → 24 blocks; the 5-arg launch builder now passes `q_abs_ref_eff` between `buf_lens_dev` and `BASE_FLAG_IDX`. Sticky-flag semantics preserved (kernel writes 1 only). Closes the diagnostic instrumentation loop for SP3 Mech 5 — slots 36-47 fire when their threshold is exceeded, providing observability for the other 4 SP3 mechanisms' effectiveness; if the SP3 fix doesn't fully resolve F1 NaN, the slot 36-47 firing pattern guides the next iteration. Zero new HtoD/DtoD/HtoH copies (one f32 by-value kernel arg = register pressure only, identical to existing `BASE_FLAG_IDX`). Zero new ISV slots (reuses `Q_ABS_REF_INDEX = 16`). Zero new buffers (per-slot threshold computed inline). Single launch covers all 24 slots; null-pointer guard handles slots 27/28/31/33-35/39/43 (deferred or IQN-inactive) without per-step Rust branching. `cargo check -p ml --lib` clean. + +SP3 Task B7 — name-table entries for slots 36-47 readback log (2026-04-30): replaced the SP1 Phase B placeholder strings (`"rsv36"`-`"rsv47"`) in both `training_loop.rs` `let names = [...]` tables (the `halt_nan` block ~L1992 and the `halt_grad_collapse` block ~L2089) with the SP3 Mech 5 diagnostic slot names per the B1/B6 accessor allocation: 36 `trunk_adam_m_max`, 37 `value_adam_m_max`, 38 `branch_adam_m_max`, 39 `iqn_adam_m_max`, 40 `trunk_adam_v_max`, 41 `value_adam_v_max`, 42 `branch_adam_v_max`, 43 `iqn_adam_v_max`, 44 `trunk_weight_max`, 45 `heads_weight_max`, 46 `target_q_post_clip`, 47 `atom_span_max`. Both name tables receive byte-identical replacement content (modulo the indentation difference between the two enclosing blocks) per `feedback_no_partial_refactor` — the post-SP2 stale-doc cleanup commit `387335e2b` already established that the two tables must remain identical, and the same shared-contract migration principle applies here. When the fused kernel (B6) sets a slot bit, the readback log line names the buffer that exceeded its ISV-derived threshold (e.g. `flagged=[42=branch_adam_v_max, 46=target_q_post_clip]`) instead of the opaque `rsv*` placeholder, providing direct observability for SP3 mechanism effectiveness. Pure name-string replacement — no logic changes, no kernel changes, no buffer changes, no ISV changes. `cargo check -p ml --lib` clean.