// crates/ml/src/cuda_pipeline/val_sharpe_delta_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04, A1.1): val-sharpe delta + variance canary producer.
//
// Reads a 2-element `val_sharpe_history` mapped-pinned buffer (the host
// writes [prev_epoch, curr_epoch] just before launch) and the prior
// VAL_SHARPE_DELTA_EMA_INDEX from ISV (set by Pearls A+D on the prior
// observation). Writes:
//   scratch_out[0] = (curr - prev)                       — raw delta
//   scratch_out[1] = (delta - prev_delta_ema)^2          — squared deviation
//
// Downstream `apply_pearls_ad_kernel` (n_slots=2) chained on the same
// stream smooths these into:
//   ISV[VAL_SHARPE_DELTA_EMA_INDEX=350]   (Δ EMA)
//   ISV[VAL_SHARPE_VAR_EMA_INDEX=351]     (variance EMA, true variance)
//
// Pearl A: sentinel 0 → first observation replaces directly. The squared-
// deviation against an as-yet-unwritten ISV slot is `delta * delta` on
// the first call (since `prev_delta_ema = 0` from the FoldReset sentinel),
// which Pearl A then bootstraps into `VAL_SHARPE_VAR_EMA_INDEX`.
//
// Pearl D: Wiener-α smoothing applied per slot inside apply_pearls_ad_kernel
// chained on the same stream — α adapts to each slot's own noise level.
//
// Single block, single thread; cheap arithmetic. No atomicAdd
// (`feedback_no_atomicadd.md`); `__threadfence_system()` after the writes
// to make the device-resident scratch visible to the chained Pearls launch.

#include <cuda_runtime.h>

extern "C" __global__ void val_sharpe_delta_compute_kernel(
    /* val_sharpe_history: 2-element mapped-pinned buffer.
     *   [0] = previous-epoch val sharpe (-Inf at first observation; the
     *         host writes the current value into [1] then rotates).
     *   [1] = current-epoch val sharpe.
     * Read-only.
     */
    const float* __restrict__ val_sharpe_history,
    /* Producer scratch buffer (mapped pinned, indexed by scratch_idx_base). */
    float* __restrict__ scratch_out,
    /* ISV signal bus (read-only). */
    const float* __restrict__ isv,
    /* ISV slot indices for Pearls A+D consumption. */
    int delta_ema_slot,         // = VAL_SHARPE_DELTA_EMA_INDEX = 350
    int var_ema_slot            // = VAL_SHARPE_VAR_EMA_INDEX = 351
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;

    /* val_sharpe_history layout (host writes from training_loop.rs val
     * boundary): [t-1, t]. The first observation has prev = 0.0 (the
     * mapped-pinned buffer is constructor-zero-initialised), so the
     * first emitted delta carries the absolute current val sharpe,
     * which Pearl A then sentinel-bootstraps into the EMA. */
    const float prev = val_sharpe_history[0];
    const float curr = val_sharpe_history[1];
    const float delta = curr - prev;

    /* Squared deviation against the PREVIOUS delta_ema. After Pearl A's
     * first-observation replacement this becomes the per-call deviation
     * from the running mean, suitable as a true-variance-EMA input. The
     * (void)var_ema_slot read keeps the kernel signature self-contained:
     * the slot index is consumed by the chained apply_pearls_ad launcher,
     * not by this kernel directly. */
    const float prev_delta_ema = isv[delta_ema_slot];
    const float dev = delta - prev_delta_ema;

    scratch_out[0] = delta;            // → ISV[delta_ema_slot] via Pearls A+D
    scratch_out[1] = dev * dev;        // → ISV[var_ema_slot]   via Pearls A+D

    (void)var_ema_slot;                // see comment above
    __threadfence_system();
}
