diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index dc603be3a..e58f18d36 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -879,6 +879,21 @@ fn main() {
         // `isv[ISV_HOLD_COST_IDX=380]` by `isv[461]`. See
         // docs/dqn-wire-up-audit.md § "SP16 Phase 2" for rationale.
         "hold_cost_scale_update_kernel.cu",
+        // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
+        // decomposition diagnostic. Reads `reward_components_per_sample`
+        // ([n_samples × 6]) + `actions` ([n_samples], factored-action
+        // ints) from the experience collector, bins per direction
+        // (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT — 4 bins), and
+        // emits 5 stats per bin (mean r_micro, mean r_opp_cost, mean
+        // r_popart, mean |reward|, fire_rate) into a 20-float
+        // mapped-pinned `reward_decomp_diag_buf`. Block tree-reduce
+        // (4 blocks × 256 threads, one block per bin); no atomicAdd
+        // per `feedback_no_atomicadd.md`. Pure observability — does
+        // NOT modify any production-path buffer; written values are
+        // consumed exclusively by the per-epoch HEALTH_DIAG line emit.
+        // Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
+        // § Phase 0 Task 0.1.
+        "reward_decomp_diag_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 e9b2a7b52..9e6c4d7f6 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -114,6 +114,15 @@ static ATOMS_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/ato
 static Q_QUANTILE_CUBIN:    &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/q_quantile_kernel.cubin"));
 /// Plan 3 Task 1 C.2: per-component |reward| EMA into ISV[63..69).
 pub(crate) static REWARD_COMPONENT_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_component_ema_kernel.cubin"));
+/// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
+/// decomposition diagnostic. Reads `reward_components_per_sample`
+/// + `actions_out` from the collector, bins per-direction
+/// (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT), emits 5 stats per bin
+/// (mean r_micro / r_opp_cost / r_popart / |reward| / fire_rate) into
+/// a 20-float mapped-pinned `reward_decomp_diag_buf`. Block tree-
+/// 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"));
 /// 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.
@@ -6803,6 +6812,25 @@ pub struct GpuDqnTrainer {
     /// host writes — kernel-owned). Lifetime matches the trainer; reset
     /// 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.1 (2026-05-08): D-leg per-action reward
+    //    decomposition diagnostic buffer ──────────────────────────────────
+    /// 20-float mapped-pinned diagnostic buffer for the SP18 D-leg
+    /// per-action reward decomposition. Layout (row-major, 4×5):
+    ///
+    ///   row 0 = Hold,   row 1 = Long,   row 2 = Short,   row 3 = Flat
+    ///   col 0 = mean(r_micro),     col 1 = mean(r_opp_cost),
+    ///   col 2 = mean(r_popart),    col 3 = mean(|reward|),
+    ///   col 4 = fire_rate
+    ///
+    /// Written by the collector's
+    /// `launch_sp18_reward_decomp_diag(...)` kernel each per-step
+    /// boundary, BEFORE `launch_reward_component_ema_inplace` zeros the
+    /// source buffer. Consumed exclusively by the per-epoch HEALTH_DIAG
+    /// emit at `training_loop.rs`. Pure observability — no production-
+    /// path consumer in this commit. Plan: SP18 Phase 0 Task 0.1.
+    pub(crate) sp18_reward_decomp_diag_buf: super::mapped_pinned::MappedF32Buffer,
+
     // SP14 Layer C Phase C.1 (2026-05-08): α-machinery struct fields deleted
     // atomically with the kernel files (alpha_grad_compute_kernel,
     // gradient_hack_detect_kernel, sp14_scale_wire_col_kernel) per
@@ -9407,6 +9435,48 @@ impl GpuDqnTrainer {
         Ok(self.read_isv_signal_at(ADVANTAGE_CLIP_BOUND_INDEX))
     }
 
+    /// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
+    /// decomposition diagnostic device pointer.
+    ///
+    /// Returns the mapped-pinned device pointer of the 20-float
+    /// `sp18_reward_decomp_diag_buf` so the collector's
+    /// `launch_sp18_reward_decomp_diag` can write into it. The host-side
+    /// HEALTH_DIAG emit reads via `read_sp18_reward_decomp_diag()` after
+    /// stream sync. Pure observability — caller is expected to be the
+    /// per-step training loop wiring point.
+    pub fn sp18_reward_decomp_diag_dev_ptr(&self) -> u64 {
+        self.sp18_reward_decomp_diag_buf.dev_ptr
+    }
+
+    /// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
+    /// decomposition diagnostic readback.
+    ///
+    /// Returns the 20-float diagnostic block as a `[f32; 20]`, row-major
+    /// 4 bins × 5 stats:
+    ///
+    ///     [Hold_micro, Hold_opp, Hold_popart, Hold_abs, Hold_fire,
+    ///      Long_micro, Long_opp, Long_popart, Long_abs, Long_fire,
+    ///      Short_micro, Short_opp, Short_popart, Short_abs, Short_fire,
+    ///      Flat_micro, Flat_opp, Flat_popart, Flat_abs, Flat_fire]
+    ///
+    /// Reads the mapped-pinned host_ptr directly — caller MUST have
+    /// stream-synced after the collector's launch (e.g. via the natural
+    /// per-step boundary that already runs `cuStreamSynchronize` between
+    /// collection and training). No HtoD/DtoH copy.
+    ///
+    /// On cold-start (before any kernel launch) returns the
+    /// constructor-initialised zero block.
+    pub fn read_sp18_reward_decomp_diag(&self) -> [f32; 20] {
+        let v = self.sp18_reward_decomp_diag_buf.read_all();
+        debug_assert_eq!(
+            v.len(), 20,
+            "sp18_reward_decomp_diag_buf size invariant: 4 bins × 5 stats"
+        );
+        let mut out = [0.0_f32; 20];
+        for (i, x) in v.iter().take(20).enumerate() { out[i] = *x; }
+        out
+    }
+
     /// Update per-branch liquid tau modulation — GPU kernel (RK4/Euler adaptive ODE).
     ///
     /// Reads per_branch_q_gaps (pinned device-mapped) and qlstm_context from device.
@@ -20568,6 +20638,27 @@ impl GpuDqnTrainer {
             "sp17_advantage_clip_scratch alloc ({sp17_clip_scratch_len} f32): {e}"
         )))?;
 
+        // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
+        // decomposition diagnostic buffer. 4 bins × 5 stats = 20 floats
+        // mapped-pinned. The collector's
+        // `launch_sp18_reward_decomp_diag` writes per-step; the host-side
+        // HEALTH_DIAG emit reads via the mapped host_ptr after stream
+        // sync. Pure observability — no production-path consumer.
+        const SP18_REWARD_DECOMP_DIAG_LEN: usize = 4 * 5;
+        let sp18_reward_decomp_diag_buf = unsafe {
+            super::mapped_pinned::MappedF32Buffer::new(SP18_REWARD_DECOMP_DIAG_LEN)
+        }
+        .map_err(|e| MLError::ModelError(format!(
+            "sp18_reward_decomp_diag_buf alloc ({SP18_REWARD_DECOMP_DIAG_LEN} f32): {e}"
+        )))?;
+        // Initialise to 0.0 so the first HEALTH_DIAG read on cold-start
+        // (before the collector has fired) is a deterministic zero block
+        // rather than indeterminate page contents. Pearl-A first-
+        // observation bootstrap fires inside the kernel itself when
+        // n_samples > 0; this just zero-fills before the kernel ever runs.
+        sp18_reward_decomp_diag_buf
+            .write_from_slice(&vec![0.0_f32; SP18_REWARD_DECOMP_DIAG_LEN]);
+
         // SP14 Layer C Phase C.1 (2026-05-08): α-machinery cubin loads
         // (alpha_grad_compute, gradient_hack_detect, sp14_scale_wire_col)
         // deleted atomically with the kernel files. Coupling A
@@ -24810,6 +24901,11 @@ impl GpuDqnTrainer {
             sp17_v_share_kernel,
             sp17_advantage_clip_bound_kernel,
             sp17_advantage_clip_scratch,
+            // 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,
+            // read host-side at HEALTH_DIAG cadence.
+            sp18_reward_decomp_diag_buf,
             sp14_dir_concat_qaux_kernel,
             sp14_dir_qaux_concat_scratch,
             // SP14 backward dX scratch — column-SH2 wire is dropped at
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index ca33ea3c0..22e989331 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -888,6 +888,19 @@ pub struct GpuExperienceCollector {
     /// Launched after each collect_experiences_gpu epoch from launch_reward_component_ema_inplace().
     reward_component_ema_kernel: CudaFunction,
 
+    /// SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action reward
+    /// decomposition diagnostic kernel. Reads
+    /// `reward_components_per_sample [n_samples × 6]` + `actions_out
+    /// [n_samples]` (factored-action ints) and bins per-direction
+    /// (DIR_HOLD / DIR_LONG / DIR_SHORT / DIR_FLAT). Writes a 20-float
+    /// per-action-bin diagnostic block (4 bins × 5 stats per bin) into
+    /// a host-supplied mapped-pinned `out_buf`. Block tree-reduce —
+    /// 4 blocks × 256 threads. Launched BEFORE
+    /// `launch_reward_component_ema_inplace` (which zeroes the source
+    /// buffer). Pure observability — no production-path consumer.
+    /// Plan: SP18 Phase 0 Task 0.1.
+    sp18_reward_decomp_diag_kernel: CudaFunction,
+
     /// SP4 GPU-only Pearls A+D applicator (2026-05-01). Loaded from the
     /// shared `APPLY_PEARLS_AD_CUBIN`; consumed by
     /// `launch_reward_component_ema_inplace` after the
@@ -1964,6 +1977,24 @@ impl GpuExperienceCollector {
                 .map_err(|e| MLError::ModelError(format!("reward_component_ema load: {e}")))?
         };
 
+        // SP18 v2 Phase 0 Task 0.1 (2026-05-08): per-action reward
+        // decomposition diagnostic kernel. Loaded on the collector's
+        // stream so the launcher chains directly after
+        // `experience_action_select` (which writes `actions_out` packed
+        // factored action_idx) and BEFORE
+        // `launch_reward_component_ema_inplace` (which zeros the
+        // reward-component buffer).
+        let sp18_reward_decomp_diag_kernel = {
+            use super::gpu_dqn_trainer::REWARD_DECOMP_DIAG_CUBIN;
+            let m = stream.context()
+                .load_cubin(REWARD_DECOMP_DIAG_CUBIN.to_vec())
+                .map_err(|e| MLError::ModelError(format!(
+                    "sp18 reward_decomp_diag cubin load: {e}")))?;
+            m.load_function("reward_decomp_diag")
+                .map_err(|e| MLError::ModelError(format!(
+                    "sp18 reward_decomp_diag load: {e}")))?
+        };
+
         // SP4 GPU-only Pearls A+D applicator (2026-05-01). Loaded
         // separately on this collector's stream so the cross-boundary
         // `launch_reward_component_ema_inplace` can apply Pearls A+D
@@ -2583,6 +2614,7 @@ impl GpuExperienceCollector {
             market_dim_bn: market_dim_cfg,
             bn_tanh_concat_fn,
             reward_component_ema_kernel,
+            sp18_reward_decomp_diag_kernel,
             apply_pearls_ad_kernel,
             // SP13 v3 P0a.T4 (2026-05-04): Hold-rate observer + fixed-α
             // EMA applicator + 1-elem mapped-pinned scratch — chained
@@ -3322,6 +3354,83 @@ impl GpuExperienceCollector {
         Ok([0.0, cf_rate, trail_r, micro_r])
     }
 
+    /// SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
+    /// decomposition diagnostic launcher.
+    ///
+    /// Launches `reward_decomp_diag` (4 blocks × 256 threads, one block
+    /// per direction-axis bin) on this collector's stream. Reads
+    /// `reward_components_per_sample [n_samples × 6]` + `actions_out
+    /// [n_samples]` written by the last `collect_experiences_gpu`
+    /// epoch and bins per-direction (DIR_HOLD / DIR_LONG / DIR_SHORT /
+    /// DIR_FLAT). Writes a 20-float diagnostic block (4 bins × 5 stats
+    /// per bin) into the host-supplied mapped-pinned `out_dev_ptr`.
+    ///
+    /// Output layout (row-major, 4×5):
+    ///
+    ///     row 0 = Hold,   row 1 = Long,   row 2 = Short,   row 3 = Flat
+    ///     col 0 = mean(r_micro)            (component idx 3)
+    ///     col 1 = mean(r_opp_cost)         (component idx 4)
+    ///     col 2 = mean(r_popart)           (component idx 0)
+    ///     col 3 = mean(|reward|)           (sum of |components|)
+    ///     col 4 = fire_rate                (frac bars with non-zero comp)
+    ///
+    /// **Launch ordering** (CRITICAL): MUST run BEFORE
+    /// `launch_reward_component_ema_inplace` because the latter
+    /// `memset_zeros` the source `reward_components_per_sample` after
+    /// consuming it. Per `pearl_canary_input_freshness_launch_order`.
+    ///
+    /// Pure observability — does NOT modify any production-path buffer.
+    /// `n_samples` should be `alloc_episodes * alloc_timesteps` (base
+    /// batch, not CF-doubled). Pass `out_dev_ptr=0` to noop the launch
+    /// (e.g. before the trainer wires the diagnostic buffer).
+    ///
+    /// Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
+    /// § Phase 0 Task 0.1.
+    pub fn launch_sp18_reward_decomp_diag(
+        &mut self,
+        n_samples: usize,
+        b1: usize,
+        b2: usize,
+        b3: usize,
+        out_dev_ptr: u64,
+    ) -> Result<(), crate::MLError> {
+        if n_samples == 0 || out_dev_ptr == 0 {
+            return Ok(());
+        }
+        let n_i32     = n_samples as i32;
+        let b1_i32    = b1 as i32;
+        let b2_i32    = b2 as i32;
+        let b3_i32    = b3 as i32;
+        let rc_ptr    = self.reward_components_per_sample.raw_ptr();
+        let act_ptr   = self.actions_out.raw_ptr();
+        // Block tree-reduce: 4 blocks × 256 threads. Dynamic shmem =
+        // 5 × BLOCK_SIZE × sizeof(f32) for the 5 reduction tiles
+        // (micro / opp_cost / popart / abs_sum / count). Matches the
+        // SP18_DECOMP_BLOCK_SIZE=256 kernel-side constant.
+        const BLOCK_DIM: u32 = 256;
+        const NUM_BINS:  u32 = 4;
+        let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+        unsafe {
+            self.stream.launch_builder(&self.sp18_reward_decomp_diag_kernel)
+                .arg(&rc_ptr)
+                .arg(&act_ptr)
+                .arg(&n_i32)
+                .arg(&b1_i32)
+                .arg(&b2_i32)
+                .arg(&b3_i32)
+                .arg(&out_dev_ptr)
+                .launch(LaunchConfig {
+                    grid_dim:  (NUM_BINS, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .map_err(|e| crate::MLError::ModelError(format!(
+                    "sp18 reward_decomp_diag launch: {e}"
+                )))?;
+        }
+        Ok(())
+    }
+
     /// C.2 Plan 3 Task 1 (spec §4.C.2); SP4 Layer A Task A13.5 retrofit
     /// (2026-05-01); GPU-Pearls refactor (2026-05-01): GPU-driven
     /// reward-component step-observation producer + GPU Pearls A+D update.
diff --git a/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu b/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu
new file mode 100644
index 000000000..8c074d2d4
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu
@@ -0,0 +1,229 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * SP18 v2 Phase 0 Task 0.1 (2026-05-08) — D-leg per-action reward
+ * decomposition diagnostic.
+ *
+ * Reads the per-sample reward-component buffer + actions buffer produced
+ * by `experience_kernels.cu`, bins per direction-axis (Hold / Long /
+ * Short / Flat), and emits 5 per-bin statistics into a 20-float
+ * mapped-pinned diagnostic buffer:
+ *
+ *   col 0: mean(r_micro)        per action bin   (component idx 3)
+ *   col 1: mean(r_opp_cost)     per action bin   (component idx 4)
+ *   col 2: mean(r_popart)       per action bin   (component idx 0)
+ *   col 3: mean(|reward|)       per action bin   (sum of absolute values
+ *                                                  across all 6 components)
+ *   col 4: fire_rate            per action bin   (frac of bars with any
+ *                                                  non-zero component)
+ *
+ * Buffer layout (row-major, 4 rows × 5 cols):
+ *
+ *     out[bin * 5 + col]
+ *
+ * with bin order:
+ *
+ *     bin 0 = Hold   (DIR_HOLD  = 1)
+ *     bin 1 = Long   (DIR_LONG  = 2)
+ *     bin 2 = Short  (DIR_SHORT = 0)
+ *     bin 3 = Flat   (DIR_FLAT  = 3)
+ *
+ * The bin order is logical (Hold first — primary observable for SP18 KILL
+ * CRITERION) and explicitly does NOT match the DIR_* enum order. Each
+ * thread reads its sample's `actions[i]`, decodes the direction via
+ * `action / (b1 * b2 * b3)`, and increments counters for the matching
+ * bin.
+ *
+ * The reward-component layout is the SP15 Wave 2 reshape per
+ * `reward_component_ema_kernel.cu`:
+ *
+ *   [c=0] popart    — final on-policy reward (PopArt input)
+ *   [c=1] cf        — counterfactual replay
+ *   [c=2] trail     — 0.0 placeholder
+ *   [c=3] micro     — dense OFI-alignment micro-reward
+ *   [c=4] opp_cost  — Flat opportunity cost
+ *   [c=5] bonus     — trade-attempt / timing bonus
+ *
+ * Pearls + invariants:
+ *
+ *   - `feedback_no_atomicadd.md` — block tree-reduce; one block per bin,
+ *     5 shmem accumulators each, no atomicAdd.
+ *
+ *   - `pearl_no_host_branches_in_captured_graph.md` — pure kernel launch,
+ *     no host-side dispatch within the capture region (the call site is
+ *     outside the captured graph anyway: per-step diagnostic at the same
+ *     cadence as `launch_reward_component_ema_inplace`).
+ *
+ *   - `feedback_no_stubs.md` — full bin-aware reduction, no zero-output
+ *     placeholder.
+ *
+ *   - `feedback_no_htod_htoh_only_mapped_pinned.md` — `out` is a mapped-
+ *     pinned `MappedF32Buffer.dev_ptr`; host reads after stream sync.
+ *
+ *   - `pearl_first_observation_bootstrap.md` — the kernel writes raw
+ *     means each launch (no Wiener-α blend at this Phase 0 cadence —
+ *     diagnostic only). Host-side EMA blend is a separate concern; the
+ *     Phase 0 KILL CRITERION reads the raw per-epoch means.
+ *
+ * Args:
+ *
+ *   rc        — `reward_components_per_sample[n_samples * 6]` device
+ *               pointer. Row-major; row i is `[popart, cf, trail, micro,
+ *               opp_cost, bonus]`.
+ *
+ *   actions   — `actions_out[n_samples]` device pointer (i32). Each
+ *               int decodes via `dir = action / (b1 * b2 * b3)`.
+ *
+ *   n_samples — total sample count (production: `alloc_episodes ×
+ *               alloc_timesteps`; the CF-doubled buffer half is NOT
+ *               consumed by this kernel — pass the base count).
+ *
+ *   b1, b2, b3 — magnitude / order / urgency branch sizes (used to
+ *                decode direction from the factored-action int).
+ *
+ *   out       — `out[20]` mapped-pinned f32, row-major 4×5 (4 bins × 5
+ *               stats per bin). See top-of-file layout comment.
+ *
+ * Launch: grid=(4, 1, 1), block=(BLOCK_SIZE, 1, 1).
+ * Shared mem: 5 × BLOCK_SIZE × sizeof(float) for the 5 reduction tiles.
+ * ══════════════════════════════════════════════════════════════════════════ */
+
+#include <cuda_runtime.h>
+#include "state_layout.cuh"
+
+#define SP18_DECOMP_BLOCK_SIZE 256
+
+/* Component indices in `reward_components_per_sample` row. */
+#define RCP_POPART_IDX   0
+#define RCP_MICRO_IDX    3
+#define RCP_OPP_COST_IDX 4
+#define RCP_NUM_COMPS    6
+
+/* Output layout: 4 bins × 5 cols = 20 floats. */
+#define SP18_DECOMP_NUM_BINS  4
+#define SP18_DECOMP_NUM_COLS  5
+#define SP18_DECOMP_OUT_SIZE  (SP18_DECOMP_NUM_BINS * SP18_DECOMP_NUM_COLS)
+
+/* Logical bin → DIR_* enum lookup. Order: Hold, Long, Short, Flat. */
+__device__ __forceinline__ int sp18_decomp_bin_to_dir(int bin) {
+    switch (bin) {
+        case 0: return DIR_HOLD;
+        case 1: return DIR_LONG;
+        case 2: return DIR_SHORT;
+        case 3: return DIR_FLAT;
+        default: return -1;   /* unreachable */
+    }
+}
+
+extern "C" __global__
+void reward_decomp_diag(
+    const float* __restrict__ rc,        /* [n_samples × 6] */
+    const int*   __restrict__ actions,   /* [n_samples] */
+    int           n_samples,
+    int           b1,
+    int           b2,
+    int           b3,
+    float*        __restrict__ out       /* [20] mapped-pinned */
+) {
+    extern __shared__ float s_buf[];
+    float* s_micro    = s_buf;                                  /* [BLOCK_SIZE] */
+    float* s_opp      = s_buf + SP18_DECOMP_BLOCK_SIZE;         /* [BLOCK_SIZE] */
+    float* s_popart   = s_buf + 2 * SP18_DECOMP_BLOCK_SIZE;     /* [BLOCK_SIZE] */
+    float* s_abs_sum  = s_buf + 3 * SP18_DECOMP_BLOCK_SIZE;     /* [BLOCK_SIZE] */
+    float* s_count    = s_buf + 4 * SP18_DECOMP_BLOCK_SIZE;     /* [BLOCK_SIZE] */
+
+    const int tid     = threadIdx.x;
+    const int bin     = blockIdx.x;
+    if (bin >= SP18_DECOMP_NUM_BINS) return;
+    const int target_dir = sp18_decomp_bin_to_dir(bin);
+    const int dir_div    = b1 * b2 * b3;
+
+    /* Per-thread accumulators — strided over n_samples. */
+    float l_micro   = 0.0f;
+    float l_opp     = 0.0f;
+    float l_popart  = 0.0f;
+    float l_abs_sum = 0.0f;
+    float l_count   = 0.0f;        /* count of bars matching this bin */
+    float l_fire    = 0.0f;        /* count of bars with non-zero component */
+
+    for (int i = tid; i < n_samples; i += SP18_DECOMP_BLOCK_SIZE) {
+        const int dir_i = (dir_div > 0) ? (actions[i] / dir_div) : -1;
+        if (dir_i != target_dir) continue;
+        const int row = i * RCP_NUM_COMPS;
+        const float r_popart = rc[row + RCP_POPART_IDX];
+        const float r_micro  = rc[row + RCP_MICRO_IDX];
+        const float r_opp    = rc[row + RCP_OPP_COST_IDX];
+
+        /* sum of |components| → mean of |reward| across all 6 components.
+         * Read all 6 explicitly so any future component addition fails
+         * loudly via the row-stride contract instead of silently skipping. */
+        float abs_total = 0.0f;
+        #pragma unroll
+        for (int c = 0; c < RCP_NUM_COMPS; ++c) {
+            abs_total += fabsf(rc[row + c]);
+        }
+
+        l_micro   += r_micro;
+        l_opp     += r_opp;
+        l_popart  += r_popart;
+        l_abs_sum += abs_total;
+        l_count   += 1.0f;
+        if (abs_total > 0.0f) {
+            l_fire += 1.0f;
+        }
+    }
+
+    /* Track count for division. We reuse s_count to reduce both the
+     * matching-bin count AND the firing-bar count with a two-phase pattern:
+     * first phase reduces matching count (denominator for means); second
+     * phase reduces firing count (numerator for fire_rate). */
+    s_micro[tid]   = l_micro;
+    s_opp[tid]     = l_opp;
+    s_popart[tid]  = l_popart;
+    s_abs_sum[tid] = l_abs_sum;
+    s_count[tid]   = l_count;
+    __syncthreads();
+
+    /* Block tree-reduce. */
+    for (int s = SP18_DECOMP_BLOCK_SIZE / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            s_micro[tid]   += s_micro[tid + s];
+            s_opp[tid]     += s_opp[tid + s];
+            s_popart[tid]  += s_popart[tid + s];
+            s_abs_sum[tid] += s_abs_sum[tid + s];
+            s_count[tid]   += s_count[tid + s];
+        }
+        __syncthreads();
+    }
+
+    const float total_micro   = s_micro[0];
+    const float total_opp     = s_opp[0];
+    const float total_popart  = s_popart[0];
+    const float total_abs_sum = s_abs_sum[0];
+    const float total_count   = s_count[0];
+    __syncthreads();
+
+    /* Phase 2: reduce firing count using the same s_count tile. */
+    s_count[tid] = l_fire;
+    __syncthreads();
+    for (int s = SP18_DECOMP_BLOCK_SIZE / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            s_count[tid] += s_count[tid + s];
+        }
+        __syncthreads();
+    }
+    const float total_fire = s_count[0];
+
+    if (tid == 0) {
+        /* Means default to 0.0 when no samples landed in the bin (rather
+         * than NaN — the consumer reads these as means and would mis-
+         * interpret NaN as a kernel failure). fire_rate is bounded
+         * [0, 1] by construction. */
+        const float denom = (total_count > 0.0f) ? total_count : 1.0f;
+        const int   o     = bin * SP18_DECOMP_NUM_COLS;
+        out[o + 0] = (total_count > 0.0f) ? (total_micro   / denom) : 0.0f;
+        out[o + 1] = (total_count > 0.0f) ? (total_opp     / denom) : 0.0f;
+        out[o + 2] = (total_count > 0.0f) ? (total_popart  / denom) : 0.0f;
+        out[o + 3] = (total_count > 0.0f) ? (total_abs_sum / denom) : 0.0f;
+        out[o + 4] = (total_count > 0.0f) ? (total_fire    / denom) : 0.0f;
+        __threadfence_system();   /* PCIe-visible to mapped-pinned host_ptr */
+    }
+}
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 916469a62..33b61799d 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -4073,6 +4073,44 @@ impl DQNTrainer {
                                     .as_ref()
                                     .map(|c| c.alloc_episodes() * c.alloc_timesteps())
                                     .unwrap_or(0);
+                                // SP18 v2 Phase 0 Task 0.1 (2026-05-08):
+                                // D-leg per-action reward decomposition
+                                // diagnostic. Runs BEFORE
+                                // `launch_reward_component_ema_inplace`
+                                // (which zeros the source `reward_
+                                // components_per_sample` buffer after
+                                // consuming it). Per
+                                // `pearl_canary_input_freshness_launch_order`.
+                                // Pure observability — writes a 20-float
+                                // mapped-pinned buffer on `fused.trainer
+                                // ()` consumed only by the per-epoch
+                                // HEALTH_DIAG emit; no production-path
+                                // consumer in this commit. ISV/dev_ptr
+                                // is 0 unless the trainer's diag buffer
+                                // was constructed (always true post-PP.2;
+                                // the launcher's `out_dev_ptr=0` guard
+                                // noops if the wire-up regresses).
+                                if let Some(ref fused) = self.fused_ctx {
+                                    if let Some(ref mut c) =
+                                        self.gpu_experience_collector
+                                    {
+                                        let cfg = fused.trainer().config();
+                                        let b1 = cfg.branch_1_size;
+                                        let b2 = cfg.branch_2_size;
+                                        let b3 = cfg.branch_3_size;
+                                        let out_ptr = fused.trainer()
+                                            .sp18_reward_decomp_diag_dev_ptr();
+                                        if let Err(e) = c.launch_sp18_reward_decomp_diag(
+                                            n_main, b1, b2, b3, out_ptr,
+                                        ) {
+                                            tracing::warn!(
+                                                "SP18 Phase 0 reward_decomp_diag \
+                                                 launch failed (non-fatal, \
+                                                 diagnostic-only): {e}"
+                                            );
+                                        }
+                                    }
+                                }
                                 if let Some(ref mut c) = self.gpu_experience_collector {
                                     if let Err(e) =
                                         c.launch_reward_component_ema_inplace(n_main)
@@ -5343,6 +5381,51 @@ impl DQNTrainer {
                 );
             }
 
+            // SP18 v2 Phase 0 Task 0.1 (2026-05-08): D-leg per-action
+            // reward decomposition diagnostic emit. Reads the 20-float
+            // mapped-pinned `sp18_reward_decomp_diag_buf` (4 bins × 5
+            // stats: micro / opp / popart / abs / fire). Bin order:
+            // Hold / Long / Short / Flat. The buffer is written by the
+            // collector kernel `launch_sp18_reward_decomp_diag` (per-step
+            // boundary, BEFORE `launch_reward_component_ema_inplace`
+            // zeroes the source buffer) and is host-readable directly via
+            // its mapped-pinned host_ptr — no DtoH copy. The KILL
+            // CRITERION at Phase 0 Task 0.4 reads:
+            //
+            //   ABORT D-leg if |micro_hold| > 2 × |micro_long| AND Hold%
+            //                  trajectory is FALLING.
+            //   PROCEED if    |micro_hold| < 0.5 × max(|popart_long|,
+            //                                          |popart_short|)
+            //                  AND Hold% trajectory is RISING.
+            //
+            // Host-side cold-start (pre-first-launch) the buffer contains
+            // the constructor-zeroed block; HEALTH_DIAG just prints
+            // zeroes — distinguishable from a real launch via the
+            // associated `episodes_collected` / step counter that other
+            // diags surface this epoch.
+            {
+                let decomp = if let Some(ref fused) = self.fused_ctx {
+                    fused.trainer().read_sp18_reward_decomp_diag()
+                } else {
+                    [0.0_f32; 20]
+                };
+                // Layout: bin*5 + col, with bin order Hold(0), Long(1),
+                // Short(2), Flat(3) and col order
+                // micro(0), opp(1), popart(2), abs(3), fire(4).
+                tracing::info!(
+                    "HEALTH_DIAG[{}]: reward_decomp \
+                     [hold(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
+                      long(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
+                      short(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4}) \
+                      flat(micro={:.4} opp={:.4} popart={:.4} abs={:.4} fire={:.4})]",
+                    epoch,
+                    decomp[0],  decomp[1],  decomp[2],  decomp[3],  decomp[4],
+                    decomp[5],  decomp[6],  decomp[7],  decomp[8],  decomp[9],
+                    decomp[10], decomp[11], decomp[12], decomp[13], decomp[14],
+                    decomp[15], decomp[16], decomp[17], decomp[18], decomp[19],
+                );
+            }
+
             // 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
new file mode 100644
index 000000000..67d5c68ea
--- /dev/null
+++ b/crates/ml/tests/sp18_hold_reward_oracle_tests.rs
@@ -0,0 +1,342 @@
+//! SP18 v2 Phase 0 — D-leg + B-leg diagnostic-emit oracle tests.
+//!
+//! Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
+//!
+//! Phase 0 is observability-only: kernels write to new ISV slots
+//! [493..505) and to a 20-float mapped-pinned diagnostic buffer; no
+//! production-path consumer is wired. Tests therefore lock the
+//! per-action reward decomposition kernel's bit-exact output against a
+//! synthetic CPU oracle and confirm the host-side V_SHARE-trend +
+//! TD-error-magnitude diagnostics emit the expected HEALTH_DIAG line
+//! formats.
+//!
+//! Run on a GPU host:
+//!
+//!     SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
+//!       cargo test -p ml --test sp18_hold_reward_oracle_tests \
+//!         --features cuda -- --ignored --nocapture
+//!
+//! All B-tests (`#[ignore = "requires GPU"]`) load cubins via the
+//! workspace `OUT_DIR` and use `MappedF32Buffer` per
+//! `feedback_no_htod_htoh_only_mapped_pinned`.
+
+// ── Layer A (CPU oracle): per-action reward decomposition math ───────────
+
+/// CPU oracle for the SP18 D-leg per-action reward decomposition.
+///
+/// Construction (4 samples, 4 directions):
+///
+///   sample 0: dir=Hold,  components=[popart=0.10, cf=0.0, trail=0.0,
+///                                    micro=0.20, opp=0.05, bonus=0.0]
+///   sample 1: dir=Long,  components=[popart=0.50, cf=0.0, trail=0.0,
+///                                    micro=0.30, opp=0.0,  bonus=0.0]
+///   sample 2: dir=Short, components=[popart=-0.40, cf=0.0, trail=0.0,
+///                                    micro=-0.20, opp=0.0,  bonus=0.0]
+///   sample 3: dir=Flat,  components=[popart=0.0,  cf=0.0, trail=0.0,
+///                                    micro=0.0,   opp=0.10, bonus=0.0]
+///
+/// Expected per-bin output (bin order Hold, Long, Short, Flat):
+///
+///   Hold:  micro=0.20, opp=0.05, popart= 0.10, |abs|=0.35, fire=1.0
+///   Long:  micro=0.30, opp=0.0,  popart= 0.50, |abs|=0.80, fire=1.0
+///   Short: micro=-0.20, opp=0.0, popart=-0.40, |abs|=0.60, fire=1.0
+///   Flat:  micro=0.0,  opp=0.10, popart= 0.0,  |abs|=0.10, fire=1.0
+///
+/// Pins the per-bin reduction math independent of GPU implementation.
+#[test]
+fn reward_decomp_per_action_cpu_oracle() {
+    // Components: [popart, cf, trail, micro, opp_cost, bonus]
+    let rc: Vec<[f32; 6]> = vec![
+        [ 0.10, 0.0, 0.0,  0.20,  0.05, 0.0],   // Hold
+        [ 0.50, 0.0, 0.0,  0.30,  0.0,  0.0],   // Long
+        [-0.40, 0.0, 0.0, -0.20,  0.0,  0.0],   // Short
+        [ 0.0,  0.0, 0.0,  0.0,   0.10, 0.0],   // Flat
+    ];
+    // Logical bin → DIR_*: 0=Hold(1), 1=Long(2), 2=Short(0), 3=Flat(3)
+    const BIN_DIRS: [i32; 4] = [1, 2, 0, 3];
+    const N: usize = 4;
+
+    // Per-sample dir assignment matches the construction above.
+    let dirs: [i32; N] = [1, 2, 0, 3];
+
+    // CPU oracle: row-major 4×5 (bin × col).
+    let mut out = [0.0_f32; 20];
+    for (bin, target_dir) in BIN_DIRS.iter().enumerate() {
+        let mut s_micro  = 0.0f32;
+        let mut s_opp    = 0.0f32;
+        let mut s_popart = 0.0f32;
+        let mut s_abs    = 0.0f32;
+        let mut s_count  = 0.0f32;
+        let mut s_fire   = 0.0f32;
+        for i in 0..N {
+            if dirs[i] != *target_dir { continue; }
+            let row = rc[i];
+            let abs_total: f32 = row.iter().map(|x| x.abs()).sum();
+            s_popart += row[0];
+            s_micro  += row[3];
+            s_opp    += row[4];
+            s_abs    += abs_total;
+            s_count  += 1.0;
+            if abs_total > 0.0 { s_fire += 1.0; }
+        }
+        let denom = if s_count > 0.0 { s_count } else { 1.0 };
+        let o = bin * 5;
+        out[o + 0] = if s_count > 0.0 { s_micro  / denom } else { 0.0 };
+        out[o + 1] = if s_count > 0.0 { s_opp    / denom } else { 0.0 };
+        out[o + 2] = if s_count > 0.0 { s_popart / denom } else { 0.0 };
+        out[o + 3] = if s_count > 0.0 { s_abs    / denom } else { 0.0 };
+        out[o + 4] = if s_count > 0.0 { s_fire   / denom } else { 0.0 };
+    }
+
+    // Hold (bin 0).
+    assert!((out[0]  - 0.20).abs() < 1e-6, "hold micro: got {}", out[0]);
+    assert!((out[1]  - 0.05).abs() < 1e-6, "hold opp:   got {}", out[1]);
+    assert!((out[2]  - 0.10).abs() < 1e-6, "hold popart: got {}", out[2]);
+    assert!((out[3]  - 0.35).abs() < 1e-6, "hold abs:   got {}", out[3]);
+    assert!((out[4]  - 1.00).abs() < 1e-6, "hold fire:  got {}", out[4]);
+    // Long (bin 1).
+    assert!((out[5]  - 0.30).abs() < 1e-6, "long micro: got {}", out[5]);
+    assert!((out[6]  - 0.00).abs() < 1e-6, "long opp:   got {}", out[6]);
+    assert!((out[7]  - 0.50).abs() < 1e-6, "long popart: got {}", out[7]);
+    assert!((out[8]  - 0.80).abs() < 1e-6, "long abs:   got {}", out[8]);
+    assert!((out[9]  - 1.00).abs() < 1e-6, "long fire:  got {}", out[9]);
+    // Short (bin 2).
+    assert!((out[10] - (-0.20)).abs() < 1e-6, "short micro: got {}", out[10]);
+    assert!((out[11] -   0.00 ).abs() < 1e-6, "short opp:   got {}", out[11]);
+    assert!((out[12] - (-0.40)).abs() < 1e-6, "short popart: got {}", out[12]);
+    assert!((out[13] -   0.60 ).abs() < 1e-6, "short abs:   got {}", out[13]);
+    assert!((out[14] -   1.00 ).abs() < 1e-6, "short fire:  got {}", out[14]);
+    // Flat (bin 3).
+    assert!((out[15] - 0.00).abs() < 1e-6, "flat micro: got {}", out[15]);
+    assert!((out[16] - 0.10).abs() < 1e-6, "flat opp:   got {}", out[16]);
+    assert!((out[17] - 0.00).abs() < 1e-6, "flat popart: got {}", out[17]);
+    assert!((out[18] - 0.10).abs() < 1e-6, "flat abs:   got {}", out[18]);
+    assert!((out[19] - 1.00).abs() < 1e-6, "flat fire:  got {}", out[19]);
+}
+
+// ── Layer B (GPU oracle): kernel matches CPU oracle bit-for-bit ──────────
+
+#[cfg(feature = "cuda")]
+#[allow(unsafe_code)]
+mod gpu {
+    use std::sync::Arc;
+
+    use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
+    use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
+
+    /// Cubin handle for the SP18 D-leg per-action reward decomposition
+    /// diagnostic kernel.
+    const SP18_REWARD_DECOMP_DIAG_CUBIN: &[u8] =
+        include_bytes!(concat!(env!("OUT_DIR"), "/reward_decomp_diag_kernel.cubin"));
+
+    fn make_test_stream() -> Arc<CudaStream> {
+        let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
+        ctx.default_stream()
+    }
+
+    fn load_reward_decomp_diag(stream: &Arc<CudaStream>) -> CudaFunction {
+        let module = stream
+            .context()
+            .load_cubin(SP18_REWARD_DECOMP_DIAG_CUBIN.to_vec())
+            .expect("load reward_decomp_diag_kernel cubin");
+        module
+            .load_function("reward_decomp_diag")
+            .expect("load reward_decomp_diag function")
+    }
+
+    /// SP18 v2 Phase 0 Task 0.1 — GPU oracle for the per-action reward
+    /// decomposition diagnostic kernel.
+    ///
+    /// Builds a 4-sample synthetic batch matching the CPU oracle above
+    /// (one sample per direction-axis bin), launches the kernel, and
+    /// asserts the 20-float mapped-pinned output matches the closed-form
+    /// expected values bit-for-bit (within f32 rounding budget).
+    ///
+    /// Branch sizes (b1=b2=b3=3) match production; the factored-action
+    /// encoding is `dir * (b1*b2*b3) + ...` so DIR_HOLD=1 → action=27,
+    /// DIR_LONG=2 → action=54, DIR_SHORT=0 → action=0, DIR_FLAT=3 →
+    /// action=81.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_decomp_per_action_gpu_oracle() {
+        let stream = make_test_stream();
+        let kernel = load_reward_decomp_diag(&stream);
+
+        const N: usize = 4;
+        const NUM_COMPS: usize = 6;
+        const NUM_BINS: usize = 4;
+        const NUM_COLS: usize = 5;
+        const OUT_LEN: usize = NUM_BINS * NUM_COLS;
+
+        // ── Inputs ───────────────────────────────────────────────────
+        // reward_components_per_sample[N × 6], row-major.
+        let rc_host: Vec<f32> = vec![
+            // sample 0 — Hold
+            0.10, 0.0, 0.0,  0.20, 0.05, 0.0,
+            // sample 1 — Long
+            0.50, 0.0, 0.0,  0.30, 0.0,  0.0,
+            // sample 2 — Short
+           -0.40, 0.0, 0.0, -0.20, 0.0,  0.0,
+            // sample 3 — Flat
+            0.0,  0.0, 0.0,  0.0,  0.10, 0.0,
+        ];
+        assert_eq!(rc_host.len(), N * NUM_COMPS);
+
+        let rc_buf = unsafe { MappedF32Buffer::new(N * NUM_COMPS) }
+            .expect("alloc rc buf");
+        rc_buf.write_from_slice(&rc_host);
+
+        // actions_out[N], factored-action ints. dir = action / (b1*b2*b3).
+        const B1: i32 = 3;
+        const B2: i32 = 3;
+        const B3: i32 = 3;
+        const DIR_DIV: i32 = B1 * B2 * B3;   // 27
+        // DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3.
+        let actions_host: [i32; N] = [
+            1 * DIR_DIV,   // Hold
+            2 * DIR_DIV,   // Long
+            0 * DIR_DIV,   // Short
+            3 * DIR_DIV,   // Flat
+        ];
+        let actions_buf = unsafe { MappedI32Buffer::new(N) }
+            .expect("alloc actions buf");
+        actions_buf.write_from_slice(&actions_host);
+
+        // Output buffer — 20 floats, mapped-pinned. Pre-fill with NaN
+        // so any unwritten lane shows up as a test failure.
+        let out_buf = unsafe { MappedF32Buffer::new(OUT_LEN) }
+            .expect("alloc out buf");
+        out_buf.write_from_slice(&vec![f32::NAN; OUT_LEN]);
+
+        // ── Launch ───────────────────────────────────────────────────
+        let n_i32 = N as i32;
+        const BLOCK_DIM: u32 = 256;
+        const NUM_BIN_BLOCKS: u32 = NUM_BINS as u32;
+        // Dynamic shmem = 5 × BLOCK_DIM × sizeof(f32) — matches the
+        // kernel's `s_buf` carve-up (5 reduction tiles).
+        let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&rc_buf.dev_ptr)
+                .arg(&actions_buf.dev_ptr)
+                .arg(&n_i32)
+                .arg(&B1)
+                .arg(&B2)
+                .arg(&B3)
+                .arg(&out_buf.dev_ptr)
+                .launch(LaunchConfig {
+                    grid_dim: (NUM_BIN_BLOCKS, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch reward_decomp_diag");
+        }
+        stream.synchronize().expect("sync after reward_decomp_diag");
+
+        // ── Verify ───────────────────────────────────────────────────
+        let out = out_buf.read_all();
+        for (i, x) in out.iter().enumerate() {
+            assert!(x.is_finite(), "out[{i}] is non-finite ({x}) — \
+                    kernel did not write this lane");
+        }
+
+        // Closed-form expected values (matches the CPU oracle above).
+        let expected: [f32; OUT_LEN] = [
+            // Hold
+             0.20,  0.05,  0.10, 0.35, 1.00,
+            // Long
+             0.30,  0.00,  0.50, 0.80, 1.00,
+            // Short
+            -0.20,  0.00, -0.40, 0.60, 1.00,
+            // Flat
+             0.00,  0.10,  0.00, 0.10, 1.00,
+        ];
+
+        let eps = 1e-6_f32;
+        for i in 0..OUT_LEN {
+            let bin = i / NUM_COLS;
+            let col = i % NUM_COLS;
+            let bin_name = match bin { 0 => "Hold", 1 => "Long", 2 => "Short", 3 => "Flat", _ => "?" };
+            let col_name = match col { 0 => "micro", 1 => "opp", 2 => "popart", 3 => "abs", 4 => "fire", _ => "?" };
+            assert!(
+                (out[i] - expected[i]).abs() < eps,
+                "out[{bin_name}.{col_name}] expected {:.6}, got {:.6} (diff {:.3e})",
+                expected[i], out[i], (out[i] - expected[i]).abs(),
+            );
+        }
+    }
+
+    /// 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
+    /// consumer relies on this contract: a zero block is interpreted as
+    /// "bin saw no traffic this epoch", whereas NaN would cascade
+    /// downstream into the KILL CRITERION arithmetic.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn reward_decomp_empty_bin_emits_zero_not_nan() {
+        let stream = make_test_stream();
+        let kernel = load_reward_decomp_diag(&stream);
+
+        const N: usize = 1;
+        const NUM_BINS: usize = 4;
+        const NUM_COLS: usize = 5;
+        const OUT_LEN: usize = NUM_BINS * NUM_COLS;
+
+        // Single Long-only sample → bins Hold/Short/Flat are all empty.
+        let rc_host: Vec<f32> = vec![0.5, 0.0, 0.0, 0.3, 0.0, 0.0];
+        let rc_buf = unsafe { MappedF32Buffer::new(rc_host.len()) }
+            .expect("alloc rc buf");
+        rc_buf.write_from_slice(&rc_host);
+
+        const B1: i32 = 3; const B2: i32 = 3; const B3: i32 = 3;
+        const DIR_DIV: i32 = B1 * B2 * B3;
+        let actions_host: [i32; N] = [2 * DIR_DIV];   // Long
+        let actions_buf = unsafe { MappedI32Buffer::new(N) }
+            .expect("alloc actions buf");
+        actions_buf.write_from_slice(&actions_host);
+
+        let out_buf = unsafe { MappedF32Buffer::new(OUT_LEN) }
+            .expect("alloc out buf");
+        out_buf.write_from_slice(&vec![f32::NAN; OUT_LEN]);
+
+        let n_i32 = N as i32;
+        const BLOCK_DIM: u32 = 256;
+        let shmem_bytes: u32 = 5 * BLOCK_DIM * std::mem::size_of::<f32>() as u32;
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&rc_buf.dev_ptr)
+                .arg(&actions_buf.dev_ptr)
+                .arg(&n_i32)
+                .arg(&B1)
+                .arg(&B2)
+                .arg(&B3)
+                .arg(&out_buf.dev_ptr)
+                .launch(LaunchConfig {
+                    grid_dim: (NUM_BINS as u32, 1, 1),
+                    block_dim: (BLOCK_DIM, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch reward_decomp_diag");
+        }
+        stream.synchronize().expect("sync");
+
+        let out = out_buf.read_all();
+        // Hold (bin 0), Short (bin 2), Flat (bin 3) — all 5 cols must be 0.0.
+        for &bin in &[0_usize, 2, 3] {
+            for col in 0..NUM_COLS {
+                let v = out[bin * NUM_COLS + col];
+                assert!(
+                    v == 0.0,
+                    "empty-bin guard: bin={bin} col={col} expected 0.0, got {v}"
+                );
+            }
+        }
+        // Long (bin 1) — non-zero (sanity).
+        assert!(out[1 * NUM_COLS + 0] > 0.0, "long micro should be > 0");
+        assert!(out[1 * NUM_COLS + 4] > 0.0, "long fire  should be > 0");
+    }
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 043829d1f..3d7562c54 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,53 @@
 
 **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.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`.
+
+**Files added/modified:**
+
+- NEW `crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu` — block tree-reduce kernel (4 blocks × 256 threads, one block per direction-axis bin) reading `reward_components_per_sample [n_samples × 6]` + `actions_out [n_samples]` and emitting 5 stats per bin (mean r_micro / mean r_opp_cost / mean r_popart / mean |reward| / fire_rate) into a 20-float row-major output. Bin order: Hold(0)→Long(1)→Short(2)→Flat(3); col order: micro→opp→popart→abs→fire. Empty-bin guard emits 0.0 (NOT NaN) per the consumer-side KILL CRITERION arithmetic contract.
+- `crates/ml/build.rs` — cubin manifest entry for `reward_decomp_diag_kernel.cu`.
+- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — `REWARD_DECOMP_DIAG_CUBIN` re-export + 20-float `MappedF32Buffer sp18_reward_decomp_diag_buf` field on `GpuDqnTrainer` + accessor pair (`sp18_reward_decomp_diag_dev_ptr` writer-side, `read_sp18_reward_decomp_diag` reader-side). Buffer is constructor-initialised to 0.0 so cold-start HEALTH_DIAG emits a deterministic zero block.
+- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — `sp18_reward_decomp_diag_kernel` field + cubin load on the collector's stream + `launch_sp18_reward_decomp_diag(n, b1, b2, b3, out_dev_ptr)` launcher. Launches AT THE PER-STEP BOUNDARY, BEFORE `launch_reward_component_ema_inplace` (which `memset_zeros` the source buffer after consuming it) per `pearl_canary_input_freshness_launch_order`.
+- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — wire the launch right before `launch_reward_component_ema_inplace`; emit a new HEALTH_DIAG line at the per-epoch boundary (right after the SP17 dueling line):
+
+  ```
+  HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z abs=W fire=F) long(...) short(...) flat(...)]
+  ```
+
+- NEW `crates/ml/tests/sp18_hold_reward_oracle_tests.rs`:
+  - `reward_decomp_per_action_cpu_oracle` — CPU oracle pinning the per-bin reduction math against a 4-sample synthetic batch (one sample per direction-axis).
+  - `reward_decomp_per_action_gpu_oracle` (`#[ignore = "requires GPU"]`) — GPU oracle: same synthetic batch, asserts kernel output matches the CPU oracle bit-for-bit (1e-6 f32 budget).
+  - `reward_decomp_empty_bin_emits_zero_not_nan` (`#[ignore]`) — empty-bin guard test.
+
+**Atomicity envelope** (per `feedback_no_partial_refactor`): kernel + cubin manifest + Rust launcher + production wire-up + HEALTH_DIAG emit + GPU oracle test all land in a single commit. The per-epoch HEALTH_DIAG line is consumer; the kernel + launcher + buffer are producer; no intermediate state where one half exists without the other.
+
+**Wire-up summary** (per Invariant 7):
+
+| Layer | Site | Direction |
+|---|---|---|
+| Kernel | `reward_decomp_diag_kernel.cu::reward_decomp_diag` | producer (block tree-reduce, no atomicAdd) |
+| Cubin | `build.rs` manifest entry | infra |
+| Cubin re-export | `gpu_dqn_trainer.rs::REWARD_DECOMP_DIAG_CUBIN` | infra |
+| Buffer | `gpu_dqn_trainer.rs::sp18_reward_decomp_diag_buf [20]` mapped-pinned | output sink |
+| Kernel handle | `gpu_experience_collector.rs::sp18_reward_decomp_diag_kernel` | launcher |
+| Launcher | `gpu_experience_collector.rs::launch_sp18_reward_decomp_diag` | producer call |
+| Wire site | `training_loop.rs` per-step `collect_experiences_gpu` boundary, BEFORE `launch_reward_component_ema_inplace` | production wire |
+| HEALTH_DIAG emit | `training_loop.rs` per-epoch boundary, after the SP17 `dueling` line | observability |
+| GPU oracle test | `tests/sp18_hold_reward_oracle_tests.rs` | regression guard |
+
+**Pearls + invariants:**
+
+- `feedback_no_atomicadd` — block tree-reduce (one block per bin, 5 shmem accumulators each).
+- `feedback_no_htod_htoh_only_mapped_pinned` — output is a `MappedF32Buffer.dev_ptr`; host reads via the mapped host_ptr after stream sync.
+- `pearl_no_host_branches_in_captured_graph` — pure kernel launch, no host-side dispatch within the capture region. The wire site is outside `exp_fwd_graph` capture (per-step boundary, alongside `launch_reward_component_ema_inplace`).
+- `pearl_canary_input_freshness_launch_order` — launches BEFORE `launch_reward_component_ema_inplace` so the per-action means are computed against the live `reward_components_per_sample` buffer (the latter zeroes it after consuming).
+- `feedback_no_stubs` — full bin-aware reduction body, no zero-output placeholder; empty-bin guard emits 0.0 explicitly per the test contract.
+
+**Plan**: `docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md` § Phase 0 Task 0.1.
+
 ## 2026-05-08 — SP18 v2 Pre-Phase Task PP.5: H100 collection-time profiling baseline (DEFERRED-TO-PHASE-4 MARKER)
 
 PP.5 is the B-DD6 headroom-check task: profile current `gpu_experience_collector::collect_experience` end-to-end ms/cycle on H100 to confirm the new target-net forward pass (B-leg Phase 4) has the budget for an estimated ~30% slowdown. The plan explicitly marks this as **reviewer-only** ("Reviewer-run profiling … `./scripts/argo-train.sh --model dqn --epochs 1 --label sp18-pp5-profile --profile nsys`"). It does NOT modify production code — it just produces a baseline measurement for Phase 4 to validate against.