/* update_utilization_ema — GPU-only adaptive-α EMA for atom utilization.
 *
 * SP4 Layer C close-out C3 (2026-05-01). Replaces the host-side arithmetic
 * block in `GpuDqnTrainer::reduce_current_q_stats`
 * (gpu_dqn_trainer.rs ~line 18234) per `feedback_no_cpu_compute_strict`.
 * Mirrors `update_iqn_readiness_kernel.cu` shape — single-thread,
 * single-block, `__threadfence_system()` writeback to mapped-pinned scalars.
 *
 * Reads:  atom_util             — host-passed scalar (current epoch's atom
 *                                 utilization from q_stats_kernel reduction
 *                                 readback at host[6]).
 *         util_ema_dev[0]       — previous EMA (mapped-pinned scalar).
 * Writes: util_ema_dev[0]       — updated EMA.
 *         homeostatic_obs_dev[1] — homeostatic observable mirror (host
 *                                 previously wrote `*self.utilization_ema`
 *                                 to this slot; kernel keeps the mirror
 *                                 fresh).
 *
 * Cold-start sentinel: if `util_ema_dev[0] > 0.99`, treat as bootstrap
 * (constructor initialised the field to 1.0 = "no observation yet"): assign
 * util directly. Otherwise apply standard adaptive-α EMA recurrence
 *   err   = |util - ema_prev|
 *   alpha = clamp(err / (err + 0.1), 0.01, 0.30)
 *   ema   = (1 - α) × ema_prev + α × util
 * Matches the deleted host-side formula bit-for-bit (same algebraic form +
 * same clamp bounds).
 *
 * Per `feedback_no_cpu_compute_strict.md`: any compute (EMA, reduction, mean)
 * belongs on GPU regardless of frequency. The cold-path scalar EMA on a
 * mapped-pinned slot qualifies — `pearl_cold_path_no_exception_to_gpu_drives.md`.
 * Per `feedback_no_atomicadd.md`: a single thread reads two scalars and
 * writes two — no atomics needed.
 *
 * Two writers to mapped-pinned slots:
 *   - util_ema_dev[0] (this kernel's own state).
 *   - homeostatic_obs_dev[1] (previously written host-side from
 *     `update_homeostatic_observables`; this kernel takes over the
 *     `[1] = utilization_ema` mirror so the host write is replaced by the
 *     kernel itself maintaining the mirror in lockstep with the EMA).
 *
 * Mapped-pinned scalars (`util_ema_dev`, `homeostatic_obs_dev`) are
 * device pointers backed by mapped-pinned host memory.
 * `__threadfence_system()` ensures the writes are PCIe-visible to the
 * mapped-pinned host_ptr accessors (`utilization_ema()`,
 * `homeostatic_obs_pinned[1]` reads).
 */

extern "C" __global__ void update_utilization_ema_kernel(
    float        atom_util,                          /* host-passed scalar */
    float*       __restrict__ util_ema_dev,          /* [1] mapped-pinned EMA */
    float*       __restrict__ homeostatic_obs_dev    /* [N_OBS] mapped-pinned;
                                                      * we touch slot 1 only */
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;

    const float prev = util_ema_dev[0];

    float new_ema;
    if (prev > 0.99f) {
        /* Bootstrap branch — constructor init sentinel. The host-side
         * formula seeded utilization_ema at 1.0 to mark "no observation
         * yet"; the first call assigns the observation directly so the EMA
         * doesn't anchor against the sentinel. */
        new_ema = atom_util;
    } else {
        const float err      = fabsf(atom_util - prev);
        const float alpha_raw = err / (err + 0.1f);
        const float alpha    = fminf(fmaxf(alpha_raw, 0.01f), 0.30f);
        new_ema              = (1.0f - alpha) * prev + alpha * atom_util;
    }

    util_ema_dev[0] = new_ema;
    /* Homeostatic mirror — slot [1] = utilization_ema. Other slots are
     * untouched (the host-side `update_homeostatic_observables` writes
     * slot [0] = q_mean and leaves [2..5] for caller-driven values). */
    homeostatic_obs_dev[1] = new_ema;

    __threadfence_system();  /* PCIe-visible to mapped-pinned host_ptr */
}
