// crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu
//
// SP5 Task A2: Pearl 3 — per-branch NoisyNet sigma scales with per-branch Q magnitude.
// Reads ATOM_V_HALF[b] (Pearl 1 from Task A1) + BRANCH_ENTROPY[b] (q_branch_stats from
// Task A1) + current SIGMA_FRACTION[b] (Pearl A sentinel 0 -> bootstrap to 0.1).
// Writes per-branch (NOISY_SIGMA, SIGMA_FRACTION) to scratch_out.
// apply_pearls_ad_kernel then smooths these via Pearls A+D into ISV.
//
// Controller: SIGMA_FRACTION[b] adapts via entropy-deficit:
//   target_entropy = log(action_counts[b]) * 0.7  (Invariant 1 anchor: 70% of max)
//   SF[b] += SF_LR * (target_entropy - actual_entropy[b])
// New sigma = SF * max(v_half, EPS_CLAMP_FLOOR).
//
// 4-thread single-block (one thread per branch).
// No atomicAdd (feedback_no_atomicadd). __threadfence_system() after writes.
//
// SF_LR=0.005 (slow) and target_entropy_factor=0.7 are Invariant 1 structural anchors.
// Envelope [0.001, 0.5] prevents sigma runaway in either direction.

#include <cuda_runtime.h>

extern "C" __global__ void pearl_3_sigma_update(
    const float* __restrict__ isv_signals,
    int   v_half_isv_base,             // ATOM_V_HALF_BASE = 178
    int   branch_entropy_isv_base,     // BRANCH_ENTROPY_BASE = 218
    int   sigma_fraction_isv_base,     // SIGMA_FRACTION_BASE = 214
    const int*   __restrict__ action_counts,  // [4] = {4, 3, 3, 3}
    float* __restrict__ scratch_out,
    int   sigma_idx_base,              // scratch_out base for sigma[4]
    int   sigma_fraction_idx_base      // scratch_out base for SF[4]
) {
    int b = threadIdx.x;
    if (blockIdx.x != 0 || b >= 4) return;

    // EPS_CLAMP_FLOOR = 1.0 is an Invariant 1 anchor (structural lower bound on
    // v_half when Pearl 1 has not yet produced a real observation).
    const float EPS_CLAMP_FLOOR = 1.0f;

    float v_half = fmaxf(isv_signals[v_half_isv_base + b], EPS_CLAMP_FLOOR);
    float entropy = isv_signals[branch_entropy_isv_base + b];
    float sigma_fraction = isv_signals[sigma_fraction_isv_base + b];
    // Pearl A sentinel: ISV slot initialised to 0 at fold boundary.
    // First observation replaces directly with the bootstrap value.
    if (sigma_fraction <= 0.0f) sigma_fraction = 0.1f;

    int n_actions = action_counts[b];
    // target_entropy_factor=0.7 is an Invariant 1 structural anchor (70% of max entropy).
    float target_entropy = logf((float)n_actions) * 0.7f;

    // SF_LR=0.005 is an Invariant 1 structural anchor (slow controller).
    const float SF_LR = 0.005f;
    float new_sf = sigma_fraction + SF_LR * (target_entropy - entropy);
    // Structural envelope [0.001, 0.5] prevents sigma runaway in either direction.
    new_sf = fmaxf(0.001f, fminf(0.5f, new_sf));

    float new_sigma = new_sf * v_half;

    scratch_out[sigma_idx_base + b]          = new_sigma;
    scratch_out[sigma_fraction_idx_base + b] = new_sf;
    __threadfence_system();
}
