// crates/ml/src/cuda_pipeline/train_active_frac_compute_kernel.cu
//
// SP8 (2026-05-03): GPU producer for train_active_frac canary signal.
//
// Computes the per-epoch fraction of training-rollout actions that picked a
// tradable direction (Long or Short, NOT Hold or Flat). This is the canary
// per `pearl_controller_anchors_isv_driven.md`: under val-Flat-collapse the
// metric trends → 0; under healthy training it trends → 1. Used by the
// downstream `loss_balance_max_budget_compute_kernel` to derive an ISV-driven
// adaptive cap on cql/c51 budgets — replacing the prior hardcoded
// `MAX_BUDGET=1.0f` regime-encoded constant in
// `loss_balance_controller_kernel.cu`.
//
// Reads: `monitoring_summary[5..17)` (the 12-bin action_counts block written
// by `monitoring_reduce` — see `monitoring_kernel.cu`). Layout per
// `monitor.action_counts[12]`:
//
//   exp_idx = dir*3 + mag, dir ∈ {Short=0, Hold=1, Long=2, Flat=3},
//   mag ∈ {Small=0, Half=1, Full=2}.
//
// Writes: 1 float to `scratch_buf[scratch_idx]`. Downstream
// `apply_pearls_ad_kernel` smooths into ISV[TRAIN_ACTIVE_FRAC_INDEX] via
// Pearl A sentinel-bootstrap + Pearl D Wiener-α. Replaces the host-side
// computation at `training_loop.rs:3392-3396` per
// `feedback_no_cpu_compute_strict.md`.
//
// Single block, single thread — the 12-element sum is trivial; no
// atomicAdd (`feedback_no_atomicadd.md`); `__threadfence_system()` after
// the write so the mapped-pinned host_ptr/dev_ptr alias sees the value
// before downstream apply_pearls runs on the same stream.

#include <cuda_runtime.h>

extern "C" __global__ void train_active_frac_compute(
    /* GPU monitoring summary buffer [27 floats]; element layout:
     *   [0..5)   = stats (mean/std/min/max/sharpe)
     *   [5..17)  = action_counts[12] as floats (kernel writes (float)t_exp[i])
     *   [17..20) = order_counts[3]
     *   [20..23) = urgency_counts[3]
     *   [23]     = total bars N
     *   [24]     = total trades
     *   [25..27) = pad
     * Read-only (`__restrict__`). */
    const float* __restrict__ monitoring_summary,
    /* Producer scratch buffer; downstream apply_pearls_ad_kernel smooths
     * scratch_buf[scratch_idx] into ISV[TRAIN_ACTIVE_FRAC_INDEX]. */
    float* __restrict__ scratch_buf,
    int   scratch_idx
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;

    /* exp_idx layout: dir*3 + mag.
     * Short = bins [0..3)  (dir=0)
     * Hold  = bins [3..6)  (dir=1, no-op — kernel forces mag=0; bins 4,5 are 0)
     * Long  = bins [6..9)  (dir=2)
     * Flat  = bins [9..12) (dir=3, no-op — kernel forces mag=0; bins 10,11 are 0)
     *
     * "Active" (per Plan C Task 5) = direction ∈ {Short, Long}; the
     * `train_active_frac` canary measures directional commitment of the
     * policy during the training rollout. Hold and Flat are excluded.
     * Mirrors host-side formula at the deleted `training_loop.rs:3392-3396`
     * site: `(dir0 + dir2) / sum(all 12 bins)`. */
    float active_count = 0.0f;
    for (int i = 0; i < 3; i++) active_count += monitoring_summary[5 + 0 + i]; // Short
    for (int i = 0; i < 3; i++) active_count += monitoring_summary[5 + 6 + i]; // Long

    float total_count = 0.0f;
    for (int i = 0; i < 12; i++) total_count += monitoring_summary[5 + i];

    /* Numerical safety: if no actions were tallied (cold rollout), output
     * 0.0 — matches host-side default at the deleted site. The downstream
     * Pearl A bootstrap will replace the slot on first observation, so
     * 0.0 is a benign sentinel for the consumer. */
    float active_frac = (total_count > 0.0f) ? (active_count / total_count) : 0.0f;

    /* Clamp into [0, 1] for safety — algebraically active <= total so this
     * is a no-op under normal data; defensive against any monitoring
     * counter overflow round-trip. */
    if (active_frac < 0.0f) active_frac = 0.0f;
    if (active_frac > 1.0f) active_frac = 1.0f;

    scratch_buf[scratch_idx] = active_frac;
    __threadfence_system();
}
