/* kelly_cap_update — effective Kelly cap, GPU-driven (Plan 1 Task 11).
 * Reads:  ISV[LEARNING_HEALTH_INDEX=12] (safety multiplier coupled to health)
 *         ISV[KELLY_PRIOR_*_INDEX=454..458) (Class A P1-Producer adaptive
 *           Bayesian priors — slow-EMA from realized fold trades)
 *         portfolio_state[n_envs, PS_STRIDE] — Kelly win/loss stats per env
 * Writes: ISV[KELLY_CAP_EFF_INDEX=47] (shifted +3 from 44 by D.2 per-branch gamma, Plan 2 Task 3)
 * Cold path (per-epoch boundary). Single-thread kernel.
 *
 * Formula: average Kelly fraction across all envs using Bayesian priors
 * (same formula as trade_physics.cuh kelly_position_cap), then multiply by
 * health-coupled safety multiplier. Result is a scalar cap in [0, 1].
 * Downstream kernel readers use this cap to limit position sizing.
 *
 * Bayesian priors (Class A P1-Producer 2026-05-08): now read from
 * ISV[454..458) where they are produced by `kelly_bayesian_priors_update_kernel`
 * at the SAME per-epoch boundary, launched RIGHT BEFORE this kernel so the
 * cold-path Kelly cap sees the freshly-blended priors. Sentinels match
 * the pre-P1-Producer hardcoded constants
 * (2.0 / 2.0 / 0.01 / 0.01) for bit-identical cold-start behavior. The
 * fallback path below applies when an ISV slot is at sentinel (Pearl-A
 * pre-bootstrap) or out of the dimensional-safety range — defensive only.
 */
#include "state_layout.cuh"

/* SP9 Fix 37 (2026-05-08 Class A wiring fix): the warmup floor `isv[330]` was
 * already produced by `kelly_warmup_floor_compute_kernel` and consumed in the
 * per-step path at `trade_physics.cuh:377-384`, but this cold-path epoch-boundary
 * kernel was hardcoding `0.0f` as the kelly_f floor at line 53 below. Result:
 * sub-breakeven WR (typical cold-start at 46-48%) → kelly_f deterministically
 * collapses to 0 → kelly_cap collapses to 0 → max position pinned to Quarter
 * for ALL 56 epochs of train-xmd6b. Per `project_magnitude_eval_collapse_kelly_capped`,
 * this is the val-mag pathology. Fix: replace `0.0f` floor with isv[330].
 *
 * Slot constant mirrors `sp5_isv_slots.rs::KELLY_WARMUP_FLOOR_INDEX = 330`. */
#define SP9_KELLY_WARMUP_FLOOR_INDEX 330

/* Class A P1-Producer cold-start fallback helper. Returns the ISV slot value
 * if it is within the dimensional-safety range, else the static default
 * (the pre-P1-Producer hardcoded value for bit-identical cold-start). */
__device__ __forceinline__ float kelly_prior_or_default(
    float isv_val,
    float fallback_default,
    float lo_bound,
    float hi_bound
) {
    if (isv_val >= lo_bound && isv_val <= hi_bound) {
        return isv_val;
    }
    return fallback_default;
}

extern "C" __global__ void kelly_cap_update(
    const float* __restrict__ isv,
    float*       __restrict__ isv_out,
    const float* __restrict__ portfolio_state,   /* [n_envs * PS_STRIDE] */
    int n_envs,
    int ps_stride
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    /* Class A P1-Producer (2026-05-08): adaptive Bayesian priors from
     * ISV[454..458). Producer kernel `kelly_bayesian_priors_update_kernel`
     * runs at the same per-epoch boundary RIGHT BEFORE this kernel.
     * Cold-start path: ISV slot at sentinel (Pearl-A pre-bootstrap) or
     * out-of-range → fall back to the static default (matches pre-P1-Producer
     * hardcoded values for bit-identical cold-start behavior). */
    const float prior_wins       = kelly_prior_or_default(
        isv[KELLY_PRIOR_WINS_INDEX],       KELLY_PRIOR_WINS_DEFAULT,       0.5f, 100.0f);
    const float prior_losses     = kelly_prior_or_default(
        isv[KELLY_PRIOR_LOSSES_INDEX],     KELLY_PRIOR_LOSSES_DEFAULT,     0.5f, 100.0f);
    const float prior_sum_wins   = kelly_prior_or_default(
        isv[KELLY_PRIOR_SUM_WINS_INDEX],   KELLY_PRIOR_SUM_WINS_DEFAULT,   0.001f, 1.0f);
    const float prior_sum_losses = kelly_prior_or_default(
        isv[KELLY_PRIOR_SUM_LOSSES_INDEX], KELLY_PRIOR_SUM_LOSSES_DEFAULT, 0.001f, 1.0f);

    float health = fminf(fmaxf(isv[12], 0.0f), 1.0f);
    /* Safety multiplier: 0.5 + 0.5*health — collapses cap during degradation */
    float safety_mult = 0.5f + 0.5f * health;

    float kelly_sum = 0.0f;
    int count = 0;
    for (int e = 0; e < n_envs; ++e) {
        const float* ps = portfolio_state + (long long)e * ps_stride;
        float win_count  = ps[PS_KELLY_WIN_COUNT];
        float loss_count = ps[PS_KELLY_LOSS_COUNT];
        float sum_wins   = ps[PS_KELLY_SUM_WINS];
        float sum_losses = ps[PS_KELLY_SUM_LOSSES];

        float eff_wins   = win_count  + prior_wins;
        float eff_losses = loss_count + prior_losses;
        float eff_total  = eff_wins + eff_losses;
        float win_rate   = eff_wins / eff_total;
        float avg_win    = (sum_wins   + prior_sum_wins)   / eff_wins;
        float avg_loss   = (sum_losses + prior_sum_losses) / eff_losses;
        float payoff     = avg_win / fmaxf(avg_loss, 0.0001f);
        float kelly_f    = (payoff * win_rate - (1.0f - win_rate)) / fmaxf(payoff, 0.0001f);
        /* Class A P0-B (2026-05-08): floor was hardcoded 0.0f, collapsing kelly_f
         * to 0 deterministically when WR < breakeven. Now reads adaptive warmup
         * floor from ISV[330] (SP9 Fix 37 producer, consumed by per-step path
         * at trade_physics.cuh:377-384, NOT here until this fix). */
        float warmup_floor = isv[SP9_KELLY_WARMUP_FLOOR_INDEX];
        kelly_f          = fmaxf(fminf(kelly_f, 1.0f), warmup_floor);

        /* Half-Kelly for risk management */
        kelly_f *= 0.5f;
        kelly_sum += kelly_f;
        ++count;
    }

    float kelly_mean = (count > 0) ? (kelly_sum / (float)count) : 0.5f;
    float kelly_cap  = fminf(fmaxf(kelly_mean * safety_mult, 0.0f), 1.0f);
    isv_out[47] = kelly_cap;
}
