From 410ab6b0ea3bead59f97cd2afd51b00e03b322b1 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 18 May 2026 19:32:03 +0200 Subject: [PATCH] =?UTF-8?q?arch(ml-alpha):=20ISV-driven=20=CF=83=20+=20ada?= =?UTF-8?q?ptive=20Z=5FSCALE=20=E2=80=94=20both=20controllers=20anchor=20o?= =?UTF-8?q?n=20loss=5Fema?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- crates/ml-alpha/cuda/horizon_lambda.cu | 129 +++++++++++++--------- crates/ml-alpha/src/trainer/perception.rs | 65 ++++++----- 2 files changed, 112 insertions(+), 82 deletions(-) diff --git a/crates/ml-alpha/cuda/horizon_lambda.cu b/crates/ml-alpha/cuda/horizon_lambda.cu index ca5aa5057..c9f49c4ba 100644 --- a/crates/ml-alpha/cuda/horizon_lambda.cu +++ b/crates/ml-alpha/cuda/horizon_lambda.cu @@ -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); } } diff --git a/crates/ml-alpha/src/trainer/perception.rs b/crates/ml-alpha/src/trainer/perception.rs index 0b36eb2f2..b624c13de 100644 --- a/crates/ml-alpha/src/trainer/perception.rs +++ b/crates/ml-alpha/src/trainer/perception.rs @@ -375,15 +375,15 @@ pub struct PerceptionTrainer { attn_bwd_fn: CudaFunction, _attn_module: Arc, - // ── 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, // [N_HORIZONS] - pub grad_log_sigma_h_d: CudaSlice, // [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, // [N_HORIZONS] — ISV output + pub grad_log_sigma_h_d: CudaSlice, // [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, /// 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, + /// 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, /// 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::(N_HORIZONS)?; let grad_log_sigma_h_d = stream.alloc_zeros::(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::(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,