// crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu
//
// SP5 Layer D Task D3 (2026-05-02): Training-metrics EMA producer kernel.
//
// Third (and final) Layer D producer kernel after D1 (PnL aggregation) and
// D2 (LearningHealth composition). Replaces the host-side EMA arithmetic
// for the three training-loop metrics that survive at the end of each epoch
// in `training_loop.rs:5039-5113` (the post-financials adversarial-regime
// block + the adaptive-DSR block) per `feedback_no_cpu_compute_strict.md`:
//
//   training_sharpe_ema   adaptive-α EMA of `epoch_sharpe`
//                         α = clamp(0.05 + 0.3 × |obs − prev|, 0, 0.5)
//                         sentinel: `_initialized` flag (host) — first
//                         observation replaces the slot directly.
//   max_dd_ema            fixed α=0.1 EMA of `epoch_max_dd`.
//                         sentinel: prev < 1e-12 ⇒ replace directly.
//   low_dd_ratio          fixed α=0.15 EMA of binary indicator
//                         `is_low = (max_dd < new_max_dd_ema × 0.5) ? 1 : 0`.
//                         Reads the *new* max_dd_ema (this kernel's own
//                         output) — thread 2 must wait until thread 1's
//                         max_dd_ema write is visible.
//                         No sentinel — initial 0.0 behaves correctly under
//                         the EMA recurrence.
//
// **Plan brief vs code reality** (per `feedback_trust_code_not_docs.md`):
// the SP5 plan §D Task D3 names the third metric `gamma_blend` and points
// at "host-aggregated eval result" — that field does not exist in the
// codebase. The actual third EMA at `training_loop.rs:5052` is
// `low_dd_ratio` (an EMA of the binary low-drawdown indicator that drives
// the adversarial-regime hysteresis at lines 5054-5070). Layer D D3 ports
// what the code does, not what the plan speculated. The slot constant is
// named after the code.
//
// Outputs (single block, 3 threads — one per metric; 3 scratch slots written):
//
//   scratch_buf[sharpe_ema_idx]  = updated training_sharpe_ema
//   scratch_buf[max_dd_ema_idx]  = updated max_dd_ema
//   scratch_buf[low_dd_ratio_idx] = updated low_dd_ratio
//
// Threading contract:
//   - Thread 0 computes training_sharpe_ema (adaptive α, in-register; no
//     dependence on the other threads).
//   - Thread 1 computes max_dd_ema (fixed α=0.1, sentinel < 1e-12).
//   - Thread 2 computes low_dd_ratio. Depends on thread 1's max_dd_ema
//     write — synchronised via __syncthreads() between passes.
//
// Cold-start sentinels:
//   - sharpe_initialized==0 ⇒ output = epoch_sharpe (replaces the slot).
//     The `_initialized` flag is host-tracked in D3 (additive mode); D4
//     will migrate the flag to a mapped-pinned scalar so the host no
//     longer needs to consult it. For now we receive it by value.
//   - max_dd_prev < 1e-12 ⇒ output = epoch_max_dd (matches the deleted
//     host formula at training_loop.rs:5043).
//   - low_dd_ratio has no host-side sentinel: starts at 0.0 from the
//     constructor (constructor.rs:710) and the EMA recurrence converges
//     correctly from there. The kernel does likewise — no special branch.
//
// Per `pearl_first_observation_bootstrap.md`: the downstream
// `apply_pearls_ad_kernel` chain provides Wiener-optimal smoothing on top
// of the kernel's host-formula EMA, with Pearl A first-observation
// replacement on the ISV slot when the corresponding wiener triple's
// `prev_x_mean=0 AND state.x_lag=0` sentinel pair fires. Both layers
// (host-formula EMA in this kernel + Pearls A+D smoothing on top) are
// preserved in lockstep with the host-side site that D4 will replace.
//
// Producer-only — D3 is additive. The host-side updates at
// `training_loop.rs:5039-5113` continue to run unchanged. D4 (the atomic
// Layer D commit) wires D1+D2+D3 to call sites in a single migration per
// `feedback_no_partial_refactor.md`. No atomicAdd
// (`feedback_no_atomicadd.md`); single-block 3-thread kernel emits one
// __threadfence_system() at the end so the mapped-pinned host_ptr / dev_ptr
// alias sees the values before the downstream apply_pearls reads them on
// the same stream.

#include <cuda_runtime.h>

extern "C" {

__global__ void training_metrics_ema_update(
    /* Raw observations (this epoch). */
    float epoch_sharpe,
    float epoch_max_dd,

    /* Host-tracked previous EMA values (will become mapped-pinned reads
     * in D4 atomic refactor; passed by value in additive D3). */
    float prev_sharpe_ema,
    float prev_max_dd_ema,
    float prev_low_dd_ratio,

    /* Sentinel flag for the sharpe EMA (== 0 ⇒ first observation; D4 will
     * derive from a mapped-pinned scalar so host doesn't track this).
     * Passed as int because cudarc's `arg(&bool_ref)` lacks a stable ABI
     * for `bool` and the kernel ABI must match the launcher exactly per
     * `feedback_cudarc_f64_f32_abi.md`. */
    int   sharpe_initialized,

    /* Output scratch slots (apply_pearls_ad reads these on the same stream). */
    float* __restrict__ scratch_buf,
    int   sharpe_ema_idx,
    int   max_dd_ema_idx,
    int   low_dd_ratio_idx
) {
    /* Single-block contract — match SP5/SP4 producer kernel grid guard. */
    if (blockIdx.x != 0) return;

    const int tid = (int)threadIdx.x;

    /* Shared scratch: thread 1's max_dd_ema output, read by thread 2.
     * Single 4-byte allocation; no shmem launch arg needed (compile-time
     * static shmem allocation). */
    __shared__ float new_max_dd_ema_shared;

    /* ── Thread 0: training_sharpe_ema (adaptive α) ─────────────────────
     * Reproduces training_loop.rs:5091-5099 verbatim:
     *   if !initialized: ema = obs; initialized = true
     *   else:
     *     err   = |obs − prev|
     *     alpha = min(0.05 + 0.3 × err, 0.5)
     *     ema   = (1 − α) × prev + α × obs
     * The `_initialized` flag transitions 0→1 on the kernel's first call
     * after a fold reset; D3 mirrors host-side via the `sharpe_initialized`
     * arg (D4 will turn this into a mapped-pinned scalar). All clamp
     * constants (0.05, 0.3, 0.5) are Invariant 1 anchors per
     * `feedback_no_quickfixes.md` — they encode the host-side formula
     * shape exactly and migrate verbatim.
     */
    if (tid == 0) {
        float new_sharpe_ema;
        if (sharpe_initialized == 0) {
            new_sharpe_ema = epoch_sharpe;
        } else {
            const float err   = fabsf(epoch_sharpe - prev_sharpe_ema);
            float alpha       = 0.05f + 0.3f * err;
            if (alpha > 0.5f) alpha = 0.5f;
            new_sharpe_ema = (1.0f - alpha) * prev_sharpe_ema + alpha * epoch_sharpe;
        }
        scratch_buf[sharpe_ema_idx] = new_sharpe_ema;
    }

    /* ── Thread 1: max_dd_ema (fixed α=0.1, sentinel<1e-12) ─────────────
     * Reproduces training_loop.rs:5043-5047 verbatim:
     *   if prev < 1e-12: ema = obs
     *   else:            ema = 0.9 × prev + 0.1 × obs
     * 1e-12 is Invariant 1 — the deleted host formula's literal bootstrap
     * sentinel (host f64 in source; cast to f32 here matches the producer
     * scratch contract). Result is also published to shmem so thread 2
     * can derive `is_low` from the *new* EMA value (host-side at
     * training_loop.rs:5050 reads `self.max_dd_ema` after the update).
     */
    if (tid == 1) {
        float new_max_dd_ema;
        if (prev_max_dd_ema < 1e-12f) {
            new_max_dd_ema = epoch_max_dd;
        } else {
            new_max_dd_ema = 0.9f * prev_max_dd_ema + 0.1f * epoch_max_dd;
        }
        scratch_buf[max_dd_ema_idx] = new_max_dd_ema;
        new_max_dd_ema_shared = new_max_dd_ema;
    }

    /* Barrier: thread 2 reads thread 1's published max_dd_ema. */
    __syncthreads();

    /* ── Thread 2: low_dd_ratio (fixed α=0.15, binary input) ────────────
     * Reproduces training_loop.rs:5050-5052 verbatim:
     *   thresh = new_max_dd_ema × 0.5
     *   is_low = (max_dd < thresh) ? 1.0 : 0.0
     *   ema    = 0.85 × prev + 0.15 × is_low
     * 0.5 is the half-EMA threshold (host-side adversarial regime
     * activation boundary at training_loop.rs:5050). All constants
     * Invariant 1 anchors — deleted host formula's literals.
     * No sentinel branch: initial 0.0 from the constructor at
     * constructor.rs:710 behaves correctly under the EMA recurrence
     * (matches host-side, which also has no special init).
     */
    if (tid == 2) {
        const float new_max_dd_ema = new_max_dd_ema_shared;
        const float thresh = new_max_dd_ema * 0.5f;
        const float is_low = (epoch_max_dd < thresh) ? 1.0f : 0.0f;
        const float new_low_dd_ratio =
            0.85f * prev_low_dd_ratio + 0.15f * is_low;
        scratch_buf[low_dd_ratio_idx] = new_low_dd_ratio;
    }

    /* Single PCIe-visible barrier covers all three scratch writes; thread
     * count is 3 so no warp-divergence cost. Mapped-pinned host_ptr aliasing
     * requires threadfence_system before the downstream apply_pearls reads
     * (matches D1 PnL kernel and D2 health composition kernel patterns). */
    __syncthreads();
    if (tid == 0) {
        __threadfence_system();
    }
}

} /* extern "C" */
