diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index baecfb8b3..6ea3dc74b 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -849,6 +849,14 @@ fn main() {
         // dh_s2_aux_accum [B, 256] only consumes first 256 cols (plan_
         // params gradient is stop-grad).
         "strided_row_saxpy_kernel.cu",
+        // SP22 H6 vNext Phase F-3 (2026-05-14): K=3 trade-outcome head's
+        // sparse-CE EMA producer for observability. Single-thread
+        // single-block direct-to-ISV writer with Pearl A bootstrap.
+        // Writes ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538] every training
+        // step alongside the existing K=2/K=5 aux EMA producers.
+        // Dead code until launcher wired (next commit step in this
+        // phase).
+        "aux_outcome_ce_ema_kernel.cu",
         // SP22 H6 vNext Phase C (2026-05-14): per-env 3-slot softmax
         // extractor for the K=3 trade-outcome aux head. Copies
         // aux_outcome_softmax [n_envs, K=3] → prev_aux_outcome_probs
diff --git a/crates/ml/src/cuda_pipeline/aux_outcome_ce_ema_kernel.cu b/crates/ml/src/cuda_pipeline/aux_outcome_ce_ema_kernel.cu
new file mode 100644
index 000000000..6fc3a51bb
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/aux_outcome_ce_ema_kernel.cu
@@ -0,0 +1,91 @@
+// crates/ml/src/cuda_pipeline/aux_outcome_ce_ema_kernel.cu
+//
+// SP22 H6 vNext Phase F-3 (2026-05-14) — K=3 trade-outcome aux head's
+// sparse-CE EMA producer for observability during smoke runs.
+//
+// Reads `aux_to_loss_scalar_buf[0]` (the K=3 head's per-batch mean CE
+// scalar, written by `aux_trade_outcome_loss_reduce` at the end of
+// `aux_heads_forward`) and EMA-blends into
+// `ISV[AUX_TRADE_OUTCOME_CE_EMA_INDEX=538]` with α=0.05. Pearl A
+// first-observation bootstrap: when the ISV slot is at sentinel 0.0,
+// the first non-zero loss observation replaces directly instead of
+// EMA-blending (matches every other ISV EMA producer's cold-start
+// contract per `pearl_first_observation_bootstrap`).
+//
+// Why a dedicated kernel (not extension of `aux_heads_loss_ema_update`)
+// ─────────────────────────────────────────────────────────────────────
+// The existing K=2 / K=5 aux-loss EMA writes through the Pearls A+D
+// 2-stage producer scratch buffer. Extending it would require
+// allocating a new scratch slot, threading the Pearls A+D mapping, and
+// touching `apply_pearls_ad_kernel` (multi-slot consumer). The K=3 head's
+// CE is for OBSERVABILITY ONLY in this commit (Phase F-3 enables Phase
+// F smoke run visibility); no Pearls A+D adaptive α is required. A
+// dedicated single-thread direct-to-ISV kernel is the minimum-scope
+// path that still gives the observability we need.
+//
+// If the K=3 CE EMA later becomes a controller anchor (e.g., for a
+// future SP11-style adaptive weight on aux-aligned reward), it can be
+// re-routed through Pearls A+D at that point. Spec dictates this is
+// observability-only for now.
+//
+// Discipline
+// ──────────
+// - Single block, single thread (kernel is O(1) — one load, one store).
+// - No HtoD per `feedback_no_htod_htoh_only_mapped_pinned.md`.
+// - Pearl A bootstrap per `pearl_first_observation_bootstrap.md`.
+// - No atomicAdd per `feedback_no_atomicadd.md` (single thread → no
+//   contention).
+//
+// Launch: grid=(1,1,1), block=(1,1,1). One thread per training step
+// (mirrors `aux_heads_loss_ema_update`'s cadence).
+//
+// Cost: ~0.5 µs (single load + EMA + store). Negligible.
+
+#include <cuda_runtime.h>
+
+extern "C" __global__ void aux_outcome_ce_ema_update(
+    /* [1] scalar mean CE from `aux_trade_outcome_loss_reduce`. */
+    const float* __restrict__ aux_to_loss_scalar_buf,
+    /* ISV signals bus pointer. NULL-tolerant (test scaffolds without
+     * ISV wiring) — kernel returns without writing. */
+    float* __restrict__ isv_signals,
+    /* AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538. Passed runtime so the
+     * kernel stays decoupled from SP slot drift (mirrors the existing
+     * aux EMA kernel's idx-arg pattern). */
+    int          ce_ema_idx,
+    /* EMA blend factor. Fixed at 0.05 per `pearl_adaptive_moe_lambda`
+     * for a slow-moving observability signal (~20-step time constant
+     * matches one epoch's worth of batches at typical sizes). Passed
+     * runtime for testability. */
+    float        alpha
+) {
+    if (blockIdx.x != 0) return;
+    if (threadIdx.x != 0) return;
+    if (isv_signals == NULL) return;
+
+    float loss = aux_to_loss_scalar_buf[0];
+    /* Defensive NaN guard — the loss reduce flooring at fmaxf(p, 1e-30)
+     * prevents -log(0) = +inf, and fmaxf(valid, 1) prevents divide-by-
+     * zero. A NaN here would indicate an upstream bug; skip the update
+     * rather than propagating NaN into the ISV slot. */
+    if (!isfinite(loss)) return;
+
+    float ema_prev = isv_signals[ce_ema_idx];
+    if (!isfinite(ema_prev)) ema_prev = 0.0f;
+
+    /* Pearl A first-observation bootstrap: cold-start sentinel 0.0
+     * → replace directly with the first non-zero observation. The
+     * 1e-9 tolerance handles fp32 zero-equality without unintended
+     * branch into bootstrap on legitimate-tiny-loss states (a
+     * converged head producing CE ~ 0.001 would NOT trigger bootstrap
+     * — only the sentinel-exactly-zero state at fold reset does). */
+    float ema_new;
+    if (fabsf(ema_prev) < 1e-9f && loss > 0.0f) {
+        ema_new = loss;  // bootstrap
+    } else {
+        ema_new = (1.0f - alpha) * ema_prev + alpha * loss;
+    }
+
+    isv_signals[ce_ema_idx] = ema_new;
+    __threadfence_system();
+}
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index ca64943af..9e23ede27 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -2736,7 +2736,7 @@ const ISV_NETWORK_DIM: usize = 23;
 /// (and the C-side mirrors in `state_layout.cuh`). Bump this constant in the
 /// SAME commit that adds the new range, and update `layout_fingerprint_seed`
 /// to register the new slot names.
-pub(crate) const ISV_TOTAL_DIM: usize = 538;  // SP22 H6 Phase 3 (2026-05-13) adds 2 slots [536..538) for the aux→policy bypass: REWARD_AUX_ALIGN_EMA_INDEX (7th reward component EMA, anchor for SP11 controller's w_aux_align) and SP22_AUX_ALIGN_SCALE_INDEX (controller-emitted adaptive scale_β consumed by β producers at training and eval). SP21 Phase 7.5 added 8 slots [528..536) for CURRICULUM_WEIGHT_{0..8} (per_insert_pa per-segment priority boost); Phase 7 added 1 slot [527..528); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525); Phase 3 added 1 slot [520..521)
+pub(crate) const ISV_TOTAL_DIM: usize = 539;  // SP22 H6 vNext Phase F-3 (2026-05-14) adds 1 slot [538..539) for AUX_TRADE_OUTCOME_CE_EMA_INDEX (K=3 trade-outcome head's sparse-CE EMA — observability of head training during smoke). SP22 H6 Phase 3 (2026-05-13) adds 2 slots [536..538) for the aux→policy bypass: REWARD_AUX_ALIGN_EMA_INDEX (7th reward component EMA, anchor for SP11 controller's w_aux_align) and SP22_AUX_ALIGN_SCALE_INDEX (controller-emitted adaptive scale_β consumed by β producers at training and eval). SP21 Phase 7.5 added 8 slots [528..536) for CURRICULUM_WEIGHT_{0..8} (per_insert_pa per-segment priority boost); Phase 7 added 1 slot [527..528); Phase 5+6 added 2 slots [525..527); Phase 4 added 4 slots [521..525); Phase 3 added 1 slot [520..521)
 /// Legacy alias preserved for call sites that haven't been audited for the
 /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight
 /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer).
diff --git a/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs
index f966e6b32..e7ec01029 100644
--- a/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp22_isv_slots.rs
@@ -74,3 +74,37 @@ pub const REWARD_AUX_ALIGN_EMA_INDEX: usize = 536;
 /// within the first epoch after the EMA producer has accumulated a
 /// non-zero signal). FoldReset.
 pub const SP22_AUX_ALIGN_SCALE_INDEX: usize = 537;
+
+/// SP22 H6 vNext Phase F-3 (2026-05-14) — K=3 trade-outcome aux head's
+/// batch-mean sparse cross-entropy EMA. Producer: extended
+/// `aux_heads_loss_ema_update` kernel reads
+/// `aux_to_loss_scalar_buf [1]` (written by
+/// `aux_trade_outcome_loss_reduce` at the end of `aux_heads_forward`)
+/// and EMA-blends with α=0.05 (matching the K=2 / K=5 sibling slots'
+/// `AUX_NEXT_BAR_MSE_EMA_INDEX=113` / `AUX_REGIME_CE_EMA_INDEX=114`
+/// contract). Sparse-CE arithmetic: with ~95-99% bars masked, the
+/// underlying loss scalar is divided by `B_valid` (the count of
+/// non-mask labels), so the EMA tracks the per-trade-close CE rather
+/// than the per-bar marginal.
+///
+/// Cold-start sentinel = 0.0. The K=3 head's loss reduce is only
+/// meaningful when `B_valid > 0` (at least one trade closed in the
+/// replay batch); the reducer floors the divisor at 1 to avoid
+/// division by zero, so the cold-start EMA bootstraps from the first
+/// non-zero loss observation per `pearl_first_observation_bootstrap`.
+///
+/// Consumers (initial):
+/// - HEALTH_DIAG console output (extends the existing `aux_loss` /
+///   `regime_ce` columns to include `outcome_ce` — Phase F-3 follow-up
+///   commit lands the snap layout extension).
+/// - Smoke validation observability: lets us see whether the head is
+///   actually learning during Phase F smoke runs. A CE that decreases
+///   from ~1.10 (cold-start uniform ~ln(3) ≈ 1.098) toward something
+///   meaningfully lower (e.g., 0.5-0.7) is evidence the head is
+///   learning real trade-outcome structure; CE pinned at ln(3) for
+///   many epochs is evidence the head can't learn (label noise or
+///   under-capacity).
+///
+/// FoldReset sentinel 0.0 — reset registry entry forthcoming alongside
+/// the producer wireup commit.
+pub const AUX_TRADE_OUTCOME_CE_EMA_INDEX: usize = 538;
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 2e6318bf7..8fe6d01d0 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -18226,3 +18226,30 @@ The K=3 head now influences policy behavior via TWO paths:
 - **Phase E**: dW backward gradient validation — the Phase D dW kernel emits 12 partials; Phase E would add unit tests confirming the gradient direction is correct + Adam updates the prior in the right direction.
 - **Phase F**: Validation smoke at structural prior — first end-to-end smoke run with the full K=3 + plan-conditioning + atom-shift stack. Decisive test of the spec's hypothesis.
 - **B5b-2** (deferred): Collector trade plan launch (resolves train/inference asymmetry from B5b).
+
+#### Phase F-3 (kernel + ISV slot) — K=3 trade-outcome CE EMA producer (2026-05-14)
+
+Adds the observability infrastructure for Phase F's smoke validation: a new ISV slot (`AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538`) tracks the K=3 head's batch-mean sparse cross-entropy EMA. Producer kernel `aux_outcome_ce_ema_update` is registered + cubin-built; launcher wireup follows in a Phase F-3b commit alongside HEALTH_DIAG console-output extension.
+
+**Changes**:
+- `sp22_isv_slots.rs`: new `pub const AUX_TRADE_OUTCOME_CE_EMA_INDEX: usize = 538`
+- `gpu_dqn_trainer.rs::ISV_TOTAL_DIM`: bumped 538 → 539 (allocates slot 538 in the bus)
+- New kernel `aux_outcome_ce_ema_kernel.cu`: single-thread single-block direct-to-ISV EMA writer with Pearl A first-observation bootstrap. Fixed α=0.05 (slow-moving observability signal). NULL-tolerant on `isv_signals`. NaN-guarded.
+- `build.rs`: kernel registered; cubin compiles (3.2 KB)
+
+**Why dedicated kernel (not extension of `aux_heads_loss_ema_update`)**:
+The existing K=2/K=5 aux-loss EMA writes through Pearls A+D's 2-stage producer scratch buffer (`producer_step_scratch_buf`). Extending it would require allocating a new scratch slot, threading Pearls A+D mapping, and touching `apply_pearls_ad_kernel`. The K=3 head's CE is for OBSERVABILITY ONLY in this commit — no Pearls A+D adaptive α needed. The dedicated direct-to-ISV kernel is the minimum-scope path. If K=3 CE later becomes a controller anchor (e.g., for a SP11-style adaptive weight on K=3 alignment reward), it can be re-routed through Pearls A+D at that point.
+
+**Smoke validation signal interpretation**:
+- Cold-start: `ISV[538] = 0.0` (FoldReset sentinel)
+- First step with `B_valid > 0`: Pearl A bootstrap fires → `ISV[538] = first_CE_observation`
+- Typical cold-start CE for uniform K=3 prediction: `ln(3) ≈ 1.098`
+- **Expected learning signal**: CE drops from ~1.098 toward something meaningfully lower (e.g., 0.5-0.7) over the first few epochs — evidence the K=3 head is learning real trade-outcome structure
+- **Spec falsification signal**: CE pinned at ~ln(3) for many epochs → head can't learn (label noise, under-capacity, or trade outcomes are genuinely market-state-unconditional)
+
+**Phase F-3 follow-up (deferred to next commit)**:
+- Trainer-side launcher call: append `aux_outcome_ce_ema_update` after the existing `aux_heads_loss_ema_update` launch in the per-step training graph
+- HEALTH_DIAG snap layout extension: add `aux_outcome_ce_ema` field + console column
+- State reset registry entry: `AUX_TRADE_OUTCOME_CE_EMA_INDEX` as FoldReset sentinel 0
+
+**Verification**: `cargo check -p ml` clean; `cargo test -p ml --lib` → **1016/0 green**.