feat(sp14): B.4 — alpha_grad_compute_kernel (EGF heart)
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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)
|
||||
|
||||
226
crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu
Normal file
226
crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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<CudaStream>) -> 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}",
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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).
|
||||
|
||||
Reference in New Issue
Block a user