// crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu
//
// SP5 Task A5: Pearl 5 — per-branch IQN τ schedule from per-branch Q-skew.
//
// Reads q_skew_per_branch[b] from scratch_buf at skew_idx_base + b (written by
// q_skew_kurtosis_update). Computes per-branch 5-quantile τ schedule:
//   - Symmetric default: {0.05, 0.25, 0.5, 0.75, 0.95}
//   - Skewed: shift each τ in the SAME DIRECTION as the skew sign — i.e.
//     left-skewed Q (skew<0, long left tail) shifts every τ toward 0, giving
//     the IQN sampler more resolution over the long left tail; right-skewed
//     Q (skew>0) shifts τ toward 1 to oversample the long right tail. This
//     leans the quantile grid INTO the asymmetry direction, which gives more
//     IQN samples where the distribution has tail mass that the symmetric
//     5-tuple under-represents.
// Clamps every τ to the structural envelope [0.01, 0.99].
//
// 4-thread single-block (one per branch); each thread emits 5 contiguous
// τ values to scratch_out at tau_idx_base + b*5 + q.
//
// SKEW_SHIFT=0.05 is an Invariant 1 anchor (modest schedule shift; preserves
// quantile spread; max envelope reach is ±0.15 at clamped skew=±3).
// kurtosis is computed by q_skew_kurtosis_update and lives in scratch at
// ex_kurt_idx_base — reserved for future heavy-tail-aware schedule extension;
// this kernel reads only skew. The excess kurtosis slots are not consumed here.

#include <cuda_runtime.h>

extern "C" __global__ void pearl_5_iqn_tau_update(
    const float* __restrict__ scratch_buf,  // scratch buffer for reading skew
    int   skew_idx_base,                    // skew[4] written by q_skew_kurtosis_update
    float* __restrict__ scratch_out,        // scratch output (same buffer, non-overlapping)
    int   tau_idx_base                      // 20 floats: branch b, quantile q at base + b*5 + q
) {
    int b = threadIdx.x;
    if (blockIdx.x != 0 || b >= 4) return;

    float skew = scratch_buf[skew_idx_base + b];
    // SKEW_SHIFT=0.05 is an Invariant 1 anchor (modest schedule shift).
    const float SKEW_SHIFT = 0.05f;

    // Symmetric default offsets (Invariant 1 anchors — quantile interpolation grid).
    float tau_defaults[5] = {0.05f, 0.25f, 0.5f, 0.75f, 0.95f};

    for (int q = 0; q < 5; ++q) {
        float t = tau_defaults[q] + skew * SKEW_SHIFT;
        // Structural envelope [0.01, 0.99] (Invariant 1 anchor — IQN requires
        // τ strictly in (0,1); clamp prevents degenerate boundary quantiles).
        t = fmaxf(0.01f, fminf(0.99f, t));
        scratch_out[tau_idx_base + b * 5 + q] = t;
    }
    __threadfence_system();
}
