diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index 1d2a8ead6..9beda1e96 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -1748,6 +1748,17 @@ extern "C" __global__ void popart_normalize_robust( // `dqn_adam_update_kernel`). Same `1e3 × q_abs_ref_eff` weight-slice // threshold as slots 44-45 — IQN online_params share the trunk weight regime. // +// SP3 Mech 10 extension: grid extended 25 → 26 blocks. Relative slot 25 +// (= absolute 49) covers post-trunk h_s2 activation max-abs — the regression +// sentinel for Mech 10's ISV-driven activation clamp. Mech 10 clamps h_s2 at +// 100 × H_S2_RMS_EMA.max(1.0); slot 49 fires at 1e3 × the same base, so it +// stays quiet under normal Mech 10 operation and only triggers if the clamp +// is disabled or `H_S2_RMS_EMA` itself drifted. Mirrors the Mech 5 family +// pattern (clamp at 100×, diagnostic at 1000×). The threshold base is +// `h_s2_rms_ema_eff = max(ISV[H_S2_RMS_EMA_INDEX=96], 1.0)` — distinct from +// `q_abs_ref_eff` because h_s2 lives in pre-readout activation space, not Q +// space. Don't conflate the two ISV bases. +// // Slot-index → threshold mapping (relative slot ≥ 12; absolute = slot + 24): // 12-15 (abs 36-39): Adam m max-abs ≥ 100 × q_abs_ref_eff // 16-19 (abs 40-43): Adam v max-abs ≥ 1e6 × q_abs_ref_eff² @@ -1760,11 +1771,15 @@ extern "C" __global__ void popart_normalize_robust( // q_abs_ref_eff) // 24 (abs 48) : IQN online_params max-abs ≥ 1e3 × q_abs_ref_eff // (SP3 close-out v2 — slot-26 GEMM regression sentinel) +// 25 (abs 49) : h_s2 post-clamp max-abs ≥ 1e3 × h_s2_rms_ema_eff +// (SP3 Mech 10 — activation-clamp regression sentinel) // -// q_abs_ref_eff = max(ISV[Q_ABS_REF=16], 1.0) computed host-side at launch, -// passed by-value as a float kernel arg. ε on multiplier per SP1 pearl: cold- -// start ISV (~0) gives q_abs_ref_eff = 1, so all thresholds floor at their -// ε-floor multiplier × 1. No per-slot threshold buffer, no new ISV slots. +// q_abs_ref_eff = max(ISV[Q_ABS_REF=16], 1.0) (slots 12-24) +// h_s2_rms_ema_eff = max(ISV[H_S2_RMS_EMA_INDEX=96], 1.0)(slot 25 only) +// Both computed host-side at launch, passed by-value as float kernel args. +// ε on multiplier per SP1 pearl: cold-start ISV (~0) gives *_eff = 1, so +// all thresholds floor at their ε-floor multiplier × 1. No per-slot +// threshold buffer, no new ISV slots. // // Invariants preserved from dqn_nan_check_f32: // - sticky-flag semantics (writes 1 only; never clears) @@ -1772,9 +1787,10 @@ extern "C" __global__ void popart_normalize_robust( // - graph-capture safe (pure kernel launch, no host ops) // ============================================================================ extern "C" __global__ void dqn_nan_check_fused_f32_kernel( - const float* const* buf_ptrs, // [N_SLOTS] (25 entries: 12 NaN-only + 13 threshold) + const float* const* buf_ptrs, // [N_SLOTS] (26 entries: 12 NaN-only + 14 threshold) const int* buf_lens, // [N_SLOTS] - float q_abs_ref_eff, // SP3 Mech 5: ISV-derived multiplier (slots ≥ 12) + float q_abs_ref_eff, // SP3 Mech 5: ISV-derived multiplier (slots 12-24) + float h_s2_rms_ema_eff, // SP3 Mech 10: ISV-derived multiplier (slot 25 only) int base_flag_idx, // first slot index (24 for backward checks) int* nan_flags_buf ) { @@ -1785,9 +1801,11 @@ extern "C" __global__ void dqn_nan_check_fused_f32_kernel( // SP3 Mech 5: per-slot threshold by relative slot index. Slots 0-11 // (= absolute 24-35) are NaN-only checks (threshold = +inf so the - // magnitude check is trivially false). Slots 12-24 (= absolute 36-48) - // check magnitude exceedance per the audit table above. All thresholds - // derive from q_abs_ref_eff = max(isv[Q_ABS_REF=16], 1.0). + // magnitude check is trivially false). Slots 12-25 (= absolute 36-49) + // check magnitude exceedance per the audit table above. Slots 12-24 + // derive thresholds from q_abs_ref_eff = max(isv[Q_ABS_REF=16], 1.0); + // slot 25 uses h_s2_rms_ema_eff = max(isv[H_S2_RMS_EMA_INDEX=96], 1.0) + // because h_s2 lives in activation space, not Q space. float thr = INFINITY; if (slot >= 12 && slot < 16) { // slots 36-39: Adam m @@ -1808,6 +1826,12 @@ extern "C" __global__ void dqn_nan_check_fused_f32_kernel( // slot 48 (SP3 close-out v2): IQN online_params — same weight-slice // regime as slots 44-45. Captures the slot-26 GEMM blind spot. thr = 1.0e3f * q_abs_ref_eff; + } else if (slot == 25) { + // slot 49 (SP3 Mech 10): post-trunk h_s2 max-abs sentinel. Mech 10 + // clamps h_s2 at 100 × h_s2_rms_ema_eff; this fires at 1e3 ×, so + // it never trips under normal operation. If it does, the clamp was + // disabled or H_S2_RMS_EMA itself drifted. + thr = 1.0e3f * h_s2_rms_ema_eff; } int flag_set = 0; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 2008609e5..261b19637 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -11479,31 +11479,36 @@ impl GpuDqnTrainer { // SP2 framework checker hooks + SP3 structural observers. SP3 close- // out v2 (slot 48): one additional slot for IQN online_params (the // separate weight buffer that drives the slot-26 cuBLAS GEMM and was - // the F1 step-3540 NaN root). Buffer size reviewable by SP2 framework - // codification. - let nan_flags_buf = stream.alloc_zeros::(49) + // the F1 step-3540 NaN root). SP3 Mech 10 (slot 49): one additional + // slot for post-trunk h_s2 activation max-abs sentinel — fires at + // 1e3 × H_S2_RMS_EMA.max(1.0) and stays quiet under normal Mech 10 + // clamp operation (clamp at 100×, diagnostic at 1000×). Buffer size + // reviewable by SP2 framework codification. + let nan_flags_buf = stream.alloc_zeros::(50) .map_err(|e| MLError::ModelError(format!("nan_flags alloc: {e}")))?; - // SP2 + SP3 Mech 5: fused NaN/threshold-check (ptr, len) tables. - // 25 entries each — slots 0-11 (= absolute flag idx 24-35) NaN-only + // SP2 + SP3 Mech 5 + Mech 10: fused NaN/threshold-check (ptr, len) tables. + // 26 entries each — slots 0-11 (= absolute flag idx 24-35) NaN-only // (SP2 Task A3) + slots 12-23 (= absolute 36-47) ISV-threshold // (SP3 Mech 5 Task B6) + slot 24 (= absolute 48) IQN online_params - // (SP3 close-out v2). Allocated as mapped-pinned + // (SP3 close-out v2) + slot 25 (= absolute 49) h_s2 activation + // sentinel (SP3 Mech 10). Allocated as mapped-pinned // (`cuMemHostAlloc(DEVICEMAP|PORTABLE)`) per // `feedback_no_htod_htoh_only_mapped_pinned` — host-side writes are // visible to the kernel through the device-mapped pointer with zero // HtoD copies. Populated once via `populate_nan_check_meta` after all // backward-path buffers are constructed AND `gpu_iqn` is known. The // kernel reads `buf_ptrs[blockIdx.x]` and `buf_lens[blockIdx.x]` per - // block; per-slot threshold is computed inline from a single - // `q_abs_ref_eff` launch arg (no separate thresholds buffer). + // block; per-slot threshold is computed inline from two ISV-derived + // launch args (`q_abs_ref_eff` for slots 12-24, `h_s2_rms_ema_eff` + // for slot 25 — distinct ISV bases, no per-slot threshold buffer). // // Safety: a CUDA context is active on this thread (the GpuDqnTrainer // constructor runs within `FusedTrainingCtx::new` which has already // initialised CUDA via the shared cublas handle). - let nan_check_buf_ptrs = unsafe { MappedU64Buffer::new(25) } + let nan_check_buf_ptrs = unsafe { MappedU64Buffer::new(26) } .map_err(|e| MLError::ModelError(format!("nan_check_buf_ptrs alloc: {e}")))?; - let nan_check_buf_lens = unsafe { MappedI32Buffer::new(25) } + let nan_check_buf_lens = unsafe { MappedI32Buffer::new(26) } .map_err(|e| MLError::ModelError(format!("nan_check_buf_lens alloc: {e}")))?; // v8: PopArt running statistics buffers @@ -15378,6 +15383,14 @@ impl GpuDqnTrainer { /// captures the slot-26 GEMM blind /// spot that hid the F1 step-3540 NaN) /// + /// SP3 Mech 10 — slot 49 (post-trunk h_s2 activation sentinel): + /// 25 (slot 49) save_h_s2 — max-abs ≥ 1e3 × ISV[H_S2_RMS_EMA_INDEX].max(1) + /// (regression sentinel for Mech 10's + /// activation clamp; mirrors Mech 5 family + /// clamp-at-100×/diagnostic-at-1000× pattern; + /// distinct ISV base from slots 36-48 because + /// h_s2 lives in activation space, not Q space) + /// /// Mapped-pinned host write — no HtoD copy. The kernel reads the same /// memory through the device-mapped pointer (`dev_ptr`) once the launch /// stream barrier ensures coherence. @@ -15436,7 +15449,14 @@ impl GpuDqnTrainer { let atom_positions_ptr = self.atom_positions_ptr(); let atom_positions_len = self.atom_positions_len() as i32; - let entries: [(u64, i32); 25] = [ + // SP3 Mech 10 — slot 49 accessor for post-trunk h_s2 activation + // sentinel. Same buffer as slot 12's NaN check (`save_h_s2`); the + // sentinel only fires at 1e3 × H_S2_RMS_EMA.max(1) so it stays quiet + // under normal Mech 10 clamp operation (clamp at 100×). + let save_h_s2_ptr = self.save_h_s2.raw_ptr(); + let save_h_s2_len = self.save_h_s2.len() as i32; + + let entries: [(u64, i32); 26] = [ // SP2 Task A3 — slot 24 — post-c51_grad value gradient (self.d_value_logits_buf_ptr(), self.d_value_logits_buf.len() as i32), // slot 25 — post-c51_grad branch advantage @@ -15506,21 +15526,26 @@ impl GpuDqnTrainer { iqn_online_params_ptr.unwrap_or(0), iqn_online_params_len.unwrap_or(0) as i32, ), + // SP3 Mech 10 — slot 49 — post-trunk h_s2 activation sentinel. + // Same buffer as slot 12 (NaN-only check). Threshold is computed + // inline in the fused kernel from `h_s2_rms_ema_eff` (distinct + // from `q_abs_ref_eff` — h_s2 lives in activation space). + (save_h_s2_ptr, save_h_s2_len), ]; // Host-side write through the mapped-pinned pages — kernel sees the // values via dev_ptr after the next stream-sync barrier (no HtoD copy // required, per `feedback_no_htod_htoh_only_mapped_pinned`). - let ptrs_view: [u64; 25] = std::array::from_fn(|i| entries[i].0); - let lens_view: [i32; 25] = std::array::from_fn(|i| entries[i].1); + let ptrs_view: [u64; 26] = std::array::from_fn(|i| entries[i].0); + let lens_view: [i32; 26] = std::array::from_fn(|i| entries[i].1); self.nan_check_buf_ptrs.write_from_slice(&ptrs_view); self.nan_check_buf_lens.write_from_slice(&lens_view); Ok(()) } - /// SP2 + SP3 Mech 5 + SP3 close-out v2: launch the fused + /// SP2 + SP3 Mech 5 + SP3 close-out v2 + SP3 Mech 10: launch the fused /// NaN/threshold-check kernel. - /// Single launch processes all 25 backward-path slots (24-48) in 25 + /// Single launch processes all 26 backward-path slots (24-49) in 26 /// blocks, in parallel. Slots with null pointer (deferred slot 31, /// optional IQN slots 27/28/39/43/48 when inactive, slots 33-35 inline /// elsewhere) no-op via the kernel's null-pointer guard. Sticky-flag @@ -15530,8 +15555,11 @@ impl GpuDqnTrainer { /// trivially false). Slots 12-23 (= absolute 36-47) are SP3 Mech 5 /// ISV-threshold checks. Slot 24 (= absolute 48) is SP3 close-out v2 — /// IQN online_params, threshold = 1e3 × q_abs_ref_eff (same regime as - /// slots 44-45). Threshold computed inline per-slot from a single - /// `q_abs_ref_eff` arg (no per-slot threshold buffer, no per-step HtoD). + /// slots 44-45). Slot 25 (= absolute 49) is SP3 Mech 10 — post-trunk + /// h_s2 activation sentinel, threshold = 1e3 × h_s2_rms_ema_eff (distinct + /// ISV base — h_s2 lives in activation space, not Q space). Thresholds + /// computed inline per-slot from two ISV-derived args (no per-slot + /// threshold buffer, no per-step HtoD). /// /// Replaces the 8-launch sequence in `run_nan_checks_post_backward`. pub(crate) fn launch_nan_check_fused_f32(&mut self) -> Result<(), MLError> { @@ -15546,7 +15574,14 @@ impl GpuDqnTrainer { // bug pearl: passing f64 to a `float` param reads only low 4 bytes). let q_abs_ref = self.read_isv_signal_at(Q_ABS_REF_INDEX); let q_abs_ref_eff: f32 = q_abs_ref.max(1.0_f32); - const BLOCKS: u32 = 25; + // SP3 Mech 10: read ISV[H_S2_RMS_EMA_INDEX=96] for slot 49 threshold + // (= 1e3 × h_s2_rms_ema_eff). Distinct ISV base from `q_abs_ref_eff` + // because h_s2 lives in activation space, not Q space — don't conflate + // the two. Same ε floor convention (`max(.., 1.0)`) so cold-start RMS + // (~0) gives h_s2_rms_ema_eff = 1 (threshold floors at 1e3). + let h_s2_rms_ema = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX); + let h_s2_rms_ema_eff: f32 = h_s2_rms_ema.max(1.0_f32); + const BLOCKS: u32 = 26; const THREADS: u32 = 256; const BASE_FLAG_IDX: i32 = 24; let cfg = LaunchConfig { @@ -15560,6 +15595,7 @@ impl GpuDqnTrainer { .arg(&buf_ptrs_dev) .arg(&buf_lens_dev) .arg(&q_abs_ref_eff) + .arg(&h_s2_rms_ema_eff) .arg(&BASE_FLAG_IDX) .arg(&nan_flags_ptr) .launch(cfg) @@ -15619,13 +15655,14 @@ impl GpuDqnTrainer { Ok(()) } - /// Read NaN flags back to CPU (synchronizes stream). Returns [48] flags. + /// Read NaN flags back to CPU (synchronizes stream). Returns [50] flags. /// Index → buffer mapping is documented in `run_nan_checks_post_forward`. /// SP1 Phase B: slots 24-35 reserved for backward-path checks (wired by /// Task 4); slots 36-47 SP3 Mech 5 ISV-threshold checks; slot 48 SP3 - /// close-out v2 IQN online_params weight diagnostic. - pub fn read_nan_flags(&self) -> Result<[i32; 49], MLError> { - let mut host = [0_i32; 49]; + /// close-out v2 IQN online_params weight diagnostic; slot 49 SP3 Mech 10 + /// post-trunk h_s2 activation max-abs sentinel (1e3 × H_S2_RMS_EMA.max(1)). + pub fn read_nan_flags(&self) -> Result<[i32; 50], MLError> { + let mut host = [0_i32; 50]; self.stream.memcpy_dtoh(&self.nan_flags_buf, &mut host) .map_err(|e| MLError::ModelError(format!("nan_flags read: {e}")))?; Ok(host) @@ -17140,6 +17177,55 @@ impl GpuDqnTrainer { self.risk_budget_forward(batch_size)?; } + // ── 1d. SP3 Mech 10: ISV-driven post-trunk activation clamp ── + // + // The F1 step-1000 NaN cascade (smoke-test-2xrxk on commit 0d7e27d61) + // root-caused to forward-pass GEMM accumulator overflow under + // Mech-9-clamped weights. With trunk weights pinned at 100 × + // Q_ABS_REF=5000 and K=256 inner-dim, accumulators compound by + // ~80000× per layer; with trunk depth ≥6 (encoder + VSN + GRN + // blocks + bottleneck) the chain hits f32_max ≈ 3.4e38. The slot 12 + // (`save_h_s2`) firing was the smoking gun — h_s2 itself goes + // non-finite mid-forward. Mech 1+2 clamp targets+atoms; Mech 9 clamps + // weights; Mech 10 closes the gap by clamping h_s2 immediately after + // the trunk + temporal pipeline finalises it, BEFORE any downstream + // consumer (loss path, IQN, value head, replay save, eval-action + // select) reads the buffer. + // + // Bound: `100 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)`. ISV slot 96 + // already tracks the per-batch RMS EMA of h_s2 (per Plan 4 Task + // 2c.3c.5 — same ISV slot consumed by `mag_concat_qdir` adaptive + // scale). Reusing the existing slot keeps Mech 10 consistent with + // `feedback_isv_for_adaptive_bounds.md` (no new ISV slots, no + // hardcoded multipliers). The `max(1.0)` ε floor on the multiplier + // honours the SP1 cold-start convention — at fold boundary the EMA + // resets to 1.0 (neutral RMS) per `state_reset_registry.rs`, so the + // clamp floor is always at least 100 × 1 = 100, never zero. + // + // Diagnostic ordering (SP1 pearl): the inline NaN check on slot 12 + // fires BEFORE the clamp sanitises, so the sentinel sees the original + // Inf/NaN before sanitization replaces it with 0. This is redundant + // with the slot 12 check inside `run_nan_checks_post_forward` (which + // runs later, post-clamp, on the sanitized buffer) — but the inline + // check at this site is the only one that captures the pre-clamp + // state. + // + // Sanitiser: `launch_clamp_finite_f32` (dqn_utility_kernels.cu:462) + // does `isfinite(v) ? clamp(v, ±max_abs) : 0.0f`. NaN→0, Inf→0, + // finite→clamped. NO new kernel — same helper used by Mech 9 sites. + // The clamp adds a single kernel launch in the captured graph + // (expected — captured graph topology grows by one node). + { + self.check_nan_f32(self.save_h_s2.raw_ptr(), self.save_h_s2.len(), 12)?; + let h_s2_rms_ema_eff = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX).max(1.0_f32); + let max_abs_h_s2 = 100.0_f32 * h_s2_rms_ema_eff; + self.launch_clamp_finite_f32( + self.save_h_s2.raw_ptr(), + self.save_h_s2.len(), + max_abs_h_s2, + )?; + } + // ── 2+3. Loss + gradient (blended MSE + C51 via c51_alpha ramp) ─ self.launch_curiosity_inference()?; diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 798f0bc05..69fe7af39 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -3750,9 +3750,10 @@ impl FusedTrainingCtx { /// Read NaN detection flags (synchronizes stream). Used by training guard on NaN halt. /// Index → buffer mapping is documented in /// `GpuDqnTrainer::run_nan_checks_post_forward`. - /// SP1 Phase B: returns 48 flags (was 24). Slots 24-35 cover backward-path - /// checks; slots 36-47 reserved for SP2/SP3 headroom. - pub(crate) fn read_nan_flags(&self) -> Result<[i32; 49], crate::MLError> { + /// SP1 Phase B: returns 50 flags. Slots 24-35 cover backward-path + /// checks; slots 36-47 SP3 Mech 5 ISV-threshold; slot 48 SP3 close-out v2 + /// IQN online_params; slot 49 SP3 Mech 10 post-trunk h_s2 sentinel. + pub(crate) fn read_nan_flags(&self) -> Result<[i32; 50], crate::MLError> { self.trainer.read_nan_flags() } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 93765ad59..a6ecfd7bc 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -2055,6 +2055,11 @@ impl DQNTrainer { // weight diagnostic — the slot-26 GEMM blind // spot that hid the F1 step-3540 NaN. "iqn_weight_max", // 48 + // SP3 Mech 10 (slot 49): post-trunk h_s2 max-abs + // sentinel — fires at 1e3 × H_S2_RMS_EMA.max(1) + // (clamp at 100×, diagnostic at 1000×). Stays + // quiet under normal Mech 10 operation. + "h_s2_max_abs", // 49 ]; let flagged: Vec = flags.iter().enumerate() .filter(|(_, &f)| f != 0) @@ -2062,7 +2067,7 @@ impl DQNTrainer { .collect(); tracing::error!( "NaN SOURCE at step {}: flagged=[{}] (empty = NaN entered via \ - loss-component buffer not in checks 0..47; expand coverage)", + loss-component buffer not in checks 0..49; expand coverage)", train_step_count, flagged.join(", ") ); } @@ -2156,6 +2161,11 @@ impl DQNTrainer { // weight diagnostic — the slot-26 GEMM blind // spot that hid the F1 step-3540 NaN. "iqn_weight_max", // 48 + // SP3 Mech 10 (slot 49): post-trunk h_s2 max-abs + // sentinel — fires at 1e3 × H_S2_RMS_EMA.max(1) + // (clamp at 100×, diagnostic at 1000×). Stays + // quiet under normal Mech 10 operation. + "h_s2_max_abs", // 49 ]; let flagged: Vec = flags.iter().enumerate() .filter(|(_, &f)| f != 0) diff --git a/docs/dqn-backward-nan-audit.md b/docs/dqn-backward-nan-audit.md index e2ab5e2d5..8a6b71730 100644 --- a/docs/dqn-backward-nan-audit.md +++ b/docs/dqn-backward-nan-audit.md @@ -824,10 +824,11 @@ If Gate 1 fails, halt before SP3 starts and re-investigate F0 cause (alignment e | ~~Mech 7~~ | ~~Per-element gradient clip in Adam~~ | reverted in `f67ede94f` | (n/a) | | ~~Mech 8~~ | ~~slow_ema fold-boundary reset~~ | reverted in `8956c2fb7` | (n/a) | | Mech 9 | Post-Adam ISV-driven weight clamp (all 5 Adam kernels) | `8956c2fb7` (dqn) + close-out v2 (iqn/iql/attn/curiosity) | dqn_utility_kernels.cu, iqn_dual_head_kernel.cu, iql_value_kernel.cu, attention_backward_kernel.cu, curiosity_training_kernel.cu, gpu_dqn_trainer.rs, gpu_iqn_head.rs, gpu_iql_trainer.rs, gpu_attention.rs, gpu_tlob.rs, gpu_curiosity_trainer.rs, gpu_experience_collector.rs, fused_training.rs, training_loop.rs, decision_transformer.rs | +| Mech 10 | ISV-driven post-trunk h_s2 activation clamp + slot 49 sentinel | (this commit) | dqn_utility_kernels.cu, gpu_dqn_trainer.rs, fused_training.rs, training_loop.rs | **Supporting:** Task B1 (`aee11f321`) — 24 diagnostic accessors (Adam m/v + weight + target_q + atom_positions). -### Mech 5 slot allocation (slots 36-48) +### Mech 5 slot allocation (slots 36-49) | Slot | Name | Threshold | Source | |---|---|---|---| @@ -844,12 +845,15 @@ If Gate 1 fails, halt before SP3 starts and re-investigate F0 cause (alignment e | 46 | target_q_post_clip | ≥ 95 × q_abs_ref_eff | denoise_target_q_buf | | 47 | atom_span_max | ≥ 190 × q_abs_ref_eff | per_sample_support | | 48 | iqn_weight_max | ≥ 1e3 × q_abs_ref_eff | IQN online_params (SP3 close-out v2 — slot-26 GEMM blind spot, fold-1 step-3540 sentinel) | +| 49 | h_s2_max_abs | ≥ 1e3 × h_s2_rms_ema_eff | save_h_s2 post-trunk (SP3 Mech 10 — activation-clamp regression sentinel; clamp at 100×, diagnostic at 1000×; mirrors Mech 5 family pattern) | + +Slot 49 uses `h_s2_rms_ema_eff = max(isv[H_S2_RMS_EMA_INDEX=96], 1.0)` — distinct ISV base from slots 36-48 because h_s2 lives in pre-readout activation space, not Q space. Don't conflate the two ISV bases. `q_abs_ref_eff = max(isv[Q_ABS_REF_INDEX=16], 1.0)` — ε on multiplier per SP1 pearl. Cold-start ISV (~0) → all thresholds floor at their multiplier × 1. ### ISV slot bindings -All SP3 mechanisms use existing `Q_ABS_REF_INDEX = 16`. NO new ISV slots. +SP3 Mechs 1-9 use existing `Q_ABS_REF_INDEX = 16`. Mech 10 reuses existing `H_S2_RMS_EMA_INDEX = 96` (already populated per Plan 4 Task 2c.3c.5; consumed by `mag_concat_qdir` adaptive scale per 2c.3c.6). NO new ISV slots. ### Expected validation outcome (Gate 2 smoke) @@ -927,3 +931,85 @@ should attack the GEMM accumulator path itself (TF32 disable, FP32 enforcement, or per-layer gain rescaling). Close SP3 honestly with the documented residual rather than tuning further. +### Mech 10 — ISV-driven post-trunk h_s2 activation clamp + +**Outcome of smoke-test-2xrxk (HEAD `0d7e27d61` = SP3 close-out v2 active).** +F0 trained clean (Best Sharpe 45.47); F1 NaN at step ~1000; F2 first +epoch reached **Best Sharpe 56.70** (above the 55.87 baseline) before +NaN epoch 2. Slot fire pattern at F1 NaN: +`[6=grad_buf, 12=save_h_s2, 26=iqn_trunk_m, 32=bn_d_concat_buf, +36-38=trunk/value/branch_adam_m, 40-42=trunk/value/branch_adam_v]`. +Crucially: **slot 48 (iqn_weight_max) clean, slots 39+43 (iqn_adam_m/v) +clean, slots 44-45 (trunk/heads weights) clean**. + +**Diagnosis.** SP3 close-out v2 hypothesis (IQN partial refactor) was +wrong. The actual cause: **forward-pass activation overflow under +Mech-9-clamped weights**. Slot 12 (`save_h_s2`) firing was the smoking +gun — `h_s2` itself goes non-finite mid-forward. With Mech 9 clamping +weights at `100 × Q_ABS_REF = 5000` and trunk inner-dim K = 256, forward +GEMM accumulators compound by `√K × |w|max ≈ 80000×` per layer. With +trunk depth ≥6 (encoder + VSN + GRN-blocks + bottleneck), the chain +hits `f32_max ≈ 3.4e38` somewhere. F0/F2 don't fail because their +gradients keep weights at natural scale (~0.5); F1's regime shift +drives weights toward the clamp ceiling and the forward chain saturates. + +**The protected surfaces before Mech 10.** Mech 1+2 clamped TARGETS + +ATOMS. Mech 9 clamped WEIGHTS. Nothing clamped ACTIVATIONS. Mech 10 +closes that gap. + +**Mech 10 mechanism.** After the trunk + temporal pipeline finalises +`save_h_s2` in `submit_forward_ops_main` (end of "1c. Temporal pipeline" +block — last writers are `mamba2_step` + `apply_regime_dropout`), the +trainer launches `launch_clamp_finite_f32` on `save_h_s2` with bound +`100 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)`. The kernel does +`isfinite(v) ? clamp(v, ±max_abs) : 0.0f` — NaN→0, Inf→0, finite +clamped. NO new kernel — same helper used by Mech 9 sites. + +**ISV-driven design.** `H_S2_RMS_EMA_INDEX = 96` already tracks the +per-batch RMS EMA of h_s2 (per Plan 4 Task 2c.3c.5; consumer wired in +2c.3c.6 for `mag_concat_qdir` adaptive scale). Reusing the existing +slot keeps Mech 10 consistent with `feedback_isv_for_adaptive_bounds.md` +(no new ISV slots, no hardcoded multipliers). The `max(1.0)` ε floor +on the multiplier honours the SP1 cold-start convention — at fold +boundary the EMA resets to 1.0 (neutral RMS) per +`state_reset_registry.rs`, so the clamp floor is always at least +`100 × 1 = 100`, never zero. + +**Diagnostic ordering (SP1 pearl).** The inline NaN check on slot 12 +fires BEFORE the clamp sanitises, so the sentinel sees the original +Inf/NaN before sanitization replaces it with 0. Redundant with the +slot 12 check inside `run_nan_checks_post_forward` (which runs later, +post-clamp, on the sanitized buffer) — but the inline check at this +site is the only one that captures the pre-clamp state. + +**Slot 49 diagnostic.** Mirrors the Mech 5 family pattern (clamp at +100×, diagnostic at 1000×). Threshold = `1e3 × h_s2_rms_ema_eff`. +Sentinel that NEVER fires under normal Mech 10 operation; if it does, +the activation clamp is disabled or `H_S2_RMS_EMA` itself drifted. + +**Implementation surface.** +- `dqn_utility_kernels.cu` — `dqn_nan_check_fused_f32_kernel` gains a + second ISV-derived float arg `h_s2_rms_ema_eff`; new branch for + relative slot 25 (= absolute 49) with threshold + `1e3 × h_s2_rms_ema_eff`. Slots 12-24 keep their `q_abs_ref_eff` + thresholds — distinct ISV bases. +- `gpu_dqn_trainer.rs` — `nan_flags_buf` 49→50; metadata buffers + (`nan_check_buf_ptrs`, `nan_check_buf_lens`) 25→26 entries; + `populate_nan_check_meta` writes `(save_h_s2.raw_ptr(), + save_h_s2.len())` at slot 49; `launch_nan_check_fused_f32` reads + `H_S2_RMS_EMA_INDEX` host-side and passes `h_s2_rms_ema_eff` as the + second ISV arg; `read_nan_flags()` returns `[i32; 50]`; fused-kernel + block count 25→26; clamp inserted at end of "1c. Temporal pipeline" + block in `submit_forward_ops_main` (before `launch_curiosity_inference`). +- `fused_training.rs` — `read_nan_flags()` wrapper return type 49→50. +- `training_loop.rs` — both name tables (halt_nan + halt_grad_collapse + formatters) extended to 50 entries with `"h_s2_max_abs"` at slot 49. + +**Constraints honoured.** No new ISV slots, no new kernels, no +graph-topology surprise (one extra launch in the captured forward +graph). Per `feedback_no_partial_refactor`, Mech 10 + slot 49 land in +the same commit. + +**Validation.** Deferred to one L40S smoke. Expected: F1 trains past +step 1000; slots 36-43 + slot 49 stay quiet. + diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 873ac4867..ea3274afa 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. +SP3 Mech 10 — ISV-driven post-trunk h_s2 activation clamp + slot 49 sentinel (2026-04-30): real root cause of the F1 step-1000 NaN identified by smoke-test-2xrxk on commit `0d7e27d61` (SP3 close-out v2 active). Trajectory: F0 Best Sharpe 45.47 (clean) → F1 NaN at step ~1000 → F2 first epoch reached Best Sharpe **56.70** (above the 55.87 baseline) before NaN epoch 2. Slot fire pattern at F1 NaN: `[6=grad_buf, 12=save_h_s2, 26=iqn_trunk_m, 32=bn_d_concat_buf, 36-38=trunk/value/branch_adam_m, 40-42=trunk/value/branch_adam_v]`. Crucially: slots 48 (iqn_weight_max), 39+43 (iqn_adam_m/v), 44-45 (trunk/heads weights) all stayed clean — disproving the SP3 close-out v2 hypothesis (IQN partial refactor). The actual cause: forward-pass GEMM accumulator overflow under Mech-9-clamped weights. With Mech 9 pinning weights at `100 × Q_ABS_REF = 5000` and trunk inner-dim K = 256, forward GEMM accumulators compound by `√K × |w|max ≈ 80000×` per layer; trunk depth ≥6 (encoder + VSN + GRN-blocks + bottleneck) hits `f32_max ≈ 3.4e38` somewhere. F0/F2 train clean because their gradients keep weights at natural scale (~0.5); F1's regime shift drives weights toward the clamp ceiling and the forward chain saturates. Mech 1+2 clamped TARGETS+ATOMS, Mech 9 clamped WEIGHTS, **nothing clamped ACTIVATIONS** — Mech 10 closes that gap. Mech 10 mechanism: after the trunk + temporal pipeline finalises `save_h_s2` in `submit_forward_ops_main` (end of "1c. Temporal pipeline" block — last writers are `mamba2_step` + `apply_regime_dropout`), the trainer launches `launch_clamp_finite_f32` on `save_h_s2` with bound `100 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)`. Reuses the existing helper at `dqn_utility_kernels.cu:462` (`isfinite(v) ? clamp(v, ±max_abs) : 0.0f` — NaN→0, Inf→0, finite clamped) — no new kernel. ISV-driven via existing `H_S2_RMS_EMA_INDEX = 96` (already populated per Plan 4 Task 2c.3c.5; consumed by `mag_concat_qdir` adaptive scale per 2c.3c.6) — no new ISV slot, no hardcoded multipliers, consistent with `feedback_isv_for_adaptive_bounds.md`. ε on multiplier per SP1 pearl (`isv.max(1.0)`) — at fold boundary the EMA resets to 1.0 (neutral RMS) per `state_reset_registry.rs`, so the clamp floor is always at least `100 × 1 = 100`, never zero. SP1 diagnostic-ordering pearl honoured: the inline `check_nan_f32(slot 12)` fires BEFORE the clamp sanitises so the sentinel sees the original Inf/NaN before sanitization replaces it with 0; redundant with the slot 12 check inside `run_nan_checks_post_forward` (which runs later, post-clamp, on the sanitized buffer) — but the inline check at this site is the only one that captures the pre-clamp state. New diagnostic slot 49 = `h_s2_max_abs` at threshold `1e3 × h_s2_rms_ema_eff` mirrors the Mech 5 family pattern (clamp at 100×, diagnostic at 1000×) — sentinel that NEVER fires under normal Mech 10 operation. The fused NaN-check kernel (`dqn_nan_check_fused_f32_kernel`) gains a second ISV-derived float arg `h_s2_rms_ema_eff` distinct from `q_abs_ref_eff` (don't conflate the two ISV bases — `h_s2` lives in pre-readout activation space, not Q space). Buffer bumps: `nan_flags_buf` 49→50, metadata buffers (`nan_check_buf_ptrs`, `nan_check_buf_lens`) 25→26 entries, fused-kernel block count 25→26. Both name tables in `training_loop.rs` (`halt_nan` formatter + `halt_grad_collapse` formatter) extended to 50 entries with `"h_s2_max_abs"` at slot 49. `read_nan_flags()` return type 49→50 in both `GpuDqnTrainer` and `FusedTrainingCtx`. Per `feedback_no_partial_refactor.md`, Mech 10 + slot 49 land in the same commit (single coordinated migration). No new ISV slots, no new kernels, no graph-topology surprise (one extra launch added to the captured forward graph at the temporal-pipeline → loss-path boundary). Touched: `cuda_pipeline/dqn_utility_kernels.cu` (kernel signature + slot 25 branch), `cuda_pipeline/gpu_dqn_trainer.rs` (alloc 49→50, metadata 25→26, `populate_nan_check_meta` + slot 49 entry, fused-kernel `BLOCKS` 25→26 + `h_s2_rms_ema_eff` arg, `read_nan_flags` 49→50, Mech 10 clamp call at end of "1c. Temporal pipeline" in `submit_forward_ops_main`), `trainers/dqn/fused_training.rs` (`read_nan_flags` wrapper 49→50), `trainers/dqn/trainer/training_loop.rs` (both name tables 49→50 entries with `"h_s2_max_abs"` at slot 49), `docs/dqn-backward-nan-audit.md` (Mech 10 row in mechanism table + slot 49 row in slot-allocation table + ISV bindings note + full Mech 10 section). cargo check -p ml --lib clean (12 pre-existing warnings, none new). Validation deferred to one L40S smoke at this HEAD; expected F1 trains past step 1000 and slots 36-43 + slot 49 stay quiet. + SP3 close-out v2 — Mech 9 to all 5 Adam kernels + IQN weight diag slot 48 (2026-04-30): real fix for the F1 step-3540 NaN. Commit `8956c2fb7` wired Mech 9 (post-Adam `|p_val|` clamp) into ONE Adam kernel (`dqn_adam_update_kernel`) out of FIVE in the codebase. The slot-26 GEMM (`apply_iqn_trunk_gradient` output) computes `grad_iqn @ W_iqn^T`; `W_iqn` lives in IQN's separate `online_params` buffer, updated by `iqn_adam_kernel` — never reached by Mech 9. Slots 44-45 only cover trunk + heads weights, leaving IQN's separate param buffer without a diagnostic. The chain: IQN weights drift via `iqn_adam_kernel` → finite gradient × non-finite weight → slot-26 NaN. Same partial-refactor class as Mech 4's missing GpuAttention reset (caught by B5 audit), corrected the same way per `feedback_no_partial_refactor.md`: every consumer of the contract migrates together. Extended Mech 9 to all four remaining Adam kernels — `iqn_adam_kernel` (`iqn_dual_head_kernel.cu`), `iql_adam_kernel` (`iql_value_kernel.cu`), `attn_adam_kernel` (`attention_backward_kernel.cu`, used by both `GpuAttention` and `GpuTlob` via the shared cubin), `curiosity_adam_step` (`curiosity_training_kernel.cu`). Same `if (weight_clamp_max_abs > 0.0f) p = fminf(fmaxf(p, -bound), bound)` pattern, last step before writeback. Same `100 × Q_ABS_REF.max(1.0)` ISV-driven bound (same slot 16, ε on multiplier per SP1 pearl). Each launch site computes the bound host-side and passes it as the trailing kernel arg, mirroring the existing dqn_adam_update_kernel pattern. IQN: `train_iqn_step_gpu` + `execute_training_pipeline` take new `weight_clamp_max_abs: f32` parameters; both fork-join arms in `fused_training.rs::submit_aux_ops` compute and pass the bound. IQL: `train_value_step` takes the bound; both `gpu_iql.train_value_step` (high-tau) and `gpu_iql_low.train_value_step` (low-tau) wired. Attention: `GpuAttention::adam_step` and `GpuTlob::adam_step` both take new `weight_clamp_max_abs: f32`; all 3 call sites in fused_training (parallel attn, sequential attn, TLOB Phase 6) wired. Curiosity: `GpuCuriosityTrainer::train_on_collector_buffers` and the inner `launch_adam_step` helper take the bound; `GpuExperienceCollector::train_curiosity_gpu` wraps and forwards; `training_loop.rs::collect_step` reads ISV from `fused_ctx` (curiosity collector doesn't own ISV) and passes through. DT launch in `decision_transformer.rs` keeps the 0.0 disable (offline-RL, outside SP3 scope). Added IQN-weight diagnostic slot 48 to plug the slot-26 GEMM blind spot: `nan_flags_buf` 48 → 49 (`stream.alloc_zeros::(49)`), fused-kernel block count 24 → 25 with new branch for absolute slot 48 = relative slot 24 at threshold `1e3 × q_abs_ref_eff` (matches slot 44-45 weight regime), metadata buffers `nan_check_buf_ptrs` / `nan_check_buf_lens` extended 24 → 25 with the new slot 48 entry populated by `populate_nan_check_meta` from `gpu_iqn.online_params_ptr()` + `online_params_len()` (new public accessors on `GpuIqnHead`; existing `cfg(test)` `online_params_slice()` accessor kept). Both name tables in `training_loop.rs` (the `halt_nan` formatter at line ~2042 and the `halt_grad_collapse` formatter at line ~2143) extended to 49 entries with `"iqn_weight_max"` at slot 48 — symmetry required because either path can format `nan_flags_buf`, and an out-of-range table would out-of-bounds index when slot 48 fires. `read_nan_flags()` return type bumped `[i32; 48] → [i32; 49]` on both `GpuDqnTrainer` and the `FusedTrainingCtx` wrapper. Audit doc `dqn-backward-nan-audit.md` Mech 5 slot table now spans 36-48 with the new row; Mech 9 row updated to span all 5 Adam kernels (file list expanded). No new ISV slots, no new mapped-pinned allocations beyond the metadata-buffer +1, no graph-topology change beyond the +1 block in the already-fused NaN check. Touched: `cuda_pipeline/iqn_dual_head_kernel.cu` (kernel signature + clamp), `cuda_pipeline/iql_value_kernel.cu` (kernel signature + clamp), `cuda_pipeline/attention_backward_kernel.cu` (kernel signature + clamp), `cuda_pipeline/curiosity_training_kernel.cu` (kernel signature + clamp), `cuda_pipeline/dqn_utility_kernels.cu` (fused NaN check: grid 24→25, slot 24 branch), `cuda_pipeline/gpu_iqn_head.rs` (new `online_params_ptr` / `online_params_len` accessors + `train_iqn_step_gpu` + `execute_training_pipeline` signature + Adam launch arg), `cuda_pipeline/gpu_iql_trainer.rs` (`train_value_step` signature + Adam launch arg), `cuda_pipeline/gpu_attention.rs` (`adam_step` signature + launch arg), `cuda_pipeline/gpu_tlob.rs` (`adam_step` signature + launch arg), `cuda_pipeline/gpu_curiosity_trainer.rs` (`launch_adam_step` + `train_on_collector_buffers` signatures + 4 launch args), `cuda_pipeline/gpu_experience_collector.rs` (`train_curiosity_gpu` signature wrapper), `cuda_pipeline/gpu_dqn_trainer.rs` (alloc 48→49, metadata 24→25, `populate_nan_check_meta` 2 new params + slot 48 entry, fused-kernel `BLOCKS` 24→25, `read_nan_flags` 48→49), `trainers/dqn/fused_training.rs` (Q_ABS_REF_INDEX import + 4 wiring sites + read_nan_flags type), `trainers/dqn/trainer/training_loop.rs` (curiosity ISV read + name table 49 entries), `docs/dqn-backward-nan-audit.md` (slot table + Mech 9 row). cargo check -p ml --lib clean (12 pre-existing warnings, none new). Validation deferred to one L40S smoke at this HEAD; expected F1 trains past step 3720 ceiling and slots 36-43 + new slot 48 stay quiet. SP3 Mech 9 — post-Adam ISV-driven weight clamp + Mech 8 revert (2026-04-30): coordinated close-out of SP3 Q-learning numerical-stability per `feedback_no_partial_refactor.md`. Mech 8 (slow_ema fold-boundary reset, commit `b8a7ac6f7`) reverted after smoke-test-rxhjh F1 NaN @ step 2300 — WORSE than the 3720-step ceiling Mech 6 alone produced. Persisting the α=0.001 `grad_norm_slow_ema` across folds (anchored at F0's smaller grad scale) was providing UNINTENTIONAL F1 protection by tightening Mech 6's `100 × slow_ema × isv` upper_bound during F1 ramp-up; resetting it loosened the bound during the very transient when F1 needed it tightest, accelerating Adam saturation. Removed `GpuDqnTrainer::reset_grad_norm_slow_ema` method and the `fused_training.rs::reset_for_fold` call site; kept the `grad_norm_slow_ema_pinned` field (still consumed by Mech 6) and Mech 6 multiplier at 100× (correct steady-state value, already restored from the 5× experiment in the Mech 8 commit). Mech 9 is the real fix for the cross-fold Adam pathology — root cause being targeted is cuBLAS sgemm f32-accumulator overflow at slots 26 (`iqn_trunk_m`) + 32 (`bn_d_concat_buf`) at F1 step ~3540 with INPUTS clean (slots 27, 35 unflagged). The matmul output saturates because WEIGHT tensors have drifted to extreme-but-finite magnitudes via Adam updates over thousands of steps. Mech 1 (gradient clamp) and Mech 6 (grad-norm clip) bound gradients, not parameters; neither prevents this drift. Mech 9 clamps `|p_val| ≤ 100 × Q_ABS_REF.max(1.0)` inside `dqn_adam_update_kernel` after the L1 proximal step (new trailing arg `weight_clamp_max_abs` on the kernel signature). ISV-driven: bound reads `Q_ABS_REF_INDEX = 16` host-side (same slot consumed by Mechs 1+2+5+6 — no new slot), ε on the multiplier per the SP1 pearl (`isv.max(1.0)`, never floor the bound itself or re-introduce the cold-start clamp pathology), 1 OOM below the slot 44-45 diagnostic threshold (`1e3 × Q_ABS_REF`) so the regression sentinel retains 10× firing headroom — symmetric with Mech 1's clamp/diagnostic ratio. All four Adam launch sites in `gpu_dqn_trainer.rs` wired with the bound (`launch_adam_update` at ~line 19470, `step_selectivity_adam` at ~line 20860, denoise-head Adam at ~line 17014, OFI-embed Adam at ~line 5758); the Decision Transformer Adam launch in `decision_transformer.rs` passes 0.0 to disable the clamp (DT is a separate offline-RL trainer outside SP3 scope, no ISV bus access; kernel branch `if (bound > 0.0f)` makes the disabled path zero-cost). No new ISV slots, no new kernels, no graph topology change. Touched: `cuda_pipeline/dqn_utility_kernels.cu` (kernel signature + clamp), `cuda_pipeline/gpu_dqn_trainer.rs` (4 launch sites + Mech 6 comment update + `reset_grad_norm_slow_ema` method removed), `cuda_pipeline/decision_transformer.rs` (DT launch site, disabled), `trainers/dqn/fused_training.rs` (Mech 8 call removed), `docs/dqn-backward-nan-audit.md` (mechanism table + Mech 8/9 subsection). cargo check clean. No fingerprint change — kernel signature change but param-tensor layout unchanged.