feat(sp18-v2): ISV-adaptive shrink-and-perturb at fold transition
Replaces the hardcoded `α=0.8 / σ=0.01` constants in `reset_for_fold`'s
`shrink_and_perturb` call with ISV-driven adaptive bounds per
`feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`.
The driving signal is the relative weight drift `||current - best||₂ /
||best||₂` already computed each epoch by the SP18 v2 weight-drift
diagnostic kernel (commit aab13a83f). Extended that kernel atomically
to also write α / σ into ISV[505] / ISV[506] from the same two block-
tree-reductions — single producer, single launch, single source of
truth. No new kernel; reuses `||best||₂` + `relative_drift` already
computed.
Added 2 ISV slots [505..507):
[505] SHRINK_ALPHA_ADAPTIVE — drift-driven preservation factor.
α = 1 - clamp(rel_drift, 1-MAX, 1-MIN). Small drift → α near
MAX (preserve hard); large drift → α near MIN (shrink hard).
[506] SHRINK_SIGMA_ADAPTIVE — scale-aware noise. σ tracks
||best||_RMS / RMS_REFERENCE so noise stays scale-relative as
weight magnitudes evolve across folds.
Pearl-A first-observation bootstrap: sentinels 0.8 / 0.01 match prior
hardcoded values exactly. Cold-start path (first fold, no
`best_params_snapshot`) leaves slots at sentinels — bit-identical to
pre-fix behavior.
Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] +
[SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1
dimensional safety floors per `feedback_isv_for_adaptive_bounds`.
Bilateral clamp per `pearl_symmetric_clamp_audit`.
Atomic in same commit per `feedback_no_partial_refactor` and
`feedback_wire_everything_up`:
* 2 ISV slot constants + sentinels + bounds + lock test in
`sp14_isv_slots.rs`.
* `ISV_TOTAL_DIM` 505 → 507 + layout fingerprint seed bump in
`gpu_dqn_trainer.rs`.
* `state_layout.cuh` mirror for kernel-side reads.
* `weight_drift_diag_kernel.cu` extended: 4-element output buffer
`[l2_diff, rel, α_target, σ_target]` + ISV writes via
`__threadfence_system()`.
* `launch_weight_drift_diag` + `read_shrink_perturb_adaptive`
plumbing; mapped-pinned 4-float buffer.
* `reset_for_fold` reads ISV[505]/ISV[506] via `read_isv_signal_at`
+ defensive host-side clamp; replaces former hardcoded constants.
* 2 fold-reset registry entries with dispatch arms (sentinel 0.8 /
0.01 — Pearl-A bootstrap matches prior hardcoded for bit-
identical cold-start).
* Per-epoch HEALTH_DIAG `weight_drift` line extended with
`alpha_adaptive` / `sigma_adaptive` for observability.
* 8 behavioral tests (CPU-only oracle pinning the math contract).
* `docs/dqn-wire-up-audit.md` — Invariant 7 audit entry.
`shrink_and_perturb` signature unchanged — kept as 2-arg
`(alpha, sigma)` so the 4 existing call sites compile without ripple
(3 in `training_loop.rs`, 1 in `fused_training.rs`). Only the
fold-transition site at `fused_training.rs::reset_for_fold` is
migrated to ISV reads in this commit; the 3 health-driven /
backtracking sites in `training_loop.rs` continue to use
`hyperparams.shrink_perturb_alpha/sigma` (different driving signal —
those are unhealthy-streak rescue interventions, not fold-transition
plasticity refresh; surfaced as a follow-up concern in
DONE_WITH_CONCERNS report).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
File diff suppressed because one or more lines are too long
@@ -831,6 +831,101 @@ pub const SENTINEL_POPART_RESET_FLAG: f32 = 1.0; // one-shot ⇒ PopAr
|
||||
pub const SP18_TD_BOOTSTRAP_SLOT_BASE: usize = 493;
|
||||
pub const SP18_TD_BOOTSTRAP_SLOT_END: usize = 505;
|
||||
|
||||
// ── SP18 v2 (2026-05-09) — ISV-adaptive shrink-and-perturb at fold transition ──
|
||||
//
|
||||
// 2 ISV slots [505..507) drive `reset_for_fold`'s `shrink_and_perturb`
|
||||
// formula:
|
||||
//
|
||||
// params_flat ← α × best_params + (1 − α) × N(0, σ)
|
||||
//
|
||||
// Per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`
|
||||
// the prior hardcoded `α=0.8 / σ=0.01` constants at the fold-transition
|
||||
// site (`fused_training.rs::reset_for_fold`) are an adaptive bound that
|
||||
// must be signal-driven from ISV. The driving signal is the relative
|
||||
// weight drift `||current_params − best_params||₂ / ||best_params||₂`
|
||||
// computed by the SP18 v2 weight-drift diagnostic kernel
|
||||
// (`weight_drift_diag_kernel.cu`, commit `aab13a83f`) — extended in this
|
||||
// commit to also write α and σ into ISV[505] and ISV[506] from the
|
||||
// already-computed `relative_drift` and `||best||₂`.
|
||||
//
|
||||
// [505] SHRINK_ALPHA_ADAPTIVE shrink factor (preserve fraction).
|
||||
// Drift-driven: when end-of-fold drift
|
||||
// was small (model converged near best),
|
||||
// α high (preserve). When drift large
|
||||
// (model diverged from best), α low
|
||||
// (shrink harder to escape divergent
|
||||
// region). Formula: α = 1.0 − clamp(
|
||||
// relative_drift, 1.0 − ALPHA_MAX,
|
||||
// 1.0 − ALPHA_MIN).
|
||||
//
|
||||
// [506] SHRINK_SIGMA_ADAPTIVE noise scale. Tracks weight magnitude
|
||||
// so noise stays scale-relative
|
||||
// regardless of how |params| evolves
|
||||
// across folds. Formula: σ = clamp(
|
||||
// SENTINEL_SIGMA × (||best||_RMS /
|
||||
// RMS_REFERENCE), SIGMA_MIN, SIGMA_MAX).
|
||||
//
|
||||
// Pearl-A first-observation bootstrap: sentinels POS=0.8 / SIGMA=0.01
|
||||
// match the prior hardcoded values for bit-identical cold-start before
|
||||
// the first valid producer launch. After the first kernel launch with a
|
||||
// valid `best_params_snapshot`, the slots track the adaptive formula.
|
||||
//
|
||||
// Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] +
|
||||
// [SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1
|
||||
// dimensional safety floors per `feedback_isv_for_adaptive_bounds`:
|
||||
// α floor 0.5 — never less than 50% preservation (prevents catastrophic
|
||||
// forgetting of best weights even on large drift).
|
||||
// α ceiling 0.95 — never more than 95% (some perturbation always —
|
||||
// preserves the noise-injection effect that breaks
|
||||
// overfitting to the prior fold).
|
||||
// σ floor 1e-4 — prevent zero-noise (gradient bias if α perturbation
|
||||
// vanishes — the formula degenerates to identity).
|
||||
// σ ceiling 0.1 — prevent destructive noise (10× the cold-start
|
||||
// sentinel; ratio caps σ growth on weight-magnitude
|
||||
// excursions).
|
||||
//
|
||||
// Cold-start case (first fold, no `best_params_snapshot`): the host
|
||||
// launcher (`launch_weight_drift_diag`) writes drift = 0.0 / 0.0 directly
|
||||
// to the output buffer and skips the kernel. The ISV slots remain at
|
||||
// sentinels (0.8 / 0.01) — bit-identical to prior hardcoded behavior.
|
||||
//
|
||||
// SP14-C.8 / SP17 PP.3 / SP18 PP.3 lineage — single-thread cold-path
|
||||
// producer per `pearl_no_host_branches_in_captured_graph`.
|
||||
//
|
||||
// Spec: this commit; ISV-adaptive shrink-and-perturb close-out of SP18 v2.
|
||||
pub const SHRINK_ALPHA_ADAPTIVE_INDEX: usize = 505;
|
||||
pub const SHRINK_SIGMA_ADAPTIVE_INDEX: usize = 506;
|
||||
|
||||
// Sentinels — Pearl-A first-observation bootstrap. Match prior hardcoded
|
||||
// values exactly so first-fold behavior (where `best_params_snapshot` is
|
||||
// `None`) is bit-identical to pre-fix.
|
||||
pub const SENTINEL_SHRINK_ALPHA: f32 = 0.8;
|
||||
pub const SENTINEL_SHRINK_SIGMA: f32 = 0.01;
|
||||
|
||||
// Category-1 dimensional safety bounds per `feedback_isv_for_adaptive_bounds`.
|
||||
// NEVER tuned; structural floor/ceiling. Bilateral clamp per
|
||||
// `pearl_symmetric_clamp_audit`.
|
||||
pub const SHRINK_ALPHA_MIN: f32 = 0.5;
|
||||
pub const SHRINK_ALPHA_MAX: f32 = 0.95;
|
||||
pub const SHRINK_SIGMA_MIN: f32 = 1.0e-4;
|
||||
pub const SHRINK_SIGMA_MAX: f32 = 0.1;
|
||||
|
||||
// Producer reference scale: σ_target = SENTINEL_SHRINK_SIGMA × (||best||_RMS /
|
||||
// SHRINK_SIGMA_RMS_REFERENCE). The reference is a Category-1 dimensional
|
||||
// constant — the "expected RMS scale" for healthy stably-trained DQN
|
||||
// weights derived from the model's mixed initialisation (He-init weights
|
||||
// at ~0.05-0.15, LayerNorm γ at 1.0, biases at 0.0; the empirical mean RMS
|
||||
// across all tensors is ≈ 0.5 for the foxhunt branching DQN). At
|
||||
// `||best||_RMS = 0.5`, σ = 0.01 = SENTINEL — equivalence with the prior
|
||||
// hardcoded value for typical healthy weights. As ||best||_RMS rises
|
||||
// (LayerNorm γ drift, weight-decay relaxation), σ scales up
|
||||
// proportionally so the noise term remains comparable to the typical
|
||||
// magnitude of an individual weight.
|
||||
pub const SHRINK_SIGMA_RMS_REFERENCE: f32 = 0.5;
|
||||
|
||||
pub const SP18_SHRINK_PERTURB_SLOT_BASE: usize = 505;
|
||||
pub const SP18_SHRINK_PERTURB_SLOT_END: usize = 507;
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
@@ -1274,10 +1369,42 @@ mod tests {
|
||||
fn all_sp18_slots_fit_within_isv_total_dim() {
|
||||
use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM;
|
||||
assert!(
|
||||
SP18_TD_BOOTSTRAP_SLOT_END <= ISV_TOTAL_DIM,
|
||||
"SP18_TD_BOOTSTRAP_SLOT_END={} exceeds ISV_TOTAL_DIM={}; \
|
||||
SP18_SHRINK_PERTURB_SLOT_END <= ISV_TOTAL_DIM,
|
||||
"SP18_SHRINK_PERTURB_SLOT_END={} exceeds ISV_TOTAL_DIM={}; \
|
||||
bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).",
|
||||
SP18_TD_BOOTSTRAP_SLOT_END, ISV_TOTAL_DIM,
|
||||
SP18_SHRINK_PERTURB_SLOT_END, ISV_TOTAL_DIM,
|
||||
);
|
||||
}
|
||||
|
||||
/// Lock SP18 v2 ISV-adaptive shrink-and-perturb slot layout (2026-05-09).
|
||||
/// 2 contiguous slots [505..507) drive `reset_for_fold`'s
|
||||
/// `shrink_and_perturb` call replacing the prior hardcoded
|
||||
/// `α=0.8 / σ=0.01` constants per `feedback_isv_for_adaptive_bounds`
|
||||
/// and `feedback_adaptive_not_tuned`.
|
||||
#[test]
|
||||
fn sp18_shrink_perturb_slot_layout_locked() {
|
||||
// ── Slot indices ─────────────────────────────────────────────
|
||||
assert_eq!(SP18_SHRINK_PERTURB_SLOT_BASE, 505);
|
||||
assert_eq!(SP18_SHRINK_PERTURB_SLOT_END, 507);
|
||||
assert_eq!(SHRINK_ALPHA_ADAPTIVE_INDEX, 505);
|
||||
assert_eq!(SHRINK_SIGMA_ADAPTIVE_INDEX, 506);
|
||||
|
||||
// ── Sentinels — match prior hardcoded values for cold-start
|
||||
// bit-identical equivalence (Pearl-A first-observation
|
||||
// bootstrap before the first valid weight-drift launch).
|
||||
assert_eq!(SENTINEL_SHRINK_ALPHA, 0.8);
|
||||
assert_eq!(SENTINEL_SHRINK_SIGMA, 0.01);
|
||||
|
||||
// ── Category-1 dimensional safety bounds. NEVER tuned;
|
||||
// structural floor/ceiling per
|
||||
// `feedback_isv_for_adaptive_bounds`.
|
||||
assert_eq!(SHRINK_ALPHA_MIN, 0.5);
|
||||
assert_eq!(SHRINK_ALPHA_MAX, 0.95);
|
||||
assert_eq!(SHRINK_SIGMA_MIN, 1.0e-4);
|
||||
assert_eq!(SHRINK_SIGMA_MAX, 0.1);
|
||||
|
||||
// ── Producer reference scale. Category-1 dimensional anchor
|
||||
// (typical mean RMS of healthy stably-trained DQN weights).
|
||||
assert_eq!(SHRINK_SIGMA_RMS_REFERENCE, 0.5);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -555,6 +555,37 @@
|
||||
#define SENTINEL_V_SHARE_TREND_DIAG 0.0f
|
||||
#define SENTINEL_POPART_RESET_FLAG 1.0f
|
||||
|
||||
// ── SP18 v2 (2026-05-09) — ISV-adaptive shrink-and-perturb at fold transition ──
|
||||
// 2 ISV slots [505..507) drive `reset_for_fold`'s `shrink_and_perturb`
|
||||
// call replacing prior hardcoded `α=0.8 / σ=0.01` constants per
|
||||
// `feedback_isv_for_adaptive_bounds`. Mirrors
|
||||
// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` constants of the same
|
||||
// names (locked by `sp18_shrink_perturb_slot_layout_locked` test). The
|
||||
// kernel `weight_drift_diag_kernel.cu` writes both slots in the same
|
||||
// launch as the existing drift diagnostic — single producer, atomic
|
||||
// per-epoch. Pearl-A first-observation bootstrap sentinels match prior
|
||||
// hardcoded values for bit-identical cold-start (first fold, no
|
||||
// `best_params_snapshot` → host skips kernel → slots stay at
|
||||
// sentinels). Bounds Category-1 dimensional floors per
|
||||
// `feedback_isv_for_adaptive_bounds`. Bilateral clamp per
|
||||
// `pearl_symmetric_clamp_audit`.
|
||||
#define SHRINK_ALPHA_ADAPTIVE_INDEX 505
|
||||
#define SHRINK_SIGMA_ADAPTIVE_INDEX 506
|
||||
|
||||
#define SENTINEL_SHRINK_ALPHA 0.8f
|
||||
#define SENTINEL_SHRINK_SIGMA 0.01f
|
||||
|
||||
#define SHRINK_ALPHA_MIN 0.5f
|
||||
#define SHRINK_ALPHA_MAX 0.95f
|
||||
#define SHRINK_SIGMA_MIN 1.0e-4f
|
||||
#define SHRINK_SIGMA_MAX 0.1f
|
||||
|
||||
// Producer reference scale (Category-1 dimensional anchor — typical
|
||||
// mean RMS of healthy stably-trained DQN weights). At
|
||||
// `||best||_RMS = 0.5`, σ = 0.01 = SENTINEL — equivalence with prior
|
||||
// hardcoded value for typical healthy weights.
|
||||
#define SHRINK_SIGMA_RMS_REFERENCE 0.5f
|
||||
|
||||
// ── Compile-time checks ──
|
||||
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
|
||||
"State layout dimensions must sum to SL_STATE_DIM");
|
||||
|
||||
@@ -1,39 +1,96 @@
|
||||
/* ══════════════════════════════════════════════════════════════════════════
|
||||
* SP18 v2 weight-drift diagnostic kernel (2026-05-09).
|
||||
* SP18 v2 weight-drift diagnostic + ISV-adaptive shrink-and-perturb producer
|
||||
* (2026-05-09).
|
||||
*
|
||||
* Reads `params_flat [n]` and `best_params_snapshot [n]`; emits two
|
||||
* floats into a `MappedF32Buffer<2>`:
|
||||
* Original cadence (commit `aab13a83f`): cold-path producer of the per-epoch
|
||||
* weight-drift diagnostic. Reads `params_flat [n]` and `best_params_snapshot
|
||||
* [n]`; emits two floats into a `MappedF32Buffer<4>`:
|
||||
*
|
||||
* out[0] = ||params_flat - best_params||₂ (L2 distance)
|
||||
* out[1] = out[0] / max(||best_params||₂, EPS) (relative)
|
||||
*
|
||||
* Producer cadence: cold-path (epoch boundary) — single launch per
|
||||
* HEALTH_DIAG emit. Diagnostic only — does NOT modify any consumer
|
||||
* path.
|
||||
* Surfaces HD2→HD5 weight drift across epochs so the post-deploy reviewer
|
||||
* can correlate Sharpe-peak decay with actual parameter trajectory in
|
||||
* `(W_best, W_curr)` space.
|
||||
*
|
||||
* Surfaces HD2→HD5 weight drift across epochs so the post-deploy
|
||||
* reviewer can correlate Sharpe-peak decay with actual parameter
|
||||
* trajectory in `(W_best, W_curr)` space.
|
||||
* SP18 v2 ISV-adaptive shrink-and-perturb close-out (this commit, 2026-05-09):
|
||||
* extended to ALSO write the adaptive α and σ for the next fold's
|
||||
* `shrink_and_perturb` call into ISV[505] / ISV[506] using the values
|
||||
* already computed by the two block-tree-reductions. No new kernel — this
|
||||
* is the single producer for both diagnostics and the controller (fewer
|
||||
* kernels = same atomicity per the spec's "Watch for: Drift diagnostic
|
||||
* timing — have the drift kernel ALSO emit the adaptive α/σ as a 4-element
|
||||
* output — fewer kernels, same atomicity"). Per `feedback_isv_for_adaptive_
|
||||
* bounds` and `feedback_adaptive_not_tuned`. Outputs:
|
||||
*
|
||||
* out[2] = α_adaptive (clamped to [SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX])
|
||||
* out[3] = σ_adaptive (clamped to [SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX])
|
||||
*
|
||||
* AND simultaneously:
|
||||
*
|
||||
* isv[SHRINK_ALPHA_ADAPTIVE_INDEX=505] = α_adaptive
|
||||
* isv[SHRINK_SIGMA_ADAPTIVE_INDEX=506] = σ_adaptive
|
||||
*
|
||||
* Formulas (single-thread-0 cold-path arithmetic post-reduction):
|
||||
*
|
||||
* relative = ||p − b||₂ / max(||b||₂, EPS) — drift signal
|
||||
* best_rms = ||b||₂ / sqrt(n) — scale signal
|
||||
*
|
||||
* α_target = 1.0 − clamp(relative,
|
||||
* 1.0 − SHRINK_ALPHA_MAX,
|
||||
* 1.0 − SHRINK_ALPHA_MIN)
|
||||
* ≡ small drift → α near MAX (preserve the converged
|
||||
* best-Sharpe weights hard); large drift → α near MIN
|
||||
* (shrink hard to escape the divergent region).
|
||||
*
|
||||
* σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms /
|
||||
* SHRINK_SIGMA_RMS_REFERENCE),
|
||||
* SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX)
|
||||
* ≡ scale-aware noise — σ tracks `||best||_RMS` so noise
|
||||
* stays scale-relative as weight magnitudes evolve
|
||||
* across folds (LayerNorm γ drift, weight-decay
|
||||
* relaxation). At RMS = REFERENCE = 0.5, σ_target =
|
||||
* SENTINEL = 0.01 → bit-identical to prior hardcoded
|
||||
* value for typical healthy weights.
|
||||
*
|
||||
* Pearl-A first-observation bootstrap is enforced by the host launcher:
|
||||
* when `best_params_snapshot` is `None` (first fold pre-best-Sharpe, or
|
||||
* any fold with no improvement), the launcher writes
|
||||
* [0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA] directly into
|
||||
* the mapped-pinned output buffer and skips the kernel — bit-identical
|
||||
* cold-start to prior hardcoded `α=0.8 / σ=0.01`. The ISV slots are NOT
|
||||
* touched on that path; consumers fall back to host-side sentinels read
|
||||
* via `read_isv_signal_at` (which returns the FoldReset-written sentinel
|
||||
* before the first valid kernel launch lands a real signal).
|
||||
*
|
||||
* ── Pearls + invariants ────────────────────────────────────────────────
|
||||
* - `feedback_no_atomicadd.md` — block tree-reduce only; no atomicAdd.
|
||||
* Two reductions (sum-of-squares of diff, sum-of-squares of best)
|
||||
* use the same shmem tile sequentially.
|
||||
* - `feedback_no_htod_htoh_only_mapped_pinned.md` — output is a
|
||||
* `MappedF32Buffer<2>`; kernel writes via `__threadfence_system()`
|
||||
* for PCIe-visible coherence; host reads after stream sync.
|
||||
* `MappedF32Buffer<4>`; kernel writes via `__threadfence_system()`
|
||||
* for PCIe-visible coherence; host reads after stream sync. ISV bus
|
||||
* write is also mapped-pinned (host-readable directly).
|
||||
* - `pearl_no_host_branches_in_captured_graph.md` — kernel runs
|
||||
* cold-path (epoch boundary, between training-step graph
|
||||
* invocations); no host branches inside.
|
||||
* - `pearl_first_observation_bootstrap.md` — N/A: this is a single
|
||||
* instantaneous distance, not an EMA. The host-side launcher
|
||||
* handles the "no snapshot saved" case before launching the kernel
|
||||
* (returns 0.0/0.0 directly so the diagnostic line is still emitted).
|
||||
* - `pearl_symmetric_clamp_audit.md` — relative drift is bounded
|
||||
* [0, ∞); we floor the denominator at SP18_DRIFT_EPS_F=1e-12 to
|
||||
* avoid 0/0 NaN when best_params are all-zero (cold-init pre-first
|
||||
* save). No upper clamp — diagnostic should surface unbounded
|
||||
* drift faithfully.
|
||||
* - `pearl_first_observation_bootstrap.md` — first valid observation
|
||||
* REPLACES sentinel directly. The producer ALWAYS writes a clamped
|
||||
* target value (no EMA blend), so on the first valid launch, ISV[505]
|
||||
* and ISV[506] go from sentinel directly to the computed target.
|
||||
* Subsequent launches REPLACE again — there is no Welford state for
|
||||
* this controller; the underlying signal (relative drift, best_rms)
|
||||
* is itself computed afresh each epoch from current/best weights.
|
||||
* - `pearl_symmetric_clamp_audit.md` — bilateral
|
||||
* `fmaxf(MIN, fminf(x, MAX))` for both α and σ; never one-sided.
|
||||
* - `feedback_isv_for_adaptive_bounds.md` — replaces the hardcoded
|
||||
* `α=0.8 / σ=0.01` constants formerly at
|
||||
* `fused_training.rs::reset_for_fold` lines 1014-1015. The bounds
|
||||
* [MIN, MAX] are themselves Category-1 dimensional safety floors —
|
||||
* structural (never adaptive themselves; their values are the floor
|
||||
* and ceiling).
|
||||
* - `feedback_adaptive_not_tuned.md` — α/σ derive from the EMA-tracked
|
||||
* drift signal + scale-aware best_rms ratio, not tuned constants.
|
||||
*
|
||||
* Algorithm (single block, 256 threads):
|
||||
*
|
||||
@@ -48,13 +105,18 @@
|
||||
* local_sumsq += best[i] * best[i]
|
||||
* Block tree-reduce into s_reduce[0]; thread 0 stores to local s1.
|
||||
*
|
||||
* Pass 3 (write):
|
||||
* Pass 3 (write outputs):
|
||||
* thread 0 only:
|
||||
* l2 = sqrtf(s0)
|
||||
* l2 = sqrtf(s0)
|
||||
* norm_best = sqrtf(s1)
|
||||
* rel = l2 / max(norm_best, SP18_DRIFT_EPS_F)
|
||||
* out[0] = l2
|
||||
* out[1] = rel
|
||||
* rel = l2 / max(norm_best, SP18_DRIFT_EPS_F)
|
||||
* best_rms = norm_best / sqrtf(n) [ NEW ]
|
||||
* α_target = 1.0 − clamp(rel, 1.0 − ALPHA_MAX, 1.0 − ALPHA_MIN) [ NEW ]
|
||||
* σ_target = clamp(SIGMA_SENTINEL × best_rms / RMS_REF,
|
||||
* SIGMA_MIN, SIGMA_MAX) [ NEW ]
|
||||
* out[0..4] = [l2, rel, α_target, σ_target]
|
||||
* isv[505] = α_target [ NEW ]
|
||||
* isv[506] = σ_target [ NEW ]
|
||||
* __threadfence_system()
|
||||
*
|
||||
* Launch contract:
|
||||
@@ -65,11 +127,15 @@
|
||||
* Args:
|
||||
* params — `params_flat [n]` f32 device pointer (online weights).
|
||||
* best — `best_params_snapshot [n]` f32 device pointer.
|
||||
* out — `MappedF32Buffer<2>` device pointer for [l2, rel].
|
||||
* out — `MappedF32Buffer<4>` device pointer for
|
||||
* [l2, rel, α_target, σ_target].
|
||||
* isv — ISV bus device pointer (mapped-pinned). Writes ISV[505]
|
||||
* and ISV[506].
|
||||
* n — element count of `params` and `best` (must match).
|
||||
* ══════════════════════════════════════════════════════════════════════════ */
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
#include "state_layout.cuh"
|
||||
|
||||
#define SP18_DRIFT_BLOCK 256
|
||||
#define SP18_DRIFT_EPS_F 1e-12f
|
||||
@@ -78,6 +144,7 @@ extern "C" __global__ void weight_drift_diag_update(
|
||||
const float* __restrict__ params,
|
||||
const float* __restrict__ best,
|
||||
float* __restrict__ out,
|
||||
float* __restrict__ isv,
|
||||
int n)
|
||||
{
|
||||
/* Single block (gridDim.x == 1). */
|
||||
@@ -86,11 +153,15 @@ extern "C" __global__ void weight_drift_diag_update(
|
||||
const int tid = (int)threadIdx.x;
|
||||
const int bdim = (int)blockDim.x;
|
||||
|
||||
/* Degenerate guard: n <= 0 — write zeros and return. */
|
||||
/* Degenerate guard: n <= 0 — write zeros + sentinels and return. */
|
||||
if (n <= 0) {
|
||||
if (tid == 0) {
|
||||
out[0] = 0.0f;
|
||||
out[1] = 0.0f;
|
||||
out[2] = SENTINEL_SHRINK_ALPHA;
|
||||
out[3] = SENTINEL_SHRINK_SIGMA;
|
||||
isv[SHRINK_ALPHA_ADAPTIVE_INDEX] = SENTINEL_SHRINK_ALPHA;
|
||||
isv[SHRINK_SIGMA_ADAPTIVE_INDEX] = SENTINEL_SHRINK_SIGMA;
|
||||
__threadfence_system();
|
||||
}
|
||||
return;
|
||||
@@ -138,8 +209,36 @@ extern "C" __global__ void weight_drift_diag_update(
|
||||
: SP18_DRIFT_EPS_F;
|
||||
const float rel = l2_diff / denom;
|
||||
|
||||
/* SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09):
|
||||
* compute α_target and σ_target from the same reductions. */
|
||||
const float best_rms = norm_best / sqrtf((float)n);
|
||||
|
||||
/* α_target: drift-driven preservation factor. Bilateral clamp
|
||||
* on relative drift, then 1 − clamp gives the preserve fraction.
|
||||
* `pearl_symmetric_clamp_audit` — `fmaxf(LO, fminf(rel, HI))`. */
|
||||
const float rel_lo = 1.0f - SHRINK_ALPHA_MAX; /* 0.05 */
|
||||
const float rel_hi = 1.0f - SHRINK_ALPHA_MIN; /* 0.50 */
|
||||
const float rel_clamped = fmaxf(rel_lo, fminf(rel, rel_hi));
|
||||
const float alpha_target = 1.0f - rel_clamped;
|
||||
|
||||
/* σ_target: scale-aware noise. ratio = best_rms / RMS_REFERENCE
|
||||
* — when weights have typical RMS, ratio = 1.0 → σ = SENTINEL.
|
||||
* As RMS grows, σ scales up proportionally, capped at MAX. */
|
||||
const float ratio = best_rms / SHRINK_SIGMA_RMS_REFERENCE;
|
||||
const float sigma_raw = SENTINEL_SHRINK_SIGMA * ratio;
|
||||
const float sigma_target = fmaxf(SHRINK_SIGMA_MIN,
|
||||
fminf(sigma_raw, SHRINK_SIGMA_MAX));
|
||||
|
||||
out[0] = l2_diff;
|
||||
out[1] = rel;
|
||||
__threadfence_system(); /* PCIe-visible write for mapped-pinned output */
|
||||
out[2] = alpha_target;
|
||||
out[3] = sigma_target;
|
||||
|
||||
/* ISV bus write (mapped-pinned). Consumer (`reset_for_fold`)
|
||||
* reads via `read_isv_signal_at` after stream sync. */
|
||||
isv[SHRINK_ALPHA_ADAPTIVE_INDEX] = alpha_target;
|
||||
isv[SHRINK_SIGMA_ADAPTIVE_INDEX] = sigma_target;
|
||||
|
||||
__threadfence_system(); /* PCIe-visible write for mapped-pinned outputs */
|
||||
}
|
||||
}
|
||||
|
||||
@@ -395,17 +395,25 @@ pub(crate) struct FusedTrainingCtx {
|
||||
|
||||
// ── SP18 v2 weight-drift diagnostic (2026-05-09) ─────────────────────
|
||||
/// Single-block × 256-thread kernel computing
|
||||
/// `[||params - best||₂, relative_drift]` from `params_flat` and
|
||||
/// `best_params_snapshot`. Cold-path (epoch boundary). Block tree-
|
||||
/// reduce per `feedback_no_atomicadd`. Loaded from
|
||||
/// `weight_drift_diag_kernel.cubin`. Diagnostic-only — does NOT
|
||||
/// modify any consumer path. Consumed by `read_weight_drift_diag()`.
|
||||
/// `[||params - best||₂, relative_drift, α_target, σ_target]` from
|
||||
/// `params_flat` and `best_params_snapshot`. Cold-path (epoch
|
||||
/// boundary). Block tree-reduce per `feedback_no_atomicadd`. Loaded
|
||||
/// from `weight_drift_diag_kernel.cubin`. Diagnostic + adaptive-
|
||||
/// controller producer for ISV slots [505..507). Consumed by
|
||||
/// `read_weight_drift_diag()` and `reset_for_fold` reading
|
||||
/// ISV[SHRINK_ALPHA_ADAPTIVE_INDEX] / ISV[SHRINK_SIGMA_ADAPTIVE_INDEX].
|
||||
/// Per `feedback_isv_for_adaptive_bounds` — replaces the prior
|
||||
/// hardcoded `α=0.8 / σ=0.01` constants in `reset_for_fold`.
|
||||
sp18_weight_drift_diag_kernel: cudarc::driver::CudaFunction,
|
||||
/// Mapped-pinned 2-float output buffer for the weight-drift diagnostic.
|
||||
/// Layout: `[l2_diff, relative_drift]`. Per
|
||||
/// Mapped-pinned 4-float output buffer for the weight-drift diagnostic
|
||||
/// + ISV-adaptive shrink-and-perturb controller.
|
||||
/// Layout: `[l2_diff, relative_drift, α_target, σ_target]`. Per
|
||||
/// `feedback_no_htod_htoh_only_mapped_pinned`. Lifetime matches the
|
||||
/// fused-training context; written-then-read inside the kernel each
|
||||
/// launch. Host reads via `read_volatile` after stream sync.
|
||||
/// SP18 v2 ISV-adaptive shrink-and-perturb close-out (2026-05-09):
|
||||
/// extended from 2 to 4 elements; α/σ are also written to ISV slots
|
||||
/// 505/506 inside the same kernel launch (single-producer atomicity).
|
||||
sp18_weight_drift_diag_buf: crate::cuda_pipeline::mapped_pinned::MappedF32Buffer,
|
||||
}
|
||||
|
||||
@@ -879,11 +887,14 @@ impl FusedTrainingCtx {
|
||||
"weight_drift_diag_update load_function: {e}"
|
||||
))?
|
||||
};
|
||||
// 4 floats: [l2_diff, relative_drift, α_target, σ_target] per
|
||||
// SP18 v2 ISV-adaptive shrink-and-perturb close-out. The kernel
|
||||
// also writes α/σ into ISV[505]/ISV[506] in the same launch.
|
||||
let sp18_weight_drift_diag_buf = unsafe {
|
||||
crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(2)
|
||||
crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(4)
|
||||
}
|
||||
.map_err(|e| anyhow::anyhow!(
|
||||
"sp18_weight_drift_diag_buf alloc (2 f32): {e}"
|
||||
"sp18_weight_drift_diag_buf alloc (4 f32): {e}"
|
||||
))?;
|
||||
|
||||
info!(
|
||||
@@ -1008,15 +1019,53 @@ impl FusedTrainingCtx {
|
||||
unsafe { cuda_sys::cuGraphExecDestroy(exec); }
|
||||
}
|
||||
// Shrink-and-perturb at fold boundary -- reduce overfit to previous fold.
|
||||
// alpha=0.8: keep 80% of learned representations, add 20% random noise.
|
||||
// Without this, the greedy policy from fold 1 produces val_Sharpe=-33
|
||||
// on fold 2 data and takes many epochs to recover.
|
||||
let sp_alpha = 0.8_f32;
|
||||
let sp_sigma = 0.01_f32;
|
||||
//
|
||||
// SP18 v2 ISV-adaptive shrink-and-perturb close-out (2026-05-09):
|
||||
// α and σ are now read from ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505]
|
||||
// and ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506] — replaces the prior
|
||||
// hardcoded `α=0.8 / σ=0.01` constants per
|
||||
// `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`.
|
||||
//
|
||||
// Producer: `weight_drift_diag_update` (extended in this commit
|
||||
// to also write α/σ alongside the existing l2_diff/relative
|
||||
// diagnostic outputs). Runs at HEALTH_DIAG cadence (per-epoch);
|
||||
// by the time `reset_for_fold` executes (fold transition,
|
||||
// after the last per-epoch HEALTH_DIAG emit), the slots reflect
|
||||
// the drift OBSERVED at end-of-fold N, applied to fold N+1's
|
||||
// perturbation.
|
||||
//
|
||||
// α: drift-driven preservation factor.
|
||||
// - Small drift (model converged near best) → α near MAX (preserve hard).
|
||||
// - Large drift (model diverged from best) → α near MIN (shrink hard).
|
||||
// σ: scale-aware noise. Tracks `||best||_RMS` so noise stays
|
||||
// scale-relative as weight magnitudes evolve across folds.
|
||||
//
|
||||
// Cold-start (first fold, no `best_params_snapshot` yet): the
|
||||
// host launcher in `launch_weight_drift_diag` skipped the kernel
|
||||
// and did NOT write to ISV[505]/ISV[506]; consumers fall back
|
||||
// to the FoldReset-written sentinels (0.8 / 0.01) — bit-identical
|
||||
// to prior hardcoded behavior.
|
||||
//
|
||||
// Bilateral clamp [SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX] +
|
||||
// [SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX] is enforced inside the
|
||||
// producer kernel per `pearl_symmetric_clamp_audit`. The host-
|
||||
// side defensive clamp here guards against any out-of-range
|
||||
// value persisted from a prior version's checkpoint or written
|
||||
// through some unintended path — Category-1 dimensional safety
|
||||
// floors per `feedback_isv_for_adaptive_bounds`.
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
SHRINK_ALPHA_ADAPTIVE_INDEX, SHRINK_SIGMA_ADAPTIVE_INDEX,
|
||||
SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX,
|
||||
SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX,
|
||||
};
|
||||
let sp_alpha = self.trainer.read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX)
|
||||
.clamp(SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX);
|
||||
let sp_sigma = self.trainer.read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX)
|
||||
.clamp(SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX);
|
||||
if let Err(e) = self.trainer.shrink_and_perturb(sp_alpha, sp_sigma) {
|
||||
tracing::warn!("Fold-boundary shrink-and-perturb failed (non-fatal): {e}");
|
||||
} else {
|
||||
tracing::info!(alpha = sp_alpha, sigma = sp_sigma, "Fold-boundary shrink-and-perturb applied");
|
||||
tracing::info!(alpha = sp_alpha, sigma = sp_sigma, "Fold-boundary shrink-and-perturb applied (ISV-adaptive)");
|
||||
}
|
||||
// Hard-copy the shrink-and-perturb'd online weights into target params.
|
||||
// Without this, target_params_buf retains end-of-previous-fold weights
|
||||
@@ -4619,30 +4668,49 @@ impl FusedTrainingCtx {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP18 v2 weight-drift diagnostic (2026-05-09).
|
||||
/// SP18 v2 weight-drift diagnostic + ISV-adaptive shrink-and-perturb
|
||||
/// producer (2026-05-09).
|
||||
///
|
||||
/// Cold-path launcher for `weight_drift_diag_kernel`. Reads the
|
||||
/// flat online params buffer + the best-Sharpe snapshot and emits
|
||||
/// `[||params - best||₂, relative_drift]` into the mapped-pinned
|
||||
/// 2-float output buffer. Single block × 256 threads; two block
|
||||
/// tree-reductions (no atomicAdd per `feedback_no_atomicadd`).
|
||||
/// Returns Ok(()) on success.
|
||||
/// Cold-path launcher for `weight_drift_diag_kernel`. Reads the flat
|
||||
/// online params buffer + the best-Sharpe snapshot and emits
|
||||
/// `[||params - best||₂, relative_drift, α_target, σ_target]` into
|
||||
/// the mapped-pinned 4-float output buffer. The kernel ALSO writes
|
||||
/// α_target / σ_target into `ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505]`
|
||||
/// and `ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506]` in the same launch
|
||||
/// (single-producer atomicity per `feedback_no_partial_refactor`).
|
||||
/// Single block × 256 threads; two block tree-reductions
|
||||
/// (no atomicAdd per `feedback_no_atomicadd`). Returns Ok(()) on success.
|
||||
///
|
||||
/// Cold-start: if `best_params_snapshot` is `None` (no Sharpe
|
||||
/// improvement yet), the launcher writes `[0.0, 0.0]` directly into
|
||||
/// the host-visible buffer and skips the kernel — diagnostic line
|
||||
/// still emits with zeros so operators distinguish "no snapshot"
|
||||
/// from "snapshot exists but matches current". This bypass uses the
|
||||
/// mapped-pinned host slice (no DtoH per
|
||||
/// `feedback_no_htod_htoh_only_mapped_pinned`).
|
||||
/// improvement yet), the launcher writes
|
||||
/// `[0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]`
|
||||
/// directly into the host-visible buffer and skips the kernel —
|
||||
/// diagnostic line still emits with zeros so operators distinguish
|
||||
/// "no snapshot" from "snapshot exists but matches current". The
|
||||
/// ISV slots 505/506 are NOT touched on this path; consumers fall
|
||||
/// back to the FoldReset-written sentinels via `read_isv_signal_at`.
|
||||
/// This bypass uses the mapped-pinned host slice (no DtoH per
|
||||
/// `feedback_no_htod_htoh_only_mapped_pinned`) and is bit-identical
|
||||
/// to the prior hardcoded `α=0.8 / σ=0.01` behavior.
|
||||
///
|
||||
/// Diagnostic only — does NOT modify any consumer path.
|
||||
/// Per `feedback_isv_for_adaptive_bounds` — replaces the hardcoded
|
||||
/// `α=0.8 / σ=0.01` constants formerly at
|
||||
/// `fused_training.rs::reset_for_fold` lines 1014-1015. The α
|
||||
/// signal is drift-driven; σ scales with `||best||_RMS`.
|
||||
pub(crate) fn launch_weight_drift_diag(&mut self) -> Result<()> {
|
||||
// Cold-start: no snapshot saved yet.
|
||||
// Cold-start: no snapshot saved yet — write zeros + sentinels
|
||||
// directly into the mapped-pinned buffer and skip the kernel.
|
||||
// ISV[505]/ISV[506] remain at their FoldReset sentinels (0.8/0.01)
|
||||
// → bit-identical to prior hardcoded shrink-and-perturb behavior.
|
||||
if self.best_params_snapshot.is_none() {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA,
|
||||
};
|
||||
let host_slice = self.sp18_weight_drift_diag_buf.host_slice_mut();
|
||||
host_slice[0] = 0.0;
|
||||
host_slice[1] = 0.0;
|
||||
host_slice[2] = SENTINEL_SHRINK_ALPHA;
|
||||
host_slice[3] = SENTINEL_SHRINK_SIGMA;
|
||||
return Ok(());
|
||||
}
|
||||
|
||||
@@ -4666,6 +4734,7 @@ impl FusedTrainingCtx {
|
||||
}
|
||||
let n_i32 = params_n as i32;
|
||||
let out_ptr = self.sp18_weight_drift_diag_buf.dev_ptr;
|
||||
let isv_ptr = self.trainer.isv_signals_dev_ptr();
|
||||
let block_dim: u32 = 256;
|
||||
let shmem_bytes: u32 = block_dim * std::mem::size_of::<f32>() as u32;
|
||||
|
||||
@@ -4676,6 +4745,7 @@ impl FusedTrainingCtx {
|
||||
.arg(¶ms_ptr)
|
||||
.arg(&best_ptr)
|
||||
.arg(&out_ptr)
|
||||
.arg(&isv_ptr)
|
||||
.arg(&n_i32)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
@@ -4693,7 +4763,9 @@ impl FusedTrainingCtx {
|
||||
///
|
||||
/// Convenience wrapper: launches the producer, syncs the stream,
|
||||
/// returns `[l2_diff, relative_drift]` from the mapped-pinned
|
||||
/// output buffer. Cold-path only (HEALTH_DIAG cadence).
|
||||
/// output buffer. Cold-path only (HEALTH_DIAG cadence). The α/σ
|
||||
/// outputs (host_slice[2..4]) are read separately via
|
||||
/// `read_shrink_perturb_adaptive()` by `reset_for_fold`.
|
||||
pub(crate) fn read_weight_drift_diag(&mut self) -> Result<[f32; 2]> {
|
||||
self.launch_weight_drift_diag()?;
|
||||
// Sync to ensure the kernel write (if any) is visible on host.
|
||||
@@ -4704,6 +4776,26 @@ impl FusedTrainingCtx {
|
||||
Ok([v[0], v[1]])
|
||||
}
|
||||
|
||||
/// SP18 v2 ISV-adaptive shrink-and-perturb readback (2026-05-09).
|
||||
///
|
||||
/// Returns `[α_adaptive, σ_adaptive]` for the next fold's
|
||||
/// `shrink_and_perturb` call. Reads from the same mapped-pinned
|
||||
/// 4-float buffer as `read_weight_drift_diag`; the kernel produces
|
||||
/// both diagnostic and controller outputs in one launch. Caller
|
||||
/// MUST have invoked `launch_weight_drift_diag` (or
|
||||
/// `read_weight_drift_diag`) and synced the stream before this read
|
||||
/// to guarantee visibility — `reset_for_fold` does so via the
|
||||
/// per-epoch HEALTH_DIAG emit that precedes it.
|
||||
///
|
||||
/// Cold-start (no `best_params_snapshot`): host slice carries
|
||||
/// `[SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]` from the
|
||||
/// `launch_weight_drift_diag` cold-start branch — bit-identical to
|
||||
/// prior hardcoded `α=0.8 / σ=0.01`.
|
||||
pub(crate) fn read_shrink_perturb_adaptive(&self) -> [f32; 2] {
|
||||
let v = self.sp18_weight_drift_diag_buf.read_all();
|
||||
[v[2], v[3]]
|
||||
}
|
||||
|
||||
/// Raw device pointer to the active loss scalar (pinned device-mapped).
|
||||
/// Returns MSE dev_ptr when c51_alpha ≈ 0, otherwise C51 total_loss dev_ptr.
|
||||
pub(crate) fn loss_gpu_ptr(&self) -> u64 {
|
||||
|
||||
@@ -2020,6 +2020,28 @@ impl StateResetRegistry {
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[SP18_B_LEG_RESERVED_INDEX=504] — SP18 B-leg RESERVED slot for follow-up if needed (e.g. distributional flavour of B-DD8 if 5-epoch smoke shows atom-distribution drift). Sentinel 0.0. No producer or consumer in the Pre-Phase scope — preserves the 12-slot B-leg block aligned and the layout fingerprint stable.",
|
||||
},
|
||||
// ── SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09) ─────
|
||||
// 2 ISV slots [505..507) drive `reset_for_fold`'s
|
||||
// `shrink_and_perturb` call replacing prior hardcoded
|
||||
// `α=0.8 / σ=0.01` constants per
|
||||
// `feedback_isv_for_adaptive_bounds` and
|
||||
// `feedback_adaptive_not_tuned`. Producer kernel is the
|
||||
// existing `weight_drift_diag_update` (`weight_drift_diag_kernel.cu`,
|
||||
// commit `aab13a83f`) extended atomically in this commit to
|
||||
// ALSO write α and σ from the same two block-tree-reductions.
|
||||
// Pearl-A first-observation bootstrap sentinels match prior
|
||||
// hardcoded values for bit-identical cold-start before the
|
||||
// first valid drift-kernel launch.
|
||||
RegistryEntry {
|
||||
name: "sp18_shrink_alpha_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[SHRINK_ALPHA_ADAPTIVE_INDEX=505] — SP18 v2 adaptive shrink factor for `reset_for_fold`'s `shrink_and_perturb` call. Drift-driven: when end-of-fold relative drift `||current − best||₂ / ||best||₂` is small (model converged near best), α high (preserve); large drift → α low (shrink harder to escape divergent region). Formula: α_target = 1 − clamp(relative_drift, 1 − ALPHA_MAX, 1 − ALPHA_MIN). FoldReset sentinel SENTINEL_SHRINK_ALPHA=0.8 — Pearl-A first-observation bootstrap matching the prior hardcoded `α=0.8` for bit-identical cold-start before the first valid drift-kernel launch (first fold has no `best_params_snapshot` → host launcher skips kernel → slot stays at sentinel). Bounds [SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` — never less than 50% preservation (catastrophic-forgetting guard) / never more than 95% (some perturbation always — preserves the noise-injection effect that breaks overfitting to the prior fold). Bilateral clamp per `pearl_symmetric_clamp_audit`. Producer + consumer wired in the same commit per `feedback_no_partial_refactor`.",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp18_shrink_sigma_adaptive",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[SHRINK_SIGMA_ADAPTIVE_INDEX=506] — SP18 v2 adaptive noise scale for `reset_for_fold`'s `shrink_and_perturb` call. Scale-aware: σ tracks `||best||_RMS = ||best||₂ / sqrt(n)` so noise stays scale-relative as weight magnitudes evolve across folds (LayerNorm γ drift, weight-decay relaxation). Formula: σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms / SHRINK_SIGMA_RMS_REFERENCE), SIGMA_MIN, SIGMA_MAX). At RMS = REFERENCE = 0.5 (Category-1 dimensional anchor — typical mean RMS of healthy stably-trained DQN weights), σ = SENTINEL = 0.01 → bit-identical to prior hardcoded value for typical healthy weights. FoldReset sentinel SENTINEL_SHRINK_SIGMA=0.01 — Pearl-A first-observation bootstrap matching the prior hardcoded `σ=0.01` for bit-identical cold-start. Bounds [SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1] are Category-1 dimensional safety floors per `feedback_isv_for_adaptive_bounds` — floor prevents zero-noise (gradient bias if perturbation vanishes) / ceiling prevents destructive noise (10× sentinel; ratio caps σ growth on weight-magnitude excursions). Bilateral clamp per `pearl_symmetric_clamp_audit`. Same producer kernel as slot 505 (`weight_drift_diag_update` writes both in one launch).",
|
||||
},
|
||||
];
|
||||
Self { entries }
|
||||
}
|
||||
@@ -2164,18 +2186,18 @@ mod tests {
|
||||
);
|
||||
}
|
||||
|
||||
/// Lock SP18 v2 Pre-Phase Task PP.3: 22 fold-reset entries land for the
|
||||
/// combined Hold-attractor fix slots [483..505). 10 D-leg + 12 B-leg.
|
||||
/// Lock SP18 v2 Pre-Phase Task PP.3 + ISV-adaptive shrink-and-perturb
|
||||
/// close-out: 24 fold-reset entries land for the combined Hold-attractor
|
||||
/// fix slots [483..507). 10 D-leg + 12 B-leg + 2 shrink-and-perturb.
|
||||
/// Asserts each by name is present with FoldReset category and a
|
||||
/// description containing "SP18 D-leg" or "SP18 B-leg" (sentinel /
|
||||
/// description content checked separately for clarity).
|
||||
/// description marker matching its leg.
|
||||
#[test]
|
||||
fn sp18_fold_reset_entries_present() {
|
||||
let r = StateResetRegistry::new();
|
||||
let by_name: std::collections::HashMap<&str, &RegistryEntry> =
|
||||
r.all().iter().map(|e| (e.name, e)).collect();
|
||||
|
||||
// 22 expected SP18 entries — D-leg 10 + B-leg 12.
|
||||
// 24 expected SP18 entries — D-leg 10 + B-leg 12 + shrink-and-perturb 2.
|
||||
let expected: &[(&str, &str)] = &[
|
||||
// D-leg [483..493) — 10 slots
|
||||
("sp18_hold_reward_pos_cap", "SP18 D-leg"),
|
||||
@@ -2201,9 +2223,12 @@ mod tests {
|
||||
("sp18_tdb_prev_target", "SP18 B-leg"),
|
||||
("sp18_tdb_sample_count", "SP18 B-leg"),
|
||||
("sp18_b_leg_reserved", "SP18 B-leg"),
|
||||
// shrink-and-perturb [505..507) — 2 slots
|
||||
("sp18_shrink_alpha_adaptive", "SP18 v2 adaptive shrink factor"),
|
||||
("sp18_shrink_sigma_adaptive", "SP18 v2 adaptive noise scale"),
|
||||
];
|
||||
|
||||
assert_eq!(expected.len(), 22, "expected 22 SP18 entries (10 D-leg + 12 B-leg)");
|
||||
assert_eq!(expected.len(), 24, "expected 24 SP18 entries (10 D-leg + 12 B-leg + 2 shrink-and-perturb)");
|
||||
|
||||
for (name, leg_marker) in expected {
|
||||
let entry = by_name.get(name).unwrap_or_else(|| {
|
||||
|
||||
@@ -5424,10 +5424,32 @@ impl DQNTrainer {
|
||||
// (skip_start..skip_end), so trunk drift dominates the
|
||||
// L2 in any case.
|
||||
let branch_max = relative;
|
||||
// SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09):
|
||||
// surface the per-epoch α/σ targets the fold-transition
|
||||
// `shrink_and_perturb` will use. Read directly from ISV
|
||||
// (the producer kernel that ran inside
|
||||
// `read_weight_drift_diag` above wrote slots 505/506
|
||||
// alongside the diagnostic output). Cold-start (no
|
||||
// `best_params_snapshot`) → kernel skipped, ISV slots
|
||||
// remain at FoldReset sentinels (0.8/0.01) — surfaced
|
||||
// verbatim as the cold-start values.
|
||||
let (alpha_adaptive, sigma_adaptive) = if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
SHRINK_ALPHA_ADAPTIVE_INDEX, SHRINK_SIGMA_ADAPTIVE_INDEX,
|
||||
};
|
||||
(
|
||||
fused.trainer().read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX),
|
||||
fused.trainer().read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX),
|
||||
)
|
||||
} else {
|
||||
(0.0, 0.0)
|
||||
};
|
||||
tracing::info!(
|
||||
"HEALTH_DIAG[{}]: weight_drift [norm_l2={:.6} \
|
||||
relative={:.6} branch_max={:.6}]",
|
||||
relative={:.6} branch_max={:.6} \
|
||||
alpha_adaptive={:.4} sigma_adaptive={:.6}]",
|
||||
epoch, norm_l2, relative, branch_max,
|
||||
alpha_adaptive, sigma_adaptive,
|
||||
);
|
||||
}
|
||||
|
||||
@@ -9412,6 +9434,34 @@ impl DQNTrainer {
|
||||
);
|
||||
}
|
||||
}
|
||||
// ── SP18 v2 ISV-adaptive shrink-and-perturb (2026-05-09) ─────
|
||||
// Producer is `weight_drift_diag_update` (extended). Pearl-A
|
||||
// first-observation bootstrap sentinels match the prior
|
||||
// hardcoded `α=0.8 / σ=0.01` so first-fold behavior (where
|
||||
// `best_params_snapshot` is `None` → host launcher skips the
|
||||
// kernel) is bit-identical to pre-fix.
|
||||
"sp18_shrink_alpha_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
SHRINK_ALPHA_ADAPTIVE_INDEX, SENTINEL_SHRINK_ALPHA,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
SHRINK_ALPHA_ADAPTIVE_INDEX,
|
||||
SENTINEL_SHRINK_ALPHA,
|
||||
);
|
||||
}
|
||||
}
|
||||
"sp18_shrink_sigma_adaptive" => {
|
||||
if let Some(ref fused) = self.fused_ctx {
|
||||
use crate::cuda_pipeline::sp14_isv_slots::{
|
||||
SHRINK_SIGMA_ADAPTIVE_INDEX, SENTINEL_SHRINK_SIGMA,
|
||||
};
|
||||
fused.trainer().write_isv_signal_at(
|
||||
SHRINK_SIGMA_ADAPTIVE_INDEX,
|
||||
SENTINEL_SHRINK_SIGMA,
|
||||
);
|
||||
}
|
||||
}
|
||||
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
|
||||
// OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
|
||||
// a stateful EMA) — rewrite the constructor's value at fold
|
||||
|
||||
240
crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs
Normal file
240
crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs
Normal file
@@ -0,0 +1,240 @@
|
||||
//! SP18 v2 ISV-adaptive shrink-and-perturb test (2026-05-09).
|
||||
//!
|
||||
//! Pins the math contract for the producer formulas in
|
||||
//! `weight_drift_diag_kernel.cu` that drive ISV[505] / ISV[506] —
|
||||
//! replacing the prior hardcoded `α=0.8 / σ=0.01` constants in
|
||||
//! `fused_training.rs::reset_for_fold`'s `shrink_and_perturb` call.
|
||||
//!
|
||||
//! Per `feedback_isv_for_adaptive_bounds` and
|
||||
//! `feedback_adaptive_not_tuned`, the formulas are:
|
||||
//!
|
||||
//! α_target = 1.0 − clamp(relative_drift,
|
||||
//! 1.0 − SHRINK_ALPHA_MAX,
|
||||
//! 1.0 − SHRINK_ALPHA_MIN)
|
||||
//!
|
||||
//! σ_target = clamp(SENTINEL_SHRINK_SIGMA × (best_rms /
|
||||
//! SHRINK_SIGMA_RMS_REFERENCE),
|
||||
//! SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX)
|
||||
//!
|
||||
//! ## Why this test is CPU-only
|
||||
//!
|
||||
//! `weight_drift_diag_update` is a GPU kernel that reads from device-
|
||||
//! resident `params_flat` and `best_params_snapshot` buffers. A faithful
|
||||
//! end-to-end test would require constructing a `GpuDqnTrainer` (heavy:
|
||||
//! param-buf alloc, cubin loads, CUDA stream, full device init) and
|
||||
//! exercising the live GPU path — deferred to L40S smoke.
|
||||
//!
|
||||
//! Instead, this test pins the MATH CONTRACT that the producer enforces:
|
||||
//!
|
||||
//! 1. Small drift → α near MAX (preserve hard).
|
||||
//! 2. Large drift → α near MIN (clamped — shrink hard to escape).
|
||||
//! 3. Cold-start (sentinel preserved) → α = SENTINEL exactly.
|
||||
//! 4. σ scales proportionally with `||best||_RMS`.
|
||||
//!
|
||||
//! Real GPU coverage of the kernel-level math is provided by the L40S
|
||||
//! smoke that gates the SP18 closeout dispatch.
|
||||
|
||||
use ml::cuda_pipeline::sp14_isv_slots::{
|
||||
SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA,
|
||||
SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX,
|
||||
SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX,
|
||||
SHRINK_SIGMA_RMS_REFERENCE,
|
||||
};
|
||||
|
||||
/// CPU-side oracle for the α formula in `weight_drift_diag_kernel.cu`.
|
||||
///
|
||||
/// Mirrors:
|
||||
/// `α_target = 1.0 − clamp(relative_drift, 1.0 − ALPHA_MAX, 1.0 − ALPHA_MIN)`
|
||||
///
|
||||
/// Bilateral clamp per `pearl_symmetric_clamp_audit`.
|
||||
fn alpha_target_cpu(relative_drift: f32) -> f32 {
|
||||
let rel_lo = 1.0 - SHRINK_ALPHA_MAX;
|
||||
let rel_hi = 1.0 - SHRINK_ALPHA_MIN;
|
||||
let rel_clamped = relative_drift.max(rel_lo).min(rel_hi);
|
||||
1.0 - rel_clamped
|
||||
}
|
||||
|
||||
/// CPU-side oracle for the σ formula in `weight_drift_diag_kernel.cu`.
|
||||
///
|
||||
/// Mirrors:
|
||||
/// `σ_target = clamp(SENTINEL_SIGMA × (best_rms / RMS_REFERENCE),
|
||||
/// SIGMA_MIN, SIGMA_MAX)`
|
||||
///
|
||||
/// Bilateral clamp per `pearl_symmetric_clamp_audit`.
|
||||
fn sigma_target_cpu(best_rms: f32) -> f32 {
|
||||
let ratio = best_rms / SHRINK_SIGMA_RMS_REFERENCE;
|
||||
let raw = SENTINEL_SHRINK_SIGMA * ratio;
|
||||
raw.max(SHRINK_SIGMA_MIN).min(SHRINK_SIGMA_MAX)
|
||||
}
|
||||
|
||||
/// Test 1: small drift → α high (near MAX, preserve hard).
|
||||
///
|
||||
/// drift = 0.02 (model converged within 2% of best) → α should be
|
||||
/// 1.0 − 0.05 = 0.95 (clamped at the floor of `1 − ALPHA_MAX`).
|
||||
#[test]
|
||||
fn small_drift_alpha_near_max() {
|
||||
let drift = 0.02_f32;
|
||||
let alpha = alpha_target_cpu(drift);
|
||||
// Bilateral clamp pins α at exactly SHRINK_ALPHA_MAX when drift is
|
||||
// below the lower bound `1 − ALPHA_MAX = 0.05`.
|
||||
assert!(
|
||||
(alpha - SHRINK_ALPHA_MAX).abs() < 1e-6,
|
||||
"small drift={} should give α=ALPHA_MAX={}, got {}",
|
||||
drift, SHRINK_ALPHA_MAX, alpha,
|
||||
);
|
||||
// Sanity: α is in the operating band [MIN, MAX].
|
||||
assert!(alpha >= SHRINK_ALPHA_MIN && alpha <= SHRINK_ALPHA_MAX);
|
||||
}
|
||||
|
||||
/// Test 2: large drift → α low (clamped at MIN).
|
||||
///
|
||||
/// drift = 0.50 (heavy divergence) → α should be 1.0 − 0.50 = 0.50,
|
||||
/// clamped exactly at SHRINK_ALPHA_MIN. drift > 0.50 stays clamped.
|
||||
#[test]
|
||||
fn large_drift_alpha_clamped_min() {
|
||||
// Exactly at the upper-clamp boundary.
|
||||
let drift_at = 1.0 - SHRINK_ALPHA_MIN;
|
||||
let alpha_at = alpha_target_cpu(drift_at);
|
||||
assert!(
|
||||
(alpha_at - SHRINK_ALPHA_MIN).abs() < 1e-6,
|
||||
"drift=1−MIN={} should give α=ALPHA_MIN={}, got {}",
|
||||
drift_at, SHRINK_ALPHA_MIN, alpha_at,
|
||||
);
|
||||
|
||||
// Beyond the upper-clamp boundary — still clamped.
|
||||
let drift_high = 0.80_f32;
|
||||
let alpha_high = alpha_target_cpu(drift_high);
|
||||
assert!(
|
||||
(alpha_high - SHRINK_ALPHA_MIN).abs() < 1e-6,
|
||||
"drift=0.80 should give α=ALPHA_MIN={}, got {} (bilateral clamp)",
|
||||
SHRINK_ALPHA_MIN, alpha_high,
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 3: medium drift → α in operating band (linear interpolation
|
||||
/// between MIN and MAX inside the unclamped region).
|
||||
///
|
||||
/// drift = 0.20 → α = 0.80 (matches the legacy hardcoded constant —
|
||||
/// equivalence point for the typical mid-fold drift signal).
|
||||
#[test]
|
||||
fn medium_drift_alpha_interpolates() {
|
||||
let drift = 0.20_f32;
|
||||
let alpha = alpha_target_cpu(drift);
|
||||
let expected = 1.0 - 0.20;
|
||||
assert!(
|
||||
(alpha - expected).abs() < 1e-6,
|
||||
"drift=0.20 should give α=0.80, got {}",
|
||||
alpha,
|
||||
);
|
||||
assert!(alpha > SHRINK_ALPHA_MIN);
|
||||
assert!(alpha < SHRINK_ALPHA_MAX);
|
||||
}
|
||||
|
||||
/// Test 4: σ scales with `||best||_RMS` — proportional in the
|
||||
/// unclamped region.
|
||||
///
|
||||
/// At RMS = REFERENCE = 0.5, σ = SENTINEL = 0.01 (bit-identical to
|
||||
/// the prior hardcoded value).
|
||||
/// At RMS = 2 × REFERENCE = 1.0, σ = 2 × SENTINEL = 0.02.
|
||||
#[test]
|
||||
fn sigma_scales_with_best_rms() {
|
||||
let sigma_at_ref = sigma_target_cpu(SHRINK_SIGMA_RMS_REFERENCE);
|
||||
assert!(
|
||||
(sigma_at_ref - SENTINEL_SHRINK_SIGMA).abs() < 1e-7,
|
||||
"RMS=REF should give σ=SENTINEL={}, got {}",
|
||||
SENTINEL_SHRINK_SIGMA, sigma_at_ref,
|
||||
);
|
||||
|
||||
let sigma_at_2x = sigma_target_cpu(2.0 * SHRINK_SIGMA_RMS_REFERENCE);
|
||||
let expected_2x = 2.0 * SENTINEL_SHRINK_SIGMA;
|
||||
assert!(
|
||||
(sigma_at_2x - expected_2x).abs() < 1e-7,
|
||||
"RMS=2×REF should give σ=2×SENTINEL={}, got {}",
|
||||
expected_2x, sigma_at_2x,
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 5: σ floor — RMS far below REFERENCE clamps σ at SIGMA_MIN.
|
||||
#[test]
|
||||
fn sigma_clamped_at_floor() {
|
||||
// RMS = 0 → ratio = 0 → raw σ = 0 → clamped at SIGMA_MIN.
|
||||
let sigma_zero = sigma_target_cpu(0.0);
|
||||
assert!(
|
||||
(sigma_zero - SHRINK_SIGMA_MIN).abs() < 1e-9,
|
||||
"RMS=0 should give σ=SIGMA_MIN={:e}, got {:e}",
|
||||
SHRINK_SIGMA_MIN, sigma_zero,
|
||||
);
|
||||
|
||||
// Tiny RMS → still clamped (ratio = 1e-3, raw = 1e-5 ≪ MIN=1e-4).
|
||||
let sigma_tiny = sigma_target_cpu(1.0e-3 * SHRINK_SIGMA_RMS_REFERENCE);
|
||||
assert!(
|
||||
sigma_tiny >= SHRINK_SIGMA_MIN,
|
||||
"tiny RMS should clamp σ ≥ SIGMA_MIN, got {:e}",
|
||||
sigma_tiny,
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 6: σ ceiling — RMS far above REFERENCE clamps σ at SIGMA_MAX.
|
||||
#[test]
|
||||
fn sigma_clamped_at_ceiling() {
|
||||
// RMS = 100 × REF → raw σ = 1.0 ≫ MAX=0.1 → clamped at MAX.
|
||||
let sigma_huge = sigma_target_cpu(100.0 * SHRINK_SIGMA_RMS_REFERENCE);
|
||||
assert!(
|
||||
(sigma_huge - SHRINK_SIGMA_MAX).abs() < 1e-7,
|
||||
"huge RMS should give σ=SIGMA_MAX={}, got {}",
|
||||
SHRINK_SIGMA_MAX, sigma_huge,
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 7: cold-start equivalence — at the (drift, RMS) point that
|
||||
/// corresponds to "model just converged near best with healthy weight
|
||||
/// magnitudes", α and σ should be in the vicinity of the prior
|
||||
/// hardcoded values 0.8 / 0.01. This pins the architectural intent of
|
||||
/// the formulas: typical operating point ≈ legacy constants, with
|
||||
/// signal-driven excursion when drift or weight scale departs from
|
||||
/// typical.
|
||||
#[test]
|
||||
fn typical_operating_point_near_legacy_constants() {
|
||||
// drift = 0.20 → α = 0.80 (exactly matches legacy SENTINEL).
|
||||
let alpha_typical = alpha_target_cpu(0.20);
|
||||
assert!(
|
||||
(alpha_typical - SENTINEL_SHRINK_ALPHA).abs() < 1e-6,
|
||||
"typical drift=0.20 should give α=SENTINEL={}, got {}",
|
||||
SENTINEL_SHRINK_ALPHA, alpha_typical,
|
||||
);
|
||||
|
||||
// RMS = REFERENCE → σ = SENTINEL (exactly matches legacy SENTINEL).
|
||||
let sigma_typical = sigma_target_cpu(SHRINK_SIGMA_RMS_REFERENCE);
|
||||
assert!(
|
||||
(sigma_typical - SENTINEL_SHRINK_SIGMA).abs() < 1e-7,
|
||||
"typical RMS=REF should give σ=SENTINEL={}, got {}",
|
||||
SENTINEL_SHRINK_SIGMA, sigma_typical,
|
||||
);
|
||||
}
|
||||
|
||||
/// Test 8: cold-start path (first fold, `best_params_snapshot=None`)
|
||||
/// — the host launcher writes `[0.0, 0.0, SENTINEL_SHRINK_ALPHA,
|
||||
/// SENTINEL_SHRINK_SIGMA]` to the mapped-pinned buffer and skips the
|
||||
/// kernel; ISV[505]/ISV[506] remain at FoldReset sentinels. The
|
||||
/// effective `shrink_and_perturb(α, σ)` call uses
|
||||
/// `α=SENTINEL_SHRINK_ALPHA=0.8` and `σ=SENTINEL_SHRINK_SIGMA=0.01` —
|
||||
/// bit-identical to the prior hardcoded constants. This test verifies
|
||||
/// the sentinel-value contract.
|
||||
#[test]
|
||||
fn cold_start_uses_legacy_constants_via_sentinel() {
|
||||
assert_eq!(
|
||||
SENTINEL_SHRINK_ALPHA, 0.8,
|
||||
"SENTINEL_SHRINK_ALPHA must equal prior hardcoded value 0.8"
|
||||
);
|
||||
assert_eq!(
|
||||
SENTINEL_SHRINK_SIGMA, 0.01,
|
||||
"SENTINEL_SHRINK_SIGMA must equal prior hardcoded value 0.01"
|
||||
);
|
||||
// The host-side defensive clamp in `reset_for_fold` (and the
|
||||
// bilateral kernel clamp) must accept the sentinel values without
|
||||
// modification — they are inside the operating band.
|
||||
let alpha_clamped = SENTINEL_SHRINK_ALPHA.clamp(SHRINK_ALPHA_MIN, SHRINK_ALPHA_MAX);
|
||||
let sigma_clamped = SENTINEL_SHRINK_SIGMA.clamp(SHRINK_SIGMA_MIN, SHRINK_SIGMA_MAX);
|
||||
assert_eq!(alpha_clamped, SENTINEL_SHRINK_ALPHA);
|
||||
assert_eq!(sigma_clamped, SENTINEL_SHRINK_SIGMA);
|
||||
}
|
||||
@@ -2,6 +2,45 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
## 2026-05-09 — SP18 v2 ISV-adaptive shrink-and-perturb at fold transition
|
||||
|
||||
Replaces the hardcoded `α=0.8 / σ=0.01` constants in `fused_training.rs::reset_for_fold`'s `shrink_and_perturb` call with ISV-driven adaptive bounds per `feedback_isv_for_adaptive_bounds` and `feedback_adaptive_not_tuned`. The driving signal is the relative weight drift `||current − best||₂ / ||best||₂` already computed each epoch by the SP18 v2 weight-drift diagnostic kernel (commit `aab13a83f`).
|
||||
|
||||
**Files added/modified:**
|
||||
|
||||
- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — 2 new ISV slot constants `SHRINK_ALPHA_ADAPTIVE_INDEX=505` + `SHRINK_SIGMA_ADAPTIVE_INDEX=506`; sentinels `SENTINEL_SHRINK_ALPHA=0.8` + `SENTINEL_SHRINK_SIGMA=0.01` (match prior hardcoded values for bit-identical cold-start); Category-1 dimensional safety bounds `[SHRINK_ALPHA_MIN=0.5, SHRINK_ALPHA_MAX=0.95]` + `[SHRINK_SIGMA_MIN=1e-4, SHRINK_SIGMA_MAX=0.1]`; producer reference scale `SHRINK_SIGMA_RMS_REFERENCE=0.5` (Category-1 dimensional anchor — typical mean RMS of healthy DQN weights). Lock test `sp18_shrink_perturb_slot_layout_locked`.
|
||||
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `ISV_TOTAL_DIM` 505 → 507 with full SP18 close-out commentary; `layout_fingerprint_seed` extended with `SLOT_505_SHRINK_ALPHA_ADAPTIVE` + `SLOT_506_SHRINK_SIGMA_ADAPTIVE`.
|
||||
- `crates/ml/src/cuda_pipeline/state_layout.cuh` — mirror `#define`s for the 2 new slots, sentinels, bounds, and `SHRINK_SIGMA_RMS_REFERENCE`.
|
||||
- `crates/ml/src/cuda_pipeline/weight_drift_diag_kernel.cu` — extended from 2-element output `[l2_diff, rel]` to 4-element `[l2_diff, rel, α_target, σ_target]`. Added 5th `isv` argument; same single-block × 256-thread launch + two block-tree-reductions. Thread-0 cold-path arithmetic computes `α_target = 1 − clamp(rel, 1−ALPHA_MAX, 1−ALPHA_MIN)` and `σ_target = clamp(SENTINEL_SIGMA × (best_rms / RMS_REF), SIGMA_MIN, SIGMA_MAX)` with bilateral clamp per `pearl_symmetric_clamp_audit`, writes to both the mapped-pinned output buffer AND `isv[505] / isv[506]` with `__threadfence_system()`. No new kernel — this is the sole producer of both diagnostic and controller outputs (single-launch atomicity per the spec's "Watch for: Drift diagnostic timing — have the drift kernel ALSO emit the adaptive α/σ as a 4-element output — fewer kernels, same atomicity").
|
||||
- `crates/ml/src/trainers/dqn/fused_training.rs` — `sp18_weight_drift_diag_buf` resized 2 → 4 floats; `launch_weight_drift_diag` plumbs `isv_signals_dev_ptr` as a new arg + cold-start path writes `[0.0, 0.0, SENTINEL_SHRINK_ALPHA, SENTINEL_SHRINK_SIGMA]` to the host-pinned slice (skip kernel; ISV slots remain at FoldReset sentinels — bit-identical to prior hardcoded behavior). New convenience reader `read_shrink_perturb_adaptive`. `reset_for_fold` migrated: replaces hardcoded `let sp_alpha = 0.8; let sp_sigma = 0.01;` with ISV reads via `read_isv_signal_at(SHRINK_ALPHA_ADAPTIVE_INDEX)` / `read_isv_signal_at(SHRINK_SIGMA_ADAPTIVE_INDEX)` + defensive host-side clamp at the same MIN/MAX bounds the kernel enforces (Category-1 dimensional safety). 4-arg `shrink_and_perturb(alpha, sigma)` signature unchanged — no ripple to the 3 other call sites in `training_loop.rs` (those continue to use `hyperparams.shrink_perturb_alpha/sigma` — different driving signal: unhealthy-streak rescue vs fold-transition plasticity refresh).
|
||||
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — 2 new `RegistryEntry` records (`sp18_shrink_alpha_adaptive`, `sp18_shrink_sigma_adaptive`) with `FoldReset` category. Sentinels match prior hardcoded values for cold-start bit-identical equivalence. `sp18_fold_reset_entries_present` lock test bumped to assert 24 SP18 entries (was 22).
|
||||
- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — 2 dispatch arms in `reset_named_state` writing the sentinels at fold boundary (so the cold-start fallback in `launch_weight_drift_diag` reads the sentinel rather than stale state). Per-epoch HEALTH_DIAG `weight_drift` line extended: `[norm_l2={} relative={} branch_max={} alpha_adaptive={} sigma_adaptive={}]` — surfaces the per-epoch α/σ targets the next fold-transition `shrink_and_perturb` will use.
|
||||
- `crates/ml/tests/sp18_shrink_perturb_adaptive_test.rs` (new) — 8 CPU-only oracle tests pinning the math contract: small drift → α near MAX, large drift → α clamped MIN, medium drift → linear interpolation, σ scales with `||best||_RMS`, σ floor/ceiling clamps, typical operating point ≡ legacy constants, cold-start = sentinel = legacy.
|
||||
|
||||
**Atomicity envelope** (per `feedback_no_partial_refactor` and `feedback_wire_everything_up`): ISV slots + producer + consumer + HEALTH_DIAG + tests all land in the same commit. The producer kernel's outputs are wired to two consumers in this commit — the existing diagnostic readback (`read_weight_drift_diag` returns `[l2, rel]` from `out[0..2]`) and the new fold-transition controller (`reset_for_fold` reads `ISV[505] / ISV[506]`).
|
||||
|
||||
**Cold-start handling** (per `pearl_first_observation_bootstrap`): on the first fold or any fold where `best_params_snapshot` is `None` (no Sharpe improvement yet), the host launcher writes zeros + sentinels to the mapped-pinned buffer and skips the kernel. ISV slots stay at FoldReset sentinels (`0.8 / 0.01`) — `reset_for_fold`'s `shrink_and_perturb(0.8, 0.01)` call is bit-identical to the prior hardcoded behavior.
|
||||
|
||||
**Producer cadence**: HEALTH_DIAG (per-epoch). The drift kernel runs at `read_weight_drift_diag()` inside the per-epoch HEALTH_DIAG block, which sync's the stream — host visibility of `ISV[505]/ISV[506]` is guaranteed before the next fold-transition. `reset_for_fold` reads the values written by the LAST epoch's drift launch — i.e., drift OBSERVED at end-of-fold N, applied to fold N+1's perturbation (per the spec).
|
||||
|
||||
**Concerns surfaced** (DONE_WITH_CONCERNS):
|
||||
|
||||
1. **3 other `shrink_and_perturb` call sites unchanged**: `training_loop.rs:1096` (Phase 3 boundary), `training_loop.rs:3549` (D3/N3 health-triggered), `training_loop.rs:3609` (D6/N6 ensemble-collapse-triggered), and `training_loop.rs:7320` (PerturbationStrategy::ShrinkPerturbBranches with hardcoded `0.9, 0.1`). These all read from `self.hyperparams.shrink_perturb_alpha/sigma` (or hardcoded `0.9, 0.1` for backtracking). They're a different driving-signal regime — health-driven rescue interventions, not fold-transition plasticity refresh. Migrating those to ISV-adaptive is a separate concern (different signal: health collapse + ensemble disagreement, not weight drift) — surfaced for follow-up but explicitly out of scope for this commit per `feedback_no_partial_refactor` (the contract change is the fold-transition site only; the others have their own driving signal).
|
||||
|
||||
2. **Cold-start drift = 0.0 → α = 1.0?**: Per the spec's "Watch for" section, when drift = 0 and we're not in cold-start, `α_target = 1.0 − clamp(0, 0.05, 0.5) = 1.0 − 0.05 = 0.95` — clamped at MAX, not 1.0. The bilateral clamp's lower bound `1 − ALPHA_MAX = 0.05` ensures we never get α=1.0 (full preservation, no perturbation). The cold-start case (no `best_params_snapshot`) is handled via the host-side bypass writing sentinels — drift never appears as 0.0 from the kernel side because the kernel doesn't run in cold-start.
|
||||
|
||||
**Verification**:
|
||||
- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace` — clean.
|
||||
- `cargo test -p ml --lib sp18` — 4 SP18 lock tests pass: `all_sp18_slots_fit_within_isv_total_dim`, `sp18_combined_slot_layout_locked`, `sp18_shrink_perturb_slot_layout_locked`, `sp18_fold_reset_entries_present`.
|
||||
- `cargo test -p ml --lib every_fold_and_soft_reset_entry_has_dispatch_arm` — pass (24 SP18 dispatch arms verified).
|
||||
- `cargo test -p ml --test sp18_shrink_perturb_adaptive_test` — 8/8 pass (CPU oracle math contract).
|
||||
- `cargo test -p ml --test sp18_fold_transition_test` — 3/3 pass (existing fold-transition restore-best test still validates).
|
||||
|
||||
**HEALTH_DIAG line format** (extended):
|
||||
```
|
||||
HEALTH_DIAG[N]: weight_drift [norm_l2=0.123456 relative=0.234567 branch_max=0.234567 alpha_adaptive=0.7654 sigma_adaptive=0.012345]
|
||||
```
|
||||
|
||||
## 2026-05-08 — SP18 v2 Phase 0 Task 0.2: B-leg V_SHARE trajectory + TD-error magnitude diagnostic
|
||||
|
||||
Phase 0 observability: kernel + Rust launcher + V_SHARE history ring buffer + per-epoch HEALTH_DIAG emit + GPU oracle tests. All in the same atomic commit per `feedback_no_partial_refactor`.
|
||||
|
||||
Reference in New Issue
Block a user