gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(sp4): migrate winsorized adaptive grad-clip update to GPU per feedback_no_cpu_compute_strict

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 15:09:57 +02:00
parent ca63158606
commit 1112abc2a4
4 changed files with 397 additions and 113 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
crates/ml/build.rs
View File

@@ -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`

351
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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::<f32>(), 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::<f32>(), 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

104
crates/ml/src/cuda_pipeline/update_adaptive_clip_kernel.cu Normal file
View File

@@ -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. */
}

44
docs/dqn-wire-up-audit.md
View File

@@ -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).