// crates/ml/src/cuda_pipeline/eval_intent_dist_compute_kernel.cu
//
// SP9 (Fix 37, 2026-05-03): GPU producer for eval-side magnitude intent
// distribution. Replaces the host-side `read_eval_intent_magnitude_distribution()`
// at `gpu_backtest_evaluator.rs:1353` per `feedback_no_cpu_compute_strict.md`
// and `feedback_no_htod_htoh_only_mapped_pinned.md` (the prior path performed
// a DtoH `read_all()` of the entire `intent_mag_pinned` mirror — both a host-
// side reduction AND an effective DtoH-style host loop over device-mapped
// memory).
//
// Reads: `intent_mag_buf` (i32, [n_windows * max_len]), the per-step
// `mag_idx` written by the action-select kernel BEFORE the Hold/Flat
// dir_idx forcing and BEFORE Kelly/margin caps inside env_step. Layout
// indexed `w * max_len + step`. Negative entries (sentinel skip) are
// excluded — this matches the host-side semantics at the deleted site.
//
// Writes: 3 floats to `scratch_buf[scratch_base..+3)`:
//   [0] = count[Quarter] / total
//   [1] = count[Half]    / total
//   [2] = count[Full]    / total
// Downstream `apply_pearls_ad_kernel` (3 launches) smooths into
// ISV[EVAL_DIST_Q/H/F_INDEX] = ISV[336..339).
//
// Single block, 3 threads — one per magnitude bin. Each thread linearly
// scans `intent_mag_buf` once counting its bin; thread 0 also counts
// total. Cheap (intent_mag_buf is at most a few thousand entries on the
// 1-window val configuration). No atomicAdd (`feedback_no_atomicadd.md`),
// no shared-memory reduce — the parallelism is across bins, not entries.
// `__threadfence_system()` after writes so the mapped-pinned alias sees
// values before downstream apply_pearls runs on the same stream.

#include <cuda_runtime.h>

extern "C" __global__ void eval_intent_dist_compute(
    /* Eval evaluator's intent_mag_buf (i32, n_windows × max_len). Read-only.
     * NULL-tolerant: when the evaluator has not yet allocated the buffer
     * (lazy init via ensure_action_select_ready), the launcher passes 0
     * for n_total and the kernel writes [0, 0, 0] — matches the host-side
     * `Ok([0.0; 3])` fallback at the deleted site. */
    const int* __restrict__ intent_mag_buf,
    int n_total,
    /* Producer scratch buffer. */
    float* __restrict__ scratch_buf,
    int scratch_base
) {
    if (blockIdx.x != 0 || threadIdx.x >= 3) return;
    int my_bin = (int)threadIdx.x;  // 0=Quarter, 1=Half, 2=Full

    int my_count   = 0;
    int total_count = 0;

    for (int i = 0; i < n_total; i++) {
        int v = intent_mag_buf[i];
        if (v < 0) continue;            // sentinel skip — pre-allocated unwritten slots
        if (v >= 0 && v < 3) {
            if (my_bin == 0) total_count++;  // thread 0 also counts total
            if (v == my_bin) my_count++;
        }
    }

    /* Thread 0 broadcasts total via shared memory; threads 1 and 2 read it. */
    __shared__ int shared_total;
    if (my_bin == 0) shared_total = total_count;
    __syncthreads();
    int t = shared_total;

    float frac = (t > 0) ? ((float)my_count / (float)t) : 0.0f;

    scratch_buf[scratch_base + my_bin] = frac;
    __threadfence_system();
}
