// crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu
//
// SP15 Phase 3.1 — r_total = α × r_quality + (1 − α) × r_discipline
// (per spec §8.2 (3.1) post-amendment-2 fix).
//
// SP15 Phase 3.2 (2026-05-06) — explicit cost in r_quality on trade-close
// events (per spec §8.2 (3.2)). The composer subtracts
// `cost_t × trade_close_indicator` from `r_quality` BEFORE the α blend so
// the gradient signal carries the same per-event cost the model is
// charged at evaluation time. `cost_t` is the precomputed per-bar cost
// produced upstream by the SP15 Phase 1.2 cost_net_sharpe kernel
// (commission + per-side spread + OFI-impact); `trade_close_indicator`
// gates the subtraction to rt_ind=1 events so non-close bars are no-ops.
// At all current call sites the cost args may be passed as
// `cost_t=0.0, trade_close_indicator=0` to preserve pre-3.2 behaviour
// (no-op bypass) — production wiring of the cost-feed lands in a
// follow-up commit per `feedback_no_partial_refactor.md`, mirroring the
// Phase 3.1 kernel-first → consumer-second precedent.
//
// ALPHA_SPLIT slot (ISV[417]) is initialized DIRECTLY to 0.5 in the
// trainer constructor (NOT derived from formula at cold-start, which
// would produce 0/0=0 from zero grad-norm EMAs). The formula
//   α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)
// only takes over once BOTH grad-norm EMAs have accumulated ≥ N_WARM
// non-zero observations. Until then, the slot remains at 0.5 and the
// composer reads the constructor-written value through the
// r_quality_discipline_split_kernel below.
//
// Two `extern "C" __global__` symbols share this single cubin (the
// established 1:1-source-to-cubin / multi-kernel-per-file pattern used
// by `baseline_kernels.cu` and `novelty_simhash_kernel.cu`):
//
//   r_quality_discipline_split_kernel — per-step composition + warm count
//     • Reads ALPHA_SPLIT_INDEX (slot 417)
//     • Reads/writes alpha_warm_count scratch buffer ([1])
//     • While `*alpha_warm_count < N_WARM`, OVERRIDES alpha to 0.5 (the
//       cold-start sentinel) regardless of what the producer kernel may
//       have written — this protects the composer from premature formula
//       activation when the grad-norm EMAs are still warming up.
//     • Writes r_total to r_total_out
//
//   alpha_split_producer_kernel — per-step ALPHA_SPLIT update from grad
//     ratio
//     • Returns immediately if `*alpha_warm_count < N_WARM` (the
//       composer alone owns the warm-count increment; the producer is a
//       silent no-op until the composer has fully warmed up)
//     • Reads GRAD_NORM_QUALITY_INDEX (slot 418), GRAD_NORM_DISCIPLINE_INDEX
//       (slot 419)
//     • Writes ALPHA_SPLIT_INDEX (slot 417) clamped to [0.05, 0.95] so
//       neither signal can totally dominate the composer
//
// Hard rules:
//   • feedback_no_atomicadd — single-thread/single-block; no reductions.
//   • feedback_isv_for_adaptive_bounds — N_WARM is the only structural
//     constant; α itself is signal-driven.
//   • feedback_first_observation_bootstrap — alpha_warm_count starts at 0;
//     the warm-up window is the EMA cold-start absorber.

#include <cuda_runtime.h>

// Mirror the named slot constants from
// `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`. Keep these literals in
// lockstep with the Rust constants — the Rust unit test
// `all_sp15_slots_fit_within_isv_total_dim` enforces ISV_TOTAL_DIM bounds
// and `sp15_slot_layout_locked` pins each named index, so any drift here
// shows up as an immediate test failure on the Rust side.
#define ALPHA_SPLIT_INDEX           417
#define GRAD_NORM_QUALITY_INDEX     418
#define GRAD_NORM_DISCIPLINE_INDEX  419

// Welford warm-up count. After both grad-norm EMAs have accumulated this
// many non-zero observations the producer kernel takes over. Until then
// the composer enforces the cold-start sentinel α=0.5.
#define N_WARM 100.0f

extern "C" __global__ void r_quality_discipline_split_kernel(
    float r_quality,
    float r_discipline,
    float cost_t,
    unsigned int trade_close_indicator,
    const float* __restrict__ isv,
    float* __restrict__ r_total_out,
    float* __restrict__ alpha_warm_count
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    // Load the constructor-written sentinel (or, post-warmup, the
    // producer-written value).
    float alpha = isv[ALPHA_SPLIT_INDEX];

    // Cold-start absorber. The composer is the SOLE owner of the warm
    // count — the producer reads it but never increments. This guarantees
    // exactly one increment per call to this kernel even though both
    // kernels alias the same buffer.
    float warm = *alpha_warm_count;
    if (warm < N_WARM) {
        alpha = 0.5f;
        *alpha_warm_count = warm + 1.0f;
    }

    // SP15 Phase 3.2 (2026-05-06): explicit cost subtraction on trade-close
    // events. `trade_close_indicator` gates the term to rt_ind=1 bars so
    // non-close bars receive a no-op. `cost_t` is the precomputed per-bar
    // cost from the Phase 1.2 cost_net_sharpe kernel (commission +
    // per-side spread + OFI-impact); the same scalar the model is charged
    // at evaluation time, so the gradient signal sees the same bill.
    float r_quality_eff = r_quality - cost_t * (float)trade_close_indicator;

    *r_total_out = alpha * r_quality_eff + (1.0f - alpha) * r_discipline;

    // Make the slot writes visible to the host on the mapped-pinned
    // buffers (same pattern used by the SP15 cost-net / dd-state kernels).
    __threadfence_system();
}

extern "C" __global__ void alpha_split_producer_kernel(
    float* __restrict__ isv,
    const float* __restrict__ alpha_warm_count
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    // Wait out the cold-start window. While the composer is overriding
    // alpha to 0.5 (warm < N_WARM), writing the formula's output to the
    // slot would be wasted work and could leak a transient value if the
    // composer somehow ran on a stale warm count — leave the slot at the
    // constructor-written 0.5 sentinel during warmup.
    if (*alpha_warm_count < N_WARM) return;

    float gn_q = isv[GRAD_NORM_QUALITY_INDEX];
    float gn_d = isv[GRAD_NORM_DISCIPLINE_INDEX];
    const float EPS = 1e-6f;

    float alpha = gn_q / (gn_q + gn_d + EPS);
    // Clamp to [0.05, 0.95] so neither head can totally dominate the
    // composer — bilateral clamp per `pearl_symmetric_clamp_audit.md`.
    alpha = fmaxf(0.05f, fminf(alpha, 0.95f));

    isv[ALPHA_SPLIT_INDEX] = alpha;

    __threadfence_system();
}
