feat(sp22-vnext): Phase B5a — input concat kernel scaffolding

Lands the plan-conditioning concat kernel for the K=3 trade-outcome
forward as reusable scaffolding. Phase B5b (integration) is deferred
with rationale: Phase C (state slots) is more critical for testing the
K=3 head's effect on policy behavior, and can land independently of
the plan-conditioning refinement.

NEW kernel aux_to_input_concat_kernel.cu:
- Writes [B, SH2+P] from h_s2_aux [B, 256] || plan_params [B, 6]
- Pure GPU map; one thread per output element, no atomicAdd
- NULL-tolerant on plan_params (zeros trailing P cols when source
  unavailable, e.g., collector cold-start where the trade plan head
  doesn't run)
- Registered in build.rs; cubin compiles (5.7 KB). Dead code at this
  commit — no Rust launcher yet.

Why B5 is split + B5b deferred:

Full Phase B5 (integration) requires three coordinated changes:
1. Forward path: bump aux_to_fwd.forward() to SH2=262 + 262-dim input
2. Backward stride mismatch: backward emits dh_s2_aux_to_buf [B, 262],
   but dh_s2_aux_accum (input to aux trunk backward) is [B, 256]. A
   direct SAXPY mismatches row strides (262 vs 256) and corrupts the
   trunk's upstream gradient. Needs a strided-SAXPY kernel.
3. Collector-path plan_params unavailability: trade plan head only runs
   trainer-side. Workarounds: zero-fill, add trade plan to collector,
   or skip K=3 forward in collector. All have trade-offs.

Phase B5b would need (1) strided-SAXPY kernel and (2) collector
plan_params decision. Real work but NOT on the critical path for
testing the K=3 head's effect on WR.

Why Phase C should land first:

The K=3 head currently trains on real labels (post-B4b) but doesn't
influence policy behavior. Phase C wires the head's softmax into state
slots [121..124) = (p_Profit, p_Stop, p_Timeout), replacing the K=2
single-slot 121 = 2*p_up - 1. WITH Phase C the policy reads aux's
outcome predictions as state features → behavior changes → testable.

Without Phase C, validation runs would show "K=3 head trains and
converges" but predictions don't reach the policy → WR signal isn't
a function of K=3 at all. We'd be testing nothing.

Recommendation: skip the full B5 for now, do Phase C next, then Phase
D (atom-shift). Phase B5b (plan-conditioning) is a refinement we add
IF Phase C/D's no-plan-params version shows promise but plateaus
below the WR ≥ 0.55 target.

Audit: docs/dqn-wire-up-audit.md Phase B5a section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

crates/ml/build.rs

@@ -833,6 +833,14 @@ fn main() {
         // the per-sample → final-grad batch-reduce step. Dead code
         // until Phase B wireup lands the Rust launcher chain.
         "aux_trade_outcome_backward_kernel.cu",
+        // SP22 H6 vNext Phase B5a (2026-05-14): input concat kernel for
+        // plan-conditioned K=3 trade-outcome forward. Writes [B, 262]
+        // from h_s2_aux [B, 256] || plan_params [B, 6]. NULL-tolerant on
+        // plan_params (zeros for the trailing P=6 cols when source
+        // unavailable, e.g., collector-path cold-start). Dead code at
+        // this commit — Phase B5b allocates buffers + wires launches +
+        // flips SH2_TOTAL=262 in forward/backward.
+        "aux_to_input_concat_kernel.cu",
         // SP22 H6 (2026-05-12): per-env p_up extractor that copies
         // `aux_softmax[env, 1]` → `prev_aux_dir_prob[env]` after each aux
         // forward. The cache buffer is read by `experience_state_gather`

crates/ml/src/cuda_pipeline/aux_to_input_concat_kernel.cu

@@ -0,0 +1,103 @@
+// 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;
+}