// crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu
//
// SP4 Layer A Task A6 — atom_pos_p99 producer × 4 branches.
//
// Single-block, 256-thread, parameterised producer kernel. Reads ONE branch's
// slice of `atom_positions_buf` (canonical layout `[4 * num_atoms]` with
// branch `b` at offset `b * num_atoms`, mirroring `atoms_update_kernel.cu`'s
// `positions_out + (long long)branch * num_atoms` write contract). Computes
// p99(|atom_positions[branch]|) via the header-only `sp4_histogram_p99<256>`
// device function (Task A4) and writes the per-step observation to
// `producer_step_scratch_buf[scratch_idx]` with `__threadfence_system()`
// so the host pinned-mapped read sees the write.
//
// The host launcher (`GpuDqnTrainer::launch_sp4_atom_pos_p99_all_branches`)
// loops `branch ∈ 0..4`, launching this same kernel four times with distinct
// `(branch_buf, scratch_idx)` pairs (scratch slots `1..5`, each → ISV slot
// `ATOM_POS_BOUND_BASE + branch ∈ {132, 133, 134, 135}`). Then a single
// GPU `apply_pearls_ad_kernel` launch (n_slots=4, scratch [1..5) → ISV
// [132..136) contiguously) applies Pearls A+D on the same stream — no
// host sync (2026-05-01 GPU-Pearls refactor, graph-capture-compatible).
//
// The kernel signature is identical in shape to `target_q_p99_update`
// (Task A5); the launcher's per-branch loop is the only structural
// addition over A5.

#include "sp4_histogram_p99.cuh"

extern "C" __global__ void atom_pos_p99_update(
    const float* __restrict__ atom_positions_branch,  /* slice for this branch: [count] */
    int          count,
    float*       __restrict__ scratch_buf,
    int          scratch_idx
) {
    if (blockIdx.x != 0) return;
    float p99 = sp4_histogram_p99<256>(atom_positions_branch, count);
    if (threadIdx.x == 0) {
        scratch_buf[scratch_idx] = p99;
        __threadfence_system();   // make host-mapped write visible
    }
}
