From 1112abc2a44f097fba714149e7aab7216e9aca42 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 15:09:57 +0200 Subject: [PATCH] fix(sp4): migrate winsorized adaptive grad-clip update to GPU per feedback_no_cpu_compute_strict MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Layer C close-out C4 — the most architecturally substantive site of the feedback_no_cpu_compute_strict sweep. GpuDqnTrainer::update_adaptive_clip was running a 6-step host-side compute chain on GPU-produced inputs: winsor (1) + cold-start sentinel + EMA (3) + scalar reduction (4) + ISV upper-bound clamp (5) + mapped-pinned write (6). Per feedback_no_partial_refactor the chain must migrate coherently — splitting EMA-only into a kernel and leaving the surrounding scalar reductions on host would be a partial migration violating the rule. Migration: new update_adaptive_clip_kernel.cu (single-thread, single-block). Takes host-passed observed_grad_norm (already a mapped-pinned readback) + 6 fixed structural constants (legacy values preserved per feedback_no_quickfixes) + 4 mapped-pinned dev_ptrs + ISV[GRAD_CLIP_BOUND_INDEX]. Writes the full output chain (adaptive_clip_pinned, grad_norm_ema_pinned, outlier_diag_pinned). Storage migration: - grad_norm_ema migrated from host-resident f32 to mapped-pinned scalar (matches C1/C2/C3 pattern). - New outlier_diag_pinned mapped-pinned slot for the GRAD_CLIP_OUTLIER warn diagnostic. The kernel writes `delta = observed - clamped` and the host reads it post-launch to format the warn log without running scalar arithmetic on the host. All structural constants preserved: EMA_BETA=0.95, CLIP_MULTIPLIER=2.0, MIN_CLIP=1.0, GRAD_CLIP_OUTLIER_K=100, EPS_CLAMP_FLOOR (SP4). Same cold-start sentinel `prev_ema <= 0.0 ⇒ assign clamped directly`; same Mech 6 (SP3) + Layer B (SP4) bound design. Same outlier-warn log format reconstructed from the mapped-pinned diagnostic slot. Host-side early-return guard preserved (the pre-existing pattern from C1 redesigned). grad_norm_emas_step_count counter unchanged (scalar control-flow metadata, not compute, per the rule's explicit carve-out). Verification: SP4 lib tests + 16 SP4 GPU producer unit tests pass on RTX 3050 Ti. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 11 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 351 ++++++++++++------ .../update_adaptive_clip_kernel.cu | 104 ++++++ docs/dqn-wire-up-audit.md | 44 +++ 4 files changed, 397 insertions(+), 113 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index fd0ff1299..c7ef314cc 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -261,6 +261,17 @@ fn main() { // lockstep. Cold-start sentinel `prev > 0.99 ⇒ assign obs directly` // (constructor init=1.0 marker) preserves the deleted host formula. "update_utilization_ema_kernel.cu", + // SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized + // adaptive grad-clip update. Single-thread, single-block — replaces + // the host-side scalar-reduction + EMA arithmetic chain in + // `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. The + // chain is winsor (1) + cold-start sentinel + EMA (3) + scalar + // reduction (4) + ISV upper-bound clamp (5) + mapped-pinned write + // (6); all six steps now run in a single kernel that reads + // host-passed `observed_grad_norm` + the three mapped-pinned + // scalars (`adaptive_clip_pinned`, `grad_norm_ema_pinned`, + // `outlier_diag_pinned`) + `isv[GRAD_CLIP_BOUND_INDEX]`. + "update_adaptive_clip_kernel.cu", // HEALTH_DIAG GPU port — Phase 2A landing (2026-04-28). Single-block // single-thread `health_diag_isv_mirror` kernel copies a curated set // of ISV signal-bus slots into the mapped-pinned `HealthDiagSnapshot` diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index c390dc506..2b0f648e1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -357,6 +357,15 @@ static UPDATE_IQN_READINESS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR" /// preserves the deleted host-side bootstrap branch (constructor init=1.0). static UPDATE_UTILIZATION_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_utilization_ema_kernel.cubin")); +/// SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized adaptive +/// grad-clip update kernel. Single-thread, single-block — replaces the +/// host-side scalar-reduction + EMA chain in `update_adaptive_clip` per +/// `feedback_no_cpu_compute_strict`. Migrates the ENTIRE chain (winsor + +/// cold-start sentinel + EMA + 2 clamps + ISV upper bound + mapped-pinned +/// write) per `feedback_no_partial_refactor` rather than splitting EMA-only +/// into the kernel and leaving the surrounding host scalar reductions. +static UPDATE_ADAPTIVE_CLIP_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_adaptive_clip_kernel.cubin")); + /// Plan 4 Task 1B-ii: per-group VSN MLP hidden dim. `Linear_1[g]` is /// `[VSN_HIDDEN_DIM, group_dim_g]`, `Linear_2[g]` is `[1, VSN_HIDDEN_DIM]`. /// 16 chosen as a tractable expansion that fits comfortably against the @@ -2949,7 +2958,32 @@ pub struct GpuDqnTrainer { /// Device pointer for adaptive_clip_pinned (from cuMemHostGetDevicePointer). adaptive_clip_dev_ptr: u64, /// EMA of gradient L2 norm for adaptive clipping. - grad_norm_ema: f32, + /// SP4 Layer C close-out C4 (2026-05-01) replaced the host-resident `f32` + /// field with `grad_norm_ema_pinned` (mapped-pinned scalar updated GPU- + /// side by `update_adaptive_clip_kernel`) per `feedback_no_cpu_compute_strict`. + /// Cold-start sentinel `prev <= 0.0 ⇒ assign clamped observation directly` + /// preserves the deleted host formula's bootstrap branch. Stores the + /// WINSORIZED grad-norm EMA (vs. fast/slow EMAs in C1 redesigned which + /// store the RAW grad-norm) — mechanism intentional per the comment + /// block at update_adaptive_clip line 22652-22657. + grad_norm_ema_pinned: *mut f32, + /// Device-mapped view of `grad_norm_ema_pinned`. Read+written in-place by + /// `update_adaptive_clip_kernel`; not consumed by other GPU kernels via + /// dev_ptr (the EMA itself is internal state; the kernel produces + /// `adaptive_clip_pinned` as the actual consumed output). + grad_norm_ema_dev_ptr: u64, + /// SP4 Layer C close-out C4 (2026-05-01): mapped-pinned diagnostic slot + /// for the GRAD_CLIP_OUTLIER warn log. `update_adaptive_clip_kernel` + /// writes `(observed - clamped)` into this slot when winsorization + /// triggers (delta > 0); the host reads it post-launch to decide + /// whether to emit `tracing::warn!` and recovers the original observed + /// value as `clamped + delta`. Mapped-pinned + `__threadfence_system()` + /// guarantees the value is visible by the time the host accessor runs + /// after the kernel launch returns. + outlier_diag_pinned: *mut f32, + /// Device-mapped view of `outlier_diag_pinned`. Written by + /// `update_adaptive_clip_kernel`; read host-side after launch. + outlier_diag_dev_ptr: u64, /// EMA of Q-divergence for adaptive tau computation. /// NOTE: only consumed by the orphan helper `compute_adaptive_tau` (zero /// production callers as of 2026-05-01); per `feedback_wire_everything_up.md` @@ -3936,6 +3970,14 @@ pub struct GpuDqnTrainer { /// `utilization_ema_pinned` + `homeostatic_obs_pinned[1]` mirror in /// lockstep. Loaded from `update_utilization_ema_kernel.cubin`. update_utilization_ema_kernel: CudaFunction, + /// SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized adaptive + /// grad-clip update kernel. Single-thread, single-block. Replaces the + /// host-side scalar-reduction + EMA chain in `update_adaptive_clip` + /// per `feedback_no_cpu_compute_strict`. Reads host-passed + /// `observed_grad_norm` + four mapped-pinned scalars + ISV[168]; writes + /// `adaptive_clip_pinned`, `grad_norm_ema_pinned`, `outlier_diag_pinned`. + /// Loaded from `update_adaptive_clip_kernel.cubin`. + update_adaptive_clip_kernel: CudaFunction, /// Mapped-pinned host pointer for the fast (α=0.1, ~10-step horizon) /// grad-norm EMA. Updated GPU-side by `update_grad_norm_emas_kernel` /// (chained on the producer's stream after `launch_grad_norm_finalize`) @@ -4789,6 +4831,12 @@ impl Drop for GpuDqnTrainer { if !self.utilization_ema_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.utilization_ema_pinned.cast()) }; } + if !self.grad_norm_ema_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.grad_norm_ema_pinned.cast()) }; + } + if !self.outlier_diag_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.outlier_diag_pinned.cast()) }; + } if !self.td_error_scratch_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.td_error_scratch_pinned.cast()) }; } @@ -11648,6 +11696,37 @@ impl GpuDqnTrainer { cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, utilization_ema_pinned.cast(), 0); dp }; + // SP4 Layer C close-out C4 (2026-05-01): mapped-pinned slots for + // `grad_norm_ema` (winsorized grad-norm EMA driving adaptive_clip) + // and `outlier_diag` (winsor-trigger diagnostic). Replaces the + // prior host-resident `grad_norm_ema: f32` field + the host-only + // `tracing::warn!` decision per `feedback_no_cpu_compute_strict`. + // grad_norm_ema_pinned init=0.0 — the cold-start sentinel that the + // GPU kernel's `prev <= 0.0 ⇒ assign clamped` branch checks. + let grad_norm_ema_pinned: *mut f32 = unsafe { + let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; + cudarc::driver::result::malloc_host(std::mem::size_of::(), flags) + .map_err(|e| MLError::ModelError(format!("pinned grad_norm_ema alloc: {e}")))? + as *mut f32 + }; + unsafe { *grad_norm_ema_pinned = 0.0; } + let grad_norm_ema_dev_ptr = unsafe { + let mut dp = 0u64; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, grad_norm_ema_pinned.cast(), 0); + dp + }; + let outlier_diag_pinned: *mut f32 = unsafe { + let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; + cudarc::driver::result::malloc_host(std::mem::size_of::(), flags) + .map_err(|e| MLError::ModelError(format!("pinned outlier_diag alloc: {e}")))? + as *mut f32 + }; + unsafe { *outlier_diag_pinned = 0.0; } + let outlier_diag_dev_ptr = unsafe { + let mut dp = 0u64; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, outlier_diag_pinned.cast(), 0); + dp + }; let total_actions = config.branch_0_size + config.branch_1_size + config.branch_2_size + config.branch_3_size; let q_out_buf = alloc_f32(&stream, b * total_actions, "q_out")?; let eval_td_snapshot = alloc_f32(&stream, b, "eval_td_snapshot")?; @@ -12163,6 +12242,20 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("update_utilization_ema_kernel load: {e}")))? }; + // SP4 Layer C close-out C4 (2026-05-01): load update_adaptive_clip + // kernel. Single-thread, single-block — replaces the host-side + // scalar-reduction + EMA chain in `update_adaptive_clip` per + // `feedback_no_cpu_compute_strict`. Updates the three coupled + // mapped-pinned scalars (`adaptive_clip_pinned`, + // `grad_norm_ema_pinned`, `outlier_diag_pinned`) plus reads + // ISV[GRAD_CLIP_BOUND_INDEX]. + let update_adaptive_clip_kernel = { + let module = stream.context().load_cubin(UPDATE_ADAPTIVE_CLIP_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("update_adaptive_clip cubin load: {e}")))?; + module.load_function("update_adaptive_clip_kernel") + .map_err(|e| MLError::ModelError(format!("update_adaptive_clip_kernel load: {e}")))? + }; + // Plan 4 Task 3 (E.3): load iqn_quantile_ema kernel (cold-path, per-step). // 4-block kernel — one block per off-median fixed quantile in // FIXED_TAUS = [0.05, 0.25, 0.50, 0.75, 0.95]. Producer for @@ -14921,7 +15014,10 @@ impl GpuDqnTrainer { aux_lr_dev_ptr, adaptive_clip_pinned, adaptive_clip_dev_ptr, - grad_norm_ema: 0.0, + grad_norm_ema_pinned, + grad_norm_ema_dev_ptr, + outlier_diag_pinned, + outlier_diag_dev_ptr, q_div_ema: 0.0, adam_step: 0, total_params, @@ -15176,6 +15272,7 @@ impl GpuDqnTrainer { update_grad_norm_emas_kernel, update_iqn_readiness_kernel, update_utilization_ema_kernel, + update_adaptive_clip_kernel, grad_norm_fast_ema_pinned, grad_norm_fast_ema_dev_ptr, grad_norm_slow_ema_pinned, @@ -22938,125 +23035,90 @@ impl GpuDqnTrainer { if !observed_grad_norm.is_finite() || observed_grad_norm <= 0.0 { return; } + + // SP4 Layer C close-out C4 (2026-05-01): the host-side scalar- + // reduction + EMA arithmetic chain (winsor + cold-start sentinel + + // EMA + 2 clamps + ISV upper-bound) now runs GPU-side via + // `update_adaptive_clip_kernel` per `feedback_no_cpu_compute_strict`. + // The full chain — winsor (1) + cold-start sentinel + EMA (3) + + // scalar reduction (4) + ISV upper-bound clamp (5) + mapped-pinned + // write (6) — was migrated coherently per + // `feedback_no_partial_refactor` rather than splitting the EMA into + // the kernel and leaving the surrounding scalar reductions on host. + // Per `pearl_cold_path_no_exception_to_gpu_drives.md`: the cold- + // path scalar arithmetic stays on GPU because all inputs (the + // winsor's `prev_clip`, the EMA's `prev_ema`, the upper-bound's ISV + // slot) already live in mapped-pinned host memory the device reads + // directly. + // + // Constants (legacy values, preserved exactly per + // `feedback_no_quickfixes.md`): + // EMA_BETA=0.95 — fixed-α (cross-fold steady-state + // behaviour the Mech 6 design depends + // on; Pearls A+D would defeat the + // winsor→EMA→clip chain's stability + // guarantees). + // CLIP_MULTIPLIER=2.0 — clip threshold = 2 × steady-state + // grad-norm (legacy value; same + // "headroom over the EMA" the SP4 + // Layer B Mech 6 anchor uses). + // MIN_CLIP=1.0 — ε-floor preventing cold-start + // `EMA × multiplier ≈ 0` from pinning + // the clip threshold at 0 (the actual + // clip's `min(grad, threshold)` would + // zero every gradient). + // GRAD_CLIP_OUTLIER_K=100 — winsor multiplier (Plan C Phase 2 + // follow-up N rationale at the deleted + // host code's comment block). + // EPS_CLAMP_FLOOR — SP4 ε-floor for the ISV upper-bound + // slot (cold-start ISV reads as 0 + // must not collapse the bound). + // + // Outlier-warn diagnostic: the kernel writes + // `outlier_diag_pinned[0] = observed - clamped` so this host-side + // post-launch readback can replicate the original + // `tracing::warn!("GRAD_CLIP_OUTLIER: clamping ...")` log without + // running scalar arithmetic on the host. The mapped-pinned slot + // is `__threadfence_system()`-visible by the time the launcher + // returns. + use crate::cuda_pipeline::sp4_isv_slots::GRAD_CLIP_BOUND_INDEX; + use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR; + + // Snapshot the prev_clip BEFORE the kernel overwrites + // `adaptive_clip_pinned[0]` — needed only for the outlier warn log + // (host log formatting is data-routing, not compute on GPU values). + let prev_clip_for_log = unsafe { *self.adaptive_clip_pinned }; + const EMA_BETA: f32 = 0.95; const CLIP_MULTIPLIER: f32 = 2.0; const MIN_CLIP: f32 = 1.0; - - // Plan C Phase 2 follow-up N (2026-04-29): Winsorized EMA input. - // Single-batch outliers (e.g., F1 ep2's grad_norm = 8.4B reported by - // smoke-test-xw4c6) pollute the adaptive_clip EMA, making the next - // clip threshold huge → subsequent extreme grads pass unclipped → - // NaN propagation. Clamp the sample to GRAD_CLIP_OUTLIER_K × current - // adaptive_clip BEFORE the EMA update so a single sample can grow - // the EMA at most K× in one step. - // - // Per `feedback_isv_for_adaptive_bounds.md`: the bound is internal - // EMA-tracked state (the previous adaptive_clip itself), not a - // hardcoded multiplier. Per `feedback_adaptive_not_tuned.md`: - // K = 100 is a numerical-stability bound (Invariant 1 carve-out) - // with explicit "single-sample-can't-fully-corrupt-EMA" semantics, - // not a tuned constant. Steady-state behaviour: typical - // grad_norm/adaptive_clip ratio < 5 → min() is a no-op → invisible - // in normal training, only kicks in on extreme outliers. - // - // The clip itself still acts on the raw grad_norm during the kernel's - // clip operation — N only modifies what's accepted INTO the - // adaptive_clip EMA, not the clip's per-step grad scaling. - // - // The fast/slow grad-norm EMAs (driving the fold-warmup factor) - // intentionally remain un-winsorized — they're a stability signal - // that SHOULD respond to outliers (transient spike → factor drops - // → extra warmup damping on the next step, providing additional - // defense). Only the adaptive_clip EMA gets winsorized because IT - // is the threshold that drives clipping; corruption is the bug. const GRAD_CLIP_OUTLIER_K: f32 = 100.0; - let prev_clip = unsafe { *self.adaptive_clip_pinned }; - let prev_clip_safe = prev_clip.max(1e-3); // numerical-stability floor - let clamped_grad_norm = observed_grad_norm.min(GRAD_CLIP_OUTLIER_K * prev_clip_safe); - if clamped_grad_norm < observed_grad_norm { + + if let Err(e) = self.launch_update_adaptive_clip( + observed_grad_norm, + EMA_BETA, + CLIP_MULTIPLIER, + MIN_CLIP, + GRAD_CLIP_OUTLIER_K, + EPS_CLAMP_FLOOR, + GRAD_CLIP_BOUND_INDEX as i32, + ) { + tracing::warn!("update_adaptive_clip launch failed: {e}"); + return; + } + + // Post-launch outlier-warn — read the diagnostic slot the kernel + // populated. Mapped-pinned + `__threadfence_system()` ⇒ the value + // is visible without any explicit host sync. + let delta = unsafe { *self.outlier_diag_pinned }; + if delta > 0.0 { + let clamped = observed_grad_norm - delta; tracing::warn!( "GRAD_CLIP_OUTLIER: clamping grad_norm {:.4e} → {:.4e} (prev_clip={:.4e}, K={})", - observed_grad_norm, clamped_grad_norm, prev_clip, GRAD_CLIP_OUTLIER_K + observed_grad_norm, clamped, prev_clip_for_log, GRAD_CLIP_OUTLIER_K ); } - if self.grad_norm_ema <= 0.0 { - self.grad_norm_ema = clamped_grad_norm; - } else { - self.grad_norm_ema = EMA_BETA * self.grad_norm_ema + (1.0 - EMA_BETA) * clamped_grad_norm; - } - let new_clip = (self.grad_norm_ema * CLIP_MULTIPLIER).max(MIN_CLIP); - - // SP3 Mech 6 (anchored): anchored upper bound on `new_clip`. The - // winsorizer above caps a SINGLE sample at K × prev_clip but does - // NOT prevent CONSECUTIVE elevated samples from compounding the - // EMA upward without bound. Over hundreds of steps the clip - // threshold ratchets to thousands while actual grad_norm tracks - // it from below → clipping becomes a no-op against - // in-distribution drift → Adam m/v EMAs poisoned → saturate. - // Original F1 NaN root cause from smoke-test-5rqzs (commit - // b9edccfc1) at step 3060 (Mech 5 diag slots 36–38, 40–42 - // firing). - // - // Bound: 100 × grad_norm_slow_ema × ISV[Q_ABS_REF=16].max(1.0) - // - `grad_norm_slow_ema` (existing α=0.001 slow EMA, updated below - // in this same function) anchors against the legitimate - // steady-state grad norm; here we read the PREVIOUS step's slow - // EMA from pinned memory (it gets overwritten further down). - // The slow EMA INTENTIONALLY persists across fold boundaries — - // anchoring at F0's grad scale provides tighter clipping during - // F1 ramp-up. Resetting the anchor (Mech 8 experiment, reverted - // 2026-04-30 per smoke-test-rxhjh F1 NaN @ 2300) loosened the - // bound during F1 transients and accelerated Adam saturation — - // worse than the 3720-step ceiling that Mech 6 alone produced. - // Cross-fold Adam saturation is now addressed by Mech 9 (post- - // Adam weight clamp), not by anchor manipulation. - // - 100× multiplier: principled per-step headroom over the - // steady-state grad norm. Earlier v2 of this bound used 5× — - // that over-clipped F0 ramp-up gradients while still letting - // F1 saturate. The 100× value provides legitimate 10× headroom - // over the slot-36 diagnostic threshold of 100 × isv. - // - ISV[Q_ABS_REF=16] multiplier: scales the cap with the - // Q-magnitude regime per `feedback_isv_for_adaptive_bounds.md`. - // Same ISV slot used by SP3 Mechs 1+2+5+9 — no new slot. - // - `.max(1.0_f32)` ε on the multiplier per the SP1 ε-floor pearl: - // cold-start ISV[16] ≈ 0 must not collapse the bound. - // - `.max(MIN_CLIP)` ε-floor on the bound itself: cold-start - // `grad_norm_slow_ema` ≈ 0 (true on F0 step 1) must not pin - // upper_bound at 0; MIN_CLIP=1.0 floor stays active for the - // brief F0 cold-start transient. - // - // F1+ smoke history (multiplier search converged on 100×): - // - smoke-test-fxvkk (Mech 6 only, 100×): F0=44, F1 NaN @ 3720 - // - smoke-test-ftdjz (Mech 6 + Mech 7, 100×): F0=38, F1 collapse - // @ 2040 — Mech 7 reverted (per-element clip misdiagnosis) - // - smoke-test-d25vq (Mech 6 v2, 5×): F0=21, F1 collapse @ 2820 - // - smoke-test-rxhjh (Mech 6 100× + Mech 8 anchor reset): F1 NaN - // @ 2300 (worse than Mech 6 alone) — Mech 8 reverted; the - // persistent slow_ema across folds was unintentional protection. - // SP3 closed the cross-fold Adam pathology with Mech 9 (post-Adam - // weight clamp via host-side multiplier × Q_ABS_REF) rather than - // further tuning the clip-anchor side of the chain. Layer B - // generalised the bound to per-group `ISV[WEIGHT_BOUND[group]]` - // (Wiener-EMA over p99(|params|), no host multiplier). - // - // SP4 Layer B (Mech 6): adaptive clip upper bound = ISV[GRAD_CLIP_BOUND=168]. - // Producer `launch_sp4_grad_norm_p99` (Layer A Task A8) writes the - // Pearl D–smoothed p99(grad_norm) into ISV[168] directly; the - // hardcoded `100 × prev_slow_ema × Q_ABS_REF.max(1.0)` bound from - // SP3 is replaced with this signal-driven upper-bound. The - // EPS_CLAMP_FLOOR ε on the bound itself preserves cold-start - // semantics (pre-first-observation ISV reads as 0). Note: - // `grad_norm_slow_ema_pinned` is no longer read by this consumer - // — its retirement is deferred to Layer C per the audit doc. - use crate::cuda_pipeline::sp4_isv_slots::GRAD_CLIP_BOUND_INDEX; - use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR; - let upper_bound = self.read_isv_signal_at(GRAD_CLIP_BOUND_INDEX).max(EPS_CLAMP_FLOOR); - let new_clip = new_clip.min(upper_bound); - - // Write directly to pinned memory — GPU sees it on next kernel read - unsafe { *self.adaptive_clip_pinned = new_clip; } - // Plan C Phase 2 follow-up K (2026-04-29) + SP4 Layer C close-out // C1 redesigned (2026-05-01): feed the SAME observation into the // fast/slow grad-norm EMAs that drive the fold-boundary warmup @@ -23116,6 +23178,63 @@ impl GpuDqnTrainer { /// the deleted host-side formula's semantics exactly. FAST_ALPHA / /// SLOW_ALPHA constants mirror the architectural ~10-step / ~1000-step /// horizons that the existing fold-warmup-factor design depends on. + /// SP4 Layer C close-out C4 (2026-05-01): GPU-only winsorized adaptive + /// grad-clip update launcher. + /// + /// Replaces the host-side scalar-reduction + EMA arithmetic chain in + /// `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. The + /// kernel runs on the trainer's stream (all four mapped-pinned slots + /// are owned by the trainer), reads the host-passed `observed_grad_norm` + /// + four mapped-pinned scalars + ISV[GRAD_CLIP_BOUND_INDEX], and + /// writes the three output mapped-pinned scalars (adaptive_clip, + /// grad_norm_ema, outlier_diag) in lockstep with `__threadfence_system()`. + fn launch_update_adaptive_clip( + &self, + observed_grad_norm: f32, + ema_beta: f32, + clip_multiplier: f32, + min_clip: f32, + outlier_k: f32, + eps_clamp_floor: f32, + grad_clip_bound_idx: i32, + ) -> Result<(), MLError> { + let clip_dev = self.adaptive_clip_dev_ptr; + let ema_dev = self.grad_norm_ema_dev_ptr; + let outlier_dev = self.outlier_diag_dev_ptr; + let isv_ptr = self.isv_signals_dev_ptr; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + + // Safety: kernel signature + // (float, float, float, float, float, float, + // float*, float*, const float*, int, float*) + // matches the eleven args below; the three writeable dev_ptrs are + // mapped-pinned scalars and `isv_signals_dev_ptr` is the ISV bus + // pointer (read-only here). + unsafe { + self.stream + .launch_builder(&self.update_adaptive_clip_kernel) + .arg(&observed_grad_norm) + .arg(&ema_beta) + .arg(&clip_multiplier) + .arg(&min_clip) + .arg(&outlier_k) + .arg(&eps_clamp_floor) + .arg(&clip_dev) + .arg(&ema_dev) + .arg(&isv_ptr) + .arg(&grad_clip_bound_idx) + .arg(&outlier_dev) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("update_adaptive_clip: {e}")))?; + } + Ok(()) + } + fn launch_update_grad_norm_emas(&self) -> Result<(), MLError> { const FAST_ALPHA: f32 = 0.1; // ~10-step horizon const SLOW_ALPHA: f32 = 0.001; // ~1000-step horizon (cross-fold steady baseline) @@ -23370,7 +23489,13 @@ impl GpuDqnTrainer { pub fn grad_norm_slow_ema_value(&self) -> f32 { unsafe { *self.grad_norm_slow_ema_pinned } } pub fn adaptive_clip_value(&self) -> f32 { unsafe { *self.adaptive_clip_pinned } } - pub fn grad_norm_ema_value(&self) -> f32 { self.grad_norm_ema } + /// SP4 Layer C close-out C4 (2026-05-01): the winsorized grad-norm EMA + /// is now stored in `grad_norm_ema_pinned` (mapped-pinned). Accessor + /// reads via host_ptr after the kernel's `__threadfence_system()`. + pub fn grad_norm_ema_value(&self) -> f32 { + if self.grad_norm_ema_pinned.is_null() { 0.0 } + else { unsafe { *self.grad_norm_ema_pinned } } + } /// Plan C T11 follow-up — F1 ep2 explosion diagnostic (2026-04-29). /// Read-only accessor over the pinned C51 (total) loss. Same pinned diff --git a/crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu b/crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu new file mode 100644 index 000000000..17469c0b0 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu @@ -0,0 +1,104 @@ +/* update_adaptive_clip — GPU-only winsorized adaptive grad-clip update. + * + * SP4 Layer C close-out C4 (2026-05-01). Replaces the host-side scalar- + * reduction + EMA arithmetic chain in `GpuDqnTrainer::update_adaptive_clip` + * (gpu_dqn_trainer.rs ~line 22660-22744) per `feedback_no_cpu_compute_strict`. + * Host code previously ran: + * 1. winsor: `clamped = min(observed, K × max(prev_clip, 1e-3))` + * 2. cold-start sentinel: `if ema ≤ 0.0 { ema = clamped }` + * 3. EMA: `ema = β × ema + (1 - β) × clamped` + * 4. scalar reduction: `new_clip = max(ema × CLIP_MULTIPLIER, MIN_CLIP)` + * 5. ISV upper bound: `new_clip = min(new_clip, max(ISV[GRAD_CLIP_BOUND_INDEX], EPS_CLAMP_FLOOR))` + * 6. mapped-pinned write: `adaptive_clip_pinned[0] = new_clip` + * — every step of which is compute on GPU-produced inputs (raw_grad_norm, + * prev_clip, ISV slot are all mapped-pinned scalars). The whole chain now + * runs in a single-thread kernel. + * + * Reads: observed_grad_norm — host-passed scalar (the caller has + * already verified `is_finite() && > 0.0` + * per the host-side early-return guard). + * clip_dev[0] — previous adaptive_clip (mapped-pinned). + * ema_dev[0] — previous winsorized grad-norm EMA + * (mapped-pinned scalar). + * isv[grad_clip_bound_idx] — ISV[GRAD_CLIP_BOUND_INDEX=168] upper + * bound from `launch_sp4_grad_norm_p99` + * (Pearl D-smoothed p99 of grad_norm). + * outlier_diag_dev[0] — diagnostic flag (mapped-pinned scalar): + * set to `observed - clamped` if winsor + * triggered, else 0.0. Host reads it + * post-launch to drive the GRAD_CLIP_OUTLIER + * warn log without cross-stream sync. + * Writes: ema_dev[0] — updated EMA. + * clip_dev[0] — updated adaptive_clip (final value). + * outlier_diag_dev[0] — winsor diagnostic delta. + * + * Cold-start sentinel: if `ema_dev[0] <= 0.0`, treat as bootstrap (cold- + * start for this run): assign `clamped` directly. Otherwise apply the + * fixed-α EMA recurrence (β=0.95, mirroring the deleted host constant). + * + * Per `feedback_no_cpu_compute_strict.md`: the entire chain is compute on + * mapped-pinned-resident inputs — belongs on GPU. Per `feedback_no_partial_refactor`: + * we migrate the COMPLETE chain (winsor + EMA + 2 clamps + ISV bound) in + * one kernel rather than a partial split that would leave host scalar + * reductions interleaved with the EMA arithmetic. + * + * Per `feedback_no_atomicadd.md`: a single thread reads four scalars and + * writes three — no atomics needed. + * + * The `__threadfence_system()` ensures the writes are PCIe-visible to the + * mapped-pinned host_ptr accessors (`grad_norm_ema_value()` HEALTH_DIAG + + * the host-side outlier-diagnostic warn check after this launch returns) + * and to other-stream kernel reads via dev_ptr (`adaptive_clip_pinned` is + * read by the Adam clip kernel during the next training step's replay). + */ + +extern "C" __global__ void update_adaptive_clip_kernel( + float observed_grad_norm, /* host-passed scalar */ + float ema_beta, /* fixed α=0.95 (legacy const) */ + float clip_multiplier, /* fixed 2.0 (legacy const) */ + float min_clip, /* fixed 1.0 (legacy const) */ + float outlier_k, /* fixed 100.0 (legacy const) */ + float eps_clamp_floor, /* SP4 EPS_CLAMP_FLOOR */ + float* __restrict__ clip_dev, /* [1] mapped-pinned adaptive_clip */ + float* __restrict__ ema_dev, /* [1] mapped-pinned grad_norm_ema */ + const float* __restrict__ isv, /* ISV bus */ + int grad_clip_bound_idx, /* GRAD_CLIP_BOUND_INDEX = 168 */ + float* __restrict__ outlier_diag_dev /* [1] winsor-delta diagnostic */ +) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; + + /* 1. Winsor: clamp observation against K × prev_clip_safe. The + * `prev_clip_safe = max(prev_clip, 1e-3)` ε floor guards against + * cold-start where prev_clip might be 0; matches deleted host code. */ + const float prev_clip = clip_dev[0]; + const float prev_clip_safe = fmaxf(prev_clip, 1e-3f); + const float upper_winsor = outlier_k * prev_clip_safe; + const float clamped = fminf(observed_grad_norm, upper_winsor); + + /* 2. Outlier diagnostic — host reads this slot post-launch to decide + * whether to emit the GRAD_CLIP_OUTLIER warn log. Encoded as the + * delta `observed - clamped` so the host can recover the original + * observed value (`clamped + delta`) for the log message. */ + outlier_diag_dev[0] = observed_grad_norm - clamped; + + /* 3. Cold-start sentinel + EMA. */ + const float prev_ema = ema_dev[0]; + float new_ema; + if (prev_ema <= 0.0f) { + new_ema = clamped; + } else { + new_ema = ema_beta * prev_ema + (1.0f - ema_beta) * clamped; + } + ema_dev[0] = new_ema; + + /* 4. Scalar reductions: clip = max(EMA × multiplier, MIN_CLIP), + * bounded by ISV[GRAD_CLIP_BOUND_INDEX].max(EPS_CLAMP_FLOOR). */ + const float new_clip_raw = fmaxf(new_ema * clip_multiplier, min_clip); + const float upper_bound = fmaxf(isv[grad_clip_bound_idx], eps_clamp_floor); + const float new_clip = fminf(new_clip_raw, upper_bound); + + clip_dev[0] = new_clip; + + __threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr + + * other-stream kernel reads via dev_ptr. */ +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index ecf45c418..215602837 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3046,3 +3046,47 @@ Discovered during the C3 sweep (deferred): - `gpu_experience_collector::set_utilization_ema` (gpu_experience_collector.rs:1886) — was originally listed as a 9th site but is just a setter `self.util_ema = util.clamp(0.0, 1.0)`, NOT EMA arithmetic (caller passes the already-EMA'd value). Misclassification in the original audit grid; no migration needed. Refs: SP4 Layer C close-out C3 — third site of the `feedback_no_cpu_compute_strict` sweep. + +### Layer C Task C4 — winsorized adaptive grad-clip update migrated to GPU (2026-05-01) + +Fourth site of the `feedback_no_cpu_compute_strict` sweep, the most complex by structure: `GpuDqnTrainer::update_adaptive_clip` was running a 6-step host-side compute chain on GPU-produced inputs: + +``` +1. winsor: clamped = min(observed, K × max(prev_clip, 1e-3)) +2. cold-start: if ema <= 0.0 { ema = clamped } +3. EMA: ema = β × ema + (1 - β) × clamped +4. scalar reduction: new_clip = max(ema × CLIP_MULTIPLIER, MIN_CLIP) +5. ISV upper bound: new_clip = min(new_clip, max(ISV[GRAD_CLIP_BOUND], EPS_CLAMP_FLOOR)) +6. mapped-pinned write: adaptive_clip_pinned[0] = new_clip +``` + +All inputs are mapped-pinned scalars (raw_grad_norm from the gpu_training_guard readback, prev_clip from `adaptive_clip_pinned`, ISV[168] from `launch_sp4_grad_norm_p99`). Per `feedback_no_cpu_compute_strict` the entire chain belongs on GPU; per `feedback_no_partial_refactor` it must migrate coherently rather than splitting EMA-only into a kernel and leaving the surrounding scalar reductions on host. + +Migrated: +- New `update_adaptive_clip_kernel.cu` — single-thread, single-block. Takes the host-passed `observed_grad_norm` (already a mapped-pinned readback) plus 6 fixed structural constants (legacy values preserved per `feedback_no_quickfixes`) and the 4 mapped-pinned dev_ptrs + ISV[GRAD_CLIP_BOUND_INDEX]; writes the full output chain (`adaptive_clip_pinned[0]`, `grad_norm_ema_pinned[0]`, `outlier_diag_pinned[0]`). +- Storage migration: `grad_norm_ema` migrated from host-resident `f32` field to mapped-pinned device-mapped scalar. New `outlier_diag_pinned` mapped-pinned slot for the GRAD_CLIP_OUTLIER warn diagnostic. Constructor allocates 2 new `cuMemHostAlloc(DEVICEMAP)` slots; Drop frees them. +- New `launch_update_adaptive_clip` Rust launcher chained on the trainer's stream — all four pinned slots are owned by the trainer. +- `update_adaptive_clip` host-side compute chain replaced with the launcher call + post-launch outlier-diag readback for the warn log (mapped-pinned + `__threadfence_system()` ⇒ value visible without explicit host sync). +- `grad_norm_ema_value()` accessor migrated to read through the pinned host_ptr. + +Preserved (no behaviour change): +- Same fixed-α formula `EMA_BETA=0.95` and structural constants (`CLIP_MULTIPLIER=2.0`, `MIN_CLIP=1.0`, `GRAD_CLIP_OUTLIER_K=100`, `EPS_CLAMP_FLOOR` from `sp4_wiener_ema`). +- Same cold-start sentinel `prev_ema <= 0.0 ⇒ assign clamped directly`. +- Same winsor formula and ISV upper-bound clamp semantics — Mech 6 (SP3) + Layer B (SP4) bound design unchanged. +- Same outlier-warn log message format; recovered host-side from the `delta = observed - clamped` mapped-pinned diagnostic slot the kernel writes. +- Host-side early-return guard `!observed_grad_norm.is_finite() || observed_grad_norm <= 0.0` — kernel only launches when value is valid (matches the pre-existing pattern that `update_grad_norm_emas_kernel` from C1 redesigned uses). +- `grad_norm_emas_step_count` host counter unchanged (scalar control-flow metadata, not compute, per the rule's explicit carve-out). +- Downstream consumer chain unchanged: `adaptive_clip_pinned` still consumed by Adam clip kernel via `adaptive_clip_dev_ptr`; the fast/slow grad-norm EMAs (driving fold-warmup factor) still updated by the C1 redesigned launcher. + +Verification: +- `cargo check -p ml --offline`: clean, same pre-existing warnings. +- `cargo test -p ml --lib --offline -- sp4 state_reset_registry`: 11 passed. +- `cargo test -p ml --test sp4_producer_unit_tests --offline --release -- --ignored`: 16/16 SP4 GPU tests pass on RTX 3050 Ti. + +Files touched (Layer C C4): +- `crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu` — new kernel. +- `crates/ml/build.rs` — 1 new kernel registration. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 1 new CUBIN static + 1 new struct field + 2 new mapped-pinned slot pairs (grad_norm_ema, outlier_diag) + 1 new constructor loader + 2 new pinned allocations + 2 new Drop entries + 1 new launcher (`launch_update_adaptive_clip`); host-side compute chain in `update_adaptive_clip` replaced with launcher call + post-launch warn readback; `grad_norm_ema_value()` accessor migrated to mapped-pinned read. +- `docs/dqn-wire-up-audit.md` — this entry. + +Refs: SP4 Layer C close-out C4 — fourth site of the `feedback_no_cpu_compute_strict` sweep. The most architecturally substantive close-out so far (full 6-step compute chain migration vs. C1/C2/C3 single-EMA migrations).