feat(sp14): B.5 — gradient_hack_detect_kernel (anti-mesa-opt circuit breaker)
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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)
|
||||
|
||||
121
crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu
Normal file
121
crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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<CudaStream>) -> 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}",
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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`.
|
||||
|
||||
Reference in New Issue
Block a user