4476 Commits

This Branch

This Branch

All Branches

Author	SHA1	Message	Date
jgrusewski	8956c2fb77	fix(dqn): SP3 close-out — Mech 9 post-Adam weight clamp + Mech 8 revert ... Coordinated close-out of SP3 Q-learning numerical-stability per feedback_no_partial_refactor. REVERT Mech 8 (slow_ema fold-boundary reset). smoke-test-rxhjh on b8a7ac6f7 showed F1 NaN at step 2300 — WORSE than the 3720 ceiling with Mech 6 alone. Persistent slow_ema across folds was providing unintentional F1 protection (anchored at F0's smaller grad scale → tighter Mech 6 upper_bound during F1 ramp-up). Resetting loosened the clip and accelerated Adam saturation. Removed: reset_grad_norm_slow_ema() method + reset_for_fold call site. Kept: grad_norm_slow_ema_pinned field (consumed by Mech 6). Kept: Mech 6 multiplier at 100× (already restored from the 5× experiment in the Mech 8 commit; correct steady-state value). ADD Mech 9 (post-Adam ISV-driven weight clamp). Root cause being targeted: 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 the WEIGHT matrices 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. ISV-driven (slot 16, ε on multiplier per SP1 pearl). 1 OOM below 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. Implementation: - dqn_adam_update_kernel: new trailing arg `weight_clamp_max_abs`. Clamp `p_val = fminf(fmaxf(p_val, -bound), bound)` after L1 step. - All Adam launch sites in gpu_dqn_trainer.rs: compute bound from read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0) × 100.0 host-side, pass as final arg. - No new ISV slots. No new kernels. No graph topology change. Validation: deferred to one L40S smoke at the new HEAD. Expected: - F1 trains past step 3720 (Mech-6-only ceiling) — Mech 9 closes SP3. - F1 still NaNs at similar steps — close SP3 honestly with documented residual pointing at SP4/SP5. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 13:14:33 +02:00
jgrusewski	b8a7ac6f70	fix(dqn): SP3 Mech 8 — slow_ema fold-boundary reset (real root cause) ... User-identified root cause: grad_norm_slow_ema (α=0.001, half-life ~693 steps) persists across fold boundaries while every other distribution-tracking signal (Adam m/v, target nets, atom positions) gets reset. Mech 6's upper_bound formula (100 × slow_ema × isv) was anchored to the WRONG scale during F1 ramp-up, driving the non-monotonic multiplier-tuning dance across smokes: - smoke-test-fxvkk (mult=100×): F0=44, F1 NaN @ 3720 - smoke-test-ftdjz (mult=100×, +Mech7): F0=38, F1 collapse @ 2040 - smoke-test-d25vq (mult=5×): F0=21, F1 collapse @ 2820 The multiplier was searching the wrong dimension — the anchor itself was stale. Three coordinated changes (per feedback_no_partial_refactor): 1. RESTORE Mech 6 multiplier 5× → 100×. The original 100× was correct for steady-state; the F1 saturation was driven by anchor staleness, not multiplier looseness. Tightening it harmed F0 (over-clip during ramp-up when slow_ema lagged grad_norm_ema). 2. ADD reset_grad_norm_slow_ema method on GpuDqnTrainer. Zeroes the mapped-pinned scalar. First step of new fold builds up slow_ema fresh, with MIN_CLIP=1.0 floor active during the brief transient. 3. WIRE Mech 8 call in fused_training.rs::reset_for_fold, alongside Mech 4's existing Adam resets. Mech 6's anchor now aligns with the new fold's grad scale from step 1 — same philosophy as Mech 3 (target net hard-sync) and Mech 4 (Adam EMA reset). Net SP3 design: Mech 6 stays at 100× multiplier (broad, principled headroom), Mech 8 keeps the anchor honest. The pair is more robust than either change alone: - Without Mech 8: anchor is stale, multiplier tuning has no winning setting (5× hurts F0 ramp, 100× lets F1 saturate at slot 36). - With Mech 8: anchor is fresh per fold; 100× multiplier provides legitimate per-step headroom over CURRENT fold's grad norm. Mech 7 stays reverted (per-element clip was misdiagnosis — over-clipped legitimate gradient outliers without addressing the saturation root cause). F0 risk: low — F0 starts with slow_ema=0 anyway (cold start), so Mech 8 is a no-op on F0. Only changes F1+F2 fold-boundary behavior. F1+F2 expectation: Mech 6's upper_bound now scales with the CURRENT fold's grad norm, providing legitimate ~10× headroom per step without allowing Adam EMA saturation. Slot 36-42 should stay quiet.	2026-04-30 12:10:46 +02:00
jgrusewski	f67ede94fb	fix(dqn): SP3 Mech 6 v2 — tighten clip multiplier 100x -> 5x; revert Mech 7 ... Smoke smoke-test-ftdjz (commit d9a4d98a3, Mech 6 + Mech 7) regressed both F0 (44 -> 38.82, per-element cap clipped legitimate outliers) and F1 (grad-collapse at step 2040 vs 3720 with Mech 6 alone). Slots 36-42 STILL fired — Mech 7's per-element clip didn't prevent Adam EMA saturation, just slowed training to grad-collapse. Re-analysis: Mech 6's 100x multiplier on the upper-bound formula was mismatched with the slot 36 threshold ratio. Per-element gradient max <= adaptive_clip = 100 x slow_ema x isv ~= 200. Adam m_X steady-state reaches 200, exceeding slot 36 threshold of 100*isv = 100. Mech 6 was doing its job but the bound was wider than the diagnostic threshold. Two coordinated changes (per feedback_no_partial_refactor): 1. REVERT Mech 7 (per-element clip in dqn_adam_update_kernel). The per-element approach was misdiagnosis — clipping post-global-clip gradients tighter than legitimate per-element variance harms F0 training without addressing Adam saturation root cause. Kernel returns to its post-Mech-6 state (blob 546feee48). 2. TIGHTEN Mech 6's upper-bound multiplier from 100x to 5x. Standard DL practice (5-10x steady-state grad norm). Per-element gradient max becomes <= 5 x slow_ema x isv ~= 10, well below slot 36 threshold of 100. Adam m_X EMA stays bounded <= 10 — slots 36-42 should not fire. Why 5x and not 10x: slot 36 threshold is 100 x isv. 5x slow_ema (~=10 absolute) leaves 10x headroom against the threshold, providing robust margin. 10x slow_ema would be 20 absolute, only 5x margin — risk of fluctuations triggering slot 36. Why not tighter (e.g., 2x): per-step gradient norms can legitimately spike to 5x slow_ema in normal training (e.g., gradient resumption after warmup); tighter bounds would over-clip. F0 risk: low — F0 typical adaptive_clip values are bounded by Mech 6's upper anchor only when the EMA-driven clip exceeds 5x slow_ema, which is rare in steady F0 training. Cap should be a no-op for F0. F1 risk: prevents the saturation pathway diagnosed by smoke smoke-test-ftdjz. Validates by running the next smoke at this commit. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 11:43:03 +02:00
jgrusewski	d9a4d98a3d	feat(dqn): SP3 Mech 7 — per-element gradient clip in Adam kernel ... Smoke smoke-test-fxvkk (commit 48c25d999, Mech 6 anchored clip) F1-NaN'd at step 3720 with slots 36-38, 40-42 still firing — Adam EMAs saturated despite Mech 6 capping the global clip threshold. Diagnosis: Mech 6 bounds the AGGREGATE L2 gradient norm, but Adam m/v EMAs are PER-ELEMENT. A gradient with one large element (e.g., element_X = 100, rest small) has L2 norm ≈ 100, passing a clip threshold of 1000 untouched. Adam m_X EMA accumulates the large element; β1=0.9 steady-state gives m_X ≈ 1000, exceeding slot 36 threshold (100 × isv). Fix: in dqn_adam_update_kernel, after the global L2 clip is applied, clip EACH gradient element to ±per_element_cap where: per_element_cap = 10 × sqrt(adaptive_clip / total_params) Rationale: - sqrt(adaptive_clip / N) is the average per-element contribution to the L2 norm budget - 10× allows legitimate per-element deviations up to 10× average - Scales with adaptive_clip — when global clip is doing its job (Mech 6), per_element_cap is tight enough to prevent saturation - When global clip is loose (post-fold warmup), per_element_cap scales with it — never tighter than the global clip's intent Together with Mech 6, prevents Adam saturation at both the aggregate (L2 norm) and per-element levels. Closes the residual pathology after Mech 6's partial 3060→3720 improvement.	2026-04-30 11:27:54 +02:00
jgrusewski	48c25d9997	feat(dqn): SP3 Mech 6 — anchored upper bound on adaptive grad clip ... Adds an upper bound to update_adaptive_clip's new_clip formula to prevent clip-drift pathology: consecutive elevated samples were ratcheting the EMA-driven clip threshold upward without bound, eventually rendering clipping a no-op against in-distribution drift. Smoke smoke-test-5rqzs (commit b9edccfc1) F1-NaN'd at step 3060 with Mech 5 diagnostic slots 36-38, 40-42 firing — DQN main Adam m + v EMAs saturated. Diagnostic confirmed Adam saturation as the root cause. Diagnostic chain identified: grad_norm_ema ratchets up via repeat- sample compounding (winsorizer caps single-sample but not consecutive elevated samples) -> clip threshold > actual grad_norm -> clip no-op -> Adam m/v poisoned -> saturate at slot 36-42 thresholds -> cuBLAS overflow at slots 26+32 -> loss=34.68 grad=inf. Fix: cap new_clip by grad_norm_slow_ema * 100 * isv[Q_ABS_REF=16].max(1). - grad_norm_slow_ema (existing alpha=0.001 slow EMA): anchors against legitimate steady-state grad norm - 100x: headroom for per-step deviations - isv[Q_ABS_REF].max(1): adaptive scale per feedback_isv_for_adaptive_bounds; epsilon on multiplier per SP1 pearl - .max(MIN_CLIP): cold-start epsilon-floor Standard DL practice — bounds the runaway EMA clip-drift while preserving adaptive behavior. F0 risk: low — slow_ema * 100 is broad headroom; F0-typical clip thresholds are well below this cap. F1+F2 benefit by preventing the ratchet pathology. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 11:06:36 +02:00
jgrusewski	b9edccfc19	docs(dqn): SP3 Phase F — audit summary of B1-B7 + Gate 2 prep	2026-04-30 10:30:37 +02:00
jgrusewski	5f7db7d8d8	feat(dqn): SP3 — name table entries for slots 36-47 diagnostic ... 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.	2026-04-30 10:29:28 +02:00
jgrusewski	45e077188f	feat(dqn): SP3 Mech 5 — fused kernel extended for slots 36-47 threshold checks ... Extends dqn_nan_check_fused_f32_kernel to handle 24 slots (12 NaN-only + 12 ISV-threshold). Per-slot thresholds computed inline in the kernel from a single q_abs_ref_eff arg (no HtoD per step; no separate thresholds buffer needed). Slot index → threshold mapping: - 12-15 (slots 36-39): Adam m ≥ 100 × q_abs_ref_eff - 16-19 (slots 40-43): Adam v ≥ 1e6 × q_abs_ref_eff² - 20-21 (slots 44-45): Weight max ≥ 1e3 × q_abs_ref_eff - 22 (slot 46): target_q ≥ 95 × q_abs_ref_eff (= 9.5 × max_abs_target_q) - 23 (slot 47): atom span ≥ 190 × q_abs_ref_eff (= 9.5 × max_atom_abs × 2) nan_check_buf_ptrs/lens resized 12→24 entries (mapped-pinned host write at construction). populate_nan_check_meta extended with 4 new args for IQN Adam m/v ptr+len (Option<u64>/Option<usize> — null when IQN inactive). Grid: 12 → 24 blocks. Single launch covers all 24 backward-path diagnostic slots. Slot 31 (deferred) + slots 33-35 (inline elsewhere) + slots with null IQN entries no-op via the kernel's null-pointer guard. q_abs_ref_eff = max(isv[Q_ABS_REF=16], 1.0) — ε on multiplier per SP1 pearl; cold-start ISV (~0) gives q_abs_ref_eff = 1, so all thresholds floor at their ε-floor multiplier × 1. SP3 Mech 5 closes the diagnostic instrumentation loop — slots 36-47 fire when their threshold is exceeded, providing observability for SP3's other 4 mechanisms' effectiveness. Fail-safe: if SP3 fix doesn't fully resolve F1 NaN, slot 36-47 firing pattern guides the next iteration. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 10:24:41 +02:00
jgrusewski	ef429c25d8	feat(dqn): SP3 Mech 4 — comprehensive Adam EMA reset (audit + GpuAttention wire) ... Audit of all Adam optimizer state in crates/ml/src/ against the existing reset_adam_state calls at the fold-boundary in fused_training.rs::reset_for_fold: - DQN main (line 971) — covers m_buf/v_buf, IQN-trunk, q_attn, sel, denoise, mamba2, ofi_embed, PopArt, Q-div EMA via GpuDqnTrainer::reset_adam_state. Also covers VSN params (in main params buffer slots) and aux next_bar/regime heads (slots 119-126). - IQN head (line 1011) - TLOB (line 1026) - IQL-high (line 1031) - IQL-low (line 1036) Missing wire identified and added: - GpuAttention (4-head feature attention on h_s2): owns a separate (attn_m, attn_v, attn_adam_step) tuple in gpu_attention.rs, called every training step via attn.adam_step at fused_training.rs:1962/1995, but was 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. Implementation: - Add GpuAttention::reset_adam_state mirroring the gpu_iqn_head and gpu_tlob pattern (memset_zeros m/v + zero step counter + write 0 through pinned t_pinned for next adam_step launch). - Wire it under `if let Some(ref mut attn) = self.gpu_attention` in reset_for_fold, immediately after the IQL-low reset, with the same warn-on-failure / info-on-success log pattern. - Append SP3 Task B5 entry to docs/dqn-wire-up-audit.md documenting the audit + fix (Invariant 7 contract). Out-of-scope optimizers (documented in audit, not changed): - DecisionTransformer: scoped within DT pretrain block, dropped at end of pretrain (no fold leak possible). - 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. Comprehensive coverage now per feedback_no_partial_refactor — every in-context Adam state buffer zeros at fold transition. SP3 Mech 4 closes the "Adam EMA accumulation drives weight pathology" pathway.	2026-04-30 10:14:40 +02:00
jgrusewski	96b77043e0	feat(dqn): SP3 Mech 2 — C51 atom-position growth bounds ... Clamps C51 atom positions during all dynamic write paths to ±10 × ISV[Q_ABS_REF=16].max(1.0) via inline fminf(fmaxf(...)): - atoms_update_kernel.cu — active GPU-driven shared atom grid (all 4 branches via grid.x = branch_id) - iql_value_kernel.cu::iql_compute_per_sample_support — per-sample [v_min, v_max, delta_z] tile (delta_z recomputed from clamped span) - experience_kernels.cu::adaptive_atom_positions — legacy single-branch entry kept in lockstep per feedback_no_partial_refactor Same ISV bound as Mech 1 (target_q clip) — atoms and target_q share the magnitude scale. ε on the multiplier (`isv.max(1.0)`) per the SP1 ε-floor pearl; bound ≥ 10 even with cold-start ISV. Closes the second of two Q-target inflation pathways: with Mech 1 capping the projection target and Mech 2 capping the atom support, the C51 expected_q (mean of atom × prob) is bounded by ±10 × ISV[16].max(1.0) regardless of probability mass distribution. Prevents the Adam EMA saturation → weight pathology → cuBLAS overflow chain. Reuses isv_signals already in scope at all three kernels — no new kernel arg, no new launch site, no new ISV slot, no new buffer, no new kernel. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 10:03:33 +02:00
jgrusewski	27ed7daa6a	feat(dqn): SP3 Mech 1 — target_q clipping at single-point source ... Clips denoise_target_q_buf to ±10 × ISV[Q_ABS_REF=16].max(1.0) after compute_denoise_target_q. Reuses dqn_clamp_finite_f32_kernel from SP1 (no new kernel). Single-point clip covers all downstream consumers (C51 / IQN / MSE). ε on multiplier per SP1 pearl; bound at least 10 even with cold-start ISV. F0 paper-review: F0-typical |target_q| ≤ 10; with Q_ABS_REF EMA ≈ 1-5 at F0, max_abs = 10-50 — F0 guard no-op. F1 inflation (thousands+) gets clamped, breaking the Q-target inflation pathway that drives Adam EMA saturation → weight pathology → cuBLAS overflow. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 09:57:41 +02:00
jgrusewski	aee11f3210	feat(dqn): SP3 — slot 36-47 diagnostic accessors (Task B1) ... Adds pub(crate) accessor methods on GpuDqnTrainer and pub accessors on GpuIqnHead exposing device pointers for the SP3 Mech 5 threshold checks at slots 36-47: GpuDqnTrainer (pub(crate)): 36: trunk_adam_m_ptr/_len — m_buf trunk slice (tensors [0..13)) 37: value_adam_m_ptr/_len — m_buf value-head slice (tensors [13..17)) 38: branch_adam_m_ptr/_len — m_buf branch-head slice (tensors [17..33)) 40: trunk_adam_v_ptr/_len — v_buf trunk slice 41: value_adam_v_ptr/_len — v_buf value-head slice 42: branch_adam_v_ptr/_len — v_buf branch-head slice 44: trunk_params_ptr/_len — params_buf trunk slice 45: heads_params_ptr/_len — params_buf value+branch concat 46: target_q_ptr/_len — denoise_target_q_buf [B, 12] 47: atom_positions_ptr/_len — atom_positions_buf [4, num_atoms] GpuIqnHead (pub): 39: adam_m_ptr/_len — IQN Adam first-moment buffer 43: adam_v_ptr/_len — IQN Adam second-moment buffer DQN Adam state is UNIFIED in m_buf/v_buf (single TOTAL_PARAMS-sized buffer covering trunk + heads at the same offsets as params_buf). Slot 36/37/38 (and 40/41/42) accessors return pointers into the same buffer at offsets computed via padded_byte_offset over the existing compute_param_sizes layout. The kernel discriminates by slot index for threshold checks. IQN Adam state lives separately on GpuIqnHead. Trunk/heads param-buf split (slots 44/45) uses the same offset helper — no new buffers, no new ISV slots, no DtoD/HtoD copies. Each ptr accessor returns self.<field>.raw_ptr() (or +offset for slices); each len accessor returns self.<field>.len() or computed from compute_param_sizes. Audit doc updated (docs/dqn-wire-up-audit.md SP3 Task B1 entry). Used by populate_nan_check_meta_v2 (Task B6) to feed the fused kernel's threshold-check entries when extended to 24 slots. Unused yet — wired in B6. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 09:53:09 +02:00
jgrusewski	05048d03eb	docs(dqn): SP2 Gate 1 partial pass — proceed to Phase B per Option B	2026-04-30 09:47:29 +02:00
jgrusewski	e270b5cc42	docs(dqn): SP2 Phase E — audit doc summary of A1-A4 + Gate 1 prep	2026-04-30 09:25:45 +02:00
jgrusewski	b5a064f6ca	feat(dqn): SP2 — replace 8 check_nan_f32 calls with fused launch ... run_nan_checks_post_backward now fires a single fused kernel launch covering slots 24-35 in 12 blocks. Reduces graph-capture overhead from 8 launches to 1 per step. Slots 33-35 inline checks at backward orchestration sites (gpu_dqn_trainer.rs:18580/:18601/:6939) unchanged — multi-point localization preserved. Method signature simplified — IQN pointers no longer per-step args. populate_nan_check_meta (called once at construction in fused_training.rs) baked Option<u64> nullity into the metadata buffer entries; the fused kernel's null-pointer guard handles slots 27/28 when IQN inactive without per-step branching at the Rust caller. Both call sites in fused_training.rs (ungraphed + graph-captured paths) updated together per feedback_no_partial_refactor. This is the F0 regression fix — Phase B's per-step kernel-launch overhead was the F0 cause across 3 SP1 smokes (F0 ~ 35 vs baseline 55.87). Gate 1 smoke (Task A6) validates F0 >= 53.08.	2026-04-30 09:22:52 +02:00
jgrusewski	fbf48df9de	feat(dqn): SP2 A3 — fused NaN check populate + launch wrapper ... Replaces A2's CudaSlice<u64>/<i32> field types with MappedU64Buffer/ MappedI32Buffer per feedback_no_htod_htoh_only_mapped_pinned. Mapped- pinned eliminates the HtoD copy entirely — the kernel reads via the device-mapped pointer (cuMemHostGetDevicePointer_v2) while the trainer writes through the same mapped pages on the host side. Adds populate_nan_check_meta on GpuDqnTrainer (one-shot construction- time write of 12 (ptr, len) tuples for slots 24-35). Slot 31 deferred (null entry); slots 27/28 nullable on Option<u64> (None when IQN inactive); slots 33-35 null (inline checks fire separately at backward orchestration phases — kept individual for entry-point localization). Adds launch_nan_check_fused_f32 (per-step kernel launch wrapper with grid_dim=12, block_dim=256, base_flag_idx=24). Registers dqn_nan_check_fused_f32_kernel in compile_training_kernels (tuple 43→44, info log 38→39 utility kernels) — same module as the per-buffer dqn_nan_check_f32 to share the captured replay group. Constructor-time wire-up lands in FusedTrainingCtx::new after gpu_iqn construction (gpu_iqn is owned by FusedTrainingCtx, not GpuDqnTrainer — mirrors the same Option<u64> arg pattern used by apply_iqn_trunk_gradient and run_nan_checks_post_backward). Wrapper unused yet — call-site replacement (8 individual check_nan_f32 calls in run_nan_checks_post_backward → single fused launch) lands in A4. Audit doc updated (Invariant 7). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 09:14:03 +02:00
jgrusewski	82b6bd369e	feat(dqn): SP2 — nan_check_buf_ptrs + nan_check_buf_lens device buffers ... Adds 12-entry u64 ptr table + 12-entry i32 len table on GpuDqnTrainer for the fused NaN-check kernel. Allocated as zeros; populated once in next commit (slot pointers stable across trainer lifetime). Slot 31 (deferred ensemble) entry will be 0 — kernel skips that block. Two separate buffers chosen over packed-stride layout for direct ABI match with the kernel signature (const float* const* buf_ptrs, const int* buf_lens). Avoids stride-alignment concerns. Unused yet — populate + wrapper + call-site replacement in next commits.	2026-04-30 09:00:39 +02:00
jgrusewski	0250722a0d	feat(dqn): SP2 — fused multi-buffer NaN check kernel ... Adds dqn_nan_check_fused_f32_kernel: single-launch replacement for 8 per-step dqn_nan_check_f32 calls in run_nan_checks_post_backward. Block N processes slot base_flag_idx + N; sticky-flag semantics preserved; graph-capture safe; null buf_ptr = no-op (deferred slot 31). Unused yet — Rust wrapper + call-site replacement in next commits.	2026-04-30 08:54:17 +02:00
jgrusewski	9b35fe9d45	plan(dqn): SP2 + SP3 implementation — fused NaN kernel + 5-mechanism Q-stability ... Implementation plan for the approved combined design spec (commit ed727b51c). Three phases: Phase A — SP2 (5 tasks A1-A5 + Gate 1 smoke task A6): - A1: Append dqn_nan_check_fused_f32_kernel to dqn_utility_kernels.cu - A2: Add nan_check_buf_ptrs / nan_check_buf_lens device buffers (separate ptr + len arrays per the design fork resolution; cleaner than packed) - A3: populate_nan_check_meta + launch_nan_check_fused_f32 wrappers - A4: Replace 8 individual check_nan_f32 calls with single fused launch - A5: Audit doc Phase E + wire-up Invariant 7 entries - A6: Gate 1 smoke — F0 ≥ 53.08 (must pass before Phase B) Phase B — SP3 (8 tasks B1-B8 + Gate 2 smoke task B9): - B1: Slot 36-47 diagnostic accessors (Adam m/v + weight slice + target_q + atoms) - B2: Mech 1 — target_q clip after compute_denoise_target_q (single-point) - B3: Mech 2 — C51 atom-position growth bounds in EMA-update kernel - B4: Mech 3 — hard target sync (DQN main + IQN) at fold boundary - B5: Mech 4 — comprehensive Adam EMA reset at fold transitions - B6: Mech 5 — fused kernel extended to 24 slots with inline threshold-by- slot-index logic (no per-step HtoD; q_abs_ref_eff passed as single arg) - B7: Name table entries for slots 36-47 (both halt_nan + halt_grad_collapse) - B8: Audit doc Phase F + wire-up Invariant 7 entries - B9: Gate 2 smoke — full SP1 7-criterion validation Phase C — Closure (2 tasks C1-C2): - C1: project_sp2_sp3_resolved.md memory entry + SP1 closure update - C2: Audit doc closure + wire-up final entry + push Operating principles enforced throughout: - ISV-driven bounds (Q_ABS_REF=16, ε on multiplier per SP1 pearl) - No new ISV slots - No DtoD/HtoD per step (inline-compute resolution for thresholds) - Permanent diagnostic always-on (no debug flags) - Per-task commits during dev; SP3 mechanisms ship together - F0 paper-review documented per ISV bound Plan covers ~1000 LOC across multiple commits, 2 L40S smokes (Gate 1 + Gate 2), 4-7 days estimated. Self-review confirms full spec coverage.	2026-04-30 08:49:36 +02:00
jgrusewski	ed727b51c0	spec(dqn): SP2 + SP3 combined — F0 regression fix + Q-learning structural stability ... Designs the fixes for the two SP1 leftover issues identified at SP1 closure (commit 047f175fb): 1. SP2 — F0 regression fix (instrumentation overhead optimization): Consolidate the 11 per-step check_nan_f32 launches in run_nan_checks_post_backward into a single fused-reduce kernel (dqn_nan_check_fused_f32_kernel). 1 launch instead of 11; expected F0 return to ~55 baseline (was regressing to ~35 across 3 SP1 smokes). 2. SP3 — Q-learning structural stability (5 mechanisms): - M1: Target-Q clipping at single-point source, ISV-driven (10 × isv[Q_ABS_REF].max(1.0)) - M2: C51 atom-position growth bounds, same ISV bound - M3: Hard target sync at fold boundary (DQN + IQN + ensemble target nets) - M4: Adam EMA reset at fold transitions (comprehensive — all Adams) - M5: Embedded Adam-centric diagnostic instrumentation (slots 36-47): Adam m/v max-abs (slots 36-43), weight max-abs (44-45), target_q post-clip (46), atom span (47). Sticky-bit threshold checks; SP2's fused kernel is extended (NOT replaced) to handle all 24 slots in 24 blocks. Architecture: SP2 lands first; Gate 1 must pass (F0 ≥ 53.08 + slot 24-35 firing topology unchanged) before SP3 starts. SP3's 5 mechanisms ship in ONE commit per feedback_no_partial_refactor (related fixes for the same root pathology — Q-target inflation drives Adam EMA saturation drives weight drift drives cuBLAS overflow). Gate 2 evaluates the full SP1 7-criterion bar. Operating principles (mandatory): - ISV-driven design for every dynamic bound (no hardcoded clip values) - ε floor on ISV multiplier per SP1 pearl (`isv.max(1.0)`) - Combined RELATED commits per feedback_no_partial_refactor - Permanent diagnostic infrastructure (always-on, no debug flags) - No deferrals — anomalies during impl get fixed within scope Validation cost: SP2 ~1-2 days + 1 smoke; SP3 ~3-5 days + 1-2 smokes. Combined ~1000 LOC across 2 commits. Spec covers: architecture, components (SP2 fused kernel + SP3 5 mechanisms), data flow, ISV slot bindings, validation gates, error handling failure modes, operating principles, anti-patterns, sequencing constraints, effort estimate, handoff conditions.	2026-04-30 08:36:32 +02:00
jgrusewski	047f175fbf	docs(dqn): SP1 closure — surgical-fix scope delivered; SP2/SP3 handoff ... SP1 (F1 NaN root-cause investigation) closes at commit ab2133463 on plan-c-phase-2-thompson with the surgical-fix-scope deliverables. Three L40S validation smokes (smoke-test-xvzgk, dr2bn, ckgv8) characterize both what SP1 CAN address (cold-start clamp pathology) and what it CANNOT (structural Adam weight pathology — SP3 scope). What SP1 produced ----------------- 1. Permanent diagnostic infrastructure verified working: - nan_flags_buf 24→48 with backward-path coverage at slots 24-35. - run_nan_checks_post_backward + 3 inline checks during backward orchestration (slots 33/34/35 multi-point bw_d_h_s2 snapshots). - Sentinel correctly fires on cuBLAS GEMM accumulator overflow. 2. F1 NaN entry kernels pinpointed (slots 26 + 32): - apply_iqn_trunk_gradient cuBLAS sgemm (slot 26 = iqn_trunk_m) - Bottleneck Linear backward sgemm (slot 32 = bn_d_concat_buf) - IQN-internal buffers (slots 27, 28, 35) stayed CLEAN — the IQN backward chain is NOT the seed; the cuBLAS consumer of clean inputs produces NaN via accumulator overflow. 3. Cold-start clamp pathology fixed (ε floor formula): - Original `(1e6 × isv).max(1e3)` clipped F1-startup gradients to ±1e3 when fold-boundary ISV reset put H_S2_RMS_EMA ≈ 0. - New `1e6 × isv.max(1.0)` guarantees max_abs ≥ 1e6 in all states. - Validated: F1 NaN moved from step 240 → 3540 (15× later). What is SP3 territory (NOT fixed in SP1) ----------------------------------------- F1 NaN at step 3540 has same `[6, 12, 26, 32]` signature with clean inputs (slots 27, 35). The cuBLAS GEMM produces NaN/Inf because the WEIGHTS being multiplied have accumulated NaN/Inf via Adam updates over 3540 training steps. Surgical clamps on gradient inputs cannot prevent this; root cause is Q-target inflation drives Adam EMA saturation drives weight pathology drives GEMM overflow. SP3 scope per spec: - Target-Q clipping or pessimistic ensemble - Atom-position growth bounds (C51) - Conservative target sync at fold boundary - Adam EMA reset frequency for fold transitions What is SP2 territory (NOT fixed in SP1) ----------------------------------------- F0 regression: 35.30 (Phase B+C+ε-fix) vs 55.87 baseline. Three consecutive smokes show consistent F0 ≈ 35 across post-Phase-B commits, indicating Phase B instrumentation has real F0 cost. SP2 framework codification opportunity: consolidate the 11 per-step check_nan_f32 launches into a single fused reduce kernel scanning all 12 backward-path slots in parallel. Expected: F0 returns to ~55 baseline, regression sentinel preserved. Pearls for memory ----------------- 1. ε-floor placement: when an ISV-driven adaptive bound has a cold-start safety floor, the ε goes on the ISV MULTIPLIER (`isv.max(1.0)`), not the BOUND (`bound.max(1e3)`). Putting ε on the bound clips legitimate signal at cold-start; putting it on the multiplier preserves the wide guard-band intent. 2. NaN diagnostic ordering: when adding sanitization (clamp, isfinite-or-zero) to a buffer that's also NaN-checked by a post-backward sentinel, the inline NaN check MUST fire BEFORE the sanitization. Otherwise the sentinel sees zero'd buffers and silently dies. References ---------- - Spec: docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md - Plan: docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md - Audit (durable for SP2/SP3): docs/dqn-backward-nan-audit.md - Memory entry: ~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_sp1_f1_nan_root_cause_resolved.md	2026-04-30 03:05:54 +02:00
jgrusewski	ab2133463e	fix(dqn): SP1 Phase C — ε floor fix-up #2 (cold-start clamp pathology) ... Smoke smoke-test-dr2bn (commit 19b008e1c) F1-NaN'd at step 240 — earlier than pre-fix smoke smoke-test-xvzgk (step 890). The fix made things worse, indicating the ε floor `(1e6 × isv).max(1e3)` is actively destabilizing F1 startup. Diagnosis: at fold boundary, ISV[H_S2_RMS_EMA_INDEX=96] and ISV[Q_DIR_ABS_REF_INDEX=21] reset to 0 (per StateResetRegistry). Formula `(1e6 × 0).max(1e3) = 1e3` makes max_abs aggressively narrow. F1 startup gradients can have natural magnitudes > 1e3 (post-fold Bellman-target shift); clamping them to ±1e3 destabilizes Adam EMAs, which then drive cuBLAS GEMM accumulators into pathological inputs that overflow → slot 26 + 32 NaN. New formula: `1e6 × isv.max(1.0)` guarantees max_abs ≥ 1e6 regardless of ISV state: - ISV = 0 → max_abs = 1e6 × max(0, 1.0) = 1e6 - ISV = 0.5 → max_abs = 1e6 × max(0.5, 1.0) = 1e6 - ISV = 2.0 → max_abs = 1e6 × max(2.0, 1.0) = 2e6 - ISV = 100 → max_abs = 1e8 F0 no-op intent preserved (F0 inputs ≪ 1e6 in all states; ISV[96]≈1.0 for converged F0 → max_abs = 1e6, well above F0-typical |iqn_d_h_s2| ≤ ~10²). F1 startup gradients ≤ 1e6 are un-clipped. The 1.0 ε floor is on the ISV multiplier (Invariant 1 carve-out per `feedback_isv_for_adaptive_bounds`), not on the bound itself — bound is still ISV-driven when ISV is meaningful (≥ 1.0). Two edit sites: - gpu_dqn_trainer.rs:6967 (apply_iqn_trunk_gradient, slot 26 path) - gpu_dqn_trainer.rs:18631 (launch_cublas_backward_to, slot 32 path) F0 regression to 35.24 still unsolved — separate investigation thread within SP1 (no deferral; Phase B instrumentation timing impact suspected).	2026-04-30 02:16:45 +02:00
jgrusewski	0c99e08002	docs(dqn): SP1 Phase D — multi-fold validation FAIL (criterion 7/7) ... Smoke smoke-test-dr2bn (commit 19b008e1c) failed all 7 SP1 pass criteria. F1 first-fire signature at step 240: flagged=[6=grad_buf, 12=save_h_s2, 26=iqn_trunk_m, 32=bn_d_concat_buf] — IDENTICAL to pre-fix smoke-xvzgk but at step 240 vs 890. The surgical fix at 19b008e1c made F1 NaN EARLIER, not prevented it. Diagnosis: the ε floor `(1e6 * isv).max(1e3)` clips legitimate F1 startup gradients to ±1e3 when fold-boundary ISV reset puts ISV[96] or ISV[21] near 0. Clipped gradients destabilize Adam EMAs → cuBLAS GEMM accumulator overflow → slot 26 + 32 NaN. F0 regression unchanged across Phase C (35.24 vs pre-fix 34.55). Confirms the F0 regression is a Phase B instrumentation side effect, not Phase C. Separate investigation thread within SP1. Next iteration: Task 6 fix-up #2 changes ε floor to `(1e6 * isv.max(1.0))` guaranteeing max_abs ≥ 1e6 regardless of ISV state. Removes cold-start clamp pathology while preserving F0 paper-review intent.	2026-04-30 01:50:33 +02:00
jgrusewski	19b008e1cb	fix(dqn): SP1 Phase C — preserve NaN diagnostics + correct F0 paper-review ... Quality-review follow-ups to commit 97f1d25f5: 1. CRITICAL: preserve regression sentinel for slots 24, 25, 26, 32. The post-clamps in apply_iqn_trunk_gradient + launch_cublas_backward_to ran BEFORE run_nan_checks_post_backward, silently zeroing the buffers the diagnostic was supposed to observe. Fix: inline check_nan_f32 calls IMMEDIATELY BEFORE each launch_clamp_finite_f32 — same pattern as the existing slot 35 check at line 6949. Now the diagnostic fires on NaN entry; the clamp then sanitizes (preventing propagation but preserving the flag — flags are sticky/OR'd in dqn_nan_check_f32, never cleared). 2. CRITICAL: correct F0 paper-review line reference + reasoning. The commit body cited "5.0 norm-clip at line 16747" — that's a memset, not a clip. Real norm-clip lives at lines 16827-16868. Also tightened the L2-vs-per-element reasoning: L2 norm <= 5.0 worst-case bounds per-element |x| <= 5.0 (single-element edge), still several orders of magnitude below the 1e6 max_abs guard. 3. Inline comment at line 6951 mislabeled the clamp target as 'slot 27 input' — actually clamps the slot 35 buffer (bw_d_h_s2 post-DtoD). Fixed. 4. Articulated pre-clamp rationale: defense-in-depth regression protection for input buffers (slots 24, 25, 35) against future pathologies that could create extreme-but-finite inputs. The F1 ep2 NaN was outputs overflowing finite inputs, but the same fix template covers both failure modes per feedback_no_partial_refactor. 5. Renamed clamp_finite_f32_kernel -> dqn_clamp_finite_f32_kernel for consistency with sibling utility kernels (dqn_nan_check_f32, dqn_zero_kernel, dqn_grad_norm_kernel). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 01:27:23 +02:00
jgrusewski	97f1d25f54	fix(dqn): SP1 Phase C — F1 NaN cuBLAS-overflow surgical fix (slots 26 + 32) ... Patches two cuBLAS GEMM backward operations identified by SP1 Phase B smoke smoke-test-xvzgk (commit f139a63ee) as the F1 ep2 NaN sources: 1. apply_iqn_trunk_gradient (gpu_dqn_trainer.rs:6885+): - Pre-call: clamp bw_d_h_s2 (the IQN DtoD-overwrite target that feeds the GRN trunk encoder's cuBLAS Linear_a/Linear_b dW/dX/dB GEMMs; symmetric with slot 27 source-side check) to ±1e6×ISV[96]. - Post-call: sanitize iqn_trunk_m (slot 26 output) — isfinite-or-zero + magnitude clamp; defence-in-depth. 2. launch_cublas_backward_to main backward path: - Pre-call: clamp d_value_logits_buf + d_adv_logits_buf (slots 24/25 inputs that feed bw.backward_full's dueling/branch backward GEMMs) to ±1e6×ISV[21]. - Post-call: sanitize bn_d_concat_buf (slot 32 output, gated on bottleneck_dim > 0). Both fixes use the new clamp_finite_f32_kernel utility (CUDA — replaces NaN/Inf with 0 + clamps finite values to ±max_abs) + launch_clamp_finite_f32 (Rust wrapper). Kernel lives in dqn_utility_kernels.cu (already in build.rs); loaded into the same module as the NaN check kernels in compile_training_kernels — both call sites land inside captured children (forward_child for slot 32, aux_child for slot 26) and use only stream- bound launch_builder, so the kernel is graph-safe. ISV-driven max_abs bound = 1e6 × isv[<slot>] with 1e3 ε floor for uninitialized state (Invariant 1 carve-out per feedback_isv_for_adaptive_bounds). Wide guard band — F0 inputs sit several orders of magnitude below the threshold (F0 ISV[96] ≈ 1.0 LayerNorm RMS, F0 ISV[21] ≈ 1.0 Q-value scale; F0 |iqn_d_h_s2| ≤ ~10², F0 |d_value_logits| ≤ ~10³ post the existing 5.0 norm-clip at line 16747), so guards are no-ops on F0; F1 overflows trigger clamp. F0 paper-review pre-smoke confirmed. Combined RELATED commit per feedback_no_partial_refactor — both kernels share the same unsafe-pattern (cuBLAS GEMM extreme-intermediate-product overflow) + the same fix template, so they ship together. Phase B slots 24-35 remain as permanent diagnostic infrastructure; will catch any future regression of this NaN class. SP1 Phase D validation smoke pending. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 01:10:46 +02:00
jgrusewski	2ba0eef718	docs(dqn): SP1 Phase B smoke result — F1 NaN topology captured (slots 26 + 32) ... Smoke smoke-test-xvzgk (commit f139a63ee) confirms Phase B instrumentation fires correctly: backward-path slots 24-35 capture the F1 ep2 explosion's NaN topology beyond the prior visible-only [6, 12] (grad_buf + save_h_s2). F1 first-fire signature: flagged=[6, 12, 26, 32] - 26 iqn_trunk_m — apply_iqn_trunk_gradient cuBLAS bwd output - 32 bn_d_concat_buf — Bottleneck Linear backward dy - 6, 12 = downstream propagation (existing visible flags) CLEAN at first-fire: slots 27 (iqn_d_h_s2_buf), 28 (d_branch_logits_buf), 33/34/35 (bw_d_h_s2 multi-point), 24/25 (d_value/adv_logits_buf), 29 (cql), 30 (aux). Source kernels identified (drives Task 6 surgical fix): 1. apply_iqn_trunk_gradient cuBLAS sgemm (gpu_dqn_trainer.rs:6843+) — matches session_2026-04-05 residual-8% finding, never closed. 2. Bottleneck Linear backward cuBLAS sgemm — NEW finding (not in audit's original top-3 ranking). cuBLAS GEMM produces NaN from clean inputs; suggests extreme intermediate products at F1's Bellman-target shift. Both are template (f) cuBLAS-wrapper fixes — input range guard + output sanitization, ISV-driven bounds (Q_ABS_REF_INDEX=16, Q_DIR_ABS_REF_INDEX=21, H_S2_RMS_EMA_INDEX=96 — all existing slots, no new ISV). F0 Sharpe regression to 34.55 (from 55.87 baseline) noted as concern. Possible stochastic variance OR Phase B instrumentation timing impact; calibrate with Task 6 fix smoke. No deferral per operating principles. F2 cascade (steps 5+): flagged adds slots 2, 3, 7, 8, 24, 25, 29, 33, 34 once F1's corrupted Adam EMAs + weights flow into F2. Slots 27, 28, 35 STAY CLEAN throughout — IQN-internal backward chain is NOT the seed; the cascade enters IQN's consumer but never the IQN producer.	2026-04-30 00:59:09 +02:00
jgrusewski	f139a63eea	feat(dqn): SP1 Phase B instrumentation — backward NaN checks (slots 24-30, 32-35) ... Wires per-step NaN checks on backward-path kernel outputs. Coverage per audit per-slot accessor table (docs/dqn-backward-nan-audit.md :530-548). GpuDqnTrainer::run_nan_checks_post_backward (NEW) — fire-once-at-end: - 24 d_value_logits_buf (post-c51_grad value gradient) - 25 d_adv_logits_buf (post-c51_grad branch advantage) - 26 iqn_trunk_m (apply_iqn_trunk_gradient cuBLAS bwd output) - 27 iqn_d_h_s2_buf (IQN backward dh_s2; arg from FusedTrainingCtx) - 28 d_branch_logits_buf (IQN production backward; arg from caller) - 29 cql_d_value_logits (CQL gradient output) - 30 aux_dh_s2_nb_buf (Aux next-bar backward dh_s2) - 32 bn_d_concat_buf (Bottleneck Linear backward dy) Inline checks during backward orchestration: - 33 bw_d_h_s2 post-main (in launch_cublas_backward_to, after backward_full + branch concat accum, BEFORE aux_heads_backward SAXPY) - 34 bw_d_h_s2 post-aux (in launch_cublas_backward_to, after aux_heads_backward SAXPY) - 35 bw_d_h_s2 post-iqn (in apply_iqn_trunk_gradient, after graph_safe_copy_f32 DtoD overwrite, BEFORE encoder_backward_chain consumes) Sequential 33→34→35 fire pattern localises the NaN entry point: - 33 alone fires → main backward chain (c51 + MSE + branch concat) - 34 fires after 33-clean → aux SAXPY (aux_heads_backward) - 35 fires after 34-clean → IQN DtoD or per-sample IQN backward Slot 31 (ensemble_d_logits_buf) cleanly deferred per Task 2 commit 387335e2b — owner is FusedDqnTraining (different struct); will be un-deferred when ensemble Phase B saxpy guards are verified. IQN pointers (slots 27, 28) passed as Option<u64> from FusedTrainingCtx because GpuDqnTrainer does NOT own GpuIqnHead (Task 3 deviation finding, audit lines 535-536). None case (when iqn_lambda == 0.0 and gpu_iqn = None) cleanly skips slots 27/28 — semantically honest "buffer doesn't exist this run" rather than false-clean signal. Pattern matches apply_iqn_trunk_gradient(iqn_d_h_s2_ptr: u64, ...) at gpu_dqn_trainer.rs:6843. Call sites (atomic — feedback_no_partial_refactor): - fused_training.rs:1519 ungraphed step path - fused_training.rs:2324 capture_training_graph closure (post_aux child) - gpu_dqn_trainer.rs::launch_cublas_backward_to (slots 33, 34 — captured in forward child via submit_forward_ops_main) - gpu_dqn_trainer.rs::apply_iqn_trunk_gradient (slot 35 — captured in aux child via submit_aux_ops) Each check uses existing check_nan_f32(buf_ptr, len, flag_idx) — single-block GPU reduce, no atomicAdd, no DtoD/HtoD/HtoH copies, no per-step DtoH. Lengths use CudaSlice::len() where possible (auto-syncs with allocator padding); inline arithmetic for slot 28 since d_branch_logits_buf.len() is private to GpuIqnHead. Flags accumulate within fold; reset at fold boundary via reset_nan_flags(). Readback flow (commit d1808df14) consumes them in BOTH halt_nan and halt_grad_collapse paths via name tables annotated in commit 387335e2b. Permanent diagnostic infrastructure — stays as production-grade regression sentinel after the surgical fix lands.	2026-04-30 00:32:35 +02:00
jgrusewski	797f8bf326	feat(dqn): SP1 Phase B accessors — backward-buffer NaN check pointers ... Adds 5 new pub(crate) accessor methods exposing backward-path buffer device pointers for Task 4's per-step NaN checks (slots 24, 25, 28, 29, 30 per audit per-slot table at docs/dqn-backward-nan-audit.md :530-548): GpuDqnTrainer (4 new): - d_value_logits_buf_ptr (slot 24 — post-c51_grad value gradient) - d_adv_logits_buf_ptr (slot 25 — post-c51_grad branch advantage) - cql_d_value_logits_ptr (slot 29 — CQL gradient output) - aux_dh_s2_nb_buf_ptr (slot 30 — aux next-bar backward dh_s2) GpuIqnHead (1 new): - d_branch_logits_buf_ptr (slot 28 — production IQN backward output, iqn_quantile_huber_loss) Slot 27 (iqn_d_h_s2_buf) reuses existing GpuIqnHead::d_h_s2_raw_ptr() at gpu_iqn_head.rs:1660 — no new method per feedback_no_legacy_aliases: the existing accessor is already public and sufficient; renaming + chasing the single call site adds churn without value. Slots 26, 32, 33-35 reuse pre-existing handles: - 26: self.ptrs.iqn_trunk_m - 32: self.bn_d_concat_buf() (existing, returns &CudaSlice<f32>) - 33-35: self.ptrs.bw_d_h_s2 (3 different Task 4 call sites) Slot 31 (ensemble_d_logits_buf) deferred per Task 2 commit 387335e2b (cross-struct on FusedDqnTraining). DEVIATION FROM PLAN: the plan called for "delegate accessors on GpuDqnTrainer for slots 27/28" — structurally invalid because GpuDqnTrainer does NOT own GpuIqnHead. The IQN head is owned by FusedTrainingCtx (fused_training.rs:289) alongside the trainer at line 234. The audit's per-slot accessor table (lines 535-536) is correct: accessors land on GpuIqnHead. Task 4's run_nan_checks_post_backward will receive IQN pointers as u64 arguments from the FusedTrainingCtx call site — same pattern already in use at gpu_dqn_trainer.rs:6843 (apply_iqn_trunk_gradient(&mut self, iqn_d_h_s2_ptr: u64, ...)). NO new scratch buffer added — the plan's bw_d_h_s2_pre_saxpy scratch + DtoD-copy approach was superseded by the audit's 3-call-site reformulation. Slots 33/34/35 are post-main / post-aux / post-iqn snapshots of the same bw_d_h_s2 (one buffer, three Task 4 invocations). Pattern follows commit e9096c7be's GRN-block accessors (concise pub(crate) fn name_ptr(&self) -> u64 with doc-comment referencing slot number + audit doc + buffer semantics). Additive — no behavioral change; new accessors consumed by Task 4's NaN check call sites. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 00:17:39 +02:00
jgrusewski	387335e2b9	fix(dqn): SP1 Phase B foundation — stale-doc cleanup + Task 4 prep ... Quality-review follow-ups to commit 53bc0bc50 (nan_flags_buf 24->48): 1. Update run_nan_checks_post_forward docstring (gpu_dqn_trainer.rs:14842-14873): replace 24-slot map with 48-slot range summary + audit-doc pointer. Was actively misleading after the 24->48 expansion; partial-refactor residue per feedback_no_partial_refactor. 2. Update '[24] system' comment in training_loop.rs (around line 2044): reflect the 48-slot post-expansion state (slots 0-23 fwd, 24-35 bwd, 36-47 reserved). Also fix stale '0..11' tracing message to '0..47'. 3. Slot 31 (ensemble_d_logits_buf) annotation: flag DEFERRED + owner on FusedDqnTraining (different struct than 24-30, 32-35). Prevents Task 4 from blanket-launching check_nan_f32 on slot 31's null accessor. 4. Both name-table header comments now reference the future run_nan_checks_post_backward method (Task 4) plus the audit's per-slot table — pre-empts contract drift when Task 4 lands a 3-way name-table dependency. Audit-doc entry appended to docs/dqn-wire-up-audit.md SP1 Phase B section. No behavioral change. Both name tables remain byte-identical. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-30 00:10:37 +02:00
jgrusewski	53bc0bc505	feat(dqn): SP1 Phase B foundation — nan_flags_buf 24→48 ... Expands the NaN flag buffer from 24 to 48 slots to make room for backward-path NaN checks (slots 24-35 per audit doc per-slot table) plus 12 reserved headroom slots (36-47) for SP2 framework + SP3 observer hooks. Touches: - gpu_dqn_trainer.rs: alloc size 24→48 (line ~11178); read_nan_flags signature [i32; 24]→[i32; 48] (line ~14982); field docstring updated (line ~2658) to reflect 48-slot layout - fused_training.rs: pub(crate) read_nan_flags signature [i32; 24]→[i32; 48] - training_loop.rs: BOTH name table sites (halt_nan + halt_grad_collapse block from commit d1808df14) updated to 48 entries Slot names use audit allocation (docs/dqn-backward-nan-audit.md per-slot table), which supersedes the plan's placeholder names per the audit's plan-supersession note. Slots 24-35 cover production buffers: d_value_logits_buf, d_adv_logits_buf, iqn_trunk_m, iqn_d_h_s2_ptr, d_branch_logits_buf, cql_d_value_logits, aux_dh_s2_nb_buf, ensemble_d_logits_buf, bn_d_concat_buf, bw_d_h_s2 (×3 call sites). Slots 36-47 are rsv36-rsv47 (headroom). No behavioral change — new slots stay at zero until Task 4 wires the kernel-output NaN checks. Buffer size reviewable by SP2.	2026-04-30 00:00:43 +02:00
jgrusewski	2406ebf2f7	docs(dqn): SP1 Phase A — slot 24/25 buffer correction + dormancy reconciliation ... Three review-driven fixes: 1. Slot 24/25 cite the actual production buffers d_value_logits_buf (gpu_dqn_trainer.rs:3123) and d_adv_logits_buf (:3125) — the f32 atomicAdd buffers that launch_c51_grad writes — not the staging buffers at :3127/:3129. Pointer expressions exist at the launch site (:17707/:17708), making new accessors optional. Task 3 summary and priority-list entries updated to match. 2. Explicit plan-supersession note inserted immediately above the per-slot table: the audit's per-slot allocation supersedes the plan Task 2 step 6 name table. The plan's allocation was placeholder; the audit's is grounded in per-kernel inspection. Task 2 should use the audit's names (d_value_logits_buf, d_adv_logits_buf, iqn_trunk_m, iqn_d_h_s2_ptr, d_branch_logits_buf, etc.). 3. Summary suspicion-ranking row #3 reconciled with slot-28 dormancy: iqn_backward_per_sample has no Rust caller (verified via grep crates/ml/src/), so row #3 is re-pointed at the production kernel iqn_quantile_huber_loss (iqn_dual_head_kernel.cu:1346-1413, loaded at gpu_iqn_head.rs:2114). Same unsafe-write pattern (no isfinite guard at line 1410's d_q_online[idx] = qw*d_huber/Q), now attributed to the live path. apply_iqn_trunk_gradient at #1 stands — its reasoning (orchestrator consuming iqn_d_h_s2_ptr) is unchanged by which specific kernel writes that buffer. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 23:56:30 +02:00
jgrusewski	8fd9cc0603	docs(dqn): SP1 Phase A — audit citation + accessor precision ... Quality-review fixes for the SP1 Phase A audit: 1. nan_flags_buf [i32;24] declaration: quad-cite (field 2659, alloc 11178, constructor 12563, read_nan_flags signature 14982). Task 2's 24->48 expansion must touch all four atomically per feedback_no_partial_refactor; consumer site in trainers/dqn/fused_training.rs and the two name-table sites in trainers/dqn/trainer/training_loop.rs are still listed in Task 2's plan body. 2. Per-slot Rust buffer-pointer expression added for slots 24-35. Already-accessible (no new accessor): slots 26, 32, 33, 34, 35 (4 via self.ptrs, 1 via existing bn_d_concat_buf accessor). Need new accessor on GpuDqnTrainer: slots 24 (d_value_logits), 25 (d_adv_logits), 30 (aux_dh_s2_nb_buf), and 29 (cql_d_value_logits, only if un-deferred). Need new accessor on GpuIqnHead: slot 28 (d_branch_logits_buf — note: production IQN backward uses iqn_quantile_huber_loss, NOT iqn_backward_per_sample which is declared but never loaded). Optional: slot 27 (d_h_s2_buf_ptr) — can be inlined inside apply_iqn_trunk_gradient instead. Need new accessor on FusedDqnTraining: slot 31 (only if un-deferred). Becomes input for Task 3. 3. Citation typo: c51_loss_kernel.cu line 274 reference removed — line 274 is __syncthreads() in the projection-reduction warp loop; the a_std=sqrtf reference belongs to c51_grad_kernel.cu:274. Loss-kernel sqrtf sites are 779 and 804. Verified by reading each cited line; SQLX_OFFLINE=true cargo check --workspace passes (docs-only).	2026-04-29 23:48:31 +02:00
jgrusewski	6994b9cf97	docs(dqn): SP1 Phase A — backward NaN audit ... Read-only γ audit of every backward-path kernel writing to bw_d_h_s2 / grad_buf / save_h_s2 accumulators. Per-kernel inventory against unsafe-pattern checklist (sqrtf-neg, 1/0, logf-≤0, expf-large, EMA variance, atomicAdd/saxpy NaN-propagation). Each section includes: identified pattern + line citation, severity, proposed guard form, ISV bound option, F0 risk, Phase B flag-slot allocation. Cross-referenced session_2026-04-05_nan_investigation.md residual 8% step-2 NaN in apply_iqn_trunk_gradient — never closed; ranked top suspect for SP1. Output drives SP1 Phase B (12 new flag slots in nan_flags_buf 24→48) and Phase C surgical fix(es). Durable artifact for SP2 framework codification + SP3 structural-fix scoping.	2026-04-29 23:35:29 +02:00
jgrusewski	05958d3a0c	plan(dqn): SP1 numerical stability — F1 NaN root-cause investigation ... Implementation plan for SP1 (Sub-project 1 of 3) of the numerical stability investigation. Follows the γ + β methodology from the spec at docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md. 8 tasks across 4 phases: - Phase A (Task 1): γ read-only audit producing docs/dqn-backward-nan-audit.md - Phase B (Tasks 2-5): always-landing β instrumentation expanding nan_flags_buf 24→48 with 12 new backward-kernel NaN check slots + 12 reserved slots for future coverage - Phase C (Task 6): surgical fix(es), content-driven by audit + smoke topology, ISV-driven for any dynamic bound (mandatory) - Phase D (Task 7): multi-fold L40S smoke validation against 7 pass criteria (F0 ≥ 95% baseline, F1+F2 monotone improvement, zero NaN-CLAMPED-TO-ZERO, all 48 NaN flag slots remain at zero) - Closure (Task 8): audit doc closure, memory entry, SP2/SP3 handoff Operating principles (mandatory per spec): - No deferrals — anomalies discovered during investigation get fixed within SP1, not punted to SP2/SP3 - Combined RELATED fixes ship as rich commits (per feedback_no_partial_refactor) - ISV-driven design for any dynamic bound Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 23:21:26 +02:00
jgrusewski	792812baa1	spec(dqn): SP1 numerical stability — revise per critical review ... 9 substantive issues addressed inline: 1. ISV-driven design elevated from 'if applicable' to MANDATORY for all dynamic bounds in SP1 fixes. Numerical-stability ε bounds are the only carve-out (Invariant 1). Hardcoded tuning constants for dynamic ranges explicitly rejected. 2. F0 Sharpe regression criterion changed from absolute (≥55) to ratio-based (≥95% of latest baseline; floor 53.08 currently). Prevents iterative erosion across multiple fix commits. 3. 'F1 trending positive' replaced with concrete monotone-improvement test: Best Sharpe at last epoch ≥ Best Sharpe at first epoch of the same fold. 4. Pass criterion distinguishes NaN-CLAMPED-TO-ZERO (failure) from 'Genuine grad collapse' (legitimate observation, permitted) per the existing infrastructure from commit d1808df14. 5. Multi-source NaN scenarios explicitly supported — γ + β may identify multiple kernels; SP1 fixes ALL within the same cycle. 6. F0-safety paper-review gate added BEFORE smoke validation. Audit doc carries 'F0 risk' (low/medium/high) per proposed fix; high- risk fixes get math-on-paper inspection before consuming L40S. 7. Audit doc structure now requires 'ISV bound option' column — forces ISV-first thinking at audit stage, not as afterthought. 8. Anti-patterns expanded: micro-clamping (per-op clamping that hides upstream causes); combining unrelated fixes (anti-pattern of the rich-commit principle); hardcoded constants for dynamic bounds. 9. 48-slot allocation explicitly justified (24 used + 12 new + 12 headroom) AND marked reviewable by SP2 if right-size differs. All 5 design sections preserved structurally; revisions integrated inline. Per brainstorming skill: spec self-review fixes applied without re-review cycle.	2026-04-29 22:59:00 +02:00
jgrusewski	dab5990287	spec(dqn): SP1 numerical stability investigation — F1 NaN root-cause ... Brainstorm session 2026-04-29 produced this spec scoping Sub-project 1 of a three-sub-project decomposition addressing the F1 ep2 NaN explosion that persists across 9+ defensive layers in Plan C Phase 2 (commits e445d07a..e9096c7be). Decomposition: - SP1 (this spec): F1 NaN root-cause; γ audit + β always-landing instrumentation + surgical fix + multi-fold validation - SP2 (future): numerical stability framework — codify guards - SP3 (future): Q-learning structural stability — target-Q clip, pessimistic ensemble, atom-range governance Operating principles: - No deferrals (anomalies fixed within SP1, not punted) - Combined fixes — rich commits (per feedback_no_partial_refactor) - Always-landing diagnostic instrumentation (24-slot nan_flags_buf expands to 48; permanent regression sentinel) - F0 Best Sharpe ≥ 55 preserved (no regression on the working path) Pass criterion: all 3 folds train 5 epochs, F0+F1+F2 Sharpe ≥ 0, zero NaN-CLAMPED-TO-ZERO log lines, all 48 flag slots remain zero. 5 design sections approved iteratively: Architecture, Components, Data Flow, Error Handling, Testing. Next: writing-plans produces implementation plan.	2026-04-29 22:53:13 +02:00
jgrusewski	e9096c7be1	fix(dqn): R1 — eliminate K Adam shrink + P — GRN-stage NaN checks ... R1: K's hardcoded shrink-and-perturb (m×0.1, v×0.01) at fold boundary violated feedback_adaptive_not_tuned (untracked tunable knobs) AND created a downstream pathology: tiny v_hat denominator → oversized Adam updates 50+ steps post-reset → trunk param overshoot → save_h_s2 NaN at F1 ~step 1745 (smoke-test-bkdx5 diagnostic). K was introduced (commit 4ef1d8ebb) BEFORE fold_warmup_factor existed in the same commit's "K + adaptive warmup" pair. With warmup_factor in place — ISV-driven, dampens lr+clip via lr_eff = lr_base × max(MIN_WARMUP_LR_FRAC, fold_warmup_factor) — K is redundant. Single mechanism, ISV-driven, no hardcoded constants. Eliminating K leaves m,v reset to 0 at fold boundary; warmup_factor handles cold-start. P: expanded nan_flags_buf 16→24 with 5 GRN-stage checks (linear_a_out, elu_out, linear_b_out, glu_sigmoid_out, layernorm_var/out) for finer-grained source identification if R1 alone doesn't fix F1. Predicted outcomes: - If K's tiny-v_hat was the cause: F1 trains successfully (R1 alone) - If different mechanism: new GRN-stage flags pinpoint which sub- stage produces NaN, enabling layer-level fix Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 22:11:46 +02:00
jgrusewski	d1808df14c	diag(dqn): read NaN flags on grad-collapse path too (find F1 ep2 explosion source) ... Existing nan_flags_buf [16] covers 13 buffers with per-step NaN checks inside the captured training graph. But read_nan_flags() was only called when training guard's halt_nan fired — which checks pinned- readback grad_norm. NaN-clamped-to-zero gradients reach the pinned scalar as 0, not NaN, so halt_nan never fires for the explosion case. The collapse path (halt_grad_collapse, grad < 0.01) was firing instead without reading the flags. Add flag readback in that path when gr.raw_grad_norm < 1e-6 (suspicious zero) — logs NaN-CLAMPED-TO-ZERO with flagged buffer names. Genuine near-zero gradients get a separate "no NaN flags set" log so we can distinguish the two cases. This should pinpoint which kernel produces the F1 ep2 NaN first (C51 KL projection? IQN aux? CQL? aux heads?). Diagnostic only — keep after fix lands; correct gate for future regressions.	2026-04-29 21:40:39 +02:00
jgrusewski	7640a681c5	diag(dqn): per-step F1 explosion diagnostic — pinpoint which signal blows first ... After 6+ intervention layers (A.1+A.2+A.3+F+H+K+warmup+N) F1 ep2 still explodes with grad_norm=2.66T while F1 ep1 trains cleanly. None of our defenses catch the explosion path. Need data to pinpoint WHICH signal explodes first. Per-step FOLD_EXPLOSION_DIAG in fold >= 1: when grad_norm > 1000 OR jumps 5x from prior guard step, fire diagnostic warn with: - loss decomposition: total / c51 / mse / iqn (pinned readback, no DtoH) - Q range: q_min / q_mean / q_max via cold-path reduce - atom positions: per_sample_support[0,d0] + [0,d2] (v_min, v_max, dz) - prev grad norm + ratio for context - trigger flag + post-trigger trajectory countdown After trigger, continues emitting for 5 guard steps so the explosion trajectory is captured (not just the first crossing). Plus an unconditional FOLD_EXPLOSION_DIAG[F1_END_EP1] baseline log at the end of fold 1 epoch 0 — healthy state immediately preceding the explosion. Compare-and-contrast with per-step explosion frames pins the runaway driver. CQL and ensemble losses are not pinned-readback (transient device buffers consumed inside the training graph) and are explicitly absent from the decomposition. If none of {c51, iqn, mse} is the runaway driver but total_loss still explodes, that implicates the unexposed CQL/ens path — the absence is itself diagnostic information. Implementation: - Three diagnostic-only fields on DQNTrainer: current_fold, last_logged_grad_norm, explosion_diag_steps_remaining. Reset last_logged + remaining in reset_for_fold; current_fold overwritten unconditionally at fold-loop entry. - Three accessors on FusedTrainingCtx: explosion_diag_loss_components (pinned), explosion_diag_atom_range (12-elem cold DtoH), explosion_diag_q_range (reduce + 28B DtoH). - Three accessors on GpuDqnTrainer: c51_loss_pinned_value, mse_loss_pinned_value, stream_for_diag. Diagnostic-only — remove once root cause identified. Not for prod. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 21:16:48 +02:00
jgrusewski	250cee124e	fix(dqn): Winsorize adaptive_clip EMA input against single-sample outliers (N) ... Smoke smoke-test-xw4c6 showed F1 ep2 grad_norm=8.4B polluting the adaptive_clip EMA -> next adaptive_clip threshold became huge -> subsequent extreme grads passed unclipped -> NaN propagation -> grad collapse -> fold 1 fails despite K+warmup fixing the first-step problem. Add Winsorized clamp on raw_grad_norm before the adaptive_clip EMA update: clamped_sample = min(raw, K * previous_adaptive_clip) where K = 100 (numerical-stability bound, not a tuned constant — explicit "single sample can grow EMA at most 100x in one step" semantics). Companion observability log emits GRAD_CLIP_OUTLIER warning per clamp event so we can see when this fires in subsequent smokes. Fast/slow grad_norm EMAs (driving warmup factor) intentionally NOT winsorized — they're a stability signal that SHOULD respond to outliers, providing extra warmup damping when the system is unstable. Predicted impact on Plan C smoke F1 ep2: 8.4B grad -> clamped to ~3000 for EMA -> next adaptive_clip caps at ~1000 -> subsequent F1 batches get bounded gradients passed to Adam -> no NaN propagation. F1+F2 should now train through, exposing whether the underlying Q-target inflation requires further structural work or whether clipping alone suffices. Composes with K+warmup (4ef1d8ebb), A.2, A.1/A.3, F+H — fourth layer of the adaptive gradient-stability scaffolding (now five total: A.2 target-sync drift dampening, K+warmup post-reset transient dampening, N persistent clip-EMA outlier defense). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 20:51:50 +02:00
jgrusewski	4ef1d8ebb7	fix(dqn): Plan C K — Adam shrink-and-perturb + adaptive fold-warmup ISV ... Closes the F1 ep1 catastrophic-overshoot gap exposed by smoke-test-s9h4h (F0 succeeded with Best Sharpe = 36.03; F1 ep1 grad_norm = 355,009 — 5 orders of magnitude larger than F0 steady-state ~10 — leading to NaN propagation and grad-clamp-to-zero early stop). Combined fix: (1) Adam shrink-and-perturb at fold boundary (replace m=0/v=0 with m*=0.1, v*=0.01 to preserve direction while damping magnitude), and (2) a single adaptive ISV signal driving BOTH lr_eff and clip_eff dampening over the fold's first ~50 steps. K — Adam shrink-and-perturb in `GpuDqnTrainer::reset_adam_state`: - m *= 0.1, v *= 0.01 via existing `dqn_scale_f32_kernel` (loaded as `scale_f32_ungraphed`); t_pinned still zeroed so bias correction restarts. Architectural constants (preserve direction / lose magnitude history) per `feedback_isv_for_adaptive_bounds.md` Invariant 1 carve-out — not tuned. Composes with existing param shrink-and-perturb (`alpha=0.8`) in `FusedTrainingCtx::reset_for_fold`. Root cause for F1 overshoot: m=0,v=0 → first Adam step ≈ lr × g / ε → 6 OoM amplification. New CUDA kernel `fold_warmup_factor_kernel.cu`: - Single-block single-thread cold-path producer mirroring `q_drift_rate_ema_kernel.cu` / `moe_lambda_eff_kernel.cu` shape. - Reads two grad-norm EMAs (fast α=0.1, slow α=0.001) plus host-passed step counter; writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1). - Bootstrap branches (steps_observed < 200, slow EMA < 1e-6) emit factor=1.0 (no damping during cold-start). No atomicAdd; no DtoH. New ISV slot Q_DRIFT_RATE_INDEX → FOLD_WARMUP_FACTOR_INDEX = 130: - ISV_TOTAL_DIM 130 → 131; layout fingerprint shifts (checkpoint- incompatible per `feedback_no_legacy_aliases.md`, expected for a real architecture change). - FoldReset entries: `isv_fold_warmup_factor` → 0.0 and companion `isv_grad_norm_fast_ema` → 0.0 (lockstep reset per `feedback_no_partial_refactor.md`); slow EMA persists across folds as the cross-fold steady-state baseline. - Two new mapped-pinned scalars on GpuDqnTrainer (grad_norm_fast_ema_pinned, grad_norm_slow_ema_pinned) fed by `update_adaptive_clip` from the same `gr.raw_grad_norm` observation source as the existing adaptive clip EMA. Two consumers, both monotone (only dampen, never excite): - lr_eff = cosine_effective_lr × max(MIN_WARMUP_LR_FRAC=0.05, factor) via `set_lr` per-step. New `cosine_effective_lr_base` field on DQNTrainer composes the cosine schedule's per-epoch baseline with the warmup factor's per-step damping (rather than overriding the cosine schedule). - clip_eff = clip_base × (MIN_CLIP_FRAC=0.1 + 0.9 × factor) via new `set_active_clip` setter on FusedTrainingCtx + GpuDqnTrainer. Composes with the EMA-derived `clip_base = grad_norm_ema × 2` that `update_adaptive_clip` just wrote to the pinned slot. Numerical-stability bounds 0.05 / 0.1 are Invariant 1 carve-outs. Steady-state behaviour unchanged: factor=1 → lr_eff=lr_base, clip_eff=clip_base. Fold-boundary behaviour: factor starts at 0 → lr_eff = 0.05 × lr_base, clip_eff ≈ 0.1 × clip_base; rises to 1 over ~50 steps as the fast EMA catches up to the slow steady-state EMA. Predicted impact on Plan C smoke F1: 355,009-magnitude transient grad clipped to ~clip_base × 0.1 ≈ 1.0 (vs 10), Adam state shrunk instead of zeroed → first step update bounded; grad recovers normally over ~50 steps. Companion to A.1 (prev_epoch_q_mean reset), A.2 (adaptive Polyak-tau), A.3 (gradient_collapse_counter reset), F+H (kill-criterion robustness) — completes the fold-boundary state-reset family. Per `pearl_adaptive_moe_lambda.md` (kernel + ISV slot + bootstrap + reset + observability template), `pearl_cold_path_no_exception_to_gpu_drives.md` (GPU-stays-on-GPU even at cold-path cadence), `pearl_blend_formulas_must_have_permanent_floor.md` (lr/clip floors are permanent minimums), `feedback_adaptive_not_tuned.md` (lr+clip ISV-driven), `feedback_isv_for_adaptive_bounds.md` (factor IS the bound; consumers compose at runtime), `feedback_no_atomicadd.md` (single-thread reduce), `feedback_cudarc_f64_f32_abi.md` (slot index passed as i32), `feedback_no_partial_refactor.md` (kernel + slot + reset + producer + 2 consumers all land together), `feedback_no_quickfixes.md` (replaces brittle full-reset with adaptive damping; not threshold relaxation). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 20:24:36 +02:00
jgrusewski	ef243e771d	fix(dqn): reset gradient_collapse_counter + training_steps at fold boundary (A.3) ... Smoke smoke-test-vh9bj revealed: after F+H Q-drift kill fix, fold 0 trains successfully (5 epochs, Best Sharpe 36.03), but fold 1 fails at epoch 1 with "Gradient collapse detected for 5 consecutive epochs" despite epoch 1 grad_norm=296,417 (healthy). Root cause: the per-step gradient_collapse_counter on DQN was already near patience (5) from fold 0's late-epoch near-zero-grad steps, and DQNTrainer::reset_for_fold didn't clear it. Plus, training_steps accumulating across folds made the `past_warmup` gate always-true in fold 1+, removing the warmup grace period for data distribution shifts. Add DQN::reset_for_fold zeroing both. Wire from DQNTrainer::reset_for_fold alongside the A.1 prev_epoch_q_mean reset. Pure additive — same fold-boundary state-reset gap pattern as A.1, A.2, F. Predicted impact: fold 1+ now starts with counter=0 and warmup window restored; gradient collapse check has its full per-fold grace period.	2026-04-29 19:43:34 +02:00
jgrusewski	cca9dd36ae	fix(dqn): replace q_mean ratio with rolling-window MAD deviation (H — kill criterion robustness) ... Current criterion uses |q_mean| / |prev_q_mean| which explodes near zero crossings (legitimate cold-start has |prev_q_mean| ≈ 0.01-0.1, making any non-tiny current trigger the ratio threshold). smoke-test-n9xzr fired with prev=−0.0795, curr=0.7647, ratio=9.62× despite this being natural cold-start growth, not geometric runaway. Replace with rolling 5-epoch window of q_means. Compute median + MAD (Median Absolute Deviation, a 50%-breakdown estimator robust to single outliers); kill condition becomes |q_mean − median(window)| > 4.0 × max(MAD(window), 0.01) AND the existing adaptive floor. Constants are declared `const` near the kill block: - Q_DRIFT_WINDOW_SIZE=5 (matches smoke fold length, gives MAD a meaningful estimator without averaging across regime shifts) - Q_DRIFT_WARMUP_SAMPLES=3 (skip until ≥ 3 priors — smaller window degenerates to half-range MAD that trips on monotonic trajectories) - Q_DRIFT_DEVIATION_THRESHOLD=4.0 (4 MADs ≈ 2.7σ Gaussian- equivalent; clear outlier without firing on every legitimate dip; literal MAD count rather than σ because q_mean is non- Gaussian during cold-start) - Q_DRIFT_MIN_DEVIATION=0.01 (numerical-stability floor when window is constant) Window resets at fold boundary alongside prev_epoch_q_mean / adaptive_tau (A.1 pattern — cross-fold q-stats are independent training runs, mixing them would inflate MAD or shift the median). Predicted impact on smoke-test-n9xzr: - Plan C smoke fold 0 ep2: window has only 2 priors, warmup gate not yet satisfied → kill stays silent (was firing on ratio=9.62×) - Genuine geometric runaway (q_mean → 62.5 from baseline 0.5 over 4 epochs): ep3 deviation = |62.5 − 2.5|/MAD=2.0 = 30 > 4 AND |62.5| > floor=3×12=36 → kill fires correctly Both conditions ANDed (floor + deviation), preserving the production-safety semantics of the original criterion. The floor check is unchanged; only the divergence detector is replaced. Per feedback_no_quickfixes.md: this is a principled robust- statistics replacement, not a threshold relaxation. Per feedback_no_partial_refactor.md: window field, constants, criterion site, and fold-boundary reset land in lockstep — the kill criterion's contract is internally consistent across trainer/mod.rs, constructor.rs, and training_loop.rs. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 19:15:30 +02:00
jgrusewski	790c100719	fix(dqn): wire Q_ABS_REF / Q_DIR_ABS_REF EMA producers per-step (F — kill criterion adaptive floor) ... The Q-drift kill criterion's adaptive floor formula is max(0.5, 3 × max(ISV[Q_ABS_REF=16], ISV[Q_DIR_ABS_REF=21])) but smoke-test-n9xzr showed both ISVs stuck at bootstrap 0.05 even when q_mean reached 0.76 — the producers (q_mag_bin_means_reduce and q_dir_bin_means_reduce) only fired inside reduce_current_q_stats at epoch boundary, while the per-step captured update_isv_signals consumed the resulting scratch buffers each training step. Concrete failure mode at α=0.05 EMA with raw ≈ 1.0: - Epoch 0: scratch is zero-initialized, all per-step EMA pulls toward zero, ISV[16,21] stay 0.0 - Epoch 0 boundary: reduce_current_q_stats fires, scratch becomes q_abs_ref ≈ 1.0, single update_isv_signals writes ISV[16] = 0.05 - Epoch 1+: with stale scratch held constant between boundaries, per-step EMA over hundreds of steps would saturate — but smoke step counts are small enough (small batch=64, buffer=256) that the slot stays near 0.05 - kill_floor degenerates to max(0.5, 3 × 0.05) = 0.5 forever, so the criterion fires on legitimate cold-start growth Fix: launch q_mag_bin_means_reduce + q_dir_bin_means_reduce in both fused_training paths (captured adam_update_child at ~line 2228 AND ungraphed step-0 fallback at ~line 1465) immediately before update_isv_signals. Per-step graph replays now read q_out_buf that forward_child populated this same step, write fresh scratch, and the captured update_isv_signals consumes it on the same stream. Per pearl_cold_path_no_exception_to_gpu_drives.md: cold-path EMAs fire alongside their consumers. Per feedback_no_partial_refactor.md: graphed and ungraphed paths migrate together — same producer- consumer contract. After this fix, the floor will adapt: ISV[16,21] saturate to the policy's actual q_abs_ref scale within ~60 steps at α=0.05, so by epoch 2 the floor becomes 3 × ~0.5 = ~1.5, no longer firing on legitimate cold-start growth. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 19:10:48 +02:00
jgrusewski	f64eeb64da	fix(reward): bound bonus shaping by ISV[Q_DIR_ABS_REF] to break optimism loop ... C.4 timing bonus (experience_kernels.cu) and D.4b regime penalty multiplied by the trade-cumulative |reward| / |final_pnl| as an unbounded multiplicand. Across trades this compounds — bonus values inflate as Q values inflate, then bonus rewards inflate Q values further (Bellman bootstraps off shaped reward). Per pearl_one_unbounded_signal_per_reward + feedback_isv_for_adaptive_bounds: replace |reward|/|final_pnl| with ISV[Q_DIR_ABS_REF_INDEX]-bounded variant. The unbounded multiplicand becomes the direction-branch Q-scale EMA (already-tracked, gradient-decoupled), not the trade-cumulative shaping output. Cap is adaptive (matches Q magnitude as it evolves) and breaks the multi-trade compounding loop. Discovered during Plan C Phase 2 smoke diagnosis (researcher report 2026-04-29 a25f669e9df953174). a52d99613 baseline also hits bonus=235 — pre-existing structural pathology, not Plan-C-specific. Other shaping sites (D.4a persist, B.2 novelty, D.4c stable) already have all factors bounded — no fix needed. Surgical change set: - experience_env_step gains final `int q_dir_abs_ref_idx` param. - C.4 (line 2521): pnl_unit = |final_pnl|/max(ISV[21], eps); pnl_capped = min(pnl_unit, 1) * ISV[21]. - D.4b (line 2581): same pattern, reward_unit/reward_capped. - gpu_experience_collector.rs:3800 wires Q_DIR_ABS_REF_INDEX as i32. No new ISV slot, no new producer kernel, no layout-fingerprint shift — ISV[21] already produced by q_stats_kernel.cu since Plan 1. Predicted impact: - bonus EMA drops O(100-256) -> O(0.05-2.0) - rc[5] -> Bellman-target optimism loop broken - Plan C smoke: Q-drift kill at F0 ep2 likely no longer fires - a52d99613 baseline: same effect; validates pre-existing fix Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:37:43 +02:00
jgrusewski	b4d4a8d046	feat(dqn): Plan C A.2 — adaptive Polyak tau coupled to q-drift rate ... Wires the q_drift_rate_ema kernel + ISV slot Q_DRIFT_RATE_INDEX=129 (landed in fc8dbb0a8) into the production training path: - StateResetRegistry FoldReset entry isv_q_drift_rate -> 0.0 (companion to A.1's prev_epoch_q_mean reset; both ensure no cross-fold leakage of drift state). - reset_named_state dispatch arm in training_loop.rs. - Per-epoch launch_q_drift_rate_ema call after the Q-drift kill check, gated on the same `prev_epoch_q_mean.abs() > 1e-6` cold-start guard. Both q_mean_curr and q_mean_prev are in scope BEFORE the `prev_epoch_q_mean = q_mean` update, so the producer sees the correct delta. - tau_update_kernel.cu multiplies the cosine-scheduled, health-coupled tau_eff by `1 / (1 + clip(ISV[129], 0, 4))` so tau ranges [tau_base/5, tau_base] — monotone dampening under drift; healthy runs (drift ≈ 0) unaffected. - reset_for_fold comment in trainer/mod.rs notes the registry entry handles the new ISV slot. Predicted impact: Plan C fold 0 ep2 with q_mean(t)=0.82, q_mean(t-1)=-0.018, ISV[16]+ISV[21]≈0.1 -> drift_rate ≈ 8.4 -> clipped to 4 -> tau_eff = tau_base × 0.2 (5× slower target sync, dampens Q-target optimism through the inflation spike). a52d99613 fold 0 with q_mean staying ~0.21 -> drift_rate ≈ 0 -> tau_eff = tau_base (unchanged). Per pearl_adaptive_moe_lambda.md — canonical "EMA-tracked diagnostic drives a controller" pattern. Per pearl_cold_path_no_exception_to_gpu_drives.md — cold-path scalar arithmetic stays on GPU. Per feedback_no_partial_refactor.md — consumer (tau_update + state_reset + producer launch) all migrate together in this commit. Layout fingerprint already shifted by fc8dbb0a8 (slot + ISV_TOTAL_DIM bump); no additional fingerprint shift in this commit. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:31:36 +02:00
jgrusewski	fc8dbb0a85	feat(dqn): q_drift_rate_ema kernel + Rust wrapper (Plan C A.2 scaffolding) ... Single-block single-thread ISV producer mirrors h_s2_rms_ema / moe_lambda_eff pattern. Computes per-epoch drift_rate = |q_mean(t) - q_mean(t-1)| / max(|q_mean(t-1)|, ISV[Q_ABS_REF] + ISV[Q_DIR_ABS_REF], 1e-6) clipped to [0, 4] and writes to ISV[Q_DRIFT_RATE_INDEX=129]. Includes the full ISV-contract shift required for the kernel to load: - ISV_TOTAL_DIM 129 -> 130 - Q_DRIFT_RATE_INDEX = 129 (tail-appended after MOE_LAMBDA_EFF=128) - Cold-start ISV[129] = 0.0 in constructor (no-op dampening factor) - layout_fingerprint_seed entry Q_DRIFT_RATE=129 + ISV_TOTAL_DIM=130 - Cubin static, kernel field, kernel load, struct assignment Per feedback_no_partial_refactor.md the ISV slot + dim + fingerprint all migrate together (the wrapper references Q_DRIFT_RATE_INDEX so they cannot be split). Tau consumer + state reset registry + per-epoch producer launch land in the next commit. Audit doc dqn-wire-up-audit.md updated with the kernel + ISV slot description per Invariant 7. No callers in this commit; layout fingerprint shifts so existing checkpoints will fail-fast at load per feedback_no_legacy_aliases.md (expected for a real architecture change). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:29:20 +02:00
jgrusewski	dfbadf2f3f	test(dqn): Phase 2 Test 2.D — real-batch e2e on synthetic non-degenerate logits ... Plan C Phase 2 T9. Verifies the production kernel's eval-mode argmax exactly matches a Rust ground-truth E[Q] argmax on a 256-sample batch with realistic per-direction non-degenerate C51 distributions, and confirms Thompson explores Long+Short ≥ 40% in training mode. The plan-prescribed approach (reuse Phase 0 Test 0.F's converged checkpoint loader) was deferred — Test 0.F itself already exercises the safetensors load + branching forward path. Replacement strategy from the dispatch brief: synthetic batch via direct buffer write. Setup: - 256 samples, 21 atoms, per-(sample, direction, atom) C51 logits drawn from a deterministic hash → uniform [-1, 1] (post-Xavier-init scale of fresh C51 head outputs) - Per-sample adaptive support [v_min ∈ [-1.0, -0.2], v_max ∈ [0.2, 1.0]] matching the layout produced by `update_per_sample_support` - Uniform q_values (mag/ord/urg fall through Boltzmann uniformly) Rust ground-truth: softmax(b_logits[i, d]) · atom_vals[i, d] computed with the numerically-stable subtract-max softmax matching the kernel's softmax_c51_inline. Assertions: - Eval mode: per-sample dir_idx EXACTLY matches ground-truth argmax E[Q] - Train mode: count(d ∈ {Long, Short}) / batch > 0.40 Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:21:45 +02:00
jgrusewski	cc55c8a25c	test(dqn): Phase 2 Test 2.C — mag/ord/urg branches unaffected (parity vs Boltzmann ref) ... Plan C Phase 2 T8. The plan-prescribed pre-T2 snapshot approach was impossible (T2 had already landed); replacement strategy (ii) from the dispatch brief — behavioral parity vs an analytical Boltzmann reference computed in Rust — is used. Setup forces direction = Long (d=2) deterministically via peaked C51 logits (Long peaked at v=+0.8 atom; other directions at v=-0.5), so the kernel's Hold/Flat → mag_idx=0 short-circuit doesn't mask the magnitude branch's Boltzmann sampling. q_values are crafted with each branch (mag/ord/urg) peaked at a single bin with magnitude 1.0: Mag Q = [0.0, 0.0, 1.0] peak at Full (mag=2) Order Q = [1.0, 0.0, 0.0] peak at Market (ord=0) Urgency Q = [0.0, 1.0, 0.0] peak at urg=1 With q_range=1.0 in all three branches, tau collapses to 1.0 and the analytical Boltzmann probabilities are: P(best) = 1/(1 + 2/e) ≈ 0.5767 P(other) = 1/e/(1 + 2/e) ≈ 0.2117 Tolerance: at batch=8192 the 1-σ Bernoulli noise is ~0.0055 for p≈0.58; ±5% absolute tolerance covers ~9σ. Algorithmic divergence (e.g. an inadvertent strict-argmax substitution) would shift P(best) to 1.0 — trivially detected by the ±5% tolerance. Assertions: - dir_idx == Long for every sample (eval argmax E[Q] over peaked C51) - mag/ord/urg histograms each within ±5% of the Boltzmann reference Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:20:59 +02:00
jgrusewski	1a80fcab10	test(dqn): Phase 2 Test 2.B — production vs standalone (KS-fallback) ... Plan C Phase 2 T7. #[ignore]-gated GPU test that runs BOTH the production experience_action_select kernel AND the standalone direction_thompson_v2_test (added in commit 5de5e546a) on identical Phase 0 Test 0.D failure-mode inputs and asserts the resulting direction histograms are statistically indistinguishable. Path (a) — bit-identical comparison via a shared pre-computed uniform array — would require editing the standalone v2 kernel's signature to accept a `float* uniforms` instead of generating its own LCG draws. Out of scope for the 1-2-hour dispatch. Path (b) — KS-style histogram comparison — is used. Setup: - Single per-direction logits tile (Phase 0 Test 0.D failure mode) - Production: tile replicated across batch=100 000 samples; Philox keyed on (i, timestep=0, ctr) per thread - Standalone v2: same tile fed to single-tile launcher with n_seeds=100 000; LCG keyed on (base_seed=0 + seed_idx) per thread Assertion: - KS distance over the {Short, Hold, Long, Flat} histograms ≤ 0.02 (n=100k empirical-CDF noise floor for matching distributions is ~4·sqrt(2/n) ≈ 0.018; algorithm divergences would shift mass by O(10%) → KS ≈ O(0.1), easily detected) - Sanity: both histograms have P(Long+Short) ≥ 0.20 (rules out coincidental shared-mode collapse with KS ≈ 0) Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>	2026-04-29 18:17:26 +02:00