// crates/ml/src/cuda_pipeline/pearl_1_ext_num_atoms_kernel.cu
//
// SP5 Task A8: Pearl 1-ext — per-branch C51 num_atoms from per-branch v_half.
//
// Reads ATOM_V_HALF[b] (ISV[178..182), populated by Pearl 1 in Task A1).
// Maps to discrete num_atoms ∈ {16, 32, 64} based on v_half thresholds:
//   v_half < 0.1  → 64 atoms (narrow Q range, high resolution)
//   v_half < 1.0  → 32 atoms (moderate Q range)
//   v_half ≥ 1.0  → 16 atoms (wide Q range, modest resolution)
//
// The discrete output is smoothed by Pearls A+D in the launcher's apply_pearls
// step. During transitions (e.g. v_half drifts across 1.0), the smoothed ISV
// value will be intermediate between {16, 32}; the C51 atoms_update consumer
// in Layer B rounds to the nearest valid count at consume time.
//
// 4-thread single-block (one per branch). No atomicAdd.
// __threadfence_system() after writes.
//
// Invariant 1 anchors:
//   THRESHOLD_NARROW = 0.1f       (high-resolution boundary)
//   THRESHOLD_MODERATE = 1.0f     (moderate-resolution boundary)
//   ATOMS_HIGH = 64.0f, ATOMS_MID = 32.0f, ATOMS_LOW = 16.0f
//   (atom counts are powers of 2; total cap 4 branches × 64 = 256 atoms,
//    within the HW shared-memory budget the C51 backward kernel uses.)

#include <cuda_runtime.h>

extern "C" __global__ void pearl_1_ext_num_atoms_update(
    const float* __restrict__ isv_signals,
    int   v_half_isv_base,             // ATOM_V_HALF_BASE = 178
    float* __restrict__ scratch_out,
    int   num_atoms_idx_base           // 4 outputs, scratch_out[base..base+4)
) {
    int b = threadIdx.x;
    if (blockIdx.x != 0 || b >= 4) return;

    const float THRESHOLD_NARROW   = 0.1f;   // Invariant 1 anchor: high-res boundary
    const float THRESHOLD_MODERATE = 1.0f;   // Invariant 1 anchor: moderate-res boundary
    const float ATOMS_HIGH = 64.0f;           // Invariant 1 anchor: narrow-range atom count
    const float ATOMS_MID  = 32.0f;           // Invariant 1 anchor: moderate-range atom count
    const float ATOMS_LOW  = 16.0f;           // Invariant 1 anchor: wide-range atom count

    float v_half = isv_signals[v_half_isv_base + b];

    float num_atoms;
    if (v_half < THRESHOLD_NARROW) {
        num_atoms = ATOMS_HIGH;
    } else if (v_half < THRESHOLD_MODERATE) {
        num_atoms = ATOMS_MID;
    } else {
        num_atoms = ATOMS_LOW;
    }

    scratch_out[num_atoms_idx_base + b] = num_atoms;
    __threadfence_system();
}
