diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 97b1504f4..a8064fafa 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -18970,6 +18970,7 @@ impl GpuDqnTrainer {
                 "launch_health_diag_isv_mirror: health_diag orchestrator not constructed".into()
             )),
         };
+        use crate::cuda_pipeline::sp22_isv_slots::AUX_TRADE_OUTCOME_CE_EMA_INDEX;
         hd.launch_isv_mirror(
             &self.stream,
             self.isv_signals_dev_ptr,
@@ -18988,6 +18989,10 @@ impl GpuDqnTrainer {
             MOE_EXPERT_UTIL_EMA_BASE       as i32,
             MOE_GATE_ENTROPY_EMA_INDEX     as i32,
             MOE_LAMBDA_EFF_INDEX           as i32,
+            // SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome
+            // head CE EMA slot. Kernel reads ISV[538] and mirrors into
+            // snap.aux_trade_outcome_ce (appended field at word 149).
+            AUX_TRADE_OUTCOME_CE_EMA_INDEX as i32,
         )
     }
 
diff --git a/crates/ml/src/cuda_pipeline/gpu_health_diag.rs b/crates/ml/src/cuda_pipeline/gpu_health_diag.rs
index decc7f320..8bb6c6f38 100644
--- a/crates/ml/src/cuda_pipeline/gpu_health_diag.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_health_diag.rs
@@ -157,6 +157,9 @@ impl GpuHealthDiag {
         moe_expert_util_base_idx: i32,
         moe_gate_entropy_idx: i32,
         moe_lambda_eff_idx: i32,
+        // SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head
+        // CE EMA ISV slot index (= AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538).
+        aux_trade_outcome_ce_idx: i32,
     ) -> Result<(), MLError> {
         let snap_dev_ptr: u64 = self.snapshot.dev_ptr();
         unsafe {
@@ -181,6 +184,9 @@ impl GpuHealthDiag {
                 .arg(&moe_expert_util_base_idx)
                 .arg(&moe_gate_entropy_idx)
                 .arg(&moe_lambda_eff_idx)
+                // SP22 H6 vNext Phase F-3c (2026-05-14): K=3 head's CE
+                // EMA ISV slot index (= 538).
+                .arg(&aux_trade_outcome_ce_idx)
                 .launch(LaunchConfig {
                     grid_dim: (1, 1, 1),
                     block_dim: (1, 1, 1),
diff --git a/crates/ml/src/cuda_pipeline/health_diag.rs b/crates/ml/src/cuda_pipeline/health_diag.rs
index 11692978e..b1ef447c0 100644
--- a/crates/ml/src/cuda_pipeline/health_diag.rs
+++ b/crates/ml/src/cuda_pipeline/health_diag.rs
@@ -269,6 +269,20 @@ pub struct HealthDiagSnapshot {
     /// aggregator scripts) read the same 12-bin layout that already lives
     /// in `MonitoringSummary`.
     pub action_counts: [u32; HEALTH_DIAG_ACTION_BINS],
+
+    // ── SP22 H6 vNext Phase F-3c (2026-05-14): trade-outcome head CE EMA ────
+    /// K=3 trade-outcome aux head's batch-mean sparse cross-entropy EMA,
+    /// mirrored from `ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]` by an
+    /// extended `health_diag_isv_mirror` kernel write. Cold-start sentinel
+    /// 0.0; Pearl A bootstrap replaces with the first non-zero loss
+    /// observation. Healthy learning: `~1.098 = ln(3)` early →
+    /// monotonic-ish decrease toward 0.5-0.7 over epochs. Pinned at
+    /// `~1.098` indicates the head can't learn (label noise / under-
+    /// capacity / outcomes genuinely unconditional). **Appended at end
+    /// of struct** per "Field order is stable" doc rule — existing
+    /// fields' offsets preserved so the kernel's pre-vNext word writes
+    /// don't need re-validation.
+    pub aux_trade_outcome_ce: f32,
 }
 
 impl Default for HealthDiagSnapshot {
@@ -400,11 +414,14 @@ mod tests {
         // = 8+8+7+2+1+10+4+5+9+2+6+6+6+3+3+4+7+3+16+4+4+6+3+10+12 = 149 fields
         // SP13 B1.0 (2026-05-05): aux block dropped from 4 → 3 floats
         // (retired `aux_label_scale` together with ISV[117]).
-        let expected_bytes = 149 * 4;
+        // SP22 H6 vNext Phase F-3c (2026-05-14): +1 field at end —
+        // `aux_trade_outcome_ce: f32` (K=3 head CE EMA mirror). Total
+        // now 150 × 4 = 600 bytes. Matches kernel WORD_TOTAL = 150.
+        let expected_bytes = 150 * 4;
         assert_eq!(
             std::mem::size_of::<HealthDiagSnapshot>(),
             expected_bytes,
-            "HealthDiagSnapshot must be {} bytes (149 × 4); update test if struct changes",
+            "HealthDiagSnapshot must be {} bytes (150 × 4); update test if struct changes",
             expected_bytes,
         );
     }
diff --git a/crates/ml/src/cuda_pipeline/health_diag_kernel.cu b/crates/ml/src/cuda_pipeline/health_diag_kernel.cu
index 6441b29bc..405a9845f 100644
--- a/crates/ml/src/cuda_pipeline/health_diag_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/health_diag_kernel.cu
@@ -197,10 +197,15 @@
 /* Action-count histogram (12 words, [137..149), all u32). */
 #define WORD_ACTION_COUNTS_BASE    137
 
-#define WORD_TOTAL                 149
+/* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head CE EMA,
+ * appended at end of struct per "Field order is stable" doc rule. Word
+ * 149 (= old WORD_TOTAL). Snap layout total bumps to 150. */
+#define WORD_AUX_TRADE_OUTCOME_CE  149
+
+#define WORD_TOTAL                 150
 
 /* Pin the layout: matches `snapshot_size_is_stable` on the Rust side. */
-static_assert(WORD_TOTAL == 149,
+static_assert(WORD_TOTAL == 150,
               "HealthDiagSnapshot layout drift — update kernel WORD_* and the "
               "Rust struct in lockstep, then bump snapshot_size_is_stable.");
 
@@ -264,7 +269,12 @@ extern "C" __global__ void health_diag_isv_mirror(
      * argument dropped from kernel signature. */
     int moe_expert_util_base_idx,          /* ISV[118], 8 contiguous slots */
     int moe_gate_entropy_idx,              /* ISV[126] */
-    int moe_lambda_eff_idx                 /* ISV[128] */
+    int moe_lambda_eff_idx,                /* ISV[128] */
+    /* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head's
+     * sparse-CE EMA. Source: ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]
+     * (producer: aux_outcome_ce_ema_update kernel, single-thread direct-
+     * to-ISV EMA with Pearl A bootstrap, α=0.05). */
+    int aux_trade_outcome_ce_idx           /* ISV[538] */
 ) {
     /* Single-block, single-thread contract. Mirrors aux_heads_loss_ema's
      * grid guard: the work is ~22 scalar EMA reads + writes. */
@@ -309,6 +319,13 @@ extern "C" __global__ void health_diag_isv_mirror(
     snap[WORD_AUX_MOE_ENT]            = isv[moe_gate_entropy_idx];
     snap[WORD_AUX_MOE_LAMBDA_EFF]     = isv[moe_lambda_eff_idx];
 
+    /* SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome head's
+     * sparse-CE EMA mirrored from ISV[538] into the appended struct
+     * field at word 149. The producer (`aux_outcome_ce_ema_update`)
+     * fires every step before this mirror runs (stream-implicit
+     * ordering), so the read picks up the just-updated EMA. */
+    snap[WORD_AUX_TRADE_OUTCOME_CE]   = isv[aux_trade_outcome_ce_idx];
+
     /* No __threadfence_system() here — the finalise kernel (Phase 2E)
      * issues a single fence covering all writes from the family. */
 }
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index e80e3e93c..cc9c2e57e 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -10766,6 +10766,17 @@ impl DQNTrainer {
                 "aux_regime_ce_ema".to_owned(),
                 trainer.read_isv_signal_at(AUX_REGIME_CE_EMA_INDEX) as f64,
             ));
+            // SP22 H6 vNext Phase F-3c (2026-05-14): K=3 trade-outcome
+            // head's batch-mean sparse-CE EMA. Healthy learning: ~ln(3)
+            // ≈ 1.098 early → decrease toward 0.5-0.7. Pinned at ln(3)
+            // for many epochs indicates the head can't learn.
+            {
+                use crate::cuda_pipeline::sp22_isv_slots::AUX_TRADE_OUTCOME_CE_EMA_INDEX;
+                out.push((
+                    "aux_trade_outcome_ce_ema".to_owned(),
+                    trainer.read_isv_signal_at(AUX_TRADE_OUTCOME_CE_EMA_INDEX) as f64,
+                ));
+            }
             out.push((
                 "isv_reward_popart_ema".to_owned(),
                 trainer.read_isv_signal_at(REWARD_POPART_EMA_INDEX) as f64,
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 80f2651eb..4e2b15427 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -18294,3 +18294,54 @@ Completes the Phase F-3 observability chain. The kernel + ISV slot landed in F-3
 - Smoke observability: F-3 + F-3b ✓
 
 **Phase F deployment** (next): push branch + `argo-train.sh` smoke run. Decisive test of the spec's hypothesis. Spec falsification criterion (per the original spec doc): if smoke produces WR ≈ 0.4346 ± 0.005, the trade-outcome hypothesis is falsified; the conclusion would be "WR is data-determined at this bar resolution + horizon, not a model architecture issue", and future work explores different data sampling / different prediction target.
+
+#### Phase F-3c — HEALTH_DIAG snap + console line for K=3 CE EMA (2026-05-14)
+
+Completes the observability story. Smoke runs now print `aux_trade_outcome_ce_ema` every epoch in the standard HEALTH_DIAG output — no ISV inspector required.
+
+**Changes**:
+
+`health_diag.rs` (struct):
+- New field `aux_trade_outcome_ce: f32` appended at END of `HealthDiagSnapshot` per the documented "Field order is stable; adding fields appends to the end" rule. Existing fields' byte offsets preserved → no kernel-side word offset re-validation needed.
+- `snapshot_size_is_stable` test pin bumped 149 × 4 → 150 × 4 = 600 bytes.
+
+`health_diag_kernel.cu`:
+- New `WORD_AUX_TRADE_OUTCOME_CE = 149` (appended at end).
+- `WORD_TOTAL = 150` (was 149); static_assert updated.
+- New kernel arg `aux_trade_outcome_ce_idx` appended after `moe_lambda_eff_idx`.
+- New mirror write `snap[WORD_AUX_TRADE_OUTCOME_CE] = isv[aux_trade_outcome_ce_idx]` after the existing MoE mirrors. Stream-implicit ordering: the `aux_outcome_ce_ema_update` producer (Phase F-3b) fires before `health_diag_isv_mirror` in the per-step training graph, so the read picks up the just-updated EMA.
+
+`gpu_health_diag.rs`:
+- `launch_isv_mirror` Rust launcher gets new `aux_trade_outcome_ce_idx: i32` arg threaded through.
+
+`gpu_dqn_trainer.rs`:
+- `launch_health_diag_isv_mirror` passes `AUX_TRADE_OUTCOME_CE_EMA_INDEX as i32` as the new arg.
+
+`training_loop.rs`:
+- Per-epoch metrics push appends `("aux_trade_outcome_ce_ema", ISV[538])` to the standard `out` vec. The downstream metrics output (CSV, console, etc.) automatically includes the new column.
+
+**End-to-end Phase F-3 chain now closed**:
+```
+K=3 forward (D + B5b) → aux_to_loss_scalar_buf[0]
+  → aux_outcome_ce_ema_update (F-3a kernel) → ISV[538]
+  → health_diag_isv_mirror (F-3c kernel ext) → snap.aux_trade_outcome_ce
+  → training_loop metrics push (F-3c Rust ext) → console output
+```
+
+Operator now sees CE every epoch:
+- Cold-start: `aux_trade_outcome_ce_ema=0.000`
+- After bootstrap: `aux_trade_outcome_ce_ema≈1.098` (= ln(3))
+- After training: ideally `aux_trade_outcome_ce_ema≈0.5-0.7`
+
+**Verification**:
+- `cargo check -p ml` clean.
+- `cargo test -p ml --lib` → **1016/0 green** (including the bumped `snapshot_size_is_stable` byte-size pin test).
+
+**Phase F end-to-end ready**. The trade-outcome vNext stack now has:
+- Full GPU kernel chain (A2-A5 + D + plan-conditioning)
+- Full Rust wireup (B0-B4, B4b-1/2, C-1/C-2, B5b)
+- Real labels reaching the trainer (B4b-2 direct-path gather)
+- K=3 → policy via TWO paths (state slots + atom-shift)
+- End-to-end CE observability (F-3 + F-3b + F-3c)
+
+Spec-falsification criterion (per docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md): if the smoke produces WR ≈ 0.4346 ± 0.005 AND CE drops meaningfully (showing the head IS learning the labels), the spec is FALSIFIED — trade-outcome prediction is learnable but doesn't influence policy effectively. If CE pinned AND WR pinned, the spec is INCONCLUSIVE — label noise or under-capacity, head couldn't learn the signal so we don't know if the signal would have helped.