/* fold_warmup_factor_update — adaptive fold-boundary warmup signal, GPU-driven.
 *
 * Plan C Phase 2 follow-up K (2026-04-29). Companion to A.2's q-drift-rate
 * adaptive Polyak-tau dampening. Single ISV signal driving BOTH lr_eff and
 * clip_eff at fold boundaries; both consumers are monotone (only dampen,
 * never excite).
 *
 * Reads:  current_grad_norm_ema   — fast EMA of `gr.raw_grad_norm` (α=0.1, ~10-step)
 *         steady_grad_norm_ema    — slow EMA of same source (α=0.001, ~1000-step)
 *         min_steps_for_ratio     — numerical bootstrap guard (cold-start delay)
 *         steps_observed          — step counter (host-driven; gates compute)
 * Writes: ISV[FOLD_WARMUP_FACTOR_INDEX=130]   — clamped factor ∈ [0, 1]
 *
 * Cold-path cadence — single-block, single-thread kernel matching
 * `q_drift_rate_ema_kernel.cu`, `moe_lambda_eff_kernel.cu`, and
 * `kelly_cap_update_kernel.cu` shapes. Launched once per training step AFTER
 * the per-step grad-norm EMAs have been updated by `update_grad_norm_emas_kernel.cu`
 * (the producer feeds these EMAs from `grad_norm_buf[0]` per
 * `feedback_no_cpu_compute_strict`; SP4 Layer C close-out C1 redesigned
 * 2026-05-01 migrated the prior host-side arithmetic block).
 *
 * Formula:
 *
 *   if steps_observed < min_steps_for_ratio  -> factor = 1.0  (no damping; pre-warmup)
 *   else if steady < 1e-6                    -> factor = 1.0  (slow EMA cold-start)
 *   else                                     -> factor = clamp(curr / steady, 0, 1)
 *
 * The `1e-6` denominator floor is a numerical-stability bound (Invariant 1
 * carve-out per `feedback_isv_for_adaptive_bounds.md`). The `[0, 1]` clamp
 * is structural — `factor > 1` would mean fast > slow (a transient spike) and
 * we explicitly do NOT excite past baseline (monotone-only-dampens contract,
 * matching A.2's `q_drift_rate_ema_update` kernel). After fold reset
 * (FoldReset: fast EMA → 0, ISV[130] → 0; slow EMA persists across folds),
 * curr starts at 0 → factor = 0 → heavy damping; as curr recovers toward
 * steady → factor → 1 → consumers return to baseline.
 *
 * Per `pearl_cold_path_no_exception_to_gpu_drives.md`: even cold-path scalar
 * arithmetic stays on GPU when its EMA inputs already live on GPU (mapped-
 * pinned host ptrs visible to the device).
 *
 * No atomicAdd (per `feedback_no_atomicadd.md`); a single thread reads two
 * scalars and writes one ISV slot. The `__threadfence_system()` mirrors the
 * other ISV producers' emission so the mapped-pinned host ptr sees the
 * fresh value before downstream consumers (lr/clip composers) read it.
 *
 * 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
 * 130 mirror).
 */

extern "C" __global__ void fold_warmup_factor_update(
    float*       __restrict__ isv,                      /* ISV bus [ISV_TOTAL_DIM] */
    const float* __restrict__ current_grad_norm_ema,    /* [1] fast EMA of grad norm */
    const float* __restrict__ steady_grad_norm_ema,     /* [1] slow EMA of grad norm */
    int          isv_fold_warmup_factor_idx,            /* FOLD_WARMUP_FACTOR_INDEX = 130 */
    int          min_steps_for_ratio,                   /* host counter floor */
    int          steps_observed                         /* host-passed step counter */
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    const float curr   = current_grad_norm_ema[0];
    const float steady = steady_grad_norm_ema[0];

    /* Bootstrap branches:
     *   (a) Pre-warmup window: not enough steps yet for the EMAs to be
     *       informative — let the system run freely (factor = 1).
     *   (b) Slow EMA still cold-start (~zero): same — no meaningful
     *       baseline to compare against, so emit no-op. The `1e-6` floor
     *       is the same Invariant 1 numerical-stability bound used by
     *       `q_drift_rate_ema_kernel.cu`.
     *
     * Otherwise compute the ratio and clamp to [0, 1]. The upper clamp
     * preserves the monotone-dampening contract (factor > 1 would EXCITE
     * lr/clip past baseline, which we explicitly forbid). The lower clamp
     * is structurally guaranteed (curr / steady ≥ 0 when both are non-
     * negative, which they are by construction) but kept defensively. */
    float factor;
    if (steps_observed < min_steps_for_ratio) {
        factor = 1.0f;
    } else if (steady < 1e-6f) {
        factor = 1.0f;
    } else {
        const float raw = curr / steady;
        factor = fmaxf(0.0f, fminf(1.0f, raw));
    }

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