// crates/ml/src/cuda_pipeline/sp4_histogram_p99_test_kernel.cu
//
// SP4 unit-test wrapper for `sp4_histogram_p99` (header-only device fn
// in `sp4_histogram_p99.cuh`). Single-block kernel that calls the device
// function and writes the resulting p99 to a host-mapped scalar.
//
// Test-only; no production callers. Producer kernels in Tasks A5-A9
// `#include "sp4_histogram_p99.cuh"` directly.

#include "sp4_histogram_p99.cuh"

extern "C" __global__ void sp4_histogram_p99_test_kernel(
    const float* __restrict__ buf,
    int count,
    float* __restrict__ out_p99      // [1] — pinned-mapped
) {
    if (blockIdx.x != 0) return;
    float p99 = sp4_histogram_p99<256>(buf, count);
    if (threadIdx.x == 0) {
        *out_p99 = p99;
        __threadfence_system();   // make write PCIe-visible to host pointer
    }
}
