diff --git a/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu b/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu
index 16afce46d..43c5a11be 100644
--- a/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu
@@ -9,7 +9,11 @@
  * Reads:
  *   spacing_raw[b]  — learnable softmax spacing params for branch b,
  *                     accessed via pre-resolved device pointers passed as args.
- *   isv_signals     — ISV bus (per-branch v-range at slots 23+2*b / 24+2*b).
+ *   isv_signals     — ISV bus:
+ *     SP5 Layer B (Pearl 1): per-branch v-range at
+ *       ATOM_V_CENTER_BASE=174: v_center[b] = isv_signals[174 + b]
+ *       ATOM_V_HALF_BASE=178:   v_half[b]   = isv_signals[178 + b]
+ *     (SP4 slots 23+2*b / 24+2*b are superseded by SP5 Pearl 1.)
  *
  * Writes:
  *   positions_out   — atom grid [4 * num_atoms], branch b at offset b*num_atoms.
@@ -17,7 +21,7 @@
  * Formula (same as adaptive_atom_positions in experience_kernels.cu):
  *   spacing = softmax(spacing_raw)
  *   positions[j] = v_min + (v_max - v_min) * cumsum(spacing)[j]
- *   with v_min = isv[23+2*b] - isv[24+2*b], v_max = isv[23+2*b] + isv[24+2*b].
+ *   with v_min = isv[174+b] - isv[178+b], v_max = isv[174+b] + isv[178+b].
  */
 
 /* SP4 Layer B: per-branch atom-position bound from ISV[ATOM_POS_BOUND_BASE + branch].
@@ -25,6 +29,34 @@
  */
 #define ATOM_POS_BOUND_BASE 132
 
+/* SP5 Layer B (Pearl 1): per-branch v_center and v_half from ISV.
+ * Mirrors sp5_isv_slots.rs: ATOM_V_CENTER_BASE=174, ATOM_V_HALF_BASE=178. */
+#define ATOM_V_CENTER_BASE 174
+#define ATOM_V_HALF_BASE   178
+
+/* SP5 Layer B (Pearl 1-ext): per-branch num_atoms from ISV.
+ * Mirrors sp5_isv_slots.rs: ATOM_NUM_ATOMS_BASE=274.
+ * Valid discrete values: {16, 32, 64}. Rounds ISV float to nearest valid
+ * count. Cold-start floor: if ISV=0 (pre-observation), fall back to the
+ * global `num_atoms` parameter (Invariant-1 carve-out). */
+#define ATOM_NUM_ATOMS_BASE 274
+
+/* Round an ISV-driven f32 num_atoms signal to the nearest valid discrete C51
+ * atom count {16, 32, 64}, bounded to [16, num_atoms_max]. */
+__device__ __forceinline__ int round_to_valid_num_atoms(float isv_val, int num_atoms_max) {
+    if (isv_val < 1.0f) return num_atoms_max;  /* cold-start floor */
+    int v = (int)(isv_val + 0.5f);
+    /* Snap to nearest of {16, 32, 64}: */
+    int d16 = abs(v - 16);
+    int d32 = abs(v - 32);
+    int d64 = abs(v - 64);
+    int best = (d16 <= d32 && d16 <= d64) ? 16 : (d32 <= d64) ? 32 : 64;
+    /* Clamp to [16, num_atoms_max]: never exceed the buffer's allocated stride */
+    if (best > num_atoms_max) best = num_atoms_max;
+    if (best < 16)            best = 16;
+    return best;
+}
+
 extern "C" __global__ void atoms_update(
     unsigned long long spacing_ptr_b0,   /* device ptr to spacing_raw[dir] */
     unsigned long long spacing_ptr_b1,   /* device ptr to spacing_raw[mag] */
@@ -39,15 +71,30 @@ extern "C" __global__ void atoms_update(
 
     int tid = threadIdx.x;
     extern __shared__ float shmem[];
+    /* Shared memory is allocated for global `num_atoms` (worst case) at launch.
+     * eff_na ≤ num_atoms, so both halves fit safely. */
     float* s_softmax = shmem;
     float* s_cumsum  = shmem + num_atoms;
 
-    /* Per-branch v-range from ISV: centre at 23+2*b, half at 24+2*b */
-    float v_center = isv_signals[23 + 2 * branch];
-    float v_half   = isv_signals[24 + 2 * branch];
+    /* SP5 Layer B (Pearl 1): per-branch v-range from ISV.
+     * v_center[b] = ISV[ATOM_V_CENTER_BASE + b] = ISV[174 + b]
+     * v_half[b]   = ISV[ATOM_V_HALF_BASE   + b] = ISV[178 + b]
+     * EPS_CLAMP_FLOOR = 1.0 (Invariant 1 anchor): pre-first-observation ISV
+     * reads as 0; fmaxf prevents zero/negative half-width collapsing the range. */
+    float v_center = isv_signals[ATOM_V_CENTER_BASE + branch];
+    float v_half   = fmaxf(isv_signals[ATOM_V_HALF_BASE + branch], 1.0f);
     float v_min    = v_center - v_half;
     float v_max    = v_center + v_half;
 
+    /* SP5 Layer B (Pearl 1-ext): per-branch effective num_atoms from ISV.
+     * ISV[ATOM_NUM_ATOMS_BASE + branch] = ISV[274 + branch].
+     * Rounded to valid discrete {16,32,64}; cold-start (ISV=0) falls back to
+     * the global `num_atoms` parameter (Invariant-1 carve-out). The global
+     * `num_atoms` is still used for shared-memory stride and output-buffer
+     * stride (buffer is allocated as 4 * config.num_atoms). */
+    int eff_na = round_to_valid_num_atoms(
+        isv_signals[ATOM_NUM_ATOMS_BASE + branch], num_atoms);
+
     /* Resolve spacing_raw pointer for this branch */
     const float* spacing_raw;
     switch (branch) {
@@ -57,16 +104,16 @@ extern "C" __global__ void atoms_update(
         default: spacing_raw = (const float*)spacing_ptr_b3; break;
     }
 
-    /* Load spacing_raw into shared memory */
-    if (tid < num_atoms) {
+    /* Load spacing_raw into shared memory (only eff_na elements are valid) */
+    if (tid < eff_na) {
         s_softmax[tid] = spacing_raw[tid];
     }
     __syncthreads();
 
-    /* Find max for numerical stability (single thread, num_atoms <= 256 is tiny) */
+    /* Find max for numerical stability (single thread, eff_na <= 256 is tiny) */
     if (tid == 0) {
         float max_val = s_softmax[0];
-        for (int j = 1; j < num_atoms; j++)
+        for (int j = 1; j < eff_na; j++)
             max_val = fmaxf(max_val, s_softmax[j]);
         s_cumsum[0] = max_val;  /* store max temporarily */
     }
@@ -75,7 +122,7 @@ extern "C" __global__ void atoms_update(
     __syncthreads();
 
     /* Compute softmax */
-    if (tid < num_atoms) {
+    if (tid < eff_na) {
         s_softmax[tid] = expf(s_softmax[tid] - max_val);
     }
     __syncthreads();
@@ -83,10 +130,10 @@ extern "C" __global__ void atoms_update(
     /* Sum for normalization */
     if (tid == 0) {
         float sum = 0.0f;
-        for (int j = 0; j < num_atoms; j++)
+        for (int j = 0; j < eff_na; j++)
             sum += s_softmax[j];
         float inv_sum = 1.0f / fmaxf(sum, 1e-8f);
-        for (int j = 0; j < num_atoms; j++)
+        for (int j = 0; j < eff_na; j++)
             s_softmax[j] *= inv_sum;
     }
     __syncthreads();
@@ -94,14 +141,17 @@ extern "C" __global__ void atoms_update(
     /* Inclusive prefix sum (sequential, tid=0 only) */
     if (tid == 0) {
         s_cumsum[0] = s_softmax[0];
-        for (int j = 1; j < num_atoms; j++)
+        for (int j = 1; j < eff_na; j++)
             s_cumsum[j] = s_cumsum[j-1] + s_softmax[j];
     }
     __syncthreads();
 
-    /* Write positions: positions[0] = v_min, positions[j>0] = v_min + range * cumsum[j-1] */
+    /* Write positions: positions[0] = v_min, positions[j>0] = v_min + range * cumsum[j-1].
+     * Output stride is the global `num_atoms` (buffer layout invariant). Positions
+     * beyond eff_na are left unchanged (they are not read — C51 consumers also
+     * use eff_na from ISV). */
     float* out = positions_out + (long long)branch * num_atoms;
-    if (tid < num_atoms) {
+    if (tid < eff_na) {
         float range = v_max - v_min;
         float frac  = (tid == 0) ? 0.0f : s_cumsum[tid - 1];
         float new_pos = v_min + range * frac;
diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
index c21af836b..b88d4fde7 100644
--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -254,7 +254,8 @@ extern "C" __global__ void backtest_env_step(
         &prev_position_sign,
         &actual_dir,
         &actual_mag,
-        &trail_triggered
+        &trail_triggered,
+        isv_signals              /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
     );
     (void)prev_position; (void)prev_position_sign; (void)trail_triggered;
 
@@ -508,7 +509,8 @@ extern "C" __global__ void backtest_env_step_batch(
             &prev_position_sign,
             &actual_dir,
             &actual_mag,
-            &trail_triggered
+            &trail_triggered,
+            isv_signals              /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
         );
         (void)prev_position; (void)prev_position_sign; (void)trail_triggered;
 
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 56fdfd169..31fa50aa9 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -2124,7 +2124,8 @@ extern "C" __global__ void experience_env_step(
         &prev_position_sign_core,
         &actual_dir_core,
         &actual_mag_core,
-        &trail_triggered
+        &trail_triggered,
+        isv_signals_ptr      /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */
     );
     /* actual_mag_core is now consumed by Task 2.X per-magnitude instrumentation
      * at the trade-lifecycle block below (~line 1760). actual_dir_core +
diff --git a/crates/ml/src/cuda_pipeline/gpu_attention.rs b/crates/ml/src/cuda_pipeline/gpu_attention.rs
index 3d107f5f0..43bb8fd9b 100644
--- a/crates/ml/src/cuda_pipeline/gpu_attention.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_attention.rs
@@ -1064,6 +1064,12 @@ impl GpuAttention {
         /* SP4 Layer B (Mech 9): ISV-driven AdamW weight_decay rate.
          * Caller reads `ISV[WD_RATE[Attn=6]]` with `.max(0.0)`. */
         weight_decay_value: f32,
+        /* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Attn=6.
+         * Caller reads ISV[ADAM_BETA1_BASE+6], ISV[ADAM_BETA2_BASE+6], ISV[ADAM_EPS_BASE+6]
+         * and applies Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
+        attn_beta1: f32,
+        attn_beta2: f32,
+        attn_epsilon: f32,
     ) -> Result<(), MLError> {
         let stream = override_stream.unwrap_or(&self.stream);
         // override_workspace is accepted for API symmetry; adam_step does not use
@@ -1111,9 +1117,11 @@ impl GpuAttention {
             shared_mem_bytes: 0,
         };
 
-        let beta1 = 0.9_f32;
-        let beta2 = 0.999_f32;
-        let eps = 1e-8_f32;
+        // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
+        // FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+6] / ISV[ADAM_EPS_BASE+6].
+        let beta1 = attn_beta1;
+        let beta2 = attn_beta2;
+        let eps   = attn_epsilon;
         // SP4 Layer B (Mech 9): AdamW weight_decay sourced from
         // ISV[WD_RATE[Attn=6]] by the caller (`FusedTrainingCtx`) and
         // threaded through `adam_step`. `.max(0.0)` upstream enforces the
diff --git a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
index 298d3395b..e087b9c79 100644
--- a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs
@@ -522,6 +522,12 @@ fn launch_adam_step(
     batch_size: usize,
     step: i32,
     weight_clamp_max_abs: f32,
+    /* SP5 Pearl 4: per-group Adam β1/β2/ε read from ISV[ADAM_BETA1_BASE+7],
+     * ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7] by caller and threaded
+     * through train_on_collector_buffers. Curiosity = ParamGroup index 7. */
+    cur_beta1: f32,
+    cur_beta2: f32,
+    cur_epsilon: f32,
     /* SP4 Task A14: per-sub-launch Pearl C engagement-counter args.
      * `engage_buf_offset` is unique per sub-buffer (W1/b1/W2/b2) so
      * the 4 sub-launches don't overwrite each other's per-block counts.
@@ -556,9 +562,9 @@ fn launch_adam_step(
             .arg(&np_i32)
             .arg(&bs_i32)
             .arg(&ADAM_LR)
-            .arg(&ADAM_BETA1)
-            .arg(&ADAM_BETA2)
-            .arg(&ADAM_EPS)
+            .arg(&cur_beta1)
+            .arg(&cur_beta2)
+            .arg(&cur_epsilon)
             .arg(&step)
             .arg(&weight_clamp_max_abs) // SP3 Mech 9 (close-out v2): post-Adam |p| clamp
             .arg(&nan_flags_ptr)        // SP4 A14: sticky engage flag store
@@ -759,6 +765,13 @@ impl GpuCuriosityTrainer {
          * `ISV[WEIGHT_BOUND[Curiosity]].max(EPS_CLAMP_FLOOR)` host-side.
          * 0=disabled (debug only). */
         weight_clamp_max_abs: f32,
+        /* SP5 Pearl 4: per-group Adam β1/β2/ε for the four
+         * `curiosity_adam_step` launches. Caller reads ISV[ADAM_BETA1_BASE+7],
+         * ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7] from FusedTrainingCtx
+         * and threads through train_curiosity_gpu → here. */
+        cur_beta1: f32,
+        cur_beta2: f32,
+        cur_epsilon: f32,
     ) -> Result<(), MLError> {
         if n_samples < 2 {
             return Ok(());
@@ -1004,24 +1017,28 @@ impl GpuCuriosityTrainer {
             stream, &self.adam_func, &mut weights.w1, &self.grad_w1,
             CUR_W1_LEN, &mut self.adam_m_w1, &mut self.adam_v_w1,
             n_train, step, weight_clamp_max_abs,
+            cur_beta1, cur_beta2, cur_epsilon,
             nan_flags_ptr, diag_slot, engage_buf_ptr, off_w1,
         )?;
         launch_adam_step(
             stream, &self.adam_func, &mut weights.b1, &self.grad_b1,
             CUR_B1_LEN, &mut self.adam_m_b1, &mut self.adam_v_b1,
             n_train, step, weight_clamp_max_abs,
+            cur_beta1, cur_beta2, cur_epsilon,
             nan_flags_ptr, diag_slot, engage_buf_ptr, off_b1,
         )?;
         launch_adam_step(
             stream, &self.adam_func, &mut weights.w2, &self.grad_w2,
             CUR_W2_LEN, &mut self.adam_m_w2, &mut self.adam_v_w2,
             n_train, step, weight_clamp_max_abs,
+            cur_beta1, cur_beta2, cur_epsilon,
             nan_flags_ptr, diag_slot, engage_buf_ptr, off_w2,
         )?;
         launch_adam_step(
             stream, &self.adam_func, &mut weights.b2, &self.grad_b2,
             CUR_B2_LEN, &mut self.adam_m_b2, &mut self.adam_v_b2,
             n_train, step, weight_clamp_max_abs,
+            cur_beta1, cur_beta2, cur_epsilon,
             nan_flags_ptr, diag_slot, engage_buf_ptr, off_b2,
         )?;
 
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index a32524b25..86bfbf8b9 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -23428,11 +23428,11 @@ impl GpuDqnTrainer {
         use crate::cuda_pipeline::sp4_isv_slots::{
             ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX, SP4_ENGAGE_OFFSET_DISABLED,
         };
+        use crate::cuda_pipeline::sp5_isv_slots::{
+            ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE,
+        };
 
         let lr_ptr = self.lr_dev_ptr;  // pinned device-mapped — graph-safe, value read at replay
-        let beta1 = self.config.beta1;
-        let beta2 = self.config.beta2;
-        let epsilon = self.config.epsilon;
         let adaptive_clip_ptr = self.ptrs.adaptive_clip_buf;
         let norm_ptr = self.ptrs.grad_norm_buf;
         let t_ptr = self.ptrs.t_buf;
@@ -23518,6 +23518,16 @@ impl GpuDqnTrainer {
                 continue;
             }
 
+            // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε from ISV[226+g],
+            // ISV[234+g], ISV[242+g].  Floors are Invariant 1 carve-outs
+            // (numerical-stability): β1≥0.5 avoids momentum inversion,
+            // β2≥0.9 prevents exploding second-moment estimates, ε≥1e-12
+            // avoids division-by-zero.  Cold-start ISV reads 0 — floors fire.
+            let g = group.idx();
+            let beta1   = self.read_isv_signal_at(ADAM_BETA1_BASE + g).max(0.5_f32).min(0.9999_f32);
+            let beta2   = self.read_isv_signal_at(ADAM_BETA2_BASE + g).max(0.9_f32).min(0.99999_f32);
+            let epsilon = self.read_isv_signal_at(ADAM_EPS_BASE   + g).max(1e-12_f32);
+
             // Per-group ISV-driven post-Adam weight clamp + AdamW
             // weight_decay rate via the SP4 Layer B helper. Trunk-extras
             // sub-launch reuses DqnTrunk's bounds (its `group` tag).
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 30d07ddc9..a9423c7d1 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -4361,11 +4361,19 @@ impl GpuExperienceCollector {
     /// bound for `curiosity_adam_step`. The collector doesn't own ISV — the
     /// caller (training_loop) reads `100 × ISV[Q_ABS_REF].max(1.0)` from the
     /// FusedTrainingCtx and threads it through here. 0=disabled (debug only).
+    ///
+    /// `cur_beta1/cur_beta2/cur_epsilon` are SP5 Pearl 4: ISV-driven Adam
+    /// hyperparameters for ParamGroup::Curiosity (index 7). Caller reads
+    /// ISV[ADAM_BETA1_BASE+7], ISV[ADAM_BETA2_BASE+7], ISV[ADAM_EPS_BASE+7]
+    /// from FusedTrainingCtx and threads through here.
     pub fn train_curiosity_gpu(
         &mut self,
         n_episodes: usize,
         timesteps: usize,
         weight_clamp_max_abs: f32,
+        cur_beta1: f32,
+        cur_beta2: f32,
+        cur_epsilon: f32,
     ) -> Result<(), MLError> {
         if let Some(ref mut trainer) = self.curiosity_trainer {
             let total = n_episodes.saturating_mul(timesteps);
@@ -4378,6 +4386,9 @@ impl GpuExperienceCollector {
                 &self.actions_out,
                 total,
                 weight_clamp_max_abs,
+                cur_beta1,
+                cur_beta2,
+                cur_epsilon,
             )?;
         }
         Ok(())
diff --git a/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs b/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs
index e8f311ecd..789dbdeaa 100644
--- a/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs
@@ -824,6 +824,12 @@ impl GpuIqnHead {
         weight_clamp_max_abs: f32,
         /* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate. */
         weight_decay_value: f32,
+        /* SP5 Pearl 4: ISV-driven Adam β1/β2/ε for ParamGroup::Iqn (index 3).
+         * Caller reads ISV[ADAM_BETA1_BASE+3], ISV[ADAM_BETA2_BASE+3],
+         * ISV[ADAM_EPS_BASE+3] and threads through here. */
+        iqn_beta1: f32,
+        iqn_beta2: f32,
+        iqn_epsilon: f32,
     ) -> Result<f32, MLError> {
         // 1-2. Decode actions + DtoD copy rewards/dones
         self.prepare_buffers(dqn_actions_buf, dqn_rewards_buf, dqn_dones_buf)?;
@@ -837,6 +843,9 @@ impl GpuIqnHead {
             None,
             weight_clamp_max_abs,
             weight_decay_value,
+            iqn_beta1,
+            iqn_beta2,
+            iqn_epsilon,
         )
     }
 
@@ -879,6 +888,12 @@ impl GpuIqnHead {
         /* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate.
          * Caller reads `ISV[WD_RATE[Iqn=3]]` with `.max(0.0)` non-negativity. */
         weight_decay_value: f32,
+        /* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Iqn=3.
+         * Caller reads ISV[ADAM_BETA1_BASE+3], ISV[ADAM_BETA2_BASE+3], ISV[ADAM_EPS_BASE+3]
+         * and applies the Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
+        iqn_beta1: f32,
+        iqn_beta2: f32,
+        iqn_epsilon: f32,
     ) -> Result<f32, MLError> {
         let effective_stream = override_stream.unwrap_or(&self.stream);
         let (lt_ws_ptr, lt_ws_size) = override_workspace
@@ -1499,9 +1514,11 @@ impl GpuIqnHead {
         };
 
         let lr = self.config.lr;
-        let beta1 = self.config.beta1;
-        let beta2 = self.config.beta2;
-        let eps = self.config.epsilon;
+        // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
+        // FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+3] / ISV[ADAM_EPS_BASE+3].
+        let beta1 = iqn_beta1;
+        let beta2 = iqn_beta2;
+        let eps   = iqn_epsilon;
         // SP4 Layer B (Mech 9): AdamW weight_decay sourced from
         // ISV[WD_RATE[Iqn=3]] by the caller (`FusedTrainingCtx`) and
         // threaded through `execute_training_pipeline` alongside
@@ -1932,6 +1949,73 @@ impl GpuIqnHead {
         self.config.num_quantiles
     }
 
+    /// SP5 Pearl 5: Update τ schedule from ISV[IQN_TAU_BASE=250..270).
+    ///
+    /// Reads 4 branches × 5 quantiles from ISV, averages across branches per
+    /// quantile slot, applies Invariant-1 cold-start floors (if ISV reads 0,
+    /// fall back to `FIXED_TAUS[q]`), then re-uploads `online_taus`,
+    /// `target_taus`, and `cos_features` in-place via mapped-pinned staging.
+    ///
+    /// Must be called once per epoch from `FusedTrainingCtx` before the IQN
+    /// training pipeline runs. Non-fatal on mapped-pinned upload error
+    /// (logs a warning).
+    pub fn refresh_taus_from_isv(
+        &mut self,
+        isv_per_branch_taus: &[f32; 20], // ISV[250..270): branch b, quantile q at [b*5+q]
+    ) {
+        let n = self.config.num_quantiles;   // 5
+        let b = self.config.batch_size;      // batch dimension for tiling
+        let d = self.config.embed_dim;       // 64
+
+        // Average 4 branches per quantile, apply cold-start floor from FIXED_TAUS.
+        let mut taus = [0.0_f32; 5];
+        for q in 0..n {
+            let mean = (isv_per_branch_taus[q]
+                + isv_per_branch_taus[5 + q]
+                + isv_per_branch_taus[10 + q]
+                + isv_per_branch_taus[15 + q]) / 4.0_f32;
+            // Invariant-1 cold-start floor: if ISV hasn't been populated yet
+            // (all four branches read 0.0), retain FIXED_TAUS[q] to avoid
+            // collapsing the quantile distribution to 0.
+            taus[q] = if mean < 1e-6_f32 { FIXED_TAUS[q] } else { mean.min(1.0_f32) };
+        }
+
+        // Recompute cosine features [D, N]: cos_feat[dim + q*D] = cos(π·(d+1)·τ_q)
+        let mut cos_feat_host = vec![0.0_f32; d * n];
+        for q in 0..n {
+            let tau_q = taus[q];
+            for dim in 0..d {
+                cos_feat_host[dim + q * d] =
+                    (std::f32::consts::PI * ((dim + 1) as f32) * tau_q).cos();
+            }
+        }
+
+        // Tile taus [B, N] — each row is the same N quantile values repeated B times.
+        let mut tiled = Vec::with_capacity(b * n);
+        for _ in 0..b {
+            tiled.extend_from_slice(&taus[..n]);
+        }
+
+        // Upload all three buffers in-place via mapped-pinned (no new allocations).
+        if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
+            &self.stream, &tiled, &mut self.online_taus,
+        ) {
+            tracing::warn!("SP5 Pearl 5 online_taus upload failed (non-fatal): {e}");
+            return;
+        }
+        if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
+            &self.stream, &tiled, &mut self.target_taus,
+        ) {
+            tracing::warn!("SP5 Pearl 5 target_taus upload failed (non-fatal): {e}");
+            return;
+        }
+        if let Err(e) = super::mapped_pinned::upload_f32_via_pinned(
+            &self.stream, &cos_feat_host, &mut self.cos_features,
+        ) {
+            tracing::warn!("SP5 Pearl 5 cos_features upload failed (non-fatal): {e}");
+        }
+    }
+
     /// Compute CVaR-based position scaling from IQN quantiles.
     ///
     /// Runs the IQN forward-only kernel on `h_s2` (trunk activation),
diff --git a/crates/ml/src/cuda_pipeline/gpu_tlob.rs b/crates/ml/src/cuda_pipeline/gpu_tlob.rs
index 3ff6164c6..ca780bab1 100644
--- a/crates/ml/src/cuda_pipeline/gpu_tlob.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_tlob.rs
@@ -659,6 +659,12 @@ impl GpuTlob {
         max_grad_norm: f32,
         weight_clamp_max_abs: f32,
         weight_decay_value: f32,
+        /* SP5 Layer B (Pearl 4): ISV-derived Adam β1/β2/ε for ParamGroup::Attn=6.
+         * TLOB shares the Attn param-group taxonomy with GpuAttention; same ISV slots.
+         * Caller applies Invariant 1 floors (β1≥0.5, β2≥0.9, ε≥1e-12). */
+        attn_beta1: f32,
+        attn_beta2: f32,
+        attn_epsilon: f32,
     ) -> Result<(), MLError> {
         let tp = TLOB_TOTAL_PARAMS;
         let tp_i32 = tp as i32;
@@ -695,9 +701,12 @@ impl GpuTlob {
         self.adam_step += 1;
         unsafe { *self.t_pinned = self.adam_step; }
         let adam_blocks = (tp + 255) / 256;
-        let beta1 = 0.9_f32;
-        let beta2 = 0.999_f32;
-        let eps   = 1e-8_f32;
+        // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε threaded from
+        // FusedTrainingCtx via ISV[ADAM_BETA{1,2}_BASE+6] / ISV[ADAM_EPS_BASE+6].
+        // TLOB shares ParamGroup::Attn=6 with GpuAttention.
+        let beta1 = attn_beta1;
+        let beta2 = attn_beta2;
+        let eps   = attn_epsilon;
         // SP4 Layer B (Mech 9): AdamW weight_decay sourced from
         // ISV[WD_RATE[Attn=6]] by the caller and threaded through
         // `adam_step`. `.max(0.0)` upstream enforces the AdamW
diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh
index 45c46b1f2..b8bfdd150 100644
--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh
+++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh
@@ -18,6 +18,19 @@
 #ifndef TRADE_PHYSICS_CUH
 #define TRADE_PHYSICS_CUH
 
+/* ── SP5 Layer B: ISV slot constants consumed in unified_env_step_core ── */
+
+/* Pearl 6 (cross-fold-persistent Kelly): smoothed Kelly fraction.
+ * ISV[KELLY_F_SMOOTH_INDEX] = ISV[280].
+ * Mirrors sp5_isv_slots.rs::KELLY_F_SMOOTH_INDEX. */
+#define SP5_KELLY_F_SMOOTH_INDEX 280
+
+/* Pearl 8 (per-direction trailing stop distance): base trail distance per dir.
+ * ISV[TRAIL_DIST_PER_DIR_BASE + dir] = ISV[270..274).
+ * dir: 0=Short, 1=Hold, 2=Long, 3=Flat.
+ * Mirrors sp5_isv_slots.rs::TRAIL_DIST_PER_DIR_BASE. */
+#define SP5_TRAIL_DIST_PER_DIR_BASE 270
+
 /* ── Factored action → exposure index ─────────────────────────────────── */
 
 __device__ __forceinline__ int decode_exposure_index(
@@ -639,7 +652,11 @@ __device__ __forceinline__ void unified_env_step_core(
     int*   prev_position_sign_out,  /* sign of position BEFORE the trade — for caller's reversal logic */
     int*   actual_dir_out,       /* post-enforcement direction (Hold / Long / Short / Flat) — for caller's actions_history */
     int*   actual_mag_out,       /* post-enforcement magnitude bucket (0/1/2 for Quarter/Half/Full; 0 if Hold/Flat) */
-    int*   trail_triggered_out   /* 1 if trailing stop forced position to 0 */
+    int*   trail_triggered_out,  /* 1 if trailing stop forced position to 0 */
+    /* ── SP5 Layer B: ISV signal bus for Pearl 6 + Pearl 8 ──
+     * NULL-tolerant: when NULL, falls back to static defaults
+     * (0.5f for kelly_f_smooth, 0.005f for trail distance). */
+    const float* __restrict__ isv_signals_ptr  /* [SP5_SLOT_END] pinned ISV signals; NULL = static defaults */
 ) {
     /* ── Step 1: Decode action (4-branch) ── */
     int dir_idx        = decode_direction_4b(action_val, b1_size, b2_size, b3_size);
@@ -704,14 +721,31 @@ __device__ __forceinline__ void unified_env_step_core(
         float dir_conv_clamped = fminf(1.0f, fmaxf(0.0f, conviction));
         float mag_conv_clamped = fminf(1.0f, fmaxf(0.0f, magnitude_conviction));
         float policy_conviction = fmaxf(dir_conv_clamped, mag_conv_clamped);
-        float safety = fmaxf(health_safety, policy_conviction);
+
+        /* SP5 Layer B (Pearl 6): cross-fold-persistent smoothed Kelly fraction
+         * from ISV[KELLY_F_SMOOTH_INDEX=280]. When ISV is available and the
+         * slot has been populated (value > 0), incorporate it as an additional
+         * floor on `health_safety` — both are fractions ∈ [0, 1] measuring
+         * long-run sizing confidence. The max() is safe: when Kelly says
+         * "go bigger" than health alone would permit, we trust the data. When
+         * health says "be conservative", health dominates. Bootstrap-safe
+         * because ISV[280] starts at 0 and Pearl 6 uses EWMA α=0.01 — it
+         * stays near zero until real Kelly data accumulate.
+         * EPS_CLAMP_FLOOR = 0.001f: prevents a very small cross-fold Kelly
+         * from being amplified by the max() in a numerically unstable way
+         * (same cold-start pattern as SP4 EPS_CLAMP_FLOOR). */
+        float kelly_f_smooth = (isv_signals_ptr != NULL)
+            ? fmaxf(0.001f, fminf(1.0f, isv_signals_ptr[SP5_KELLY_F_SMOOTH_INDEX]))
+            : 0.5f;
+        float health_safety_sp5 = fmaxf(health_safety, kelly_f_smooth);
+        float safety = fmaxf(health_safety_sp5, policy_conviction);
         target_position = apply_kelly_cap(
             target_position,
             *win_count, *loss_count,
             *sum_wins,  *sum_losses,
             max_position_physics, safety,
             policy_conviction,
-            health_safety
+            health_safety_sp5
         );
     }
 
@@ -723,9 +757,19 @@ __device__ __forceinline__ void unified_env_step_core(
             float unrealized = (*position) * (close - *entry_price);
             trade_ret = unrealized / prev_equity;
         }
+        /* SP5 Layer B (Pearl 8): per-direction trailing stop base distance from
+         * ISV[TRAIL_DIST_PER_DIR_BASE + dir_idx] = ISV[270 + dir_idx].
+         * dir_idx ∈ {0=Short,1=Hold,2=Long,3=Flat}; Hold/Flat are excluded by
+         * the enclosing !is_hold_action guard so only Short (0) and Long (2)
+         * reach here in practice. Bootstrap floor: when ISV is NULL or the
+         * slot reads as 0 (pre-first-observation), fall back to the SP4
+         * static value of 0.005f (0.5% trail). */
+        float base_trail_dist = (isv_signals_ptr != NULL)
+            ? fmaxf(0.001f, isv_signals_ptr[SP5_TRAIL_DIST_PER_DIR_BASE + dir_idx])
+            : 0.005f;
         if (check_trailing_stop(
                 *hold_time, *max_equity, prev_equity,
-                trade_ret, 0.005f, vol_scale, trend_scale
+                trade_ret, base_trail_dist, vol_scale, trend_scale
             )) {
             target_position = 0.0f;
             trail_triggered = 1;
diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs
index 5990a4179..e69c81ef5 100644
--- a/crates/ml/src/trainers/dqn/fused_training.rs
+++ b/crates/ml/src/trainers/dqn/fused_training.rs
@@ -39,6 +39,13 @@ use crate::cuda_pipeline::gpu_tlob::GpuTlob;
 use crate::cuda_pipeline::gpu_dqn_trainer::{
     GpuDqnTrainConfig, GpuDqnTrainer, TAU_EFF_INDEX,
 };
+use crate::cuda_pipeline::sp5_isv_slots::{
+    BUDGET_C51_BASE, BUDGET_IQN_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE,
+    NOISY_SIGMA_BASE,
+    ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE,
+    IQN_TAU_BASE,
+    ATOM_NUM_ATOMS_BASE,
+};
 use crate::cuda_pipeline::gpu_her::{GpuHer, GpuHerConfig, parse_her_strategy};
 use crate::cuda_pipeline::gpu_iql_trainer::{GpuIqlConfig, GpuIqlTrainer};
 use crate::cuda_pipeline::gpu_iqn_head::{GpuIqnConfig, GpuIqnHead};
@@ -1624,12 +1631,20 @@ impl FusedTrainingCtx {
                 // non-negativity.
                 let (tlob_weight_clamp_max_abs, tlob_weight_decay_value) =
                     self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
+                // SP5 Layer B (Pearl 4): TLOB shares ParamGroup::Attn=6.
+                let tlob_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
+                let tlob_beta1   = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + tlob_g).max(0.5_f32).min(0.9999_f32);
+                let tlob_beta2   = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + tlob_g).max(0.9_f32).min(0.99999_f32);
+                let tlob_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE   + tlob_g).max(1e-12_f32);
                 self.gpu_tlob
                     .adam_step(
                         self.trainer.config().lr,
                         self.trainer.config().max_grad_norm,
                         tlob_weight_clamp_max_abs,
                         tlob_weight_decay_value,
+                        tlob_beta1,
+                        tlob_beta2,
+                        tlob_epsilon,
                     )
                     .map_err(|e| anyhow::anyhow!("TLOB Adam: {e}"))?;
             }
@@ -1982,6 +1997,17 @@ impl FusedTrainingCtx {
         // ── IQN prep: decode actions + DtoD copies (main stream, before fork) ──
         // Must complete before execute_training_pipeline runs on iqn_stream.
         if let Some(ref mut iqn) = self.gpu_iqn {
+            // SP5 Pearl 5: refresh τ schedule from ISV[IQN_TAU_BASE=250..270).
+            // Read 4 branches × 5 quantiles, average per quantile, apply cold-start
+            // floor from FIXED_TAUS, re-upload online_taus/target_taus/cos_features.
+            // Must run before prepare_buffers so the kernel sees updated tau values.
+            {
+                let mut isv_taus = [0.0_f32; 20];
+                for i in 0..20 {
+                    isv_taus[i] = self.trainer.read_isv_signal_at(IQN_TAU_BASE + i);
+                }
+                iqn.refresh_taus_from_isv(&isv_taus);
+            }
             iqn.set_cached_target_h_s2(self.trainer.tg_h_s2_ptr());
             let dqn_actions = self.trainer.actions_buf();
             let dqn_rewards = self.trainer.rewards_buf();
@@ -2031,12 +2057,20 @@ impl FusedTrainingCtx {
                 // non-negativity).
                 let (weight_clamp_max_abs, weight_decay_value) =
                     self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
+                // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Iqn=3.
+                let iqn_g = crate::cuda_pipeline::ParamGroup::Iqn.idx();
+                let iqn_beta1   = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + iqn_g).max(0.5_f32).min(0.9999_f32);
+                let iqn_beta2   = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + iqn_g).max(0.9_f32).min(0.99999_f32);
+                let iqn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE   + iqn_g).max(1e-12_f32);
 
                 match iqn.execute_training_pipeline(
                     online_h_s2, next_states_buf, &self.target_dueling,
                     Some(&iqn_stream_ref), ws_override,
                     weight_clamp_max_abs,
                     weight_decay_value,
+                    iqn_beta1,
+                    iqn_beta2,
+                    iqn_epsilon,
                 ) {
                     Ok(_) => { iqn_ok = true; }
                     Err(e) => {
@@ -2113,8 +2147,14 @@ impl FusedTrainingCtx {
                 // non-negativity).
                 let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
                     self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
+                // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Attn=6.
+                let attn_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
+                let attn_beta1   = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + attn_g).max(0.5_f32).min(0.9999_f32);
+                let attn_beta2   = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + attn_g).max(0.9_f32).min(0.99999_f32);
+                let attn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE   + attn_g).max(1e-12_f32);
                 attn.adam_step(lr, mgn, Some(&attn_stream_ref), ws_override,
-                               attn_weight_clamp_max_abs, attn_weight_decay_value)
+                               attn_weight_clamp_max_abs, attn_weight_decay_value,
+                               attn_beta1, attn_beta2, attn_epsilon)
                     .map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
 
                 // Record completion on attn_stream.
@@ -2151,8 +2191,14 @@ impl FusedTrainingCtx {
                 // Sequential-fallback path.
                 let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
                     self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
+                // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Attn=6.
+                let attn_g = crate::cuda_pipeline::ParamGroup::Attn.idx();
+                let attn_beta1   = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + attn_g).max(0.5_f32).min(0.9999_f32);
+                let attn_beta2   = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + attn_g).max(0.9_f32).min(0.99999_f32);
+                let attn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE   + attn_g).max(1e-12_f32);
                 attn.adam_step(lr, mgn, None, None,
-                               attn_weight_clamp_max_abs, attn_weight_decay_value)
+                               attn_weight_clamp_max_abs, attn_weight_decay_value,
+                               attn_beta1, attn_beta2, attn_epsilon)
                     .map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
             }
 
@@ -2165,12 +2211,20 @@ impl FusedTrainingCtx {
                 // parallel arm above. Mirror the dqn_adam pattern.
                 let (weight_clamp_max_abs, weight_decay_value) =
                     self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
+                // SP5 Layer B (Pearl 4): per-group Adam β1/β2/ε for Iqn=3.
+                let iqn_g = crate::cuda_pipeline::ParamGroup::Iqn.idx();
+                let iqn_beta1   = self.trainer.read_isv_signal_at(ADAM_BETA1_BASE + iqn_g).max(0.5_f32).min(0.9999_f32);
+                let iqn_beta2   = self.trainer.read_isv_signal_at(ADAM_BETA2_BASE + iqn_g).max(0.9_f32).min(0.99999_f32);
+                let iqn_epsilon = self.trainer.read_isv_signal_at(ADAM_EPS_BASE   + iqn_g).max(1e-12_f32);
 
                 match iqn.execute_training_pipeline(
                     online_h_s2, next_states_buf, &self.target_dueling,
                     None, None,
                     weight_clamp_max_abs,
                     weight_decay_value,
+                    iqn_beta1,
+                    iqn_beta2,
+                    iqn_epsilon,
                 ) {
                     Ok(_) => { iqn_ok = true; }
                     Err(e) => {
@@ -3214,20 +3268,38 @@ impl FusedTrainingCtx {
     /// Returns (c51_budget, iqn_budget, cql_budget, ens_budget). Caches to trainer fields
     /// for HEALTH_DIAG logging.
     ///
-    /// Formula:
-    ///   iqn_budget = 0.10 + 0.30 × health
-    ///   cql_budget = 0.10 × (1 − regime_stability) × health
-    ///   ens_budget = 0.05  (constant)
-    ///   c51_budget = 1.0 − iqn_budget − cql_budget − ens_budget
+    /// SP5 Layer B (Pearl 2): reads per-branch budget signals from ISV[190..210)
+    /// and collapses them to a single effective scale per loss component by
+    /// averaging across the 4 branches.  The Layer A producer (SP5 Task A3) runs
+    /// the Wiener-optimal EMA (Pearls A+D) so the values already encode
+    /// health/regime dynamics — no second-hand formula is needed here.
     ///
-    /// At health=1: iqn=0.40, cql variable (depends on regime_stability), ens=0.05, c51=rest.
-    /// At health=0: iqn=0.10, cql=0, ens=0.05, c51=0.85 (C51 takes over).
+    /// Cold-start floor: ISV reads 0 before the first observation; the
+    /// normalisation `/ 4.0` and clamp ensure we never pass 0 to the loss
+    /// scalers.  The floor values (0.05 C51, 0.05 IQN, 0.02 CQL, 0.02 Ens)
+    /// are Invariant 1 carve-outs (numerical-stability), not tuned constants.
     pub(crate) fn compute_adaptive_budgets(&mut self) -> (f32, f32, f32, f32) {
-        let (health, regime_stability) = self.trainer.read_isv_health_and_regime();
-        let iqn_budget = 0.10_f32 + 0.30_f32 * health;
-        let cql_budget = 0.10_f32 * (1.0 - regime_stability) * health;
-        let ens_budget = 0.05_f32;
-        let c51_budget = (1.0_f32 - iqn_budget - cql_budget - ens_budget).max(0.0);
+        // Mean of 4 per-branch slots → single effective scale.
+        let c51_budget = ((0..4_usize)
+            .map(|b| self.read_isv_signal_at(BUDGET_C51_BASE + b))
+            .sum::<f32>()
+            / 4.0_f32)
+            .max(0.05_f32);
+        let iqn_budget = ((0..4_usize)
+            .map(|b| self.read_isv_signal_at(BUDGET_IQN_BASE + b))
+            .sum::<f32>()
+            / 4.0_f32)
+            .max(0.05_f32);
+        let cql_budget = ((0..4_usize)
+            .map(|b| self.read_isv_signal_at(BUDGET_CQL_BASE + b))
+            .sum::<f32>()
+            / 4.0_f32)
+            .max(0.02_f32);
+        let ens_budget = ((0..4_usize)
+            .map(|b| self.read_isv_signal_at(BUDGET_ENS_BASE + b))
+            .sum::<f32>()
+            / 4.0_f32)
+            .max(0.02_f32);
 
         self.trainer.last_iqn_budget_eff = iqn_budget;
         self.trainer.last_cql_budget_eff = cql_budget;
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 1ad926aee..c087b9993 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -26,6 +26,12 @@ use crate::cuda_pipeline::gpu_dqn_trainer::{
     EPOCH_IDX_INDEX, EPSILON_EFF_INDEX, GAMMA_DIR_EFF_INDEX, Q_ABS_REF_INDEX,
     Q_DIR_ABS_REF_INDEX, TLOB_REGIME_FOCUS_EMA_INDEX,
 };
+use crate::cuda_pipeline::sp5_isv_slots::{
+    NOISY_SIGMA_BASE,
+    ADAM_BETA1_BASE as SP5_ADAM_BETA1_BASE,
+    ADAM_BETA2_BASE as SP5_ADAM_BETA2_BASE,
+    ADAM_EPS_BASE   as SP5_ADAM_EPS_BASE,
+};
 use crate::dqn::logging::{log_epoch_start, log_epoch_end, log_training_progress};
 use crate::dqn::target_update::convergence_half_life;
 use crate::evaluation::metrics::calculate_var_cvar;
@@ -1734,7 +1740,17 @@ impl DQNTrainer {
             max_trace_length: self.hyperparams.max_trace_length as i32,
             hindsight_fraction: self.hyperparams.hindsight_fraction as f32,
             hindsight_lookahead: self.hyperparams.hindsight_lookahead as i32,
-            noise_sigma: self.hyperparams.noise_sigma as f32,
+            // SP5 Layer B (Pearl 3): read per-branch NoisyNet σ from ISV[210..214)
+            // and collapse to a single effective scale by averaging across 4 branches.
+            // Falls back to the static hyperparam when fused_ctx is unavailable
+            // (pre-initialization path).  0.01f floor = Invariant 1 carve-out.
+            noise_sigma: self.fused_ctx.as_ref().map(|ctx| {
+                let mean = (0..4_usize)
+                    .map(|b| ctx.read_isv_signal_at(NOISY_SIGMA_BASE + b))
+                    .sum::<f32>()
+                    / 4.0_f32;
+                mean.max(0.01_f32)
+            }).unwrap_or(self.hyperparams.noise_sigma as f32),
             ..Default::default()
         };
 
@@ -1793,9 +1809,24 @@ impl DQNTrainer {
                 // unbounded clamps.
                 0.0_f32
             };
+            // SP5 Pearl 4: per-group Adam β1/β2/ε for Curiosity (ParamGroup
+            // index 7). ISV[ADAM_BETA1_BASE+7=233], ISV[ADAM_BETA2_BASE+7=241],
+            // ISV[ADAM_EPS_BASE+7=249]. Curiosity collector doesn't own ISV —
+            // read here and thread through to `train_on_collector_buffers`.
+            let cur_g = crate::cuda_pipeline::ParamGroup::Curiosity.idx();
+            let (cur_beta1, cur_beta2, cur_epsilon) = if let Some(ref fused) = self.fused_ctx {
+                (
+                    fused.read_isv_signal_at(SP5_ADAM_BETA1_BASE + cur_g).max(0.5_f32).min(0.9999_f32),
+                    fused.read_isv_signal_at(SP5_ADAM_BETA2_BASE + cur_g).max(0.9_f32).min(0.99999_f32),
+                    fused.read_isv_signal_at(SP5_ADAM_EPS_BASE   + cur_g).max(1e-12_f32),
+                )
+            } else {
+                (0.9_f32, 0.999_f32, 1e-8_f32)
+            };
             if let Err(e) = collector.train_curiosity_gpu(
                 gpu_batch.n_episodes, gpu_batch.timesteps,
                 weight_clamp_max_abs,
+                cur_beta1, cur_beta2, cur_epsilon,
             ) {
                 debug!("GPU curiosity training failed (non-fatal): {e}");
             }
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 8fa188090..5bb912b1f 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -3190,3 +3190,42 @@ Refs: SP4 Layer C close-out follow-up — sites #2 + #9 of the sweep audit grid
 - 2 unit tests pass (slot layout invariants + cross-fold-persistent Kelly carve-out)
 
 No behavior change. No producer/consumer wired yet. Subsequent Layer A commits (A1–A8) populate slots via per-pearl producer kernels.
+
+### SP5 Layer B — atomic consumer migration to ISV-driven adaptive signals
+
+All 11 consumer sites migrated in one atomic commit per `feedback_no_partial_refactor`. Every site that reads from ISV slots 174..286 now does so at runtime via `read_isv_signal_at` (Rust, CPU-side mapped-pinned read) or direct ISV pointer dereference (CUDA kernel). No static constants remain for any of the 8 pearls covered.
+
+**Pearl 1 (C51 atom span — v_center/v_half):**
+- `atoms_update_kernel.cu`: reads `ISV[174+b]` (v_center) and `ISV[178+b]` (v_half, floor 1.0) per branch.
+
+**Pearl 1-ext (per-branch num_atoms):**
+- `atoms_update_kernel.cu`: reads `ISV[274+b]`, calls new `round_to_valid_num_atoms()` device inline to snap to valid {16,32,64} with cold-start fallback to global `num_atoms`. Uses `eff_na` for all softmax/cumsum iteration; output stride remains global `num_atoms` (buffer invariant).
+
+**Pearl 2 (per-branch loss budgets):**
+- `fused_training.rs` `compute_adaptive_budgets()`: reads ISV[190..194), ISV[194..198), ISV[198..202), ISV[202..206) per branch, averages 4 branches per loss type. Invariant-1 floors: C51=0.05, IQN=0.05, CQL=0.02, Ens=0.02.
+
+**Pearl 3 (per-branch NoisyNet sigma):**
+- `training_loop.rs`: `ExperienceCollectorConfig.noise_sigma` reads `mean(ISV[210..214])` with 0.01 floor; falls back to `hyperparams.noise_sigma` when `fused_ctx` absent.
+
+**Pearl 4 (per-group Adam β1/β2/ε — 5 sites):**
+- `gpu_dqn_trainer.rs` `launch_adam_update`: reads ISV[226+g], ISV[234+g], ISV[242+g] per group `g` inside the existing per-group loop. Static `ADAM_BETA1/BETA2/EPS` remain declared but are unused (dead code left for reference; may be removed in cleanup).
+- `gpu_iqn_head.rs` `execute_training_pipeline` + `train_iqn_step_gpu`: 3 new `iqn_beta1/beta2/epsilon` parameters threaded from callers.
+- `gpu_attention.rs` `adam_step`: 3 new `attn_beta1/beta2/epsilon` parameters replacing hardcoded 0.9/0.999/1e-8.
+- `gpu_tlob.rs` `adam_step`: same pattern as GpuAttention (shares ParamGroup::Attn=6).
+- `fused_training.rs`: 4 call sites (2 IQN, 1 GpuAttention, 1 TLOB) read ISV[226+g..242+g] and thread values through.
+- `gpu_curiosity_trainer.rs` `launch_adam_step` + `train_on_collector_buffers`: `cur_beta1/beta2/epsilon` params replace constant references.
+- `gpu_experience_collector.rs` `train_curiosity_gpu`: threads `cur_beta1/beta2/epsilon`.
+- `training_loop.rs`: reads ISV[233/241/249] for `ParamGroup::Curiosity=7` at the `train_curiosity_gpu` call site.
+
+**Pearl 5 (per-branch IQN τ schedule):**
+- `gpu_iqn_head.rs` new `refresh_taus_from_isv(&[f32; 20])`: averages 4 branches per quantile slot (ISV[250+b*5+q]), applies cold-start floor from `FIXED_TAUS[q]` (ISV=0 → retain original), recomputes `cos_features [D,N]` from updated taus, re-uploads `online_taus`, `target_taus`, `cos_features` in-place via `upload_f32_via_pinned`.
+- `fused_training.rs`: reads ISV[250..270) into `[f32; 20]` array, calls `iqn.refresh_taus_from_isv(&isv_taus)` before `prepare_buffers` each training step.
+
+**Pearl 6 (cross-fold Kelly cap) + Pearl 8 (per-direction trail distance):**
+- `trade_physics.cuh`, `experience_kernels.cu`, `backtest_env_kernel.cu`: done in prior session.
+
+**Compilation and tests:**
+- `SQLX_OFFLINE=true cargo check -p ml --offline` → clean (12 warnings, all pre-existing)
+- `cargo test -p ml --lib` → 930 pass, 14 pre-existing failures unchanged
+
+Refs: `feedback_no_partial_refactor`, `feedback_isv_for_adaptive_bounds`, `feedback_adaptive_not_tuned`