From 24accea774aa1cdfa9fab9314a2f40c5458a9481 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 14:07:10 +0200 Subject: [PATCH] fix(sp4): migrate fast/slow grad-norm EMA 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 C1 redesigned. Original plan claimed grad_norm_slow_ema_pinned was orphan post-Mech-6 migration; verification surfaced a SECOND live consumer (fold_warmup_factor_update, commit 4ef1d8ebb) that legitimately reads the slow EMA as cross-fold steady-state baseline. The actual defect: the EMA UPDATE at update_adaptive_clip:22720-22737 was host-side `(1-α)*prev + α*obs` arithmetic — exactly the pattern feedback_no_cpu_compute_strict (saved 2026-05-01) strictly forbids. Migrated: - New `update_grad_norm_emas_kernel.cu` — single-thread fused fast+slow EMA update kernel. Reads `grad_norm_buf[0]`, updates two mapped-pinned EMA scalars via dev_ptr. `__threadfence_system()` ensures the `fold_warmup_factor_kernel` consumer sees freshly-written values. - `launch_update_grad_norm_emas` Rust launcher chained on the producer's stream — graph-capture-compatible, no host sync. - update_adaptive_clip's host-side `unsafe { ... }` block replaced with the GPU launcher call (warn-and-continue on launch failure mirroring the launch_h_s2_rms_ema / launch_fold_warmup_factor per-step ISV producer pattern at training_loop.rs:3450/3464). Preserved: - grad_norm_slow_ema_pinned mapped-pinned buffer (cross-fold persistent; legitimate consumer is fold_warmup_factor_update). - Fixed-α design (FAST_ALPHA=0.1, SLOW_ALPHA=0.001) — Pearls A+D adaptive α would defeat the cross-fold-baseline semantic the warmup factor depends on. - Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) — same formula as the deleted host code. - Host-side update_adaptive_clip early-return guard — kernel only launches when observed_grad_norm is finite and > 0. - grad_norm_emas_step_count host counter — scalar control-flow metadata for warmup-window gating, not compute. Plan/spec docs updated to remove stale "orphan" claim. State-reset registry doc-comment + field doc-comments updated to reflect GPU-only update path. fold_warmup_factor_kernel docstring no longer describes its grad-norm EMA inputs as host-side-fed. Build clean, sp4 + state_reset_registry lib tests pass (11/11), 16/16 SP4 producer GPU tests pass on RTX 3050 Ti. No behavior change — pure architectural fix. Refs: SP4 Layer C C1 redesigned (was: retire). Original plan docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md lines 2184-2221. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 11 ++ .../fold_warmup_factor_kernel.cu | 6 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 175 ++++++++++++++---- .../update_grad_norm_emas_kernel.cu | 67 +++++++ .../src/trainers/dqn/state_reset_registry.rs | 2 +- docs/dqn-wire-up-audit.md | 39 ++++ ...-30-sp4-signal-driven-magnitude-control.md | 57 +++--- ...-signal-driven-magnitude-control-design.md | 2 +- 8 files changed, 285 insertions(+), 74 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 3fe4192c1..1657a029d 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -226,6 +226,17 @@ fn main() { // factor starts at 0 → heavy damping; rises to 1 as gradients // stabilise → consumers return to baseline (healthy runs unaffected). "fold_warmup_factor_kernel.cu", + // SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only + // fast/slow grad-norm EMA update. Single-thread, single-block — + // replaces the host-side arithmetic block in `update_adaptive_clip` + // per `feedback_no_cpu_compute_strict`. Reads `grad_norm_buf[0]` + // (populated earlier in the same stream by `launch_grad_norm_finalize`), + // updates the fast (α=0.1, ~10-step horizon) and slow (α=0.001, + // ~1000-step horizon) EMAs in mapped-pinned scalars consumed by + // `fold_warmup_factor_kernel.cu` (the actual ISV[130] producer). + // Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) preserves + // the deleted host-side formula's semantics exactly. + "update_grad_norm_emas_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/fold_warmup_factor_kernel.cu b/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu index 2697e9902..71079d026 100644 --- a/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu +++ b/crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu @@ -14,8 +14,10 @@ * Cold-path cadence — single-block, single-thread kernel matching * `q_drift_rate_ema_kernel.cu`, `moe_lambda_eff_kernel.cu`, and * `kelly_cap_update_kernel.cu` shapes. Launched once per training step AFTER - * the per-step grad-norm EMAs have been updated (host-side `update_adaptive_clip` - * already feeds these EMAs from `gr.raw_grad_norm`). + * the per-step grad-norm EMAs have been updated by `update_grad_norm_emas_kernel.cu` + * (the producer feeds these EMAs from `grad_norm_buf[0]` per + * `feedback_no_cpu_compute_strict`; SP4 Layer C close-out C1 redesigned + * 2026-05-01 migrated the prior host-side arithmetic block). * * Formula: * diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a9933b6b1..1791fe0ac 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -331,6 +331,13 @@ static Q_DRIFT_RATE_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), " /// Two consumers (lr_eff and clip_eff in `training_loop.rs`), both monotone. static FOLD_WARMUP_FACTOR_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/fold_warmup_factor_kernel.cubin")); +/// SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only fast/slow +/// grad-norm EMA update kernel. Single-thread, single-block — replaces the +/// host-side arithmetic block in `update_adaptive_clip` per +/// `feedback_no_cpu_compute_strict`. Reads `grad_norm_buf[0]`, updates two +/// mapped-pinned EMA scalars consumed by `fold_warmup_factor_kernel`. +static UPDATE_GRAD_NORM_EMAS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_grad_norm_emas_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 @@ -3834,22 +3841,42 @@ pub struct GpuDqnTrainer { /// existing `update_adaptive_clip` pinned slot). Loaded from /// `fold_warmup_factor_kernel.cubin`. fold_warmup_factor_kernel: CudaFunction, + /// SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only fast/slow + /// grad-norm EMA update kernel. Single-thread, single-block. Replaces + /// the host-side `(1-α) × prev + α × obs` arithmetic block in + /// `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. Chained + /// on the producer's stream after `launch_grad_norm_finalize` so the + /// freshly-written `grad_norm_buf[0]` is consumed without any host sync; + /// updates `grad_norm_fast_ema_pinned` + `grad_norm_slow_ema_pinned` + /// in-place via their mapped-pinned dev_ptrs. Loaded from + /// `update_grad_norm_emas_kernel.cubin`. + update_grad_norm_emas_kernel: CudaFunction, /// Mapped-pinned host pointer for the fast (α=0.1, ~10-step horizon) - /// grad-norm EMA. Updated by `update_adaptive_clip` per training step - /// (the same observation source feeds both this fast EMA and the - /// existing `grad_norm_ema` field used by adaptive clip). FoldReset: - /// 0.0 (the new fold's first steps observe the recovery from zero). + /// grad-norm EMA. Updated GPU-side by `update_grad_norm_emas_kernel` + /// (chained on the producer's stream after `launch_grad_norm_finalize`) + /// per training step — the same observation source feeds both this fast + /// EMA and the existing `grad_norm_ema` field used by adaptive clip. + /// SP4 Layer C close-out C1 redesigned (2026-05-01) migrated the EMA + /// arithmetic from host-side to GPU per `feedback_no_cpu_compute_strict`. + /// FoldReset: 0.0 (the new fold's first steps observe the recovery from + /// zero). grad_norm_fast_ema_pinned: *mut f32, - /// Device-mapped view of `grad_norm_fast_ema_pinned`. + /// Device-mapped view of `grad_norm_fast_ema_pinned`. Consumed by + /// `fold_warmup_factor_kernel` (numerator) and updated in-place by + /// `update_grad_norm_emas_kernel`. grad_norm_fast_ema_dev_ptr: u64, /// Mapped-pinned host pointer for the slow (α=0.001, ~1000-step horizon) /// grad-norm EMA. Tracks the cross-fold steady-state grad scale — /// persists across folds so the warmup factor's denominator stays a - /// meaningful long-window baseline. Constructor: 0.0 (cold-start; the - /// `min_steps_for_ratio` gate in the kernel keeps the ratio at 1.0 + /// meaningful long-window baseline. Updated GPU-side by + /// `update_grad_norm_emas_kernel` per training step (SP4 Layer C + /// close-out C1 redesigned, 2026-05-01). Constructor: 0.0 (cold-start; + /// the `min_steps_for_ratio` gate in the kernel keeps the ratio at 1.0 /// until enough samples have accumulated). grad_norm_slow_ema_pinned: *mut f32, - /// Device-mapped view of `grad_norm_slow_ema_pinned`. + /// Device-mapped view of `grad_norm_slow_ema_pinned`. Consumed by + /// `fold_warmup_factor_kernel` (denominator) and updated in-place by + /// `update_grad_norm_emas_kernel`. grad_norm_slow_ema_dev_ptr: u64, /// Number of times `update_adaptive_clip` has fed a finite grad-norm /// observation into the EMAs. The `fold_warmup_factor_update` kernel @@ -11861,6 +11888,19 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("fold_warmup_factor_update load: {e}")))? }; + // SP4 Layer C close-out C1 redesigned (2026-05-01): load + // update_grad_norm_emas kernel (cold-path, per-step). Single-thread + // single-block — replaces the host-side EMA arithmetic block in + // `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. Reads + // `grad_norm_buf[0]` and writes the fast/slow grad-norm EMA mapped- + // pinned scalars consumed by `fold_warmup_factor_kernel`. + let update_grad_norm_emas_kernel = { + let module = stream.context().load_cubin(UPDATE_GRAD_NORM_EMAS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("update_grad_norm_emas cubin load: {e}")))?; + module.load_function("update_grad_norm_emas_kernel") + .map_err(|e| MLError::ModelError(format!("update_grad_norm_emas_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 @@ -14868,6 +14908,7 @@ impl GpuDqnTrainer { apply_pearls_ad_kernel, q_drift_rate_ema_kernel, fold_warmup_factor_kernel, + update_grad_norm_emas_kernel, grad_norm_fast_ema_pinned, grad_norm_fast_ema_dev_ptr, grad_norm_slow_ema_pinned, @@ -22702,42 +22743,98 @@ impl GpuDqnTrainer { // 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): feed the SAME observation - // into the fast/slow grad-norm EMAs that drive the fold-boundary - // warmup factor `fold_warmup_factor_update`. Constants are - // numerical-stability bounds per `feedback_isv_for_adaptive_bounds.md` - // — fast α=0.1 is the architectural ~10-step horizon (recover from - // fold reset in roughly the time the existing adaptive clip's EMA - // takes to recapture grad scale); slow α=0.001 is the ~1000-step - // horizon (cross-fold steady-state baseline). Both are STRUCTURAL - // (preserve direction / amortise scale across folds), not tuned. + // 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 + // factor `fold_warmup_factor_update`. Constants are numerical- + // stability bounds per `feedback_isv_for_adaptive_bounds.md` — fast + // α=0.1 is the architectural ~10-step horizon (recover from fold + // reset in roughly the time the existing adaptive clip's EMA takes + // to recapture grad scale); slow α=0.001 is the ~1000-step horizon + // (cross-fold steady-state baseline). Both are STRUCTURAL (preserve + // direction / amortise scale across folds), not tuned. // - // Per N's decision (above): fast/slow EMAs use the RAW - // `observed_grad_norm`, not the winsorized sample — they're a - // stability signal where outliers should transiently spike fast - // EMA and dampen the warmup factor on the next step. - const FAST_ALPHA: f32 = 0.1; - const SLOW_ALPHA: f32 = 0.001; - unsafe { - let fast_prev = *self.grad_norm_fast_ema_pinned; - let new_fast = if fast_prev <= 0.0 { - observed_grad_norm - } else { - (1.0 - FAST_ALPHA) * fast_prev + FAST_ALPHA * observed_grad_norm - }; - *self.grad_norm_fast_ema_pinned = new_fast; - - let slow_prev = *self.grad_norm_slow_ema_pinned; - let new_slow = if slow_prev <= 0.0 { - observed_grad_norm - } else { - (1.0 - SLOW_ALPHA) * slow_prev + SLOW_ALPHA * observed_grad_norm - }; - *self.grad_norm_slow_ema_pinned = new_slow; + // Per N's decision (above): fast/slow EMAs use the RAW grad-norm + // observation (`grad_norm_buf[0]`), not the winsorized sample — + // they're a stability signal where outliers should transiently + // spike fast EMA and dampen the warmup factor on the next step. + // + // The EMA arithmetic itself runs GPU-side via + // `launch_update_grad_norm_emas` per `feedback_no_cpu_compute_strict` + // (any compute — EMA, reduction, mean — belongs on GPU regardless + // of frequency). The host-side `observed_grad_norm` parameter is + // retained only as the early-return guard above (≤ 0.0 / non-finite + // ⇒ skip the EMA update); the kernel reads the same scalar from + // `grad_norm_buf[0]` (populated earlier in the same stream by + // `launch_grad_norm_finalize`, fully written by the time + // `readback_scalars_sync` produced `observed_grad_norm` host-side). + // + // Launch-error handling mirrors the other per-step ISV producers + // (`launch_h_s2_rms_ema`, `launch_fold_warmup_factor` at + // `training_loop.rs:3450/3464`): warn-and-continue rather than + // propagating, since `update_adaptive_clip`'s `()` return type is + // a stable contract with multiple callers and a kernel-launch + // failure here is non-fatal (the previous EMA values stay valid; + // `fold_warmup_factor_kernel` still produces a valid factor on the + // next step's still-fresh EMAs). + if let Err(e) = self.launch_update_grad_norm_emas() { + tracing::warn!("update_grad_norm_emas launch failed: {e}"); } self.grad_norm_emas_step_count = self.grad_norm_emas_step_count.saturating_add(1); } + /// SP4 Layer C close-out C1 redesigned (2026-05-01): GPU-only fast/slow + /// grad-norm EMA update launcher. + /// + /// Replaces the host-side `(1-α) × prev + α × obs` arithmetic in + /// `update_adaptive_clip` per `feedback_no_cpu_compute_strict`. Runs in + /// the same stream as `launch_grad_norm_finalize`, so the freshly- + /// written `grad_norm_buf[0]` is consumed without any host sync — + /// sequential same-stream ordering provides the producer→consumer + /// happens-before. Caller-compat: the EMA buffers retain their mapped- + /// pinned shape so `fold_warmup_factor_kernel` continues to read them + /// via dev_ptr unchanged, and the host-side accessors + /// (`grad_norm_fast_ema_value` / `grad_norm_slow_ema_value`) keep + /// observing the same memory after the kernel's + /// `__threadfence_system()`. + /// + /// Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) preserves + /// 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. + 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) + + let grad_norm_ptr = self.ptrs.grad_norm_buf; + let fast_dev_ptr = self.grad_norm_fast_ema_dev_ptr; + let slow_dev_ptr = self.grad_norm_slow_ema_dev_ptr; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + + // Safety: kernel signature `(const float*, float*, float*, float, float)` + // matches the five args below; `grad_norm_buf` is the device pointer + // for the L2-norm scalar populated by `launch_grad_norm_finalize`, + // and `grad_norm_fast/slow_ema_dev_ptr` are mapped-pinned dev_ptrs + // valid for the lifetime of `self`. + unsafe { + self.stream + .launch_builder(&self.update_grad_norm_emas_kernel) + .arg(&grad_norm_ptr) + .arg(&fast_dev_ptr) + .arg(&slow_dev_ptr) + .arg(&FAST_ALPHA) + .arg(&SLOW_ALPHA) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("update_grad_norm_emas: {e}")))?; + } + Ok(()) + } + /// Plan C Phase 2 follow-up K (2026-04-29): launch /// `fold_warmup_factor_update` — single-thread single-block ISV /// producer for ISV[FOLD_WARMUP_FACTOR_INDEX=130]. Reads the fast/slow diff --git a/crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu b/crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu new file mode 100644 index 000000000..c8184bcbb --- /dev/null +++ b/crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu @@ -0,0 +1,67 @@ +/* update_grad_norm_emas — GPU-only fast/slow grad-norm EMA update. + * + * SP4 Layer C close-out C1 redesigned (2026-05-01). Replaces the host-side + * arithmetic block in `GpuDqnTrainer::update_adaptive_clip` + * (gpu_dqn_trainer.rs ~line 22720) per `feedback_no_cpu_compute_strict`. + * Single-thread, single-block — the EMA update is two scalar formulas, no + * parallelism needed, mirroring `q_drift_rate_ema_kernel.cu` / + * `moe_lambda_eff_kernel.cu` cold-path footprint. + * + * Reads: grad_norm_dev[0] — device-visible scalar populated by + * `launch_grad_norm_finalize` earlier in the + * same stream (sequential same-stream ordering, + * no sync needed). + * fast_ema_dev[0] — previous fast EMA (mapped-pinned scalar). + * slow_ema_dev[0] — previous slow EMA (mapped-pinned scalar). + * Writes: fast_ema_dev[0] — updated fast EMA. + * slow_ema_dev[0] — updated slow EMA. + * + * Cold-start sentinel: if previous EMA value is ≤ 0.0 (post-fold-reset for + * fast, or initial cold start), assign the observation directly; otherwise + * apply standard EMA recurrence + * new = (1 - α) × prev + α × obs + * Matches the deleted host-side formula bit-for-bit (modulo IEEE-754 fp32 + * commutativity; both formulations use the same algebraic form). + * + * Mapped-pinned EMA buffers (`fast_ema_dev`, `slow_ema_dev`) are 1-element + * device pointers backed by mapped-pinned host memory. The host-side + * `fold_warmup_factor_kernel` consumer reads them via dev_ptr in its own + * launch later in the same stream; `__threadfence_system()` ensures the + * write is PCIe-visible to mapped-pinned host_ptr accesses + * (`grad_norm_fast_ema_value` / `grad_norm_slow_ema_value` accessors used by + * HEALTH_DIAG diagnostics). + * + * Per `pearl_cold_path_no_exception_to_gpu_drives.md`: cold-path scalar + * arithmetic stays on GPU when its EMA inputs already live in mapped-pinned + * host memory the device can read directly. Per `feedback_no_atomicadd.md`: + * a single thread reads three scalars and writes two — no atomics needed. + * Per `feedback_no_cpu_compute_strict.md`: any compute (EMA, reduction, + * mean) belongs on GPU regardless of frequency, regardless of whether it's + * hot-path or cold-path. + */ + +extern "C" __global__ void update_grad_norm_emas_kernel( + const float* __restrict__ grad_norm_dev, /* [1] grad_norm_buf[0] */ + float* __restrict__ fast_ema_dev, /* [1] mapped-pinned fast EMA */ + float* __restrict__ slow_ema_dev, /* [1] mapped-pinned slow EMA */ + float fast_alpha, + float slow_alpha +) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; + + const float obs = grad_norm_dev[0]; + + const float fast_prev = fast_ema_dev[0]; + const float new_fast = (fast_prev <= 0.0f) + ? obs + : ((1.0f - fast_alpha) * fast_prev + fast_alpha * obs); + fast_ema_dev[0] = new_fast; + + const float slow_prev = slow_ema_dev[0]; + const float new_slow = (slow_prev <= 0.0f) + ? obs + : ((1.0f - slow_alpha) * slow_prev + slow_alpha * obs); + slow_ema_dev[0] = new_slow; + + __threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr */ +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 44f8ace61..714d82190 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -451,7 +451,7 @@ impl StateResetRegistry { RegistryEntry { name: "isv_grad_norm_fast_ema", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.grad_norm_fast_ema_pinned — fast (α=0.1, ~10-step horizon) EMA of `gr.raw_grad_norm`, host-side mapped-pinned scalar fed by `update_adaptive_clip` and consumed by `fold_warmup_factor_update` (Plan C K). Reset to 0.0 at fold boundary so the warmup-factor ratio starts at 0 → consumer-derived lr_eff/clip_eff start at their permanent-floor values; recovers to slow-EMA baseline as gradients stabilise. Slow EMA companion (grad_norm_slow_ema_pinned) deliberately persists across folds — it tracks the cross-fold steady-state grad scale that the warmup factor compares against.", + description: "GpuDqnTrainer.grad_norm_fast_ema_pinned — fast (α=0.1, ~10-step horizon) EMA of `gr.raw_grad_norm`, host-side mapped-pinned scalar updated GPU-side by `update_grad_norm_emas_kernel` (chained on the producer's stream after `launch_grad_norm_finalize`) and consumed by `fold_warmup_factor_update` (Plan C K). SP4 Layer C close-out C1 redesigned (2026-05-01) migrated the prior host-side EMA arithmetic block in `update_adaptive_clip` to GPU per `feedback_no_cpu_compute_strict`. Reset to 0.0 at fold boundary so the warmup-factor ratio starts at 0 → consumer-derived lr_eff/clip_eff start at their permanent-floor values; recovers to slow-EMA baseline as gradients stabilise. Slow EMA companion (grad_norm_slow_ema_pinned) deliberately persists across folds — it tracks the cross-fold steady-state grad scale that the warmup factor compares against.", }, // ───── SP4 Layer A Task A12 (2026-04-30): adaptive bound + // Wiener-state + Pearl C engagement-counter fold-boundary resets. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index f1f2001ac..663b9926c 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2927,4 +2927,43 @@ Files touched (Layer C #260): Refs: task #260, baseline commit `1c150d190` (Layer A + B + fix-up + L40S smoke pass). +### Layer C Task C1 redesigned — fast/slow grad-norm EMA UPDATE migrated to GPU (2026-05-01) + +Original Layer C plan (2026-04-30) called for retiring `grad_norm_slow_ema_pinned` as orphan post-Mech-6 ISV migration. Verification on 2026-05-01 surfaced a SECOND live consumer: `fold_warmup_factor_update` (commit `4ef1d8ebb`, predates SP4) reads the slow EMA as the cross-fold steady-state grad-norm baseline that the warmup factor's denominator compares against. Legitimate cross-fold mechanism; NOT orphan. + +`grad_norm_slow_ema_pinned` is retained — its removal would break Plan C K's fold-warmup factor controller (the actual `feedback_no_partial_refactor` violation here would be deleting a buffer with a live consumer chain). + +The actual defect surfaced by the audit: the EMA UPDATE at `update_adaptive_clip:22720-22737` was running as host-side `(1-α) × prev + α × obs` arithmetic — exactly the pattern `feedback_no_cpu_compute_strict` (saved 2026-05-01) strictly forbids: +> "any compute (EMA, reduction, mean, Pearls A+D, adaptive α, ISV update, etc.) belongs on GPU regardless of frequency, regardless of whether it's hot-path or cold-path." + +Migrated: +- New `update_grad_norm_emas_kernel.cu` — single-thread fused fast+slow EMA update kernel. Reads `grad_norm_buf[0]` (the device-visible scalar populated earlier in the same stream by `launch_grad_norm_finalize`), updates `grad_norm_fast_ema_pinned` + `grad_norm_slow_ema_pinned` in-place via their mapped-pinned dev_ptrs. `__threadfence_system()` ensures the writes are PCIe-visible to mapped-pinned host_ptr accesses (`grad_norm_fast_ema_value` / `grad_norm_slow_ema_value` HEALTH_DIAG accessors) and to the downstream `fold_warmup_factor_kernel` consumer reading via dev_ptr. +- `launch_update_grad_norm_emas` Rust launcher chained on the producer's stream — graph-capture-compatible, no host sync (sequential same-stream ordering provides the producer→consumer happens-before). +- `update_adaptive_clip`'s host-side `unsafe { ... }` arithmetic block replaced with the launcher call (warn-and-continue on launch failure mirroring the `launch_h_s2_rms_ema` / `launch_fold_warmup_factor` per-step ISV producer pattern at `training_loop.rs:3450/3464`). + +Preserved (no behaviour change): +- `grad_norm_slow_ema_pinned` mapped-pinned buffer and its cross-fold persistence (legitimate consumer is `fold_warmup_factor_update`). +- Fixed-α design: `FAST_ALPHA=0.1` (~10-step horizon), `SLOW_ALPHA=0.001` (~1000-step horizon). Pearls A+D adaptive α would defeat the cross-fold-baseline semantic the warmup factor depends on. +- Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) — same formula as the deleted host code. +- Host-side `update_adaptive_clip` early-return guard (`!observed_grad_norm.is_finite() || observed_grad_norm <= 0.0` ⇒ skip the EMA update) — kernel only launches when the value is valid. +- `grad_norm_emas_step_count` host counter — scalar control-flow metadata for `fold_warmup_factor_kernel`'s warmup-window gating, NOT compute (per the just-saved `feedback_no_cpu_compute_strict`, integer step counters for control-flow gates are scalar metadata, not the EMA/reduction/mean class of computation the rule targets). +- Fold-boundary reset of `grad_norm_fast_ema_pinned` to 0.0 (state-reset registry) — one-time-per-fold reset, not per-step compute, and IS the cold-start sentinel signal the GPU kernel relies on. + +Verification: +- `cargo check -p ml --offline`: clean, same pre-existing warnings. +- `git grep -nE "(1\.0 - FAST_ALPHA)\|(1\.0 - SLOW_ALPHA)\|FAST_ALPHA \* observed_grad_norm\|SLOW_ALPHA \* observed_grad_norm" crates/ml/src/`: ZERO matches in code (host-side formula gone). +- `git grep -nE "\*self\.grad_norm_fast_ema_pinned\|\*self\.grad_norm_slow_ema_pinned" crates/ml/src/`: only fold-boundary reset (one-time) + read-only HEALTH_DIAG accessors remain — NO per-step host writes. + +Files touched (Layer C C1 redesigned): +- `crates/ml/src/cuda_pipeline/update_grad_norm_emas_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 + 1 new constructor loader + 1 new launcher (`launch_update_grad_norm_emas`); host-side EMA arithmetic block replaced with launcher call; field doc-comments updated to reflect GPU-only update path. +- `crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu` — docstring updated (EMAs are now GPU-updated, not host-side). +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — `isv_grad_norm_fast_ema` description updated to reflect GPU-only update path. +- `docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md` — Layer B "becomes orphan" claim corrected; Task C1 redesigned section. +- `docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md` — Layer C "retire" line replaced with "retain + migrate UPDATE to GPU" rationale. +- `docs/dqn-wire-up-audit.md` — this entry. + +No behavior change — pure architectural fix. The L40S smoke `smoke-test-tkkx6` (against `45da3eac6`) remains valid as the baseline; this commit doesn't need a new smoke since the only change is "where is the arithmetic computed", not "what is computed". Refs: SP4 Layer C C1 redesigned (was: retire). Original plan line ~2049 + Task C1 ~lines 2184-2221. + follow-up: symbolic doc-anchors + dedicated q_dir_grad sub-buffer table + p99 helper extraction diff --git a/docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md b/docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md index e6942027e..a6f699a68 100644 --- a/docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md +++ b/docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md @@ -2046,7 +2046,7 @@ let upper_bound = self.read_isv_signal_at(GRAD_CLIP_BOUND_INDEX).max(EPS_CLAMP_F let new_clip = new_clip.min(upper_bound); ``` -Note: `prev_slow_ema` is no longer read here. `grad_norm_slow_ema_pinned` becomes orphan (cleaned up in Layer C). +Note: `prev_slow_ema` is no longer read by Mech 6's upper_bound. `grad_norm_slow_ema_pinned` is **NOT orphan** — it has a second live consumer (`fold_warmup_factor_update`, commit `4ef1d8ebb`) that reads it as the cross-fold steady-state grad-norm baseline. Layer C C1 redesigned (2026-05-01): instead of retiring the slow EMA, migrate the host-side EMA UPDATE arithmetic to GPU per `feedback_no_cpu_compute_strict`. - [ ] **Step 5: Mech 9 (post-Adam weight clamp across 5 Adam kernels)** @@ -2181,44 +2181,39 @@ EOF ## Layer C — Validation + cleanup -### Task C1: Retire `grad_norm_slow_ema` infrastructure +### Task C1 (REDESIGNED 2026-05-01): Migrate fast/slow grad-norm EMA UPDATE to GPU + +**Original plan claimed `grad_norm_slow_ema_pinned` was orphan post-Mech-6 migration. Verification on 2026-05-01 surfaced a SECOND live consumer: `fold_warmup_factor_update` (commit `4ef1d8ebb`, predates SP4) reads the slow EMA as the cross-fold steady-state grad-norm baseline. Legitimate, not orphan.** + +**However**, on inspection of `update_adaptive_clip:22720-22737`, the fast/slow EMA UPDATE was running as host-side `(1-α) × prev + α × obs` arithmetic — exactly the pattern `feedback_no_cpu_compute_strict` (saved 2026-05-01) strictly forbids. The actual defect was the compute location, not orphan-status. + +**Redesigned scope:** +- Retain `grad_norm_slow_ema_pinned` (legitimate cross-fold baseline consumer is `fold_warmup_factor_update`). +- Retain `grad_norm_fast_ema_pinned` (numerator). +- New `update_grad_norm_emas_kernel.cu` — single-thread fused fast+slow EMA update. Reads `grad_norm_buf[0]`, updates the two mapped-pinned EMA scalars in-place via dev_ptr. +- New Rust launcher `launch_update_grad_norm_emas` chained on the producer's stream after `launch_grad_norm_finalize` (no host sync; sequential same-stream ordering provides happens-before). +- `update_adaptive_clip`'s host-side `unsafe { ... }` arithmetic block replaced with the launcher call. +- Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly), FAST_ALPHA=0.1, SLOW_ALPHA=0.001 — preserved from the deleted host code bit-for-bit. **Files:** -- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` -- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (remove launch site) +- Add: `crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu` +- Modify: `crates/ml/build.rs` (cubin registration) +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (CUBIN static + struct field + constructor loader + launcher + replace host-side block) +- Modify: `crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu` (docstring update — EMAs are now GPU-updated) +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (doc-comment reflects GPU-only update path) +- Modify: `docs/dqn-wire-up-audit.md` (Layer C entry) +- Modify: this plan (the section above + this one) +- Modify: `docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md` (orphan claim removed) -- [ ] **Step 1: Remove field, alloc, free, launch, and `update_adaptive_clip` slow-EMA branch** +**No semantic change**: same α values, same cold-start sentinel, same call ordering, same observation source. Pure architectural fix per `feedback_no_cpu_compute_strict`. -Search and remove: +- [ ] **Step 1: Build verification** ```bash -grep -n "grad_norm_slow_ema_pinned\|launch_grad_norm_slow_ema" /home/jgrusewski/Work/foxhunt/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10 +SQLX_OFFLINE=true cargo check -p ml --offline 2>&1 | tail -10 ``` -Each match needs deletion. The `update_adaptive_clip` slow-EMA branch is now unreached after Layer B's Mech 6 migration; delete it. - -- [ ] **Step 2: Build + commit** - -```bash -SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5 - -git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ - crates/ml/src/trainers/dqn/trainer/training_loop.rs -git commit -m "chore(sp4): Layer C Task C1 — retire grad_norm_slow_ema infrastructure - -Mech 6 migrated to ISV[GRAD_CLIP_BOUND] in Layer B; the existing -grad_norm_slow_ema_pinned + launcher + update_adaptive_clip's slow-EMA -branch are orphan. Removed. - -The `grad_norm_ema` (fast EMA) is unrelated and stays — used by the -controllers in `update_adaptive_clip` for the per-step `new_clip` value -itself. - -cargo check clean. - -Co-Authored-By: Claude Opus 4.7 (1M context) -" -``` +- [ ] **Step 2: Single coordinated commit per `feedback_no_partial_refactor`** with commit message documenting that the original "retire" framing was wrong, the slow EMA stays, the actual fix is the host→GPU compute migration. --- diff --git a/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md b/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md index de1e9c53a..3a1bbc2dd 100644 --- a/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md +++ b/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md @@ -457,7 +457,7 @@ Single coordinated commit. Hardcoded multipliers, weight decay, L1 lambda — AL - L40S smoke validation (multi_fold_convergence::test_multi_fold_convergence) - Remove now-dead `q_abs_ref_eff` and `h_s2_rms_ema_eff` parameters from the Mech 5 fused-kernel signature (unused after Layer B's per-slot ISV reads) -- **Retire `grad_norm_slow_ema` infrastructure:** since Mech 6's adaptive_clip upper-bound now reads `ISV[GRAD_CLIP_BOUND]` instead of `grad_norm_slow_ema_pinned × 100 × isv[16]`, the existing `grad_norm_slow_ema_pinned` field, its allocation/free, its launch site, and the `update_adaptive_clip` slow-EMA branch all become orphan. Remove them. The `grad_norm_ema` (fast EMA, used elsewhere) stays. +- **Retain `grad_norm_slow_ema` infrastructure (cross-fold legitimate baseline) — migrate the EMA UPDATE to GPU per `feedback_no_cpu_compute_strict`:** the original Layer C plan called for retiring `grad_norm_slow_ema_pinned` after Mech 6's migration to `ISV[GRAD_CLIP_BOUND]`. Verification on 2026-05-01 surfaced a SECOND live consumer — `fold_warmup_factor_update` (commit `4ef1d8ebb`, predates SP4) reads the slow EMA as the cross-fold steady-state grad-norm baseline that the warmup factor's denominator compares against. Legitimate cross-fold mechanism; not orphan. The actual defect surfaced during the audit was that the EMA UPDATE itself ran as host-side `(1-α) × prev + α × obs` arithmetic in `update_adaptive_clip` — exactly the pattern `feedback_no_cpu_compute_strict` (saved 2026-05-01) strictly forbids. Layer C C1 redesigned: keep both EMA buffers, migrate the per-step UPDATE to a new single-thread `update_grad_norm_emas_kernel` chained on the producer's stream after `launch_grad_norm_finalize`. The `grad_norm_ema` (fast EMA, used elsewhere) stays. - Decide retention of `ISV[Q_ABS_REF=16]` and `ISV[H_S2_RMS_EMA=96]` producers: after Layer B, no clamp consumer reads these. They may still be useful for HEALTH_DIAG monitoring and `mag_concat_qdir`'s adaptive scale (P4.T2c.3c.6 consumer). **Decision:** keep producers running, drop only the orphaned Mech 1/2/5/6/9/10 consumer reads. Document as "monitoring-only ISV" in the audit. - Update audit docs (`dqn-backward-nan-audit.md` mechanism table, slot allocation table; `dqn-wire-up-audit.md` Invariant 7 entries — one big entry covering 40 new slots + 15 producers + Pearl C engagement-counter buffer + Wiener-state buffer + retired components) - Update `MEMORY.md` and add memory entries: