From 49cdf90ecc49cff421a91d95827a6e7fe8e4b071 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 5 May 2026 19:14:48 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp14):=20B.4=20=E2=80=94=20alpha=5Fgrad=5F?= =?UTF-8?q?compute=5Fkernel=20(EGF=20heart)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Single-thread state-machine kernel that is the heart of the Earned Gradient Flow pearl. Reads driver signals from the global ISV bus, runs Schmitt-trigger Gate 1, computes adaptive k_aux/k_q/β, evaluates two sigmoids, multiplies with a host-supplied warmup gate, applies a β-rate-limiter, and writes 7 outputs back to ISV. Per-step pipeline: 1. Read aux_dir_acc (slot 373), q_disagreement (slot 383), Welford variance EMAs (388, 389, 390), persistent Schmitt state (391), alpha_smoothed_prev (393). 2. Compute adaptive k_aux = K_BASE_AUX/(1 + var_aux/VARIANCE_REF_AUX) and k_q analogously (B.2.5; floor at K_MIN = 1.0). 3. Run Schmitt-trigger Gate 1 state update (open at target+0.03, close at target-0.03; intentional discontinuity at transition is smoothed by the β rate-limiter downstream). 4. Evaluate Gate 1 sigmoid (aux competence, distance from threshold) and Gate 2 sigmoid (Q-aux disagreement vs analytic 0.5 baseline). 5. alpha_grad_raw = gate1 × gate2 × warmup_gate (structurally bounded to [0, 1] per pearl_bounded_modifier_outputs_require_structural_ activation; no runtime clamp). 6. Update Welford variance of alpha_grad_raw → adaptive β (B.2.8; floor BETA_BASE = 0.5, ceiling BETA_MAX = 0.95). 7. alpha_grad_smoothed = β × prev + (1-β) × raw (rate-limited). 8. Write back 7 outputs: k_aux (385), k_q (386), β (387), var_alpha (390), gate1_state (391), alpha_raw (392), alpha_smoothed (393). Sigmoid arguments clipped to [-30, 30] before __expf for fp32 overflow guard (precision-neutral; sigmoid saturates bit-equal at those bounds). Per pearl_bounded_modifier_outputs_require_structural_activation: sigmoid composition produces structurally-bounded [0, 1] output. KNOWN LIMITATION: as of B.4 landing, NO upstream kernel writes ISV[388] (AUX_DIR_ACC_VARIANCE_EMA). The grep at status-report time finds only the sp14_isv_slots.rs declaration. Effect: var_aux stays at sentinel 0.0 forever, so k_aux is degenerate-but-non-fatal at K_BASE_AUX (constant). Gate 1 still works, the sigmoid just doesn't soften under noisy aux_dir_acc. To be resolved in B.11 producer- chain orchestrator OR a separate fix-up task that adds a Welford- variance update next to the existing AUX_DIR_ACC_SHORT_EMA producer. var_q (389) IS written by q_disagreement_update_kernel (B.3), so adaptive k_q is fully functional from B.4 onward. Slot indices hardcoded inside the kernel via const int locals — must match crates/ml/src/cuda_pipeline/sp14_isv_slots.rs (and 372/373 from sp13_isv_slots.rs). The plan originally documented 381/383/ 384/385/386/387/388/389/390/391 for SP14 slots; the actual values are +2 because SP13 closeout added HOLD_RATE_TARGET=381 + HOLD_RATE_OBSERVED_EMA=382 after the plan was written. Tests (RTX 3050 Ti pass; B.3's 2 tests still pass — no regression): - alpha_grad_schmitt_hysteresis: 4-step trajectory verifies the closed→open→open→closed transition. Closed at aux=0.55 (below open=0.58); opens at aux=0.60; stays open at aux=0.54 (in hysteresis band [close=0.52, open=0.58]); finally closes at aux=0.50 (below close=0.52). - alpha_grad_adaptive_beta: 20-oscillation regime verifies β grows above β_base=0.5 and remains bounded by β_max=0.95. docs/dqn-wire-up-audit.md updated per Invariant 7 with full B.4 behaviour contract, per-step pipeline, single-thread launch convention, sigmoid clip rationale, Schmitt discontinuity note, and the var_aux Known Limitation. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 24 ++ .../alpha_grad_compute_kernel.cu | 226 +++++++++++++++ crates/ml/tests/sp14_oracle_tests.rs | 268 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 76 +++++ 4 files changed, 594 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f5042b804..0e5288093 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -756,6 +756,30 @@ fn main() { // (currently 383, 384, 389; the layout fingerprint regression // test catches drift). "q_disagreement_update_kernel.cu", + // SP14 Layer B Task B.4 (2026-05-05): Earned Gradient Flow + // consumer kernel — the heart of the EGF pearl. Single-thread + // state-machine kernel that reads driver signals (target dir- + // acc, aux dir-acc EMA, Q-disagreement EMA, three Welford + // variance EMAs, persistent Gate-1 state) from ISV, runs a + // Schmitt-trigger Gate 1 (open at target+0.03, close at + // target-0.03), computes adaptive k_aux/k_q/β from the + // variance EMAs, evaluates two sigmoids (Gate 1 = aux + // competence, Gate 2 = Q-head agreement vs random-alignment + // baseline 0.5), multiplies with the host-supplied warmup + // gate to produce α_grad_raw, applies a β-rate-limiter to + // produce α_grad_smoothed, and writes 7 outputs (k_aux, k_q, + // β, var_α, gate1_state, α_raw, α_smoothed) to ISV slots + // 385-393. Sigmoid arguments clipped to [-30, 30] for fp32 + // overflow guard. Per pearl_bounded_modifier_outputs_require_ + // structural_activation: sigmoid composition is structurally + // bounded to [0, 1], no runtime clamp. KNOWN LIMITATION: the + // var_aux producer (slot 388) is not yet wired as of B.4 — + // until B.11 / fix-up adds it, k_aux is degenerate-but-non- + // fatal at K_BASE_AUX (constant). Filed in the B.4 status + // report. var_q (slot 389) is written by + // q_disagreement_update_kernel (B.3), so adaptive k_q is + // fully functional. + "alpha_grad_compute_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu b/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu new file mode 100644 index 000000000..ce12b7f8a --- /dev/null +++ b/crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu @@ -0,0 +1,226 @@ +// crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu +// +// SP14 Layer B Task B.4 (2026-05-05): Earned Gradient Flow consumer kernel — +// the heart of the EGF pearl. Single-thread state-machine kernel that reads +// driver signals from the global ISV bus, runs Schmitt-trigger Gate 1, +// computes adaptive k_aux/k_q/β from Welford variance EMAs, evaluates two +// sigmoids, multiplies with a host-supplied warmup gate, applies a β-rate- +// limiter on the result, and writes 7 outputs back to ISV. +// +// ── Single-thread launch ──────────────────────────────────────────────── +// The body runs only on (threadIdx.x == 0, blockIdx.x == 0) — there is no +// parallelism to exploit (this is a state machine consuming O(1) inputs +// and producing O(1) outputs). Caller MUST still launch with +// blockDim ≥ (1,1,1) and gridDim ≥ (1,1,1); a 32-thread block is fine +// because the early-return masks all but lane 0. No shared memory, no +// reductions, no atomicAdd (per `feedback_no_atomicadd.md`). +// +// ── Driver inputs (read-only on this kernel) ──────────────────────────── +// ISV[TARGET_DIR_ACC = 372] target directional accuracy (SP13) +// ISV[AUX_DIR_ACC_SHORT_EMA = 373] fast EMA of aux dir-acc (SP13) +// ISV[Q_DISAGREEMENT_SHORT_EMA = 383] fast EMA of Q-aux disagreement (SP14 B.3) +// ISV[VAR_AUX = 388] Welford variance EMA of aux dir-acc +// ISV[VAR_Q = 389] Welford variance EMA of Q disagreement (SP14 B.3) +// ISV[GATE1_OPEN = 391] persistent Schmitt state (read+written) +// ISV[VAR_ALPHA = 390] Welford variance EMA of α_grad_raw (read+written) +// ISV[ALPHA_SMOOTHED = 393] α_grad_smoothed previous value +// +// ── Outputs (written by this kernel) ──────────────────────────────────── +// ISV[K_AUX_ADAPTIVE = 385] adaptive Gate-1 sigmoid steepness +// ISV[K_Q_ADAPTIVE = 386] adaptive Gate-2 sigmoid steepness +// ISV[BETA_ADAPTIVE = 387] adaptive β rate-limiter coefficient +// ISV[VAR_ALPHA = 390] updated α_grad_raw variance EMA +// ISV[GATE1_OPEN = 391] updated Schmitt state (0 = closed, 1 = open) +// ISV[ALPHA_RAW = 392] α_grad_raw = gate1 × gate2 × warmup_gate +// ISV[ALPHA_SMOOTHED = 393] α_grad_smoothed = β × prev + (1-β) × raw +// +// ── Slot indices: must match `sp14_isv_slots.rs` ──────────────────────── +// The kernel hardcodes ISV slot indices via `const int` locals below. +// They MUST match the constants declared in +// `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (and 372/373 from +// `sp13_isv_slots.rs`). Any edit to the .rs constants requires a matching +// edit here. The plan originally documented 381/383/384/385/386/387/388/ +// 389/390/391 for SP14 slots — the actual values are +2 because SP13 +// closeout added HOLD_RATE_TARGET=381 + HOLD_RATE_OBSERVED_EMA=382 after +// the plan was written. See `sp14_isv_slots.rs` lines 17-20. +// +// ── Schmitt-trigger hysteresis intentional discontinuity ──────────────── +// At the moment Gate 1 transitions (closed → open or open → closed), the +// sigmoid argument switches between `(aux_short - threshold_open)` and +// `(aux_short - threshold_close)`, producing a small DISCONTINUITY in +// α_grad_raw. This is intentional, not a bug — Schmitt hysteresis +// requires asymmetric thresholds between the rising and falling edges to +// suppress chatter in noisy signals. The discontinuity is smoothed +// downstream by the β rate-limiter, and consumers downstream see the +// β-smoothed signal (α_grad_smoothed = ISV[393]), so the discontinuity +// is invisible to the gradient flow it modulates. +// +// ── Bounded modifier per pearl_bounded_modifier_outputs_require_… ─────── +// α_grad_raw = sigmoid(·) × sigmoid(·) × warmup_gate is structurally +// bounded to [0, 1] — the two sigmoids guarantee [0, 1] each, the warmup +// gate is supplied in [0, 1] by the host. No runtime clamp needed. Per +// `pearl_bounded_modifier_outputs_require_structural_activation.md`. +// +// ── Numerical stability: sigmoid arg clipping ─────────────────────────── +// Sigmoid arguments are clipped to [-30, 30] before `__expf`. At |arg| > +// 30, fp32 sigmoid saturates to 0 or 1 with bit-equal output, so the +// clip is a precision-neutral overflow guard (`__expf(30) ≈ 1.07e13`, +// well within fp32 range; `__expf(38) ≈ 3.18e16` and `__expf(89)` +// overflows). The clip prevents NaN propagation from pathological +// k_aux × (aux_short - threshold) products if k_aux is corrupted. +// +// ── KNOWN LIMITATION: var_aux producer not yet wired ──────────────────── +// This kernel READS ISV[388] (AUX_DIR_ACC_VARIANCE_EMA) but does NOT +// write it. As of B.4 landing, NO upstream kernel writes slot 388 — +// `sp14_isv_slots.rs` is the only file referencing that constant. The +// effect: var_aux stays at sentinel 0.0 forever, so +// k_aux = max(K_BASE_AUX / (1 + 0/VARIANCE_REF_AUX), K_MIN) +// = max(K_BASE_AUX, K_MIN) = K_BASE_AUX (constant). +// The adaptive-k_aux mechanism is degenerate-but-non-fatal: Gate 1 still +// works, the sigmoid just doesn't soften under noisy aux_dir_acc. This +// is to be resolved in the B.11 producer chain orchestrator OR a +// separate fix-up task that adds a Welford-variance update next to the +// existing AUX_DIR_ACC_SHORT_EMA producer. Filed in the B.4 status report. +// +// var_q (slot 389) IS written — by `q_disagreement_update_kernel` (B.3), +// which lands the Welford variance EMA in the same launch as the mean +// EMAs. So adaptive k_q is fully functional from B.4 onward. + +extern "C" __global__ +void alpha_grad_compute_kernel( + /* Global ISV bus. Slots 385, 386, 387, 390, 391, 392, 393 are written; + * slots 372, 373, 383, 388, 389 are read; all other slots untouched. + */ + float* __restrict__ isv, + /* Per-epoch warmup ramp value in [0, 1]. Host-computed (B.7+ wires + * this from the trainer's per-epoch step counter against the + * configured WARMUP_STEPS). 0.0 during cold-start (no gradient + * flow); 1.0 once warmup completes (gates fully control the flow). + */ + float warmup_gate, + /* EMA decay coefficient for the α_grad_raw variance update (typical + * 0.05; matches the `alpha_var` argument to + * `q_disagreement_update_kernel`). + */ + float alpha_var) +{ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // ── ISV slot indices (must match sp14_isv_slots.rs and sp13_isv_slots.rs) ── + const int TARGET_DIR_ACC = 372; + const int AUX_DIR_ACC_SHORT_EMA = 373; + const int Q_DISAGREEMENT_SHORT_EMA = 383; + const int K_AUX_ADAPTIVE = 385; + const int K_Q_ADAPTIVE = 386; + const int BETA_ADAPTIVE = 387; + const int VAR_AUX = 388; + const int VAR_Q = 389; + const int VAR_ALPHA = 390; + const int GATE1_OPEN = 391; + const int ALPHA_RAW = 392; + const int ALPHA_SMOOTHED = 393; + + // ── Structural constants (must match sp14_isv_slots.rs) ──────────── + const float K_BASE_AUX = 20.0f; + const float K_BASE_Q = 15.0f; + const float K_MIN = 1.0f; + const float VARIANCE_REF_AUX = 0.01f; + const float VARIANCE_REF_Q = 0.05f; + const float VARIANCE_REF_ALPHA = 0.01f; + const float BETA_BASE = 0.5f; + const float BETA_MAX = 0.95f; + const float SCHMITT_BAND = 0.03f; + const float Q_DISAGREEMENT_BASELINE = 0.5f; + + // ── Read drivers ─────────────────────────────────────────────────── + const float target = isv[TARGET_DIR_ACC]; + const float aux_short = isv[AUX_DIR_ACC_SHORT_EMA]; + const float q_dis = isv[Q_DISAGREEMENT_SHORT_EMA]; + const float var_aux = isv[VAR_AUX]; + const float var_q = isv[VAR_Q]; + const float var_alpha_prev = isv[VAR_ALPHA]; + const float gate1_state_prev = isv[GATE1_OPEN]; + const float alpha_smoothed_prev = isv[ALPHA_SMOOTHED]; + + // ── Adaptive sigmoid steepness (B.2.5) ───────────────────────────── + // Higher variance → smaller k → flatter sigmoid (more diffusion at + // the edge). Floor at K_MIN prevents the sigmoid from collapsing to + // a flat 0.5 line under unbounded noise. + const float k_aux = fmaxf(K_BASE_AUX / (1.0f + var_aux / VARIANCE_REF_AUX), K_MIN); + const float k_q = fmaxf(K_BASE_Q / (1.0f + var_q / VARIANCE_REF_Q), K_MIN); + + // ── Schmitt-trigger Gate 1 state update (B.2.4) ──────────────────── + // Open at target+0.03 (rising edge); close at target-0.03 (falling + // edge). Persistent across calls via ISV[GATE1_OPEN]. + const float threshold_open = target + SCHMITT_BAND; + const float threshold_close = target - SCHMITT_BAND; + float gate1_state_new = gate1_state_prev; + if (gate1_state_prev < 0.5f && aux_short > threshold_open) { + gate1_state_new = 1.0f; + } else if (gate1_state_prev >= 0.5f && aux_short < threshold_close) { + gate1_state_new = 0.0f; + } + + // ── Sigmoid arg for Gate 1 (uses appropriate threshold per state) ── + // Open state: distance from CLOSE threshold (must drop further to + // flip); the kernel evaluates the soft sigmoid above the close + // threshold so α stays high while the gate is open. + // Closed state: distance from OPEN threshold (must rise further to + // flip); the sigmoid stays near 0 below the open threshold. + // The intentional discontinuity at the transition is documented in + // the header comment block above. + const float gate1_arg_raw = (gate1_state_new >= 0.5f) + ? (aux_short - threshold_close) + : (aux_short - threshold_open); + const float gate1_arg = fmaxf(-30.0f, fminf(30.0f, k_aux * gate1_arg_raw)); + const float gate1 = 1.0f / (1.0f + __expf(-gate1_arg)); + + // ── Sigmoid for Gate 2 (B.2.3) ───────────────────────────────────── + // q_dis above the analytic random-alignment baseline (0.5) is the + // "Q-head agrees more with aux than chance" signal; sigmoid maps it + // into [0, 1]. No Schmitt on Gate 2 — q_dis is already an EMA so + // its high-frequency noise is filtered by upstream smoothing. + const float gate2_arg_raw = q_dis - Q_DISAGREEMENT_BASELINE; + const float gate2_arg = fmaxf(-30.0f, fminf(30.0f, k_q * gate2_arg_raw)); + const float gate2 = 1.0f / (1.0f + __expf(-gate2_arg)); + + // ── α_grad_raw (B.2.1) ───────────────────────────────────────────── + // gate1 ∈ [0, 1], gate2 ∈ [0, 1], warmup_gate ∈ [0, 1] → product is + // structurally bounded to [0, 1]. No runtime clamp per + // `pearl_bounded_modifier_outputs_require_structural_activation.md`. + const float alpha_raw = gate1 * gate2 * warmup_gate; + + // ── Welford-style variance EMA on α_grad_raw ─────────────────────── + // Reference is α_smoothed (the rate-limited tracking estimate); under + // stable α_raw the diff stays small and var_alpha decays toward 0. + // Under chatter (Schmitt hysteresis ringing, sigmoid edge oscillation) + // diff grows and var_alpha rises — driving β toward β_max via the + // adaptive-β formula below. + const float diff_alpha = alpha_raw - alpha_smoothed_prev; + const float var_alpha_new = alpha_var * (diff_alpha * diff_alpha) + + (1.0f - alpha_var) * var_alpha_prev; + + // ── Adaptive β (B.2.8) ───────────────────────────────────────────── + // Floor BETA_BASE prevents β collapsing below light-smoothing under + // very-stable α; ceiling BETA_MAX prevents lockup under pathological + // noise. The variance ratio drives the slide between the two. + const float beta = fmaxf(BETA_BASE, + fminf(BETA_MAX, + BETA_BASE + var_alpha_new / VARIANCE_REF_ALPHA)); + + // ── α_grad_smoothed (rate-limited) ───────────────────────────────── + // β = 0 → no smoothing (α_smoothed = α_raw); β = 1 → frozen + // (α_smoothed = α_smoothed_prev). The kernel writes the smoothed + // value to ISV[393] for downstream consumers (the backward path's + // gradient scaler in B.7+). + const float alpha_smoothed = beta * alpha_smoothed_prev + (1.0f - beta) * alpha_raw; + + // ── Write back ───────────────────────────────────────────────────── + isv[K_AUX_ADAPTIVE] = k_aux; + isv[K_Q_ADAPTIVE] = k_q; + isv[BETA_ADAPTIVE] = beta; + isv[VAR_ALPHA] = var_alpha_new; + isv[GATE1_OPEN] = gate1_state_new; + isv[ALPHA_RAW] = alpha_raw; + isv[ALPHA_SMOOTHED] = alpha_smoothed; +} diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index 79389e34a..064887a64 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -81,7 +81,19 @@ mod gpu { use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + use ml::cuda_pipeline::sp13_isv_slots::{ + AUX_DIR_ACC_SHORT_EMA_INDEX, + TARGET_DIR_ACC_INDEX, + }; use ml::cuda_pipeline::sp14_isv_slots::{ + ALPHA_GRAD_RAW_INDEX, + ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX, + ALPHA_GRAD_SMOOTHED_INDEX, + AUX_DIR_ACC_VARIANCE_EMA_INDEX, + BETA_RATE_LIMITER_ADAPTIVE_INDEX, + GATE1_OPEN_STATE_INDEX, + K_AUX_ADAPTIVE_INDEX, + K_Q_ADAPTIVE_INDEX, Q_DISAGREEMENT_LONG_EMA_INDEX, Q_DISAGREEMENT_SHORT_EMA_INDEX, Q_DISAGREEMENT_VARIANCE_EMA_INDEX, @@ -307,4 +319,260 @@ mod gpu { "all-Hold: short_ema should be in [0.0, 0.5]; got {short_ema}", ); } + + // ── B.4 cubin handle ──────────────────────────────────────────────────── + + const SP14_ALPHA_GRAD_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/alpha_grad_compute_kernel.cubin")); + + fn load_alpha_grad(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP14_ALPHA_GRAD_CUBIN.to_vec()) + .expect("load alpha_grad_compute_kernel cubin"); + module + .load_function("alpha_grad_compute_kernel") + .expect("load alpha_grad_compute_kernel function") + } + + /// Block-dim convention for the alpha_grad kernel: the kernel runs as a + /// state machine on lane 0 only (`if (threadIdx.x != 0 || blockIdx.x != + /// 0) return;`), so any blockDim ≥ 1 and gridDim ≥ 1 works. We launch + /// at 32 threads × 1 block to mirror the production wiring; no shared + /// memory required. + const ALPHA_BLOCK: u32 = 32; + + // ── Test B.4a: Schmitt-trigger hysteresis ────────────────────────────── + + /// 4-step trajectory verifying Gate 1 hysteresis: below open-threshold + /// the gate stays closed; once aux_dir_acc rises above open, the gate + /// opens; while open, drops between close and open thresholds keep the + /// gate open (hysteresis band); only a drop below close-threshold + /// finally closes it. + #[test] + #[ignore = "requires GPU"] + fn alpha_grad_schmitt_hysteresis() { + let stream = make_test_stream(); + let kernel = load_alpha_grad(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // SP13 driver slots (no shift). + isv[TARGET_DIR_ACC_INDEX] = 0.55; + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.55; // exactly at target → below threshold_open=0.58 + // SP14 driver slots (post-+2 shift). + isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.55; // > 0.5 baseline → gate2 ≈ 0.66 + isv[K_AUX_ADAPTIVE_INDEX] = 20.0; + isv[K_Q_ADAPTIVE_INDEX] = 15.0; + isv[BETA_RATE_LIMITER_ADAPTIVE_INDEX] = 0.5; + isv[AUX_DIR_ACC_VARIANCE_EMA_INDEX] = 0.0; // → k_aux = K_BASE_AUX = 20.0 + isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = 0.0; // → k_q = K_BASE_Q = 15.0 + isv[ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX] = 0.0; // → β = β_base = 0.5 + isv[GATE1_OPEN_STATE_INDEX] = 0.0; // closed (sentinel) + isv[ALPHA_GRAD_SMOOTHED_INDEX] = 0.0; // sentinel + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) } + .expect("alloc isv buffer"); + isv_buf.write_from_slice(&isv); + + let warmup_gate: f32 = 1.0; + let alpha_var: f32 = 0.05; + + // ── Step 1: aux_dir_acc = 0.55 < threshold_open = 0.58 → gate stays closed ── + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (ALPHA_BLOCK, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch alpha_grad_compute_kernel (step 1)"); + } + stream.synchronize().expect("sync after alpha_grad step 1"); + + let result = isv_buf.read_all(); + let alpha_raw_below = result[ALPHA_GRAD_RAW_INDEX]; + let gate1_state_below = result[GATE1_OPEN_STATE_INDEX]; + // Closed-state sigmoid arg: k_aux × (0.55 - 0.58) = 20 × (-0.03) = -0.6 + // → sigmoid(-0.6) ≈ 0.354 → α_raw ≈ 0.354 × gate2 × 1.0. Always < 0.5. + assert!( + alpha_raw_below < 0.5, + "Below threshold_open=0.58 (aux=0.55), gate1 stays closed; α should be < 0.5; got {alpha_raw_below}", + ); + assert_eq!( + gate1_state_below, 0.0, + "gate1_open_state should remain closed (0.0); got {gate1_state_below}", + ); + + // ── Step 2: aux_dir_acc rises to 0.60 > threshold_open=0.58 → gate opens ── + // Re-read full ISV state from device so the kernel sees its own outputs + // as the next step's inputs (state-machine semantics). + let mut isv_step2 = isv_buf.read_all(); + isv_step2[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.60; + isv_buf.write_from_slice(&isv_step2); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (ALPHA_BLOCK, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch alpha_grad_compute_kernel (step 2)"); + } + stream.synchronize().expect("sync after alpha_grad step 2"); + + let result = isv_buf.read_all(); + let alpha_raw_above = result[ALPHA_GRAD_RAW_INDEX]; + let gate1_state_above = result[GATE1_OPEN_STATE_INDEX]; + assert!( + alpha_raw_above > alpha_raw_below, + "Above threshold_open, gate1 opens → α_raw should rise; got below={alpha_raw_below} above={alpha_raw_above}", + ); + assert_eq!( + gate1_state_above, 1.0, + "gate1_open_state should be open (1.0) after rising above threshold_open; got {gate1_state_above}", + ); + + // ── Step 3: aux_dir_acc = 0.54 in hysteresis band [0.52, 0.58] → stays open ── + let mut isv_step3 = isv_buf.read_all(); + isv_step3[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.54; + isv_buf.write_from_slice(&isv_step3); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (ALPHA_BLOCK, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch alpha_grad_compute_kernel (step 3)"); + } + stream.synchronize().expect("sync after alpha_grad step 3"); + + let result = isv_buf.read_all(); + let gate1_state_between = result[GATE1_OPEN_STATE_INDEX]; + assert_eq!( + gate1_state_between, 1.0, + "Schmitt hysteresis: gate stays open in band [close=0.52, open=0.58] (aux=0.54); got {gate1_state_between}", + ); + + // ── Step 4: aux_dir_acc = 0.50 < threshold_close = 0.52 → gate closes ── + let mut isv_step4 = isv_buf.read_all(); + isv_step4[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.50; + isv_buf.write_from_slice(&isv_step4); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (ALPHA_BLOCK, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch alpha_grad_compute_kernel (step 4)"); + } + stream.synchronize().expect("sync after alpha_grad step 4"); + + let result = isv_buf.read_all(); + let gate1_state_closed = result[GATE1_OPEN_STATE_INDEX]; + assert_eq!( + gate1_state_closed, 0.0, + "Below threshold_close=0.52 (aux=0.50), gate finally closes; got {gate1_state_closed}", + ); + } + + // ── Test B.4b: Adaptive β rate limiter ───────────────────────────────── + + /// Drive 20 oscillations of aux_dir_acc above and below the Schmitt + /// thresholds. The rate-limiter β should grow above β_base = 0.5 and + /// stay bounded by β_max = 0.95 as the α_grad_raw variance EMA rises. + /// This exercises the full adaptive-β feedback loop end-to-end. + #[test] + #[ignore = "requires GPU"] + fn alpha_grad_adaptive_beta() { + let stream = make_test_stream(); + let kernel = load_alpha_grad(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + isv[TARGET_DIR_ACC_INDEX] = 0.55; + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.55; + isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.6; + isv[K_AUX_ADAPTIVE_INDEX] = 20.0; + isv[K_Q_ADAPTIVE_INDEX] = 15.0; + isv[BETA_RATE_LIMITER_ADAPTIVE_INDEX] = 0.5; + isv[AUX_DIR_ACC_VARIANCE_EMA_INDEX] = 0.0; + isv[Q_DISAGREEMENT_VARIANCE_EMA_INDEX] = 0.0; + isv[ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX] = 0.0; + isv[GATE1_OPEN_STATE_INDEX] = 0.0; + isv[ALPHA_GRAD_RAW_INDEX] = 0.0; + isv[ALPHA_GRAD_SMOOTHED_INDEX] = 0.0; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) } + .expect("alloc isv buffer"); + isv_buf.write_from_slice(&isv); + + // 20 steps with aux_dir_acc oscillating above and below thresholds. + // Each oscillation flips Gate 1 state, producing α_grad_raw chatter + // that grows the variance EMA → adaptive β rises toward β_max. + let oscillation: [f32; 20] = [ + 0.60, 0.50, 0.60, 0.50, 0.60, 0.50, 0.60, 0.50, + 0.60, 0.50, 0.60, 0.50, 0.60, 0.50, 0.60, 0.50, + 0.60, 0.50, 0.60, 0.50, + ]; + + for &acc in oscillation.iter() { + // Re-read the full ISV state, mutate the input slot, write back. + // Mirrors the production state-machine semantics where each step + // observes the previous step's outputs. + let mut current = isv_buf.read_all(); + current[AUX_DIR_ACC_SHORT_EMA_INDEX] = acc; + isv_buf.write_from_slice(¤t); + + let warmup_gate: f32 = 1.0; + let alpha_var: f32 = 0.05; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .arg(&warmup_gate) + .arg(&alpha_var) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (ALPHA_BLOCK, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch alpha_grad_compute_kernel"); + } + stream.synchronize().expect("sync after alpha_grad oscillation step"); + } + + let result = isv_buf.read_all(); + let beta_after = result[BETA_RATE_LIMITER_ADAPTIVE_INDEX]; + assert!( + beta_after > 0.5, + "After 20 oscillations, adaptive β should grow above β_base=0.5; got {beta_after}", + ); + assert!( + beta_after <= 0.95, + "β should remain bounded by β_max=0.95; got {beta_after}", + ); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index bf2186d6f..685dc5857 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -6307,3 +6307,79 @@ The original SP14 plan documented the slots as 381/382/387, but Phase 0 verifica - **Consumers**: ISV[383], ISV[384] read by the `α_grad` gate logic (B.4); ISV[389] read by the adaptive `k_q` sigmoid steepness in B.4. None active until B.4+. This is a known-orphan kernel for the duration of the B.3..B.6 producer chain. The orphan is acknowledged here in the audit doc per `feedback_wire_everything_up.md` (the rule's "same-commit wire-up" requirement is relaxed for atomic chained-producer-consumer landings as long as every producer is documented as orphaned and every consumer's wire-up commit is in the same chain). + +## SP14 Layer B — Commit B.4: alpha_grad_compute_kernel + cubin registration + 2 GPU oracle tests (2026-05-05) + +**Why this commit.** B.4 is the heart of the Earned Gradient Flow pearl: a single-thread state-machine kernel that consumes the producer outputs landed by B.3 (`q_disagreement_update_kernel` writing slots 383/384/389) plus SP13's `aux_dir_acc` EMA (slot 373) and target (slot 372), runs Schmitt-trigger Gate 1 (aux competence) and a baseline-comparing Gate 2 (Q-aux disagreement), composes both with a host-supplied warmup gate to produce `α_grad_raw`, and writes a β-rate-limited `α_grad_smoothed` for downstream consumers. After B.4, the EGF gate is computable end-to-end given the variance EMAs are populated; B.5/B.6 add the gradient-hack circuit breaker and the `dir_concat_qaux` consumer wiring respectively, B.7+ wires the producer-chain into the captured CUDA Graph. + +This commit lands the kernel + cubin registration + GPU oracle tests; it does NOT add a host-side launcher (no calls into the new cubin from production code). The kernel is dormant after this commit — only the oracle tests exercise it. Same isolation rationale as B.3. + +### Behaviour contract (B.4 isolated) + +| Side | Owns | Reads | Writes | +|------|------|-------|--------| +| `alpha_grad_compute_kernel.cu` | Single-thread state machine: Schmitt-trigger Gate 1 + sigmoid composition + Welford variance EMA + adaptive β + rate-limited smoothing | ISV[372] `TARGET_DIR_ACC`, ISV[373] `AUX_DIR_ACC_SHORT_EMA`, ISV[383] `Q_DISAGREEMENT_SHORT_EMA`, ISV[388] `AUX_DIR_ACC_VARIANCE_EMA` (read but not produced — see Known Limitation), ISV[389] `Q_DISAGREEMENT_VARIANCE_EMA`, ISV[390] `ALPHA_GRAD_RAW_VARIANCE_EMA` (read+written), ISV[391] `GATE1_OPEN_STATE` (read+written), ISV[393] `ALPHA_GRAD_SMOOTHED` (read+written) | ISV[385] `K_AUX_ADAPTIVE`, ISV[386] `K_Q_ADAPTIVE`, ISV[387] `BETA_RATE_LIMITER_ADAPTIVE`, ISV[390] `ALPHA_GRAD_RAW_VARIANCE_EMA`, ISV[391] `GATE1_OPEN_STATE`, ISV[392] `ALPHA_GRAD_RAW`, ISV[393] `ALPHA_GRAD_SMOOTHED` | +| `build.rs` kernel manifest | Cubin emission via shared `try_compile_kernel` path | `alpha_grad_compute_kernel.cu` source | `OUT_DIR/alpha_grad_compute_kernel.cubin` | +| `tests/sp14_oracle_tests.rs::gpu` module | Oracle correctness — Schmitt hysteresis trajectory + adaptive β under chatter | The cubin + driver-signal-populated ISV slots | Pass/fail on Schmitt 4-step open/close transitions + β > β_base after 20 oscillations | + +### Per-step pipeline + +1. Read drivers (lines: `target = isv[372]`, `aux_short = isv[373]`, `q_dis = isv[383]`, `var_aux = isv[388]`, `var_q = isv[389]`, `var_alpha_prev = isv[390]`, `gate1_state_prev = isv[391]`, `alpha_smoothed_prev = isv[393]`). +2. Compute adaptive sigmoid steepness (B.2.5): `k_aux = max(K_BASE_AUX / (1 + var_aux/VARIANCE_REF_AUX), K_MIN)` and analogously for `k_q`. Higher variance → flatter sigmoid; floor at K_MIN = 1.0 prevents collapse to a flat-0.5 line. +3. Schmitt-trigger Gate 1 state update (B.2.4): open at `target + 0.03`, close at `target - 0.03`. Persistent state in ISV[391] survives across calls. The intentional discontinuity at the transition (sigmoid argument flips between `aux - threshold_close` and `aux - threshold_open`) is smoothed downstream by the β rate-limiter. +4. Evaluate Gate 1 sigmoid (`1 / (1 + exp(-k_aux × (aux_short - threshold_·)))`) using the appropriate threshold per state. Argument clipped to [-30, 30] for fp32 overflow guard (precision-neutral: sigmoid saturates bit-equal at those bounds). +5. Evaluate Gate 2 sigmoid (B.2.3): `1 / (1 + exp(-k_q × (q_dis - 0.5)))` against the analytic random-alignment baseline. No Schmitt — `q_dis` is already an EMA so its high-frequency noise is filtered upstream. +6. Compose: `alpha_raw = gate1 × gate2 × warmup_gate`. Structurally bounded to [0, 1] per `pearl_bounded_modifier_outputs_require_structural_activation.md`; no runtime clamp. +7. Update Welford-style variance EMA on `alpha_raw` against the rate-limited tracking estimate `alpha_smoothed_prev`: `var_alpha_new = α_var × diff² + (1-α_var) × var_alpha_prev`. +8. Compute adaptive β (B.2.8): `beta = clamp(BETA_BASE + var_alpha_new / VARIANCE_REF_ALPHA, BETA_BASE, BETA_MAX)`. Floors at BETA_BASE = 0.5 (light smoothing); ceilings at BETA_MAX = 0.95 (lockup guard). +9. Rate-limit: `alpha_smoothed = β × alpha_smoothed_prev + (1-β) × alpha_raw`. +10. Write back 7 outputs to ISV[385], ISV[386], ISV[387], ISV[390], ISV[391], ISV[392], ISV[393]. + +### Single-thread launch contract + +The kernel runs only on `(threadIdx.x == 0, blockIdx.x == 0)` — there is no parallelism to exploit (state machine consumes O(1) inputs and produces O(1) outputs). Caller MUST still launch with `blockDim ≥ (1,1,1)` and `gridDim ≥ (1,1,1)`; a 32-thread block is fine because the early-return masks all but lane 0. No shared memory, no reductions, no atomicAdd (per `feedback_no_atomicadd.md`). The tests launch at 32 × 1 × 1 with `shared_mem_bytes = 0`; B.7+'s production launcher will pass the same. + +### Sigmoid arg clipping for fp32 numerical stability + +Both gate sigmoid arguments are clipped to `[-30, 30]` before `__expf`. At |arg| > 30, fp32 sigmoid saturates to 0 or 1 bit-equally, so the clip is precision-neutral. `__expf(30) ≈ 1.07e13` (within fp32 range); `__expf(38) ≈ 3.18e16` and `__expf(89)` overflows to inf. The clip prevents NaN propagation from pathological `k_aux × (aux_short - threshold)` products if `k_aux` gets corrupted upstream. + +### Schmitt-trigger intentional discontinuity + +At the moment Gate 1 transitions (closed → open or open → closed), the sigmoid argument switches between `(aux_short - threshold_open)` and `(aux_short - threshold_close)`, producing a small DISCONTINUITY in `α_grad_raw`. This is intentional, not a bug — Schmitt hysteresis requires asymmetric thresholds between the rising and falling edges to suppress chatter in noisy signals. Downstream consumers see only `α_grad_smoothed` (ISV[393]), which is β-filtered, so the discontinuity is invisible to the gradient flow it modulates. + +### Known limitation: var_aux producer not yet wired + +This kernel READS ISV[388] (`AUX_DIR_ACC_VARIANCE_EMA`) but does NOT write it. As of B.4 landing, NO upstream kernel writes slot 388 — `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` is the only file in `crates/ml/src/cuda_pipeline/` referencing that constant (verified via grep `AUX_DIR_ACC_VARIANCE_EMA_INDEX|signal_at(388|isv\[388\]` over `crates/ml/src/cuda_pipeline/*.cu` + `*.rs` excluding tests). The effect: `var_aux` stays at sentinel `0.0` forever, so + +``` +k_aux = max(K_BASE_AUX / (1 + 0/VARIANCE_REF_AUX), K_MIN) + = max(K_BASE_AUX, K_MIN) = K_BASE_AUX (constant). +``` + +The adaptive-`k_aux` mechanism is degenerate-but-non-fatal: Gate 1 still works, the sigmoid just doesn't soften under noisy `aux_dir_acc`. To be resolved in the B.11 producer-chain orchestrator OR a separate fix-up task that adds a Welford-variance update next to the existing `AUX_DIR_ACC_SHORT_EMA` producer (slot 373's writer in the SP13 chain). Filed in this audit + the B.4 status report. Symmetrical observation for `var_q` (slot 389) does NOT apply — `q_disagreement_update_kernel` (B.3) DOES write slot 389, so adaptive `k_q` is fully functional from B.4 onward. + +### Files changed + +| Type | Files (count) | +|------|---------------| +| New kernel file | `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` (NEW) | +| `build.rs` registration | 1 entry added to `kernels_with_common` array immediately after `q_disagreement_update_kernel.cu` | +| Test append | `crates/ml/tests/sp14_oracle_tests.rs` — 2 GPU oracle tests appended to `mod gpu` (`alpha_grad_schmitt_hysteresis`, `alpha_grad_adaptive_beta`); imports of 8 new ISV slot constants from `sp14_isv_slots` + 2 from `sp13_isv_slots` | +| Audit doc | This section | + +### Build + test verification + +- `SQLX_OFFLINE=true cargo check -p ml --features cuda` — clean (only the 18 pre-existing warnings) +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests alpha_grad --features cuda -- --ignored --nocapture` — 2 PASS: + - `gpu::alpha_grad_schmitt_hysteresis` — 4-step trajectory: aux=0.55 (closed: α<0.5, gate1=0); aux=0.60 (opens: α↑, gate1=1); aux=0.54 (in band [0.52, 0.58]: gate1=1, hysteresis); aux=0.50 (below close: gate1=0, finally closes) + - `gpu::alpha_grad_adaptive_beta` — 20-oscillation regime with aux flipping 0.60 ↔ 0.50 across the Schmitt thresholds; final β > β_base = 0.5 and ≤ β_max = 0.95 +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests q_disagreement --features cuda -- --ignored --nocapture` — 2 PASS (B.3 regression check; no impact) + +### Wire-up status + +- **Producer**: kernel exists; cubin built; tests green. +- **Launcher**: NOT YET WIRED. The Rust launcher (`launch_alpha_grad_compute` or similar) will land in B.7+ together with the captured-graph integration of the full producer chain. +- **Consumers**: ISV[392] (`ALPHA_GRAD_RAW`) is HEALTH_DIAG-visibility only; ISV[393] (`ALPHA_GRAD_SMOOTHED`) is consumed by the `α_grad`-multiplied SAXPY in the Q-head backward path (B.7+). +- **Reverse dependencies**: this kernel depends on B.3's writes to ISV[383] and ISV[389], plus SP13's writes to ISV[372] and ISV[373]. ISV[388] is not yet produced (Known Limitation above). + +This kernel is the second known-orphan in the B.3..B.6 producer chain (B.3 was the first). The orphan is acknowledged here per `feedback_wire_everything_up.md` (the rule's "same-commit wire-up" requirement is relaxed for atomic chained-producer-consumer landings as long as every producer is documented as orphaned and every consumer's wire-up commit is in the same chain).