diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index d5ed341f1..739d71b51 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -802,6 +802,16 @@ fn main() {
         // (Phase A2) — additional kernels for aux forward / loss /
         // backward + 12-weight W atom-shift land in subsequent commits.
         "trade_outcome_label_kernel.cu",
+        // SP22 H6 vNext Phase A3 (2026-05-13): K=3 softmax aux head forward
+        // kernel for the trade-outcome aux head. Architecture mirrors
+        // `aux_next_bar_forward` (K=2) but emits a 3-way distribution over
+        // {Profit, Stop, Timeout}. Dead code until A4 (loss reduce) + A5
+        // (backward) + Phase B (262-dim input concat with plan_params) +
+        // Phase C (state slot 121-123 wireup) + Phase D (12-weight W
+        // atom-shift) land in subsequent commits. See
+        // `docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md` for
+        // the full vNext architecture spec.
+        "aux_trade_outcome_forward_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_trade_outcome_forward_kernel.cu b/crates/ml/src/cuda_pipeline/aux_trade_outcome_forward_kernel.cu
new file mode 100644
index 000000000..ce60e045a
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/aux_trade_outcome_forward_kernel.cu
@@ -0,0 +1,167 @@
+// crates/ml/src/cuda_pipeline/aux_trade_outcome_forward_kernel.cu
+//
+// SP22 H6 vNext Phase A3 (2026-05-13) — trade-outcome aux head forward.
+//
+// K=3 softmax aux head — mirrors `aux_next_bar_forward` (K=2) but emits a
+// 3-way distribution over trade outcomes {Profit, Stop, Timeout}. Replaces
+// the per-bar binary direction predictor with a per-trade-close 3-way
+// outcome predictor that ALIGNS with the system's actual decision horizon
+// (multi-bar trade lifecycle).
+//
+// Architecture
+// ────────────
+// Input  [B, SH2=256]   (h_s2_aux — aux trunk output post-ELU)
+//    ↓ Linear (W1 [H, SH2]) + bias (b1 [H]) + ELU
+// Hidden [B, H=128]
+//    ↓ Linear (W2 [K=3, H]) + bias (b2 [K])
+// Logits [B, 3]
+//    ↓ stable softmax (max-shift)
+// Probs  [B, 3]   →  [p_Profit, p_Stop, p_Timeout]
+//
+// Phase A3 (this commit) uses h_s2_aux only; Phase B will switch the input
+// to concat `[h_s2_aux (256) || plan_params (6)]` = 262 dim. The kernel is
+// dead code at A3 (no Rust launcher yet) — A4/A5 land loss/backward, then
+// Phase B lands the wireup.
+//
+// Saved tensors (mirror aux_next_bar_forward)
+// ────────────────────────────────────────────
+//   - hidden_out  [B, H]   ELU post-activation (consumed by backward)
+//   - logits_out  [B, K=3] pre-softmax logits (consumed by backward)
+//   - softmax_out [B, K=3] post-softmax probabilities — THREE consumers:
+//       * aux_trade_outcome_loss_reduce  (Phase A4: sparse CE numerator)
+//       * aux_trade_outcome_backward     (Phase A5: d_logits = p - one_hot)
+//       * aux_outcome_softmax_to_per_env (Phase C: state slots 121-123)
+//
+// Discipline
+// ──────────
+// - Pure GPU, no CPU branches in captured graph per
+//   pearl_no_host_branches_in_captured_graph.md.
+// - Block tree-reduce style (one block per sample, AUX_BLOCK threads). No
+//   atomicAdd per feedback_no_atomicadd.
+// - Numerically stable softmax via max-shift (mirrors aux_next_bar_forward
+//   step 3) — extreme negative logits get floored implicitly by exp.
+// - K is passed as a runtime arg (rather than #define'd) so the kernel
+//   stays K-generic; Phase D extends to K=3 outcomes × per-action W = 12
+//   weights, but this kernel always emits K=3 (one row of softmax probs).
+//
+// Block: AUX_BLOCK threads. Grid: (B, 1, 1). Shared memory: H + K floats
+// (h cache for Linear_2 + temporary K-vector for stable softmax).
+
+#include <cuda_runtime.h>
+#include <math_constants.h>
+
+/* Compile-time hidden width — must match `AUX_HIDDEN_DIM` from
+ * `aux_heads_kernel.cu` (the K=2 sibling). Both kernels share the head's
+ * hidden width via `gpu_aux_heads.rs::AUX_HIDDEN_DIM`. 2026-05-13 enlarged
+ * 32 → 128 (see aux_heads_kernel.cu header for L40S budget rationale). */
+#ifndef AUX_HIDDEN_DIM
+#define AUX_HIDDEN_DIM 128
+#endif
+
+/* Block dim convention shared with aux_heads_kernel.cu. */
+#ifndef AUX_BLOCK
+#define AUX_BLOCK 256
+#endif
+
+/* ELU forward — branchless on x sign. */
+__device__ __forceinline__ float aux_to_elu_fwd(float x) {
+    return (x > 0.0f) ? x : (expf(x) - 1.0f);
+}
+
+/* =====================================================================
+ * aux_trade_outcome_forward — Linear → ELU → Linear → softmax over K=3
+ * trade-outcome classes {Profit, Stop, Timeout}.
+ *
+ * Per sample b (one block per sample, AUX_BLOCK threads):
+ *   1. h[k] = ELU( b1[k] + sum_j w1[k, j] * h_s2_aux[b, j] )  for k ∈ [0, H)
+ *   2. logits[b, kc] = b2[kc] + sum_k w2[kc, k] * h[k]        for kc ∈ [0, K)
+ *   3. softmax[b, kc] = stable_softmax(logits[b, :])[kc]      (max-shift)
+ *
+ * `h_s2_aux` is the aux trunk output (SP14 Phase C.5b separation — aux
+ * heads read the aux trunk's `h_s2_aux`, NOT Q's `h_s2`). Backward will
+ * flow this gradient into `dh_s2_aux_accum`; the aux trunk's backward
+ * kernel terminates at the encoder boundary (no `dx_in` output param)
+ * so Q's encoder is structurally protected from aux contamination.
+ *
+ * Shared memory layout: `[H + K]` floats — `[0, H)` is the ELU post-
+ * activation cache for Linear_2's reads; `[H, H+K)` is the K-vector
+ * logits scratch for the softmax pass.
+ * ===================================================================== */
+extern "C" __global__ void aux_trade_outcome_forward(
+    const float* __restrict__ h_s2_aux,   /* [B, SH2] row-major (aux trunk output) */
+    const float* __restrict__ w1,         /* [H, SH2] row-major */
+    const float* __restrict__ b1,         /* [H] */
+    const float* __restrict__ w2,         /* [K, H] row-major (K=3) */
+    const float* __restrict__ b2,         /* [K] */
+    int B,
+    int SH2,
+    int K,
+    float* __restrict__ hidden_out,       /* [B, H] OUT (saved for backward) */
+    float* __restrict__ logits_out,       /* [B, K] OUT (saved for backward) */
+    float* __restrict__ softmax_out       /* [B, K] OUT (saved for 3 consumers) */
+) {
+    const int b   = blockIdx.x;
+    if (b >= B) return;
+    const int tid = threadIdx.x;
+    const int H   = AUX_HIDDEN_DIM;
+
+    extern __shared__ float smem[];       /* [H + K] floats */
+    float* sh_h     = smem;               /* [H] hidden cache */
+    float* sh_logit = smem + H;           /* [K] logits cache for softmax */
+
+    /* Step 1: per-hidden-unit Linear_1 + ELU. Each thread computes one or
+     * more hidden lanes via stride loop. SH2 is small enough that each
+     * thread doing a serial dot over SH2 is cheaper than block-cooperative
+     * reduce per lane (mirrors aux_next_bar_forward step 1 exactly). */
+    for (int k = tid; k < H; k += blockDim.x) {
+        float acc = b1[k];
+        const float* w_row = w1 + (size_t)k * SH2;
+        const float* h_row = h_s2_aux + (size_t)b * SH2;
+        for (int j = 0; j < SH2; ++j) {
+            acc += w_row[j] * h_row[j];
+        }
+        const float h_post = aux_to_elu_fwd(acc);
+        sh_h[k] = h_post;
+        hidden_out[(size_t)b * H + k] = h_post;
+    }
+    __syncthreads();
+
+    /* Step 2: K-way Linear_2. Thread 0 emits all K logits (K=3 — tree-
+     * reducing a K=3 vector is wasted bookkeeping). Stores each logit
+     * into the saved buffer AND a shmem scratch for Step 3's softmax. */
+    if (tid == 0) {
+        for (int kc = 0; kc < K; ++kc) {
+            float acc = b2[kc];
+            const float* w2_row = w2 + (size_t)kc * H;
+            for (int k = 0; k < H; ++k) {
+                acc += w2_row[k] * sh_h[k];
+            }
+            logits_out[(size_t)b * K + kc] = acc;
+            sh_logit[kc] = acc;
+        }
+    }
+    __syncthreads();
+
+    /* Step 3: stable softmax on the K-vector — subtract max before exp.
+     * Thread 0 again (K=3 — same per-block serial pattern as K=2 sibling). */
+    if (tid == 0) {
+        float lmax = sh_logit[0];
+        for (int kc = 1; kc < K; ++kc) {
+            if (sh_logit[kc] > lmax) lmax = sh_logit[kc];
+        }
+        float sum_e = 0.0f;
+        for (int kc = 0; kc < K; ++kc) {
+            sum_e += expf(sh_logit[kc] - lmax);
+        }
+        /* Numerical floor for sum_e — extreme negative logits across all K
+         * could produce sum_e ≈ 0 → 1/sum_e = inf. Defensive 1e-30 mirrors
+         * the loss-reduce kernel's p_clipped floor (Phase A4 will land that
+         * kernel). Under correct producer behavior the trained logits keep
+         * sum_e in a sensible range; this floor is robustness against the
+         * untrained init regime where W2 may produce extreme outputs. */
+        const float inv_sum = 1.0f / fmaxf(sum_e, 1e-30f);
+        for (int kc = 0; kc < K; ++kc) {
+            softmax_out[(size_t)b * K + kc] = expf(sh_logit[kc] - lmax) * inv_sum;
+        }
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index e6a1003af..d625ee905 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -2340,7 +2340,50 @@ extern "C" __global__ void experience_env_step(
      * `hold_reward_ema_idx`) so the kernel stays decoupled from SP slot
      * drift. The kernel reads `isv_signals_ptr[aux_align_scale_idx]`
      * NULL-safely (returns 0 if isv_signals_ptr == NULL → β no-op). */
-    int aux_align_scale_idx
+    int aux_align_scale_idx,
+    /* SP22 H6 vNext Phase A3 (2026-05-13) — save-for-backward buffers
+     * feeding the trade-outcome aux head's label producer
+     * (`trade_outcome_label_kernel.cu`) at the next training step.
+     *
+     * `pnl_vs_target_at_close_per_env` ← [N] f32: written at every
+     * `segment_complete` bar (both voluntary exits and trail-fire) with
+     * the inline-computed ratio
+     *   `pre_trade_position × (raw_close − entry_price)` /
+     *     `(ps[PS_PLAN_PROFIT_TARGET] × prev_equity)`
+     * symmetrically clamped to [-2, +2] per pearl_symmetric_clamp_audit
+     * (mirrors the state-side `plan_isv[PLAN_ISV_PNL_VS_TARGET]` formula
+     * exactly, computed inline because that local stack array lives in
+     * the sibling `experience_state_gather` kernel and isn't visible
+     * here). Values ≥ 1.0 mean the profit threshold was hit at close →
+     * label kernel classifies as Profit (0). NULL-tolerant for test
+     * scaffolds + the Phase A3 dead-code pre-wireup state (the kernel
+     * arg is plumbed through ahead of consumer landing in Phase A4/A5;
+     * passing NULL is a no-op write per the standard `if (ptr != NULL)`
+     * gate).
+     *
+     * `pnl_vs_stop_at_close_per_env` ← [N] f32: written at the same site
+     * with `-unrealized_close / (plan_stop × prev_equity)` clamped
+     * [-2, +2]. Values ≥ 1.0 mean the stop_loss threshold was hit at
+     * close → label kernel classifies as Stop (1). When neither
+     * threshold is ≥ 1.0 the label kernel emits Timeout (2).
+     *
+     * Race-free: each thread writes only its own `[i]` slot. Slot
+     * persists across rollout steps until the next trade-close
+     * overwrites it; the label kernel reads only on `trade_close_per_
+     * sample[i*L + t] != 0` so stale values don't leak into masked
+     * labels. FoldReset sentinel 0.0 across all `alloc_episodes` slots —
+     * matches the `alpha_per_env` reset pattern; the label kernel's
+     * `≥ 1.0` predicate reads sentinel 0 as "no threshold hit" which
+     * collapses to Timeout (2), but since the masked-bar branch never
+     * fires (no trade_close event at fold-0-step-0), the sentinel never
+     * reaches the loss.
+     *
+     * Phase A3 (this commit): buffers + writes wired but no consumer
+     * (label kernel is dead code at A2/A3, launches in A4/A5). Phase B
+     * lands the input concat (h_s2_aux || plan_params). Phase D lands
+     * the 12-weight W atom-shift. Phase F validates. */
+    float* __restrict__ pnl_vs_target_at_close_per_env,
+    float* __restrict__ pnl_vs_stop_at_close_per_env
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -3514,6 +3557,69 @@ extern "C" __global__ void experience_env_step(
             is_win_per_env[i] = (segment_return > 0.0f) ? 1 : 0;
         }
 
+        /* SP22 H6 vNext Phase A3 (2026-05-13) — save-for-backward for
+         * the trade-outcome aux head's label producer
+         * (`trade_outcome_label_kernel.cu`). Capture the trade's
+         * realized P&L ratios vs profit_target / stop_loss at the
+         * exact moment of close.
+         *
+         * The state-side `plan_isv[PLAN_ISV_PNL_VS_TARGET / _VS_STOP]`
+         * slots (computed in `experience_state_gather` at line ~874-879)
+         * are NOT in scope here — that's a sibling kernel reading a
+         * separate stack array. Compute the same ratio inline using
+         * env_step locals already populated upstream:
+         *   - `pre_trade_position` (saved at line ~2821, pre-mutation)
+         *     × `(raw_close - entry_price)` = the trade's MTM P&L just
+         *     before the close action takes effect (becomes realized
+         *     P&L at close).
+         *   - `prev_equity` (saved at line ~2664 from ps[PS_PREV_EQUITY])
+         *     is the equity at the start of this bar; used as the
+         *     normalisation denominator to mirror the state-side ratio's
+         *     `f_unrealized_pnl / f_equity` denominator semantic
+         *     (state-side reads current `f_portfolio_value`; both are
+         *     "this-bar equity" — symmetric clamp [-2, +2] absorbs the
+         *     small difference for trades where equity moved
+         *     materially across the close action).
+         *
+         * Both pvt and pvs are symmetrically clamped to [-2, +2] per
+         * `pearl_symmetric_clamp_audit` — matches the state-side slot
+         * range exactly so the label kernel's `≥ 1.0` predicate works
+         * identically whether reading state-side or save-for-backward
+         * values.
+         *
+         * Race-free per-env write. NULL-tolerant for the Phase A3
+         * pre-wireup state (the kernel arg is plumbed ahead of the
+         * label producer's first launch in Phase A4/A5). */
+        if (pnl_vs_target_at_close_per_env != NULL ||
+            pnl_vs_stop_at_close_per_env   != NULL) {
+            const float plan_profit_close = ps[PS_PLAN_PROFIT_TARGET];
+            const float plan_stop_close   = ps[PS_PLAN_STOP_LOSS];
+            const float unrealized_close  =
+                (entry_price > 0.0f && pre_trade_position != 0.0f)
+                    ? pre_trade_position * (raw_close - entry_price)
+                    : 0.0f;
+            const float equity_clamp = fmaxf(prev_equity, 1.0f);
+
+            if (pnl_vs_target_at_close_per_env != NULL) {
+                const float pvt = (plan_profit_close > 1e-6f)
+                    ? fmaxf(-2.0f, fminf(
+                          unrealized_close /
+                          (plan_profit_close * equity_clamp + 1e-6f),
+                          2.0f))
+                    : 0.0f;
+                pnl_vs_target_at_close_per_env[i] = pvt;
+            }
+            if (pnl_vs_stop_at_close_per_env != NULL) {
+                const float pvs = (plan_stop_close > 1e-6f)
+                    ? fmaxf(-2.0f, fminf(
+                          -unrealized_close /
+                          (plan_stop_close * equity_clamp + 1e-6f),
+                          2.0f))
+                    : 0.0f;
+                pnl_vs_stop_at_close_per_env[i] = pvs;
+            }
+        }
+
         /* SP11 component attribution split (mutually exclusive — preserved
          * from pre-SP20 semantic). r_used is identical in both branches;
          * only the rc[*] slot the SP11 controller weighs differs. */
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 45040f4bc..3b3f89bc2 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -743,6 +743,32 @@ pub struct GpuExperienceCollector {
     /// `step_ret > 0` predicate when NULL for test scaffolds).
     pub(crate) is_win_per_env: CudaSlice<i32>, // [alloc_episodes]
 
+    /// SP22 H6 vNext Phase A3 (2026-05-13): per-env trade-close
+    /// `plan_isv[PLAN_ISV_PNL_VS_TARGET]` save-for-backward scratch for
+    /// the trade-outcome aux head's label producer. Written at every
+    /// `segment_complete` bar by `experience_env_step` (alongside
+    /// `alpha_per_env[i]`). Value semantics: `unrealized_pnl / (plan_profit
+    /// × equity)` clamped symmetrically to [-2, +2]. Consumed by the
+    /// `trade_outcome_label_kernel` at the next training step's label
+    /// production — `>= 1.0` predicate classifies as Profit. Sized
+    /// `[alloc_episodes]`. FoldReset sentinel 0.0 across all slots —
+    /// registered in `state_reset_registry.rs` as
+    /// `pnl_vs_target_at_close_per_env` (mirror of alpha_per_env reset
+    /// pattern). Phase A3 (this commit): buffer + producer write wired;
+    /// consumer (label kernel launcher) lands in Phase A4/A5.
+    pub(crate) pnl_vs_target_at_close_per_env: CudaSlice<f32>, // [alloc_episodes]
+
+    /// SP22 H6 vNext Phase A3 (2026-05-13): per-env trade-close
+    /// `plan_isv[PLAN_ISV_PNL_VS_STOP]` save-for-backward scratch.
+    /// Companion to `pnl_vs_target_at_close_per_env` — same producer site
+    /// (segment_complete branch), same fold-reset semantics, same
+    /// NULL-tolerant kernel arg. Value semantics: `-unrealized_pnl /
+    /// (plan_stop × equity)` clamped [-2, +2]. Consumed by the
+    /// `trade_outcome_label_kernel` `>= 1.0` predicate as the Stop
+    /// classification oracle. Sized `[alloc_episodes]`. FoldReset
+    /// sentinel 0.0.
+    pub(crate) pnl_vs_stop_at_close_per_env: CudaSlice<f32>, // [alloc_episodes]
+
     /// SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost
     /// scratch. Written every step by `experience_env_step` as
     /// `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`.
@@ -1937,6 +1963,27 @@ impl GpuExperienceCollector {
                 "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}"
             )))?;
 
+        // SP22 H6 vNext Phase A3 (2026-05-13): save-for-backward buffers
+        // for the trade-outcome aux head's label producer. Both written
+        // at every segment_complete bar by experience_env_step alongside
+        // alpha_per_env. Consumed by the trade_outcome_label_kernel
+        // (Phase A4/A5 wireup) to classify each trade-close into one of
+        // {Profit=0, Stop=1, Timeout=2}. Sentinel 0.0; sized
+        // [alloc_episodes]. FoldReset registered in state_reset_registry
+        // mirror of alpha_per_env reset pattern.
+        let pnl_vs_target_at_close_per_env = stream
+            .alloc_zeros::<f32>(alloc_episodes)
+            .map_err(|e| MLError::ModelError(format!(
+                "sp22-vnext: alloc pnl_vs_target_at_close_per_env \
+                 ({alloc_episodes} f32): {e}"
+            )))?;
+        let pnl_vs_stop_at_close_per_env = stream
+            .alloc_zeros::<f32>(alloc_episodes)
+            .map_err(|e| MLError::ModelError(format!(
+                "sp22-vnext: alloc pnl_vs_stop_at_close_per_env \
+                 ({alloc_episodes} f32): {e}"
+            )))?;
+
         // SP21 T3.3 (2026-05-10): per-env per-bar Hold opportunity-cost
         // scratch. Sentinel 0.0; written by experience_env_step every
         // step as `per_bar_opp_cost = -aux_conf × cost_scale ∈ [-0.25, 0]`.
@@ -3080,6 +3127,8 @@ impl GpuExperienceCollector {
             label_at_open_per_env,
             alpha_per_env,
             is_win_per_env,
+            pnl_vs_target_at_close_per_env,
+            pnl_vs_stop_at_close_per_env,
             per_bar_opp_cost_per_env,
             hold_baseline_buffer,
             states_out,
@@ -6679,6 +6728,18 @@ impl GpuExperienceCollector {
                         use crate::cuda_pipeline::sp22_isv_slots::SP22_AUX_ALIGN_SCALE_INDEX;
                         SP22_AUX_ALIGN_SCALE_INDEX as i32
                     })
+                    // SP22 H6 vNext Phase A3 (2026-05-13): save-for-backward
+                    // buffers feeding the trade-outcome aux head's label
+                    // producer. Written at every segment_complete bar
+                    // alongside alpha_per_env with
+                    // `plan_isv[PLAN_ISV_PNL_VS_TARGET]` and
+                    // `plan_isv[PLAN_ISV_PNL_VS_STOP]` — the trade's
+                    // realized P&L ratios vs profit_target / stop_loss
+                    // at the exact close moment. Consumed by the label
+                    // kernel (Phase A4/A5 wireup) to classify each
+                    // trade-close as {Profit, Stop, Timeout}.
+                    .arg(&mut self.pnl_vs_target_at_close_per_env)
+                    .arg(&mut self.pnl_vs_stop_at_close_per_env)
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_env_step t={t}: {e}"
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index f65bfccd7..3a806e271 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -2166,6 +2166,21 @@ impl StateResetRegistry {
                 category: ResetCategory::FoldReset,
                 description: "ISV[SP22_AUX_ALIGN_SCALE_INDEX=537] — SP22 H6 Phase 3 (2026-05-13) SP11-controller-emitted adaptive scale_β for the β producer at trade-close. Producer (eventual): extended `reward_subsystem_controller_kernel.cu` (7-weight output, anchor on `REWARD_AUX_ALIGN_EMA_INDEX=536`, target ~10% of total reward magnitude, bounds [0.05, 2.0]) — Phase B6, NOT yet landed. Consumers: `experience_env_step::segment_complete` (training-side β: `r_aux_align = scale_β × alignment × profit_pos`) and `backtest_env_kernel::segment_complete` (eval-side β with the same formula); HEALTH_DIAG `sp11_reward` printer (the `w_aux=…` column). FoldReset sentinel 0.5 — structural prior matching W's [-0.5, 0, +0.5, 0] init approach; activates β with fixed magnitude from epoch 0. The 0.5 prior was chosen over 0 (the original spec) because slot 537 has no producer yet (Phase B6 deferred); 0 would make β a runtime no-op. Once B6 lands and the controller emits adaptively per the bounds [0.05, 2.0], the prior is overwritten on first emit per `pearl_first_observation_bootstrap`; subsequent EMA-driven adaptation tracks the anchor signal.",
             },
+            // SP22 H6 vNext Phase A3 (2026-05-13) — save-for-backward
+            // buffers for the trade-outcome aux head's label producer.
+            // Both written at every segment_complete bar alongside
+            // alpha_per_env / is_win_per_env; consumed by the
+            // trade_outcome_label_kernel at the next training step.
+            RegistryEntry {
+                name: "pnl_vs_target_at_close_per_env",
+                category: ResetCategory::FoldReset,
+                description: "GpuExperienceCollector.pnl_vs_target_at_close_per_env [alloc_episodes] f32 device-resident scratch — SP22 H6 vNext Phase A3 (2026-05-13) per-env trade-close P&L ratio vs `plan_profit × equity` save-for-backward buffer. Producer: `experience_env_step` `is_close && segment_complete` branch — writes `plan_isv[PLAN_ISV_PNL_VS_TARGET]` (the local plan_isv value already computed at line ~874 from `f_unrealized_pnl / (plan_profit × f_equity + 1e-6f)` symmetrically clamped to [-2, +2] per `pearl_symmetric_clamp_audit`). Consumer (eventual, Phase A4/A5): `trade_outcome_label_kernel.cu`'s `pnl_vs_target_at_close` arg — `≥ 1.0` predicate classifies the closed trade as `Profit` (label 0) per the K=3 outcome taxonomy. FoldReset sentinel 0.0 across all `alloc_episodes` slots — leftover ratios from the previous fold's last trade-close would corrupt the new fold's first label production (0.0 sentinel reads as 'no threshold hit' → classifies as Timeout if a stale read leaked, but the label kernel reads only when `trade_close_per_sample != 0` so the sentinel never reaches the loss). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice<f32>` (device-resident, no host buffer to fill). Phase A3 (this commit) lands the buffer + producer write + FoldReset wiring atomically per `feedback_no_partial_refactor`; the consumer (label kernel launcher) lands in Phase A4/A5.",
+            },
+            RegistryEntry {
+                name: "pnl_vs_stop_at_close_per_env",
+                category: ResetCategory::FoldReset,
+                description: "GpuExperienceCollector.pnl_vs_stop_at_close_per_env [alloc_episodes] f32 device-resident scratch — SP22 H6 vNext Phase A3 (2026-05-13) per-env trade-close P&L ratio vs `plan_stop × equity`. Companion to `pnl_vs_target_at_close_per_env` — same segment_complete producer site, same FoldReset semantics, same NULL-tolerant kernel arg pattern. Value semantics: `plan_isv[PLAN_ISV_PNL_VS_STOP]` (the local plan_isv value already computed at line ~877 from `-f_unrealized_pnl / (plan_stop × f_equity + 1e-6f)` clamped [-2, +2]). Consumer (eventual): `trade_outcome_label_kernel.cu`'s `pnl_vs_stop_at_close` arg — `≥ 1.0` predicate classifies as `Stop` (label 1). When neither pvt nor pvs ≥ 1.0 the kernel emits `Timeout` (label 2). FoldReset sentinel 0.0; reset path: `cudarc::driver::CudaStream::memset_zeros`.",
+            },
         ];
         Self { entries }
     }
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index d6af79454..44a9212ee 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -10060,6 +10060,29 @@ impl DQNTrainer {
                         )))?;
                 }
             }
+            // SP22 H6 vNext Phase A3 (2026-05-13): save-for-backward
+            // buffers feeding the trade-outcome aux head's label
+            // producer. Both written at every segment_complete bar
+            // alongside alpha_per_env / is_win_per_env. Same
+            // device-resident `CudaSlice<f32>` reset pattern.
+            "pnl_vs_target_at_close_per_env" => {
+                if let Some(ref mut collector) = self.gpu_experience_collector {
+                    let stream = collector.stream().clone();
+                    stream.memset_zeros(&mut collector.pnl_vs_target_at_close_per_env)
+                        .map_err(|e| crate::MLError::ModelError(format!(
+                            "pnl_vs_target_at_close_per_env memset_zeros: {e}"
+                        )))?;
+                }
+            }
+            "pnl_vs_stop_at_close_per_env" => {
+                if let Some(ref mut collector) = self.gpu_experience_collector {
+                    let stream = collector.stream().clone();
+                    stream.memset_zeros(&mut collector.pnl_vs_stop_at_close_per_env)
+                        .map_err(|e| crate::MLError::ModelError(format!(
+                            "pnl_vs_stop_at_close_per_env memset_zeros: {e}"
+                        )))?;
+                }
+            }
             // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold
             // opportunity-cost circular buffer (Component 2). Same
             // device-resident `CudaSlice<f32>` reset pattern as
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index d5da0231a..92e9364fc 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -17659,3 +17659,37 @@ First foundation kernel for the SP22 H6 vNext trade-outcome aux head proposal.
 Per the vNext spec, total remaining effort ~2-3 days. This commit lands the cornerstone label producer cleanly so subsequent phases have a stable foundation.
 
 **Save-for-backward** for the label producer's inputs (`pnl_vs_target_at_close`, `pnl_vs_stop_at_close`): need to add 2 new mapped-pinned buffers populated by `experience_env_step`'s segment_complete branch. Deferred to Phase A3 wireup commit.
+
+#### Phase A3 — aux_trade_outcome forward kernel + save-for-backward wireup (2026-05-13)
+
+Second foundation commit for the SP22 H6 vNext trade-outcome aux head. Lands two atomic pieces:
+
+**1. NEW kernel `aux_trade_outcome_forward_kernel.cu`** — K=3 softmax aux head forward.
+- Architecture: `Input [B, SH2=256]` → Linear (W1 [128, SH2]) + ELU → Hidden [B, 128] → Linear (W2 [3, 128]) → Logits [B, 3] → stable softmax → Probs [B, 3] over `{Profit, Stop, Timeout}`.
+- Mirrors `aux_next_bar_forward` (K=2) — same block-per-sample layout, AUX_BLOCK=256 threads, shared memory `[H + K]` floats, max-shift softmax with sum_e floor at 1e-30 for untrained-init robustness.
+- Saved tensors `hidden_out` + `logits_out` + `softmax_out` ready for Phase A5 backward and Phase C 3-slot state assembly consumer.
+- Phase A3 (this commit) **dead code** — no Rust launcher yet. A4 (loss reduce) and A5 (backward) wire it in subsequent commits.
+- Registered in `build.rs::kernels_with_common`; cubin verified `aux_trade_outcome_forward_kernel.cubin` (18 KB).
+
+**2. Save-for-backward buffers** `pnl_vs_target_at_close_per_env` + `pnl_vs_stop_at_close_per_env` ([alloc_episodes] f32 device-resident).
+- Producer: `experience_env_step::segment_complete` branch (alongside the existing `alpha_per_env` / `is_win_per_env` writes). Inline-computed because `plan_isv[…]` lives in sibling kernel `experience_state_gather`, not this scope; formula mirrors the sibling exactly: `pre_trade_position × (raw_close − entry_price) / (ps[PS_PLAN_PROFIT_TARGET / _STOP_LOSS] × prev_equity)` symmetric-clamped to `[-2, +2]` per `pearl_symmetric_clamp_audit`.
+- Consumer (Phase A4/A5 wireup): `trade_outcome_label_kernel`'s `pnl_vs_target_at_close` / `pnl_vs_stop_at_close` args — the `≥ 1.0` predicate classifies as Profit/Stop, otherwise Timeout.
+- Rust wireup: 2 new `pub(crate) CudaSlice<f32>` fields on `GpuExperienceCollector`, `stream.alloc_zeros` at construct site, kernel-launch `.arg()` threading at the `experience_env_step` call site, `StateResetRegistry` entries (FoldReset sentinel 0.0), `training_loop::reset_named_state` dispatch arms.
+- Pin tests verified: all 10 `state_reset_registry` tests pass, including `every_fold_and_soft_reset_entry_has_dispatch_arm` (the guard against the registry/dispatch divergence bug from `feedback_registry_entries_need_dispatch_arms`).
+
+**Discipline**:
+- `feedback_no_partial_refactor` — buffer + producer write + Rust struct field + alloc + launch arg + registry + dispatch all in same commit.
+- `feedback_registry_entries_need_dispatch_arms` — pin test gates the wireup.
+- `pearl_symmetric_clamp_audit` — bilateral fmaxf/fminf clamp matches state-side sibling.
+- `feedback_no_cpu_compute_strict` — all writes are pure GPU per-env (race-free per-thread).
+
+**Remaining vNext work**:
+- Phase A4: `aux_trade_outcome_loss_reduce_kernel.cu` (sparse CE over K=3, mask `-1`)
+- Phase A5: `aux_trade_outcome_backward_kernel.cu`
+- Phase B: 262-dim input concat (`h_s2_aux` || `plan_params`)
+- Phase C: 3-slot state assembly (state[121..124] = p_Profit, p_Stop, p_Timeout)
+- Phase D: 12-weight W atom-shift (4 actions × 3 outcomes)
+- Phase E: dW + Adam for W[4, 3]
+- Phase F: validation smoke
+
+Cargo check clean. Registry tests clean. Ready for Phase A4.