/* trade_attempt_rate_ema_update — GPU-driven Flat→Positioned transition rate EMA.
 *
 * Plan 3 Task 3 B.2, spec §4.B.2. GPU-drives-CPU-reads (spec §4.C.6).
 *
 * Reads `flat_to_pos_flags [n_samples]` (0/1 per (i,t) slot, written by
 * experience_env_step at the entering_trade check), reduces the sum on-GPU
 * to `rate = count / n_samples`, then applies an adaptive-rate EMA into
 * ISV[TRADE_ATTEMPT_RATE_EMA_INDEX] using:
 *
 *   alpha = alpha_base * (1 + 0.5 * |clamp(isv[sharpe_ema_idx], -2, 2)|)
 *
 * Higher absolute Sharpe → faster update so the novelty estimate tracks
 * regime shifts promptly. Sharpe ∈ [-2, 2] keeps alpha bounded in
 * [alpha_base, 2 * alpha_base] regardless of the underlying training signal.
 *
 * Single-block, single-thread kernel — cold-path (once per epoch). The
 * reduction loop stays GPU-side per `feedback_no_atomicadd.md` (no atomicAdd)
 * and `feedback_gpu_cpu_roundtrip.md` (host never sees the flag buffer).
 */

extern "C" __global__ void trade_attempt_rate_ema_update(
    const int* __restrict__ flat_to_pos_flags, /* [n_samples] 0/1 per-(i,t) flag */
    int          n_samples,
    float*       __restrict__ isv_out,          /* ISV bus [ISV_TOTAL_DIM] */
    int          trade_attempt_rate_ema_idx,    /* TRADE_ATTEMPT_RATE_EMA_INDEX = 71 */
    int          sharpe_ema_idx,                /* SHARPE_EMA_INDEX = 22 */
    float        alpha_base                     /* base EMA rate (e.g. 0.05) */
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    /* Sum the per-sample flags on-GPU (no atomicAdd, no DtoH). */
    int count = 0;
    for (int idx = 0; idx < n_samples; ++idx) {
        count += flat_to_pos_flags[idx];
    }
    const float rate   = (float)count / fmaxf(1.0f, (float)n_samples);
    const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx]));
    const float alpha  = alpha_base * (1.0f + 0.5f * fabsf(sharpe));
    const float prev   = isv_out[trade_attempt_rate_ema_idx];
    isv_out[trade_attempt_rate_ema_idx] = (1.0f - alpha) * prev + alpha * rate;
}
