fix(sp4): migrate fast/slow grad-norm EMA update to GPU per feedback_no_cpu_compute_strict
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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`
|
||||
|
||||
@@ -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:
|
||||
*
|
||||
|
||||
@@ -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
|
||||
|
||||
67
crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu
Normal file
67
crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu
Normal file
@@ -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 */
|
||||
}
|
||||
@@ -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.
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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) <noreply@anthropic.com>
|
||||
"
|
||||
```
|
||||
- [ ] **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.
|
||||
|
||||
---
|
||||
|
||||
|
||||
@@ -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:
|
||||
|
||||
Reference in New Issue
Block a user