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

arch(ml-alpha): ISV-driven σ + adaptive Z_SCALE — both controllers anchor on loss_ema

Browse Source 
Per pearl_controller_anchors_isv_driven, every controller anchor/target/
cap derives from a tracked signal, not hardcoded constants. The σ-only
revert kept Kendall σ as a free Adam-learned scalar — that violated ISV
discipline and fought Adam's m/√v normalization
(pearl_adam_normalizes_loss_weights).

Single source of truth for both per-horizon controllers:

  log_sigma_h[h]  ← max(log(0.5), 0.5 * log(loss_ema[h]))
                    Kendall equilibrium (∂L/∂log σ = 0 ⟹ σ_h² = mean_bce_h)
                    in closed form. Asymmetric floor at log(0.5) prevents
                    collapse. No Adam state, no gradient delay.

  lambda[h]       ← clamp(1.0, 2.0, 1.0 + Z_SCALE_ISV * z_h)
                    Z_SCALE_ISV = (LAMBDA_CEILING - LAMBDA_FLOOR) / z_max_ema
                    Adaptive scale auto-uses the full clamp envelope:
                    the historical-max-z horizon maps exactly to
                    LAMBDA_CEILING. Replaces hardcoded Z_SCALE=0.5 which
                    rarely engaged on real data (max observed λ ~1.04).

Both anchor on the same ISV (loss_ema). z_max_ema is a new single-scalar
EMA state tracking max |z| across horizons, with first-obs bootstrap.

Removes:
  - opt_log_sigma AdamW optimizer (σ no longer learned)
  - grad_log_sigma_h_d memset (BCE kernel writes; output ignored — kept
    only to preserve BCE kernel signature)

Kernel signature change (horizon_ema_and_lambda):
  +z_max_ema [1]       (read+write EMA state)
  +log_sigma_h [5]     (closed-form output, overwrite)

Discipline:
  - First-obs bootstrap (sentinel <= 0) per pearl_first_observation_bootstrap
  - Permanent floor (max(real, floor)) per pearl_blend_formulas_must_have_permanent_floor
  - Asymmetric clamp per pearl_audit_unboundedness_for_implicit_asymmetry
  - Z-score normalisation per pearl_zscore_normalization_for_magnitude_asymmetric_signals
  - No nvrtc, no atomicAdd, no host branches in graph capture

All 9 perception_overfit tests pass — including
horizon_ema_and_lambda_track_after_training which validates the kernel
end-to-end through 64 K-loop iterations of capture/replay.

Submit local smoke; cluster A/B vs σ-only baseline (0.7506/0.7519) and
vs Phase 1+2+3 (0.7749/0.7591) follows once the perf-only 3-fold A/B
confirms no regression at b23f8f2ef.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-18 19:32:03 +02:00
parent b23f8f2efa
commit 410ab6b0ea
2 changed files with 112 additions and 82 deletions
Download Patch File Download Diff File Expand all files Collapse all files

129
crates/ml-alpha/cuda/horizon_lambda.cu
View File

@@ -1,81 +1,68 @@
// horizon_lambda.cu — ISV-driven per-horizon gradient scaler.
// horizon_lambda.cu — ISV-driven per-horizon gradient scaler AND closed-form Kendall σ.
//
// Step 1: maintain an EMA of the UNWEIGHTED per-horizon BCE loss that
// the BCE kernel emits each training step into
// `loss_per_horizon[N_HORIZONS]`.
// Step 2: convert the EMA into a per-horizon multiplicative lambda
// used by the backward path to scale how strongly each
// horizon influences the shared trunk gradient.
// Single source of truth for the two per-horizon controllers:
//
// Why ISV: the current static `auto-horizon-weights` formula
// (`min(1, K/h)`) is a closed-form heuristic that ignores actual
// per-horizon learning difficulty. Empirically mhzs7 spent most of
// training over-weighting short horizons (whose label correlation
// within the K-snapshot window dominates the gradient signal) while
// h6000 — the deployment-relevant multi-minute horizon — stayed at
// AUC≈0.69. Tracking per-horizon BCE directly lets the lambda boost
// the horizons that the model is currently failing to learn, without
// hand-tuned constants.
// 1) lambda[h] ∈ [LAMBDA_FLOOR, LAMBDA_CEILING] — multiplier on the
// trunk grad_h contribution in heads_grn_bwd. Boosts horizons
// that the model is currently failing to learn (per their
// unweighted BCE EMA).
//
// Why not just lift BCE coefficients: per
// `pearl_adam_normalizes_loss_weights.md`, Adam's m/sqrt(v) cancels
// per-loss weight lifts (SP13: 13× aux_w produced only 0.6%/epoch
// divergence). The effective lever is to scale the GRADIENT into the
// shared trunk, not the loss aggregate. heads_bwd will multiply the
// per-horizon `d_z` contribution by lambda[h] before accumulating
// into `grad_h`, bypassing Adam normalization.
// 2) log_sigma_h[h] — Kendall σ for the BCE forward kernel. Computed
// in closed form from the same loss_ema signal (Kendall's
// gradient equilibrium ⟹ σ_h = sqrt(mean_bce_h)). Replaces
// Adam-learned σ, eliminating the m/√v normalization conflict
// (pearl_adam_normalizes_loss_weights).
//
// First-observation bootstrap: loss_ema is zero-initialised; the
// kernel detects `prev <= 0` and replaces (rather than blends) on
// the first step. After that it uses a fixed α — Wiener-optimal α is
// a Phase 3 follow-up; for now a conservative 0.1 keeps the EMA
// stable across training noise.
// ISV anchors per `pearl_controller_anchors_isv_driven`: every anchor/
// target/cap derives from a tracked signal, never from hardcoded
// constants outside the bootstrap epsilons:
//
// Lambda safety: ASYMMETRIC clamp `[1.0, LAMBDA_CEILING]` per
// `pearl_audit_unboundedness_for_implicit_asymmetry.md` — boost-only.
// Lambda derivation: Z-SCORE NORMALISED, not raw ratio. The ratio
// formula `ema_h / mean(ema)` produced lambdas in 0.97-1.03 in our
// data because per-horizon BCE clusters tightly (range 0.04 abs)
// while the mean is ~0.65 — so the controller barely engaged
// (asymmetric ceiling 2.0 was never approached, max observed
// lambda ~1.04). Z-score `z_h = (ema_h - mean) / std(ema)` makes
// the spread scale-invariant per
// `pearl_zscore_normalization_for_magnitude_asymmetric_signals.md`,
// then `lambda = 1.0 + Z_SCALE * z` fills the clamp envelope: a
// horizon 1σ above mean gets lambda 1.5 (with Z_SCALE=0.5), 2σ
// saturates at the ceiling. Boost-only floor at 1.0 still applies.
// loss_ema[h] — EMA(unweighted BCE per horizon)
// z_max_ema — EMA(max |z_h|) across horizons → drives
// adaptive Z_SCALE so the OBSERVED-max-z
// horizon maps exactly to LAMBDA_CEILING.
// With the prior hardcoded Z_SCALE=0.5 the
// controller rarely engaged on real data
// (max observed lambda ~1.04); adaptive
// Z_SCALE uses the full envelope.
//
// Sentinel bootstrap (`prev <= 0`) per
// `pearl_first_observation_bootstrap.md`; permanent floor
// (`max(real, floor)`) per `pearl_blend_formulas_must_have_permanent_floor.md`;
// asymmetric clamp (floor-only on σ, floor + ceiling on λ) per
// `pearl_audit_unboundedness_for_implicit_asymmetry.md`. Z-score
// normalization per `pearl_zscore_normalization_for_magnitude_asymmetric_signals.md`.
#define N_HORIZONS_LAMBDA 5
#define ALPHA_FIXED 0.1f
#define LAMBDA_FLOOR 1.0f // boost-only: never demote a horizon below uniform
#define LAMBDA_FLOOR 1.0f
#define LAMBDA_CEILING 2.0f
#define Z_SCALE 0.5f // 1σ above mean → lambda 1.5; 2σ → ceiling
#define LOG_SIGMA_FLOOR (-0.6931472f) // log(0.5) — σ never collapses below 0.5
#define Z_MAX_FLOOR 0.1f // guards Z_SCALE_ISV when z_max_ema is tiny
extern "C" __global__ void horizon_ema_and_lambda(
const float* __restrict__ loss_per_horizon, // [5] — current step UNWEIGHTED BCE
float* __restrict__ loss_ema, // [5] — EMA state (read + write)
float* __restrict__ lambda // [5] — output multiplier
float* __restrict__ z_max_ema, // [1] — EMA(max|z|) state (read + write)
float* __restrict__ lambda, // [5] — output: λ_h
float* __restrict__ log_sigma_h // [5] — output: closed-form Kendall σ
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
// -------- 1) EMA update on loss_per_horizon, with first-obs bootstrap.
float new_ema[N_HORIZONS_LAMBDA];
float sum = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS_LAMBDA; ++h) {
const float cur = loss_per_horizon[h];
const float prev = loss_ema[h];
// Sentinel = 0 ⇒ first observation replaces directly.
// `prev <= 0` is safer than `== 0` under --use_fast_math.
new_ema[h] = (prev <= 0.0f) ? cur : (prev + ALPHA_FIXED * (cur - prev));
loss_ema[h] = new_ema[h];
sum += new_ema[h];
}
const float mean = sum / (float)N_HORIZONS_LAMBDA;
// Per-horizon std (population, not sample — N=5 is small + fixed).
// EPSILON guards the first-step case where all EMAs are equal
// (z would be 0/0); under that condition lambda falls back to 1.0
// via the asymmetric floor.
// -------- 2) Z-score normalize across horizons.
float ssq = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS_LAMBDA; ++h) {
@@ -84,10 +71,44 @@ extern "C" __global__ void horizon_ema_and_lambda(
}
const float std = sqrtf(ssq / (float)N_HORIZONS_LAMBDA + 1e-12f);
const float inv_std = (std > 1e-6f) ? (1.0f / std) : 0.0f;
float z[N_HORIZONS_LAMBDA];
float z_max = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS_LAMBDA; ++h) {
const float z = (new_ema[h] - mean) * inv_std;
const float raw = 1.0f + Z_SCALE * z;
lambda[h] = fmaxf(LAMBDA_FLOOR, fminf(LAMBDA_CEILING, raw));
z[h] = (new_ema[h] - mean) * inv_std;
const float az = fabsf(z[h]);
if (az > z_max) z_max = az;
}
// -------- 3) z_max_ema: tracks the typical spread, drives adaptive Z_SCALE.
// Sentinel = 0 ⇒ first-obs bootstrap. Fixed α matches loss_ema's
// smoothing horizon.
const float prev_zmax = z_max_ema[0];
const float new_zmax = (prev_zmax <= 0.0f)
? z_max
: (prev_zmax + ALPHA_FIXED * (z_max - prev_zmax));
z_max_ema[0] = new_zmax;
// -------- 4) Adaptive Z_SCALE: maps the historical-max-z to LAMBDA_CEILING.
// When z_max_ema is large (spread regime), Z_SCALE_ISV shrinks
// to avoid saturating early; when small (tight regime), it
// grows to amplify the controller into engagement.
const float z_anchor = fmaxf(new_zmax, Z_MAX_FLOOR);
const float z_scale_isv = (LAMBDA_CEILING - LAMBDA_FLOOR) / z_anchor;
// -------- 5) Per-horizon outputs.
#pragma unroll
for (int h = 0; h < N_HORIZONS_LAMBDA; ++h) {
// λ_h: boost-only, asymmetric clamped.
const float raw_lambda = 1.0f + z_scale_isv * z[h];
lambda[h] = fmaxf(LAMBDA_FLOOR, fminf(LAMBDA_CEILING, raw_lambda));
// log σ_h: Kendall equilibrium from the same loss_ema signal.
// ∂L/∂log σ = 0 ⟹ σ_h² = mean_bce_h ⟹ log σ_h = ½·log(loss_ema_h).
// Floor at log(0.5) so σ can't collapse to 0 (which would explode w_h).
const float ema_h = fmaxf(new_ema[h], 1e-12f);
const float raw_log_sigma = 0.5f * logf(ema_h);
log_sigma_h[h] = fmaxf(LOG_SIGMA_FLOOR, raw_log_sigma);
}
}

65
crates/ml-alpha/src/trainer/perception.rs
View File

@@ -375,15 +375,15 @@ pub struct PerceptionTrainer {
attn_bwd_fn: CudaFunction,
_attn_module: Arc<CudaModule>,
// ── Kendall σ-weighted BCE (axis A) ──
// Per-horizon learnable `log σ_h` scalars. Weighted BCE per horizon
// is `(1/(2·exp(2·log σ_h))) · mean_bce_h + log σ_h`. AdamW updates
// log_sigma_h at 1/4 the head LR (slow update). σ on its own without
// C/D/E is the keeper from the v2 sweep (see
// project_ml_alpha_v2_ab_verdict.md).
pub log_sigma_h_d: CudaSlice<f32>, // [N_HORIZONS]
pub grad_log_sigma_h_d: CudaSlice<f32>, // [N_HORIZONS]
pub opt_log_sigma: AdamW,
// ── Kendall σ-weighted BCE (axis A, ISV-driven) ──
// Per-horizon `log σ_h` is computed in closed form by the
// horizon_ema_and_lambda kernel from loss_ema (Kendall equilibrium
// σ_h = sqrt(mean_bce_h)). No Adam state — replaces the prior
// Adam-learned variant which fought m/√v normalization
// (pearl_adam_normalizes_loss_weights). Single source of truth
// shared with λ_h: both anchor on loss_ema.
pub log_sigma_h_d: CudaSlice<f32>, // [N_HORIZONS] — ISV output
pub grad_log_sigma_h_d: CudaSlice<f32>, // [N_HORIZONS] — BCE kernel writes here; ignored downstream
// ── K-loop parallelization (Phase B) ──
// Per-batch grad scratch buffers for cfc_step_backward_batched.
@@ -422,11 +422,14 @@ pub struct PerceptionTrainer {
/// `pearl_first_observation_bootstrap.md`.
loss_ema_d: CudaSlice<f32>,
/// Per-horizon multiplicative gradient scaler. Updated each step
/// by horizon_ema_and_lambda. Clamped to [0.5, 2.0]. Wired into
/// heads_bwd in a follow-up commit; until then this buffer is
/// computed but unused — keeps the EMA infrastructure runnable
/// and validatable without changing training math.
/// by horizon_ema_and_lambda. Asymmetric clamp [1.0, 2.0].
/// Consumed by heads_grn_bwd as `s_lambda` multiplier on the
/// trunk grad_h contribution.
lambda_d: CudaSlice<f32>,
/// EMA(max |z_h|) state for adaptive Z_SCALE in horizon_ema_and_lambda.
/// Single scalar; tracks the typical inter-horizon spread so
/// Z_SCALE_ISV maps the observed max-z to LAMBDA_CEILING.
z_max_ema_d: CudaSlice<f32>,
/// Cached function handle for the horizon_ema_and_lambda kernel.
horizon_lambda_fn: CudaFunction,
/// Module that owns `horizon_lambda_fn`; kept alive so the function
@@ -819,7 +822,9 @@ impl PerceptionTrainer {
// updates at 1/4 the head LR per the spec.
let log_sigma_h_d = stream.alloc_zeros::<f32>(N_HORIZONS)?;
let grad_log_sigma_h_d = stream.alloc_zeros::<f32>(N_HORIZONS)?;
let opt_log_sigma = AdamW::new(dev, N_HORIZONS, cfg.lr_cfc * 0.25)?;
// ISV-driven: z_max_ema tracks EMA(max |z_h|) across horizons to set
// adaptive Z_SCALE for λ. Sentinel = 0 ⇒ first-obs bootstrap.
let z_max_ema_d = stream.alloc_zeros::<f32>(1)?;
// V1 (2026-05-18): per-horizon Q_h trainer state (C24/C25) removed.
// A/B sweep at commit 83546b5c3 falsified the approach (mean_auc
@@ -886,7 +891,7 @@ impl PerceptionTrainer {
_attn_module: attn_module,
log_sigma_h_d,
grad_log_sigma_h_d,
opt_log_sigma,
z_max_ema_d,
// Phase B: cfc per-batch grad scratch + reducer.
cfc_grad_w_in_scratch_d,
cfc_grad_w_rec_scratch_d,
@@ -1678,9 +1683,8 @@ impl PerceptionTrainer {
}
drop((_g_hpk, _g_probs, _g_z1, _g_a1, _g_z2, _g_gate, _g_main, _g_logit));
// ── 4.5. Zero σ grad scratch each step (capture-safe).
self.stream.memset_zeros(&mut self.grad_log_sigma_h_d)
.map_err(|e| anyhow::anyhow!("zero grad_log_sigma_h: {e}"))?;
// σ grad scratch: BCE kernel writes overwrite-style; no zeroing needed.
// The output is unused downstream (σ is ISV-driven via horizon_ema_and_lambda).
// ── 5. Fused multi-horizon BCE (Kendall σ-weighted, axis A).
// The kernel doesn't distinguish position-vs-batch
@@ -1710,13 +1714,15 @@ impl PerceptionTrainer {
// ── 5a. (v2 reserved) — per-horizon Q_h backward removed at V1.
// v2 backward path will live here in commit V10.
// ── 5b. ISV-driven per-horizon EMA + lambda. Updates the EMA
// of unweighted per-horizon BCE and emits a clamped
// multiplier `lambda_d[h]` per
// `pearl_adam_normalizes_loss_weights.md` strategy of
// scaling effective gradient instead of loss weight.
// Currently captured into the graph; lambda_d will be
// consumed by heads_bwd in the next commit.
// ── 5b. ISV-driven per-horizon controllers. Single kernel emits
// both:
// lambda_d[h] — λ multiplier on trunk grad_h
// log_sigma_h_d[h] — Kendall σ for next-step BCE loss
// Both anchor on loss_ema (EMA of unweighted per-horizon
// BCE). σ in closed form (Kendall equilibrium), λ as
// z-score-scaled multiplier with adaptive Z_SCALE_ISV
// (derived from z_max_ema). No Adam involvement in σ
// (per pearl_adam_normalizes_loss_weights).
{
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
@@ -1727,7 +1733,9 @@ impl PerceptionTrainer {
launch
.arg(&self.loss_per_horizon_d)
.arg(&mut self.loss_ema_d)
.arg(&mut self.lambda_d);
.arg(&mut self.z_max_ema_d)
.arg(&mut self.lambda_d)
.arg(&mut self.log_sigma_h_d);
unsafe { launch.launch(cfg).context("horizon_ema_and_lambda launch")?; }
}
@@ -2224,8 +2232,9 @@ impl PerceptionTrainer {
self.opt_vsn_w.step(&mut self.vsn_w_d, &self.grad_vsn_w_d)?;
self.opt_vsn_b.step(&mut self.vsn_b_d, &self.grad_vsn_b_d)?;
self.opt_attn_q.step(&mut self.attn_q_d, &self.grad_attn_q_d)?;
// Kendall σ step (axis A; 1/4 head LR).
self.opt_log_sigma.step(&mut self.log_sigma_h_d, &self.grad_log_sigma_h_d)?;
// Kendall σ is ISV-driven (closed form from loss_ema in
// horizon_ema_and_lambda) — no Adam step. grad_log_sigma_h_d
// is still written by the BCE kernel but intentionally ignored.
// (v2 reserved) — per-horizon AdamW step removed at V1; v2's six
// optimizer groups (horizon_tokens, Q_inv, w_fuse + b_fuse,