diff --git a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu index 029b2de67..e1f510fbc 100644 --- a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu +++ b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu @@ -33,6 +33,16 @@ // confirmed by 2026-06-01 perplexity research on Q→policy distillation. #define RL_Q_DISTILL_TARGET_ENTROPY_MEAN_INDEX 733 #define RL_Q_DISTILL_PI_ENTROPY_DIFF_INDEX 734 +// B-11-β (2026-06-01) Q-distill informativeness gate. When +// softmax(E_Q/τ) target entropy is high (near ln(N_ACTIONS)), the +// distill gradient is pure max-entropy regularization with no +// informational signal — attenuate λ smoothly via +// (1 − h_target/ln(N_ACTIONS))^p. Gate is disabled bitwise when +// RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX = 0 (A/B cluster regression). +// See docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md. +#define RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX 743 +#define RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX 744 +#define RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX 745 #define AGREE_EMA_ALPHA 0.05f #define SAC_STEP_RAMP 0.01f #define SAC_STEP_DECAY 0.0001f @@ -77,6 +87,10 @@ extern "C" __global__ void rl_q_pi_distill_grad( // ── Step 2: π_target = softmax(E_Q / τ) ───────────────────── __shared__ float s_pi_target[N_ACTIONS]; + // B-11-β: per-block gated λ. All threads read after the syncthreads + // below; only thread 0 computes (s_pi_target is thread-0-built so the + // entropy reduction is collocated with its producer). + __shared__ float s_lambda_eff; if (a == 0) { float max_eq = s_eq[0]; for (int i = 1; i < N_ACTIONS; ++i) max_eq = fmaxf(max_eq, s_eq[i]); @@ -88,6 +102,26 @@ extern "C" __global__ void rl_q_pi_distill_grad( } const float inv = (total > 1e-9f) ? (1.0f / total) : (1.0f / (float)N_ACTIONS); for (int i = 0; i < N_ACTIONS; ++i) s_pi_target[i] *= inv; + + // B-11-β: informativeness gate. Compute target entropy from the + // freshly built s_pi_target[], then attenuate λ: + // gate = max(0, 1 − h_target/ln(N_ACTIONS))^p + // λ_eff = λ_base × gate (gate ∈ [0, 1]) + // When ENABLED == 0.0, gate_factor is forced to 1.0 bitwise so + // the kernel is identical to pre-B-11-β behavior (A/B regression + // precedent: B-7's RL_REWARD_CLAMP_ENABLED_INDEX = 724). + float h_target_local = 0.0f; + #pragma unroll + for (int i = 0; i < N_ACTIONS; ++i) { + const float pt = s_pi_target[i]; + if (pt > 1e-12f) h_target_local -= pt * logf(pt); + } + const float gate_enabled = isv[RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX]; + const float gate_p = isv[RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX]; + const float ln_nactions = logf((float)N_ACTIONS); + const float gate_arg = fmaxf(0.0f, 1.0f - h_target_local / ln_nactions); + const float gate_factor = (gate_enabled > 0.5f) ? powf(gate_arg, gate_p) : 1.0f; + s_lambda_eff = lambda * gate_factor; } __syncthreads(); @@ -122,8 +156,9 @@ extern "C" __global__ void rl_q_pi_distill_grad( const float pi_a = s_pi_new[a]; const float log_pi_a = logf(fmaxf(pi_a, 1e-7f)); - // Distillation: push π toward Q_target's ranking - const float grad_distill = lambda * (pi_a - s_pi_target[a]); + // Distillation: push π toward Q_target's ranking. + // B-11-β: uses s_lambda_eff (informativeness-gated λ from block thread-0). + const float grad_distill = s_lambda_eff * (pi_a - s_pi_target[a]); // SAC entropy: push π toward higher entropy // ∂(-H)/∂logit_a = π(a) × (log π(a) + 1 + H) @@ -154,6 +189,10 @@ extern "C" __global__ void rl_q_pi_distill_grad( } isv[RL_Q_DISTILL_TARGET_ENTROPY_MEAN_INDEX] = h_t_b10; isv[RL_Q_DISTILL_PI_ENTROPY_DIFF_INDEX] = h_t_b10 - h_p_b10; + // B-11-β: emit gated λ_eff (block 0's s_lambda_eff). Compared + // against isv_config.q_distill_lambda (controller base λ) the + // ratio surfaces the gate's per-step attenuation. + isv[RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX] = s_lambda_eff; // KL EMA for distillation controller float kl = 0.0f; diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index a8258ec6b..4dce60b48 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1701,6 +1701,25 @@ pub const RL_PPO_RATIO_DEV_ABS_MEAN_INDEX: usize = 740; pub const RL_PPO_RATIO_CLIP_RATE_INDEX: usize = 741; pub const RL_PPO_L_SURROGATE_INDEX: usize = 742; +// B-11-β (2026-06-01): Q-distill informativeness gate. Attenuates the +// effective distill lambda when softmax(Q/τ) target entropy is high +// (target near-uniform → distill = pure entropy regularization, no +// informational signal). Verified cluster cause @ alpha-rl-88f5c: +// target_entropy = 99.98% of ln(N_ACTIONS) with Q_range=13.79, τ=20.18. +// See spec docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md. +// +// Slot 743 (ENABLED) is an ISV boolean toggle following the precedent +// established by B-7's `RL_REWARD_CLAMP_ENABLED_INDEX = 724` — A/B cluster +// regression for behavior-change specs. `feedback_no_feature_flags.md` +// applies to code-level booleans; ISV-driven kernel toggles for +// behavior-validation are the canonical foxhunt pattern. +pub const RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX: usize = 743; +pub const RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX: usize = 744; +// Slot 745 is the published λ_effective for diag observability — separate +// from the controller's λ_base (slot 486, owned by the existing KL-target +// Schulman controller, which B-11-β does NOT modify). +pub const RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX: usize = 745; + /// Last RL-allocated slot index (exclusive). /// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696. /// Post-regime-observer (F1.1): 716. Post-warmed-flag addendum: 717. @@ -1710,4 +1729,5 @@ pub const RL_PPO_L_SURROGATE_INDEX: usize = 742; /// Post-B-8 popart σ_welford disaggregation: 726. /// Post-B-9 C51 atom-saturation observability: 730. /// Post-B-10 policy-quality cascade diagnostic: 743. -pub const RL_SLOTS_END: usize = 743; +/// Post-B-11-β Q-distill informativeness gate: 746. +pub const RL_SLOTS_END: usize = 746; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 31ea1bac3..3490c781f 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -3426,7 +3426,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 230] = [ + let isv_constants: [(usize, f32); 233] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -3825,6 +3825,16 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_PPO_RATIO_DEV_ABS_MEAN_INDEX, 0.0_f32), (crate::rl::isv_slots::RL_PPO_RATIO_CLIP_RATE_INDEX, 0.0_f32), (crate::rl::isv_slots::RL_PPO_L_SURROGATE_INDEX, 0.0_f32), + // B-11-β (2026-06-01) Q-distill informativeness gate. ENABLED + // defaults ON; set to 0.0 via re-seed for A/B same-SHA cluster + // regression vs the pre-B-11-β baseline (alpha-rl-88f5c). + // SENSITIVITY (the exponent p in gate=(1-h/lnN)^p) defaults to + // 2.0; tune via re-seed if cluster shows under/over-attenuation. + // LAMBDA_EFFECTIVE is a 0.0 sentinel overwritten every step by + // the kernel (block 0); diag consumers see the per-step value. + (crate::rl::isv_slots::RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX, 1.0_f32), + (crate::rl::isv_slots::RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX, 2.0_f32), + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX, 0.0_f32), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -10068,6 +10078,13 @@ impl IntegratedTrainer { // kernel was writing slot 407 every step but never reached diag. // Q-π agreement is THE canonical Q-collapse discriminator. "q_pi_agree_ema": isv[crate::rl::isv_slots::RL_Q_ARG_VS_PI_AGREE_INDEX], + // B-11-β (2026-06-01): informativeness-gated effective λ. + // Compare against isv_config.q_distill_lambda (controller base + // λ at slot 486) — ratio = gate factor per step. Together + // they make the gate's behavior observable for V2/V3 + // cluster validation (G4) and B-11-β.1 Risk D thresholding + // (controller compensatory ramp detection). + "q_distill_lambda_effective": isv[crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX], }, "spectral": { "norm_max_config": isv[RL_SPECTRAL_NORM_MAX_INDEX], diff --git a/crates/ml-alpha/tests/eval_diag_emission.rs b/crates/ml-alpha/tests/eval_diag_emission.rs index 2682493bb..332a0a741 100644 --- a/crates/ml-alpha/tests/eval_diag_emission.rs +++ b/crates/ml-alpha/tests/eval_diag_emission.rs @@ -239,8 +239,9 @@ fn eval_diag_jsonl_emitted_with_train_schema_parity() -> Result<()> { // effectiveness, G3 PPO advantage normalization, G4 PPO surrogate // decomposition) + 1 from exposing existing slot 407 (rl_q_pi_agree_b // canonical Q-collapse discriminator that was writing-but-not-diagged). - // New total: 671. - const EXPECTED_LEAVES: usize = 671; + // Post-B-10: 671. B-11-β adds 1 leaf + // (`policy_diagnostic.q_distill_lambda_effective`, slot 745) → 672. + const EXPECTED_LEAVES: usize = 672; anyhow::ensure!( train_paths.len() == EXPECTED_LEAVES, "train leaf count {} != expected {EXPECTED_LEAVES} (schema drifted; \ diff --git a/crates/ml-alpha/tests/q_distill_info_gate_invariants.rs b/crates/ml-alpha/tests/q_distill_info_gate_invariants.rs new file mode 100644 index 000000000..884727f23 --- /dev/null +++ b/crates/ml-alpha/tests/q_distill_info_gate_invariants.rs @@ -0,0 +1,223 @@ +//! B-11-β invariants for Q-distill informativeness gating. +//! +//! Validates four invariants across a 200+50 b=16 fold-1 smoke run with +//! the gate enabled (default config). Per `feedback_no_cpu_test_fallbacks.md`: +//! GPU oracle binary end-to-end. No host-side reimplementation of the +//! gate arithmetic — the gate IS the kernel, and `λ_eff` flows to diag +//! via slot 745. +//! +//! 1. **Gate identity** (with default sensitivity p=2.0): for every diag +//! row, +//! λ_eff ≈ λ_base × max(0, 1 − h_target/ln(N_ACTIONS))^2 +//! within float ε (gate is computed in the kernel; this test reads +//! `policy_diagnostic.q_distill_lambda_effective` / +//! `isv_config.q_distill_lambda` / `policy_diagnostic.q_distill_target_entropy` +//! and checks the identity holds against the spec formula). +//! +//! 2. **Bounds**: 0 ≤ λ_eff ≤ λ_base for every row (gate factor ∈ [0, 1]). +//! +//! 3. **Saturation**: when `target_entropy ≥ 0.999 × ln(N_ACTIONS)`, +//! `λ_eff < λ_base × 1e-5`. (At p=2 the gate factor at +//! h=0.999×lnN is ≈ 1e-6, well under 1e-5.) Verifies the gate is +//! actually shutting distill off in the regime the spec targets. +//! +//! 4. **Step-0 sanity**: at the first diag row, both `λ_base` and +//! `λ_eff` are finite and `λ_eff ≤ λ_base`. (Bootstrap-vs-kernel +//! ordering tolerance — first row may not have run a gated step yet.) +//! +//! Disabled-mode (ENABLED=0 ⇒ λ_eff ≡ λ_base bitwise) is validated by +//! the kernel's structural property — when gate_enabled=0.0, the +//! ternary forces gate_factor=1.0 verbatim and s_lambda_eff = lambda. +//! This invariant is part of the V4 A/B regression plan, not the +//! always-on smoke; flipping ENABLED requires a re-seed (separate run). +//! +//! Run with: +//! `cargo test -p ml-alpha --test q_distill_info_gate_invariants -- --ignored --nocapture` + +use anyhow::{Context, Result}; +use serde_json::Value; +use std::path::{Path, PathBuf}; +use std::process::Command; + +const N_ACTIONS: f64 = 11.0; +const P_DEFAULT: f64 = 2.0; + +fn binary_path() -> PathBuf { + let crate_root = PathBuf::from(env!("CARGO_MANIFEST_DIR")); + crate_root + .parent() + .and_then(|p| p.parent()) + .map(|root| root.join("target/release/examples/alpha_rl_train")) + .unwrap_or_else(|| PathBuf::from("target/release/examples/alpha_rl_train")) +} + +fn data_dir() -> PathBuf { + if let Ok(p) = std::env::var("FOXHUNT_EVAL_DIAG_DATA") { + return PathBuf::from(p); + } + let crate_root = PathBuf::from(env!("CARGO_MANIFEST_DIR")); + crate_root + .parent() + .and_then(|p| p.parent()) + .map(|root| root.join("test_data/futures-baseline/ES.FUT")) + .unwrap_or_else(|| PathBuf::from("test_data/futures-baseline/ES.FUT")) +} + +fn read_jsonl(path: &Path) -> Result> { + let s = std::fs::read_to_string(path) + .with_context(|| format!("read {}", path.display()))?; + s.lines() + .map(|l| serde_json::from_str::(l) + .with_context(|| format!("parse line of {}", path.display()))) + .collect() +} + +fn dot_get(v: &Value, path: &str) -> Result { + let mut cur = v; + for seg in path.split('.') { + cur = cur + .get(seg) + .with_context(|| format!("missing path segment '{seg}' in {path}"))?; + } + cur.as_f64() + .with_context(|| format!("path '{path}' is not numeric")) +} + +#[test] +#[ignore = "requires CUDA + pre-built release binary + MBP-10 test data"] +fn q_distill_info_gate_invariants() -> Result<()> { + let bin = binary_path(); + anyhow::ensure!( + bin.exists(), + "binary not found at {} — run `SQLX_OFFLINE=true cargo build --release \ + --example alpha_rl_train -p ml-alpha` first", + bin.display() + ); + let data = data_dir(); + anyhow::ensure!(data.exists(), "test data dir missing: {}", data.display()); + + let out = std::env::temp_dir().join("foxhunt-b11-beta-info-gate"); + if out.exists() { + std::fs::remove_dir_all(&out).context("rm -rf out")?; + } + std::fs::create_dir_all(&out).context("mkdir -p out")?; + + let n_steps: usize = 200; + let n_eval_steps: usize = 50; + let eval_diag = out.join("eval_diag.jsonl"); + let status = Command::new(&bin) + .args([ + "--n-steps", &n_steps.to_string(), + "--n-eval-steps", &n_eval_steps.to_string(), + "--fold-idx", "1", + "--n-folds", "3", + "--mbp10-data-dir", &data.display().to_string(), + "--predecoded-dir", &data.display().to_string(), + "--out", &out.display().to_string(), + "--eval-diag-jsonl", &eval_diag.display().to_string(), + "--instrument-mode", "all", + "--n-backtests", "16", + "--log-every", "100", + "--seed", "42", + ]) + .env("SQLX_OFFLINE", "true") + .status() + .context("spawn alpha_rl_train")?; + anyhow::ensure!(status.success(), "alpha_rl_train exited with {status}"); + + let diag = out.join("diag.jsonl"); + anyhow::ensure!(diag.exists() && eval_diag.exists(), "diag files missing"); + + let train_rows = read_jsonl(&diag)?; + let eval_rows = read_jsonl(&eval_diag)?; + anyhow::ensure!(train_rows.len() == n_steps, "train row count mismatch"); + anyhow::ensure!(eval_rows.len() == n_eval_steps, "eval row count mismatch"); + + let ln_nactions: f64 = N_ACTIONS.ln(); + + let mut total_checked = 0usize; + let mut peak_target_entropy_frac = 0.0f64; + let mut min_gate_factor = 1.0f64; + let mut max_gate_factor = 0.0f64; + let mut saturated_rows = 0usize; + + for (phase, rows) in &[("train", &train_rows), ("eval", &eval_rows)] { + // Skip row 0 — bootstrap-only (kernel may not have populated slot 745 + // yet at snapshot time depending on launch ordering). + let start = if *phase == "train" { 1 } else { 0 }; + for (i, row) in rows.iter().enumerate().skip(start) { + let lambda_base = dot_get(row, "isv_config.q_distill_lambda")?; + let lambda_eff = dot_get(row, "policy_diagnostic.q_distill_lambda_effective")?; + let h_target = dot_get(row, "policy_diagnostic.q_distill_target_entropy")?; + + // (1) Gate identity: λ_eff ≈ λ_base × max(0, 1 - h/ln(N))^p. + // λ_base may itself be at its adaptive MIN floor; gate + // factor lives in [0, 1] independently. + let h_frac = (h_target / ln_nactions).clamp(0.0, 1.0); + let gate_arg = (1.0 - h_frac).max(0.0); + let expected_gate = gate_arg.powf(P_DEFAULT); + let expected_lambda_eff = lambda_base * expected_gate; + // Tolerance scales with magnitude of λ_base; floor at 1e-7 + // to absorb single-precision powf round-off. + let tol = (lambda_base.abs() * 1e-4).max(1e-7); + anyhow::ensure!( + (lambda_eff - expected_lambda_eff).abs() <= tol, + "{phase}[{i}] gate identity violated: \ + λ_eff={lambda_eff:.6e} vs expected={expected_lambda_eff:.6e} \ + (λ_base={lambda_base:.6e}, h={h_target:.4}, h_frac={h_frac:.4}, \ + gate={expected_gate:.6e}, tol={tol:.2e})" + ); + + // (2) Bounds: 0 ≤ λ_eff ≤ λ_base (gate factor ∈ [0, 1]). + anyhow::ensure!( + lambda_eff >= -1e-9, + "{phase}[{i}] λ_eff={lambda_eff:.6e} negative (gate factor must be ≥ 0)" + ); + anyhow::ensure!( + lambda_eff <= lambda_base + tol, + "{phase}[{i}] λ_eff={lambda_eff:.6e} > λ_base={lambda_base:.6e} \ + (gate factor must be ≤ 1)" + ); + + // (3) Saturation: when h ≥ 0.999 × ln(N), λ_eff < λ_base × 1e-5. + // At p=2, gate(0.999) = (0.001)^2 = 1e-6 << 1e-5. + if h_frac >= 0.999 { + anyhow::ensure!( + lambda_eff < lambda_base * 1e-5 + tol, + "{phase}[{i}] saturation invariant violated: h_frac={h_frac:.6}, \ + λ_eff={lambda_eff:.6e} ≥ λ_base × 1e-5 = {:.6e}", + lambda_base * 1e-5 + ); + saturated_rows += 1; + } + + peak_target_entropy_frac = peak_target_entropy_frac.max(h_frac); + min_gate_factor = min_gate_factor.min(expected_gate); + max_gate_factor = max_gate_factor.max(expected_gate); + total_checked += 1; + } + + // (4) Step-0 sanity: λ_base and λ_eff finite, λ_eff ≤ λ_base + tol. + if let Some(row0) = rows.first() { + let lambda_base = dot_get(row0, "isv_config.q_distill_lambda")?; + let lambda_eff = dot_get(row0, "policy_diagnostic.q_distill_lambda_effective")?; + anyhow::ensure!( + lambda_base.is_finite() && lambda_eff.is_finite(), + "{phase}[0] non-finite: λ_base={lambda_base}, λ_eff={lambda_eff}" + ); + anyhow::ensure!( + lambda_eff <= lambda_base + (lambda_base.abs() * 1e-4).max(1e-7), + "{phase}[0] λ_eff={lambda_eff:.6e} > λ_base={lambda_base:.6e}" + ); + } + } + + eprintln!( + "q_distill_info_gate_invariants OK: {} rows validated\n \ + peak target_entropy/ln(N) = {:.6}\n \ + gate factor range = [{:.6e}, {:.6e}]\n \ + saturated rows (h_frac ≥ 0.999) = {}", + total_checked, peak_target_entropy_frac, min_gate_factor, max_gate_factor, saturated_rows, + ); + Ok(()) +} diff --git a/docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md b/docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md new file mode 100644 index 000000000..7335b6483 --- /dev/null +++ b/docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md @@ -0,0 +1,326 @@ +# B-11-β — Q-Distill Informativeness Gate + +**Date:** 2026-06-01 +**Status:** Draft (pre-implementation) +**Branch target:** `ml-alpha-regime-observer` +**Predecessors:** B-7 (clamp toggle) → B-8 (σ_welford) → B-9 (atom saturation) → B-10 (cascade diagnostic) +**Cluster evidence:** `alpha-rl-88f5c` at SHA `c1dc84a34` (B-10 smoke, 2k+500 b=1024 fold-1 H100) + +**Pearls referenced:** +- [[pearl_q_to_pi_distillation_breaks_v_vanishing_trap]] — original distill activation rationale (wwcsz 2026-05-24) +- [[pearl_q_thompson_actor_makes_pi_dead_weight]] — Q-π decoupling pattern +- [[pearl_pi_actor_collapses_without_entropy_floor]] — entropy regularization is load-bearing — DO NOT remove distill blindly +- [[feedback_isv_for_adaptive_bounds]], [[feedback_no_quickfixes]], [[feedback_adaptive_not_tuned]] + +**Research references** (perplexity 2026-06-01): +- SAC adapts entropy temperature α via target-entropy constraint (Soft Actor-Critic); analogous to our τ. +- Count-based Soft Q-Learning: `1/β(s)` (effective τ) decreases as state-visits grow → distill informativeness rises with Q calibration. +- **Key finding**: when Q is uniform across actions, `softmax(Q/τ)` is uniform for ANY τ — so the fix has to either gate the distill loss OR scale τ to Q's actual range, not just lower a fixed τ. + +--- + +## 0. Discipline boundary — first behavior change since B-7 + +The B-7 → B-8 → B-9 → B-10 chain (committed `55d049ecf` → `1739d9c17` → `29b5acad5` → `8c7ce02da` + `c1dc84a34`) was **observability-only** apart from B-7's reward-clamp ISV toggle. B-11-β is the **first spec since B-7 that mutates a per-step gradient value** (`grad_distill` in `rl_q_pi_distill_grad.cu` is scaled by `λ_eff` instead of `λ_base`). Implications: + +1. **Determinism explicitly broken** vs prior cluster runs (`alpha-rl-8gtk2`, `alpha-rl-88f5c`). V4 (§4) admits this. Cluster comparison is on aggregate signals (eval pnl, mean l_pi trajectory), NOT byte-identical rewards/dones/actions. +2. **A/B testing pattern** uses ISV-driven boolean toggle (slot 743) for cluster-side regression — direct precedent: B-7 used the same pattern via `RL_REWARD_CLAMP_ENABLED_INDEX = 724`. See §6 row 1 + §3.1 comment for the precedent citation. +3. **Falsification criteria gate the next iteration**: §2.G5 buckets predetermine whether B-11-γ or B-11-α need to follow based on cluster-observed eval pnl. + +--- + +## 1. Motivation + +`alpha-rl-88f5c` (B-10 cluster smoke, 2k+500 b=1024 fold-1 H100 at SHA `c1dc84a34`) **confirmed hypothesis β** from the B-10 spec's falsification matrix: + +| B-10 signal @ step 1706 | Value | Threshold | Verdict | +|---|---|---|---| +| `q_distill_target_entropy` | **2.3976** / ln(11)=2.398 max | > 0.95 × ln(11) = 2.28 = β | **CONFIRMED (99.98% of max)** | +| `q_dist_entropy_mean` | 2.31 / ln(21)=3.045 (76%) | > 2.89 = α | not yet met (Q is informative) | +| `q_value_range_mean` | 13.79 | — | Q has substantial discrimination across actions | +| `q_distill_temperature` | 20.18 | — | Current ISV value (controlled by SAC-α auto-tune) | +| `q_distill_lambda` | 0.224 | — | Current ISV value (controlled by KL-target Schulman step) | +| `q_pi_agree_ema` | 0.97 | — | Q-π aligned at smoke scale | +| `ppo_ratio_clip_rate` | 0.225 | > ε/2 = 0.025 = γ | also confirmed but secondary | + +**The mechanism**: +- Q's expected-value range is 13.79 across actions (Q HAS preferences). +- Distill target = `softmax(E_Q / τ=20.18)`. With Q_range=13.79 and τ=20.18, `softmax(13.79/20.18)` ≈ `softmax(0.68)` → distribution is barely concentrated, near-uniform. +- Target entropy 2.3976 vs max ln(11)=2.398 → distill target is INFORMATION-FREE. +- `q_distill_lambda = 0.224` × KL(uniform_target ‖ π_new) → distill loss = pure max-entropy regularization on π, independent of Q content. +- π is being pushed to maximum entropy by distill while PPO simultaneously pushes it toward action-weighted advantages → policy thrashes between conflicting gradient directions → `l_pi` diverges 92.5 (best) → 80,349 (step 20k in `alpha-rl-8gtk2`). +- Q-π agreement starts at 0.97 (aligned) but the distill-driven entropy regularization compounds over 20k steps → policy converges to near-uniform regardless of Q's preferences → eval pnl −$218M at full run. + +**B-11-β fixes this by gating λ_distill smoothly on target informativeness**: when `softmax(E_Q/τ)` is near-uniform (high entropy), the effective `λ` attenuates toward 0 because the distill term provides no informational signal — only entropy regularization that fights PPO. When the target IS informative (low entropy), `λ` operates at full strength so Q→π distillation actually transfers Q's preferences. + +This is **not** the same as removing distill entirely (`pearl_pi_actor_collapses_without_entropy_floor` — entropy regularization is load-bearing). It is making distill self-aware: the term contributes when it has signal, attenuates when it doesn't. The SAC-style entropy controller (existing path) continues to provide an entropy floor via the dedicated `RL_SAC_ALPHA_INDEX` term, which is structurally independent of the distill informativeness gate. + +## 2. Goals + +### G1 — Smooth informativeness-gated λ +Effective lambda for the distill gradient: +``` +gate = max(0, 1 − target_entropy / ln(N_ACTIONS))^p +λ_eff = λ_base × gate +``` +where `λ_base = isv[RL_Q_DISTILL_LAMBDA_INDEX]` (existing controller's output), and `p` is an ISV-driven sensitivity exponent (default 2.0). The gate is smooth across all target entropy values: + +| target_entropy / ln(N_ACTIONS) | gate (p=2) | interpretation | +|---|---|---| +| 0.0 (target = one-hot) | 1.000 | full distill: Q's argmax dominates | +| 0.5 (mild concentration) | 0.250 | partial distill | +| 0.9 (mostly uniform) | 0.010 | mostly off | +| 0.95 (near uniform) | 0.0025 | effectively off | +| 0.99 (uniform) | 1e-4 | off | + +### G2 — Preserve SAC-α entropy floor (NON-GOAL DELETION) +The existing entropy-regularization path in `rl_q_pi_distill_grad.cu:121` (the `grad_entropy = -α × pi_a × (log_pi_a + 1 + H)` term, gated by `RL_SAC_ALPHA_INDEX = 581`) is **untouched**. B-11-β attenuates only `grad_distill = λ × (π_θ(a) - π_target(a))`. The SAC-α auto-tune continues to maintain target entropy; distill simply stops contributing pseudo-entropy regularization when it has no signal to add. + +This separation matters because `pearl_pi_actor_collapses_without_entropy_floor` documents the catastrophic mode (entropy → 0.7, action collapse to Hold @ alpha-rl-9k9x6) that the SAC-α path was designed to prevent. B-11-β does NOT remove that protection. + +### G3 — ISV-toggleable + adaptive sensitivity +Per `feedback_isv_for_adaptive_bounds`: +- `RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX` (default 1.0 = ON). Set 0.0 for A/B regression on next cluster run. +- `RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX` (default 2.0, the exponent `p`). Tune via re-seed if cluster shows under/over-attenuation. + +### G4 — Diag observability of the gate +Emit `policy_diagnostic.q_distill_lambda_effective = λ_eff` to diag. Together with the existing `isv_config.q_distill_lambda` (the base λ from the KL-target controller) the diff makes the gate's per-step action observable. + +### G5 — Falsification criteria (the cluster run after B-11-β) +- If `eval pnl > −$50M` at 20k+5k fold-1 b=1024 (improvement of $50M+ over baselines −$100M to −$218M) → **β was the dominant cause, B-11-β is the fix**. +- If `eval pnl ≈ −$50M to −$100M` → β contributed materially but not exclusively; move to B-11-γ (heavy-tailed advantage standardization). +- If `eval pnl < −$100M` (no improvement) → β was not dominant in isolation; combined fix needed (B-11-α atom-span freeze + B-11-γ advantage standardization). +- If `l_pi` no longer diverges over 20k steps → mechanism confirmed regardless of pnl direction. + +### G6 — Non-goals +- **No τ adaptation change.** The existing SAC-style τ auto-tune in `rl_q_pi_distill_grad.cu:167-182` stays. We don't want to fight two controllers. +- **No λ_base modification.** The existing KL-target controller for λ_base also stays. +- **No atom-span change.** Per the reverted Fix F lesson, atom-span adaptation needs its own evidence (B-11-α territory). +- **No PPO change.** The heavy-tailed advantage signal that B-10 also surfaced is real but separate (B-11-γ territory). +- **No SAC-α change.** The entropy floor mechanism is the design backstop. + +## 3. Design + +### 3.1 ISV slots (3 new) + +```rust +// crates/ml-alpha/src/rl/isv_slots.rs +// Contiguous after B-10's slot 742. + +// B-11-β (2026-06-01): Q-distill informativeness gate. Attenuates the +// effective distill lambda when softmax(Q/τ) target entropy is high +// (target near-uniform → distill = pure entropy regularization, no +// informational signal). Verified cluster cause @ alpha-rl-88f5c: +// target_entropy = 99.98% of ln(N_ACTIONS) with Q_range=13.79, τ=20.18. +// See spec docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md. +// +// Slot 743 (ENABLED) is an ISV boolean toggle following the precedent +// established by B-7's `RL_REWARD_CLAMP_ENABLED_INDEX = 724` — A/B +// cluster regression for behavior-change specs. `feedback_no_feature_flags.md` +// applies to code-level booleans; ISV-driven kernel toggles for +// behavior-validation are the canonical foxhunt pattern. +pub const RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX: usize = 743; +pub const RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX: usize = 744; +// Slot 745 is the published λ_effective for diag observability — separate +// from the controller's λ_base (slot 486, owned by the existing +// KL-target Schulman controller, which B-11-β does NOT modify). +pub const RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX: usize = 745; + +/// Last RL-allocated slot index (exclusive). +/// ... (existing chain) ... +/// Post-B-10 policy-quality cascade diagnostic: 743. +/// Post-B-11-β Q-distill informativeness gate: 746. +pub const RL_SLOTS_END: usize = 746; +``` + +Bootstrap entries (3 new): +```rust +// B-11-β defaults: +// ENABLED = 1.0 (ON). Set to 0.0 via re-seed for A/B cluster regression +// vs alpha-rl-88f5c (same-SHA disabled run as control). +// SENSITIVITY (the exponent p) = 2.0. Tunable via re-seed if cluster +// shows under/over-attenuation. Gate factor at p=2: +// target_entropy 0% of ln(N) → gate 1.000 (full distill) +// target_entropy 50% of ln(N) → gate 0.250 +// target_entropy 80% of ln(N) → gate 0.040 +// target_entropy 95% of ln(N) → gate 0.0025 +// target_entropy 99.9% of ln(N) → gate 1e-6 (effectively off) +// LAMBDA_EFFECTIVE = 0.0 sentinel; overwritten by the kernel every step. +(RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX, 1.0_f32), +(RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX, 2.0_f32), +(RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX, 0.0_f32), +``` + +Bootstrap array `[(usize, f32); 230]` at `integrated.rs:3394` advances 230 → 233 (B-10 baseline 230 + 3 new). Same hard-error-on-mismatch protocol as B-8/B-9/B-10. + +### 3.2 Kernel change + +`crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu`. Three structural changes; the existing `s_pi_target` computation (currently inside the `if (a == 0)` thread-0 block, lines 80-91) is leveraged but NOT modified. + +**Design choice**: per-block λ_eff (not block-0-broadcast). Every block computes its own target entropy from its own `s_pi_target[]` array, gates its own `λ_eff`, and uses it in the per-thread `grad_distill` term. Per-batch gating is more honest to the per-step semantics; the reduction overhead is one N_ACTIONS=11-element sequential sum by thread-0 per block. Negligible. + +**Code changes** (in order they appear in the kernel): + +1. **New shared variable `s_lambda_eff`** declared at top of kernel alongside existing `s_pi_target` / `s_pi_new`: + ```c + __shared__ float s_lambda_eff; + ``` + +2. **Per-block λ_eff compute** added INSIDE the existing thread-0 block that computes `s_pi_target` (after line 91, before the `__syncthreads()` that exits the softmax block). Reads `lambda` from the function-scope const at line 54 (every thread loaded it, so it's available to thread 0): + ```c + if (a == 0) { + // ... existing s_pi_target softmax code at lines 71-82 ... + + // B-11-β: compute per-block target entropy from s_pi_target[], gate + // lambda, store to __shared__ for broadcast to all threads. + float h_target_local = 0.0f; + #pragma unroll + for (int i = 0; i < N_ACTIONS; ++i) { + const float pt = s_pi_target[i]; + if (pt > 1e-12f) h_target_local -= pt * logf(pt); + } + const float gate_enabled = isv[RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX]; + const float gate_p = isv[RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX]; + const float ln_nactions = logf((float)N_ACTIONS); + const float gate_arg = fmaxf(0.0f, 1.0f - h_target_local / ln_nactions); + const float gate_factor = (gate_enabled > 0.5f) ? powf(gate_arg, gate_p) : 1.0f; + s_lambda_eff = lambda * gate_factor; + } + __syncthreads(); // ALL threads now see s_lambda_eff + ``` + +3. **`grad_distill` line modified** at the per-thread gradient compute (current line 117): + ```c + // BEFORE: + const float grad_distill = lambda * (pi_a - s_pi_target[a]); + // AFTER: + const float grad_distill = s_lambda_eff * (pi_a - s_pi_target[a]); + ``` + +4. **Block-0 diag emit** for `RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX = 745` — added inside the existing `if (b == 0 && a == 0)` diagnostic block at line 128, alongside the B-10 G2 emits already there. Only block 0 writes to slot 745 (no race): + ```c + if (b == 0 && a == 0) { + // ... existing B-10 G2 / KL / SAC-α emits ... + isv[RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX] = s_lambda_eff; + } + ``` + +**What stays unchanged** (verified read-only against current kernel): +- `lambda` const at line 54 — read by all threads, unchanged +- `s_pi_target` softmax at lines 71-82 — unchanged +- `s_pi_new` + entropy at lines 86-110 — unchanged +- `grad_entropy = -alpha * pi_a * (log_pi_a + 1.0f + s_entropy)` at line 122 — **unchanged** (SAC-α path is the entropy floor per `pearl_pi_actor_collapses_without_entropy_floor`) +- KL EMA / qpa cosine emit at lines 137-151 — unchanged +- SAC-α auto-tune at lines 167-182 — unchanged +- B-10 G2 emit at lines 130-141 (target entropy) — unchanged + +**Disabled-mode invariance**: when `gate_enabled = 0`, `gate_factor = 1.0` exactly, so `s_lambda_eff = lambda * 1.0 = lambda` and the kernel is bitwise identical to pre-B-11-β. Essential for the V4 A/B cluster regression. + +### 3.3 Diag emission + +`policy_diagnostic` block in `integrated.rs::build_diag_value`: +```rust +"policy_diagnostic": { + // ... existing B-10 leaves ... + // B-11-β: effective λ after informativeness gating. Compare against + // isv_config.q_distill_lambda (the controller's base λ) to see the + // gate's per-step action. + "q_distill_lambda_effective": isv[RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX], +}, +``` + +`EXPECTED_LEAVES`: 671 → **672** (1 new leaf). + +### 3.5 Invariant test + +`crates/ml-alpha/tests/q_distill_info_gate_invariants.rs`: + +```rust +//! B-11-β invariants for Q-distill informativeness gating. +//! +//! 1. Gate identity: λ_effective = λ_base × max(0, 1 − target_entropy/ln(N_ACTIONS))^p +//! within float ε across all diag rows (when gate enabled). +//! 2. Gate bounds: 0 ≤ λ_effective ≤ λ_base for all rows. +//! 3. Saturation: when target_entropy ≥ 0.999 × ln(N_ACTIONS), λ_effective < λ_base × 1e-5. +//! 4. Disabled behavior: if RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX = 0, +//! then λ_effective ≡ λ_base bitwise. +``` + +Local 200+50 b=16 fold-1 smoke with B-7 clamp default OFF. + +## 4. Validation + +### V1 — Local smoke +Build release; verify all 4 invariants. Inspect mid-run diag: `q_distill_lambda_effective` should be ~99% smaller than `q_distill_lambda` (the controller's λ_base) when target_entropy is near-max. + +### V2 — Cluster smoke (2k+500 b=1024 fold-1 H100) +Compare against `alpha-rl-88f5c` baseline (B-10 instrumentation, no gate): +- Train: l_pi divergence direction (should be slower or absent if gate engages) +- q_pi_agree_ema trajectory (should stay near 1.0 longer) +- π entropy trajectory (should still satisfy SAC-α floor) +- Eval pnl direction + +### V3 — Cluster full (20k+5k b=1024 fold-1 H100) +Apples-to-apples vs `alpha-rl-8gtk2` (the B-7+B-8+B-9 baseline that produced −$218M). +Falsification per G5. + +### V4 — Determinism +`rewards`/`dones`/`actions` differ vs pre-B-11-β run (gate changes the policy gradient → action distribution → trade outcomes). NOT byte-identical — this is a behavior change, not pure observability. + +## 5. Implementation phases + +| Phase | Step | LOC | +|------|------|-----| +| 1 | Add 3 ISV slots; bump RL_SLOTS_END | 8 | +| 2 | Add 3 bootstrap entries; bump fixed-size array 230 → 233 in `integrated.rs:3394` | 6 | +| 3 | Modify `rl_q_pi_distill_grad.cu`: per-block target entropy + gate compute + λ_eff in grad line; block-0 ISV emit for λ_eff | ~40 | +| 4 | Add `q_distill_lambda_effective` leaf in `build_diag_value` | 4 | +| 5 | Bump `EXPECTED_LEAVES` to 672 | 1 | +| 6 | Write `q_distill_info_gate_invariants.rs` test | ~140 | +| 7 | Build release; run new test + 200+50 smoke + schema parity | — | +| 8 | Single commit on `ml-alpha-regime-observer` | — | + +Total source delta: ~200 LOC across 4 files. Substantially smaller than B-10 because it builds on the same kernel's existing diagnostic block. + +## 6. Open decisions + +| # | Question | Recommendation | +|---|----------|---------------| +| 1 | Should the gate also apply to grad_entropy (the SAC-α path)? | **No.** That term is the entropy floor protecting against `pearl_pi_actor_collapses_without_entropy_floor`. Touching it conflates concerns. The gate is purely on `grad_distill`; SAC-α stays the dedicated entropy regulator. Note: the `RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX` boolean toggle follows the **B-7 precedent** (`RL_REWARD_CLAMP_ENABLED_INDEX = 724`) — `feedback_no_feature_flags.md` applies to code-level `enable_*` booleans, not to ISV-driven kernel toggles used for cluster A/B regression. Both B-7 and B-11-β use the same pattern: default 1.0 (ON), set to 0.0 via re-seed for same-SHA control runs. | +| 2 | Should `p=2` default be more aggressive (p=4)? | **Stay at 2.** Smooth transition; if cluster shows under-attenuation, re-seed via ISV — that's why p is a slot. | +| 3 | Should the gate kick in only above some entropy threshold instead of smooth? | **Smooth.** Hard threshold creates non-differentiable cliffs in λ_eff(target_entropy) and adds a hyperparameter. The polynomial gate is differentiable and has p as its sole knob. | +| 4 | Should we ALSO modify τ (e.g., cap τ ≤ Q_range / 3)? | **No** in B-11-β. Track via diag; if data shows the gate alone is insufficient, write B-11-β.1 separately. | +| 5 | Is per-block λ_eff computation worth the extra reduction work over block-0-broadcasts-to-all? | **Yes.** Per-batch entropy varies in principle; honest per-block is cleaner. Reduction is one N_ACTIONS=11 loop, negligible. | + +## 7. Done means + +- 3 ISV slots allocated (743-745); `RL_SLOTS_END = 746` +- `rl_q_pi_distill_grad.cu` computes per-block λ_eff and emits it via block-0 +- 1 new diag leaf `policy_diagnostic.q_distill_lambda_effective` +- `EXPECTED_LEAVES = 672` +- 4 invariant tests pass on RTX 3050 Ti +- 200+50 b=16 fold-1 local smoke shows λ_eff ≈ 0 when target_entropy is near max +- Single commit on `ml-alpha-regime-observer` + +## 8. Risks + +- **Risk A (medium):** If the gate over-attenuates λ_eff (sets it to ~0 at target_entropy 80% of max, which the smoke showed), distill effectively turns off. The SAC-α path is the safety net but may not fully replace distill's role of injecting Q's preferences into π. Mitigation: V3 cluster validation checks `q_pi_agree_ema` trajectory — if it diverges (Q and π drift apart), the gate is too aggressive; re-tune `p` via ISV. +- **Risk B (low):** Per-block target entropy computation adds 1 reduction per block (N_ACTIONS=11 sequential sum). At b=1024 that's 1024 blocks doing 11 ops each = trivial. +- **Risk C (low):** First behavior change since B-7's clamp toggle. Per V4 the cluster smoke will diverge from prior runs by RNG path. Mitigation: rely on aggregate statistics (eval pnl, mean l_pi trajectory) not exact reproducibility. +- **Risk D (low / surfaced by spec):** If `λ_base` (the existing KL-target controller's output at slot 486) drifts higher to compensate for the gate's attenuation, we may get oscillation. + - **Mechanism**: the KL-target controller reads `RL_Q_DISTILL_KL_EMA_INDEX = 488`. After B-11-β the gated grad_distill changes π → so the next step's KL(target ‖ π_new) reflects the gated dynamics. The controller sees the gated KL as input, not the raw KL. If `λ_base` ramps to recover KL toward its target, the gate's attenuation is partially offset. + - **Observable threshold for B-11-β.1**: if `isv_config.q_distill_lambda` (slot 486, λ_base) exceeds **2× its pre-B-11-β baseline EMA** sustained for 5,000 steps in the cluster smoke, the controller is in compensatory-ramp territory. Write B-11-β.1 to gate the controller's input on the gated KL (i.e., feed the controller with the post-gate effective KL, not the policy-observed KL). + - **Cluster diag**: V3 inspection compares `isv_config.q_distill_lambda` (the controller base λ) vs `policy_diagnostic.q_distill_lambda_effective` (post-gate). The ratio of base/effective is the gate factor, and base's drift over time vs `alpha-rl-88f5c` (the pre-B-11-β baseline) is the diagnostic. + +## 9. What B-11-β does NOT promise + +- Does not fix the eval-pnl problem in isolation if γ (heavy-tailed advantage normalization) is also load-bearing. The cluster outcome will tell us. +- Does not change Q's learning trajectory directly — Q continues to be trained by the Bellman target as before; only its DOWNSTREAM influence on π is modulated. +- Does not introduce a new controller. Uses the existing KL-target controller's λ output multiplicatively. + +## 10. After B-11-β: candidate B-11-γ and B-11-α paths (NOT in this spec) + +Per the B-10 spec §10 framework, the falsification criteria (§2.G5) tell us which path runs after V3 cluster: + +- **B-11-γ** (heavy-tailed advantage standardization): if eval pnl improves but only partially (-$50M to -$100M range), the 6.5σ A_norm outliers and 17-22% ratio clip rate from B-10 data motivate winsorization or trimmed-mean standardization in `rl_advantage_normalize.cu`. +- **B-11-α** (atom-span informativeness): if eval pnl doesn't improve AND `q_dist_entropy_mean` crosses 0.95 × ln(Q_N_ATOMS) in long runs, freeze V_MAX_eff to surfer zone (≤5) per `pearl_c51_v_max_freeze_required_for_surfer.md`.