diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index a703d8fb1..dc603be3a 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -941,6 +941,38 @@ fn main() {
         // `read_a_var_ema_per_branch()` in gpu_dqn_trainer.rs. See
         // docs/dqn-wire-up-audit.md § "SP17 Phase 3.1" for rationale.
         "sp17_a_var_ema_kernel.cu",
+        // SP17 Phase 3.2 (2026-05-08): per-branch V_share EMA producer.
+        // 4 blocks × 256 threads. Reads `on_v_logits_buf [B, NA]` and
+        // BRANCH-MAJOR `on_b_logits_buf`; computes per-branch
+        // `|E[V]| / (|E[V]| + |E[A_centered, picked]|)` where picked is
+        // argmax_a Σ_z A_raw[i, a, z] (max-Q action — tractable
+        // per-batch without depending on actions_history_buf). Writes
+        // to ISV[V_SHARE_DIR/MAG/ORD/URG_INDEX] (slots 478..482) with
+        // Pearl-A bootstrap (sentinel 0.5 = healthy 50/50 cold-start) +
+        // α=WELFORD_ALPHA_MIN=0.4 EMA blend. Bilateral clamp [0, 1]
+        // before write per `pearl_symmetric_clamp_audit`. Block
+        // tree-reduce (no atomicAdd) per `feedback_no_atomicadd`.
+        // Diagnostic observability — does NOT modify any consumer path.
+        // Consumed by HEALTH_DIAG `dueling [v_share=...]` via
+        // `read_v_share_per_branch()`. See
+        // docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
+        "sp17_v_share_kernel.cu",
+        // SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound
+        // producer. Single block × 256 threads. Reads BRANCH-MAJOR
+        // `on_b_logits_buf`; for each branch d, computes per-atom mean
+        // of A across actions and writes |A_raw - mean[z]| into a
+        // pre-allocated flat scratch buffer of size Σ_d B×b_d×NA;
+        // calls `sp4_histogram_p99` over the full flat scratch (block
+        // tree-reduce + per-warp tile binning, no atomicAdd per
+        // `pearl_fused_per_group_statistics_oracle` and
+        // `feedback_no_atomicadd`); blends `bound = p99 × 1.5` into
+        // ISV[ADVANTAGE_CLIP_BOUND_INDEX=482] with Pearl-A bootstrap
+        // (sentinel 1.0 = no effective clipping cold-start) + α=0.01
+        // slow per-fold cadence + bilateral clamp [0.1, 100.0] per
+        // `pearl_symmetric_clamp_audit`. Observability-only in this
+        // commit — the actual clip wire-up is Phase 5 follow-up. See
+        // docs/dqn-wire-up-audit.md § "SP17 Phase 3.2" for rationale.
+        "sp17_advantage_clip_bound_kernel.cu",
         // SP14 Layer C Phase C.1 (2026-05-08): cubin entries for
         // alpha_grad_compute_kernel.cu, gradient_hack_detect_kernel.cu,
         // and sp14_scale_wire_col_kernel.cu deleted atomically with
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index d024197bf..1e0fb46fe 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -686,6 +686,30 @@ pub(crate) static SP17_V_A_MEANS_DIAG_CUBIN: &[u8] =
 pub(crate) static SP17_A_VAR_EMA_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/sp17_a_var_ema_kernel.cubin"));
 
+/// SP17 Phase 3.2 (2026-05-08): per-branch V_share EMA. 4 blocks × 256
+/// threads. Cold-path producer for HEALTH_DIAG `dueling [v_share=...]`.
+/// V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]|) per branch,
+/// picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic). Pearl-A
+/// bootstrap (sentinel 0.5) + α=WELFORD_ALPHA_MIN=0.4 + bilateral [0, 1]
+/// clamp. Block tree-reduce per `feedback_no_atomicadd`. Loaded from
+/// `sp17_v_share_kernel.cubin`. Consumed by
+/// `read_v_share_per_branch()`.
+pub(crate) static SP17_V_SHARE_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/sp17_v_share_kernel.cubin"));
+
+/// SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound producer.
+/// Single block × 256 threads. Cold-path producer for HEALTH_DIAG
+/// `dueling [clip=K]`. Computes p99(|A_centered|) × 1.5 over the full
+/// flat batch via `sp4_histogram_p99` (no atomicAdd per
+/// `pearl_fused_per_group_statistics_oracle`); EMA-blends with α=0.01
+/// + bilateral clamp [0.1, 100.0] per `pearl_symmetric_clamp_audit`.
+/// Pearl-A bootstrap on sentinel 1.0. Observability-only in Phase 3 —
+/// the actual clip wire-up is Phase 5 follow-up. Loaded from
+/// `sp17_advantage_clip_bound_kernel.cubin`. Consumed by
+/// `read_advantage_clip_bound()`.
+pub(crate) static SP17_ADVANTAGE_CLIP_BOUND_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/sp17_advantage_clip_bound_kernel.cubin"));
+
 // SP14 Layer C Phase C.1 (2026-05-08): α-machinery cubin statics deleted
 // atomically with the alpha_grad_compute / gradient_hack_detect /
 // sp14_scale_wire_col kernel files. SP14_ALPHA_GRAD_CUBIN,
@@ -6733,6 +6757,30 @@ pub struct GpuDqnTrainer {
     /// Loaded from `sp17_a_var_ema_kernel.cubin`. Consumed by
     /// `read_a_var_ema_per_branch()`.
     sp17_a_var_ema_kernel: CudaFunction,
+    // ── SP17 Phase 3.2 (2026-05-08): per-branch V_share producer ──────────
+    /// 4-block × 256-thread kernel computing per-branch
+    /// `|E[V]| / (|E[V]| + |E[A_centered, picked]|)` (picked = argmax E[Q])
+    /// and EMA-blending into ISV[V_SHARE_DIR/MAG/ORD/URG_INDEX] (slots
+    /// 478..482). Pearl-A bootstrap on sentinel 0.5; α=0.4 post-bootstrap.
+    /// Block tree-reduce per `feedback_no_atomicadd`. Loaded from
+    /// `sp17_v_share_kernel.cubin`. Consumed by
+    /// `read_v_share_per_branch()`.
+    sp17_v_share_kernel: CudaFunction,
+    // ── SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound ────────
+    /// Single-block × 256-thread kernel using `sp4_histogram_p99` to
+    /// compute `p99(|A_centered|) × 1.5` over the full flat batch and
+    /// EMA-blend (α=0.01) with bilateral clamp [0.1, 100.0] into
+    /// ISV[ADVANTAGE_CLIP_BOUND_INDEX=482]. Pearl-A bootstrap on sentinel
+    /// 1.0. Observability-only in Phase 3 — clip wire-up is Phase 5.
+    /// Loaded from `sp17_advantage_clip_bound_kernel.cubin`. Consumed by
+    /// `read_advantage_clip_bound()`.
+    sp17_advantage_clip_bound_kernel: CudaFunction,
+    /// SP17 Phase 3.2 scratch buffer for `sp17_advantage_clip_bound`. Sized
+    /// to the maximum possible flat |A_centered| tensor: B × Σ_d b_d × NA.
+    /// Mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned` (no
+    /// 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,
     // 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
@@ -9136,6 +9184,207 @@ impl GpuDqnTrainer {
         ])
     }
 
+    /// SP17 Phase 3.2 (2026-05-08) — per-branch V_share producer launch.
+    ///
+    /// Cold-path launcher for `sp17_v_share_kernel`. Reads
+    /// `on_v_logits_buf` and `on_b_logits_buf`; computes per-branch
+    /// `V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]|)` where
+    /// picked = argmax_a Σ_z A_raw[i, a, z] (max-Q semantic — tractable
+    /// per-batch without depending on `actions_history_buf` which is
+    /// collector-time state stale relative to the cuBLAS forward driving
+    /// the on_*_logits buffers at HEALTH_DIAG cadence).
+    ///
+    /// Pearl-A bootstrap on sentinel `SENTINEL_V_SHARE=0.5` (cold-start
+    /// healthy 50/50 split); α=`WELFORD_ALPHA_MIN=0.4` post-bootstrap.
+    /// Bilateral [0, 1] clamp before ISV write per
+    /// `pearl_symmetric_clamp_audit`. 4 blocks × 256 threads.
+    pub fn launch_v_share_update(&self) -> Result<(), MLError> {
+        use super::sp14_isv_slots::{
+            V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
+            V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
+            SENTINEL_V_SHARE, WELFORD_ALPHA_MIN,
+        };
+        let b = self.config.batch_size;
+        let na = self.config.num_atoms;
+        let b0 = self.config.branch_0_size;
+        let b1 = self.config.branch_1_size;
+        let b2 = self.config.branch_2_size;
+        let b3 = self.config.branch_3_size;
+        if b == 0 || na == 0 || (b0 + b1 + b2 + b3) == 0 {
+            return Ok(());
+        }
+        let v_ptr = self.on_v_logits_buf.raw_ptr();
+        let adv_ptr = self.on_b_logits_buf.raw_ptr();
+        let isv_ptr = self.isv_signals_dev_ptr;
+        let b_i32 = b as i32;
+        let na_i32 = na as i32;
+        let b0_i32 = b0 as i32;
+        let b1_i32 = b1 as i32;
+        let b2_i32 = b2 as i32;
+        let b3_i32 = b3 as i32;
+        let share_idx_dir = V_SHARE_DIR_INDEX as i32;
+        let share_idx_mag = V_SHARE_MAG_INDEX as i32;
+        let share_idx_ord = V_SHARE_ORD_INDEX as i32;
+        let share_idx_urg = V_SHARE_URG_INDEX as i32;
+        let alpha: f32 = WELFORD_ALPHA_MIN;
+        let sentinel: f32 = SENTINEL_V_SHARE;
+        // Dynamic shmem = (BLOCK_SIZE + NA + 1) × sizeof(f32) — tree-reduce
+        // tile + per-atom mean cache + 1 float for E[V] broadcast.
+        let block_dim: u32 = 256;
+        let shmem_bytes: u32 =
+            (block_dim + na as u32 + 1) * std::mem::size_of::<f32>() as u32;
+        unsafe {
+            self.stream
+                .launch_builder(&self.sp17_v_share_kernel)
+                .arg(&v_ptr)
+                .arg(&adv_ptr)
+                .arg(&isv_ptr)
+                .arg(&b_i32)
+                .arg(&na_i32)
+                .arg(&b0_i32)
+                .arg(&b1_i32)
+                .arg(&b2_i32)
+                .arg(&b3_i32)
+                .arg(&share_idx_dir)
+                .arg(&share_idx_mag)
+                .arg(&share_idx_ord)
+                .arg(&share_idx_urg)
+                .arg(&alpha)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (4, 1, 1),
+                    block_dim: (block_dim, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .map_err(|e| MLError::ModelError(format!(
+                    "sp17_v_share_update launch: {e}"
+                )))?;
+        }
+        Ok(())
+    }
+
+    /// SP17 Phase 3.2 (2026-05-08) — per-branch V_share readback.
+    /// Convenience wrapper: launches → syncs → reads ISV slots 478..482.
+    pub fn read_v_share_per_branch(&self) -> Result<[f32; 4], MLError> {
+        use super::sp14_isv_slots::{
+            V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
+            V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
+        };
+        self.launch_v_share_update()?;
+        self.stream
+            .synchronize()
+            .map_err(|e| MLError::ModelError(format!("sp17_v_share sync: {e}")))?;
+        Ok([
+            self.read_isv_signal_at(V_SHARE_DIR_INDEX),
+            self.read_isv_signal_at(V_SHARE_MAG_INDEX),
+            self.read_isv_signal_at(V_SHARE_ORD_INDEX),
+            self.read_isv_signal_at(V_SHARE_URG_INDEX),
+        ])
+    }
+
+    /// SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound launch.
+    ///
+    /// Cold-path launcher for `sp17_advantage_clip_bound_kernel`. Computes
+    /// `p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5` over the
+    /// full flat batch via `sp4_histogram_p99` (no atomicAdd per
+    /// `pearl_fused_per_group_statistics_oracle`); EMA-blends into
+    /// ISV[ADVANTAGE_CLIP_BOUND_INDEX=482] with α=`ADVANTAGE_CLIP_EMA_ALPHA=0.01`
+    /// + bilateral clamp `[ADVANTAGE_CLIP_BOUND_MIN=0.1, _MAX=100.0]`
+    /// per `pearl_symmetric_clamp_audit`. Pearl-A bootstrap on sentinel
+    /// 1.0.
+    ///
+    /// Observability-only in Phase 3 — the actual clipping wire-up is
+    /// Phase 5 follow-up. This launcher is the producer end of the
+    /// chain that feeds the HEALTH_DIAG `dueling [clip=K]` line.
+    ///
+    /// Single block × 256 threads. Dynamic shmem = `(BLOCK / 32) × 256 × 4
+    /// = 8192 bytes` for `sp4_histogram_p99`'s per-warp tile array.
+    pub fn launch_advantage_clip_bound_update(&self) -> Result<(), MLError> {
+        use super::sp14_isv_slots::{
+            ADVANTAGE_CLIP_BOUND_INDEX,
+            ADVANTAGE_CLIP_BOUND_MIN, ADVANTAGE_CLIP_BOUND_MAX,
+            ADVANTAGE_CLIP_EMA_ALPHA, ADVANTAGE_CLIP_SAFETY_FACTOR,
+            SENTINEL_ADVANTAGE_CLIP_BOUND,
+        };
+        let b = self.config.batch_size;
+        let na = self.config.num_atoms;
+        let b0 = self.config.branch_0_size;
+        let b1 = self.config.branch_1_size;
+        let b2 = self.config.branch_2_size;
+        let b3 = self.config.branch_3_size;
+        if b == 0 || na == 0 || (b0 + b1 + b2 + b3) == 0 {
+            return Ok(());
+        }
+        let total = b * (b0 + b1 + b2 + b3) * na;
+        let scratch_capacity = self.sp17_advantage_clip_scratch.len;
+        if total > scratch_capacity {
+            return Err(MLError::ModelError(format!(
+                "sp17_advantage_clip_scratch too small: need {total}, have {scratch_capacity}"
+            )));
+        }
+        let adv_ptr = self.on_b_logits_buf.raw_ptr();
+        let scratch_ptr = self.sp17_advantage_clip_scratch.dev_ptr;
+        let isv_ptr = self.isv_signals_dev_ptr;
+        let scratch_cap_i32 = scratch_capacity as i32;
+        let b_i32 = b as i32;
+        let na_i32 = na as i32;
+        let b0_i32 = b0 as i32;
+        let b1_i32 = b1 as i32;
+        let b2_i32 = b2 as i32;
+        let b3_i32 = b3 as i32;
+        let bound_idx = ADVANTAGE_CLIP_BOUND_INDEX as i32;
+        let alpha: f32 = ADVANTAGE_CLIP_EMA_ALPHA;
+        let safety: f32 = ADVANTAGE_CLIP_SAFETY_FACTOR;
+        let bound_min: f32 = ADVANTAGE_CLIP_BOUND_MIN;
+        let bound_max: f32 = ADVANTAGE_CLIP_BOUND_MAX;
+        let sentinel: f32 = SENTINEL_ADVANTAGE_CLIP_BOUND;
+        // Dynamic shmem for sp4_histogram_p99: (BLOCK / 32) × SP4_HIST_BINS
+        // × sizeof(int) = 8 × 256 × 4 = 8192 bytes for BLOCK_SIZE=256.
+        let block_dim: u32 = 256;
+        let warps = block_dim / 32;
+        let shmem_bytes: u32 = warps * 256 * std::mem::size_of::<i32>() as u32;
+        unsafe {
+            self.stream
+                .launch_builder(&self.sp17_advantage_clip_bound_kernel)
+                .arg(&adv_ptr)
+                .arg(&scratch_ptr)
+                .arg(&scratch_cap_i32)
+                .arg(&isv_ptr)
+                .arg(&b_i32)
+                .arg(&na_i32)
+                .arg(&b0_i32)
+                .arg(&b1_i32)
+                .arg(&b2_i32)
+                .arg(&b3_i32)
+                .arg(&bound_idx)
+                .arg(&alpha)
+                .arg(&safety)
+                .arg(&bound_min)
+                .arg(&bound_max)
+                .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!(
+                    "sp17_advantage_clip_bound_update launch: {e}"
+                )))?;
+        }
+        Ok(())
+    }
+
+    /// SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound readback.
+    /// Convenience wrapper: launches → syncs → reads ISV slot 482.
+    pub fn read_advantage_clip_bound(&self) -> Result<f32, MLError> {
+        use super::sp14_isv_slots::ADVANTAGE_CLIP_BOUND_INDEX;
+        self.launch_advantage_clip_bound_update()?;
+        self.stream
+            .synchronize()
+            .map_err(|e| MLError::ModelError(format!("sp17_advantage_clip_bound sync: {e}")))?;
+        Ok(self.read_isv_signal_at(ADVANTAGE_CLIP_BOUND_INDEX))
+    }
+
     /// 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.
@@ -20255,6 +20504,48 @@ impl GpuDqnTrainer {
             .load_function("sp17_a_var_ema_update")
             .map_err(|e| MLError::ModelError(format!("sp17_a_var_ema_update: {e}")))?;
 
+        // SP17 Phase 3.2 (2026-05-08): per-branch V_share producer.
+        // Cold-path (epoch boundary). 4 blocks × 256 threads; Pearl-A
+        // bootstrap on slots 478..482 (sentinel 0.5 cold-start). Bilateral
+        // [0, 1] clamp before ISV write per `pearl_symmetric_clamp_audit`.
+        let sp17_v_share_module = stream.context()
+            .load_cubin(SP17_V_SHARE_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!("sp17 v_share cubin: {e}")))?;
+        let sp17_v_share_kernel = sp17_v_share_module
+            .load_function("sp17_v_share_update")
+            .map_err(|e| MLError::ModelError(format!("sp17_v_share_update: {e}")))?;
+
+        // SP17 Phase 3.2 (2026-05-08): adaptive advantage_clip_bound
+        // producer + flat |A_centered| scratch buffer. Single block × 256
+        // threads. Pearl-A bootstrap on slot 482 (sentinel 1.0). Uses
+        // `sp4_histogram_p99` (block tree-reduce + per-warp tile binning,
+        // no atomicAdd per `pearl_fused_per_group_statistics_oracle`).
+        // Scratch sized to max possible flat |A_centered| =
+        //   B × (b0 + b1 + b2 + b3) × NA  floats.
+        let sp17_advantage_clip_bound_module = stream.context()
+            .load_cubin(SP17_ADVANTAGE_CLIP_BOUND_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!("sp17 advantage_clip_bound cubin: {e}")))?;
+        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}")))?;
+        let sp17_clip_scratch_capacity = config.batch_size
+            .saturating_mul(
+                config.branch_0_size + config.branch_1_size
+                    + config.branch_2_size + config.branch_3_size
+            )
+            .saturating_mul(config.num_atoms);
+        // Defensive minimum of 1 element so MappedF32Buffer::new() doesn't
+        // fail on a zero-sized config (constructor invariants normally
+        // prevent this; the saturating_mul above could in principle floor
+        // to 0 in degenerate test configs).
+        let sp17_clip_scratch_len = sp17_clip_scratch_capacity.max(1);
+        let sp17_advantage_clip_scratch = unsafe {
+            super::mapped_pinned::MappedF32Buffer::new(sp17_clip_scratch_len)
+        }
+        .map_err(|e| MLError::ModelError(format!(
+            "sp17_advantage_clip_scratch alloc ({sp17_clip_scratch_len} f32): {e}"
+        )))?;
+
         // 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
@@ -24494,6 +24785,9 @@ impl GpuDqnTrainer {
             sp17_v_a_means_diag_kernel,
             sp17_v_a_means_diag_buf,
             sp17_a_var_ema_kernel,
+            sp17_v_share_kernel,
+            sp17_advantage_clip_bound_kernel,
+            sp17_advantage_clip_scratch,
             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/sp17_advantage_clip_bound_kernel.cu b/crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu
new file mode 100644
index 000000000..857fb795d
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu
@@ -0,0 +1,200 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * SP17 Phase 3.2 (2026-05-08) — adaptive advantage_clip_bound producer.
+ *
+ * Computes an adaptive upper-bound for |A_centered| by tracking
+ *   advantage_clip_bound ← p99(|A_centered|) × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5
+ *
+ * Mirrors the SP14 P0-A REWARD_POS_CAP producer pattern (also p99 × 1.5).
+ * The kernel uses `sp4_histogram_p99` (the canonical block-tree-reduce
+ * percentile estimator from SP4) — NO atomicAdd per
+ * `pearl_fused_per_group_statistics_oracle` and `feedback_no_atomicadd`.
+ *
+ * The bound is observability-only in this commit (Phase 3 is the
+ * diagnostic chain); the actual clipping wire-up is Phase 5 follow-up.
+ *
+ * Cadence: cold-path (epoch boundary). EMA α = ADVANTAGE_CLIP_EMA_ALPHA
+ * = 0.01 (slow per-fold — host-passed).
+ *
+ * ── Pearls + invariants ────────────────────────────────────────────────
+ *  - `feedback_no_atomicadd.md` — uses sp4_histogram_p99 (block
+ *    tree-reduce + per-warp tile binning, no atomicAdd).
+ *  - `pearl_first_observation_bootstrap.md` — sentinel
+ *    SENTINEL_ADVANTAGE_CLIP_BOUND = 1.0; cold-start REPLACE on first
+ *    observation.
+ *  - `pearl_symmetric_clamp_audit.md` — bilateral clamp to
+ *    [ADVANTAGE_CLIP_BOUND_MIN=0.1, ADVANTAGE_CLIP_BOUND_MAX=100.0]
+ *    after EMA blend.
+ *  - `feedback_isv_for_adaptive_bounds.md` — bound is producer-tracked,
+ *    NEVER hardcoded.
+ *  - `pearl_no_host_branches_in_captured_graph.md` — pure on-device.
+ *
+ * Algorithm (single-block, BLOCK=256):
+ *
+ *   For each branch d in [0, 4):
+ *     Pass 1: per-atom mean of raw A over (B × n_branch) — written to
+ *             `s_a_mean[NA]` (block-shared, recomputed per branch).
+ *     Pass 2: write |A_raw[i, a, z] - a_mean[z]| into
+ *             `abs_a_centered_scratch` at the BRANCH-MAJOR offset (so
+ *             the final scratch is a flat |A_centered| over the full
+ *             (B × Σ_d n_d × NA) tensor).
+ *
+ *   p99: call `sp4_histogram_p99<256>(scratch, total_count)` →
+ *        p99(|A_centered|) over the full flat tensor.
+ *
+ *   Blend: bound = p99 × SAFETY × α + cur × (1-α), bilateral clamp,
+ *          ISV write.
+ *
+ * Args:
+ *   b_logits                  — BRANCH-MAJOR advantage logits.
+ *   abs_a_centered_scratch    — pre-allocated scratch buffer of size
+ *                                ≥ Σ_d B × b_d × NA floats. Caller-owned;
+ *                                produced/consumed only inside this kernel.
+ *   scratch_capacity           — sentinel (host-passed) for defensive
+ *                                guard; kernel returns early if total >
+ *                                capacity.
+ *   isv                        — ISV bus.
+ *   B, NA                      — shapes.
+ *   b0..b3 sizes               — per-branch action counts.
+ *   bound_idx                  — ADVANTAGE_CLIP_BOUND_INDEX = 482.
+ *   alpha                      — ADVANTAGE_CLIP_EMA_ALPHA = 0.01.
+ *   safety_factor              — ADVANTAGE_CLIP_SAFETY_FACTOR = 1.5.
+ *   bound_min, bound_max       — bilateral clamp rails [0.1, 100.0].
+ *   sentinel_bound             — SENTINEL_ADVANTAGE_CLIP_BOUND = 1.0.
+ *
+ * Launch: grid=(1, 1, 1), block=(256, 1, 1). Dynamic shared memory =
+ * `(BLOCK_SIZE / 32) × 256 × sizeof(int) = 8192 bytes` for
+ * `sp4_histogram_p99`'s per-warp tile array (extern __shared__ int
+ * s_warp_tiles[]). Plus the static shared memory declared inside the
+ * kernel (`s_a_mean[NA_MAX]` and `s_reduce[BLOCK_SIZE]`) — those are
+ * `__shared__` static, NOT extern, so they don't compete with the
+ * dynamic sp4 allocation.
+ * ══════════════════════════════════════════════════════════════════════════ */
+
+#include <cuda_runtime.h>
+#include "sp4_histogram_p99.cuh"
+
+#define SP17_VA_BLOCK_SIZE 256
+#define SP17_NA_MAX        51       /* matches NUM_ATOMS_MAX in state_layout.cuh */
+#define SP17_EPS_F         1e-6f
+
+extern "C" __global__ void sp17_advantage_clip_bound_update(
+    const float* __restrict__ b_logits,
+    float*       __restrict__ abs_a_centered_scratch,
+    int          scratch_capacity,
+    float*       __restrict__ isv,
+    int          B,
+    int          NA,
+    int          b0_size,
+    int          b1_size,
+    int          b2_size,
+    int          b3_size,
+    int          bound_idx,
+    float        alpha,
+    float        safety_factor,
+    float        bound_min,
+    float        bound_max,
+    float        sentinel_bound)
+{
+    /* Single-block kernel — guard against accidental multi-block launch. */
+    if (blockIdx.x != 0) return;
+
+    const int tid  = (int)threadIdx.x;
+    const int bdim = (int)blockDim.x;
+
+    const long long Bll  = (long long)B;
+    const long long NAll = (long long)NA;
+    const long long n0   = (long long)b0_size;
+    const long long n1   = (long long)b1_size;
+    const long long n2   = (long long)b2_size;
+    const long long n3   = (long long)b3_size;
+    const long long off0 = 0;
+    const long long off1 = off0 + Bll * n0 * NAll;
+    const long long off2 = off1 + Bll * n1 * NAll;
+    const long long off3 = off2 + Bll * n2 * NAll;
+    const long long total = off3 + Bll * n3 * NAll;
+
+    if (B <= 0 || NA <= 0 || total <= 0) return;
+    if ((long long)scratch_capacity < total) return;   /* defensive guard */
+
+    /* Static shared memory: per-atom mean cache (small, fixed) +
+     * per-pass tree-reduce tile. These do NOT consume the extern
+     * `__shared__` dynamic allocation that sp4_histogram_p99 owns. */
+    __shared__ float s_a_mean[SP17_NA_MAX];
+    __shared__ float s_reduce[SP17_VA_BLOCK_SIZE];
+
+    /* Branch passes — for each branch d, compute per-atom A mean, then
+     * write |A_raw - mean[z]| into scratch at the branch's offset. After
+     * all 4 branches, the scratch buffer is a flat |A_centered| tensor of
+     * size `total` floats. */
+    for (int branch = 0; branch < 4; ++branch) {
+        int n_branch;
+        long long off_branch;
+        if (branch == 0)      { n_branch = b0_size; off_branch = off0; }
+        else if (branch == 1) { n_branch = b1_size; off_branch = off1; }
+        else if (branch == 2) { n_branch = b2_size; off_branch = off2; }
+        else                  { n_branch = b3_size; off_branch = off3; }
+        if (n_branch <= 0) continue;
+
+        /* Pass 1: per-atom mean over (B × n_branch) at each atom z. */
+        for (int z = 0; z < NA; ++z) {
+            float local = 0.0f;
+            const long long count_ia = Bll * (long long)n_branch;
+            for (long long idx = (long long)tid; idx < count_ia; idx += (long long)bdim) {
+                const long long off = off_branch + idx * NAll + (long long)z;
+                local += b_logits[off];
+            }
+            s_reduce[tid] = local;
+            __syncthreads();
+            for (int s = bdim / 2; s > 0; s >>= 1) {
+                if (tid < s) s_reduce[tid] += s_reduce[tid + s];
+                __syncthreads();
+            }
+            if (tid == 0) {
+                const float total_count = (float)count_ia;
+                s_a_mean[z] = (total_count > 0.0f) ? (s_reduce[0] / total_count) : 0.0f;
+            }
+            __syncthreads();
+        }
+
+        /* Pass 2: write |A_raw - mean[z]| into scratch at BRANCH-MAJOR
+         * offset. Each thread handles a strided slice. */
+        const long long branch_total = Bll * (long long)n_branch * NAll;
+        for (long long idx = (long long)tid; idx < branch_total; idx += (long long)bdim) {
+            const long long ia  = idx / NAll;
+            const long long z   = idx - ia * NAll;
+            const long long off = off_branch + idx;
+            const float a       = b_logits[off];
+            abs_a_centered_scratch[off] = fabsf(a - s_a_mean[z]);
+        }
+        __syncthreads();   /* ensure all writes visible before next branch reuses s_reduce */
+    }
+
+    /* ── p99(|A_centered|) over the full flat scratch ─────────────────
+     * sp4_histogram_p99 uses its OWN extern __shared__ int per-warp
+     * tiles + 256-int static + 1-float static (total 1028 bytes static
+     * + 8192 bytes dynamic for BLOCK=256). Our static `s_a_mean` and
+     * `s_reduce` arrays do not collide because they are statically
+     * declared (different `__shared__` sub-region). The extern dynamic
+     * shmem region is the caller-passed `shared_mem_bytes` and is
+     * exclusively owned by sp4_histogram_p99.
+     */
+    const float p99 = sp4_histogram_p99<SP17_VA_BLOCK_SIZE>(
+        abs_a_centered_scratch, (int)total);
+
+    /* ── Pearl-A bootstrap + α blend + bilateral clamp → ISV write ── */
+    if (tid == 0) {
+        const float target = p99 * safety_factor;
+        const float current = isv[bound_idx];
+        float next;
+        if (fabsf(current - sentinel_bound) < SP17_EPS_F) {
+            /* Pearl-A: cold-start sentinel → REPLACE directly. */
+            next = target;
+        } else {
+            next = (1.0f - alpha) * current + alpha * target;
+        }
+        /* `pearl_symmetric_clamp_audit` — bilateral clamp. */
+        next = fmaxf(bound_min, fminf(next, bound_max));
+        isv[bound_idx] = next;
+        __threadfence_system();
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu b/crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu
new file mode 100644
index 000000000..fa5fa0874
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu
@@ -0,0 +1,220 @@
+/* ══════════════════════════════════════════════════════════════════════════
+ * SP17 Phase 3.2 (2026-05-08) — V_share producer kernel.
+ *
+ * Per-branch magnitude share
+ *   V_share[d] = |E[V]| / (|E[V]| + |E[A_centered, max-Q action]|)
+ *
+ * The picked-action defaults to argmax E[Q] across the b_d actions of the
+ * branch (max-Q semantic — tractable per-batch without depending on the
+ * Thompson-sampled action stored in actions_history_buf, which is
+ * collector-time state and is stale relative to the cuBLAS forward
+ * driving on_v_logits_buf / on_b_logits_buf at HEALTH_DIAG cadence).
+ *
+ * Argmax simplification: argmax_a E[Q[i, a]] = argmax_a (1/NA) Σ_z
+ * (V[i, z] + A_centered[i, a, z]). Since V[i, z] is the same across
+ * actions and the per-atom mean of A is also action-invariant, the
+ * ranking depends only on Σ_z A[i, a, z]. So:
+ *   a*(i) = argmax_a Σ_z A_raw[i, a, z]
+ *
+ * For each sample i in branch d:
+ *   v_i  = (1/NA) Σ_z V[i, z]                            (mean V over atoms)
+ *   a_i  = (1/NA) Σ_z A_centered[i, a*(i), z]            (mean A_centered over atoms)
+ *
+ * Aggregated across the batch (one block per branch):
+ *   |E[V]|_batch    = | (1/B) Σ_i v_i |
+ *   |E[A_c]|_batch  = | (1/B) Σ_i a_i |
+ *
+ * V_share[d] = |E[V]|_batch / (|E[V]|_batch + |E[A_c]|_batch + EPS)
+ *
+ * Then EMA-blend into ISV[V_SHARE_*_INDEX] with the same Pearl-A bootstrap +
+ * Wiener-α-floor pattern as A_var_ema.
+ *
+ * ── Pearls + invariants ────────────────────────────────────────────────
+ *  - `feedback_no_atomicadd.md` — block tree-reduce only.
+ *  - `pearl_first_observation_bootstrap.md` — sentinel SENTINEL_V_SHARE
+ *    = 0.5 ("healthy 50/50 cold-start"); first observation REPLACES.
+ *  - `pearl_wiener_alpha_floor_for_nonstationary` — α floored at 0.4
+ *    (caller passes WELFORD_ALPHA_MIN).
+ *  - `pearl_no_host_branches_in_captured_graph.md` — pure on-device
+ *    branching only (per-branch dispatch via blockIdx.x).
+ *  - `pearl_symmetric_clamp_audit.md` — V_share is structurally bounded
+ *    in [0, 1]; defense-in-depth bilateral clamp before ISV write.
+ *
+ * Launch contract:
+ *   grid_dim  = (4, 1, 1)             — one block per branch
+ *   block_dim = (SP17_VA_BLOCK_SIZE, 1, 1)  — 256 threads
+ *   shared mem (dynamic): (BLOCK_SIZE + NA + 1) × sizeof(float)
+ *                         = tree-reduce tile + per-atom mean cache
+ *                         + 1 float for batch-aggregate E[V] broadcast.
+ *
+ * Args:
+ *   v_logits         — [B, NA] online value logits, sample-major.
+ *   b_logits         — BRANCH-MAJOR online advantage logits.
+ *   isv              — ISV bus.
+ *   B, NA            — batch size, atoms.
+ *   b0..b3 sizes     — per-branch action counts.
+ *   share_idx_*      — V_SHARE_*_INDEX per branch (478..482).
+ *   alpha            — Wiener-α blend rate (caller passes WELFORD_ALPHA_MIN).
+ *   sentinel_v_share — SENTINEL_V_SHARE = 0.5.
+ * ══════════════════════════════════════════════════════════════════════════ */
+
+#include <cuda_runtime.h>
+
+#define SP17_VA_BLOCK_SIZE 256
+#define SP17_EPS_F         1e-6f
+
+extern "C" __global__ void sp17_v_share_update(
+    const float* __restrict__ v_logits,
+    const float* __restrict__ b_logits,
+    float*       __restrict__ isv,
+    int          B,
+    int          NA,
+    int          b0_size,
+    int          b1_size,
+    int          b2_size,
+    int          b3_size,
+    int          share_idx_dir,
+    int          share_idx_mag,
+    int          share_idx_ord,
+    int          share_idx_urg,
+    float        alpha,
+    float        sentinel_v_share)
+{
+    /* One block per branch (gridDim.x == 4). */
+    const int branch = (int)blockIdx.x;
+    if (branch >= 4) return;
+
+    int n_branch, share_idx;
+    long long off_branch;
+    const long long Bll  = (long long)B;
+    const long long NAll = (long long)NA;
+    const long long n0   = (long long)b0_size;
+    const long long n1   = (long long)b1_size;
+    const long long n2   = (long long)b2_size;
+    const long long off0 = 0;
+    const long long off1 = off0 + Bll * n0 * NAll;
+    const long long off2 = off1 + Bll * n1 * NAll;
+    const long long off3 = off2 + Bll * n2 * NAll;
+    if (branch == 0)      { n_branch = b0_size; share_idx = share_idx_dir; off_branch = off0; }
+    else if (branch == 1) { n_branch = b1_size; share_idx = share_idx_mag; off_branch = off1; }
+    else if (branch == 2) { n_branch = b2_size; share_idx = share_idx_ord; off_branch = off2; }
+    else                  { n_branch = b3_size; share_idx = share_idx_urg; off_branch = off3; }
+
+    const int tid  = (int)threadIdx.x;
+    const int bdim = (int)blockDim.x;
+
+    if (B <= 0 || NA <= 0 || n_branch <= 0) return;
+
+    /* Shared layout (dynamic):
+     *   shmem[0..bdim)       — block tree-reduce tile
+     *   shmem[bdim..bdim+NA) — per-atom mean of raw A across actions
+     *   shmem[bdim+NA]       — batch-aggregate E[V] broadcast scalar
+     */
+    extern __shared__ float shmem[];
+    float* s_reduce = shmem;
+    float* s_a_mean = shmem + bdim;
+    float* s_e_v    = shmem + bdim + NA;
+
+    /* ── Pass A: per-atom mean of raw A (same as A_var_ema's pass 1) ── */
+    for (int z = 0; z < NA; ++z) {
+        float local = 0.0f;
+        const long long count_ia = Bll * (long long)n_branch;
+        for (long long idx = (long long)tid; idx < count_ia; idx += (long long)bdim) {
+            const long long off = off_branch + idx * NAll + (long long)z;
+            local += b_logits[off];
+        }
+        s_reduce[tid] = local;
+        __syncthreads();
+        for (int s = bdim / 2; s > 0; s >>= 1) {
+            if (tid < s) s_reduce[tid] += s_reduce[tid + s];
+            __syncthreads();
+        }
+        if (tid == 0) {
+            const float total_count = (float)count_ia;
+            s_a_mean[z] = (total_count > 0.0f) ? (s_reduce[0] / total_count) : 0.0f;
+        }
+        __syncthreads();
+    }
+
+    /* ── Pass B: per-sample v_i = (1/NA) Σ_z V[i, z] aggregated to E[V] ──
+     * v_logits is [B, NA] sample-major (same buffer used by all branches —
+     * the value head is shared across branches). */
+    float local_v_sum = 0.0f;
+    for (long long i = (long long)tid; i < Bll; i += (long long)bdim) {
+        float v_i = 0.0f;
+        const long long base = i * NAll;
+        for (long long z = 0; z < NAll; ++z) {
+            v_i += v_logits[base + z];
+        }
+        v_i /= (float)NA;
+        local_v_sum += v_i;
+    }
+    s_reduce[tid] = local_v_sum;
+    __syncthreads();
+    for (int s = bdim / 2; s > 0; s >>= 1) {
+        if (tid < s) s_reduce[tid] += s_reduce[tid + s];
+        __syncthreads();
+    }
+    if (tid == 0) {
+        s_e_v[0] = s_reduce[0] / (float)B;
+    }
+    __syncthreads();
+    const float e_v = s_e_v[0];
+
+    /* ── Pass C: per-sample picked-action = argmax_a Σ_z A_raw[i, a, z] ──
+     * Build a_i = (1/NA) Σ_z (A_raw[i, a*, z] − a_mean[z]) and aggregate
+     * to E[A_centered, picked]. */
+    float local_a_sum = 0.0f;
+    for (long long i = (long long)tid; i < Bll; i += (long long)bdim) {
+        float best_score   = -INFINITY;
+        long long best_a   = 0;
+        for (long long a = 0; a < (long long)n_branch; ++a) {
+            float s_score = 0.0f;
+            const long long base = off_branch + (i * (long long)n_branch + a) * NAll;
+            for (long long z = 0; z < NAll; ++z) {
+                s_score += b_logits[base + z];
+            }
+            if (s_score > best_score) {
+                best_score = s_score;
+                best_a     = a;
+            }
+        }
+        float a_i_centered = 0.0f;
+        const long long base_pick = off_branch + (i * (long long)n_branch + best_a) * NAll;
+        for (long long z = 0; z < NAll; ++z) {
+            a_i_centered += b_logits[base_pick + z] - s_a_mean[z];
+        }
+        a_i_centered /= (float)NA;
+        local_a_sum += a_i_centered;
+    }
+    s_reduce[tid] = local_a_sum;
+    __syncthreads();
+    for (int s = bdim / 2; s > 0; s >>= 1) {
+        if (tid < s) s_reduce[tid] += s_reduce[tid + s];
+        __syncthreads();
+    }
+
+    /* ── Pass D: V_share + Pearl-A bootstrap + α blend → ISV write ── */
+    if (tid == 0) {
+        const float e_a_centered = s_reduce[0] / (float)B;
+        const float abs_v = fabsf(e_v);
+        const float abs_a = fabsf(e_a_centered);
+        const float denom = abs_v + abs_a + SP17_EPS_F;
+        const float v_share = abs_v / denom;
+
+        const float current = isv[share_idx];
+        float next;
+        if (fabsf(current - sentinel_v_share) < SP17_EPS_F) {
+            /* Pearl-A: cold-start sentinel → REPLACE directly. */
+            next = v_share;
+        } else {
+            next = (1.0f - alpha) * current + alpha * v_share;
+        }
+        /* `pearl_symmetric_clamp_audit` — V_share is structurally
+         * bounded in [0, 1] by the formula, but apply bilateral
+         * defense-in-depth clamp before ISV write. */
+        next = fmaxf(0.0f, fminf(next, 1.0f));
+        isv[share_idx] = next;
+        __threadfence_system();   /* PCIe-visible write for mapped-pinned ISV */
+    }
+}
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index b56edbb1d..562188e9d 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -5273,20 +5273,26 @@ impl DQNTrainer {
                 );
             }
 
-            // SP17 Phase 3.1 (2026-05-08) — per-branch A_var EMA diagnostic.
-            // Cold-path (epoch boundary) producer launch; emits the
-            // dueling-Q identifiability variance signal so the post-deploy
+            // SP17 Phase 3.1+3.2 (2026-05-08) — dueling-Q identifiability
+            // diagnostic chain. Cold-path (epoch boundary) producer
+            // launches that emit the per-branch A_var, per-branch V_share,
+            // and adaptive advantage_clip_bound signals so the post-deploy
             // reviewer can confirm `A_centered` is doing real work
-            // (Var > 0 = the policy discriminates between actions) vs.
-            // regression to the pre-SP17 V-dominated regime (Var → 0 = "all
+            // (Var > 0 with V_share ≈ 0.5 = healthy) vs. regression to the
+            // pre-SP17 V-dominated regime (Var → 0 + V_share → 1.0 = "all
             // actions equivalent, V dominates"). Per `feedback_no_atomicadd`
-            // the kernel uses 4-block tree-reduce; per
-            // `pearl_first_observation_bootstrap` slots 474..478 use
-            // sentinel SENTINEL_A_VAR_EMA=0.0 with REPLACE on first
-            // observation. α post-bootstrap is WELFORD_ALPHA_MIN=0.4 per
-            // `pearl_wiener_alpha_floor_for_nonstationary`. Plan:
+            // the kernels use block tree-reduce (and sp4_histogram_p99 for
+            // the clip bound — per-warp tile binning, no atomicAdd per
+            // `pearl_fused_per_group_statistics_oracle`); per
+            // `pearl_first_observation_bootstrap` all slots use Pearl-A
+            // REPLACE on first observation; α post-bootstrap is
+            // WELFORD_ALPHA_MIN=0.4 (a_var, v_share) and
+            // ADVANTAGE_CLIP_EMA_ALPHA=0.01 (slow per-fold for the clip
+            // bound) per `pearl_wiener_alpha_floor_for_nonstationary`.
+            // Bilateral clamp on V_share [0, 1] and advantage_clip_bound
+            // [0.1, 100.0] per `pearl_symmetric_clamp_audit`. Plan:
             // docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md
-            // Phase 3.1.
+            // Phase 3.1 + 3.2.
             {
                 let a_var = if let Some(ref fused) = self.fused_ctx {
                     match fused.trainer().read_a_var_ema_per_branch() {
@@ -5301,9 +5307,39 @@ impl DQNTrainer {
                 } else {
                     [0.0_f32; 4]
                 };
+                let v_share = if let Some(ref fused) = self.fused_ctx {
+                    match fused.trainer().read_v_share_per_branch() {
+                        Ok(v) => v,
+                        Err(e) => {
+                            tracing::warn!(
+                                "HEALTH_DIAG read_v_share_per_branch refresh failed: {e}"
+                            );
+                            [0.0_f32; 4]
+                        }
+                    }
+                } else {
+                    [0.0_f32; 4]
+                };
+                let clip_bound = if let Some(ref fused) = self.fused_ctx {
+                    match fused.trainer().read_advantage_clip_bound() {
+                        Ok(v) => v,
+                        Err(e) => {
+                            tracing::warn!(
+                                "HEALTH_DIAG read_advantage_clip_bound refresh failed: {e}"
+                            );
+                            0.0_f32
+                        }
+                    }
+                } else {
+                    0.0_f32
+                };
                 tracing::info!(
-                    "HEALTH_DIAG[{}]: dueling [a_var=(d={:.4} m={:.4} o={:.4} u={:.4})]",
-                    epoch, a_var[0], a_var[1], a_var[2], a_var[3],
+                    "HEALTH_DIAG[{}]: dueling [v_share=(d={:.4} m={:.4} o={:.4} u={:.4})] \
+                     [a_var=(d={:.4} m={:.4} o={:.4} u={:.4})] [clip={:.4}]",
+                    epoch,
+                    v_share[0], v_share[1], v_share[2], v_share[3],
+                    a_var[0], a_var[1], a_var[2], a_var[3],
+                    clip_bound,
                 );
             }
 
diff --git a/crates/ml/tests/sp17_dueling_oracle_tests.rs b/crates/ml/tests/sp17_dueling_oracle_tests.rs
index 75320697a..49a4951be 100644
--- a/crates/ml/tests/sp17_dueling_oracle_tests.rs
+++ b/crates/ml/tests/sp17_dueling_oracle_tests.rs
@@ -2296,4 +2296,349 @@ mod gpu {
              (diff {:.3e})", (got_urg - exp_urg).abs(),
         );
     }
+
+    // ── F.2: V_share producer matches closed-form formula ──────────────
+    /// Phase 3.2 (2026-05-08) — synthetic V/A constructed so each branch
+    /// has a closed-form V_share. Launch `sp17_v_share_update`; assert
+    /// ISV[V_SHARE_*] readback matches expected to ε=1e-4.
+    ///
+    /// Construction:
+    ///   - V[i, z] = V_CONST = 2.0 (so |E[V]| = 2.0).
+    ///   - For each branch d, action a, atom z:
+    ///       A[i, a, z] = K_d × (a − (n_d − 1)/2)
+    ///     (mean-zero per atom, picked = max-a = n_d-1).
+    ///   - a_i_centered for picked = K_d × ((n_d-1)/2)
+    ///   - E[A_centered, picked] = K_d × (n_d-1)/2 (constant across i).
+    ///   - V_share[d] = 2.0 / (2.0 + |K_d × (n_d-1)/2|).
+    ///
+    /// With K_d = {1.0, 2.0, 3.0, 0.5} and n_d = {4, 3, 3, 3}:
+    ///   - dir: |a_part| = 1.5 → V_share = 2/(2+1.5) = 4/7 ≈ 0.5714
+    ///   - mag: |a_part| = 2.0 → V_share = 2/(2+2.0) = 0.5
+    ///   - ord: |a_part| = 3.0 → V_share = 2/(2+3.0) = 0.4
+    ///   - urg: |a_part| = 0.5 → V_share = 2/(2+0.5) = 0.8
+    #[test]
+    #[ignore = "requires GPU"]
+    fn v_share_per_branch_matches_closed_form() {
+        use ml::cuda_pipeline::sp14_isv_slots::{
+            V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
+            V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX,
+            SENTINEL_V_SHARE, WELFORD_ALPHA_MIN,
+        };
+
+        const SP17_V_SHARE_CUBIN: &[u8] =
+            include_bytes!(concat!(env!("OUT_DIR"), "/sp17_v_share_kernel.cubin"));
+
+        let stream = make_test_stream();
+        let module = stream
+            .context()
+            .load_cubin(SP17_V_SHARE_CUBIN.to_vec())
+            .expect("load sp17_v_share_kernel cubin");
+        let kernel = module
+            .load_function("sp17_v_share_update")
+            .expect("load sp17_v_share_update function");
+
+        const B: usize = 8;
+        const NA: usize = 51;
+        const N0: usize = 4;
+        const N1: usize = 3;
+        const N2: usize = 3;
+        const N3: usize = 3;
+        const V_CONST: f32 = 2.0;
+        const K_DIR: f32 = 1.0;
+        const K_MAG: f32 = 2.0;
+        const K_ORD: f32 = 3.0;
+        const K_URG: f32 = 0.5;
+
+        let v_logits = vec![V_CONST; B * NA];
+        let n_per_branch = [N0, N1, N2, N3];
+        let k_per_branch = [K_DIR, K_MAG, K_ORD, K_URG];
+        let mut b_logits = Vec::new();
+        for branch in 0..4 {
+            let n = n_per_branch[branch];
+            let k = k_per_branch[branch];
+            let center = (n as f32 - 1.0) / 2.0;
+            for _i in 0..B {
+                for a in 0..n {
+                    for _z in 0..NA {
+                        b_logits.push(k * (a as f32 - center));
+                    }
+                }
+            }
+        }
+
+        let v_buf = unsafe { MappedF32Buffer::new(v_logits.len()) }
+            .expect("alloc v_logits buf");
+        v_buf.write_from_slice(&v_logits);
+        let b_buf = unsafe { MappedF32Buffer::new(b_logits.len()) }
+            .expect("alloc b_logits buf");
+        b_buf.write_from_slice(&b_logits);
+
+        // ISV: pre-populate slots 478..482 with sentinel 0.5 so Pearl-A
+        // bootstrap fires on first launch.
+        const ISV_SPAN: usize = 483;
+        let mut isv_host = vec![0.0_f32; ISV_SPAN];
+        for &slot in [V_SHARE_DIR_INDEX, V_SHARE_MAG_INDEX,
+                      V_SHARE_ORD_INDEX, V_SHARE_URG_INDEX].iter() {
+            isv_host[slot] = SENTINEL_V_SHARE;
+        }
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }
+            .expect("alloc isv buf");
+        isv_buf.write_from_slice(&isv_host);
+
+        let b_i32: i32 = B as i32;
+        let na_i32: i32 = NA as i32;
+        let n0_i32: i32 = N0 as i32;
+        let n1_i32: i32 = N1 as i32;
+        let n2_i32: i32 = N2 as i32;
+        let n3_i32: i32 = N3 as i32;
+        let share_idx_dir: i32 = V_SHARE_DIR_INDEX as i32;
+        let share_idx_mag: i32 = V_SHARE_MAG_INDEX as i32;
+        let share_idx_ord: i32 = V_SHARE_ORD_INDEX as i32;
+        let share_idx_urg: i32 = V_SHARE_URG_INDEX as i32;
+        let alpha: f32 = WELFORD_ALPHA_MIN;
+        let sentinel: f32 = SENTINEL_V_SHARE;
+        let block_dim: u32 = 256;
+        let shmem_bytes: u32 =
+            (block_dim + NA as u32 + 1) * std::mem::size_of::<f32>() as u32;
+
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&v_buf.dev_ptr)
+                .arg(&b_buf.dev_ptr)
+                .arg(&isv_buf.dev_ptr)
+                .arg(&b_i32)
+                .arg(&na_i32)
+                .arg(&n0_i32)
+                .arg(&n1_i32)
+                .arg(&n2_i32)
+                .arg(&n3_i32)
+                .arg(&share_idx_dir)
+                .arg(&share_idx_mag)
+                .arg(&share_idx_ord)
+                .arg(&share_idx_urg)
+                .arg(&alpha)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (4, 1, 1),
+                    block_dim: (block_dim, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch sp17_v_share_update");
+        }
+        stream.synchronize().expect("sync after sp17_v_share_update");
+
+        let isv_after = isv_buf.read_all();
+
+        fn expected_share(n: usize, k: f32, v_const: f32) -> f32 {
+            let a_part = k.abs() * (n as f32 - 1.0) / 2.0;
+            v_const.abs() / (v_const.abs() + a_part + 1e-6_f32)
+        }
+        let exp_dir = expected_share(N0, K_DIR, V_CONST);
+        let exp_mag = expected_share(N1, K_MAG, V_CONST);
+        let exp_ord = expected_share(N2, K_ORD, V_CONST);
+        let exp_urg = expected_share(N3, K_URG, V_CONST);
+
+        let got_dir = isv_after[V_SHARE_DIR_INDEX];
+        let got_mag = isv_after[V_SHARE_MAG_INDEX];
+        let got_ord = isv_after[V_SHARE_ORD_INDEX];
+        let got_urg = isv_after[V_SHARE_URG_INDEX];
+
+        let eps = 1e-4_f32;
+        assert!(
+            (got_dir - exp_dir).abs() < eps,
+            "dir V_share expected {exp_dir:.6}, got {got_dir:.6}",
+        );
+        assert!(
+            (got_mag - exp_mag).abs() < eps,
+            "mag V_share expected {exp_mag:.6}, got {got_mag:.6}",
+        );
+        assert!(
+            (got_ord - exp_ord).abs() < eps,
+            "ord V_share expected {exp_ord:.6}, got {got_ord:.6}",
+        );
+        assert!(
+            (got_urg - exp_urg).abs() < eps,
+            "urg V_share expected {exp_urg:.6}, got {got_urg:.6}",
+        );
+    }
+
+    // ── F.3: advantage_clip_bound producer ≈ p99(|A_centered|) × 1.5 ──
+    /// Phase 3.2 (2026-05-08) — synthetic A constructed so |A_centered|
+    /// has a known distribution, then verify the kernel's p99 estimator
+    /// settles on `≈ p99(|A_centered|) × 1.5` after Pearl-A bootstrap.
+    ///
+    /// Construction:
+    ///   For all 4 branches with shapes (4, 3, 3, 3) × B × NA:
+    ///     A[i, a, z] = a × NOISE_SCALE + 0.001 × (i × NA + z) × per-(i,z) noise
+    ///   The action-component dominates so per-atom mean of A across
+    ///   actions is approximately a_center × NOISE_SCALE; |A_c| ranges
+    ///   continuously through [0, ~max] without identical-value clusters
+    ///   that would break `sp4_histogram_p99`'s per-warp tile binning
+    ///   (non-atomic warp tiles undercount when many lanes hit the same
+    ///   bin in lockstep, which the sp4 comment qualifies as "1/(256×32)
+    ///   loss for uniformly distributed signals" — the qualifier matters).
+    ///
+    /// With NOISE_SCALE = 1.0:
+    ///   - dir n=4, max |A_c| ≈ 1.5 (action 0/3 endpoints), with the
+    ///     per-(i,z) noise spreading the values into adjacent bins.
+    ///   - p99 of |A_c| ≈ 1.5 (top of action-3 cluster) ± noise.
+    ///
+    /// Tolerance: ε=0.15 absolute (the per-(i,z) noise spreads the top
+    /// bin and shifts p99 estimate; we still verify the bound is in the
+    /// right order of magnitude — the test is asserting "kernel produces
+    /// a p99-like statistic, NOT a stuck constant"). Phase 4 L40S smoke
+    /// validates the bound's behavioural utility on real |A_centered|.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn advantage_clip_bound_tracks_p99_safety() {
+        use ml::cuda_pipeline::sp14_isv_slots::{
+            ADVANTAGE_CLIP_BOUND_INDEX,
+            ADVANTAGE_CLIP_BOUND_MIN, ADVANTAGE_CLIP_BOUND_MAX,
+            ADVANTAGE_CLIP_EMA_ALPHA, ADVANTAGE_CLIP_SAFETY_FACTOR,
+            SENTINEL_ADVANTAGE_CLIP_BOUND,
+        };
+
+        const SP17_CLIP_CUBIN: &[u8] = include_bytes!(concat!(
+            env!("OUT_DIR"), "/sp17_advantage_clip_bound_kernel.cubin"));
+
+        let stream = make_test_stream();
+        let module = stream
+            .context()
+            .load_cubin(SP17_CLIP_CUBIN.to_vec())
+            .expect("load sp17_advantage_clip_bound cubin");
+        let kernel = module
+            .load_function("sp17_advantage_clip_bound_update")
+            .expect("load sp17_advantage_clip_bound_update function");
+
+        const B: usize = 8;
+        const NA: usize = 51;
+        const N0: usize = 4;
+        const N1: usize = 3;
+        const N2: usize = 3;
+        const N3: usize = 3;
+        const NOISE_SCALE: f32 = 1.0;
+        const PER_IZ_JITTER: f32 = 0.01;
+
+        // Build A[i, a, z] = a × NOISE_SCALE + jitter[i, z]
+        // The action component dominates so per-atom mean is ≈ a_center; the
+        // per-(i, z) jitter spreads the |A_centered| values continuously and
+        // breaks the lockstep-per-warp pathology that undercounts in
+        // sp4_histogram_p99's non-atomic per-warp tile binning.
+        let n_per_branch = [N0, N1, N2, N3];
+        // Deterministic "noise" so the test is reproducible without an RNG
+        // crate: hash (i × NA + z) into a small float in [-PER_IZ_JITTER, PER_IZ_JITTER].
+        fn jitter(i: usize, z: usize) -> f32 {
+            let h = ((i.wrapping_mul(7919) ^ z.wrapping_mul(31337)) % 2003) as f32;
+            (h / 1001.5_f32 - 1.0_f32) * PER_IZ_JITTER
+        }
+        let mut b_logits = Vec::new();
+        for branch in 0..4 {
+            let n = n_per_branch[branch];
+            for i in 0..B {
+                for a in 0..n {
+                    for z in 0..NA {
+                        b_logits.push((a as f32) * NOISE_SCALE + jitter(i, z));
+                    }
+                }
+            }
+        }
+        let total_len = b_logits.len();
+        assert_eq!(total_len, B * (N0 + N1 + N2 + N3) * NA);
+
+        let b_buf = unsafe { MappedF32Buffer::new(total_len) }
+            .expect("alloc b_logits buf");
+        b_buf.write_from_slice(&b_logits);
+        let scratch_buf = unsafe { MappedF32Buffer::new(total_len) }
+            .expect("alloc scratch buf");
+
+        const ISV_SPAN: usize = 483;
+        let mut isv_host = vec![0.0_f32; ISV_SPAN];
+        isv_host[ADVANTAGE_CLIP_BOUND_INDEX] = SENTINEL_ADVANTAGE_CLIP_BOUND;
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_SPAN) }
+            .expect("alloc isv buf");
+        isv_buf.write_from_slice(&isv_host);
+
+        let scratch_cap_i32: i32 = total_len as i32;
+        let b_i32: i32 = B as i32;
+        let na_i32: i32 = NA as i32;
+        let n0_i32: i32 = N0 as i32;
+        let n1_i32: i32 = N1 as i32;
+        let n2_i32: i32 = N2 as i32;
+        let n3_i32: i32 = N3 as i32;
+        let bound_idx: i32 = ADVANTAGE_CLIP_BOUND_INDEX as i32;
+        let alpha: f32 = ADVANTAGE_CLIP_EMA_ALPHA;
+        let safety: f32 = ADVANTAGE_CLIP_SAFETY_FACTOR;
+        let bound_min: f32 = ADVANTAGE_CLIP_BOUND_MIN;
+        let bound_max: f32 = ADVANTAGE_CLIP_BOUND_MAX;
+        let sentinel: f32 = SENTINEL_ADVANTAGE_CLIP_BOUND;
+        let block_dim: u32 = 256;
+        let warps = block_dim / 32;
+        let shmem_bytes: u32 = warps * 256 * std::mem::size_of::<i32>() as u32;
+
+        unsafe {
+            stream
+                .launch_builder(&kernel)
+                .arg(&b_buf.dev_ptr)
+                .arg(&scratch_buf.dev_ptr)
+                .arg(&scratch_cap_i32)
+                .arg(&isv_buf.dev_ptr)
+                .arg(&b_i32)
+                .arg(&na_i32)
+                .arg(&n0_i32)
+                .arg(&n1_i32)
+                .arg(&n2_i32)
+                .arg(&n3_i32)
+                .arg(&bound_idx)
+                .arg(&alpha)
+                .arg(&safety)
+                .arg(&bound_min)
+                .arg(&bound_max)
+                .arg(&sentinel)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (block_dim, 1, 1),
+                    shared_mem_bytes: shmem_bytes,
+                })
+                .expect("launch sp17_advantage_clip_bound_update");
+        }
+        stream.synchronize().expect("sync after sp17_advantage_clip_bound");
+
+        let isv_after = isv_buf.read_all();
+        let got = isv_after[ADVANTAGE_CLIP_BOUND_INDEX];
+
+        // p99 of |A_centered|: per-action raw values lie in
+        //   {0, 1, 2, 3} × NOISE_SCALE  (action index × scale)
+        // After centering with per-branch action mean, the dir branch's
+        // |A_c| max ≈ 1.5 (action 0/3 endpoints) ± jitter. The per-(i, z)
+        // jitter spreads each value into ~PER_IZ_JITTER × 256 / max bins.
+        //
+        // p99 target ≈ max(|A_c|) × ADVANTAGE_CLIP_SAFETY_FACTOR; Pearl-A
+        // bootstrap REPLACES on first observation, so got ≈ target.
+        // Tolerance is wide because the jitter shifts the upper-tail bin
+        // and the histogram is a linear-bin estimator (not exact p99).
+        let expected_max: f32 = 1.5 + PER_IZ_JITTER;
+        let expected: f32 = expected_max * ADVANTAGE_CLIP_SAFETY_FACTOR;
+        let eps = 0.20_f32;   // wide — jitter + linear histogram quantization
+        assert!(
+            (got - expected).abs() < eps,
+            "advantage_clip_bound expected ≈{expected:.4} (= {expected_max:.3} \
+             × ADVANTAGE_CLIP_SAFETY_FACTOR), got {got:.4} (diff {:.3e}, \
+             eps {eps:.3e})",
+            (got - expected).abs(),
+        );
+        // Also verify the bound is order-of-magnitude correct (NOT stuck
+        // at sentinel and NOT massively over-clamped).
+        assert!(
+            got > 1.0_f32,
+            "advantage_clip_bound stuck near sentinel — got {got:.4}, \
+             expected target ≈{expected:.4} after Pearl-A REPLACE",
+        );
+        // Also sanity-check the bilateral clamp didn't fire.
+        assert!(
+            got >= ADVANTAGE_CLIP_BOUND_MIN && got <= ADVANTAGE_CLIP_BOUND_MAX,
+            "advantage_clip_bound out of bounds: got {got}, must be in \
+             [{ADVANTAGE_CLIP_BOUND_MIN}, {ADVANTAGE_CLIP_BOUND_MAX}]",
+        );
+    }
 }
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 8cb54a268..a925143b6 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,74 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+## 2026-05-08 — SP17 Phase 3.2: V_share + advantage_clip_bound producers (additive observability)
+
+Phase 3.2 lands the remaining two SP17 dueling-Q diagnostic producers atomically with kernel + launcher + Rust wrapper + extended HEALTH_DIAG emit + GPU oracle tests per `feedback_wire_everything_up`.
+
+**New kernels:**
+- `crates/ml/src/cuda_pipeline/sp17_v_share_kernel.cu` — `sp17_v_share_update`
+- `crates/ml/src/cuda_pipeline/sp17_advantage_clip_bound_kernel.cu` — `sp17_advantage_clip_bound_update`
+
+### V_share producer
+
+**Per branch d:**
+1. Pass A — per-atom mean of raw A across n_d actions (same as A_var_ema's pass 1).
+2. Pass B — per-sample `v_i = (1/NA) Σ_z V[i, z]` aggregated to `E[V] = (1/B) Σ_i v_i` (broadcast via shmem).
+3. Pass C — per-sample `a*(i) = argmax_a Σ_z A_raw[i, a, z]` (max-Q semantic — picked-action default; tractable per-batch without depending on `actions_history_buf`); then `a_i = (1/NA) Σ_z (A_raw[i, a*(i), z] − a_mean[z])` aggregated to `E[A_centered, picked]`.
+4. Pass D — `V_share = |E[V]| / (|E[V]| + |E[A_centered, picked]| + EPS)`. Pearl-A bootstrap (sentinel 0.5) + α=`WELFORD_ALPHA_MIN=0.4` blend + bilateral [0, 1] clamp → ISV.
+
+Grid `(4, 1, 1)` × block `(256, 1, 1)`. Dynamic shmem = `(BLOCK_SIZE + NA + 1) × sizeof(f32)` = tree-reduce tile + per-atom mean cache + 1 float for E[V] broadcast.
+
+### advantage_clip_bound producer
+
+**Single block:**
+1. For each branch d (sequential within the block):
+   - Pass 1 — per-atom mean of raw A (cached in static `__shared__ s_a_mean[NA_MAX]`).
+   - Pass 2 — write `|A_raw - mean[z]|` into the kernel-owned `abs_a_centered_scratch` buffer at the BRANCH-MAJOR offset.
+2. Call `sp4_histogram_p99<256>(scratch, total)` per `pearl_fused_per_group_statistics_oracle` — block tree-reduce + per-warp tile binning, NO atomicAdd.
+3. `target = p99 × ADVANTAGE_CLIP_SAFETY_FACTOR=1.5`. Pearl-A bootstrap (sentinel 1.0) + α=`ADVANTAGE_CLIP_EMA_ALPHA=0.01` slow per-fold blend + bilateral clamp `[0.1, 100.0]` per `pearl_symmetric_clamp_audit` → ISV[ADVANTAGE_CLIP_BOUND_INDEX=482].
+
+Grid `(1, 1, 1)` × block `(256, 1, 1)`. Dynamic shmem = `(BLOCK_SIZE/32) × 256 × sizeof(int) = 8192 bytes` for `sp4_histogram_p99`'s per-warp tile array. Static shared memory adds `s_a_mean[51] + s_reduce[256] + sp4's s_bins[256] + sp4's s_step_max` ≈ 2256 bytes — total well within 48KB SM 8.6 limit.
+
+**Scratch buffer:** `MappedF32Buffer` sized to `B × Σ_d b_d × NA` (max possible flat |A_centered|). Allocated once at trainer construct; kernel-owned (no host writes). Per `feedback_no_htod_htoh_only_mapped_pinned`.
+
+### Rust integration (`crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`)
+
+- `SP17_V_SHARE_CUBIN`, `SP17_ADVANTAGE_CLIP_BOUND_CUBIN` static cubins.
+- `sp17_v_share_kernel: CudaFunction`, `sp17_advantage_clip_bound_kernel: CudaFunction` field on `GpuDqnTrainer`.
+- `sp17_advantage_clip_scratch: MappedF32Buffer` field (capacity = max flat |A_c|).
+- `launch_v_share_update()` + `read_v_share_per_branch() -> [f32; 4]` (launches → syncs → reads ISV slots 478..482).
+- `launch_advantage_clip_bound_update()` + `read_advantage_clip_bound() -> f32` (launches → syncs → reads ISV slot 482).
+
+### HEALTH_DIAG emit (canonical line)
+
+```
+HEALTH_DIAG[N]: dueling [v_share=(d=X m=Y o=Z u=W)] [a_var=(d=A m=B o=C u=D)] [clip=K]
+```
+
+Emitted once per epoch alongside the existing `v_a_means` line. Phase 3.3 will polish the format if needed; the format above is the canonical Phase 3 output.
+
+### State reset
+
+ISV slots 478..482 (V_SHARE) and 482 (ADVANTAGE_CLIP_BOUND) are FoldReset entries (already wired in PP.2 commit `a225926e5`). At fold boundary slots are REPLACED with their respective sentinels (`SENTINEL_V_SHARE=0.5`, `SENTINEL_ADVANTAGE_CLIP_BOUND=1.0`) so Pearl-A bootstrap fires fresh on the new fold's first epoch.
+
+### GPU oracle tests
+
+Both in `crates/ml/tests/sp17_dueling_oracle_tests.rs`:
+
+**`v_share_per_branch_matches_closed_form`:** synthetic V[i, z] = 2.0 + linear A per branch with known per-action range. Closed-form `V_share[d] = 2.0 / (2.0 + |K_d × (n_d-1)/2|)`. With `K_d ∈ {1.0, 2.0, 3.0, 0.5}` and `n_d ∈ {4, 3, 3, 3}`:
+- dir: 0.5714, mag: 0.5, ord: 0.4, urg: 0.8
+
+Tolerance ε = 1e-4. Pearl-A bootstrap REPLACES on first launch.
+
+**`advantage_clip_bound_tracks_p99_safety`:** synthetic A with action-component dominant + per-(i, z) jitter (`PER_IZ_JITTER = 0.01`). Critical implementation note: the jitter is REQUIRED — pathologically lockstep values (e.g. all-ones-and-twos with no spread) trigger non-atomic warp-tile collisions in `sp4_histogram_p99`'s per-warp binning, producing severe undercount and a wrong p99. The kernel comment `1/(256×32) loss` qualifies "for uniformly distributed signals" — concentrated values violate that assumption. Real |A_centered| in production is continuous, so the issue is test-data-only, not kernel-correctness. Tolerance ε = 0.20 (jitter spreads top bin + linear histogram quantization).
+
+### Atomicity envelope
+
+Phase 3.2 is observability-only — no consumer path is modified. The clip bound is producer-tracked but NOT yet wired as an actual clip on any kernel — that's Phase 5 follow-up. The Phase 1 mean-zero contract (commits `eabcf8d52` through `6f53d676f`) is what makes `A_centered` a meaningful signal to measure; this commit observes it.
+
+**Plan:** `docs/superpowers/plans/2026-05-08-sp17-dueling-q-network.md` Phase 3.2.
+
 ## 2026-05-08 — SP17 Phase 3.1: A_var_ema per-branch producer kernel (additive observability)
 
 Phase 3 of the SP17 dueling-Q identifiability plan introduces the diagnostic chain that lets the post-deploy reviewer confirm the mean-zero centering (Phase 1-2) is doing real work. This commit lands the first of three producers: per-branch EMA of the over-actions variance of `A_centered`.