// 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;
}
