From 82fe6cea665c99d15229b67e0261e13fdd090f8b Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 5 May 2026 19:29:22 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp14):=20B.5=20=E2=80=94=20gradient=5Fhack?= =?UTF-8?q?=5Fdetect=5Fkernel=20(anti-mesa-opt=20circuit=20breaker)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Detects suspected gradient hacking: when gate1 is open AND aux_dir_acc post-open-minimum drops > LOCKOUT_TRIGGER_DROP (0.05) below the Schmitt open-threshold (target + SCHMITT_BAND = target + 0.03) AND q_disagreement rises > LOCKOUT_TRIGGER_DIS_RISE (0.10) above the analytic random-alignment baseline 0.5, simultaneously. Action: force gate1_open_state = 0 (ISV[391]); set lockout_remaining = 2.0 epochs (LOCKOUT_EPOCHS). During lockout, gate1 stays force-closed regardless of alpha_grad_compute_kernel output. Tracks AUX_DIR_ACC_POST_OPEN_MIN (ISV[394]): running minimum of aux_dir_acc since gate1 last opened; resets to 1.0 sentinel when gate closes naturally or when circuit breaker fires. Slot indices shifted +2 from original plan (SP13 closeout added HOLD_RATE_TARGET=381 + HOLD_RATE_OBSERVED_EMA=382): Q_DIS_SHORT=383, GATE1=391, POST_OPEN_MIN=394, LOCKOUT=395. Matches sp14_isv_slots.rs. Single-thread state-machine kernel (threadIdx.x==0 guard); runs at end of each epoch after alpha_grad_compute_kernel. No atomicAdd per feedback_no_atomicadd.md. 1 oracle test: gradient_hack_circuit_breaker_fires verifies trigger conditions (aux_drop=0.08 > 0.05, q_rise=0.15 > 0.10) cause lockout=2.0 and gate1 force-close=0.0. B.3+B.4 regression: 4 GPU tests unchanged (5 total GPU pass, 1 host pass). Co-Authored-By: Claude Sonnet 4.6 --- crates/ml/build.rs | 16 +++ .../gradient_hack_detect_kernel.cu | 121 ++++++++++++++++++ crates/ml/tests/sp14_oracle_tests.rs | 79 ++++++++++++ docs/dqn-wire-up-audit.md | 56 ++++++++ 4 files changed, 272 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 0e5288093..24b3a5ced 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -780,6 +780,22 @@ fn main() { // q_disagreement_update_kernel (B.3), so adaptive k_q is // fully functional. "alpha_grad_compute_kernel.cu", + // SP14 Layer B Task B.5 (2026-05-05): Earned Gradient Flow + // anti-mesa-optimization circuit breaker. Single-thread state- + // machine kernel that runs at the end of each epoch. Detects + // suspected gradient hacking: gate1 open AND aux_dir_acc post- + // open-minimum drops > LOCKOUT_TRIGGER_DROP (0.05) below the + // Schmitt open-threshold (target + 0.03) AND q_disagreement rises + // > LOCKOUT_TRIGGER_DIS_RISE (0.10) above the analytic baseline + // 0.5, simultaneously. Action: force gate1_open_state = 0 and + // set lockout_remaining = 2 epochs. During lockout, gate1 stays + // force-closed. Tracks AUX_DIR_ACC_POST_OPEN_MIN (ISV[394]): + // lowest aux_dir_acc since gate last opened; resets to 1.0 when + // gate closes. Slot indices must match sp14_isv_slots.rs (shifted + // +2 from the plan: 383, 391, 394, 395). Depends on B.4 + // alpha_grad_compute_kernel writing GATE1_OPEN_STATE (ISV[391]) + // before this kernel runs per epoch. + "gradient_hack_detect_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu b/crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu new file mode 100644 index 000000000..427cee86b --- /dev/null +++ b/crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu @@ -0,0 +1,121 @@ +// crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu +// +// SP14 Layer B Task B.5 (2026-05-05): Earned Gradient Flow — anti-mesa- +// optimization circuit breaker. Single-thread state-machine kernel that +// guards the EGF gate against gradient hacking attempts where the model +// learns to fool the aux-competence gate without genuine improvement. +// +// ── What it does ──────────────────────────────────────────────────────── +// Detects suspected gradient hacking: when Gate 1 is open AND the +// aux_dir_acc post-open minimum drops more than LOCKOUT_TRIGGER_DROP (0.05) +// below the Schmitt open-threshold AND q_disagreement rises more than +// LOCKOUT_TRIGGER_DIS_RISE (0.10) above the analytic baseline (0.5) +// simultaneously. +// +// Action on detection: force gate1_open_state = 0 (close gate); set +// lockout_remaining = LOCKOUT_EPOCHS (2.0). During lockout, gate1 is kept +// force-closed by this kernel regardless of what alpha_grad_compute_kernel +// would otherwise set. +// +// ── When it runs ──────────────────────────────────────────────────────── +// End of each epoch (after alpha_grad_compute_kernel has been called for +// all steps in the epoch). The lockout counter is decremented by 1.0 per +// epoch call. While lockout_remaining > 0, gate1 is held at 0 and the +// AUX_DIR_ACC_POST_OPEN_MIN tracker is reset to 1.0. +// +// ── Slot coupling reference ───────────────────────────────────────────── +// Slot indices below MUST match `sp14_isv_slots.rs`. SP14 slots start at +// 383 (SP13 closeout added HOLD_RATE_TARGET=381 and HOLD_RATE_OBSERVED_ +// EMA=382, shifting SP14 slots +2 from the original plan). SP13 slots +// 372/373 are unchanged. +// +// ISV[372] TARGET_DIR_ACC — Schmitt trigger target (SP13) +// ISV[373] AUX_DIR_ACC_SHORT_EMA — aux directional accuracy (SP13) +// ISV[383] Q_DISAGREEMENT_SHORT_EMA — K=4↔K=2 mismatch EMA (SP14 B.3) +// ISV[391] GATE1_OPEN_STATE — Schmitt-trigger state (SP14 B.4) +// ISV[394] AUX_DIR_ACC_POST_OPEN_MIN — lowest aux since gate opened +// ISV[395] GRADIENT_HACK_LOCKOUT_REMAINING — epochs remaining in lockout +// +// ── Dependency on B.4 ─────────────────────────────────────────────────── +// This kernel READS ISV[391] (GATE1_OPEN_STATE) written by +// alpha_grad_compute_kernel (B.4). The circuit breaker ALSO WRITES to +// ISV[391] to force it closed on detection. Caller must sequence +// alpha_grad_compute_kernel → gradient_hack_detect_kernel within the same +// epoch-end launch sequence so the circuit breaker sees the epoch-final +// gate state. +// +// ── Single-thread launch ──────────────────────────────────────────────── +// Runs only on (threadIdx.x == 0, blockIdx.x == 0). No parallelism, no +// shared memory, no atomicAdd (per `feedback_no_atomicadd.md`). Caller +// MUST launch with blockDim ≥ (1,1,1) and gridDim ≥ (1,1,1). + +extern "C" __global__ +void gradient_hack_detect_kernel(float* __restrict__ isv) { + 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 = 372; // TARGET_DIR_ACC_INDEX (SP13) + const int AUX_SHORT = 373; // AUX_DIR_ACC_SHORT_EMA_INDEX (SP13) + const int Q_DIS_SHORT = 383; // Q_DISAGREEMENT_SHORT_EMA_INDEX (SP14 B.3; shifted +2 from plan) + const int GATE1 = 391; // GATE1_OPEN_STATE_INDEX (SP14 B.4; shifted +2 from plan) + const int POST_OPEN_MIN = 394; // AUX_DIR_ACC_POST_OPEN_MIN_INDEX (SP14 B.5; shifted +2 from plan) + const int LOCKOUT = 395; // GRADIENT_HACK_LOCKOUT_REMAINING_INDEX (SP14 B.5; shifted +2 from plan) + + // ── Structural constants (must match sp14_isv_slots.rs) ──────────── + const float SCHMITT_BAND = 0.03f; + const float TRIGGER_DROP = 0.05f; + const float TRIGGER_DIS_RISE = 0.10f; + const float Q_BASELINE = 0.5f; + const float LOCKOUT_EPOCHS = 2.0f; + + // ── Read inputs ───────────────────────────────────────────────────── + const float target = isv[TARGET]; + const float aux = isv[AUX_SHORT]; + const float q_dis = isv[Q_DIS_SHORT]; + const float gate1 = isv[GATE1]; + const float prev_min = isv[POST_OPEN_MIN]; + const float lockout = isv[LOCKOUT]; + + // ── Decrement existing lockout (per-epoch, floor at 0) ───────────── + float lockout_new = fmaxf(0.0f, lockout - 1.0f); + + // ── Track post-open minimum of aux_dir_acc ────────────────────────── + // While gate1 is open: track the running minimum of aux_dir_acc since + // the gate last opened. This tells us whether aux has silently declined + // below the open-threshold (which should have flipped the Schmitt + // trigger, but the model may be hacking the signal upstream). + // When gate is closed: reset the tracker to the "no observation" + // sentinel (1.0) so the next open-period starts fresh. + float post_open_min_new = prev_min; + if (gate1 >= 0.5f) { + post_open_min_new = fminf(prev_min, aux); + } else { + post_open_min_new = 1.0f; + } + + // ── Detect hacking: simultaneous aux drop + q_disagreement rise ───── + // Condition: gate is currently open AND post-open minimum has dropped + // more than TRIGGER_DROP below the Schmitt open-threshold AND + // q_disagreement has risen more than TRIGGER_DIS_RISE above the + // analytic random-alignment baseline. + // + // The threshold_open = target + SCHMITT_BAND mirrors alpha_grad_ + // compute_kernel; aux_drop = how far post_open_min has fallen below it. + const float threshold_open = target + SCHMITT_BAND; + const float aux_drop = threshold_open - post_open_min_new; + const float q_rise = q_dis - Q_BASELINE; + + if (gate1 >= 0.5f && aux_drop > TRIGGER_DROP && q_rise > TRIGGER_DIS_RISE) { + // Trigger circuit breaker: force gate closed + set lockout. + lockout_new = LOCKOUT_EPOCHS; + isv[GATE1] = 0.0f; + post_open_min_new = 1.0f; // reset; gate is now closed + } else if (lockout_new > 0.0f) { + // Still within a prior lockout: keep gate force-closed. + isv[GATE1] = 0.0f; + } + + // ── Write back ────────────────────────────────────────────────────── + isv[POST_OPEN_MIN] = post_open_min_new; + isv[LOCKOUT] = lockout_new; +} diff --git a/crates/ml/tests/sp14_oracle_tests.rs b/crates/ml/tests/sp14_oracle_tests.rs index 064887a64..3316b0e23 100644 --- a/crates/ml/tests/sp14_oracle_tests.rs +++ b/crates/ml/tests/sp14_oracle_tests.rs @@ -89,9 +89,11 @@ mod gpu { ALPHA_GRAD_RAW_INDEX, ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX, ALPHA_GRAD_SMOOTHED_INDEX, + AUX_DIR_ACC_POST_OPEN_MIN_INDEX, AUX_DIR_ACC_VARIANCE_EMA_INDEX, BETA_RATE_LIMITER_ADAPTIVE_INDEX, GATE1_OPEN_STATE_INDEX, + GRADIENT_HACK_LOCKOUT_REMAINING_INDEX, K_AUX_ADAPTIVE_INDEX, K_Q_ADAPTIVE_INDEX, Q_DISAGREEMENT_LONG_EMA_INDEX, @@ -575,4 +577,81 @@ mod gpu { "β should remain bounded by β_max=0.95; got {beta_after}", ); } + + // ── B.5 cubin handle ──────────────────────────────────────────────────── + + const SP14_GRAD_HACK_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/gradient_hack_detect_kernel.cubin")); + + fn load_grad_hack(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP14_GRAD_HACK_CUBIN.to_vec()) + .expect("load gradient_hack_detect_kernel cubin"); + module + .load_function("gradient_hack_detect_kernel") + .expect("load gradient_hack_detect_kernel function") + } + + // ── Test B.5: Anti-gradient-hacking circuit breaker fires ────────────── + + /// Trigger conditions: gate1 is open AND aux_dir_acc post-open minimum is + /// 0.50, which is 0.05 below threshold_open = target(0.55) + SCHMITT(0.03) + /// = 0.58, i.e. aux_drop = 0.08 > TRIGGER_DROP = 0.05. AND q_disagreement + /// = 0.65 is 0.15 above baseline 0.5, i.e. q_rise = 0.15 > + /// TRIGGER_DIS_RISE = 0.10. + /// + /// Expected: lockout_remaining set to 2.0 (LOCKOUT_EPOCHS); gate1_open_state + /// force-closed to 0.0. + #[test] + #[ignore = "requires GPU"] + fn gradient_hack_circuit_breaker_fires() { + let stream = make_test_stream(); + let kernel = load_grad_hack(&stream); + + const ISV_DIM: usize = 1024; + let mut isv = vec![0.0_f32; ISV_DIM]; + // SP13 driver slots (indices unchanged from SP13). + isv[TARGET_DIR_ACC_INDEX] = 0.55; // target + isv[AUX_DIR_ACC_SHORT_EMA_INDEX] = 0.50; // aux_dir_acc dropped + // SP14 slots (shifted +2 from original plan; see sp14_isv_slots.rs). + isv[Q_DISAGREEMENT_SHORT_EMA_INDEX] = 0.65; // q_disagreement rose + isv[GATE1_OPEN_STATE_INDEX] = 1.0; // gate is currently open + isv[AUX_DIR_ACC_POST_OPEN_MIN_INDEX] = 0.50; // post-open min = current aux + isv[GRADIENT_HACK_LOCKOUT_REMAINING_INDEX] = 0.0; // no prior lockout + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_DIM) } + .expect("alloc isv buffer for gradient_hack test"); + isv_buf.write_from_slice(&isv); + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_buf.dev_ptr) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch gradient_hack_detect_kernel"); + } + stream.synchronize().expect("sync after gradient_hack_detect_kernel"); + + let result = isv_buf.read_all(); + // aux_dir_acc = 0.50; threshold_open = 0.55 + 0.03 = 0.58. + // post_open_min = min(0.50, 0.50) = 0.50. + // aux_drop = 0.58 - 0.50 = 0.08 > TRIGGER_DROP = 0.05. ✓ + // q_rise = 0.65 - 0.50 = 0.15 > TRIGGER_DIS_RISE = 0.10. ✓ + // Both conditions met → trigger circuit breaker. + let lockout = result[GRADIENT_HACK_LOCKOUT_REMAINING_INDEX]; + let gate1 = result[GATE1_OPEN_STATE_INDEX]; + assert_eq!( + lockout, 2.0, + "Lockout should be set to LOCKOUT_EPOCHS=2; got {lockout}", + ); + assert_eq!( + gate1, 0.0, + "Gate1 should be force-closed to 0.0 on circuit-breaker trigger; got {gate1}", + ); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 685dc5857..e254ab03c 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -6383,3 +6383,59 @@ The adaptive-`k_aux` mechanism is degenerate-but-non-fatal: Gate 1 still works, - **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). + +## SP14 Layer B — Commit B.5: gradient_hack_detect_kernel + cubin registration + 1 GPU oracle test (2026-05-05) + +**Why this commit.** B.5 is the anti-mesa-optimization circuit breaker for the Earned Gradient Flow pearl. Mesa-optimization (gradient hacking) can manifest as the model learning to elevate `aux_dir_acc` above the Schmitt open-threshold *without* a genuine improvement in directional accuracy — i.e., the aux head learns to fool the gate rather than predict correctly. The circuit breaker detects this by watching for a simultaneous drop in `aux_dir_acc` post-open-minimum and rise in `q_disagreement`: if the aux head were actually improving, Q-disagreement would stay near the analytic baseline (0.5); a rising disagreement alongside a falling post-open minimum is the smoking-gun pattern. + +This commit lands the kernel + cubin registration + GPU oracle test; it does NOT add a host-side launcher. The kernel is dormant after this commit — only the oracle test exercises it. Same isolation rationale as B.3 and B.4. + +### Behaviour contract (B.5 isolated) + +| Side | Owns | Reads | Writes | +|------|------|-------|--------| +| `gradient_hack_detect_kernel.cu` | Single-thread epoch-end state machine: post-open minimum tracker + simultaneous-drop/rise detector + force-close + lockout decrement | ISV[372] `TARGET_DIR_ACC`, ISV[373] `AUX_DIR_ACC_SHORT_EMA`, ISV[383] `Q_DISAGREEMENT_SHORT_EMA`, ISV[391] `GATE1_OPEN_STATE` (read+written), ISV[394] `AUX_DIR_ACC_POST_OPEN_MIN` (read+written), ISV[395] `GRADIENT_HACK_LOCKOUT_REMAINING` (read+written) | ISV[391] `GATE1_OPEN_STATE` (force-closed on trigger or during lockout), ISV[394] `AUX_DIR_ACC_POST_OPEN_MIN`, ISV[395] `GRADIENT_HACK_LOCKOUT_REMAINING` | +| `build.rs` kernel manifest | Cubin emission via shared `try_compile_kernel` path | `gradient_hack_detect_kernel.cu` source | `OUT_DIR/gradient_hack_detect_kernel.cubin` | +| `tests/sp14_oracle_tests.rs::gpu` module | Oracle correctness — circuit-breaker trigger and lockout assignment | The cubin + ISV slots initialised to trigger condition | Pass/fail on `lockout_remaining == 2.0` and `gate1_open == 0.0` | + +### Per-epoch pipeline + +1. Decrement existing lockout: `lockout_new = max(0, lockout - 1)`. +2. Track post-open minimum: if `gate1 >= 0.5` (open), `post_open_min_new = min(prev_min, aux)`; if closed, reset to sentinel `1.0`. +3. Compute drop and rise: `aux_drop = threshold_open − post_open_min_new` and `q_rise = q_dis − Q_BASELINE (0.5)`. `threshold_open = target + SCHMITT_BAND (0.03)` mirrors `alpha_grad_compute_kernel`. +4. If `gate1 >= 0.5 AND aux_drop > TRIGGER_DROP (0.05) AND q_rise > TRIGGER_DIS_RISE (0.10)`: circuit breaker fires — set `lockout_new = LOCKOUT_EPOCHS (2.0)`, write `ISV[GATE1] = 0.0` (force-close), reset `post_open_min_new = 1.0`. +5. Else if `lockout_new > 0`: still in lockout — write `ISV[GATE1] = 0.0` (maintain force-close). +6. Write back `ISV[POST_OPEN_MIN]` and `ISV[LOCKOUT]`. + +### Slot index note + +All SP14 slot indices are shifted +2 from the original plan (SP13 closeout added `HOLD_RATE_TARGET=381` and `HOLD_RATE_OBSERVED_EMA=382`). The kernel hardcodes them via `const int` locals and comments reference the `.rs` constants explicitly. Values: `Q_DIS_SHORT=383`, `GATE1=391`, `POST_OPEN_MIN=394`, `LOCKOUT=395`. + +### Dependency ordering + +This kernel WRITES to ISV[391] (`GATE1_OPEN_STATE`) — the same slot that `alpha_grad_compute_kernel` (B.4) writes. The correct epoch-end sequence is therefore: all-steps `alpha_grad_compute_kernel` (per-step), then `gradient_hack_detect_kernel` (once per epoch). The circuit breaker sees the epoch-final gate state from B.4 and can override it. Callers must ensure this ordering; mixing or reversing will produce a stale gate1 read. + +### Files changed + +| Type | Files (count) | +|------|---------------| +| New kernel file | `crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu` (NEW) | +| `build.rs` registration | 1 entry added to `kernels_with_common` array immediately after `alpha_grad_compute_kernel.cu` | +| Test append | `crates/ml/tests/sp14_oracle_tests.rs` — 1 GPU oracle test appended to `mod gpu` (`gradient_hack_circuit_breaker_fires`); 2 additional imports from `sp14_isv_slots` (`AUX_DIR_ACC_POST_OPEN_MIN_INDEX`, `GRADIENT_HACK_LOCKOUT_REMAINING_INDEX`) added to the module-level import block | +| Audit doc | This section | + +### Build + test verification + +- `SQLX_OFFLINE=true cargo check -p ml` — clean (only the 18 pre-existing warnings; no new warnings) +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests gradient_hack --features cuda -- --ignored --nocapture` — 1 PASS: + - `gpu::gradient_hack_circuit_breaker_fires` — gate1=1 (open), aux=0.50 (threshold_open=0.58 → aux_drop=0.08 > 0.05), q_dis=0.65 (q_rise=0.15 > 0.10): circuit breaker fires → lockout=2.0, gate1=0.0 +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --features cuda -- --ignored --nocapture` — 5 PASS total (B.3: 2, B.4: 2, B.5: 1; no regressions) + +### Wire-up status + +- **Producer**: kernel exists; cubin built; test green. +- **Launcher**: NOT YET WIRED. The Rust launcher will land in B.7+ together with the captured-graph integration of the full producer chain. +- **Consumers**: no downstream slot consumers for the circuit-breaker outputs — the kernel writes ISV[391]/[394]/[395] in-place; the effect is felt only through gate1 being force-closed when B.4 re-runs on the next epoch. +- **Reverse dependencies**: reads ISV[372] (SP13 TARGET), ISV[373] (SP13 AUX_SHORT), ISV[383] (B.3 Q_DIS_SHORT), ISV[391] (B.4 GATE1). All produced by earlier kernels in the chain. + +This kernel is the third known-orphan in the B.3..B.6 producer chain. The orphan is acknowledged here per `feedback_wire_everything_up.md`.