/* q_drift_rate_ema_update — adaptive Polyak-tau dampening signal, GPU-driven.
 *
 * Plan C Phase 2 follow-up A.2 (researcher report 2026-04-29). Replaces the
 * static `tau_base` Polyak EMA target sync with an ISV-driven adaptive
 * `tau_eff` coupled to per-epoch q-drift rate. `tau_update_kernel.cu` reads
 * ISV[Q_DRIFT_RATE_INDEX] (this kernel's output) at every epoch boundary and
 * multiplies the cosine-scheduled tau by `1 / (1 + clip(drift, 0, 4))`,
 * yielding a monotone dampening factor in `[1/5, 1]` — tau only ever
 * decreases under drift; healthy runs are unaffected.
 *
 * Reads:  ISV[Q_ABS_REF_INDEX]      — magnitude-branch |Q| EMA (slot 16)
 *         ISV[Q_DIR_ABS_REF_INDEX]  — direction-branch |Q| EMA (slot 21)
 *         q_mean_curr               — host-passed scalar (epoch's q_mean)
 *         q_mean_prev               — host-passed scalar (previous epoch's q_mean)
 * Writes: ISV[Q_DRIFT_RATE_INDEX]   — clipped drift rate ∈ [0, 4] (slot 129)
 *
 * Cold-path cadence — single-block, single-thread kernel (matches
 * `tau_update_kernel.cu`, `kelly_cap_update_kernel.cu`, and
 * `moe_lambda_eff_kernel.cu` shape). Launched once per epoch boundary AFTER
 * the per-step ISV producers (`h_s2_rms_ema`, `aux_heads_loss_ema`, etc.)
 * but BEFORE the next epoch's `tau_update` consumer reads ISV[129].
 *
 * Formula:
 *
 *   denom    = max(|q_mean_prev|, ISV[Q_ABS_REF] + ISV[Q_DIR_ABS_REF], 1e-6)
 *   drift    = |q_mean_curr - q_mean_prev| / denom
 *   ISV[129] = min(drift, 4.0)
 *
 * The `1e-6` in `denom` is a numerical-stability bound (Invariant 1 carve-out
 * per `feedback_isv_for_adaptive_bounds.md`). The `4.0` upper clip is an
 * architectural "drift starts here" bound matching the kill-criterion's
 * 3.0× ratio threshold (in `training_loop.rs` Q-drift kill); a drift of 4
 * means the next-epoch tau_eff drops to `tau_base / 5` (5× slower target
 * sync). `feedback_no_quickfixes.md` carve-out — both bounds are structural,
 * not tuned.
 *
 * Per `pearl_cold_path_no_exception_to_gpu_drives.md`: even cold-path scalar
 * arithmetic stays on GPU when its inputs already live on GPU — ISV[16,21]
 * are produced by `q_stats_kernel.cu` and consumed here without DtoH. The
 * two host-passed q-mean scalars are CPU-born monitor outputs (schedule
 * inputs per spec §4.C.6), legal as kernel by-value arguments.
 *
 * No atomicAdd (per `feedback_no_atomicadd.md`); a single-thread reads three
 * slots and writes one. The `__threadfence_system()` matches
 * `moe_lambda_eff_update`'s emission so the mapped-pinned host_ptr sees the
 * fresh value before the next `launch_tau_update` reads ISV[129].
 *
 * Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic
 * drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold-
 * start) + reset (0.0 fold boundary) + observability (HEALTH_DIAG via slot
 * 129 mirror).
 */

extern "C" __global__ void q_drift_rate_ema_update(
    float*       __restrict__ isv,                  /* ISV bus [ISV_TOTAL_DIM] */
    int          isv_q_abs_ref_idx,                 /* Q_ABS_REF_INDEX = 16 */
    int          isv_q_dir_abs_ref_idx,             /* Q_DIR_ABS_REF_INDEX = 21 */
    int          isv_q_drift_rate_idx,              /* Q_DRIFT_RATE_INDEX = 129 */
    float        q_mean_curr,                       /* current epoch's q_mean (host-passed) */
    float        q_mean_prev                        /* previous epoch's q_mean (host-passed) */
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    const float q_abs_ref     = isv[isv_q_abs_ref_idx];
    const float q_dir_abs_ref = isv[isv_q_dir_abs_ref_idx];

    /* Denominator picks the largest of: previous q_mean magnitude (so drift
     * is meaningfully scale-invariant when q has grown), the two per-branch
     * |Q| EMAs (so the ratio uses the policy's currently-observed Q-scale,
     * matching the kill-criterion's adaptive floor), and a 1e-6 numerical
     * floor (Invariant 1 carve-out — guards against the cold-start case
     * where every term is 0 before any epoch has produced statistics). */
    const float prev_abs   = fabsf(q_mean_prev);
    const float ref_sum    = q_abs_ref + q_dir_abs_ref;
    const float denom      = fmaxf(fmaxf(prev_abs, ref_sum), 1e-6f);
    const float diff_abs   = fabsf(q_mean_curr - q_mean_prev);
    const float drift_raw  = diff_abs / denom;

    /* Clip to [0, 4]. Lower bound is structurally guaranteed (fabsf output
     * is non-negative, division by positive denom preserves sign). Upper
     * bound 4 matches the kill-criterion's 3.0× ratio threshold; beyond 4
     * the dampening factor 1/(1+drift) saturates at 0.2 = tau_base/5. */
    const float drift_clipped = fminf(drift_raw, 4.0f);

    isv[isv_q_drift_rate_idx] = drift_clipped;
    __threadfence_system();  /* PCIe-visible to mapped pinned host_ptr */
}
