// crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu
//
// SP14 Layer B Task B.10 (2026-05-05): Earned Gradient Flow backward gate.
// Scales the wire column of `dL/dx_concat [B, SH2 + 1]` (the gradient
// flowing back FROM the direction Q-head's first FC SGEMM TO
// `aux_softmax_diff`) by `α_grad_smoothed` (read from ISV[393]).
//
// ── What this gates ─────────────────────────────────────────────────────
// dL/dx_concat shape is [B, SH2 + 1]. Columns [0..SH2) propagate back to
// `h_s2` via the existing trunk-encoder backward chain. Column SH2 is the
// gradient w.r.t. the appended `aux_softmax_diff` value, which downstream
// (B.11) propagates to the aux head's softmax → CE backward chain. THIS
// kernel multiplies ONLY column SH2 by `ALPHA_GRAD_SMOOTHED = ISV[393]`.
//
// Per the EGF pearl design (spec §B.4):
//   - α_grad ≈ 0  → wire effectively closed (Q-loss can't pull aux head)
//   - α_grad ≈ 1  → full co-training enabled
//
// Critical: dW (the Q-head's own weight gradient) is NOT touched. Q-head
// learns to USE the wire freely from Q-loss; only the gradient PROPAGATING
// BACK to aux head's representation is gated. The dW and dX SGEMMs are
// independent (`dW = dY^T × x_concat`; `dX = dY × W^T`), so scaling dx[:, SH2]
// AFTER both SGEMMs have completed leaves dW unaffected.
//
// ── ISV slot index (must match `sp14_isv_slots.rs`) ─────────────────────
// ALPHA_GRAD_SMOOTHED_INDEX = 393 (sentinel 0.0 pre-B.11 producer wires;
// when sentinel-active the wire is force-closed, the conservative state).
//
// ── Pure per-thread map ─────────────────────────────────────────────────
// One thread handles one batch row's wire-column element. No reductions,
// no shared memory, no atomicAdd (per `feedback_no_atomicadd.md`). Read of
// `isv[ALPHA_GRAD_SMOOTHED_INDEX]` is identical across all threads — safe
// concurrent read.
//
// Grid:  ceil(B / 256), Block: 256.

extern "C" __global__
void sp14_scale_wire_col_kernel(
    /* dL/dx_concat [B, SH2 + 1] — IN-PLACE scaled at column SH2 only */
    float* __restrict__ dx_concat,
    /* Global ISV bus. Reads ISV[ALPHA_GRAD_SMOOTHED_INDEX]; writes nothing. */
    const float* __restrict__ isv,
    int B,
    int SH2)
{
    /* ISV slot index — must match `sp14_isv_slots.rs::ALPHA_GRAD_SMOOTHED_INDEX`.
     * The layout fingerprint regression test catches drift. */
    const int ALPHA_GRAD_SMOOTHED = 393;

    const int b = blockIdx.x * blockDim.x + threadIdx.x;
    if (b >= B) return;

    /* Wire column lives at offset SH2 in each row of stride (SH2+1). */
    const int wire_idx = b * (SH2 + 1) + SH2;

    /* Pre-B.11 producer is unwired → ISV[393] holds sentinel 0.0 →
     * wire force-closed (gradient zeroed). Post-B.11, the value is the
     * EGF gate output ∈ [0, 1].
     *
     * Match `dqn_scale_f32_kernel`'s NaN-safe contract: 0×NaN=NaN under
     * IEEE 754; explicit branch yields a clean 0 when α=0. */
    const float alpha = isv[ALPHA_GRAD_SMOOTHED];
    dx_concat[wire_idx] = (alpha == 0.0f) ? 0.0f : (dx_concat[wire_idx] * alpha);
}
