/* epsilon_update — effective exploration epsilon, GPU-driven (Plan 1 Task 14).
 * Reads:  ISV[EPOCH_IDX_INDEX=39], ISV[TOTAL_EPOCHS_INDEX=40],
 *         ISV[LEARNING_HEALTH_INDEX=12]
 * Writes: ISV[EPSILON_EFF_INDEX=41]
 * Cold path (per-epoch boundary). Single-thread kernel.
 *
 * Formula: cosine schedule (eps_start -> eps_end) over training, plus
 * a volatility boost of up to +0.1 when health=0 (full collapse).
 */
extern "C" __global__ void epsilon_update(
    const float* __restrict__ isv,
    float*       __restrict__ isv_out,
    float eps_start,
    float eps_end
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;
    float epoch    = isv[39];
    float total_ep = fmaxf(isv[40], 1.0f);
    float health   = fminf(fmaxf(isv[12], 0.0f), 1.0f);
    float progress = fminf(fmaxf(epoch / total_ep, 0.0f), 1.0f);
    float cos_f    = 0.5f * (1.0f + cosf(3.14159265358979f * progress));
    float eps_cos  = eps_end + (eps_start - eps_end) * cos_f;
    float boost    = (1.0f - health) * 0.1f;  /* up to +0.1 when health=0 */
    float eps_eff  = eps_cos + boost;
    eps_eff        = fminf(fmaxf(eps_eff, 0.0f), 1.0f);
    isv_out[41] = eps_eff;
}
