/* kelly_cap_update — effective Kelly cap, GPU-driven (Plan 1 Task 11).
 * Reads:  ISV[LEARNING_HEALTH_INDEX=12] (safety multiplier coupled to health)
 *         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 (matching trade_physics.cuh):
 *   prior_wins=2, prior_losses=2, prior_sum_wins=0.01, prior_sum_losses=0.01
 * This ensures the cap never collapses to zero on cold start.
 */
#include "state_layout.cuh"

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;

    const float prior_wins       = 2.0f;
    const float prior_losses     = 2.0f;
    const float prior_sum_wins   = 0.01f;
    const float prior_sum_losses = 0.01f;

    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);
        kelly_f          = fmaxf(fminf(kelly_f, 1.0f), 0.0f);

        /* 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;
}
