/* moe_lambda_eff_update — adaptive MoE load-balance λ controller, GPU-driven.
 *
 * Reads:  ISV[MOE_GATE_ENTROPY_EMA_INDEX=126] (gate-entropy EMA, α=0.05;
 *         producer is moe_expert_util_ema_update which co-EMAs entropy with
 *         per-expert util into [118..127)).
 * Writes: ISV[MOE_LAMBDA_EFF_INDEX=128] (effective load-balance weight).
 *
 * Cold-path cadence — single-block, single-thread kernel (matches
 * `kelly_cap_update_kernel.cu` and `cql_alpha_seed_update_kernel.cu`
 * shape). Launched after `moe_expert_util_ema_update` so the entropy EMA
 * fed to the controller is fresh for the NEXT step's
 * `moe_load_balance_loss` consumer.
 *
 * Formula (per `pearl_blend_formulas_must_have_permanent_floor.md`):
 *
 *   target  = entropy_target_frac × ln(K)
 *   deficit = clamp((target − ent_ema) / max(target, 1e-6), 0, 1)
 *   λ_eff   = λ_floor + λ_max_extra × deficit
 *
 * `max(real, floor)` composition keeps λ_floor as a permanent minimum
 * (λ_eff ≥ λ_floor for all entropy values). At cold-start (ISV[128]
 * bootstrapped to λ_floor in the constructor) and at fold boundary (state
 * reset registry rewrites ISV[128] = λ_floor), the controller starts from
 * the floor and rises adaptively as gate-entropy decays below target.
 *
 * Motivation: the L40S validation run train-multi-seed-fgg9x exposed
 * single-expert specialization in fold 1 (expert 4 → 81% util, 7/8
 * experts < 5%) under static λ=0.01. The static load-balance push was too
 * weak to prevent winner-take-all once the gate had a strong preference.
 * Adaptive λ rises proportional to the entropy deficit, restoring
 * diversity pressure when the gate concentrates.
 *
 * Per `feedback_isv_for_adaptive_bounds.md`: signal flows through ISV[128];
 * consumer kernel `moe_load_balance_loss` reads the slot at runtime; no
 * DtoH on hot path.
 */

extern "C" __global__ void moe_lambda_eff_update(
    float*       __restrict__ isv,                   /* ISV bus [ISV_TOTAL_DIM] */
    int          isv_lambda_eff_idx,                 /* MOE_LAMBDA_EFF_INDEX = 128 */
    int          isv_entropy_ema_idx,                /* MOE_GATE_ENTROPY_EMA_INDEX = 126 */
    float        lambda_floor,                       /* permanent minimum (= old static 0.01) */
    float        lambda_max_extra,                   /* max additional λ at full collapse */
    float        entropy_target_frac                 /* fraction of ln(K) treated as "diverse enough" */
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    /* ln(MOE_NUM_EXPERTS) — for K=8 this is 2.07944. Hard-coded because
     * MOE_NUM_EXPERTS is a compile-time constant on the Rust side and any
     * change to K already requires a kernel recompile. */
    const float ln_K = 2.0794415f;  /* ln(8); update if K changes */

    float ent_ema = isv[isv_entropy_ema_idx];
    float target  = entropy_target_frac * ln_K;

    /* Normalised entropy deficit ∈ [0, 1]:
     *   - 0 when ent_ema ≥ target (gate diverse enough → λ_eff = λ_floor)
     *   - 1 when ent_ema = 0      (gate fully collapsed → λ_eff = λ_floor + λ_max_extra)
     *
     * `target_safe` floor of 1e-6 guards against degenerate config
     * (entropy_target_frac=0). */
    float target_safe = fmaxf(target, 1e-6f);
    float deficit = fmaxf(0.0f, target - ent_ema) / target_safe;
    deficit = fminf(1.0f, deficit);

    float lambda_eff = lambda_floor + lambda_max_extra * deficit;

    isv[isv_lambda_eff_idx] = lambda_eff;
    __threadfence_system();  /* PCIe-visible to mapped pinned host_ptr */
}
