diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index e58f18d36..73a56cbc8 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -894,6 +894,18 @@ fn main() {
         // Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
         // § Phase 0 Task 0.1.
         "reward_decomp_diag_kernel.cu",
+        // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error
+        // magnitude EMA producer. Single-block 256-thread kernel
+        // computing mean(|td_errors[b]|) over a B-element batch and
+        // EMA-blending into ISV[TD_ERROR_MAG_EMA_INDEX=493] via
+        // Pearl-A first-observation bootstrap (sentinel 0.0) + fixed
+        // α=WELFORD_ALPHA_MIN=0.4 (per
+        // `pearl_wiener_alpha_floor_for_nonstationary` — Welford-α
+        // accumulators at [498..504) are reserved for Phase 4
+        // q_next_target chain). Block tree-reduce (no atomicAdd).
+        // Pure observability: pre-fix baseline for the B-DD9 ratio
+        // gate. Plan: SP18 Phase 0 Task 0.2.
+        "td_error_mag_ema_kernel.cu",
         // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
         // Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
         // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 9e6c4d7f6..64fc8753f 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -123,6 +123,13 @@ pub(crate) static REWARD_COMPONENT_EMA_CUBIN: &[u8] = include_bytes!(concat!(env
 /// reduce — 4 blocks × 256 threads. Pure observability — no
 /// production-path consumer in this commit. Plan: SP18 Phase 0 Task 0.1.
 pub(crate) static REWARD_DECOMP_DIAG_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_decomp_diag_kernel.cubin"));
+/// SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error magnitude
+/// EMA producer. Single-block 256-thread kernel computing
+/// mean(|td_errors[b]|) over [B] and EMA-blending into
+/// ISV[TD_ERROR_MAG_EMA_INDEX=493] via Pearl-A first-observation
+/// bootstrap + fixed α=0.4. Pure observability: pre-fix baseline for
+/// the B-DD9 ratio gate. Plan: SP18 Phase 0 Task 0.2.
+pub(crate) static TD_ERROR_MAG_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/td_error_mag_ema_kernel.cubin"));
 /// Plan 3 Task 3 B.2: Flat→Positioned transition rate EMA into ISV[71].
 pub(crate) static TRADE_RATE_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/trade_rate_ema_kernel.cubin"));
 /// Plan 3 Task 4 B.4: per-batch readiness EMA + derived plan_threshold.
@@ -6813,6 +6820,20 @@ pub struct GpuDqnTrainer {
     /// once at construct, written-then-read inside the kernel each launch.
     sp17_advantage_clip_scratch: super::mapped_pinned::MappedF32Buffer,
 
+    // ── SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error magnitude
+    //    EMA producer kernel handle ─────────────────────────────────────
+    /// Kernel handle for `td_error_mag_ema_update`. Reads
+    /// `td_errors_buf [B]` (post-train-step, populated by `c51_loss` /
+    /// `mse_loss` per backward configuration) and EMA-blends
+    /// `mean(|td_errors[b]|)` into `ISV[TD_ERROR_MAG_EMA_INDEX=493]`
+    /// via Pearl-A first-observation bootstrap + fixed α=0.4. Loaded
+    /// from `td_error_mag_ema_kernel.cubin`. Consumed by
+    /// `launch_sp18_td_error_mag_ema_update()` (cold-path post-train-
+    /// step launch from `training_loop.rs`). Pure observability —
+    /// pre-fix baseline for the B-DD9 ratio gate. Plan: SP18 Phase 0
+    /// Task 0.2.
+    sp18_td_error_mag_ema_kernel: CudaFunction,
+
     // ── SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
     //    decomposition diagnostic buffer ──────────────────────────────────
     /// 20-float mapped-pinned diagnostic buffer for the SP18 D-leg
@@ -9435,6 +9456,82 @@ impl GpuDqnTrainer {
         Ok(self.read_isv_signal_at(ADVANTAGE_CLIP_BOUND_INDEX))
     }
 
+    /// SP18 v2 Phase 0 Task 0.2 (2026-05-08) — B-leg TD-error magnitude
+    /// EMA producer launch.
+    ///
+    /// Cold-path (post-train-step) launcher for
+    /// `td_error_mag_ema_update`. Reads the trainer-owned
+    /// `td_errors_buf [B]` (populated by the C51/MSE loss kernel each
+    /// training step, consumed in-place by `seg_tree_update` for PER
+    /// priority recomputation), block tree-reduces
+    /// `mean(|td_errors[b]|)`, applies Pearl-A first-observation
+    /// bootstrap (sentinel 0.0 → REPLACE) and fixed α=0.4 EMA blend,
+    /// and writes the result to `ISV[TD_ERROR_MAG_EMA_INDEX=493]`.
+    ///
+    /// Pure observability: pre-fix baseline for the B-DD9 ratio gate
+    /// (`avg(|TD-error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`).
+    /// Phase 0 does NOT yet use the Welford-derived α from the TDB_*
+    /// accumulators in slots [498..504) — those are reserved for Phase
+    /// 4 q_next_target Wiener-α blending.
+    ///
+    /// Returns Ok(()) on launch success; the host-side HEALTH_DIAG emit
+    /// reads the slot via `read_isv_signal_at(TD_ERROR_MAG_EMA_INDEX)`
+    /// after stream sync.
+    pub fn launch_sp18_td_error_mag_ema_update(&self) -> Result<(), MLError> {
+        use super::sp14_isv_slots::{
+            TD_ERROR_MAG_EMA_INDEX, SENTINEL_TD_ERROR_MAG_EMA, WELFORD_ALPHA_MIN,
+        };
+        let b = self.config.batch_size;
+        if b == 0 || self.isv_signals_dev_ptr == 0 {
+            return Ok(());
+        }
+        let td_ptr = self.td_errors_buf.raw_ptr();
+        let isv_ptr = self.isv_signals_dev_ptr;
+        let b_i32 = b as i32;
+        let ema_idx = TD_ERROR_MAG_EMA_INDEX as i32;
+        let alpha: f32 = WELFORD_ALPHA_MIN;
+        let sentinel: f32 = SENTINEL_TD_ERROR_MAG_EMA;
+        // Single-block, 256 threads. Dynamic shmem = BLOCK_SIZE × f32.
+        const BLOCK_DIM: u32 = 256;
+        let shmem_bytes: u32 = BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+        unsafe {
+            self.stream
+                .launch_builder(&self.sp18_td_error_mag_ema_kernel)
+                .arg(&td_ptr)
+                .arg(&b_i32)
+                .arg(&isv_ptr)
+                .arg(&ema_idx)
+                .arg(&alpha)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .map_err(|e| MLError::ModelError(format!(
+                    "sp18 td_error_mag_ema_update launch: {e}"
+                )))?;
+        }
+        Ok(())
+    }
+
+    /// SP18 v2 Phase 0 Task 0.2 (2026-05-08) — B-leg TD-error magnitude
+    /// EMA readback.
+    ///
+    /// Convenience wrapper: launches the producer, syncs the stream,
+    /// returns the post-blend ISV slot 493 value. Cold-path only
+    /// (HEALTH_DIAG cadence).
+    pub fn read_sp18_td_error_mag_ema(&self) -> Result<f32, MLError> {
+        use super::sp14_isv_slots::TD_ERROR_MAG_EMA_INDEX;
+        self.launch_sp18_td_error_mag_ema_update()?;
+        self.stream
+            .synchronize()
+            .map_err(|e| MLError::ModelError(format!(
+                "sp18 td_error_mag_ema sync: {e}"
+            )))?;
+        Ok(self.read_isv_signal_at(TD_ERROR_MAG_EMA_INDEX))
+    }
+
     /// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
     /// decomposition diagnostic device pointer.
     ///
@@ -20620,6 +20717,20 @@ impl GpuDqnTrainer {
         let sp17_advantage_clip_bound_kernel = sp17_advantage_clip_bound_module
             .load_function("sp17_advantage_clip_bound_update")
             .map_err(|e| MLError::ModelError(format!("sp17_advantage_clip_bound_update: {e}")))?;
+
+        // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error
+        // magnitude EMA producer load. Single block × 256 threads.
+        // Pearl-A bootstrap on slot 493 (sentinel 0.0) + fixed
+        // α=WELFORD_ALPHA_MIN=0.4. Reads `td_errors_buf [B]` post-
+        // train-step. Pure observability.
+        let sp18_td_error_mag_ema_module = stream.context()
+            .load_cubin(TD_ERROR_MAG_EMA_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!(
+                "sp18 td_error_mag_ema cubin: {e}")))?;
+        let sp18_td_error_mag_ema_kernel = sp18_td_error_mag_ema_module
+            .load_function("td_error_mag_ema_update")
+            .map_err(|e| MLError::ModelError(format!(
+                "td_error_mag_ema_update load: {e}")))?;
         let sp17_clip_scratch_capacity = config.batch_size
             .saturating_mul(
                 config.branch_0_size + config.branch_1_size
@@ -24901,6 +25012,11 @@ impl GpuDqnTrainer {
             sp17_v_share_kernel,
             sp17_advantage_clip_bound_kernel,
             sp17_advantage_clip_scratch,
+            // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg TD-error
+            // magnitude EMA producer kernel handle. Trainer-stream-
+            // resident; launched from `launch_sp18_td_error_mag_ema_
+            // update` post-train-step.
+            sp18_td_error_mag_ema_kernel,
             // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action
             // reward decomposition diagnostic buffer. Mapped-pinned 20
             // floats (4 bins × 5 stats). Written by collector kernel,
diff --git a/crates/ml/src/cuda_pipeline/td_error_mag_ema_kernel.cu b/crates/ml/src/cuda_pipeline/td_error_mag_ema_kernel.cu
new file mode 100644
index 000000000..f29f340e3
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/td_error_mag_ema_kernel.cu
@@ -0,0 +1,126 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * SP18 v2 Phase 0 Task 0.2 (2026-05-08) — B-leg TD-error magnitude EMA
+ * producer.
+ *
+ * Computes mean(|td_errors[b]|) over a B-element batch and EMA-blends
+ * into ISV[TD_ERROR_MAG_EMA_INDEX=493] via Pearl-A first-observation
+ * bootstrap + fixed α=WELFORD_ALPHA_MIN=0.4 (per
+ * `pearl_wiener_alpha_floor_for_nonstationary` — non-stationary control
+ * loop floor; Phase 0 does NOT yet wire the Wiener-α Welford
+ * accumulators that the Phase 4 q_next_target chain will use).
+ *
+ * Cadence: post-train-step. The `td_errors_buf` is populated by
+ * `c51_loss` (or `mse_loss` per backward configuration) on each
+ * training step, and consumed in-place by `seg_tree_update` for PER
+ * priority recomputation. Reading it for the magnitude EMA is purely
+ * additive — the buffer is not modified, only summed.
+ *
+ * The EMA is the **pre-fix** baseline for the B-DD9 ratio gate: at the
+ * post-fix epoch (Phase 5), the same producer fires on the new
+ * `q_next = Q_target(s', argmax_a Q_online(s', a))` bootstrap path and
+ * the ratio `EMA_post / EMA_pre ∈ [0.5, 5.0]` confirms the bootstrap
+ * change did not blow up the loss magnitude.
+ *
+ * ── Pearls + invariants ────────────────────────────────────────────────
+ *  - `feedback_no_atomicadd.md` — block tree-reduce on `|td_errors|`,
+ *    no atomicAdd.
+ *  - `pearl_first_observation_bootstrap.md` — sentinel
+ *    SENTINEL_TD_ERROR_MAG_EMA = 0.0; cold-start REPLACE on first
+ *    observation. Identifies sentinel by exact-zero match (consistent
+ *    with the SP14 P0-A `REWARD_POS_CAP_ADAPTIVE` precedent).
+ *  - `pearl_wiener_alpha_floor_for_nonstationary.md` — α floored at
+ *    0.4. Phase 0 does NOT use the Welford-derived α (the TDB_*
+ *    accumulator slots [498..504) are reserved for Phase 4
+ *    q_next_target Wiener-α). The fixed α=0.4 keeps the EMA
+ *    responsive across the 1-epoch dispatch the reviewer runs at
+ *    Task 0.3.
+ *  - `pearl_no_host_branches_in_captured_graph.md` — pure single-
+ *    block kernel; runs outside `exp_fwd_graph` capture (post-train-
+ *    step boundary, similar to SP17 dueling diag launches).
+ *  - `feedback_no_stubs.md` — full reduction body, no zero-output
+ *    placeholder.
+ *  - `feedback_no_htod_htoh_only_mapped_pinned.md` — `td_errors`
+ *    is a device buffer (`CudaSlice<f32>::raw_ptr`); `isv` is the
+ *    mapped-pinned ISV bus device pointer.
+ *
+ * Algorithm (single-block, BLOCK_SIZE=256):
+ *
+ *   Pass 1: each thread strides over `td_errors`, accumulating
+ *           `|td_errors[i]|` into a local register.
+ *   Pass 2: shared-memory tree-reduce → block sum.
+ *   Pass 3 (thread 0 only):
+ *           mean = block_sum / max(B, 1)
+ *           if (current ≈ sentinel) blended = mean (Pearl-A bootstrap)
+ *           else                    blended = (1-α) × current + α × mean
+ *           ISV[ema_idx] = blended
+ *           __threadfence_system()
+ *
+ * Args:
+ *   td_errors      — `[B]` f32 device pointer (typically
+ *                    `td_errors_buf.raw_ptr()` from
+ *                    `GpuDqnTrainer`). Read-only.
+ *   B              — sample count.
+ *   isv            — `[ISV_TOTAL_DIM]` device f32 pointer (mapped-
+ *                    pinned). Modified in place at `ema_idx`.
+ *   ema_idx        — TD_ERROR_MAG_EMA_INDEX (493).
+ *   alpha          — EMA blend factor (host-passed; production uses
+ *                    WELFORD_ALPHA_MIN=0.4).
+ *   sentinel       — SENTINEL_TD_ERROR_MAG_EMA = 0.0 — Pearl-A
+ *                    direct-replace marker.
+ *
+ * Launch: grid=(1, 1, 1), block=(256, 1, 1).
+ * Dynamic shared memory: BLOCK_SIZE × sizeof(float) for the reduce tile.
+ * ══════════════════════════════════════════════════════════════════════════ */
+
+#include <cuda_runtime.h>
+
+#define SP18_TD_ERR_BLOCK_SIZE 256
+#define SP18_TD_ERR_EPS_F      1e-6f
+
+extern "C" __global__ void td_error_mag_ema_update(
+    const float* __restrict__ td_errors,   /* [B] read-only */
+    int          B,
+    float*       __restrict__ isv,
+    int          ema_idx,
+    float        alpha,
+    float        sentinel)
+{
+    extern __shared__ float s_reduce[];   /* [BLOCK_SIZE] */
+    const int tid = (int)threadIdx.x;
+
+    if (blockIdx.x != 0) return;
+    if (B <= 0) return;
+
+    /* Pass 1: per-thread strided sum of |td_errors[i]|. */
+    float local = 0.0f;
+    for (int i = tid; i < B; i += SP18_TD_ERR_BLOCK_SIZE) {
+        local += fabsf(td_errors[i]);
+    }
+    s_reduce[tid] = local;
+    __syncthreads();
+
+    /* Pass 2: tree-reduce. */
+    for (int s = SP18_TD_ERR_BLOCK_SIZE / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            s_reduce[tid] += s_reduce[tid + s];
+        }
+        __syncthreads();
+    }
+
+    /* Pass 3: thread 0 finalises, applies Pearl-A bootstrap + EMA blend. */
+    if (tid == 0) {
+        const float mean = s_reduce[0] / (float)B;
+        const float current = isv[ema_idx];
+        float blended;
+        /* Pearl-A first-observation bootstrap: sentinel exact-zero match.
+         * Any value within EPS of sentinel is treated as "first
+         * observation" and REPLACES directly. */
+        if (fabsf(current - sentinel) < SP18_TD_ERR_EPS_F) {
+            blended = mean;
+        } else {
+            blended = (1.0f - alpha) * current + alpha * mean;
+        }
+        isv[ema_idx] = blended;
+        __threadfence_system();   /* PCIe-visible to mapped-pinned host_ptr */
+    }
+}
diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs
index 77af9608c..2014afc16 100644
--- a/crates/ml/src/trainers/dqn/trainer/constructor.rs
+++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs
@@ -588,6 +588,12 @@ impl DQNTrainer {
             training_sharpe_ema_initialized: false,
             last_action_entropy: None,
             last_magnitude_dist: [0.0_f32; 3],
+            // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg V_SHARE
+            // trajectory ring buffer initialised to NaN. Entries are
+            // pushed each epoch by the HEALTH_DIAG emit; the slope
+            // calculation skips the emit until 5 epochs of history
+            // have accumulated (older entries still NaN ⇒ skip).
+            sp18_v_share_history: [[f32::NAN; 4]; 5],
             // SP8 (Fix 36, 2026-05-03): `last_train_active_frac` host-cached
             // field deleted; the canary now lives at
             // `ISV[TRAIN_ACTIVE_FRAC_INDEX]` and is read via the accessor
diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs
index 42ab86dc0..b8b1a5335 100644
--- a/crates/ml/src/trainers/dqn/trainer/mod.rs
+++ b/crates/ml/src/trainers/dqn/trainer/mod.rs
@@ -331,6 +331,19 @@ pub struct DQNTrainer {
     /// summed across all direction bins, normalized to [0, 1].
     pub(crate) last_magnitude_dist: [f32; 3],
 
+    /// SP18 v2 Phase 0 Task 0.2 (2026-05-08) — B-leg V_SHARE trajectory
+    /// history. Fixed-size ring buffer of the last 5 epochs of
+    /// per-branch V_SHARE readings (slot indices [478..482) — dir / mag /
+    /// ord / urg). Index 0 = oldest, index 4 = most recent. Used by
+    /// the per-epoch HEALTH_DIAG emit to compute slope-across-4-epochs
+    /// `(EMA[now] - EMA[now-4]) / 4` per branch and write the dir-branch
+    /// slope to `ISV[V_SHARE_TREND_DIAG_INDEX=496]` for the SP18 KILL
+    /// CRITERION. Initialised to NaN; entries marked NaN before the
+    /// fifth epoch indicate the slope is not yet computable. Plan:
+    /// docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
+    /// § Phase 0 Task 0.2.
+    pub(crate) sp18_v_share_history: [[f32; 4]; 5],
+
     /// Plan C Task 5 — fraction of training-rollout actions where direction ∈
     /// {Short, Long} (i.e. the policy committed to a directional bet rather
     /// than Hold/Flat). Counterpart to the kernel-derived val_active_frac
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 33b61799d..65a39a6d7 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -5426,6 +5426,118 @@ impl DQNTrainer {
                 );
             }
 
+            // SP18 v2 Phase 0 Task 0.2 (2026-05-08): B-leg V_SHARE
+            // trajectory + TD-error magnitude HEALTH_DIAG emit.
+            //
+            // Two diagnostic lines:
+            //
+            //   v_share_traj  [dir_slope=X mag_slope=Y ord_slope=Z urg_slope=W]
+            //   td_error_pre  [magnitude_ema=X]
+            //
+            // V_SHARE trajectory:
+            //   - V_SHARE EMA already exists at slots [478..482) per SP17
+            //     Phase 3.2 (`launch_v_share_update` fires per-epoch and
+            //     blends `|V| / (|V| + |A_centered|)` per branch into
+            //     these slots).
+            //   - This block reads the just-updated values, pushes onto
+            //     a 5-entry ring buffer (`sp18_v_share_history`), and
+            //     emits slope-across-4-epochs `(EMA[now] - EMA[now-4]) / 4`
+            //     per branch.
+            //   - The dir-branch slope is also written to
+            //     `ISV[V_SHARE_TREND_DIAG_INDEX=496]` so the host-side
+            //     KILL CRITERION + future Phase 5 PopArt-reset gate can
+            //     read it via the standard ISV bus.
+            //   - Until the ring buffer fills (epochs 0–3), the slope is
+            //     emitted as `nan` and no ISV write fires (Pearl-A
+            //     direct-replace on first valid observation kicks in at
+            //     epoch 4).
+            //
+            // TD-error magnitude EMA:
+            //   - Producer kernel `td_error_mag_ema_update` reads
+            //     `td_errors_buf [B]` (already populated by C51/MSE loss
+            //     for PER priority recomputation), block tree-reduces
+            //     `mean(|td_errors|)`, applies Pearl-A bootstrap +
+            //     fixed α=0.4 EMA blend, writes
+            //     `ISV[TD_ERROR_MAG_EMA_INDEX=493]`.
+            //   - The emit reads the post-blend slot value.
+            //   - Pre-fix baseline: this is the `q_next = rewards`
+            //     self-bootstrap regime. Phase 5 swaps to
+            //     `q_next = Q_target(s', argmax_a Q_online(s', a))`
+            //     and the same producer fires; the B-DD9 ratio gate
+            //     compares post-fix / pre-fix at HD0–HD1.
+            //
+            // Both observables are POST-PRODUCER reads — the V_SHARE
+            // producers fire above (in the dueling block) and the
+            // TD-error producer fires inside `read_sp18_td_error_mag_ema()`.
+            {
+                use crate::cuda_pipeline::sp14_isv_slots::{
+                    V_SHARE_TREND_DIAG_INDEX,
+                };
+                // Push the most-recent V_SHARE readings into the ring.
+                // The SP17 V_SHARE producer is launched above (in the
+                // dueling diag block via `read_v_share_per_branch`); we
+                // re-read the freshly-updated values here. Shift-left
+                // (oldest at index 0 drops; newest at index 4).
+                let v_share_now = if let Some(ref fused) = self.fused_ctx {
+                    fused.trainer().read_v_share_per_branch()
+                        .unwrap_or([f32::NAN; 4])
+                } else {
+                    [f32::NAN; 4]
+                };
+                self.sp18_v_share_history[0] = self.sp18_v_share_history[1];
+                self.sp18_v_share_history[1] = self.sp18_v_share_history[2];
+                self.sp18_v_share_history[2] = self.sp18_v_share_history[3];
+                self.sp18_v_share_history[3] = self.sp18_v_share_history[4];
+                self.sp18_v_share_history[4] = v_share_now;
+                let now    = self.sp18_v_share_history[4];
+                let now_m4 = self.sp18_v_share_history[0];
+                let mut slope = [f32::NAN; 4];
+                let mut all_finite = true;
+                for b in 0..4 {
+                    if now[b].is_finite() && now_m4[b].is_finite() {
+                        slope[b] = (now[b] - now_m4[b]) / 4.0;
+                    } else {
+                        all_finite = false;
+                    }
+                }
+                // Write dir-branch slope to ISV[496] when finite (Pearl-A
+                // direct-replace on first valid observation: sentinel
+                // `SENTINEL_V_SHARE_TREND_DIAG=0.0` is the cold-start; we
+                // overwrite with the real slope each epoch the ring is
+                // full).
+                if all_finite {
+                    if let Some(ref fused) = self.fused_ctx {
+                        fused.trainer().write_isv_signal_at(
+                            V_SHARE_TREND_DIAG_INDEX, slope[0],
+                        );
+                    }
+                }
+                tracing::info!(
+                    "HEALTH_DIAG[{}]: v_share_traj [dir_slope={:.6} mag_slope={:.6} \
+                     ord_slope={:.6} urg_slope={:.6}]",
+                    epoch, slope[0], slope[1], slope[2], slope[3],
+                );
+
+                // TD-error magnitude EMA: launch producer + read ISV slot.
+                let td_mag = if let Some(ref fused) = self.fused_ctx {
+                    match fused.trainer().read_sp18_td_error_mag_ema() {
+                        Ok(v) => v,
+                        Err(e) => {
+                            tracing::warn!(
+                                "HEALTH_DIAG read_sp18_td_error_mag_ema refresh failed: {e}"
+                            );
+                            f32::NAN
+                        }
+                    }
+                } else {
+                    f32::NAN
+                };
+                tracing::info!(
+                    "HEALTH_DIAG[{}]: td_error_pre [magnitude_ema={:.6}]",
+                    epoch, td_mag,
+                );
+            }
+
             // SP7 observability (grad_decomp_pinned): surface the exact contents of
             // grad_decomp_result_pinned[0..3, 12..15, 36..48] (iqn/cql_raw/c51 ×
             // [mag, dir, trunk]) that the SP7 loss-balance controller read in the
diff --git a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs
index 67d5c68ea..12b718e66 100644
--- a/crates/ml/tests/sp18_hold_reward_oracle_tests.rs
+++ b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs
@@ -268,6 +268,156 @@ mod gpu {
         }
     }
 
+    /// Cubin handle for the SP18 B-leg TD-error magnitude EMA producer.
+    const SP18_TD_ERROR_MAG_EMA_CUBIN: &[u8] =
+        include_bytes!(concat!(env!("OUT_DIR"), "/td_error_mag_ema_kernel.cubin"));
+
+    fn load_td_error_mag_ema(stream: &Arc<CudaStream>) -> CudaFunction {
+        let module = stream
+            .context()
+            .load_cubin(SP18_TD_ERROR_MAG_EMA_CUBIN.to_vec())
+            .expect("load td_error_mag_ema_kernel cubin");
+        module
+            .load_function("td_error_mag_ema_update")
+            .expect("load td_error_mag_ema_update function")
+    }
+
+    /// SP18 v2 Phase 0 Task 0.2 — TD-error magnitude EMA Pearl-A
+    /// bootstrap test.
+    ///
+    /// Construct a synthetic td_errors buffer with known magnitudes,
+    /// pre-populate the ISV slot with the sentinel (0.0), launch the
+    /// kernel, and assert the slot equals the closed-form mean(|td|)
+    /// (Pearl-A direct-replace on first observation).
+    ///
+    /// td_errors = [-1.0, +2.0, -0.5, +0.5, -1.0, +1.0, 0.0, +3.0]
+    /// mean(|td|) = (1.0+2.0+0.5+0.5+1.0+1.0+0.0+3.0) / 8 = 9.0/8 = 1.125
+    #[test]
+    #[ignore = "requires GPU"]
+    fn td_error_mag_ema_pearl_a_bootstrap() {
+        use ml::cuda_pipeline::sp14_isv_slots::{
+            TD_ERROR_MAG_EMA_INDEX, SENTINEL_TD_ERROR_MAG_EMA, WELFORD_ALPHA_MIN,
+        };
+
+        let stream = make_test_stream();
+        let kernel = load_td_error_mag_ema(&stream);
+
+        const B: usize = 8;
+        let td_errors_host: [f32; B] = [-1.0, 2.0, -0.5, 0.5, -1.0, 1.0, 0.0, 3.0];
+        let expected_mean: f32 = (1.0 + 2.0 + 0.5 + 0.5 + 1.0 + 1.0 + 0.0 + 3.0) / B as f32;
+
+        let td_buf = unsafe { MappedF32Buffer::new(B) }.expect("alloc td buf");
+        td_buf.write_from_slice(&td_errors_host);
+
+        // ISV bus: pre-populate slot 493 with sentinel 0.0 so Pearl-A
+        // bootstrap fires (first observation REPLACES directly).
+        const ISV_SPAN: usize = 505;   // matches ISV_TOTAL_DIM post-PP.2
+        let mut isv_host = vec![0.0_f32; ISV_SPAN];
+        isv_host[TD_ERROR_MAG_EMA_INDEX] = SENTINEL_TD_ERROR_MAG_EMA;
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }.expect("alloc isv buf");
+        isv_buf.write_from_slice(&isv_host);
+
+        let b_i32 = B as i32;
+        let ema_idx_i32: i32 = TD_ERROR_MAG_EMA_INDEX as i32;
+        let alpha: f32 = WELFORD_ALPHA_MIN;
+        let sentinel: f32 = SENTINEL_TD_ERROR_MAG_EMA;
+        const BLOCK_DIM: u32 = 256;
+        let shmem_bytes: u32 = BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&td_buf.dev_ptr)
+                .arg(&b_i32)
+                .arg(&isv_buf.dev_ptr)
+                .arg(&ema_idx_i32)
+                .arg(&alpha)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch td_error_mag_ema_update");
+        }
+        stream.synchronize().expect("sync");
+
+        let isv_after = isv_buf.read_all();
+        let got = isv_after[TD_ERROR_MAG_EMA_INDEX];
+        let eps = 1e-6_f32;
+        assert!(
+            (got - expected_mean).abs() < eps,
+            "TD-error mag EMA Pearl-A bootstrap: expected {expected_mean:.6} \
+             (== mean|td| since slot was at sentinel), got {got:.6}"
+        );
+    }
+
+    /// SP18 v2 Phase 0 Task 0.2 — TD-error magnitude EMA blend math.
+    ///
+    /// Pre-populate the ISV slot with a non-sentinel value (1.0); launch
+    /// the kernel against `td_errors = [+0.5, -0.5, +0.5, -0.5]` so
+    /// `mean(|td|) = 0.5`; assert the result equals the closed-form
+    /// blend `(1 - α) × 1.0 + α × 0.5` with α=0.4.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn td_error_mag_ema_blend_post_bootstrap() {
+        use ml::cuda_pipeline::sp14_isv_slots::{
+            TD_ERROR_MAG_EMA_INDEX, SENTINEL_TD_ERROR_MAG_EMA, WELFORD_ALPHA_MIN,
+        };
+
+        let stream = make_test_stream();
+        let kernel = load_td_error_mag_ema(&stream);
+
+        const B: usize = 4;
+        let td_errors_host: [f32; B] = [0.5, -0.5, 0.5, -0.5];
+        let mean_abs: f32 = 0.5;
+        let prev: f32 = 1.0;
+        let alpha = WELFORD_ALPHA_MIN;
+        let expected = (1.0 - alpha) * prev + alpha * mean_abs;
+
+        let td_buf = unsafe { MappedF32Buffer::new(B) }.expect("alloc td buf");
+        td_buf.write_from_slice(&td_errors_host);
+
+        const ISV_SPAN: usize = 505;
+        let mut isv_host = vec![0.0_f32; ISV_SPAN];
+        isv_host[TD_ERROR_MAG_EMA_INDEX] = prev;   // NOT sentinel — blend path
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }.expect("alloc isv buf");
+        isv_buf.write_from_slice(&isv_host);
+
+        let b_i32 = B as i32;
+        let ema_idx_i32: i32 = TD_ERROR_MAG_EMA_INDEX as i32;
+        let sentinel: f32 = SENTINEL_TD_ERROR_MAG_EMA;
+        const BLOCK_DIM: u32 = 256;
+        let shmem_bytes: u32 = BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&td_buf.dev_ptr)
+                .arg(&b_i32)
+                .arg(&isv_buf.dev_ptr)
+                .arg(&ema_idx_i32)
+                .arg(&alpha)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch td_error_mag_ema_update");
+        }
+        stream.synchronize().expect("sync");
+
+        let isv_after = isv_buf.read_all();
+        let got = isv_after[TD_ERROR_MAG_EMA_INDEX];
+        let eps = 1e-6_f32;
+        assert!(
+            (got - expected).abs() < eps,
+            "TD-error mag EMA post-bootstrap blend: expected {expected:.6} \
+             (= (1 - {alpha}) × {prev} + {alpha} × {mean_abs}), got {got:.6}"
+        );
+    }
+
     /// Empty-bin guard test — when no samples land in a given bin,
     /// the kernel must emit 0.0 for every column (NOT NaN, which would
     /// be the natural division-by-zero default). The HEALTH_DIAG
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 3d7562c54..4818669e3 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,73 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+## 2026-05-08 — SP18 v2 Phase 0 Task 0.2: B-leg V_SHARE trajectory + TD-error magnitude diagnostic
+
+Phase 0 observability: kernel + Rust launcher + V_SHARE history ring buffer + per-epoch HEALTH_DIAG emit + GPU oracle tests. All in the same atomic commit per `feedback_no_partial_refactor`.
+
+**Files added/modified:**
+
+- NEW `crates/ml/src/cuda_pipeline/td_error_mag_ema_kernel.cu` — single-block 256-thread kernel reading `td_errors_buf [B]` (populated by C51/MSE loss for PER priority recomputation, post-train-step), block tree-reducing `mean(|td_errors[b]|)` (no atomicAdd), and EMA-blending into `ISV[TD_ERROR_MAG_EMA_INDEX=493]` via Pearl-A first-observation bootstrap (sentinel 0.0 → REPLACE) + fixed α=`WELFORD_ALPHA_MIN=0.4` per `pearl_wiener_alpha_floor_for_nonstationary`.
+- `crates/ml/build.rs` — cubin manifest entry for `td_error_mag_ema_kernel.cu`.
+- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`:
+  - `TD_ERROR_MAG_EMA_CUBIN` re-export.
+  - `sp18_td_error_mag_ema_kernel: CudaFunction` field + cubin load on the trainer's stream.
+  - `launch_sp18_td_error_mag_ema_update()` cold-path launcher (single-block × 256 threads, dynamic shmem = `BLOCK × sizeof(f32)`).
+  - `read_sp18_td_error_mag_ema()` convenience wrapper (launch → sync → read ISV slot 493).
+- `crates/ml/src/trainers/dqn/trainer/mod.rs`:
+  - New `sp18_v_share_history: [[f32; 4]; 5]` field — fixed-size ring buffer of the last 5 epochs of per-branch V_SHARE EMA readings (slots [478..482) per SP17 Phase 3.2). Used to compute the `(EMA[now] - EMA[now-4]) / 4` slope per branch.
+- `crates/ml/src/trainers/dqn/trainer/constructor.rs` — initialise the ring buffer to `[[NaN; 4]; 5]`. Slope emit guards against NaN entries until the ring fills (epochs 0–3 emit `nan` as the slope, no ISV write fires; epoch 4 onward writes the dir-branch slope to `ISV[V_SHARE_TREND_DIAG_INDEX=496]`).
+- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — two new HEALTH_DIAG lines at the per-epoch boundary (right after the SP18 reward_decomp line):
+
+  ```
+  HEALTH_DIAG[N]: v_share_traj [dir_slope=X mag_slope=Y ord_slope=Z urg_slope=W]
+  HEALTH_DIAG[N]: td_error_pre [magnitude_ema=X]
+  ```
+
+  V_SHARE slope is host-side computation against the ring buffer (`(now - now_m4) / 4` per branch). TD-error magnitude is post-blend ISV slot 493 read (producer fires inside `read_sp18_td_error_mag_ema()`).
+
+- `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
+  - NEW `td_error_mag_ema_pearl_a_bootstrap` (`#[ignore = "requires GPU"]`) — synthetic `td_errors=[-1, +2, -0.5, +0.5, -1, +1, 0, +3]`; pre-populate ISV slot with sentinel 0.0; assert post-launch slot equals the closed-form `mean(|td|) = 9.0/8 = 1.125` (Pearl-A direct-replace).
+  - NEW `td_error_mag_ema_blend_post_bootstrap` (`#[ignore]`) — synthetic `td_errors=[+0.5, -0.5, +0.5, -0.5]`; pre-populate slot with non-sentinel `1.0`; assert blend equals `(1 - 0.4) × 1.0 + 0.4 × 0.5 = 0.8`.
+
+**Atomicity envelope** (per `feedback_no_partial_refactor`): kernel + cubin manifest + Rust launcher + V_SHARE ring buffer field + HEALTH_DIAG emit + GPU oracle tests all land in a single commit. No intermediate state where the ring buffer drifts without the slope emit consuming it.
+
+**Wire-up summary**:
+
+| Layer | Site | Direction |
+|---|---|---|
+| Kernel | `td_error_mag_ema_kernel.cu::td_error_mag_ema_update` | producer (single-block tree-reduce, no atomicAdd) |
+| Cubin | `build.rs` manifest entry | infra |
+| Cubin re-export | `gpu_dqn_trainer.rs::TD_ERROR_MAG_EMA_CUBIN` | infra |
+| Kernel handle | `gpu_dqn_trainer.rs::sp18_td_error_mag_ema_kernel` | launcher |
+| Launcher | `gpu_dqn_trainer.rs::launch_sp18_td_error_mag_ema_update` | producer call |
+| Reader | `gpu_dqn_trainer.rs::read_sp18_td_error_mag_ema` | launch+sync+read |
+| Ring buffer | `DQNTrainer::sp18_v_share_history [[f32; 4]; 5]` | host-side V_SHARE history |
+| HEALTH_DIAG emit | `training_loop.rs` per-epoch boundary, after `reward_decomp` line | observability (2 lines: `v_share_traj` + `td_error_pre`) |
+| ISV write (slope) | `ISV[V_SHARE_TREND_DIAG_INDEX=496]` (dir-branch slope, when ring filled) | downstream KILL CRITERION reader |
+| GPU oracle tests | `tests/sp18_hold_reward_oracle_tests.rs` (Pearl-A + blend) | regression guards |
+
+**Pearls + invariants:**
+
+- `feedback_no_atomicadd` — block tree-reduce over `[B]`; thread 0 finalises EMA blend.
+- `feedback_no_htod_htoh_only_mapped_pinned` — `td_errors` is a device buffer (`CudaSlice<f32>::raw_ptr`); `isv` is the mapped-pinned ISV bus device pointer.
+- `pearl_first_observation_bootstrap` — sentinel `SENTINEL_TD_ERROR_MAG_EMA=0.0`; cold-start REPLACE on exact-match (within 1e-6 of sentinel).
+- `pearl_wiener_alpha_floor_for_nonstationary` — α=0.4 (`WELFORD_ALPHA_MIN`). Phase 0 does NOT yet use the Welford-derived α from the TDB_* accumulators in slots [498..504) — those are reserved for the Phase 4 q_next_target Wiener-α blending.
+- `pearl_no_host_branches_in_captured_graph` — pure single-block kernel; runs at the per-epoch boundary outside `exp_fwd_graph` capture (mirroring the SP17 `read_v_share_per_branch` cadence).
+- `feedback_no_stubs` — full reduction body, no zero-output placeholder.
+
+**Phase 0 deliverables status (post-Task 0.2):**
+
+| Plan task | Status |
+|---|---|
+| 0.1 D-leg reward decomposition kernel + emit | ✅ landed (previous commit) |
+| 0.2 B-leg V_SHARE trajectory + TD-error magnitude emit | ✅ this commit |
+| 0.3 Reviewer L40S 1-epoch dispatch | reviewer-only — awaits push |
+| 0.4 KILL CRITERION evaluation | reviewer-only — awaits 0.3 |
+| 0.5 Pearl candidate draft | next commit |
+
+**Plan**: `docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md` § Phase 0 Task 0.2.
+
 ## 2026-05-08 — SP18 v2 Phase 0 Task 0.1: D-leg per-action reward decomposition diagnostic
 
 Phase 0 observability: kernel + Rust launcher + per-epoch HEALTH_DIAG emit + GPU oracle test, all in the same atomic commit per `feedback_no_partial_refactor`.