// crates/ml/src/cuda_pipeline/dir_concat_qaux_kernel.cu
//
// SP14 Layer B Task B.6 (2026-05-05): Earned Gradient Flow — pre-SGEMM
// concat that forward-wires `aux_softmax_diff` into the direction Q-head
// input. Mirrors `mag_concat_qdir` precedent (experience_kernels.cu:4560).
//
// ── What it does ────────────────────────────────────────────────────────
// Concatenates h_s2 with a single aux-derived scalar to produce a wider
// scratch buffer consumed by the direction Q-head's first FC SGEMM:
//
//   out[b, 0..SH2]  = h_s2[b, 0..SH2]   (contiguous copy)
//   out[b, SH2]     = aux_softmax[b, 1] - aux_softmax[b, 0]
//
// The appended scalar is the softmax probability difference (up_prob -
// down_prob), which lies in [-1, +1] by the construction of softmax. This
// is a structural bound per `pearl_bounded_modifier_outputs_require_
// structural_activation` — no runtime clamp is needed.
//
// ── Launch-order constraint ─────────────────────────────────────────────
// Per `pearl_canary_input_freshness_launch_order`: the aux head forward
// pass MUST complete before this kernel reads aux_nb_softmax_buf. The
// graph capture orchestrator (B.10/B.11) enforces the serial dependency.
// For the GPU oracle test, aux_softmax is populated directly via mapped-
// pinned write before launch, so no ordering issue arises there.
//
// ── Pure per-thread map ─────────────────────────────────────────────────
// One thread handles one element of the output tensor (row × col). No
// reduction, no shared memory, no atomicAdd (per `feedback_no_atomicadd.md`).
//
// Grid:  ceil(B * (SH2 + 1) / 256), Block: 256.

extern "C" __global__
void dir_concat_qaux_kernel(
    const float* __restrict__ h_s2,        /* [B, SH2] trunk activation */
    const float* __restrict__ aux_softmax, /* [B, 2]   aux head softmax output */
    float* __restrict__ out,               /* [B, SH2 + 1] output scratch buffer */
    int B,
    int SH2)
{
    const int tid   = blockIdx.x * blockDim.x + threadIdx.x;
    const int total = B * (SH2 + 1);
    if (tid >= total) return;

    const int row = tid / (SH2 + 1);
    const int col = tid % (SH2 + 1);

    if (col < SH2) {
        /* h_s2 portion: direct copy */
        out[tid] = h_s2[row * SH2 + col];
    } else {
        /* col == SH2: aux softmax diff = up_prob - down_prob ∈ [-1, +1] */
        const float s0 = aux_softmax[row * 2 + 0];  /* down probability */
        const float s1 = aux_softmax[row * 2 + 1];  /* up   probability */
        out[tid] = s1 - s0;
    }
}
