// crates/ml/src/cuda_pipeline/kelly_warmup_floor_compute_kernel.cu
//
// SP9 (Fix 37, 2026-05-03): GPU producer for the Kelly cold-start adaptive
// warmup floor. The main link in the SP9 chain — combines the three
// cold-start exit signals (statistical / behavioral / temporal) per
// `pearl_cold_start_exit_signal_or.md` and the self-relative base_floor
// per `pearl_controller_anchors_isv_driven.md`.
//
// Math:
//   base_floor = clamp(q_var_mag / max(q_var_mag_ema, EPS_DIV),
//                      KELLY_FLOOR_MIN_RATIO, KELLY_FLOOR_MAX_RATIO)
//
//   statistical_conf = clamp(kelly_sample_count / max(sample_count_target, EPS_DIV), 0, 1)
//   behavioral_conf  = clamp(1 − intent_eval_divergence / max(divergence_target, EPS_DIV), 0, 1)
//   temporal_conf    = clamp(epoch_idx / max(temporal_target, EPS_DIV), 0, 1)
//
//   combined_confidence = max(statistical_conf, behavioral_conf, temporal_conf)
//                         (OR semantics — any single axis firing exits cold-start)
//
//   warmup_floor = base_floor × (1 − combined_confidence)
//                  (decays toward 0 as confidence accumulates)
//
// Reads (all from ISV):
//   - q_var_mag             ← ISV[FLATNESS_BASE+1=207]   (Pearl 2 magnitude flatness)
//   - q_var_mag_ema         ← ISV[Q_VAR_MAG_EMA_INDEX=331] (SP9 EMA-tracked baseline)
//   - kelly_sample_count    ← ISV[KELLY_SAMPLE_COUNT_INDEX=283] (Pearl 6 cumulative)
//   - sample_count_target   ← ISV[KELLY_SAMPLE_COUNT_TARGET_INDEX=333] (Fix 37.1 Invariant-1 anchor = 100.0)
//   - intent_eval_divergence← ISV[INTENT_EVAL_DIVERGENCE_INDEX=332] (SP9 producer)
//   - divergence_target     ← ISV[KELLY_DIVERGENCE_TARGET_INDEX=334] (Fix 37.1 Invariant-1 anchor = 2.0)
//   - epoch_idx             ← ISV[EPOCH_IDX_INDEX=39]    (existing per-fold counter)
//   - temporal_target       ← ISV[KELLY_TEMPORAL_TARGET_INDEX=335] (Fix 37.1 Invariant-1 anchor = 5.0)
//
// Writes: 1 float to `scratch_buf[scratch_idx]`. Downstream
// `apply_pearls_ad_kernel` smooths into ISV[KELLY_WARMUP_FLOOR_INDEX=330].
//
// Constants (Invariant 1 numerical-stability anchors only):
//   KELLY_FLOOR_MIN_RATIO = 0.25f   (structural envelope — base_floor never
//                                     drops below 25% of target_position)
//   KELLY_FLOOR_MAX_RATIO = 1.0f    (structural envelope — full position cap)
//   EPS_DIV               = 1e-6f   (divide-by-zero protection)
//
// Single block, single thread — cheap arithmetic. No atomicAdd. No host-side
// computation. `__threadfence_system()` after the write.

#include <cuda_runtime.h>

extern "C" __global__ void kelly_warmup_floor_compute(
    /* ISV signal bus (read-only). */
    const float* __restrict__ isv_signals,
    /* All ISV slot indices passed by the launcher so the kernel doesn't
     * need to import sp5_isv_slots constants — preserves the launcher-as-
     * single-source-of-truth pattern from train_active_frac_compute. */
    int q_var_mag_isv_index,                // FLATNESS_BASE+1 = 207
    int q_var_mag_ema_isv_index,            // Q_VAR_MAG_EMA_INDEX = 331
    int kelly_sample_count_isv_index,       // KELLY_SAMPLE_COUNT_INDEX = 283
    int sample_count_target_isv_index,      // KELLY_SAMPLE_COUNT_TARGET_INDEX = 333
    int intent_eval_divergence_isv_index,   // INTENT_EVAL_DIVERGENCE_INDEX = 332
    int divergence_target_isv_index,        // KELLY_DIVERGENCE_TARGET_INDEX = 334
    int epoch_idx_isv_index,                // EPOCH_IDX_INDEX = 39
    int temporal_target_isv_index,          // KELLY_TEMPORAL_TARGET_INDEX = 335
    /* Producer scratch buffer. */
    float* __restrict__ scratch_buf,
    int   scratch_idx
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;

    /* Invariant 1 anchors. */
    const float EPS_DIV               = 1e-6f;
    const float KELLY_FLOOR_MIN_RATIO = 0.25f;
    const float KELLY_FLOOR_MAX_RATIO = 1.0f;

    /* ── base_floor: self-relative q_var_mag ratio ─────────────────────────
     * When current mag-Q variance exceeds its historical EMA, the network
     * is learning to distinguish magnitudes — relax the floor. When below,
     * stay conservative. Per `pearl_controller_anchors_isv_driven.md`,
     * this eliminates the prior 0.5 hardcoded half-Kelly floor. */
    float q_var_mag     = isv_signals[q_var_mag_isv_index];
    float q_var_mag_ema = isv_signals[q_var_mag_ema_isv_index];
    if (q_var_mag < 0.0f) q_var_mag = 0.0f;            // defensive — variance >= 0 always
    if (q_var_mag_ema < 0.0f) q_var_mag_ema = 0.0f;
    float base_ratio = q_var_mag / fmaxf(q_var_mag_ema, EPS_DIV);
    if (base_ratio < KELLY_FLOOR_MIN_RATIO) base_ratio = KELLY_FLOOR_MIN_RATIO;
    if (base_ratio > KELLY_FLOOR_MAX_RATIO) base_ratio = KELLY_FLOOR_MAX_RATIO;
    float base_floor = base_ratio;

    /* ── statistical_conf: Kelly samples accumulated relative to target ───
     * EMA-tracked target (`pearl_kelly_cap_signal_driven_floors`) — when
     * we've seen enough trades, statistical confidence rises. */
    float kelly_sample_count   = isv_signals[kelly_sample_count_isv_index];
    float sample_count_target  = isv_signals[sample_count_target_isv_index];
    if (kelly_sample_count < 0.0f) kelly_sample_count = 0.0f;
    float statistical_conf = kelly_sample_count / fmaxf(sample_count_target, EPS_DIV);
    if (statistical_conf < 0.0f) statistical_conf = 0.0f;
    if (statistical_conf > 1.0f) statistical_conf = 1.0f;

    /* ── behavioral_conf: intent vs eval divergence under target ─────────
     * Healthy training: divergence ≈ 1 (intent matches eval) → conf rises.
     * Val-Flat-collapse: divergence ≫ target → conf drops to 0. */
    float intent_eval_divergence = isv_signals[intent_eval_divergence_isv_index];
    float divergence_target      = isv_signals[divergence_target_isv_index];
    if (intent_eval_divergence < 0.0f) intent_eval_divergence = 0.0f;
    float divergence_ratio = intent_eval_divergence / fmaxf(divergence_target, EPS_DIV);
    float behavioral_conf  = 1.0f - divergence_ratio;
    if (behavioral_conf < 0.0f) behavioral_conf = 0.0f;
    if (behavioral_conf > 1.0f) behavioral_conf = 1.0f;

    /* ── temporal_conf: time since fold reset relative to target ──────────
     * Safety net per pearl Risk 3: if statistical and behavioral both stay
     * dormant (e.g. no validation runs early in fold 0), temporal alone
     * eventually fires once epoch_idx >= temporal_target. */
    float epoch_idx       = isv_signals[epoch_idx_isv_index];
    float temporal_target = isv_signals[temporal_target_isv_index];
    if (epoch_idx < 0.0f) epoch_idx = 0.0f;
    float temporal_conf = epoch_idx / fmaxf(temporal_target, EPS_DIV);
    if (temporal_conf < 0.0f) temporal_conf = 0.0f;
    if (temporal_conf > 1.0f) temporal_conf = 1.0f;

    /* ── OR semantics: any single axis firing exits cold-start ─────────── */
    float combined_confidence = fmaxf(fmaxf(statistical_conf, behavioral_conf), temporal_conf);

    /* ── warmup_floor: decays as confidence accumulates ─────────────────── */
    float warmup_floor = base_floor * (1.0f - combined_confidence);

    /* Defensive clamp into [0, KELLY_FLOOR_MAX_RATIO] — algebraically
     * already satisfied given the upstream clamps; cheap protection. */
    if (warmup_floor < 0.0f) warmup_floor = 0.0f;
    if (warmup_floor > KELLY_FLOOR_MAX_RATIO) warmup_floor = KELLY_FLOOR_MAX_RATIO;

    scratch_buf[scratch_idx] = warmup_floor;
    __threadfence_system();
}
