diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index b66d1e191..5b1d6ba50 100644
--- a/crates/ml/build.rs
+++ b/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`
diff --git a/crates/ml/src/cuda_pipeline/aux_to_input_concat_kernel.cu b/crates/ml/src/cuda_pipeline/aux_to_input_concat_kernel.cu
new file mode 100644
index 000000000..d7e798dbf
--- /dev/null
+++ b/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;
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 86e608e1a..45a5fa98a 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -17980,3 +17980,33 @@ The "degraded predict-Profit-everywhere" cold-start state from Phase B4 is now r
 - **Phase D**: 12-weight W atom-shift (4 actions × 3 outcomes). Largest remaining commit — extends the existing 4-weight Phase 3 atom-shift mechanism.
 - **Phase E**: dW backward + Adam for W[4, 3].
 - **Phase F**: validation smoke at structural prior (β=0.5).
+
+#### Phase B5a — Input concat kernel scaffolding (2026-05-14)
+
+Lands the input concat kernel for the plan-conditioned K=3 forward path as reusable scaffolding. Phase B5b (integration) is **deferred** with documented rationale below; Phase C (state slots) is more critical for testing the K=3 head's effect on policy behavior, and can land independently of B5b.
+
+**NEW kernel `aux_to_input_concat_kernel.cu`**:
+- Writes `[B, SH2+P]` from `h_s2_aux [B, SH2=256]` || `plan_params [B, P=6]`
+- Pure GPU map; one thread per output element, no reduction, no atomicAdd
+- NULL-tolerant on `plan_params` (writes zeros for the trailing P columns when source unavailable — needed because the collector path has no plan_params source)
+- Registered in `build.rs::kernels_with_common`; cubin compiles (5.7 KB)
+- Dead code at this commit — no Rust launcher yet
+
+**Why B5 is split + B5b deferred**:
+
+The full Phase B5 (integration) requires three coordinated changes:
+1. **Forward path**: bump `aux_to_fwd.forward()` calls to pass SH2=262 + new 262-dim input buffer
+2. **Backward stride mismatch**: backward kernel emits `dh_s2_aux_to_buf [B, 262]`, but `dh_s2_aux_accum` (input to aux trunk backward) is sized `[B, 256]`. A direct SAXPY of `n = B * 262` from `[B, 262]` to `[B, 256]` would mismatch row strides (`262 vs 256`) and corrupt the aux trunk's upstream gradient.
+3. **Collector-path plan_params unavailability**: the trade plan head only runs in the trainer path; the collector path's rollout-step forward has no plan_params source. Workarounds: (a) zero-fill, (b) add trade plan launch to collector (significant scope), (c) skip K=3 forward in collector entirely.
+
+Phase B5b would need: (1) a new strided-SAXPY kernel for the dh_s2_aux row-truncated accumulate, and (2) a decision on collector-path plan_params handling. Both are real engineering work but **not on the critical path** for testing the K=3 head's effect.
+
+**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 head's 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 the head's predictions don't reach the policy → WR signal isn't a function of the K=3 head 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.
+
+Verification: `cargo check -p ml` clean. Lib test suite unchanged from B4b-2.