diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 3a75d54ce..f3ecb024c 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -430,6 +430,18 @@ fn main() {
         // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[190..210).
         // Must run AFTER pearl_1_atom (Q_VAR) AND pearl_3_sigma (NOISY_SIGMA).
         "pearl_2_budget_kernel.cu",
+        // SP5 Task A4 (2026-05-01): per-group gradient cosine similarity + L2 norm.
+        // Single-block 8-thread kernel (one per param group: DqnTrunk/Value/Branches/
+        // IQN/IqlHigh/IqlLow/Attn/Curiosity). Reads grad_buf + grad_prev_buf;
+        // writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
+        // Also writebacks grad_curr → grad_prev_buf for the next step's comparison.
+        "grad_cosine_sim_kernel.cu",
+        // SP5 Task A4 (2026-05-01): Pearl 4 per-param-group Adam β1, β2, ε.
+        // Single-block 8-thread kernel reads cosine_sim + l2_norm from scratch;
+        // writes β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
+        // apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=5e-4)
+        // into ISV[226..250). Must run AFTER grad_cosine_sim_kernel.
+        "pearl_4_adam_hparams_kernel.cu",
         // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread
         // device-side loop applies Pearls A+D to N consecutive ISV slots
         // (each with its own Wiener-state triple) on the producer's
diff --git a/crates/ml/src/cuda_pipeline/apply_pearls_kernel.cu b/crates/ml/src/cuda_pipeline/apply_pearls_kernel.cu
index 91ac7498d..b5a43eeca 100644
--- a/crates/ml/src/cuda_pipeline/apply_pearls_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/apply_pearls_kernel.cu
@@ -20,8 +20,13 @@
 // pattern broke graph capture at `aux_label_scale_ema` mid-step
 // (CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED — smoke `smoke-test-v9kjv`).
 //
-// Constants (`ALPHA_META`, `EPS_DIV`) MUST match
-// `cuda_pipeline::sp4_wiener_ema::{ALPHA_META, EPS_DIV}` bit-for-bit.
+// `ALPHA_META` is now passed as a kernel parameter so Pearl 4 can use
+// a halved rate (5e-4) while all SP4 and other SP5 producers continue
+// to use the SP4 default (1e-3). `EPS_DIV` remains a compile-time
+// constant — it is numerical-stability-only and must never vary.
+//
+// Constants (`EPS_DIV`) MUST match
+// `cuda_pipeline::sp4_wiener_ema::EPS_DIV` bit-for-bit.
 // Single-source-of-truth lives in `sp4_wiener_ema.rs` for documentation;
 // this duplicate is required because CUDA C++ can't include the Rust
 // const directly. The new GPU oracle test asserts agreement within
@@ -34,14 +39,15 @@ extern "C" __global__ void apply_pearls_ad_kernel(
     int          isv_idx_base,
     float*       __restrict__ wiener_state_buf, /* [wiener_offset_base + 3*n_slots) packed [sv, dv, xl] triples */
     int          wiener_offset_base,
-    int          n_slots
+    int          n_slots,
+    float        alpha_meta               /* meta-EMA rate for Wiener variance tracking;
+                                             SP4 default = 1e-3; Pearl 4 uses 5e-4 */
 ) {
     // Single-thread device-side loop. Every slot has its own state; no
     // cross-slot parallelism is meaningful. Guard avoids redundant work
     // if the launcher accidentally requests >1 thread.
     if (blockIdx.x != 0 || threadIdx.x != 0) return;
 
-    constexpr float ALPHA_META = 1.0e-3f;
     constexpr float EPS_DIV    = 1.0e-8f;
 
     for (int i = 0; i < n_slots; ++i) {
@@ -85,9 +91,9 @@ extern "C" __global__ void apply_pearls_ad_kernel(
         const float x_lag_diff  = step_obs - x_lag;
 
         const float new_sample_var =
-            (1.0f - ALPHA_META) * sample_var + ALPHA_META * sample_diff * sample_diff;
+            (1.0f - alpha_meta) * sample_var + alpha_meta * sample_diff * sample_diff;
         const float new_diff_var =
-            (1.0f - ALPHA_META) * diff_var   + ALPHA_META * x_lag_diff  * x_lag_diff;
+            (1.0f - alpha_meta) * diff_var   + alpha_meta * x_lag_diff  * x_lag_diff;
 
         // EPS_DIV protects 0/0 at t=1 when both variances start near 0.
         const float alpha_opt =
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 06f38818f..348d48d39 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -318,6 +318,24 @@ static SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
 static SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));
 
+/// SP5 Task A4 (2026-05-01): auxiliary gradient cosine similarity cubin.
+/// Single-block 8-thread kernel (one thread per SP4 param group). Reads
+/// grad_buf + grad_prev_buf_per_group + group_param_offsets_buf [9 i32];
+/// computes per-group cos_sim and l2_norm; writes to scratch_buf[131..147).
+/// Also writes grad_curr → grad_prev_buf_writeback for next step's comparison.
+/// No atomicAdd (feedback_no_atomicadd).
+static SP5_GRAD_COSINE_SIM_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/grad_cosine_sim_kernel.cubin"));
+
+/// SP5 Task A4 (2026-05-01): Pearl 4 per-group Adam β1/β2/ε cubin.
+/// Single-block 8-thread kernel (one thread per param group). Reads
+/// scratch_buf[131..147) (cosine_sim + l2_norm); computes β1/β2/ε within
+/// structural envelopes (Invariant 1 anchors); writes to scratch_buf[147..171).
+/// Chained apply_pearls_ad_kernel (24 launches, ALPHA_META=5e-4) smooths into
+/// ISV[ADAM_BETA1_BASE..ADAM_EPS_BASE+8). No atomicAdd.
+static SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_kernel.cubin"));
+
 /// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103).
 /// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched
 /// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN
@@ -901,7 +919,7 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
     (SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT)
     * SP4_WIENER_FLOATS_PER_SLOT;
 
-/// SP5 Task A1+A2+A3: producer_step_scratch_buf total slot count after growth.
+/// SP5 Task A1+A2+A3+A4: producer_step_scratch_buf total slot count after growth.
 /// SP4 had 71 scratch slots [0..71); Task A1 adds:
 ///   16 floats for q_branch_stats output (SCRATCH_Q_STATS=71, 4 outputs × 4 branches)
 ///   16 floats for pearl_1_atom output  (SCRATCH_ATOM_OUT=87, 4 outputs × 4 branches)
@@ -915,8 +933,14 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
 ///     budget_cql[4]  → SCRATCH_PEARL_2_CQL=119
 ///     budget_ens[4]  → SCRATCH_PEARL_2_ENS=123
 ///     flatness[4]    → SCRATCH_PEARL_2_FLATNESS=127
-/// Combined: 71 + 16 + 16 + 4 + 4 + 20 = 131 scratch slots [0..131).
-pub const SP5_SCRATCH_TOTAL: usize = 131;
+/// Task A4 adds:
+///    8 floats for grad_cosine_sim output  (SCRATCH_PEARL_4_COSINE=131, cos_sim × 8 groups)
+///    8 floats for grad_l2_norm output     (SCRATCH_PEARL_4_L2=139, l2_norm × 8 groups)
+///    8 floats for pearl_4_beta1 output    (SCRATCH_PEARL_4_BETA1=147, β1 × 8 groups)
+///    8 floats for pearl_4_beta2 output    (SCRATCH_PEARL_4_BETA2=155, β2 × 8 groups)
+///    8 floats for pearl_4_eps output      (SCRATCH_PEARL_4_EPS=163, ε × 8 groups)
+/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 = 171 scratch slots [0..171).
+pub const SP5_SCRATCH_TOTAL: usize = 171;
 
 /// SP5 Task A2: scratch index base for pearl_3_sigma sigma[4] output block.
 /// Slots [103..107): per-branch NoisyNet sigma.
@@ -947,6 +971,34 @@ pub const SCRATCH_PEARL_2_ENS: usize = 123;
 /// Slots [127..131): per-branch flatness signal (var_q / (σ² + EPS_DIV), clamped [0,1]).
 pub const SCRATCH_PEARL_2_FLATNESS: usize = 127;
 
+/// SP5 Task A4: scratch index base for grad_cosine_sim_kernel cosine_sim[8] output block.
+/// Slots [131..139): per-param-group gradient direction cosine similarity vs previous step.
+/// Written by `grad_cosine_sim_update`; consumed by `pearl_4_adam_hparams_update`.
+pub const SCRATCH_PEARL_4_COSINE: usize = 131;
+
+/// SP5 Task A4: scratch index base for grad_cosine_sim_kernel l2_norm[8] output block.
+/// Slots [139..147): per-param-group L2 norm of current-step gradient.
+/// Written by `grad_cosine_sim_update`; consumed by `pearl_4_adam_hparams_update`.
+pub const SCRATCH_PEARL_4_L2: usize = 139;
+
+/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel β1[8] output block.
+/// Slots [147..155): per-param-group adaptive Adam β1 ∈ [0.85, 0.95].
+/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
+/// ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8).
+pub const SCRATCH_PEARL_4_BETA1: usize = 147;
+
+/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel β2[8] output block.
+/// Slots [155..163): per-param-group adaptive Adam β2 ∈ [0.99, 0.9995].
+/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
+/// ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8).
+pub const SCRATCH_PEARL_4_BETA2: usize = 155;
+
+/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel ε[8] output block.
+/// Slots [163..171): per-param-group adaptive Adam ε ∈ [1e-10, 1e-6].
+/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
+/// ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8).
+pub const SCRATCH_PEARL_4_EPS: usize = 163;
+
 /// SP5 Task A1: scratch index base for q_branch_stats 16-float output block.
 /// Slots [71..87): mean, max_abs_dev, var, entropy × 4 branches.
 /// SP4 slot 70 = q_dir_grad_p99; SP5 starts at 71 (no collision).
@@ -4095,6 +4147,40 @@ pub struct GpuDqnTrainer {
     /// No atomicAdd (feedback_no_atomicadd), no CPU compute (feedback_no_cpu_compute_strict).
     /// Loaded from `pearl_2_budget_kernel.cubin`.
     pearl_2_budget_kernel: CudaFunction,
+    // ── SP5 Task A4: Pearl 4 per-group Adam β1/β2/ε kernels ─────────────
+    /// SP5 Task A4 (2026-05-01): auxiliary gradient cosine similarity kernel.
+    /// Single-block 8-thread kernel (one thread per SP4 param group).
+    /// Reads `grad_buf [total_params]` + `grad_prev_buf_per_group [total_params]`
+    /// + `group_param_offsets_buf [9]` (prefix sums in element units); computes
+    /// per-group cos_sim = dot(g,g_prev)/(|g|×|g_prev|+EPS_COS) and l2_norm = |g|;
+    /// writes to `producer_step_scratch_buf[131..147)`. Also writebacks g_curr →
+    /// grad_prev_buf_per_group for the next step.
+    /// No atomicAdd (feedback_no_atomicadd), no CPU compute (feedback_no_cpu_compute_strict).
+    /// Loaded from `grad_cosine_sim_kernel.cubin`.
+    grad_cosine_sim_kernel: CudaFunction,
+    /// SP5 Task A4 (2026-05-01): Pearl 4 per-group Adam β1/β2/ε controller kernel.
+    /// Single-block 8-thread kernel (one thread per param group). Reads per-group
+    /// cosine_sim + l2_norm from scratch_buf[131..147); computes
+    /// β1=0.85+0.10×stability ∈ [0.85,0.95],
+    /// β2=0.99+0.0095×stability ∈ [0.99,0.9995],
+    /// ε=clamp(grad_norm×1e-7, 1e-10, 1e-6); writes to scratch_buf[147..171).
+    /// Chained apply_pearls_ad_kernel (24 single-slot launches, ALPHA_META=5e-4)
+    /// smooths via Pearls A+D into ISV[ADAM_BETA1_BASE..ADAM_EPS_BASE+8).
+    /// Structural envelopes (Invariant 1 anchors per spec line 88).
+    /// No atomicAdd, no CPU compute.
+    /// Loaded from `pearl_4_adam_hparams_kernel.cubin`.
+    pearl_4_adam_hparams_kernel: CudaFunction,
+    /// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity.
+    /// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned).
+    /// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0
+    /// on first step → β1/β2 start at envelope midpoints).
+    pub(crate) grad_prev_buf_per_group: MappedF32Buffer,
+    /// SP5 Task A4 (2026-05-01): param-group boundary prefix sums for cosine-sim kernel.
+    /// [9 i32] with group g in element range [offsets[g], offsets[g+1]).
+    /// Values are in f32-element units (not bytes): divide padded byte offsets by 4.
+    /// Host-written once at construction; static for the lifetime of the trainer.
+    pub(crate) group_param_offsets_buf: MappedI32Buffer,
+
     /// SP5 Task A1 (2026-05-01): per-branch clip-count buffer for Pearl 1
     /// headroom controller. [4 i32], one element per branch. Zero-initialized
     /// at construction; populated by `atoms_update_kernel` in Layer B.
@@ -9491,6 +9577,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 (SCRATCH_IDX * 3) as i32,
                 1,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -9575,6 +9662,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 wiener_offset,
                 1,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -9734,6 +9822,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 ((ATOM_POS_BOUND_BASE - SP4_SLOT_BASE) * 3) as i32,
                 SP4_BRANCH_COUNT as i32,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -9976,6 +10065,7 @@ impl GpuDqnTrainer {
                     self.wiener_state_buf.dev_ptr,
                     ((isv_idx - SP4_SLOT_BASE) * 3) as i32,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )
             }
         };
@@ -10169,6 +10259,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 ((isv_idx - SP4_SLOT_BASE) * 3) as i32,
                 1,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -10384,6 +10475,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 ((isv_idx - SP4_SLOT_BASE) * 3) as i32,
                 1,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -10524,6 +10616,7 @@ impl GpuDqnTrainer {
                         isv_dev,     isv_idx,
                         wiener_dev,  wiener_off,
                         1,
+                        crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                     )?;
                 }
             }
@@ -10545,6 +10638,7 @@ impl GpuDqnTrainer {
                     isv_dev,     isv_entropy,
                     wiener_dev,  wiener_entropy,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )?;
             }
             // q_var → ISV Q_VAR_PER_BRANCH_BASE + b
@@ -10559,6 +10653,7 @@ impl GpuDqnTrainer {
                     isv_dev,     isv_var,
                     wiener_dev,  wiener_var,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )?;
             }
         }
@@ -10643,6 +10738,7 @@ impl GpuDqnTrainer {
                     isv_dev,     isv_sigma,
                     wiener_dev,  wiener_sigma,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )?;
             }
             // SF → ISV SIGMA_FRACTION_BASE + b
@@ -10657,6 +10753,7 @@ impl GpuDqnTrainer {
                     isv_dev,     isv_sf,
                     wiener_dev,  wiener_sf,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )?;
             }
         }
@@ -10756,6 +10853,7 @@ impl GpuDqnTrainer {
                         isv_dev,     isv_idx,
                         wiener_dev,  wiener_off,
                         1,
+                        crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                     )?;
                 }
             }
@@ -10764,6 +10862,156 @@ impl GpuDqnTrainer {
         Ok(())
     }
 
+    /// SP5 Task A4 (2026-05-01): launch the two-kernel Pearl 4 chain +
+    /// 24 `apply_pearls_ad_kernel` launches to smooth β1/β2/ε into ISV.
+    ///
+    /// Execution sequence (all on `self.stream`):
+    ///   1. `grad_cosine_sim_update` (8 threads, 1 block): reads
+    ///      `grad_buf [total_params]` + `grad_prev_buf_per_group [total_params]`
+    ///      + `group_param_offsets_buf [9]`; writes per-group cosine_sim[8]
+    ///      to scratch[131..139) and l2_norm[8] to scratch[139..147); also
+    ///      writes grad_curr → grad_prev_buf_per_group for the next step's comparison.
+    ///   2. `pearl_4_adam_hparams_update` (8 threads, 1 block): reads
+    ///      scratch[131..147); writes β1[8]→scratch[147..155),
+    ///      β2[8]→scratch[155..163), ε[8]→scratch[163..171).
+    ///   3. 24 `apply_pearls_ad_kernel` launches (ALPHA_META=5e-4, half the
+    ///      SP4 default): smooths β1[8]/β2[8]/ε[8] via Pearls A+D into
+    ///      ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8),
+    ///      ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8),
+    ///      ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8).
+    ///
+    /// Pearl 4 is independent of Pearls 1/2/3 outputs — depends only on
+    /// grad_buf state. Can run in any order relative to other SP5 launchers.
+    ///
+    /// Structural envelopes (Invariant 1 anchors per spec line 88):
+    ///   β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6].
+    /// Fold-boundary reset zeroes grad_prev_buf_per_group (Pearl A sentinel):
+    /// first step cosine_sim=0 → β1=0.85, β2=0.99 (low-stability envelope).
+    /// Layer A only — Adam optimizer consumer migration deferred to Layer B.
+    pub(crate) fn launch_sp5_pearl_4_adam_hparams(&self) -> Result<(), MLError> {
+        use crate::cuda_pipeline::sp5_isv_slots::{
+            ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE, SP5_SLOT_BASE,
+        };
+        use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
+
+        const ALPHA_META_PEARL_4: f32 = 5.0e-4_f32; // half SP4 default → slow β change rate
+
+        let scratch_dev = self.producer_step_scratch_buf.dev_ptr;
+        let isv_dev     = self.isv_signals_dev_ptr;
+        let wiener_dev  = self.wiener_state_buf.dev_ptr;
+        // Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 Wiener slots
+        // already consumed; SP5 slots start at offset (SP4_PRODUCER_COUNT * 3).
+        let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32;
+
+        let grad_buf_ptr        = self.ptrs.grad_buf;
+        let grad_prev_ptr       = self.grad_prev_buf_per_group.dev_ptr;
+        let group_offsets_ptr   = self.group_param_offsets_buf.dev_ptr;
+        let cosine_idx_base     = SCRATCH_PEARL_4_COSINE as i32;  // 131
+        let l2_norm_idx_base    = SCRATCH_PEARL_4_L2     as i32;  // 139
+        let beta1_idx_base      = SCRATCH_PEARL_4_BETA1  as i32;  // 147
+        let beta2_idx_base      = SCRATCH_PEARL_4_BETA2  as i32;  // 155
+        let eps_idx_base        = SCRATCH_PEARL_4_EPS    as i32;  // 163
+
+        // Step 1: grad_cosine_sim_update — 8 threads, 1 block.
+        // Reads grad_buf + grad_prev_buf_per_group + group_param_offsets;
+        // writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
+        // Also writebacks grad_curr → grad_prev_buf_per_group.
+        let cfg_8t = LaunchConfig {
+            grid_dim:         (1, 1, 1),
+            block_dim:        (8, 1, 1),
+            shared_mem_bytes: 0,
+        };
+        unsafe {
+            self.stream
+                .launch_builder(&self.grad_cosine_sim_kernel)
+                .arg(&grad_buf_ptr)
+                .arg(&grad_prev_ptr)
+                .arg(&group_offsets_ptr)
+                .arg(&scratch_dev)
+                .arg(&cosine_idx_base)
+                .arg(&l2_norm_idx_base)
+                .arg(&grad_prev_ptr)  // writeback = same buffer
+                .launch(cfg_8t)
+                .map_err(|e| MLError::ModelError(format!("grad_cosine_sim_update launch: {e}")))?;
+        }
+
+        // Step 2: pearl_4_adam_hparams_update — 8 threads, 1 block.
+        // Reads scratch[131..147) (cosine_sim + l2_norm); writes
+        // β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
+        // Stream-ordered with Step 1 (same stream, no explicit sync needed).
+        let cosine_ptr = scratch_dev + (SCRATCH_PEARL_4_COSINE * std::mem::size_of::<f32>()) as u64;
+        let l2_ptr     = scratch_dev + (SCRATCH_PEARL_4_L2     * std::mem::size_of::<f32>()) as u64;
+        unsafe {
+            self.stream
+                .launch_builder(&self.pearl_4_adam_hparams_kernel)
+                .arg(&cosine_ptr)
+                .arg(&l2_ptr)
+                .arg(&scratch_dev)
+                .arg(&beta1_idx_base)
+                .arg(&beta2_idx_base)
+                .arg(&eps_idx_base)
+                .launch(cfg_8t)
+                .map_err(|e| MLError::ModelError(format!("pearl_4_adam_hparams_update launch: {e}")))?;
+        }
+
+        // Step 3: 24 apply_pearls_ad_kernel launches (ALPHA_META=5e-4).
+        // β1[8]: scratch[147..155) → ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8)
+        for g in 0..8_usize {
+            let scratch_idx = (SCRATCH_PEARL_4_BETA1 + g) as i32;
+            let isv_idx     = (ADAM_BETA1_BASE + g) as i32;
+            let wiener_off  = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
+            unsafe {
+                launch_apply_pearls(
+                    &self.stream,
+                    &self.apply_pearls_ad_kernel,
+                    scratch_dev, scratch_idx,
+                    isv_dev,     isv_idx,
+                    wiener_dev,  wiener_off,
+                    1,
+                    ALPHA_META_PEARL_4,
+                )?;
+            }
+        }
+
+        // β2[8]: scratch[155..163) → ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8)
+        for g in 0..8_usize {
+            let scratch_idx = (SCRATCH_PEARL_4_BETA2 + g) as i32;
+            let isv_idx     = (ADAM_BETA2_BASE + g) as i32;
+            let wiener_off  = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
+            unsafe {
+                launch_apply_pearls(
+                    &self.stream,
+                    &self.apply_pearls_ad_kernel,
+                    scratch_dev, scratch_idx,
+                    isv_dev,     isv_idx,
+                    wiener_dev,  wiener_off,
+                    1,
+                    ALPHA_META_PEARL_4,
+                )?;
+            }
+        }
+
+        // ε[8]: scratch[163..171) → ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8)
+        for g in 0..8_usize {
+            let scratch_idx = (SCRATCH_PEARL_4_EPS + g) as i32;
+            let isv_idx     = (ADAM_EPS_BASE + g) as i32;
+            let wiener_off  = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
+            unsafe {
+                launch_apply_pearls(
+                    &self.stream,
+                    &self.apply_pearls_ad_kernel,
+                    scratch_dev, scratch_idx,
+                    isv_dev,     isv_idx,
+                    wiener_dev,  wiener_off,
+                    1,
+                    ALPHA_META_PEARL_4,
+                )?;
+            }
+        }
+
+        Ok(())
+    }
+
     /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update`
     /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129].
     ///
@@ -10889,6 +11137,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 (SCRATCH_FIRST * 3) as i32,
                 num_groups as i32,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -10965,6 +11214,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 (SCRATCH_IDX_NEXT_BAR * 3) as i32,
                 2,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -11231,6 +11481,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 (SCRATCH_IDX_UTIL_BASE * 3) as i32,
                 (MOE_NUM_EXPERTS + 1) as i32,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
 
@@ -11555,6 +11806,7 @@ impl GpuDqnTrainer {
                     self.wiener_state_buf.dev_ptr,
                     (SCRATCH_IDX_LABEL_SCALE * 3) as i32,
                     1,
+                    crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
                 )?;
             }
         }
@@ -11869,6 +12121,7 @@ impl GpuDqnTrainer {
                 self.wiener_state_buf.dev_ptr,
                 (SCRATCH_FIRST * 3) as i32,
                 4,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
         Ok(())
@@ -12801,6 +13054,66 @@ impl GpuDqnTrainer {
             .map_err(|e| MLError::ModelError(format!("SP5 action_offsets_buf alloc: {e}")))?;
         action_offsets_buf.write_from_slice(&[0_i32, 4, 7, 10]);
 
+        // SP5 Task A4 (2026-05-01): load grad_cosine_sim_kernel (auxiliary
+        // gradient cosine similarity). Single-block 8-thread, one thread per
+        // SP4 param group. Reads grad_buf + grad_prev_buf_per_group.
+        let grad_cosine_sim_kernel = {
+            let module = stream.context().load_cubin(SP5_GRAD_COSINE_SIM_CUBIN.to_vec())
+                .map_err(|e| MLError::ModelError(format!("sp5 grad_cosine_sim cubin load: {e}")))?;
+            module.load_function("grad_cosine_sim_update")
+                .map_err(|e| MLError::ModelError(format!("grad_cosine_sim_update load: {e}")))?
+        };
+
+        // SP5 Task A4 (2026-05-01): load pearl_4_adam_hparams_kernel.
+        // Single-block 8-thread, one thread per SP4 param group. Reads
+        // cosine_sim + l2_norm from scratch; writes β1/β2/ε to scratch.
+        let pearl_4_adam_hparams_kernel = {
+            let module = stream.context().load_cubin(SP5_PEARL_4_ADAM_HPARAMS_CUBIN.to_vec())
+                .map_err(|e| MLError::ModelError(format!("sp5 pearl_4_adam_hparams cubin load: {e}")))?;
+            module.load_function("pearl_4_adam_hparams_update")
+                .map_err(|e| MLError::ModelError(format!("pearl_4_adam_hparams_update load: {e}")))?
+        };
+
+        // SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32].
+        // Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel:
+        // zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints).
+        // `total_params` was computed at constructor entry (line ~12337).
+        let grad_prev_buf_per_group = unsafe { MappedF32Buffer::new(total_params) }
+            .map_err(|e| MLError::ModelError(format!("SP5 grad_prev_buf_per_group alloc: {e}")))?;
+
+        // SP5 Task A4 (2026-05-01): allocate group_param_offsets_buf [9 i32].
+        // Prefix sums in f32-element units (not bytes): [offsets[0], ..., offsets[8]]
+        // where group g covers [offsets[g], offsets[g+1]).
+        // Groups match SP4_PARAM_GROUP_COUNT=8:
+        //   0=DqnTrunk [0..trunk_param_count)
+        //   1=DqnValue [trunk_param_count..trunk+value_count)
+        //   2=DqnBranches [trunk+value..trunk+value+branch_count)
+        //   3=Iqn, 4=IqlHigh, 5=IqlLow, 6=Attn, 7=Curiosity
+        // For aux groups 3-7: they are separate grad buffers (not slices of grad_buf)
+        // so their offsets are set to [total_params, total_params] — zero-width range
+        // signals the kernel to skip those groups (thread loop produces dot=0,
+        // norm=0, cosine_sim=0 → β1/β2 snap to stable envelope midpoints).
+        // Only DqnTrunk/Value/Branches (groups 0-2) cover the contiguous main grad_buf.
+        let param_sizes = compute_param_sizes(&config);
+        let trunk_end   = (padded_byte_offset(&param_sizes, 13) / 4) as i32;
+        let value_end   = (padded_byte_offset(&param_sizes, 17) / 4) as i32;
+        let branch_end  = (padded_byte_offset(&param_sizes, 33) / 4) as i32;
+        // Aux groups 3-7 sit outside the contiguous grad_buf slab → zero-width sentinel.
+        let tp = total_params as i32;
+        let group_param_offsets_buf = unsafe { MappedI32Buffer::new(9) }
+            .map_err(|e| MLError::ModelError(format!("SP5 group_param_offsets_buf alloc: {e}")))?;
+        group_param_offsets_buf.write_from_slice(&[
+            0,          // group 0 DqnTrunk start
+            trunk_end,  // group 1 DqnValue start / DqnTrunk end
+            value_end,  // group 2 DqnBranches start / DqnValue end
+            branch_end, // group 3 Iqn start / DqnBranches end
+            tp,         // group 4 IqlHigh start — zero-width sentinel for aux groups
+            tp,         // group 5 IqlLow start
+            tp,         // group 6 Attn start
+            tp,         // group 7 Curiosity start
+            tp,         // sentinel: offsets[8] = end of group 7
+        ]);
+
         // Plan C Phase 2 follow-up A.2: load q_drift_rate_ema kernel
         // (cold-path, per-epoch). Single-thread single-block ISV producer
         // for ISV[Q_DRIFT_RATE_INDEX=129]; consumer is `tau_update_kernel`.
@@ -15906,6 +16219,10 @@ impl GpuDqnTrainer {
             pearl_1_atom_kernel,
             pearl_3_sigma_kernel,
             pearl_2_budget_kernel,
+            grad_cosine_sim_kernel,
+            pearl_4_adam_hparams_kernel,
+            grad_prev_buf_per_group,
+            group_param_offsets_buf,
             atoms_clip_count_buf,
             action_counts_buf,
             action_offsets_buf,
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 651079d13..30d07ddc9 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -2399,6 +2399,7 @@ impl GpuExperienceCollector {
                 self.reward_component_pearls_wiener_dev_ptr,
                 (SCRATCH_FIRST * 3) as i32,
                 REWARD_COMPONENT_COUNT as i32,
+                crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
             )?;
         }
 
diff --git a/crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu b/crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu
new file mode 100644
index 000000000..2702366ea
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu
@@ -0,0 +1,62 @@
+// crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu
+//
+// SP5 Task A4 auxiliary: per-group gradient cosine similarity.
+// Computes cos_sim(grad_curr, grad_prev) for each of 8 SP4 param groups.
+// Also computes per-group L2 norm of grad_curr.
+// Writes grad_curr direction as grad_prev for the next step.
+//
+// Single-block, 8-thread (one per group). Each thread reduces over its
+// group's parameter range to compute:
+//   cos_sim = (g_curr · g_prev) / (|g_curr| × |g_prev| + EPS_COS)
+//   l2_norm = |g_curr|
+//
+// No atomicAdd per feedback_no_atomicadd — one thread per group with
+// independent writes. grad_prev_buf_writeback is written by the same
+// thread that reads it (no race).
+
+#include <cuda_runtime.h>
+
+extern "C" __global__ void grad_cosine_sim_update(
+    const float* __restrict__ grad_buf,               // [total_params] current step's gradient
+    const float* __restrict__ grad_prev_buf,          // [total_params] previous step's gradient
+    const int*   __restrict__ group_param_offsets,    // [9] prefix sums: group g in [offsets[g], offsets[g+1])
+    float* __restrict__ scratch_out,                  // cosine_sim[8] at [cosine_idx_base..+8)
+                                                      // l2_norm[8] at [l2_norm_idx_base..+8)
+    int   cosine_idx_base,
+    int   l2_norm_idx_base,
+    float* __restrict__ grad_prev_buf_writeback       // [total_params] overwrite prev with curr for next step
+) {
+    int g = threadIdx.x;
+    if (blockIdx.x != 0 || g >= 8) return;
+
+    int start = group_param_offsets[g];
+    int end   = group_param_offsets[g + 1];
+
+    const float EPS_COS = 1e-12f;
+    float dot = 0.0f, norm_curr_sq = 0.0f, norm_prev_sq = 0.0f;
+
+    for (int i = start; i < end; ++i) {
+        float gc = grad_buf[i];
+        float gp = grad_prev_buf[i];
+        dot          += gc * gp;
+        norm_curr_sq += gc * gc;
+        norm_prev_sq += gp * gp;
+    }
+
+    float norm_curr = sqrtf(norm_curr_sq);
+    float norm_prev = sqrtf(norm_prev_sq);
+    float cos_sim   = dot / (norm_curr * norm_prev + EPS_COS);
+
+    scratch_out[cosine_idx_base + g]  = cos_sim;
+    scratch_out[l2_norm_idx_base + g] = norm_curr;
+
+    // Writeback grad_curr to grad_prev for the next step's comparison.
+    // Same thread that read grad_prev writes back grad_curr — no race.
+    for (int i = start; i < end; ++i) {
+        grad_prev_buf_writeback[i] = grad_buf[i];
+    }
+
+    // Make all writes globally visible before pearl_4_adam_hparams_update
+    // reads scratch_out and before the next step reads grad_prev_buf.
+    __threadfence_system();
+}
diff --git a/crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu
new file mode 100644
index 000000000..379a55703
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu
@@ -0,0 +1,54 @@
+// crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu
+//
+// SP5 Task A4: Pearl 4 — per-param-group Adam β1, β2, ε from gradient
+// direction-stability + L2 norm.
+// 8-thread single-block (one per param group: DqnTrunk/Value/Branches/IQN/
+// IqlHigh/IqlLow/Attn/Curiosity).
+//
+// Producer signal: per-group cosine similarity of current vs previous-step
+// gradient direction (written by grad_cosine_sim_kernel). High stability
+// (cosine→1) → β2 → 0.9995 (long memory); low stability (cosine→0 or →-1)
+// → β2 → 0.99 (short memory). β1 likewise scales with stability.
+//
+// Structural envelopes (Invariant 1 anchors per spec line 88):
+//   β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6]
+//
+// THEORETICAL CAVEAT: adaptive β breaks Adam's constant-β convergence
+// proof. Mitigations:
+//   - Structural β envelopes (Invariant 1 anchors)
+//   - ALPHA_META=5e-4 in apply_pearls_ad_kernel (half the SP4 default → slow change rate)
+//   - Pearl A sentinel bootstrap on fold boundary (cosine_sim=0 → envelope midpoints)
+// If Layer C destabilization observed, fall-back path: revert to ε-only adaptive.
+
+#include <cuda_runtime.h>
+
+extern "C" __global__ void pearl_4_adam_hparams_update(
+    const float* __restrict__ grad_cosine_sim_per_group,  // [8] from grad_cosine_sim_kernel
+    const float* __restrict__ grad_l2_norm_per_group,     // [8] from grad_cosine_sim_kernel
+    float* __restrict__ scratch_out,
+    int   beta1_idx_base,
+    int   beta2_idx_base,
+    int   eps_idx_base
+) {
+    int g = threadIdx.x;
+    if (blockIdx.x != 0 || g >= 8) return;
+
+    float stability = fmaxf(0.0f, fminf(1.0f, grad_cosine_sim_per_group[g]));
+    float grad_norm = grad_l2_norm_per_group[g];
+
+    // Structural envelopes (Invariant 1 anchors per spec line 88).
+    // β1 scales linearly with stability: [0.85, 0.95]
+    float beta1 = 0.85f + 0.10f * stability;
+    // β2 scales linearly with stability: [0.99, 0.9995]
+    float beta2 = 0.99f + 0.0095f * stability;
+    // ε scales with grad norm to maintain numerical headroom; envelope clamp.
+    // grad_norm * 1e-7 maps typical grad norms [1e-3, 10.0] to [1e-10, 1e-6].
+    float eps = fmaxf(1e-10f, fminf(1e-6f, grad_norm * 1e-7f));
+
+    scratch_out[beta1_idx_base + g] = beta1;
+    scratch_out[beta2_idx_base + g] = beta2;
+    scratch_out[eps_idx_base + g]   = eps;
+
+    // Make writes globally visible before apply_pearls_ad_kernel consumes them.
+    __threadfence_system();
+}
diff --git a/crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs b/crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs
index 3f1d67c44..d1dcf72fd 100644
--- a/crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs
+++ b/crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs
@@ -146,6 +146,10 @@ pub fn pearls_ad_update(
 /// typically 1..9. Single-thread device-side loop is correct because each
 /// slot has its own state — no cross-slot parallelism is meaningful.
 ///
+/// `alpha_meta`: meta-EMA rate for Wiener variance tracking. SP4 default
+/// and all SP5 producers except Pearl 4 use [`ALPHA_META`] (1e-3). Pearl 4
+/// uses 5e-4 (half the default per spec line 89 to limit β change rate).
+///
 /// # Safety
 ///
 /// All three buffer pointers MUST be valid device pointers. Index ranges
@@ -168,6 +172,7 @@ pub(crate) unsafe fn launch_apply_pearls(
     wiener_state_dev: u64,
     wiener_offset_base: i32,
     n_slots: i32,
+    alpha_meta: f32,
 ) -> Result<(), crate::MLError> {
     use cudarc::driver::{LaunchConfig, PushKernelArg};
 
@@ -197,6 +202,7 @@ pub(crate) unsafe fn launch_apply_pearls(
         .arg(&wiener_state_dev)
         .arg(&wiener_offset_base)
         .arg(&n_slots)
+        .arg(&alpha_meta)
         .launch(cfg)
         .map_err(|e| {
             crate::MLError::ModelError(format!("apply_pearls_ad_kernel launch: {e}"))
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index ddb743384..b2e42e689 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -627,6 +627,30 @@ impl StateResetRegistry {
                 category: ResetCategory::FoldReset,
                 description: "GpuDqnTrainer.wiener_state_buf SP5 block [330 floats = 110 SP5 producers × {sample_var, diff_var, x_lag}] at offsets [213..543). SP5 Pearl D Wiener-EMA state for Task A1's 24 producers (16 Pearl 1 + 8 shared signal). Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..543) (SP4 + SP5 together) so Pearl A's sentinel branch fires on the new fold's first producer launch. Both SP4 and SP5 blocks reset atomically — wiener_state_buf grew from 213 to 543 floats in Task A1.",
             },
+            // SP5 Task A4: per-group Adam β1/β2/ε ISV slots (Pearl 4).
+            // ISV[226..250) = 24 floats (8 groups × 3 hyperparams).
+            // Pearl A sentinel 0 at fold boundary → first step fires first-observation
+            // replacement (cosine_sim=0 → low-stability envelope midpoints).
+            RegistryEntry {
+                name: "sp5_adam_beta1",
+                category: ResetCategory::FoldReset,
+                description: "ISV[ADAM_BETA1_BASE=226..234) — per-group Adam β1 (gradient direction stability). Pearl 4 producer: cosine_sim→[0,1] → β1 = 0.85 + 0.10×stability ∈ [0.85, 0.95]. SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
+            },
+            RegistryEntry {
+                name: "sp5_adam_beta2",
+                category: ResetCategory::FoldReset,
+                description: "ISV[ADAM_BETA2_BASE=234..242) — per-group Adam β2 (gradient direction stability). Pearl 4 producer: cosine_sim→[0,1] → β2 = 0.99 + 0.0095×stability ∈ [0.99, 0.9995]. SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
+            },
+            RegistryEntry {
+                name: "sp5_adam_eps",
+                category: ResetCategory::FoldReset,
+                description: "ISV[ADAM_EPS_BASE=242..250) — per-group Adam ε (grad L2 norm scaled). Pearl 4 producer: ε = clamp(grad_norm×1e-7, 1e-10, 1e-6). SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
+            },
+            RegistryEntry {
+                name: "sp5_grad_prev_buf",
+                category: ResetCategory::FoldReset,
+                description: "GpuDqnTrainer.grad_prev_buf_per_group [total_params f32s] — previous-step gradient direction for cosine similarity computation (Pearl 4). Zero at fold boundary so the new fold's first grad_cosine_sim_update sees cosine_sim=0 (low-stability bootstrap) and Pearl A fires first-observation replacement (Task A4).",
+            },
         ];
         Self { entries }
     }
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 0a5084607..f8b86f1cf 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -3603,6 +3603,13 @@ impl DQNTrainer {
                     if let Err(e) = fused.trainer().launch_sp5_pearl_2_budget() {
                         tracing::warn!(error = %e, "SP5 Pearl 2 producer launch failed");
                     }
+                    // SP5 Task A4: pearl_4_adam_hparams → apply_pearls_ad ×24
+                    // Reads grad_buf (populated in this step's backward pass) + grad_prev_buf.
+                    // No ordering dependency on Pearls 1/2/3 — reads gradient state, not ISV.
+                    // Writes ISV[226..250) = β1[8] + β2[8] + ε[8] (ALPHA_META_PEARL_4=5e-4).
+                    if let Err(e) = fused.trainer().launch_sp5_pearl_4_adam_hparams() {
+                        tracing::warn!(error = %e, "SP5 Pearl 4 producer launch failed");
+                    }
                 }
             }
 
diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs
index 2fd8615d9..c0ba714e8 100644
--- a/crates/ml/tests/sp5_producer_unit_tests.rs
+++ b/crates/ml/tests/sp5_producer_unit_tests.rs
@@ -33,6 +33,9 @@ const SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
 const SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));
 
+const SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_kernel.cubin"));
+
 // ── Helpers ───────────────────────────────────────────────────────────────────
 
 fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
@@ -55,6 +58,16 @@ fn load_pearl_2_budget_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
         .expect("load pearl_2_budget_update function")
 }
 
+fn load_pearl_4_adam_hparams_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
+    let module = stream
+        .context()
+        .load_cubin(SP5_PEARL_4_ADAM_HPARAMS_CUBIN.to_vec())
+        .expect("load pearl_4_adam_hparams_kernel cubin");
+    module
+        .load_function("pearl_4_adam_hparams_update")
+        .expect("load pearl_4_adam_hparams_update function")
+}
+
 fn make_test_stream() -> Arc<CudaStream> {
     let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
     ctx.default_stream()
@@ -861,3 +874,234 @@ fn pearl_2_budget_sharp_regime_iqn_dominates() {
         );
     }
 }
+
+// ── Tests 7 + 8: pearl_4_adam_hparams_update ──────────────────────────────────
+
+/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with high stability input.
+///
+/// Synthetic input: cosine_sim=[1.0; 8] (maximum stability), grad_norm=[1.0; 8].
+///
+/// Analytical expected:
+///   stability = fmaxf(0, fminf(1, cosine_sim)) = 1.0
+///   β1 = 0.85 + 0.10 × 1.0 = 0.95   (long-memory ceiling)
+///   β2 = 0.99 + 0.0095 × 1.0 = 0.9995 (long-memory ceiling)
+///   ε  = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7 → clamped to 1e-7 (within [1e-10, 1e-6])
+///
+/// Structural envelopes (Invariant 1 anchors): β1∈[0.85,0.95], β2∈[0.99,0.9995].
+/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
+#[test]
+#[ignore = "requires GPU"]
+fn pearl_4_adam_hparams_high_stability_yields_long_memory() {
+    // Scratch layout (Pearl 4 outputs):
+    //   [0..8)  = β1[8] (beta1_idx_base=0)
+    //   [8..16) = β2[8] (beta2_idx_base=8)
+    //   [16..24) = ε[8] (eps_idx_base=16)
+    const BETA1_BASE:    usize = 0;
+    const BETA2_BASE:    usize = 8;
+    const EPS_BASE:      usize = 16;
+    const SCRATCH_SIZE:  usize = 24;
+    const N_GROUPS:      usize = 8;
+
+    let stream = make_test_stream();
+    let kernel = load_pearl_4_adam_hparams_kernel(&stream);
+
+    // cosine_sim = 1.0 for all 8 groups → maximum stability.
+    let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
+        .expect("alloc cosine_buf");
+    cosine_buf.write_from_slice(&[1.0_f32; N_GROUPS]);
+
+    // grad_norm = 1.0 for all 8 groups → ε = 1e-7 (within envelope).
+    let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
+        .expect("alloc l2_norm_buf");
+    l2_norm_buf.write_from_slice(&[1.0_f32; N_GROUPS]);
+
+    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
+        .expect("alloc scratch_buf");
+
+    let cosine_dev      = cosine_buf.dev_ptr;
+    let l2_dev          = l2_norm_buf.dev_ptr;
+    let scratch_dev     = scratch_buf.dev_ptr;
+    let beta1_base_i32  = BETA1_BASE as i32;
+    let beta2_base_i32  = BETA2_BASE as i32;
+    let eps_base_i32    = EPS_BASE   as i32;
+
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&cosine_dev)
+            .arg(&l2_dev)
+            .arg(&scratch_dev)
+            .arg(&beta1_base_i32)
+            .arg(&beta2_base_i32)
+            .arg(&eps_base_i32)
+            .launch(LaunchConfig {
+                grid_dim:         (1, 1, 1),
+                block_dim:        (8, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .expect("launch pearl_4_adam_hparams_update (high stability)");
+    }
+    stream.synchronize().expect("sync after pearl_4_adam_hparams_update (high stability)");
+
+    let scratch = scratch_buf.read_all();
+
+    // Analytical values for stability = 1.0:
+    let expected_beta1: f32 = 0.85 + 0.10 * 1.0;   // = 0.95
+    let expected_beta2: f32 = 0.99 + 0.0095 * 1.0;  // = 0.9995
+    // ε = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7
+    let expected_eps: f32   = 1.0_f32 * 1e-7_f32;   // within [1e-10, 1e-6]
+
+    for g in 0..N_GROUPS {
+        let beta1 = scratch[BETA1_BASE + g];
+        let beta2 = scratch[BETA2_BASE + g];
+        let eps   = scratch[EPS_BASE   + g];
+
+        println!("group={g} beta1={beta1:.6} beta2={beta6:.7} eps={eps:.2e}",
+            beta6 = beta2);
+
+        // β1 ceiling: 0.95 ± 1 ULP.
+        let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
+        assert!(
+            rel_b1 < 1e-5,
+            "group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
+        );
+        // β2 ceiling: 0.9995 ± 1 ULP.
+        let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
+        assert!(
+            rel_b2 < 1e-5,
+            "group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
+        );
+        // ε must be within structural envelope.
+        assert!(
+            eps >= 1e-10_f32 && eps <= 1e-6_f32,
+            "group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
+        );
+        let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
+        assert!(
+            rel_eps < 1e-4,
+            "group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
+        );
+        // β1 must be at structural ceiling for maximum stability.
+        assert!(
+            (beta1 - 0.95_f32).abs() < 1e-5,
+            "group={g}: beta1={beta1:.6} should be at ceiling 0.95 for stability=1.0"
+        );
+        // β2 must be at structural ceiling for maximum stability.
+        assert!(
+            (beta2 - 0.9995_f32).abs() < 1e-5,
+            "group={g}: beta2={beta2:.7} should be at ceiling 0.9995 for stability=1.0"
+        );
+    }
+}
+
+/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with low stability input.
+///
+/// Synthetic input: cosine_sim=[0.0; 8] (minimum stability), grad_norm=[0.01; 8].
+///
+/// Analytical expected:
+///   stability = fmaxf(0, fminf(1, 0.0)) = 0.0
+///   β1 = 0.85 + 0.10 × 0.0 = 0.85   (low-stability floor)
+///   β2 = 0.99 + 0.0095 × 0.0 = 0.99  (low-stability floor)
+///   ε  = clamp(0.01 × 1e-7, 1e-10, 1e-6) = 1e-9 → clamped to 1e-10
+///
+/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
+#[test]
+#[ignore = "requires GPU"]
+fn pearl_4_adam_hparams_low_stability_yields_short_memory() {
+    const BETA1_BASE:    usize = 0;
+    const BETA2_BASE:    usize = 8;
+    const EPS_BASE:      usize = 16;
+    const SCRATCH_SIZE:  usize = 24;
+    const N_GROUPS:      usize = 8;
+
+    let stream = make_test_stream();
+    let kernel = load_pearl_4_adam_hparams_kernel(&stream);
+
+    // cosine_sim = 0.0 for all 8 groups → minimum stability.
+    let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
+        .expect("alloc cosine_buf");
+    cosine_buf.write_from_slice(&[0.0_f32; N_GROUPS]);
+
+    // grad_norm = 0.01 → ε = 0.01 × 1e-7 = 1e-9 → clamped to 1e-10.
+    let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
+        .expect("alloc l2_norm_buf");
+    l2_norm_buf.write_from_slice(&[0.01_f32; N_GROUPS]);
+
+    let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
+        .expect("alloc scratch_buf");
+
+    let cosine_dev      = cosine_buf.dev_ptr;
+    let l2_dev          = l2_norm_buf.dev_ptr;
+    let scratch_dev     = scratch_buf.dev_ptr;
+    let beta1_base_i32  = BETA1_BASE as i32;
+    let beta2_base_i32  = BETA2_BASE as i32;
+    let eps_base_i32    = EPS_BASE   as i32;
+
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&cosine_dev)
+            .arg(&l2_dev)
+            .arg(&scratch_dev)
+            .arg(&beta1_base_i32)
+            .arg(&beta2_base_i32)
+            .arg(&eps_base_i32)
+            .launch(LaunchConfig {
+                grid_dim:         (1, 1, 1),
+                block_dim:        (8, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .expect("launch pearl_4_adam_hparams_update (low stability)");
+    }
+    stream.synchronize().expect("sync after pearl_4_adam_hparams_update (low stability)");
+
+    let scratch = scratch_buf.read_all();
+
+    // Analytical values for stability = 0.0:
+    let expected_beta1: f32 = 0.85;   // low-stability floor
+    let expected_beta2: f32 = 0.99;   // low-stability floor
+    // ε = clamp(0.01 × 1e-7, 1e-10, 1e-6) = clamp(1e-9, 1e-10, 1e-6) = 1e-9
+    // Actually 1e-9 is within [1e-10, 1e-6], so no clamping needed.
+    let expected_eps: f32   = (0.01_f32 * 1e-7_f32).clamp(1e-10, 1e-6);  // = 1e-9
+
+    for g in 0..N_GROUPS {
+        let beta1 = scratch[BETA1_BASE + g];
+        let beta2 = scratch[BETA2_BASE + g];
+        let eps   = scratch[EPS_BASE   + g];
+
+        println!("group={g} beta1={beta1:.6} beta2={beta2:.7} eps={eps:.2e}");
+
+        // β1 floor: 0.85 ± 1 ULP.
+        let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
+        assert!(
+            rel_b1 < 1e-5,
+            "group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
+        );
+        // β2 floor: 0.99 ± 1 ULP.
+        let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
+        assert!(
+            rel_b2 < 1e-5,
+            "group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
+        );
+        // ε must be within structural envelope.
+        assert!(
+            eps >= 1e-10_f32 && eps <= 1e-6_f32,
+            "group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
+        );
+        let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
+        assert!(
+            rel_eps < 1e-4,
+            "group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
+        );
+        // β1 must be at structural floor for zero stability.
+        assert!(
+            (beta1 - 0.85_f32).abs() < 1e-5,
+            "group={g}: beta1={beta1:.6} should be at floor 0.85 for stability=0.0"
+        );
+        // β2 must be at structural floor for zero stability.
+        assert!(
+            (beta2 - 0.99_f32).abs() < 1e-5,
+            "group={g}: beta2={beta2:.7} should be at floor 0.99 for stability=0.0"
+        );
+    }
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 292a73927..c92a51710 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,8 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+SP5 Task A4 — Pearl 4 per-group Adam β1/β2/ε GPU producers (Layer A, 2026-05-01): two new CUDA kernels (`grad_cosine_sim_kernel.cu`, `pearl_4_adam_hparams_kernel.cu`) and one `MappedF32Buffer` (`grad_prev_buf_per_group`) land as Layer A additive producers that feed ISV slots [226..250) with per-param-group Adam hyperparameters. `grad_cosine_sim_update`: single-block 8-thread kernel (one thread per SP4 param group: DqnTrunk/Value/Branches/IQN/IqlHigh/IqlLow/Attn/Curiosity), reads `grad_buf[total_params]` + `grad_prev_buf[total_params]` + `group_param_offsets[9]` (prefix sums in element units); computes `dot = Σ gc·gp`, `norm_curr_sq = Σ gc²`, `norm_prev_sq = Σ gp²` per group; writes `cos_sim[8]→scratch[131..139)` and `l2_norm[8]→scratch[139..147)`; writebacks `grad_curr → grad_prev_buf` for the next step's comparison (same thread — no race). `pearl_4_adam_hparams_update`: single-block 8-thread kernel, reads `cos_sim[8]` + `l2_norm[8]` from scratch; clips stability to `[0, 1]`; computes `β1 = 0.85 + 0.10×stability ∈ [0.85, 0.95]`, `β2 = 0.99 + 0.0095×stability ∈ [0.99, 0.9995]`, `ε = clamp(grad_norm × 1e-7, 1e-10, 1e-6)`; writes `β1[8]→scratch[147..155)`, `β2[8]→scratch[155..163)`, `ε[8]→scratch[163..171)`. Structural envelopes (Invariant 1 anchors per spec): β1∈[0.85, 0.95], β2∈[0.99, 0.9995], ε∈[1e-10, 1e-6]. `launch_sp5_pearl_4_adam_hparams` fires 24 `launch_apply_pearls` calls: β1[8] → ISV[226..234), β2[8] → ISV[234..242), ε[8] → ISV[242..250). Uses `ALPHA_META_PEARL_4=5e-4` (half SP4 default 1e-3) per theoretical-caveat mitigation — slow β change rate reduces instability risk from adaptive β breaking Adam's constant-β convergence proof. Pearl A sentinel 0 at fold boundary fires first-observation replacement (cosine_sim=0 → low-stability envelope endpoints). Aux groups 3-7 (IQN/IqlHigh/IqlLow/Attn/Curiosity) have non-contiguous grad allocations; zero-width sentinel offsets `[tp, tp, tp, tp, tp]` in `group_param_offsets_buf` prevent the kernel inner loop from executing for those groups, yielding cosine_sim=0 → β1/β2 start at low-stability floors. `apply_pearls_kernel.cu` signature gained `float alpha_meta` parameter (migration of all 18 call sites atomic per `feedback_no_partial_refactor.md`: 1 in `gpu_experience_collector.rs`, 17 in `gpu_dqn_trainer.rs`). Buffer growth: `producer_step_scratch_buf` 131→171 slots (SP5_SCRATCH_TOTAL); 5 new scratch constants `SCRATCH_PEARL_4_{COSINE=131, L2=139, BETA1=147, BETA2=155, EPS=163}`. StateResetRegistry gains 4 new FoldReset entries: `sp5_adam_beta1` (ISV[226..234)), `sp5_adam_beta2` (ISV[234..242)), `sp5_adam_eps` (ISV[242..250)), `sp5_grad_prev_buf` (grad_prev_buf_per_group — zero at fold boundary so cosine_sim=0 fires Pearl A bootstrap). Wire-up in training_loop.rs: `launch_sp5_pearl_4_adam_hparams()` called immediately after `launch_sp5_pearl_2_budget()` with `tracing::warn!` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (7) high-stability (cosine_sim=1.0 → β1=0.95, β2=0.9995, ε within envelope); (8) low-stability (cosine_sim=0.0 → β1=0.85, β2=0.99, ε within envelope). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/grad_cosine_sim_kernel.cu` (new), `cuda_pipeline/pearl_4_adam_hparams_kernel.cu` (new), `cuda_pipeline/apply_pearls_kernel.cu` (+alpha_meta param), `cuda_pipeline/sp4_wiener_ema.rs` (+alpha_meta arg), `cuda_pipeline/gpu_experience_collector.rs` (+1 call site ALPHA_META arg), `build.rs` (+2 cubin entries), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 131→171 + 5 new constants + 2 static cubins + 4 struct fields + constructor allocs + struct initializer + 17 call-site ALPHA_META args + launch method), `trainers/dqn/state_reset_registry.rs` (+4 FoldReset entries), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 2 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + cubin constant + loader helper). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (apply_pearls 18 call sites all migrated).
+
 SP5 Task A3 fix-up — Pearl 2 budget tests now assert sum-to-1 normalization invariant (2026-05-01): code-quality review caught that the two A3 unit tests verified each output (c51, iqn, cql, ens) against its analytical expected value within 1% relative tolerance, but did NOT assert the structural invariant `c51 + iqn + cql + ens ≈ 1.0` that the kernel maintains by construction (`ens = max(0, 1 - iqn - c51 - cql)`). A coefficient typo such as `BASE_IQN = 0.111` instead of `0.11` would produce individually-plausible per-component values that all still pass the relative checks while quietly summing to 0.97 or 1.04. Fix adds `assert!((c51+iqn+cql+ens - 1.0).abs() < 1e-4)` inside both per-branch loops. Touched: `crates/ml/tests/sp5_producer_unit_tests.rs` (+12 LOC). cargo test --no-run clean.
 
 SP5 Task A3 — Pearl 2 per-branch loss budget GPU producer (Layer A, 2026-05-01): one new CUDA kernel (`pearl_2_budget_kernel.cu`) lands as a Layer A additive producer that feeds ISV slots [190..210) with per-branch loss budget and flatness signals. No loss-budget consumer migration in this commit — Layer B wires consumers. `pearl_2_budget_update`: single-block 4-thread kernel (one thread per branch), reads Q_VAR_PER_BRANCH[b] (Pearl 1's q_branch_stats output, ISV[222..226)) + NOISY_SIGMA[b] (Pearl 3 output, ISV[210..214)); computes `flatness[b] = clamp(var_q[b] / (σ[b]² + EPS_DIV), 0, 1)` where EPS_DIV=1e-8 is an Invariant 1 anchor (numerical stability); derives `budget_iqn = BASE_IQN + (1 - flatness) * (1 - BASE_IQN)` (IQN dominates when Q is flat relative to noise), `budget_c51 = flatness * (1 - BASE_IQN) * BASE_C51_SHARE` (C51 takes proportional remainder when Q is well-separated), `budget_cql = 0` (gated by existing regime_stability allocator), `budget_ens = max(0, 1 - iqn - c51 - cql)`; writes 20 floats to scratch[111..131) (SCRATCH_PEARL_2_C51=111, SCRATCH_PEARL_2_IQN=115, SCRATCH_PEARL_2_CQL=119, SCRATCH_PEARL_2_ENS=123, SCRATCH_PEARL_2_FLATNESS=127). BASE_IQN=0.11 and BASE_C51_SHARE=0.84/(1-BASE_IQN) are Invariant 1 anchors (preserve SP4-baseline budget under non-flat regime). `launch_sp5_pearl_2_budget` fires 20 `launch_apply_pearls` calls: 5 output groups × 4 branches each, mapping to ISV[BUDGET_C51_BASE=190..194), ISV[BUDGET_IQN_BASE=194..198), ISV[BUDGET_CQL_BASE=198..202), ISV[BUDGET_ENS_BASE=202..206), ISV[FLATNESS_BASE=206..210). Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) * 3` where `base_wiener_offset = SP4_PRODUCER_COUNT * 3 = 213`. Dependencies: Pearl 2 must chain AFTER Pearl 1 (reads Q_VAR_PER_BRANCH) AND AFTER Pearl 3 (reads NOISY_SIGMA). Buffer growth: `producer_step_scratch_buf` 111→131 slots (SP5_SCRATCH_TOTAL updated 111→131; 5 new module-level constants SCRATCH_PEARL_2_{C51,IQN,CQL,ENS,FLATNESS}). wiener_state_buf already at 543 (A1 sized for entire SP5 block — no further growth). StateResetRegistry gains 5 new FoldReset entries: sp5_budget_c51 (ISV[190..194)), sp5_budget_iqn (ISV[194..198)), sp5_budget_cql (ISV[198..202)), sp5_budget_ens (ISV[202..206)), sp5_flatness (ISV[206..210)) — Pearl A sentinel 0 at fold boundary triggers first-observation replacement on new fold's first launch. Wire-up in training_loop.rs: `launch_sp5_pearl_2_budget()` called immediately after `launch_sp5_pearl_3_sigma()` with `tracing::warn!(error = %e, ...)` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (5) flat-Q regime (σ=1.0, var_q=2.0 → flatness=1.0 → iqn=0.11, c51≈0.84); (6) sharp-Q regime (σ=2.0, var_q=0.04 → flatness≈0.01 → iqn>0.9, c51≈0.0084). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/pearl_2_budget_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (static cubin + SP5_SCRATCH_TOTAL 111→131 + 5 new scratch constants + docstring update + struct field + cubin loader + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+5 FoldReset entries), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 3 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + cubin loader helper). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_stubs, feedback_no_partial_refactor (all consumers of existing budget allocator untouched).