// crates/ml/src/cuda_pipeline/aux_to_input_concat_kernel.cu
//
// SP22 H6 vNext Phase B5a (2026-05-14) — input concat for the trade-
// outcome aux head's plan-conditioned forward path.
//
// Produces a 262-dim input by concatenating:
//   `h_s2_aux  [B, SH2 = 256]`  (aux trunk output) ||
//   `plan_params [B, P = 6]`     (trade plan head output)
//                              → `out      [B, SH2 + P = 262]`
//
// Both inputs are row-major. Output is row-major with the same B dim.
// For each row b ∈ [0, B):
//   out[b * (SH2+P) + j] = h_s2_aux[b * SH2 + j]            for j ∈ [0, SH2)
//   out[b * (SH2+P) + j] = plan_params[b * P + (j - SH2)]   for j ∈ [SH2, SH2+P)
//
// Phase B5a status
// ────────────────
// Kernel is REGISTERED (cubin compiled + embedded) but has no Rust
// launcher yet. Phase B5b will:
//   (a) Allocate `aux_to_input_buf [B, SH2+P]` on trainer + collector
//   (b) Launch this concat kernel before each aux head forward call
//   (c) Flip the forward + backward kernel launches to use SH2_TOTAL =
//       SH2 + P (the kernel reads SH2 as a runtime arg so no surgery)
//   (d) Resize `aux_dh_s2_to_buf` to [B, SH2+P] and use a strided SAXPY
//       to feed only the first SH2 columns into `dh_s2_aux_accum [B,
//       SH2]` (plan_params gradient is stop-grad — not fed back through
//       the trade plan head)
//   (e) Bump `sizes[163]` and `fan_dims[163]` to grow W1 from
//       [H, SH2] to [H, SH2+P]
//
// Plan-conditioned semantics
// ──────────────────────────
// Per the spec at `docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-
// aux.md`: the K=3 aux head answers "will THIS trade — with these plan
// parameters — close as Profit / Stop / Timeout?" The 6 plan_params are
// {target_bars, profit_target, stop_loss, conviction, asymmetry,
// trail_distance} per the trade plan head's output ABI. Without
// plan_params the head answers "what's the typical trade outcome at this
// market state, regardless of which trade is being proposed" — less
// informative but still a useful baseline.
//
// Discipline
// ──────────
// - Pure GPU map; one thread per output element, no reduction, no
//   atomicAdd per `feedback_no_atomicadd.md`
// - GPU-only per `feedback_cpu_is_read_only.md`
// - No host branches in captured graph per
//   `pearl_no_host_branches_in_captured_graph.md`
// - NULL-tolerant on `plan_params` only — when NULL the kernel writes
//   zeros for the trailing P columns (degraded but well-defined, mirrors
//   the trainer's `plan_params_buf` cold-start at alloc_zeros and the
//   collector's "no plan-params source available" state). When
//   `h_s2_aux` is NULL the kernel returns without writing anything (the
//   forward path is undefined without h_s2_aux).
//
// Launch
// ──────
// grid = (ceil((B * (SH2+P)) / 256)), block = (256). One thread per
// output element. The thread maps idx → (b, j) via integer div/mod and
// reads from h_s2_aux or plan_params based on `j < SH2`.

#include <cuda_runtime.h>

extern "C" __global__ void aux_to_input_concat(
    /* `[B, SH2]` row-major aux trunk output. */
    const float* __restrict__ h_s2_aux,
    /* `[B, P]` row-major plan_params. NULL-tolerant — when 0/null the
     * kernel writes zeros for the trailing P columns. */
    const float* __restrict__ plan_params,
    int B,
    int SH2,
    int P,
    /* `[B, SH2 + P]` row-major output. */
    float* __restrict__ out
) {
    if (h_s2_aux == NULL || out == NULL) return;
    const int total_cols = SH2 + P;
    const long long total = (long long)B * (long long)total_cols;
    long long idx = (long long)blockIdx.x * (long long)blockDim.x
                  + (long long)threadIdx.x;
    if (idx >= total) return;

    /* Decompose flat idx → (b, j) via integer div/mod. Both div and mod
     * are cheap at compile-time for known total_cols, but we use the
     * runtime arg for kernel generality (SH2 and P can vary across
     * future configs). */
    const int b = (int)(idx / (long long)total_cols);
    const int j = (int)(idx % (long long)total_cols);

    float v;
    if (j < SH2) {
        /* First SH2 columns: h_s2_aux. */
        v = h_s2_aux[(long long)b * (long long)SH2 + (long long)j];
    } else {
        /* Trailing P columns: plan_params (NULL-tolerant). */
        if (plan_params != NULL) {
            v = plan_params[(long long)b * (long long)P + (long long)(j - SH2)];
        } else {
            v = 0.0f;  /* degraded cold-start: no plan signal → zero column */
        }
    }
    out[idx] = v;
}
