// crates/ml/src/cuda_pipeline/aux_outcome_ce_ema_kernel.cu
//
// SP22 H6 vNext Phase F-3 (2026-05-14) — K=3 trade-outcome aux head's
// sparse-CE EMA producer for observability during smoke runs.
//
// Reads `aux_to_loss_scalar_buf[0]` (the K=3 head's per-batch mean CE
// scalar, written by `aux_trade_outcome_loss_reduce` at the end of
// `aux_heads_forward`) and EMA-blends into
// `ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]` with α=0.05. Pearl A
// first-observation bootstrap: when the ISV slot is at sentinel 0.0,
// the first non-zero loss observation replaces directly instead of
// EMA-blending (matches every other ISV EMA producer's cold-start
// contract per `pearl_first_observation_bootstrap`).
//
// Why a dedicated kernel (not extension of `aux_heads_loss_ema_update`)
// ─────────────────────────────────────────────────────────────────────
// The existing K=2 / K=5 aux-loss EMA writes through the Pearls A+D
// 2-stage producer scratch buffer. Extending it would require
// allocating a new scratch slot, threading the Pearls A+D mapping, and
// touching `apply_pearls_ad_kernel` (multi-slot consumer). The K=3 head's
// CE is for OBSERVABILITY ONLY in this commit (Phase F-3 enables Phase
// F smoke run visibility); no Pearls A+D adaptive α is required. A
// dedicated single-thread direct-to-ISV kernel is the minimum-scope
// path that still gives the observability we need.
//
// If the K=3 CE EMA later becomes a controller anchor (e.g., for a
// future SP11-style adaptive weight on aux-aligned reward), it can be
// re-routed through Pearls A+D at that point. Spec dictates this is
// observability-only for now.
//
// Discipline
// ──────────
// - Single block, single thread (kernel is O(1) — one load, one store).
// - No HtoD per `feedback_no_htod_htoh_only_mapped_pinned.md`.
// - Pearl A bootstrap per `pearl_first_observation_bootstrap.md`.
// - No atomicAdd per `feedback_no_atomicadd.md` (single thread → no
//   contention).
//
// Launch: grid=(1,1,1), block=(1,1,1). One thread per training step
// (mirrors `aux_heads_loss_ema_update`'s cadence).
//
// Cost: ~0.5 µs (single load + EMA + store). Negligible.

#include <cuda_runtime.h>

extern "C" __global__ void aux_outcome_ce_ema_update(
    /* [1] scalar mean CE from `aux_trade_outcome_loss_reduce`. */
    const float* __restrict__ aux_to_loss_scalar_buf,
    /* ISV signals bus pointer. NULL-tolerant (test scaffolds without
     * ISV wiring) — kernel returns without writing. */
    float* __restrict__ isv_signals,
    /* AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538. Passed runtime so the
     * kernel stays decoupled from SP slot drift (mirrors the existing
     * aux EMA kernel's idx-arg pattern). */
    int          ce_ema_idx,
    /* EMA blend factor. Fixed at 0.05 per `pearl_adaptive_moe_lambda`
     * for a slow-moving observability signal (~20-step time constant
     * matches one epoch's worth of batches at typical sizes). Passed
     * runtime for testability. */
    float        alpha
) {
    if (blockIdx.x != 0) return;
    if (threadIdx.x != 0) return;
    if (isv_signals == NULL) return;

    float loss = aux_to_loss_scalar_buf[0];
    /* Defensive NaN guard — the loss reduce flooring at fmaxf(p, 1e-30)
     * prevents -log(0) = +inf, and fmaxf(valid, 1) prevents divide-by-
     * zero. A NaN here would indicate an upstream bug; skip the update
     * rather than propagating NaN into the ISV slot. */
    if (!isfinite(loss)) return;

    float ema_prev = isv_signals[ce_ema_idx];
    if (!isfinite(ema_prev)) ema_prev = 0.0f;

    /* Pearl A first-observation bootstrap: cold-start sentinel 0.0
     * → replace directly with the first non-zero observation. The
     * 1e-9 tolerance handles fp32 zero-equality without unintended
     * branch into bootstrap on legitimate-tiny-loss states (a
     * converged head producing CE ~ 0.001 would NOT trigger bootstrap
     * — only the sentinel-exactly-zero state at fold reset does). */
    float ema_new;
    if (fabsf(ema_prev) < 1e-9f && loss > 0.0f) {
        ema_new = loss;  // bootstrap
    } else {
        ema_new = (1.0f - alpha) * ema_prev + alpha * loss;
    }

    isv_signals[ce_ema_idx] = ema_new;
    __threadfence_system();
}
