// crates/ml/src/cuda_pipeline/apply_fixed_alpha_ema_kernel.cu
//
// SP13 Phase 0a P0a.T4 (2026-05-04): fixed-alpha EMA applicator.
//
// Sibling of `apply_pearls_ad_kernel.cu` (the SP4 Wiener-optimal Pearls
// A+D applicator) for cases where the EMA needs DISTINCT TIMESCALES
// instead of a Wiener-optimal blend. SP13 Phase 0a needs both a SHORT
// (α=0.3) and a SLOW (α=0.05) EMA of the aux-head directional accuracy:
// the stagnation detector compares the two to decide whether `aux_w`
// should keep ramping or relax back toward the base. Pearl D would
// collapse both EMAs to the same value because the underlying signal is
// the same and the variance ratio doesn't depend on α — so the slot
// pair would report `short - long ≈ 0` permanently and the controller
// could never read "improving" or "stalled". Fixed-α EMAs preserve the
// timescale separation by construction.
//
// First-observation bootstrap per `pearl_first_observation_bootstrap`:
// when `prev_state == sentinel` the kernel REPLACES `prev_state` with
// the new sample directly (rather than blending `(1-α)·sentinel + α·obs`
// which would carry the sentinel into the EMA forever for the
// fold's first non-zero observation). The sentinel is passed as a
// kernel arg because each consumer slot has its own (slot 373 / 374
// use 0.5 — the random-guessing baseline for binary directional
// accuracy; slot 382 uses 0.0 — the empty-batch fallback for the
// hold-rate observer). The kernel has no compile-time knowledge of
// the slot's reset semantics; the launcher passes the same sentinel
// the registry entry uses on fold reset.
//
// One launch updates N consecutive ISV slots, each with its own
// (sample[i], prev_isv[isv_offset+i]) pair. For SP13 P0a: N=1 per
// EMA (one scalar dir-acc fed into both slots 373 and 374; two
// launches in lockstep). The kernel parallelises across slots —
// each thread updates one slot — so larger N (e.g. per-batch tile
// reductions) just scales the grid; no inter-thread coordination
// needed.
//
// Single-block, BLOCK_SIZE up to 256 (sized at launch from N).
// Per-thread bootstrap test + blend; no shared memory, no atomicAdd
// (per `feedback_no_atomicadd.md`). Pure GPU compute per
// `feedback_no_cpu_compute_strict.md`. Stream-ordered with the
// producer that wrote `sample[]`; the launcher MUST issue this
// kernel on the SAME stream so the producer's `__threadfence_system()`
// orders against the read here.

#include <cuda_runtime.h>

extern "C" __global__ void apply_fixed_alpha_ema_kernel(
    /* New per-slot observations [N]. The producer kernel that wrote
     * this buffer MUST have issued `__threadfence_system()` before
     * returning, so this thread's read is well-ordered against the
     * write. */
    const float* __restrict__ sample,
    /* Number of slots to update (the launcher's grid covers `N`
     * threads). N=1 is the common SP13 case (a single dir-acc scalar
     * fanned out into two ISV slots via two launches with different α). */
    int n,
    /* Offset into the ISV array where this kernel begins writing. The
     * caller's slot indices are `[isv_offset, isv_offset + N)`. */
    int isv_offset,
    /* Fixed blend rate. SP13: 0.3 for the short EMA (slot 373), 0.05
     * for the slow EMA (slot 374) and the hold-rate EMA (slot 382). */
    float alpha,
    /* Sentinel value for first-observation bootstrap. When
     * `isv[isv_offset + i] == sentinel` (per `pearl_first_observation_bootstrap`),
     * the kernel replaces directly with `sample[i]` instead of blending.
     * 0.5 for the SP13 dir-acc EMAs (slots 373/374), 0.0 for the
     * hold-rate EMA (slot 382). */
    float sentinel,
    /* ISV array — both read (prev state) and write (new state) at
     * `[isv_offset, isv_offset + N)`. */
    float* __restrict__ isv)
{
    const int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;

    const int slot       = isv_offset + i;
    const float prev     = isv[slot];
    const float new_obs  = sample[i];

    /* Pearl A first-observation bootstrap: replace directly when the
     * stored state still matches the registry sentinel. For an EMA
     * whose sentinel is the new fold's "no-observation-yet" baseline
     * (0.5 for dir-acc, 0.0 for hold-rate), blending α·obs with the
     * sentinel would keep the sentinel's bias in the EMA across the
     * fold boundary. Direct replacement on the first real observation
     * gives the new fold a clean start. */
    const float new_state = (prev == sentinel)
        ? new_obs
        : (1.0f - alpha) * prev + alpha * new_obs;

    isv[slot] = new_state;
    __threadfence_system();
}
