diff --git a/crates/ml/src/cuda_pipeline/attention_backward_kernel.cu b/crates/ml/src/cuda_pipeline/attention_backward_kernel.cu
index fd8e8d255..2f30e61d8 100644
--- a/crates/ml/src/cuda_pipeline/attention_backward_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/attention_backward_kernel.cu
@@ -77,10 +77,10 @@ extern "C" __global__ void attn_adam_kernel(
     float weight_decay, float max_grad_norm,
     const int* __restrict__ adam_t_buf,
     int total_params,
-    /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
+    /* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[Attn]].max(EPS_CLAMP_FLOOR)`)
      * — same pattern wired into all five Adam kernels. Used by both
-     * GpuAttention and GpuTlob (TLOB shares this kernel via the same cubin).
-     * 0=disabled (debug only). */
+     * GpuAttention and GpuTlob (TLOB shares this kernel via the same cubin
+     * and co-lives in the Attn group taxonomy). 0=disabled (debug only). */
     float weight_clamp_max_abs,
     /* SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+
      * block tree-reduce (no atomicAdd). GpuAttention + GpuTlob both pass
diff --git a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
index 36549664b..4e84aca2f 100644
--- a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
@@ -321,7 +321,7 @@ extern "C" __global__ void curiosity_adam_step(
     float beta2,
     float eps,
     int step,                         /* 1-based step counter          */
-    /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
+    /* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`)
      * — same pattern wired into all five Adam kernels. 0=disabled (debug only). */
     float weight_clamp_max_abs,
     /* SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+
diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
index 1c3aa1779..5e77032b0 100644
--- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu
@@ -142,7 +142,7 @@ extern "C" __global__ void dqn_adam_update_kernel(
     const float* __restrict__ weight_decay_mask,  /* [TOTAL_PARAMS] 1.0=decay, 0.0=no decay */
     int l1_end,                                    /* param index where L1 stops (end of w_s1) */
     float l1_lambda,                               /* L1 penalty (1e-3). 0=disabled */
-    float weight_clamp_max_abs,                    /* SP3 Mech 9: ISV-driven |p_val| bound (100 * Q_ABS_REF.max(1.0)). 0=disabled. */
+    float weight_clamp_max_abs,                    /* SP4 Layer B: ISV-driven |p_val| bound (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`). 0=disabled. */
     /* SP4 Task A14: Pearl C engagement counter. */
     int* __restrict__ nan_flags_buf,               /* sticky overflow flag store; nan_flags_buf[diag_slot]=1 on first thread that overshoots */
     int diag_slot,                                 /* index into nan_flags_buf (one per Adam kernel) */
@@ -196,13 +196,15 @@ extern "C" __global__ void dqn_adam_update_kernel(
             p_val = sign_val * fmaxf(fabsf(p_val) - lr * l1_lambda, 0.0f);
         }
 
-        /* SP3 Mech 9 + SP4 Task A14: post-Adam weight clamp + Pearl C
-         * engagement counter. ISV-driven bound (100 × Q_ABS_REF.max(1.0))
-         * computed host-side and passed in. Targets cuBLAS sgemm f32
-         * accumulator overflow at slots 26+32 — bounds weights 1 OOM below
-         * slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so the
-         * diagnostic retains regression-sentinel headroom. Disabled when
-         * bound==0 (debug fast-path; production never sets 0).
+        /* SP4 Layer B + SP4 Task A14: post-Adam weight clamp + Pearl C
+         * engagement counter. ISV-driven bound
+         * (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`) computed host-side
+         * via `read_group_adam_bounds` and passed in. Targets cuBLAS
+         * sgemm f32 accumulator overflow at slots 26+32 — Wiener-EMA over
+         * p99(|params|) keeps the bound just above legitimate weight
+         * magnitudes; the diagnostic retains regression-sentinel
+         * headroom. Disabled when bound==0 (debug fast-path; production
+         * never sets 0).
          *
          * `local_engage = 1` when |p_val| would overshoot — captured into
          * the per-block tree-reduce below. The sticky `nan_flags_buf` write
diff --git a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
index 7c8117eee..298d3395b 100644
--- a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
@@ -505,10 +505,11 @@ pub struct GpuCuriosityTrainer {
 
 /// Launch Adam optimizer step for one parameter group.
 ///
-/// `weight_clamp_max_abs` is SP3 Mech 9 (close-out v2): ISV-driven |p|
-/// bound for `curiosity_adam_step`. Caller threads `100 × Q_ABS_REF.max(1.0)`
-/// from training_loop down through `train_on_collector_buffers`. 0=disabled
-/// (debug only).
+/// `weight_clamp_max_abs` is SP4 Layer B: ISV-driven |p| bound for
+/// `curiosity_adam_step`. Caller threads
+/// `ISV[WEIGHT_BOUND[ParamGroup::Curiosity]].max(EPS_CLAMP_FLOOR)`
+/// from training_loop down through `train_on_collector_buffers`.
+/// 0=disabled (debug only).
 #[allow(clippy::too_many_arguments)]
 fn launch_adam_step(
     stream: &CudaStream,
@@ -753,9 +754,10 @@ impl GpuCuriosityTrainer {
         states: &CudaSlice<f32>,
         actions: &CudaSlice<i32>,
         n_samples: usize,
-        /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound for the four
+        /* SP4 Layer B: ISV-driven |p| bound for the four
          * `curiosity_adam_step` launches (W1, b1, W2, b2). Caller computes
-         * `100 × ISV[Q_ABS_REF].max(1.0)` host-side. 0=disabled (debug only). */
+         * `ISV[WEIGHT_BOUND[Curiosity]].max(EPS_CLAMP_FLOOR)` host-side.
+         * 0=disabled (debug only). */
         weight_clamp_max_abs: f32,
     ) -> Result<(), MLError> {
         if n_samples < 2 {
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 6117f91fc..d57ac7ba8 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -15228,6 +15228,36 @@ impl GpuDqnTrainer {
         unsafe { *self.isv_signals_pinned.add(index) }
     }
 
+    /// SP4 Layer B helper: read both Adam bounds
+    /// (`weight_clamp_max_abs`, `weight_decay_value`) for a given param
+    /// group from ISV slots in one call. Replaces the verbose 2-line
+    /// `read_isv_signal_at(weight_bound(...)).max(EPS_CLAMP_FLOOR);
+    ///  read_isv_signal_at(wd_rate(...)).max(0.0)` pattern repeated at
+    /// every per-group Adam launcher (DQN main, IQN, IQL high+low,
+    /// Attention, TLOB).
+    ///
+    /// `weight_clamp_max_abs` enforces the SP1 cold-start ε floor
+    /// (`EPS_CLAMP_FLOOR = 1.0`) on the post-Adam |p| bound. The
+    /// returned value is always ≥ EPS_CLAMP_FLOOR; cold-start ISV reads
+    /// (which can be 0.0 before the producer has run) are saturated to
+    /// the floor so the kernel never silently disables clamping.
+    ///
+    /// `weight_decay_value` enforces the AdamW non-negativity contract
+    /// (`max(0.0)`) — the kernel applies `wd * p` decoupled from Adam
+    /// state, so a negative wd would create a positive feedback loop.
+    /// Returns `0.0` for cold-start.
+    #[inline]
+    pub fn read_group_adam_bounds(
+        &self,
+        group: crate::cuda_pipeline::sp4_isv_slots::ParamGroup,
+    ) -> (f32, f32) {
+        use crate::cuda_pipeline::sp4_isv_slots::{weight_bound, wd_rate};
+        use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
+        let clamp = self.read_isv_signal_at(weight_bound(group.idx())).max(EPS_CLAMP_FLOOR);
+        let wd    = self.read_isv_signal_at(wd_rate(group.idx())).max(0.0);
+        (clamp, wd)
+    }
+
     /// Read the ISV layout fingerprint from pinned memory and compare to the
     /// compile-time `LAYOUT_FINGERPRINT_CURRENT`.
     ///
@@ -21092,30 +21122,50 @@ impl GpuDqnTrainer {
     /// `t_ptr` is a device buffer (not scalar) so the CUDA Graph can be
     /// replayed with an updated step counter.
     ///
-    /// **SP4 Layer B (2026-05-01)**: split into 3 sub-launches per group
-    /// (DqnTrunk / DqnValue / DqnBranches). Each sub-launch reads its own
-    /// `ISV[WEIGHT_BOUND[group]]` clamp + `ISV[WD_RATE[group]]` AdamW rate.
-    /// L1 lambda (`ISV[L1_LAMBDA_TRUNK_INDEX]`) applies to the trunk only;
-    /// value/branches pass `l1_end=0 + l1_lambda=0`. Per-group split
-    /// resolves the Layer A Pearl C engagement-counter accounting deferral
-    /// — each sub-launch writes per-block counts at its own group offset.
+    /// **SP4 Layer B (2026-05-01)**: split into 4 sub-launches over the
+    /// 3-group SP4 taxonomy + a "trunk-extras" tail. The first 3
+    /// (DqnTrunk / DqnValue / DqnBranches) cover tensors `[0..33)` with
+    /// per-group ISV-driven `WEIGHT_BOUND[group]` clamp and `WD_RATE[group]`
+    /// AdamW rate. L1 lambda (`ISV[L1_LAMBDA_TRUNK_INDEX]`) applies to
+    /// the trunk only; value/branches pass `l1_end=0 + l1_lambda=0`.
+    /// Per-group split also resolves the Layer A Pearl C engagement-
+    /// counter accounting deferral — each sub-launch writes per-block
+    /// counts at its own group offset.
+    ///
+    /// **Layer B fix-up (2026-05-01)**: the 4th sub-launch covers the
+    /// tail tensors `[33..163)` (bottleneck, VSN bottleneck, GLU gates,
+    /// KAN, regime gate, spacing, multi-horizon value heads, risk, ISV
+    /// encoder, VSN per-group MLP, aux heads, MoE) — collectively
+    /// "trunk-extras". The original 3-way split silently no-op'd Adam
+    /// for these ~70% of weight tensors (grads computed, m/v allocated,
+    /// no parameter step). The fix-up co-locates trunk-extras under
+    /// `ParamGroup::DqnTrunk` for the WEIGHT_BOUND / WD_RATE contract
+    /// (these tensors share the trunk's gradient regime: pre-bottleneck,
+    /// attention, encoder paths) and uses `SP4_ENGAGE_OFFSET_DISABLED`
+    /// for Pearl C engagement (rolls into trunk's accounting since they
+    /// share bounds). A future task may introduce dedicated SP4
+    /// producers for sub-trunk regions (VSN, MoE, etc.) once warranted
+    /// by signal analysis. Per `feedback_no_partial_refactor`: trunk-
+    /// extras' single (pragmatic) consumer ships with the launch fix in
+    /// this commit.
+    ///
+    /// All parameters `[0..total_params)` are covered across the 4
+    /// sub-launches:
+    ///   - trunk     = `params[0..trunk_param_count]` (tensors [0..13))
+    ///   - value     = byte-slice for tensors [13..17)
+    ///   - branches  = byte-slice for tensors [17..33)
+    ///   - trunk-tail = byte-slice for tensors [33..163)
+    /// A coverage debug_assert at the bottom of the function verifies
+    /// `trunk + value + branches + trunk_extras == total_params` to
+    /// prevent regressions of this kind.
+    ///
     /// Per `feedback_no_quickfixes`: NOT a max/min-of-3 single-launch
     /// shortcut. Per `feedback_no_partial_refactor`: every consumer of
     /// the SP4 contract migrates in this same commit.
-    ///
-    /// All parameters `[0..total_params)` are covered across the 3
-    /// sub-launches (trunk = `params[0..trunk_param_count]`, value =
-    /// indices [13..17), branches = indices [17..33)). VSN tensors at
-    /// indices `[95..119)` fall inside the trunk slice; gradient noise
-    /// from aux paths is attenuated upstream by the in-place
-    /// `VSN_DW_DAMP` scaler in `aux_bottleneck_vsn_backward_dispatch`
-    /// (1B-iv-rc) before it reaches `grad_buf[95..119)` — no Adam-state
-    /// isolation is required.
     fn launch_adam_update(&self) -> Result<(), MLError> {
         use crate::cuda_pipeline::sp4_isv_slots::{
-            weight_bound, wd_rate, ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX,
+            ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX, SP4_ENGAGE_OFFSET_DISABLED,
         };
-        use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
 
         let lr_ptr = self.lr_dev_ptr;  // pinned device-mapped — graph-safe, value read at replay
         let beta1 = self.config.beta1;
@@ -21134,24 +21184,23 @@ impl GpuDqnTrainer {
         // single fused launch. Each group reads its own ISV-driven weight
         // clamp (`ISV[WEIGHT_BOUND[group]]`) and AdamW weight_decay
         // (`ISV[WD_RATE[group]]`); L1 lambda is read from
-        // `ISV[L1_LAMBDA_TRUNK_INDEX]` for the trunk only (groups 1+2 pass 0).
-        // The split also resolves Layer A's Pearl C engagement-counter
-        // limitation (engagement was previously accounted under group 0
-        // only — A14/A15 deferral). Each sub-launch now writes per-block
-        // counts at distinct offsets `group_idx × MAX_BLOCKS_PER_ADAM`,
-        // and `pearl_c_post_adam_engagement_check` reads each group's
-        // dedicated range. Per `feedback_no_quickfixes`: NOT a
-        // max/min-of-3 single-launch shortcut. Per `feedback_no_partial_refactor`:
-        // every consumer of the WEIGHT_BOUND/WD_RATE/L1_LAMBDA_TRUNK
-        // contract migrates in this same commit.
+        // `ISV[L1_LAMBDA_TRUNK_INDEX]` for the trunk only (groups 1+2+
+        // trunk-extras pass 0). The split also resolves Layer A's Pearl C
+        // engagement-counter limitation (engagement was previously
+        // accounted under group 0 only — A14/A15 deferral). Each
+        // SP4-canonical sub-launch writes per-block counts at distinct
+        // offsets `group_idx × MAX_BLOCKS_PER_ADAM`; trunk-extras passes
+        // `SP4_ENGAGE_OFFSET_DISABLED` because it shares trunk's
+        // WEIGHT_BOUND / WD_RATE bounds (Pearl C engagement is a per-
+        // bound diagnostic, not a per-launch one).
         let nan_flags_ptr = self.nan_flags_buf.raw_ptr();
         let diag_slot: i32 = SP4_DQN_ADAM_DIAG_SLOT;
         let engage_buf_ptr: u64 = self.clamp_engage_per_block_buf.dev_ptr;
 
-        // Per-group descriptor: (ParamGroup, byte_offset, count). Mirrors
-        // `param_group_buffers` for the DQN-internal groups exactly (trunk
-        // = full main params, value = bytes [13..17], branches = bytes
-        // [17..33]). We compute byte-offsets here directly to avoid
+        // Per-group descriptor: (ParamGroup, byte_offset, count, engage_offset).
+        // Mirrors `param_group_buffers` for the DQN-internal groups exactly
+        // for groups 0/1/2; the trunk-extras tail extends through end of
+        // params buffer. We compute byte-offsets here directly to avoid
         // threading `SP4AuxBuffers` through this launcher (DQN-internal
         // groups never need aux pointers).
         let f32_sz = std::mem::size_of::<f32>() as u64;
@@ -21163,33 +21212,63 @@ impl GpuDqnTrainer {
         let branches_byte_off = padded_byte_offset(&param_sizes, 17);
         let branches_byte_end = padded_byte_offset(&param_sizes, 33);
         let branches_count = ((branches_byte_end - branches_byte_off) / f32_sz) as usize;
+        // Trunk-extras tail: tensors [33..163). Spans from end of branches
+        // to end of params buffer.
+        let trunk_extras_byte_off = padded_byte_offset(&param_sizes, 33);
+        let trunk_extras_count = self.total_params
+            .saturating_sub((trunk_extras_byte_off / f32_sz) as usize);
 
-        let groups: [(SP4Group, u64, usize); 3] = [
-            (SP4Group::DqnTrunk,    trunk_byte_off,    trunk_count),
-            (SP4Group::DqnValue,    value_byte_off,    value_count),
-            (SP4Group::DqnBranches, branches_byte_off, branches_count),
+        // (group_tag, byte_off, count, engage_offset). The 4th entry is
+        // tagged as DqnTrunk (pragma: shares trunk's WEIGHT_BOUND/WD_RATE
+        // contract) but uses SP4_ENGAGE_OFFSET_DISABLED so its block
+        // engagement does not overwrite trunk's slot in the Pearl C buf.
+        let trunk_engage_off:    i32 = (SP4Group::DqnTrunk.idx()    * MAX_BLOCKS_PER_ADAM) as i32;
+        let value_engage_off:    i32 = (SP4Group::DqnValue.idx()    * MAX_BLOCKS_PER_ADAM) as i32;
+        let branches_engage_off: i32 = (SP4Group::DqnBranches.idx() * MAX_BLOCKS_PER_ADAM) as i32;
+        let groups: [(SP4Group, u64, usize, i32); 4] = [
+            (SP4Group::DqnTrunk,    trunk_byte_off,         trunk_count,         trunk_engage_off),
+            (SP4Group::DqnValue,    value_byte_off,         value_count,         value_engage_off),
+            (SP4Group::DqnBranches, branches_byte_off,      branches_count,      branches_engage_off),
+            // Trunk-extras tail — shares DqnTrunk bounds; Pearl C disabled
+            // for this sub-launch (engagement folds into trunk's
+            // accounting since the same WEIGHT_BOUND drives both clamps).
+            (SP4Group::DqnTrunk,    trunk_extras_byte_off,  trunk_extras_count,  SP4_ENGAGE_OFFSET_DISABLED),
         ];
 
+        // Coverage invariant: every of the 163 weight tensors must be
+        // covered by exactly one of the 4 sub-launches above. Layer B
+        // fix-up: prevents regression of the silent ~70% Adam freeze
+        // that the original 3-way split introduced.
+        debug_assert_eq!(
+            trunk_count + value_count + branches_count + trunk_extras_count,
+            self.total_params,
+            "SP4 Layer B fix-up: 4-way Adam coverage must span all 163 weight tensors \
+             (got trunk={} + value={} + branches={} + trunk_extras={} = {}, expected total_params={})",
+            trunk_count, value_count, branches_count, trunk_extras_count,
+            trunk_count + value_count + branches_count + trunk_extras_count,
+            self.total_params
+        );
+
         // Safety: argument order matches the extern "C" kernel signature exactly.
         // grad_norm_buf contains L2 norm from launch_grad_norm_finalize().
-        for (group, byte_off, count) in groups {
+        for (group, byte_off, count, engage_buf_offset) in groups {
             if count == 0 {
                 continue;
             }
 
             // Per-group ISV-driven post-Adam weight clamp + AdamW
-            // weight_decay rate. EPS_CLAMP_FLOOR=1.0 ε on the bound itself
-            // (SP1 cold-start pearl); `.max(0.0)` on weight_decay enforces
-            // the non-negativity contract on the AdamW rate.
-            let weight_clamp_max_abs: f32 =
-                self.read_isv_signal_at(weight_bound(group.idx())).max(EPS_CLAMP_FLOOR);
-            let weight_decay_value: f32 =
-                self.read_isv_signal_at(wd_rate(group.idx())).max(0.0);
+            // weight_decay rate via the SP4 Layer B helper. Trunk-extras
+            // sub-launch reuses DqnTrunk's bounds (its `group` tag).
+            let (weight_clamp_max_abs, weight_decay_value) =
+                self.read_group_adam_bounds(group);
 
             // L1 proximal step is configured for the trunk only (G8 boundary
-            // at end of `w_s1`). Value/branches pass `l1_end=0 + l1_lambda=0`
-            // which deactivates the in-kernel L1 path.
-            let (l1_end_arg, l1_lambda_arg): (i32, f32) = if group == SP4Group::DqnTrunk {
+            // at end of `w_s1`). Value/branches/trunk-extras pass
+            // `l1_end=0 + l1_lambda=0` which deactivates the in-kernel L1
+            // path. Trunk-extras tensors [33..163) are well past `w_s1`'s
+            // boundary and don't get L1 — only the canonical trunk
+            // sub-launch (byte_off == 0) sets L1.
+            let (l1_end_arg, l1_lambda_arg): (i32, f32) = if byte_off == 0 {
                 let l1_lambda = self.read_isv_signal_at(L1_LAMBDA_TRUNK_INDEX).max(0.0);
                 (trunk_l1_end, l1_lambda)
             } else {
@@ -21209,10 +21288,10 @@ impl GpuDqnTrainer {
             let blocks = ((count + 255) / 256) as u32;
             debug_assert!(
                 (blocks as usize) <= MAX_BLOCKS_PER_ADAM,
-                "dqn_adam[{:?}] blocks={} exceeds MAX_BLOCKS_PER_ADAM={} — Pearl C \
-                 clamp_engage_per_block_buf cannot accommodate; resize buffer \
+                "dqn_adam[{:?} byte_off={}] blocks={} exceeds MAX_BLOCKS_PER_ADAM={} — Pearl C \
+                 clamp_engage_per_block_buf cannot accommodate; grow MAX_BLOCKS_PER_ADAM \
                  or chunk the launch.",
-                group, blocks, MAX_BLOCKS_PER_ADAM
+                group, byte_off, blocks, MAX_BLOCKS_PER_ADAM
             );
             let cfg = LaunchConfig {
                 grid_dim: (blocks, 1, 1),
@@ -21220,8 +21299,6 @@ impl GpuDqnTrainer {
                 shared_mem_bytes: 0,
             };
 
-            let engage_buf_offset: i32 = (group.idx() * MAX_BLOCKS_PER_ADAM) as i32;
-
             unsafe {
                 self.stream
                     .launch_builder(&self.adam_update_kernel)
@@ -21245,10 +21322,13 @@ impl GpuDqnTrainer {
                     .arg(&nan_flags_ptr)        // SP4 A14: sticky engage flag store
                     .arg(&diag_slot)            // SP4 A14: diag slot (50 for dqn_adam)
                     .arg(&engage_buf_ptr)       // SP4 A14: per-block engagement buf
-                    .arg(&engage_buf_offset)    // SP4 Layer B: per-group offset (resolves A14/A15 deferral)
+                    .arg(&engage_buf_offset)    // SP4 Layer B: per-group offset; SP4_ENGAGE_OFFSET_DISABLED for trunk-extras tail
                     .launch(cfg)
                     .map_err(|e| {
-                        MLError::ModelError(format!("dqn_adam_update_kernel[{:?}] launch: {e}", group))
+                        MLError::ModelError(format!(
+                            "dqn_adam_update_kernel[{:?} byte_off={}] launch: {e}",
+                            group, byte_off
+                        ))
                     })?;
             }
         }
@@ -21275,13 +21355,25 @@ impl GpuDqnTrainer {
     /// `[group_idx × MAX_BLOCKS_PER_ADAM, +MAX_BLOCKS_PER_ADAM)` range.
     ///
     /// **Layer B (2026-05-01) — split sub-launches landed.** The main
-    /// DQN Adam now runs as 3 sub-launches (Trunk/Value/Branches), each
-    /// writing per-block engagement counts at its own group offset.
-    /// `pearl_c_post_adam_engagement_check(group, count)` reads the
-    /// dedicated range for any of the 8 SP4 groups directly. The Layer A
-    /// "groups 1/2 accounted under group 0" limitation is resolved by
-    /// the same Layer B atomic commit that introduces the
+    /// DQN Adam now runs as 3 SP4-canonical sub-launches (Trunk / Value
+    /// / Branches), each writing per-block engagement counts at its own
+    /// group offset. `pearl_c_post_adam_engagement_check(group, count)`
+    /// reads the dedicated range for any of the 8 SP4 groups directly.
+    /// The Layer A "groups 1/2 accounted under group 0" limitation is
+    /// resolved by the same Layer B atomic commit that introduces the
     /// `WEIGHT_BOUND[group]` consumer migration.
+    ///
+    /// **Layer B fix-up (2026-05-01)** — the 4th sub-launch (trunk-
+    /// extras tail covering tensors `[33..163)`) is tagged DqnTrunk for
+    /// the WEIGHT_BOUND/WD_RATE bounds contract but uses
+    /// `SP4_ENGAGE_OFFSET_DISABLED` for Pearl C. Engagement for those
+    /// tensors **folds into the canonical DqnTrunk count** (same bound
+    /// drives both clamps; Pearl C is a per-bound diagnostic). Callers
+    /// passing `ParamGroup::DqnTrunk` here only see the canonical
+    /// trunk's engagement counts in the buffer; trunk-extras' clamp
+    /// engagements do not affect the rate-deficit EMA. This keeps
+    /// engagement accounting faithful to the WEIGHT_BOUND[DqnTrunk]
+    /// contract Pearl C is regulating.
     pub(crate) fn pearl_c_post_adam_engagement_check(
         &mut self,
         group: ParamGroup,
@@ -22253,9 +22345,11 @@ impl GpuDqnTrainer {
         // - smoke-test-rxhjh (Mech 6 100× + Mech 8 anchor reset): F1 NaN
         //   @ 2300 (worse than Mech 6 alone) — Mech 8 reverted; the
         //   persistent slow_ema across folds was unintentional protection.
-        // SP3 closes the cross-fold Adam pathology with Mech 9 (post-Adam
-        // weight clamp at 100 × Q_ABS_REF.max(1.0)) rather than further
-        // tuning the clip-anchor side of the chain.
+        // SP3 closed the cross-fold Adam pathology with Mech 9 (post-Adam
+        // weight clamp via host-side multiplier × Q_ABS_REF) rather than
+        // further tuning the clip-anchor side of the chain. Layer B
+        // generalised the bound to per-group `ISV[WEIGHT_BOUND[group]]`
+        // (Wiener-EMA over p99(|params|), no host multiplier).
         //
         // SP4 Layer B (Mech 6): adaptive clip upper bound = ISV[GRAD_CLIP_BOUND=168].
         // Producer `launch_sp4_grad_norm_p99` (Layer A Task A8) writes the
diff --git a/crates/ml/src/cuda_pipeline/iql_value_kernel.cu b/crates/ml/src/cuda_pipeline/iql_value_kernel.cu
index 36977fc8e..4cb749299 100644
--- a/crates/ml/src/cuda_pipeline/iql_value_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/iql_value_kernel.cu
@@ -208,7 +208,7 @@ void iql_adam_kernel(
     float max_grad_norm,
     const int* __restrict__ t_buf,  /* [1] Adam step on device */
     int total_params,
-    /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
+    /* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[IqlHigh|IqlLow]].max(EPS_CLAMP_FLOOR)`)
      * — same pattern wired into all five Adam kernels. 0=disabled (debug only). */
     float weight_clamp_max_abs,
     /* SP4 Task A14 (2026-04-30): Pearl C engagement counter. Same warp+
diff --git a/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu b/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
index 88c536946..cf5ff541f 100644
--- a/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
@@ -844,13 +844,14 @@ void iqn_grad_norm_phase2(
 }
 
 /* ── Adam Update Kernel ─────────────────────────────────────────────── */
-/* SP3 Mech 9 (close-out v2): trailing `weight_clamp_max_abs` arg matches the
+/* SP4 Layer B: trailing `weight_clamp_max_abs` arg matches the
  * pattern wired into all five Adam kernels (dqn/iqn/iql/attn/curiosity).
- * Bound = 100 × ISV[Q_ABS_REF].max(1.0) computed host-side. Targets the
- * IQN-internal weight buffer that the slot-26 cuBLAS GEMM reads —
- * `apply_iqn_trunk_gradient` saw NaN at slot 26 with finite inputs because
- * `online_params` here drifted via this Adam, never reached by the
- * single-kernel Mech 9 in `dqn_adam_update_kernel`. Disabled when bound==0.
+ * Bound = `ISV[WEIGHT_BOUND[ParamGroup::Iqn]].max(EPS_CLAMP_FLOOR)` computed
+ * host-side via `read_group_adam_bounds`. Targets the IQN-internal weight
+ * buffer that the slot-26 cuBLAS GEMM reads — `apply_iqn_trunk_gradient`
+ * saw NaN at slot 26 with finite inputs because `online_params` here
+ * drifted via this Adam, never reached by the single-kernel Mech 9 in
+ * `dqn_adam_update_kernel`. Disabled when bound==0.
  *
  * SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+block
  * tree-reduce mirroring `dqn_grad_norm_kernel` (no atomicAdd). Single
@@ -876,7 +877,7 @@ void iqn_adam_kernel(
     int b1_size,                   /* runtime: unused, for consistent interface */
     int b2_size,                   /* runtime: unused, for consistent interface */
     int b3_size,                   /* runtime: unused, for consistent interface */
-    float weight_clamp_max_abs,    /* SP3 Mech 9: ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0)). 0=disabled. */
+    float weight_clamp_max_abs,    /* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[Iqn]].max(EPS_CLAMP_FLOOR)`). 0=disabled. */
     /* SP4 Task A14: Pearl C engagement counter. */
     int* __restrict__ nan_flags_buf,              /* sticky overflow flag */
     int diag_slot,                                /* nan_flags_buf index for this Adam kernel */
@@ -918,11 +919,12 @@ void iqn_adam_kernel(
             float p = params[tid];
             p = p - lr * (m_hat / (sqrtf(v_hat) + eps) + weight_decay * p);
 
-            /* SP3 Mech 9 + SP4 Task A14: post-Adam |p| clamp +
+            /* SP4 Layer B + SP4 Task A14: post-Adam |p| clamp +
              * engagement-counter increment. Bound is host-supplied
-             * (100 × Q_ABS_REF.max(1.0)). One OOM below slot 48
-             * diagnostic threshold (1e3 × Q_ABS_REF) so the regression
-             * sentinel keeps 10× firing headroom. */
+             * (`ISV[WEIGHT_BOUND[Iqn]].max(EPS_CLAMP_FLOOR)`). The
+             * ISV producer (Wiener-EMA over p99(|params|)) keeps the
+             * bound naturally above legitimate weight magnitudes;
+             * `EPS_CLAMP_FLOOR=1.0` is a cold-start floor only. */
             if (weight_clamp_max_abs > 0.0f) {
                 if (fabsf(p) > weight_clamp_max_abs) {
                     if (nan_flags_buf != nullptr && diag_slot >= 0) {
diff --git a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
index 53e2f2b8d..bc553b218 100644
--- a/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
@@ -52,22 +52,30 @@ pub const SP4_SLOT_END: usize = 171;
 
 /// Maximum block grid size per Adam-kernel launch — the engagement
 /// counter buffer reserves `MAX_BLOCKS_PER_ADAM` slots per param group,
-/// indexed by `blockIdx.x`. 256 covers up to ≈65k params per group at
-/// blockDim=256 threads, comfortably above all current production
-/// param-group sizes (DQN trunk≈30k, DQN value≈4k, branches≈80k —
-/// branch group is the only one exceeding 65k, but the launcher caps
-/// `blocks` at `MAX_BLOCKS_PER_ADAM` defensively in launch code).
-pub const MAX_BLOCKS_PER_ADAM: usize = 256;
+/// indexed by `blockIdx.x`. 4096 covers up to ≈1M params per group at
+/// blockDim=256 threads.
+///
+/// **Layer B fix-up (2026-05-01)**: bumped 256 → 4096 to accommodate
+/// the trunk-extras tail sub-launch (tensors `[33..163)` co-located
+/// under DqnTrunk). Production cfg `shared_h2=256, num_atoms=51,
+/// adv_h=128, MOE_NUM_EXPERTS=8, MOE_EXPERT_BOTTLENECK=64` puts
+/// trunk-extras at ≈600k floats → ≈2400 blocks at blockDim=256, so
+/// 4096 leaves a comfortable safety margin (1.7×) for future config
+/// growth without wasted memory (each engage slot is 4 bytes × 11
+/// groups × 4096 = 180 KiB total — negligible vs. params buffer).
+/// Other groups (DqnValue ≈ 4k params, DqnBranches ≈ 80k, IQN/IQL/
+/// Attn/Curiosity bounded similarly) easily fit within 4096 blocks.
+pub const MAX_BLOCKS_PER_ADAM: usize = 4096;
 
 /// Total length of the Pearl C engagement-counter mapped-pinned buffer
 /// (`clamp_engage_per_block_buf`). The 8 main param groups occupy
-/// offsets [0..2048) (8 × 256). Curiosity has 4 sub-launches (W1/b1/
-/// W2/b2) sharing the same kernel; to avoid block-index collisions
-/// between sub-launches, each gets its own offset slot. Sub-launch 0
-/// reuses Curiosity's main group offset (group_idx=7 → 1792). Sub-
-/// launches 1/2/3 occupy extra slots at offsets 2048/2304/2560 — i.e.
-/// 3 extra logical "groups" beyond the canonical 8. Total length:
-/// (8 + 3) × 256 = 11 × 256 = 2816 i32s. Host-side
+/// offsets `[0..8 × MAX_BLOCKS_PER_ADAM)`. Curiosity has 4 sub-launches
+/// (W1/b1/W2/b2) sharing the same kernel; to avoid block-index
+/// collisions between sub-launches, each gets its own offset slot.
+/// Sub-launch 0 reuses Curiosity's main group offset (group_idx=7).
+/// Sub-launches 1/2/3 occupy extra slots — 3 extra logical "groups"
+/// beyond the canonical 8. Total length: (8 + 3) × MAX_BLOCKS_PER_ADAM
+/// = 11 × 4096 = 45056 i32s. Host-side
 /// `pearl_c_post_adam_engagement_check` for `ParamGroup::Curiosity`
 /// sums all 4 sub-launch ranges.
 pub const SP4_ENGAGE_EXTRA_CURIOSITY_SUBLAUNCHES: usize = 3;
@@ -212,18 +220,22 @@ mod tests {
     #[test]
     fn pearl_c_engage_buf_layout() {
         // 8 main groups + 3 extra curiosity sub-launch slots.
-        assert_eq!(MAX_BLOCKS_PER_ADAM, 256);
+        // Layer B fix-up (2026-05-01): MAX_BLOCKS_PER_ADAM grown 256 → 4096
+        // to fit trunk-extras tail (≈2400 blocks at production cfg).
+        assert_eq!(MAX_BLOCKS_PER_ADAM, 4096);
         assert_eq!(SP4_PARAM_GROUP_COUNT, 8);
         assert_eq!(SP4_ENGAGE_EXTRA_CURIOSITY_SUBLAUNCHES, 3);
-        assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * 256);
-        assert_eq!(SP4_ENGAGE_BUF_LEN, 2816);
+        assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * MAX_BLOCKS_PER_ADAM);
+        assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * 4096);
+        assert_eq!(SP4_ENGAGE_BUF_LEN, 45056);
 
         // Curiosity sub-launch offsets are distinct + contiguous starting
-        // from the canonical Curiosity group offset.
-        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W1, 7 * 256);
-        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B1, 8 * 256);
-        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W2, 9 * 256);
-        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B2, 10 * 256);
+        // from the canonical Curiosity group offset. They're derived from
+        // `MAX_BLOCKS_PER_ADAM` so they auto-track its growth.
+        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W1, 7 * MAX_BLOCKS_PER_ADAM as i32);
+        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B1, 8 * MAX_BLOCKS_PER_ADAM as i32);
+        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W2, 9 * MAX_BLOCKS_PER_ADAM as i32);
+        assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B2, 10 * MAX_BLOCKS_PER_ADAM as i32);
 
         // Highest sub-launch end must fit within the buffer.
         let highest_end = (SP4_ENGAGE_OFFSET_CURIOSITY_B2 as usize) + MAX_BLOCKS_PER_ADAM;
diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs
index a0f1fff37..5990a4179 100644
--- a/crates/ml/src/trainers/dqn/fused_training.rs
+++ b/crates/ml/src/trainers/dqn/fused_training.rs
@@ -1619,16 +1619,11 @@ impl FusedTrainingCtx {
                 // (cubin-shared with GpuAttention). TLOB params co-live with
                 // attention in the Attn (group 6) param-group taxonomy, so
                 // the same ISV-driven WEIGHT_BOUND[Attn] / WD_RATE[Attn]
-                // contract applies. ε floor (EPS_CLAMP_FLOOR = 1.0) handles
-                // cold-start; `.max(0.0)` enforces wd non-negativity.
-                let tlob_weight_clamp_max_abs: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-                let tlob_weight_decay_value: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(0.0);
+                // contract applies. `read_group_adam_bounds` enforces the
+                // EPS_CLAMP_FLOOR=1.0 ε for cold-start and AdamW
+                // non-negativity.
+                let (tlob_weight_clamp_max_abs, tlob_weight_decay_value) =
+                    self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
                 self.gpu_tlob
                     .adam_step(
                         self.trainer.config().lr,
@@ -1924,26 +1919,14 @@ impl FusedTrainingCtx {
             //
             // SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
             // `iql_adam_kernel`. Per-group bounds: IqlHigh (group 4) and
-            // IqlLow (group 5) each read `ISV[WEIGHT_BOUND[group]]`. ε
-            // floor (EPS_CLAMP_FLOOR = 1.0) handles cold-start ISV. AdamW
-            // weight_decay sourced from `ISV[WD_RATE[group]]` with
-            // `.max(0.0)` non-negativity contract.
-            let iql_high_clamp: f32 = self.trainer
-                .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                    crate::cuda_pipeline::ParamGroup::IqlHigh.idx()))
-                .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-            let iql_high_wd: f32 = self.trainer
-                .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                    crate::cuda_pipeline::ParamGroup::IqlHigh.idx()))
-                .max(0.0);
-            let iql_low_clamp: f32 = self.trainer
-                .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                    crate::cuda_pipeline::ParamGroup::IqlLow.idx()))
-                .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-            let iql_low_wd: f32 = self.trainer
-                .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                    crate::cuda_pipeline::ParamGroup::IqlLow.idx()))
-                .max(0.0);
+            // IqlLow (group 5) each read `ISV[WEIGHT_BOUND[group]]` and
+            // `ISV[WD_RATE[group]]` via `read_group_adam_bounds`, which
+            // applies the EPS_CLAMP_FLOOR=1.0 ε for cold-start ISV and
+            // the AdamW non-negativity contract.
+            let (iql_high_clamp, iql_high_wd) =
+                self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::IqlHigh);
+            let (iql_low_clamp, iql_low_wd) =
+                self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::IqlLow);
             let states_f32 = self.trainer.states_buf();
             self.gpu_iql.train_value_step(states_f32, iql_high_clamp, iql_high_wd)
                 .map_err(|e| anyhow::anyhow!("IQL high train: {e}"))?;
@@ -2042,18 +2025,12 @@ impl FusedTrainingCtx {
 
                 // SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
                 // `iqn_adam_kernel`. IQN params live in group 3
-                // (`ParamGroup::Iqn`); read `ISV[WEIGHT_BOUND[Iqn]]` with
-                // EPS_CLAMP_FLOOR = 1.0 ε for cold-start. AdamW
-                // weight_decay sourced from `ISV[WD_RATE[Iqn]]` with
-                // `.max(0.0)` non-negativity contract.
-                let weight_clamp_max_abs: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                        crate::cuda_pipeline::ParamGroup::Iqn.idx()))
-                    .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-                let weight_decay_value: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                        crate::cuda_pipeline::ParamGroup::Iqn.idx()))
-                    .max(0.0);
+                // (`ParamGroup::Iqn`); `read_group_adam_bounds` returns
+                // `ISV[WEIGHT_BOUND[Iqn]]` (with EPS_CLAMP_FLOOR=1.0 ε
+                // for cold-start) and `ISV[WD_RATE[Iqn]]` (with AdamW
+                // non-negativity).
+                let (weight_clamp_max_abs, weight_decay_value) =
+                    self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
 
                 match iqn.execute_training_pipeline(
                     online_h_s2, next_states_buf, &self.target_dueling,
@@ -2130,17 +2107,12 @@ impl FusedTrainingCtx {
                 let mgn = self.trainer.config().max_grad_norm;
                 // SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
                 // `attn_adam_kernel`. Attention params live in group 6
-                // (`ParamGroup::Attn`); read `ISV[WEIGHT_BOUND[Attn]]` with
-                // EPS_CLAMP_FLOOR = 1.0 ε for cold-start. AdamW weight_decay
-                // sourced from `ISV[WD_RATE[Attn]]` with `.max(0.0)`.
-                let attn_weight_clamp_max_abs: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-                let attn_weight_decay_value: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(0.0);
+                // (`ParamGroup::Attn`); `read_group_adam_bounds` returns
+                // `ISV[WEIGHT_BOUND[Attn]]` (with EPS_CLAMP_FLOOR=1.0 ε
+                // for cold-start) and `ISV[WD_RATE[Attn]]` (with AdamW
+                // non-negativity).
+                let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
+                    self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
                 attn.adam_step(lr, mgn, Some(&attn_stream_ref), ws_override,
                                attn_weight_clamp_max_abs, attn_weight_decay_value)
                     .map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
@@ -2177,14 +2149,8 @@ impl FusedTrainingCtx {
                 let mgn = self.trainer.config().max_grad_norm;
                 // SP4 Layer B (Mech 9): same ISV-driven bound as parallel arm.
                 // Sequential-fallback path.
-                let attn_weight_clamp_max_abs: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-                let attn_weight_decay_value: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                        crate::cuda_pipeline::ParamGroup::Attn.idx()))
-                    .max(0.0);
+                let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
+                    self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
                 attn.adam_step(lr, mgn, None, None,
                                attn_weight_clamp_max_abs, attn_weight_decay_value)
                     .map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
@@ -2197,14 +2163,8 @@ impl FusedTrainingCtx {
 
                 // SP4 Layer B (Mech 9): same ISV-driven bound as the
                 // parallel arm above. Mirror the dqn_adam pattern.
-                let weight_clamp_max_abs: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::weight_bound(
-                        crate::cuda_pipeline::ParamGroup::Iqn.idx()))
-                    .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
-                let weight_decay_value: f32 = self.trainer
-                    .read_isv_signal_at(crate::cuda_pipeline::wd_rate(
-                        crate::cuda_pipeline::ParamGroup::Iqn.idx()))
-                    .max(0.0);
+                let (weight_clamp_max_abs, weight_decay_value) =
+                    self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
 
                 match iqn.execute_training_pipeline(
                     online_h_s2, next_states_buf, &self.target_dueling,
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 87f6f9e42..0f854eb89 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2780,3 +2780,57 @@ Verification (local, RTX 3050 Ti, post-commit):
   ZERO matches.
 - `cargo test -p ml --lib --offline`: 928 passed, 14 failed (all 14
   pre-existing on HEAD `1389d1c81`; no new failures introduced).
+
+**Layer B fix-up (2026-05-01)**: code-quality reviewer surfaced a silent
+correctness regression in the original 3-way `launch_adam_update` split —
+tensors `[33..163)` (bottleneck, VSN bottleneck, GLU gates, KAN, regime
+gate, spacing, multi-horizon value heads, risk, ISV encoder, VSN per-group,
+aux heads, MoE — ~70% of the DQN's 163 weight tensors) were silently no
+longer Adam-updated (grads computed, m/v allocated, no parameter step).
+The 3-group SP4 taxonomy (DqnTrunk / DqnValue / DqnBranches via
+`param_group_buffers`) only maps tensors `[0..33)`. Resolution:
+- 4th sub-launch added — trunk-extras tail covering `[33..163)`,
+  tagged `ParamGroup::DqnTrunk` for shared `WEIGHT_BOUND/WD_RATE`
+  bounds, with `SP4_ENGAGE_OFFSET_DISABLED` so Pearl C engagement
+  rolls into trunk's accounting (same bound drives both clamps).
+- `MAX_BLOCKS_PER_ADAM` grown 256 → 4096 to fit trunk-extras (~2400
+  blocks at production cfg `shared_h2=256`, `MOE_NUM_EXPERTS=8`).
+  `SP4_ENGAGE_BUF_LEN` auto-grows: 11 × 4096 = 45056 ints (≈180 KiB
+  total; engage offsets are derived from `MAX_BLOCKS_PER_ADAM` so
+  they auto-track).
+- `read_group_adam_bounds(group) -> (clamp, wd)` helper introduced
+  on `GpuDqnTrainer`; 6 verbose 2-line ISV-read patterns in
+  `fused_training.rs` collapsed to one-liners (TLOB + IQL high +
+  IQL low + IQN parallel + Attn parallel + Attn sequential + IQN
+  sequential = 7 call sites). The 4-way Adam loop in
+  `launch_adam_update` also calls the helper.
+- Coverage `debug_assert_eq!` added: `trunk + value + branches +
+  trunk_extras == total_params`. Prevents future regressions of
+  this kind.
+- 9 stale doc-comments referring to "100 × Q_ABS_REF.max(1.0)"
+  replaced by current "ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)"
+  contract:
+    - `curiosity_training_kernel.cu:324`
+    - `gpu_curiosity_trainer.rs:509,758`
+    - `iqn_dual_head_kernel.cu:847,879,923`
+    - `iql_value_kernel.cu:211`
+    - `attention_backward_kernel.cu:80`
+    - `dqn_utility_kernels.cu:145,200`
+    - `gpu_dqn_trainer.rs:22349` (historical SP3 reference)
+
+Per `feedback_no_partial_refactor`: trunk-extras consumers (= no
+dedicated SP4 producer yet) and bound (= reused trunk's
+`WEIGHT_BOUND/WD_RATE`) ship together with the launch fix in this
+commit. A future task may split sub-trunk regions (VSN, MoE, etc.)
+into dedicated SP4 producers once warranted by signal analysis.
+
+Verification (post-fix-up):
+- `cargo check -p ml --offline`: clean, same 11 pre-existing warnings.
+- `cargo test -p ml --lib --offline -- sp4_wiener_ema sp4_isv_slots
+  state_reset_registry`: 11 passed (incl. updated
+  `pearl_c_engage_buf_layout` for MAX_BLOCKS_PER_ADAM=4096).
+- `cargo test -p ml --lib --offline`: 928 passed, 14 failed (no new
+  failures vs. Layer B initial run).
+- `git grep -nE "100 × Q_ABS_REF\\.max\\(1\\.0\\)|100\\.0f \\*
+  q_abs_ref|100 \\* Q_ABS_REF" crates/ml/src/cuda_pipeline/
+  crates/ml/src/trainers/`: ZERO matches.