// crates/ml/src/cuda_pipeline/q_dir_grad_p99_kernel.cu
//
// SP4 Layer C #260 (2026-05-01) — q_dir_grad_p99 producer for
// ISV[Q_DIR_GRAD_BOUND_INDEX=172].
//
// Single-block, 256-thread kernel reading the union of two sub-buffers
// (`d_value_logits` ∪ `d_adv_logits`) as one logical distribution, computes
// p99(|union|) via `sp4_histogram_p99_multi<256>`, and writes the per-step
// observation to `producer_step_scratch_buf[scratch_idx]` with
// `__threadfence_system()` so the chained `apply_pearls_ad_kernel` reads the
// fresh value.
//
// Migrated from SP1-Phase-C hardcoded formula
// `1e6 × ISV[Q_DIR_ABS_REF_INDEX].max(1.0)` per
// `feedback_isv_for_adaptive_bounds`. The same `max_abs_q` bound was applied
// to BOTH `d_value_logits` AND `d_adv_logits` at the SP1 site, so the
// migrated bound is necessarily a single ISV slot whose p99 reflects the
// joint distribution of both buffers.
//
// The host launcher (`GpuDqnTrainer::launch_sp4_q_dir_grad_p99`) builds a
// 2-element device-pointer table + 2-element count table in dedicated
// mapped-pinned buffers (`q_dir_grad_subbuf_table_buf` /
// `q_dir_grad_subbuf_counts_buf`, allocated in the constructor — separate
// from `oracle_subbuf_table_buf` to eliminate cross-launcher temporal
// coupling). Then chains the GPU `apply_pearls_ad_kernel` on the same
// stream, which writes the new x_mean to ISV[Q_DIR_GRAD_BOUND_INDEX] +
// Wiener state at the slot's scratch index plus a wiener offset of
// `(Q_DIR_GRAD_BOUND_INDEX - SP4_SLOT_BASE) * 3` floats. Graph-capture-
// compatible (no host sync; same-stream ordering with the producer kernel).
//
// Mirrors `param_group_oracle_kernel.cu`'s multi-sub-buffer signature
// convention (`unsigned long long*` ptr arrays + `int*` count array), so
// the shared `sp4_histogram_p99_multi<256>` device template can be reused.

#include "sp4_histogram_p99.cuh"

extern "C" __global__ void q_dir_grad_p99_update(
    const unsigned long long* __restrict__ sub_buf_ptrs,    /* [n_sub] device pointers */
    const int*                __restrict__ sub_counts,      /* [n_sub] per-sub-buffer element counts */
    int          n_sub,                /* always 2 in this caller */
    int          total_count,          /* sum of sub_counts (used for p99 normalisation) */
    float*       __restrict__ scratch_buf,
    int          scratch_idx
) {
    if (blockIdx.x != 0) return;
    float p99 = sp4_histogram_p99_multi<256>(sub_buf_ptrs, sub_counts, n_sub, total_count);
    if (threadIdx.x == 0) {
        scratch_buf[scratch_idx] = p99;
        __threadfence_system();   // make host-mapped write visible
    }
}
