/* tau_update — Polyak-EMA tau coefficient, GPU-driven (Plan 1 Task 13).
 * Reads:  ISV[EPOCH_IDX_INDEX=39], ISV[TOTAL_EPOCHS_INDEX=40],
 *         ISV[LEARNING_HEALTH_INDEX=12], ISV[Q_DRIFT_RATE_INDEX=129]
 * Writes: ISV[TAU_EFF_INDEX=42]
 * Cold path (per-epoch boundary). Single-thread kernel.
 *
 * Base formula: cosine schedule (tau_base -> tau_final) over training, with
 * health-coupled floor that rises to 0.01 at full collapse.
 *
 * Plan C Phase 2 follow-up A.2 (2026-04-29) — adaptive Polyak-tau drift
 * dampening:
 *
 *   drift_factor = 1.0 / (1.0 + clip(ISV[Q_DRIFT_RATE_INDEX], 0, 4))
 *   tau_eff      = base_clamped × drift_factor
 *
 * The drift signal is produced by `q_drift_rate_ema_update` once per epoch
 * boundary BEFORE this kernel fires. With drift saturated at 4, tau_eff
 * drops to `tau_base / 5` (5× slower target sync); with drift = 0,
 * `drift_factor = 1` and the schedule passes through unchanged. Monotone —
 * tau only ever decreases under drift; healthy runs are unaffected.
 *
 * Per `feedback_no_partial_refactor.md`: when ISV[129] is freshly written by
 * the producer kernel each epoch boundary, every consumer of the tau output
 * (Polyak EMA target sync) sees the dampened value in the same step. Plan C
 * fold 0 ep2's predicted impact: `q_mean(t)=0.82, q_mean(t-1)=-0.018,
 * ISV[16]+ISV[21]≈0.1` → drift_rate ≈ 8.4 → clipped to 4 → drift_factor =
 * 0.2 → tau_eff = tau_base × 0.2 (5× slower target sync, dampens Q-target
 * optimism through the inflation spike).
 */
extern "C" __global__ void tau_update(
    const float* __restrict__ isv,
    float*       __restrict__ isv_out,
    float tau_base,
    float tau_final
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;
    float epoch    = isv[39];
    float total_ep = fmaxf(isv[40], 1.0f);
    float health   = fminf(fmaxf(isv[12], 0.0f), 1.0f);
    float progress = fminf(fmaxf(epoch / total_ep, 0.0f), 1.0f);
    float cos_f    = 0.5f * (1.0f + cosf(3.14159265358979f * progress));
    float tau_cos  = tau_final + (tau_base - tau_final) * cos_f;
    float floor_v  = 0.01f * (1.0f - health);
    float tau_eff  = fmaxf(tau_cos, floor_v);
    tau_eff        = fminf(fmaxf(tau_eff, tau_final), tau_base);

    /* Plan C Phase 2 follow-up A.2: adaptive drift-rate dampening.
     * Multiplicative factor in [1/5, 1] applied AFTER the [tau_final,
     * tau_base] clamp so dampening can drive tau_eff below tau_final on
     * runaway drift (this is the explicit intent — tau_final is the
     * design floor for healthy runs, but a runaway requires further
     * suppression to stop the bootstrap target chasing the online Q).
     * The producer kernel `q_drift_rate_ema_update` clips ISV[129] to
     * [0, 4]; the `fminf(..., 4.0f)` here is a defensive belt-and-braces
     * bound (Invariant 1 carve-out) in case a future producer-kernel bug
     * lets a larger value through. */
    float drift_raw    = isv[129];
    float drift_clip   = fminf(fmaxf(drift_raw, 0.0f), 4.0f);
    float drift_factor = 1.0f / (1.0f + drift_clip);
    tau_eff            = tau_eff * drift_factor;

    isv_out[42] = tau_eff;
}
