fix(dqn): SP3 Mech 10 — ISV-driven h_s2 activation clamp + slot 49 diag
Real root cause of the F1 step-1000 NaN, identified by smoke-test-2xrxk diagnostic. With Mech 9 clamping weights at 100×Q_ABS_REF=5000, F0/F2 trained clean (weights stay at natural scale ~0.5) but F1's regime shift drove weights toward the clamp ceiling. With weights pinned at 5000 and K=256 inner-dim, forward GEMM accumulators compound by ~80000× per layer; trunk depth ≥6 (encoder + VSN + GRN-blocks + bottleneck) hits f32_max ≈ 3.4e38 in the chain. Slot 12 (save_h_s2) firing was the smoking gun — h_s2 itself goes non-finite mid-forward. Slots 39+43+48 (IQN-side) stayed clean, disproving the close-out v2 hypothesis (IQN partial refactor) and confirming the surface that wasn't protected: ACTIVATIONS. Mech 1+2 clamp targets+atoms. Mech 9 clamps weights. Nothing clamped the trunk's intermediate hidden state. Mech 10 closes that gap by clamping h_s2 at 100 × H_S2_RMS_EMA.max(1.0) immediately after the trunk + temporal pipeline finalises it, before any downstream consumer reads it (branches, IQN, value head, replay save). Implementation: - Reuses existing launch_clamp_finite_f32 (dqn_utility_kernels.cu:462) which does isfinite(v) ? clamp(v) : 0.0f — NaN→0, Inf→0, finite clamped. No new kernel. - ISV-driven via H_S2_RMS_EMA_INDEX=96 (already consumed by mag_concat adaptive scale per P4.T2c.3c.6 — no new slot). - ε on multiplier per SP1 pearl: isv[96].max(1.0). Cold-start RMS near 0 floors the bound at 100×1=100, not at 0. - Inline NaN check (slot 12) preserved BEFORE the clamp per SP1 diagnostic-ordering pearl — sentinel sees the original Inf/NaN before sanitization replaces it with 0. - Insertion site: end of submit_forward_ops_main 1c temporal-pipeline block (after mamba2_step + apply_regime_dropout — the last writers to save_h_s2), immediately before launch_curiosity_inference and the loss path. Captured in the same forward child graph. New diagnostic slot 49 = h_s2_max_abs at threshold 1e3 × H_S2_RMS_EMA.max(1.0). Mirrors Mech 5 family pattern (clamp at 100×, diagnostic at 1000×) — sentinel that NEVER fires under normal Mech 10 operation. If it does fire, the activation clamp is disabled or H_S2_RMS_EMA itself drifted. Buffer bumps: nan_flags_buf 49→50, metadata 25→26, fused-kernel block count 25→26. Kernel signature gains second ISV-derived float param h_s2_rms_ema_eff (distinct from q_abs_ref_eff — h_s2 lives in activation space, not Q space; don't conflate the two ISV bases). Both name tables in training_loop.rs (halt_nan + halt_grad_collapse) extended to 50 entries. No new ISV slots, no new kernels, no graph-topology surprise (one extra launch in captured graph). cargo check clean. Audit docs docs/dqn-wire-up-audit.md (Mech 10 entry prepended) and 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) updated per Invariant 7. Validation: deferred to one L40S smoke. Expected: F1 trains past step 1000; slots 36-43 + slot 49 stay quiet. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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;
|
||||
|
||||
@@ -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::<i32>(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::<i32>(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()?;
|
||||
|
||||
|
||||
@@ -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()
|
||||
}
|
||||
|
||||
|
||||
@@ -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<String> = 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<String> = flags.iter().enumerate()
|
||||
.filter(|(_, &f)| f != 0)
|
||||
|
||||
@@ -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.
|
||||
|
||||
|
||||
File diff suppressed because one or more lines are too long
Reference in New Issue
Block a user