// crates/ml/src/cuda_pipeline/intent_eval_divergence_compute_kernel.cu
//
// SP9 (Fix 37, 2026-05-03): GPU producer for intent vs eval Full-magnitude
// divergence — the behavioral-confidence axis of the OR'd cold-start exit
// per `pearl_cold_start_exit_signal_or.md`.
//
// SP10 (Fix 38, 2026-05-03): extended to write a second output — the eval
// Thompson selector temperature derived from the same divergence ratio.
// Per `pearl_thompson_for_distributional_action_selection.md` (amended)
// + `pearl_controller_anchors_isv_driven.md` + `pearl_blend_formulas_must_have_permanent_floor.md`,
// the temperature is `clamp(divergence/divergence_target, MIN_TEMP=0.5,
// MAX_TEMP=2.0)`. MIN_TEMP=0.5 enforces the permanent-stochasticity floor.
//
// Computes:
//   intent_f    = monitoring action_counts[Full bins] / total_action_counts
//   eval_f      = ISV[EVAL_DIST_F_INDEX]                  (Pearls A+D smoothed)
//   divergence  = max(intent_f, EPS_DIV) / max(eval_f, EPS_DIV)
//   temp        = clamp(divergence / max(div_target, EPS_DIV), 0.5, 2.0)
//
// Where the "Full bins" in the 12-bin action_counts layout are
// `exp_idx = dir*3 + mag` with mag=2 (Full): bins [2, 5, 8, 11]. Hold and
// Flat force mag=0 in the kernel (so bins 5, 11 are always 0); algebraically
// only bins 2 (Short Full) and 8 (Long Full) carry mass — but we sum all 4
// for layout-symmetry / future Hold-Full extensions.
//
// Reads:
//   - `monitoring_summary[5..17)` (12-bin action_counts as floats)
//   - `isv_signals[EVAL_DIST_F_INDEX=338]` (Pearls A+D smoothed eval Full ratio)
//   - `isv_signals[KELLY_DIVERGENCE_TARGET_INDEX=334]` (Invariant-1 anchor 2.0)
// Writes:
//   - 1 float to `scratch_buf[scratch_div_idx]`     (divergence) — downstream
//     `apply_pearls_ad_kernel` smooths into ISV[INTENT_EVAL_DIVERGENCE_INDEX=332].
//   - 1 float to `scratch_buf[scratch_temp_idx]`    (temperature) — downstream
//     `apply_pearls_ad_kernel` smooths into ISV[EVAL_THOMPSON_TEMP_INDEX=339].
//
// Pearl semantics: under healthy training (intent_f matches eval_f) the
// divergence ratio → 1; under val-Flat-collapse (intent_f >> eval_f) the
// ratio rises sharply, signalling the controller. The behavioral_conf axis
// of the warmup-floor formula maps `clamp(1 − divergence/divergence_target, 0, 1)`
// — divergence above target → confidence = 0 (cap stays closed); divergence
// at or below target → confidence rises (cap relaxes). The divergence
// target ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] is an Invariant-1 numerical
// anchor (2.0) per Fix 37.1 (`feedback_isv_for_adaptive_bounds.md`),
// constructor-written once and rewritten on FoldReset — see audit Fix 37.1
// for why the original Fix 37 EMA-tracked target was abandoned.
//
// Single block, single thread; cheap arithmetic. No atomicAdd
// (`feedback_no_atomicadd.md`); `__threadfence_system()` after the writes.

#include <cuda_runtime.h>

extern "C" __global__ void intent_eval_divergence_compute(
    /* Monitoring summary buffer (27 floats); element layout per
     * `train_active_frac_compute_kernel.cu` docstring. Read-only. */
    const float* __restrict__ monitoring_summary,
    /* ISV signal bus (read-only). */
    const float* __restrict__ isv_signals,
    int   eval_dist_f_isv_index,         // EVAL_DIST_F_INDEX = 338
    int   divergence_target_isv_index,   // KELLY_DIVERGENCE_TARGET_INDEX = 334
    /* Producer scratch buffer. */
    float* __restrict__ scratch_buf,
    int   scratch_div_idx,               // SCRATCH_SP9_INTENT_EVAL_DIVERGENCE
    int   scratch_temp_idx               // SCRATCH_SP10_THOMPSON_TEMP
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;

    /* Numerical-stability anchor (Invariant 1) — distinguishes a freshly-
     * fold-reset / never-populated `eval_f` (Pearl A sentinel = 0) from a
     * tiny-but-real `eval_f` value. Anything strictly below this threshold
     * is treated as "uninitialised" because `eval_dist` writes are
     * naturally bounded by [0, 1] from `eval_intent_dist_compute_kernel`'s
     * normalised output, and 1e-5 sits well below any realisable value
     * after even a single observation enters the Pearl A bootstrap path. */
    const float SENTINEL_THRESHOLD = 1e-5f;

    /* Sum action_counts over Full magnitude bins: exp_idx = dir*3 + 2 for
     * dir ∈ {0..4}. Bins 2 (Short Full), 5 (Hold Full), 8 (Long Full),
     * 11 (Flat Full). Hold and Flat force mag=0 so 5 and 11 are 0 in
     * practice; we sum all 4 for layout-symmetry. */
    float intent_f_count = 0.0f;
    intent_f_count += monitoring_summary[5 +  2];  // Short Full
    intent_f_count += monitoring_summary[5 +  5];  // Hold Full   (algebraically 0)
    intent_f_count += monitoring_summary[5 +  8];  // Long Full
    intent_f_count += monitoring_summary[5 + 11];  // Flat Full   (algebraically 0)

    float total_count = 0.0f;
    for (int i = 0; i < 12; i++) total_count += monitoring_summary[5 + i];

    float intent_f = (total_count > 0.0f) ? (intent_f_count / total_count) : 0.0f;
    /* Defensive clamp [0, 1] — algebraically intent_f_count <= total_count. */
    if (intent_f < 0.0f) intent_f = 0.0f;
    if (intent_f > 1.0f) intent_f = 1.0f;

    float eval_f = isv_signals[eval_dist_f_isv_index];
    if (eval_f < 0.0f) eval_f = 0.0f;
    if (eval_f > 1.0f) eval_f = 1.0f;

    /* Fix 37.1 (2026-05-03): handle Pearl A sentinel for `eval_f`. Before
     * the first val-window has populated `ISV[EVAL_DIST_F_INDEX=338]`,
     * the slot still carries its sentinel 0 (per
     * `pearl_first_observation_bootstrap.md`). Reading it as a real
     * denominator and flooring at any `EPS_DIV` produced a synthetic
     * `divergence ≈ intent_f / 1e-6 ≈ 70 000`, which the warmup-floor
     * kernel then mapped to `behavioral_conf = clamp(1 − 70000/2, 0, 1) = 0`
     * — algebraically the right answer "no behavioral signal yet", but
     * coupled to a magic-floor `EPS_DIV` and easily missed in
     * smoke logs (the synthetic 70 000 looked like a runaway).
     * Instead: explicitly emit `SENTINEL_DIVERGENCE` (large enough that
     * `divergence / divergence_target ≫ 1` regardless of target) so the
     * downstream warmup-floor kernel's `behavioral_conf` clamp lands
     * cleanly at 0 with provenance traceable in HEALTH_DIAG. Once
     * `eval_dist_compute` fires its first val-window observation,
     * Pearl A's first-observation-replacement raises `eval_f` above the
     * sentinel threshold and the real ratio takes over. */
    const float SENTINEL_DIVERGENCE = 1e6f;
    float divergence;
    if (eval_f < SENTINEL_THRESHOLD) {
        divergence = SENTINEL_DIVERGENCE;
    } else {
        divergence = intent_f / eval_f;
    }

    scratch_buf[scratch_div_idx] = divergence;

    /* SP10 (Fix 38, 2026-05-03): eval Thompson selector temperature.
     *
     * Per `pearl_thompson_for_distributional_action_selection.md` (amended),
     * `pearl_controller_anchors_isv_driven.md`, and
     * `pearl_blend_formulas_must_have_permanent_floor.md`:
     *   temp = clamp(divergence / div_target, MIN_TEMP, MAX_TEMP)
     *
     * MIN_TEMP=0.5 — Invariant 1: permanent-stochasticity floor; the eval
     *                selector is never fully deterministic.
     * MAX_TEMP=2.0 — Invariant 1: cap on exploration boost; prevents
     *                pathological all-uniform sampling.
     * EPS_DIV=1e-6 — Invariant 1: divide-by-zero protection on the target
     *                anchor (which is constructor-written to 2.0 in
     *                production but defended here so the kernel is
     *                self-contained).
     *
     * Behaviour:
     *   - Healthy (divergence ≈ target=2.0): temp ≈ 1.0 → pure Thompson
     *   - Collapsed (divergence ≫ target, e.g. 1e6 sentinel): temp = 2.0
     *     → exaggerated exploration boost (corrective during val-Flat-
     *     collapse)
     *   - Healthy & converging (divergence < target): temp = 0.5 → closer
     *     to argmax behaviour, but never fully deterministic
     */
    const float THOMPSON_TEMP_MIN = 0.5f;
    const float THOMPSON_TEMP_MAX = 2.0f;
    const float EPS_DIV           = 1e-6f;

    float div_target = isv_signals[divergence_target_isv_index];
    float temp = divergence / fmaxf(div_target, EPS_DIV);
    temp = fminf(THOMPSON_TEMP_MAX, fmaxf(THOMPSON_TEMP_MIN, temp));
    scratch_buf[scratch_temp_idx] = temp;

    __threadfence_system();
}
