diff --git a/crates/ml-dqn/src/gpu_replay_buffer.rs b/crates/ml-dqn/src/gpu_replay_buffer.rs
index 620bb6466..d237dd7ce 100644
--- a/crates/ml-dqn/src/gpu_replay_buffer.rs
+++ b/crates/ml-dqn/src/gpu_replay_buffer.rs
@@ -67,6 +67,8 @@ struct ReplayKernels {
     per_insert_pa: CudaFunction,
     per_prefix_scan: CudaFunction,
     per_sample: CudaFunction,
+    /// C1/P1: Diversity-weighted priority kernel (health < 0.8 path).
+    pow_alpha_diverse_f32: CudaFunction,
 }
 
 impl ReplayKernels {
@@ -108,6 +110,7 @@ impl ReplayKernels {
             per_insert_pa: per_ld("per_insert_pa")?,
             per_prefix_scan: per_ld("per_prefix_scan")?,
             per_sample: per_ld("per_sample")?,
+            pow_alpha_diverse_f32: ld("pow_alpha_diverse_f32")?,
         })
     }
 }
@@ -196,6 +199,9 @@ pub struct GpuReplayBuffer {
     trainer_dones_ptr: u64,
     trainer_is_weights_ptr: u64,
     trainer_state_dim_padded: usize,
+    /// C1/P1: Cached learning_health value from the trainer (updated once per epoch).
+    /// Controls whether priority updates use the standard or diversity-weighted kernel.
+    learning_health_cache: f32,
 }
 
 impl Drop for GpuReplayBuffer {
@@ -323,6 +329,7 @@ impl GpuReplayBuffer {
             trainer_dones_ptr: 0,
             trainer_is_weights_ptr: 0,
             trainer_state_dim_padded: 0,
+            learning_health_cache: 1.0,
         })
     }
 
@@ -350,6 +357,11 @@ impl GpuReplayBuffer {
     pub const fn epsilon(&self) -> f32 { self.config.epsilon }
     pub const fn state_dim(&self) -> usize { ml_core::state_layout::STATE_DIM }
 
+    /// C1/P1: Update the cached learning_health value. Called from trainer at epoch boundary.
+    pub fn set_learning_health(&mut self, health: f32) {
+        self.learning_health_cache = health.clamp(0.0, 1.0);
+    }
+
     /// Wire trainer destination buffer pointers for direct-to-trainer gather.
     /// Called once after GpuDqnTrainer is constructed. Pointers are stable
     /// (CudaSlice allocations never move), so this is safe for graph capture.
@@ -818,21 +830,79 @@ impl GpuReplayBuffer {
 
         let (al, ep, bsi) = (self.config.alpha, self.config.epsilon, bs as i32);
         let cap_i = self.config.capacity as i32;
+        let health = self.learning_health_cache;
 
         stream.memset_zeros(&mut self.update_batch_max)
             .map_err(|e| MLError::ModelError(format!("zero batch_max: {e}")))?;
 
-        // per_update_pa: reads f32 td_errors directly, writes priorities + priorities_pa
-        unsafe {
-            stream.launch_builder(&self.kernels.per_update_pa)
-                .arg(&self.priorities_pa)
-                .arg(&self.priorities)
-                .arg(&self.update_batch_max)
-                .arg(indices)
-                .arg(td_errors)
-                .arg(&al).arg(&ep).arg(&cap_i).arg(&bsi)
-                .launch(lcfg(bs))
-                .map_err(|e| MLError::ModelError(format!("per_update_pa: {e}")))?;
+        if health < 0.8 {
+            // C1/P1: during collapse, boost priorities of experiences whose action
+            // deviates from the batch mean. Requires a host-side mean readback to
+            // avoid a full reduction kernel — this path runs once per epoch at epoch boundary,
+            // so the DtoH sync cost is acceptable.
+            //
+            // Re-gather actions from the ring buffer using the current batch indices
+            // (sample_indices_i64 from the last sample_proportional call). This avoids
+            // depending on trainer_actions_ptr validity after graph replay.
+            let mbs = self.config.max_batch_size.max(1);
+            let gather_n = bs.min(mbs);
+            let gather_ni = gather_n as i32;
+            // Reuse sample_actions (CudaSlice<u32>) as scratch. gather_u32 gathers
+            // the action values at the sampled indices from the ring buffer.
+            let cap_i32 = self.config.capacity as i32;
+            unsafe {
+                stream.launch_builder(&self.kernels.gather_u32)
+                    .arg(&mut self.sample_actions)
+                    .arg(&self.actions)
+                    .arg(&self.sample_indices_i64)
+                    .arg(&gather_ni)
+                    .arg(&cap_i32)
+                    .launch(lcfg(gather_n))
+                    .map_err(|e| MLError::ModelError(format!("gather_u32 for diversity: {e}")))?;
+            }
+            // Synchronize to allow DtoH readback (epoch-boundary, not on hot path).
+            unsafe {
+                cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream());
+            }
+            let mut actions_host = vec![0u32; gather_n];
+            stream.memcpy_dtoh(&self.sample_actions.slice(..gather_n), &mut actions_host)
+                .map_err(|e| MLError::ModelError(format!("actions dtoh diversity: {e}")))?;
+            let sum_actions: i64 = actions_host.iter().map(|&a| a as i64).sum();
+            let mean_action_scaled: i32 = ((sum_actions * 1000) / (gather_n as i64).max(1)) as i32;
+
+            // Reinterpret sample_actions (u32) as i32 for the kernel argument.
+            // The pointer type pun is safe: same device memory, same bit-width, kernel reads int*.
+            let actions_dev_ptr = self.sample_actions.raw_ptr() as u64;
+            unsafe {
+                stream.launch_builder(&self.kernels.pow_alpha_diverse_f32)
+                    .arg(&self.priorities_pa)
+                    .arg(&self.priorities)
+                    .arg(&self.update_batch_max)
+                    .arg(indices)
+                    .arg(td_errors)
+                    .arg(&actions_dev_ptr)
+                    .arg(&mean_action_scaled)
+                    .arg(&al)
+                    .arg(&ep)
+                    .arg(&health)
+                    .arg(&cap_i)
+                    .arg(&bsi)
+                    .launch(lcfg(bs))
+                    .map_err(|e| MLError::ModelError(format!("pow_alpha_diverse_f32: {e}")))?;
+            }
+        } else {
+            // Standard PER (unchanged): reads f32 td_errors directly, writes priorities + priorities_pa
+            unsafe {
+                stream.launch_builder(&self.kernels.per_update_pa)
+                    .arg(&self.priorities_pa)
+                    .arg(&self.priorities)
+                    .arg(&self.update_batch_max)
+                    .arg(indices)
+                    .arg(td_errors)
+                    .arg(&al).arg(&ep).arg(&cap_i).arg(&bsi)
+                    .launch(lcfg(bs))
+                    .map_err(|e| MLError::ModelError(format!("per_update_pa: {e}")))?;
+            }
         }
 
         // Merge batch max into pending (unchanged logic)
diff --git a/crates/ml-dqn/src/replay_buffer_kernels.cu b/crates/ml-dqn/src/replay_buffer_kernels.cu
index 04800c927..9de823983 100644
--- a/crates/ml-dqn/src/replay_buffer_kernels.cu
+++ b/crates/ml-dqn/src/replay_buffer_kernels.cu
@@ -471,6 +471,60 @@ void f32_idx_to_u32(
     out[i] = (unsigned int)input[i];
 }
 
+/* C1/P1: Health-weighted PER priority — boosts priorities of diverse-action
+ * experiences during collapse (health < 0.8).
+ *
+ *   new_prio = clamp((|td_error[i]| + epsilon)^alpha
+ *                    × (1 + 2×(1−health) × |action[i] − mean_action|),
+ *                    epsilon, 1e6)
+ *   priorities[indices[i]]    = new_prio
+ *   priorities_pa[indices[i]] = new_prio^alpha
+ *   atomicMax(batch_max, new_prio)  (IEEE 754 int trick)
+ *
+ * `mean_action_scaled` is the batch mean of actions multiplied by 1000 (passed
+ * as int to avoid precision issues when reducing host-side).
+ *
+ * Drop-in replacement for per_update_pa when health < 0.8.
+ */
+extern "C" __global__ void pow_alpha_diverse_f32(
+    float* __restrict__ priorities_pa,
+    float* __restrict__ priorities,
+    float* __restrict__ batch_max,
+    const unsigned int* __restrict__ indices,
+    const float* __restrict__ td_errors,
+    const int* __restrict__ actions,
+    int mean_action_scaled,
+    float alpha,
+    float epsilon,
+    float health,
+    int capacity,
+    int batch_size
+) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= batch_size) return;
+
+    unsigned int idx = indices[i];
+    if (idx >= (unsigned int)capacity) return;
+
+    float base = powf(fabsf(td_errors[i]) + epsilon, alpha);
+
+    float mean_action = (float)mean_action_scaled * 0.001f;
+    float action_diff = fabsf((float)actions[i] - mean_action);
+    float diversity_mult = 1.0f + 2.0f * (1.0f - health) * action_diff;
+
+    float new_prio = base * diversity_mult;
+    if (new_prio < epsilon) new_prio = epsilon;
+    if (new_prio > 1e6f) new_prio = 1e6f;
+
+    priorities[idx] = new_prio;
+
+    int ival = __float_as_int(new_prio);
+    atomicMax((int*)batch_max, ival);
+
+    float pa = powf(new_prio, alpha);
+    priorities_pa[idx] = pa;
+}
+
 // ── 21. Regime weight combination kernel ────────────────────────────────────
 //    out[i] = trending[i] * t_scale + ranging[i] * r_scale + volatile[i] * v_scale
 
diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
index c486a2d77..1c6b16804 100644
--- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
@@ -161,6 +161,12 @@ __device__ void block_bellman_project_f(
     }
 }
 
+/* B3/G4: Read learning_health from ISV buffer at index 12.
+ * Returns 0.5 if pointer is null (safe fallback). */
+__device__ __forceinline__ float get_learning_health(const float* __restrict__ isv_signals_ptr) {
+    return (isv_signals_ptr != NULL) ? isv_signals_ptr[12] : 0.5f;
+}
+
 /* ══════════════════════════════════════════════════════════════════════
  * MAIN KERNEL: c51_loss_batched (float arithmetic, BF16 I/O)
  * ══════════════════════════════════════════════════════════════════════ */
@@ -232,7 +238,10 @@ extern "C" __global__ void c51_loss_batched(
     const float* __restrict__ atom_positions,      /* [4, num_atoms] adaptive positions. NULL = linear. */
 
     /* ── Per-sample CVaR alpha from learned risk branch ── */
-    const float* __restrict__ cvar_alpha_buf       /* [B] per-sample CVaR alpha from risk branch. NULL = use iqn_readiness. */
+    const float* __restrict__ cvar_alpha_buf,      /* [B] per-sample CVaR alpha from risk branch. NULL = use iqn_readiness. */
+
+    /* ── B3/G4: ISV signals for health-scaled Expected SARSA temperature ── */
+    const float* __restrict__ isv_signals          /* [ISV_DIM=13] pinned device-mapped. isv_signals[12] = learning_health. NULL = 0.5 fallback. */
 ) {
     extern __shared__ float shmem_f[];
 
@@ -543,7 +552,14 @@ extern "C" __global__ void c51_loss_batched(
              * meaningful at any Q-value scale. 1% of |mean Q| or 1e-6. */
             float mean_q = (max_eq + min_eq) * 0.5f;
             float tau_floor = fmaxf(fabsf(mean_q) * 0.01f, 1e-6f);
-            float tau = fmaxf(q_gap_local, tau_floor);
+            float tau_base = fmaxf(q_gap_local, tau_floor);
+
+            // B3/G4: health-coupled temperature scaling.
+            // health=1 (healthy): factor=1.0 → sharp softmax (near argmax target, deterministic).
+            // health=0 (collapsed): factor=6.0 → wide softmax (stochastic sampling breaks attractor).
+            float health = get_learning_health(isv_signals);
+            float tau_health_factor = 1.0f + 5.0f * (1.0f - health);
+            float tau = tau_base * tau_health_factor;
             float sum_exp = 0.0f;
             for (int a = 0; a < n_d; a++) {
                 action_weights[a] = expf((eq_per_action[a] - max_eq) / tau);
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index df6a5c572..42051e5bb 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -761,7 +761,8 @@ extern "C" __global__ void experience_action_select(
     float eps_urg_mult,    /* per-branch epsilon multiplier: urgency */
     int timestep,          /* current timestep for stateless RNG */
     const float* __restrict__ per_sample_epsilon,  /* [N] IQL expectile gap epsilon, NULL=use cosine schedule */
-    const float* __restrict__ isv_signals_ptr      /* [8] pinned ISV signals for adaptive hold. NULL = static. */
+    const float* __restrict__ isv_signals_ptr,     /* [8] pinned ISV signals for adaptive hold. NULL = static. */
+    int contrarian_active  /* D7/N7: when non-zero, negate Q values before Boltzmann → argmin sampling */
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -804,6 +805,14 @@ extern "C" __global__ void experience_action_select(
     const float* q_b2 = q_b1 + b1_size;
     const float* q_b3 = q_b2 + b2_size;
 
+    /* D7/N7: contrarian sign flip.
+     * When contrarian_active != 0, negate Q values before Boltzmann softmax.
+     * softmax(-q/tau) is argmin-favoring — the model picks its least-preferred
+     * action, breaking systematic anti-correlation attractors.
+     * When contrarian_active == 0, q_sign = +1.0f → bit-identical to previous
+     * behavior (identity multiplication). */
+    float q_sign = (contrarian_active != 0) ? -1.0f : 1.0f;
+
     int dir_idx, mag_idx, a2, a3;
 
     /* Hold enforcement removed — replaced by cost-driven hold timing.
@@ -833,10 +842,10 @@ extern "C" __global__ void experience_action_select(
          * Q-values differentiate (converges to argmax deterministically).
          * When Q-values are flat, spreads evenly instead of flipping on noise. */
         /* Training mode: Boltzmann softmax over direction Q-values */
-        float q_max_d = (q_b0[0]);
+        float q_max_d = q_sign * (q_b0[0]);
         float q_min_d = q_max_d;
         for (int a = 1; a < b0_size; a++) {
-            float qv = (q_b0[a]);
+            float qv = q_sign * (q_b0[a]);
             q_max_d = fmaxf(q_max_d, qv);
             q_min_d = fminf(q_min_d, qv);
         }
@@ -845,7 +854,7 @@ extern "C" __global__ void experience_action_select(
             float sum_e = 0.0f;
             float exps_d[4];  /* b0_size=4: Short/Hold/Long/Flat */
             for (int a = 0; a < b0_size; a++) {
-                float qv = (q_b0[a]);
+                float qv = q_sign * (q_b0[a]);
                 exps_d[a] = expf((qv - q_max_d) / tau_d);
                 sum_e += exps_d[a];
             }
@@ -889,10 +898,10 @@ extern "C" __global__ void experience_action_select(
         } else {
             /* Adaptive temperature: scale by Q-range so Boltzmann is meaningful
              * regardless of absolute Q magnitude. Floor at 0.01 to avoid div/0. */
-            float q_max_m = (q_b1[0]);
+            float q_max_m = q_sign * (q_b1[0]);
             float q_min_m = q_max_m;
             for (int a = 1; a < b1_size; a++) {
-                float qv = (q_b1[a]);
+                float qv = q_sign * (q_b1[a]);
                 q_max_m = fmaxf(q_max_m, qv);
                 q_min_m = fminf(q_min_m, qv);
             }
@@ -906,7 +915,7 @@ extern "C" __global__ void experience_action_select(
             float sum_e = 0.0f;
             float exps[MAX_MAGNITUDE_ACTIONS];
             for (int a = 0; a < b1_size && a < MAX_MAGNITUDE_ACTIONS; a++) {
-                float qv = (q_b1[a]);
+                float qv = q_sign * (q_b1[a]);
                 exps[a] = expf((qv - q_max_m) / tau);
                 sum_e += exps[a];
             }
@@ -925,10 +934,11 @@ extern "C" __global__ void experience_action_select(
         a2 = (r >= b2_size) ? b2_size - 1 : r;
     } else {
         /* Boltzmann softmax over order Q-values */
-        float q_max_ord = q_b2[0], q_min_ord = q_b2[0];
+        float q_max_ord = q_sign * q_b2[0], q_min_ord = q_max_ord;
         for (int a = 1; a < b2_size; a++) {
-            q_max_ord = fmaxf(q_max_ord, q_b2[a]);
-            q_min_ord = fminf(q_min_ord, q_b2[a]);
+            float qv_ord = q_sign * q_b2[a];
+            q_max_ord = fmaxf(q_max_ord, qv_ord);
+            q_min_ord = fminf(q_min_ord, qv_ord);
         }
         /* Higher temperature floor (0.5 vs 0.01) for order: Q-values for different
          * order types start nearly identical (same reward signal). The low 0.01 floor
@@ -939,7 +949,7 @@ extern "C" __global__ void experience_action_select(
         float sum_e = 0.0f;
         float exps_ord[3];
         for (int a = 0; a < b2_size; a++) {
-            exps_ord[a] = expf((q_b2[a] - q_max_ord) / tau_ord);
+            exps_ord[a] = expf((q_sign * q_b2[a] - q_max_ord) / tau_ord);
             sum_e += exps_ord[a];
         }
         float ro = philox_uniform(i, timestep, rng_ctr++) * sum_e;
@@ -957,16 +967,17 @@ extern "C" __global__ void experience_action_select(
         a3 = (r >= b3_size) ? b3_size - 1 : r;
     } else {
         /* Boltzmann softmax over urgency Q-values */
-        float q_max_urg = q_b3[0], q_min_urg = q_b3[0];
+        float q_max_urg = q_sign * q_b3[0], q_min_urg = q_max_urg;
         for (int a = 1; a < b3_size; a++) {
-            q_max_urg = fmaxf(q_max_urg, q_b3[a]);
-            q_min_urg = fminf(q_min_urg, q_b3[a]);
+            float qv_urg = q_sign * q_b3[a];
+            q_max_urg = fmaxf(q_max_urg, qv_urg);
+            q_min_urg = fminf(q_min_urg, qv_urg);
         }
         float tau_urg = fmaxf(q_max_urg - q_min_urg, 0.5f);
         float sum_e = 0.0f;
         float exps_urg[3];
         for (int a = 0; a < b3_size; a++) {
-            exps_urg[a] = expf((q_b3[a] - q_max_urg) / tau_urg);
+            exps_urg[a] = expf((q_sign * q_b3[a] - q_max_urg) / tau_urg);
             sum_e += exps_urg[a];
         }
         float ru = philox_uniform(i, timestep, rng_ctr++) * sum_e;
@@ -999,13 +1010,22 @@ extern "C" __global__ void experience_action_select(
 
     /* Output Q-gap for conviction-based position sizing.
      * Q-gap = Q(best_direction) - Q(Flat). Higher = more conviction.
-     * Used by env_step to scale target_position. */
+     * Used by env_step to scale target_position.
+     *
+     * D7/N7: During contrarian override, the action selected is argmin rather
+     * than argmax. The raw Q(best) - Q(Flat) gap would mislead the position
+     * sizer into full conviction for what's actually the "deliberately wrong"
+     * action. Zero out the gap so contrarian trades size minimal. */
     if (out_q_gaps != NULL) {
-        int flat_idx = 3; /* Flat = index 3 for direction(4): Short(0)/Hold(1)/Long(2)/Flat(3) */
-        float q_best = q_b0[argmax_n(q_b0, b0_size)];
-        float q_flat = q_b0[flat_idx];
-        float gap = q_best - q_flat;
-        out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f;
+        if (contrarian_active != 0) {
+            out_q_gaps[i] = 0.0f;
+        } else {
+            int flat_idx = 3; /* Flat = index 3 for direction(4): Short(0)/Hold(1)/Long(2)/Flat(3) */
+            float q_best = q_b0[argmax_n(q_b0, b0_size)];
+            float q_flat = q_b0[flat_idx];
+            float gap = q_best - q_flat;
+            out_q_gaps[i] = (gap > 0.0f) ? gap : 0.0f;
+        }
     }
 }
 
@@ -1111,7 +1131,8 @@ extern "C" __global__ void experience_env_step(
     float price_confirm_weight,      /* micro-reward: price confirmation weight (default 0.5) */
     float book_aggression_weight,    /* micro-reward: book aggression weight (default 0.3) */
     float hold_quality_weight,       /* micro-reward: retrospective hold quality weight (default 0.2) */
-    float micro_reward_temp          /* micro-reward: tanh temperature (default 3.0) */
+    float micro_reward_temp,         /* micro-reward: tanh temperature (default 3.0) */
+    float cf_ratio                   /* D4/N4: adaptive counterfactual flip ratio [0,1]; 0.5=healthy, 0.8=collapsed */
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -1158,7 +1179,7 @@ extern "C" __global__ void experience_env_step(
     int do_flip = 0;
     {
         float flip_u = philox_uniform(i, bar_idx + 8888, 7777);
-        do_flip = (flip_u < 0.5f) ? 1 : 0;
+        do_flip = (flip_u < cf_ratio) ? 1 : 0;
     }
     if (do_flip) {
         for (int f = 0; f < 16 && f < state_dim; f++) {
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index bc6fedef0..2e12e2a1f 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -83,7 +83,8 @@ const OFI_EMBED_DIM: usize = 10;     // OFI embedding width appended to history
 
 /// Introspective State Vector (ISV) configuration.
 const ISV_K: usize = 4;   // Temporal ISV history length
-const ISV_DIM: usize = 12; // ISV signal count (8 core + 4 regime awareness)
+const ISV_DIM: usize = 13; // ISV signal count (8 core + 4 regime awareness + 1 learning_health)
+pub const LEARNING_HEALTH_INDEX: usize = 12;
 const ISV_EMB_DIM: usize = 8; // ISV embedding output dimension (FC2 output, gate/gamma input)
 /// First ISV tensor index in the flat param buffer.
 /// ISV weights (68-79) are online-only — NOT synced to the target network.
@@ -1562,11 +1563,11 @@ pub struct GpuDqnTrainer {
     atom_util_scratch_dev_ptr: u64,
 
     // ── ISV core buffers (pinned device-mapped, GPU read/write) ──
-    isv_signals_pinned: *mut f32,        // [ISV_DIM = 12]
+    isv_signals_pinned: *mut f32,        // [ISV_DIM = 13]
     isv_signals_dev_ptr: u64,
-    isv_history_pinned: *mut f32,        // [ISV_K * ISV_DIM = 48]
+    isv_history_pinned: *mut f32,        // [ISV_K * ISV_DIM = 52]
     isv_history_dev_ptr: u64,
-    isv_decay_pinned: *mut f32,          // [ISV_DIM = 12] learned decay weights
+    isv_decay_pinned: *mut f32,          // [ISV_DIM = 13] learned decay weights
     isv_decay_dev_ptr: u64,
     lagged_td_error_pinned: *mut f32,    // [1] recursive confidence target
     lagged_td_error_dev_ptr: u64,
@@ -1658,6 +1659,41 @@ pub struct GpuDqnTrainer {
     tau_final: f32,
     /// Total number of steps over which tau is cosine-annealed.
     tau_anneal_steps: i32,
+
+    /// Last effective cql_alpha applied (base × (1−regime_stability) × health).
+    /// Populated by apply_cql_gradient. Initialized to 0.0.
+    pub(crate) last_cql_alpha_eff: f32,
+
+    /// Last effective tau value (cosine-annealed, potentially floored by 0.01×(1-health)).
+    /// Populated by `apply_health_coupled_tau_floor`. Initialized to 0.0.
+    pub(crate) last_tau_eff: f32,
+
+    /// B3/G4: Last health-scaled Expected SARSA tau factor (1.0 at full health, 6.0 at collapse).
+    /// Populated by `launch_c51_loss`. Initialized to 0.0.
+    pub(crate) last_sarsa_tau_factor: f32,
+
+    /// B4/G5: Last adaptive gradient budget for IQN (0.10 + 0.30×health).
+    /// Populated by `FusedTrainingCtx::compute_adaptive_budgets`. Default: health=1 → 0.40.
+    pub(crate) last_iqn_budget_eff: f32,
+    /// B4/G5: Last adaptive gradient budget for CQL (0.10×(1−regime)×health).
+    /// Populated by `FusedTrainingCtx::compute_adaptive_budgets`. Default: 0.00 (regime_stability=1).
+    pub(crate) last_cql_budget_eff: f32,
+    /// B4/G5: Last adaptive gradient budget for C51 (1−iqn−cql−ens).
+    /// Populated by `FusedTrainingCtx::compute_adaptive_budgets`. Default: health=1 → 0.55.
+    pub(crate) last_c51_budget_eff: f32,
+    /// B4/G5: Adaptive gradient budget for ensemble (constant 0.05).
+    /// Populated by `FusedTrainingCtx::compute_adaptive_budgets`.
+    pub(crate) last_ens_budget_eff: f32,
+
+    /// C4/P4: Last effective gamma applied to C51 Bellman projection.
+    /// Formula: gamma_base + 0.005×(regime_stability − 0.5) − 0.05×(1 − health),
+    /// clamped to [0.9, 0.995]. Populated by `apply_adaptive_gamma`. Default: 0.99.
+    pub(crate) last_gamma_eff: f32,
+
+    /// D1/N1: Ring buffer of best-health weight snapshots for temporal self-distillation.
+    pub(crate) q_snapshots: crate::cuda_pipeline::q_snapshot::SnapshotRing,
+    /// D1/N1: Whether distillation gradient was applied in the most recent epoch boundary.
+    pub(crate) last_distill_active: bool,
 }
 
 impl GpuDqnTrainer {
@@ -4051,6 +4087,7 @@ impl GpuDqnTrainer {
         &mut self,
         iqn_d_h_s2_ptr: u64,
         _online_dueling: &mut DuelingWeightSet,
+        iqn_budget: f32,
     ) -> Result<(), MLError> {
         let b = self.config.batch_size;
         let _eg = EventTrackingGuard::new(self.stream.context());
@@ -4171,13 +4208,14 @@ impl GpuDqnTrainer {
             )?;
         }
 
-        // ── 7. Plain SAXPY: grad_buf[trunk] += iqn_lambda * scratch ──
+        // ── 7. Plain SAXPY: grad_buf[trunk] += iqn_lambda * iqn_budget * scratch ──
         // Adaptive lambda: scales with IQN loss readiness (0 when uncertain, 1 when converged).
+        // B4/G5: further scaled by iqn_budget (0.10+0.30×health).
         // Uses aux_child-specific handle — saxpy_f32_kernel is captured in forward_child.
         {
             let scratch_ptr = self.ptrs.iqn_trunk_m;
             let grad_ptr = self.ptrs.grad_buf;
-            let scale = self.config.iqn_lambda * self.iqn_readiness;
+            let scale = self.config.iqn_lambda * self.iqn_readiness * iqn_budget;
             let n_i32 = trunk_grad_total as i32;
             let blocks = ((trunk_grad_total + 255) / 256) as u32;
             unsafe {
@@ -4573,7 +4611,25 @@ impl GpuDqnTrainer {
         let d_v_ptr = self.cql_d_value_logits.raw_ptr();
         let d_adv_ptr = self.cql_d_adv_logits.raw_ptr();
 
-        let cql_alpha = self.config.cql_alpha;
+        // B1/G2: Adaptive cql_alpha.
+        //  - health from ISV[LEARNING_HEALTH_INDEX] (0.5 fallback if pinned ptr null)
+        //  - regime_stability from ISV[11]
+        //  - cql_alpha_eff = base × (1 − regime_stability) × health
+        //    Collapse → 0 (CQL off); volatile+healthy → full; stable+healthy → 0.
+        let (health, regime_stability) = if self.isv_signals_pinned.is_null() {
+            (0.5_f32, 0.5_f32)
+        } else {
+            unsafe {
+                let h = (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0);
+                let s = (*self.isv_signals_pinned.add(11)).clamp(0.0, 1.0);
+                (h, s)
+            }
+        };
+        let cql_alpha = self.config.cql_alpha * (1.0 - regime_stability) * health;
+
+        // Cache for logging via FusedTrainingCtx / DQNTrainer.
+        self.last_cql_alpha_eff = cql_alpha;
+
         let n_i32 = b as i32;
         let na_i32 = na as i32;
         let b0_i32 = b0 as i32;
@@ -4740,14 +4796,15 @@ impl GpuDqnTrainer {
     ///
     /// Called after `apply_cql_gradient` populated `cql_grad_scratch`.
     /// No per-component clip — single global clip in Adam handles safety.
-    pub fn apply_cql_saxpy(&mut self) -> Result<(), MLError> {
+    pub fn apply_cql_saxpy(&mut self, cql_budget: f32) -> Result<(), MLError> {
         let total = self.total_params as i32;
         let blocks = self.grad_norm_blocks as u32;
         let scratch_ptr = self.ptrs.cql_grad_scratch;
         let grad_ptr = self.ptrs.grad_buf;
-        let alpha = 1.0_f32;
+        // B4/G5: scale CQL gradient contribution by adaptive budget (0.10×(1−regime)×health).
+        let alpha = cql_budget;
 
-        // Plain SAXPY: grad_buf += 1.0 * cql_scratch
+        // Plain SAXPY: grad_buf += cql_budget * cql_scratch
         // Uses aux_child-specific handle — saxpy_f32_kernel is captured in forward_child.
         unsafe {
             self.stream
@@ -6924,6 +6981,9 @@ impl GpuDqnTrainer {
         let tau_final = 0.001_f32;
         let tau_anneal_steps: i32 = 100_000;
 
+        // D1/N1: clone Arc<CudaStream> before it is moved into the struct literal.
+        let stream_for_snapshots = Arc::clone(&stream);
+
         Ok(Self {
             config,
             stream,
@@ -7418,6 +7478,19 @@ impl GpuDqnTrainer {
             tau_init,
             tau_final,
             tau_anneal_steps,
+            last_cql_alpha_eff: 0.0,
+            last_tau_eff: 0.0,
+            last_sarsa_tau_factor: 0.0,
+            last_iqn_budget_eff: 0.40,
+            last_cql_budget_eff: 0.00,
+            last_c51_budget_eff: 0.55,
+            last_ens_budget_eff: 0.05,
+            last_gamma_eff: 0.99,
+            q_snapshots: crate::cuda_pipeline::q_snapshot::SnapshotRing::new(
+                stream_for_snapshots,
+                total_params,
+            ),
+            last_distill_active: false,
         })
     }
 
@@ -7599,6 +7672,223 @@ impl GpuDqnTrainer {
         }
     }
 
+    /// B4/G5: Read ISV health index and regime stability from pinned host memory.
+    /// Returns (health [0,1], regime_stability [0,1]). Falls back to (0.5, 0.5) if
+    /// pinned pointer is null.
+    pub(crate) fn read_isv_health_and_regime(&self) -> (f32, f32) {
+        if self.isv_signals_pinned.is_null() { return (0.5_f32, 0.5_f32); }
+        unsafe {
+            let h = (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0);
+            let s = (*self.isv_signals_pinned.add(11)).clamp(0.0, 1.0);
+            (h, s)
+        }
+    }
+
+    /// B2/G3: Apply a health-coupled minimum to a scheduled tau value.
+    /// Returns `max(tau_scheduled, 0.01 * (1.0 - health))`. During collapse
+    /// (health≈0), the floor rises to 0.01 to accelerate target-network adaptation.
+    /// Caches the effective value for HEALTH_DIAG logging.
+    pub fn apply_health_coupled_tau_floor(&mut self, tau_scheduled: f64) -> f64 {
+        let health = if self.isv_signals_pinned.is_null() {
+            0.5_f32
+        } else {
+            unsafe { (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0) }
+        };
+        let tau_floor = (0.01 * (1.0 - health)) as f64;
+        let tau_eff = tau_scheduled.max(tau_floor);
+        self.last_tau_eff = tau_eff as f32;
+        tau_eff
+    }
+
+    /// C4/P4: Apply temporal-coupled adjustment to the scheduled gamma.
+    /// Formula: gamma_base + 0.005×(regime_stability − 0.5) − 0.05×(1 − health),
+    /// clamped to [0.9, 0.995]. Updates `adaptive_gamma` and caches the effective
+    /// value in `last_gamma_eff` for HEALTH_DIAG logging.
+    pub fn apply_adaptive_gamma(&mut self, gamma_base: f32) -> f32 {
+        let (health, regime_stability) = self.read_isv_health_and_regime();
+        let gamma_eff = gamma_base
+            + 0.005_f32 * (regime_stability - 0.5)
+            - 0.05_f32 * (1.0 - health);
+        let gamma_clamped = gamma_eff.clamp(0.9, 0.995);
+        self.adaptive_gamma = gamma_clamped;
+        self.last_gamma_eff = gamma_clamped;
+        gamma_clamped
+    }
+
+    // ── D1/N1: Temporal self-distillation ─────────────────────────────────
+
+    /// D1/N1: Whether distillation gradient was applied in the most recent epoch.
+    pub fn last_distill_active(&self) -> bool {
+        self.last_distill_active
+    }
+
+    /// D1/N1: Snapshot current params_buf if health and q_gap qualify.
+    /// Avoids split-borrow between params_buf and q_snapshots.
+    pub fn maybe_snapshot_params(&mut self, health: f32, q_gap: f32, epoch: u32) -> Result<bool, MLError> {
+        use cudarc::driver::{DevicePtr, DevicePtrMut};
+        use crate::cuda_pipeline::q_snapshot::SNAPSHOT_HEALTH_THRESHOLD;
+
+        if health < SNAPSHOT_HEALTH_THRESHOLD {
+            return Ok(false);
+        }
+
+        let ring = &mut self.q_snapshots;
+        if ring.snapshots.len() >= crate::cuda_pipeline::q_snapshot::MAX_SNAPSHOTS {
+            let worst_q_gap = ring.snapshots.iter().map(|s| s.q_gap).fold(f32::INFINITY, f32::min);
+            if q_gap <= worst_q_gap {
+                return Ok(false);
+            }
+        }
+
+        let n = ring.param_count();
+        let num_bytes = n * std::mem::size_of::<f32>();
+        let mut new_weights = self.stream.alloc_zeros::<f32>(n)
+            .map_err(|e| MLError::ModelError(format!("snapshot alloc: {e}")))?;
+        {
+            let (src_ptr, _sg) = self.params_buf.device_ptr(&self.stream);
+            let (dst_ptr, _dg) = new_weights.device_ptr_mut(&self.stream);
+            #[allow(unsafe_code)]
+            unsafe {
+                cudarc::driver::result::memcpy_dtod_async(
+                    dst_ptr, src_ptr, num_bytes, self.stream.cu_stream(),
+                ).map_err(|e| MLError::ModelError(format!("snapshot dtod: {e}")))?;
+            }
+        }
+
+        let ring = &mut self.q_snapshots;
+        if ring.snapshots.len() >= crate::cuda_pipeline::q_snapshot::MAX_SNAPSHOTS {
+            let worst_idx = ring.snapshots.iter()
+                .enumerate()
+                .min_by(|(_, a), (_, b)| a.q_gap.partial_cmp(&b.q_gap).unwrap())
+                .map(|(i, _)| i)
+                .unwrap();
+            ring.snapshots.swap_remove(worst_idx);
+        }
+
+        ring.snapshots.push(crate::cuda_pipeline::q_snapshot::QSnapshot {
+            weights: new_weights,
+            health,
+            q_gap,
+            epoch,
+        });
+        Ok(true)
+    }
+
+    /// D1/N1: Apply distillation gradient — pulls current weights toward the best
+    /// snapshot with strength `0.1 × (1 − health)`. Returns true if applied.
+    ///
+    /// Implemented as two SAXPY calls into grad_buf (ungraphed, stream-ordered):
+    ///   grad += alpha * params      (adds current)
+    ///   grad += -alpha * best_snap  (subtracts best)
+    /// Net effect: grad += alpha * (params - best_snap)
+    ///
+    /// Uses `saxpy_f32_aux` (dqn_saxpy_f32_kernel, aux_child handle) — same handle
+    /// used by `apply_cql_saxpy`, safe for ungraphed injection before graph_adam.
+    pub fn apply_distillation_gradient(&mut self) -> Result<bool, MLError> {
+        let (health, _) = self.read_isv_health_and_regime();
+        let distill_weight = 0.1_f32 * (1.0 - health).max(0.0);
+        if distill_weight < 0.01 {
+            self.last_distill_active = false;
+            return Ok(false);
+        }
+        let best_ptr = match self.q_snapshots.best_raw_ptr() {
+            Some(p) => p,
+            None => {
+                self.last_distill_active = false;
+                return Ok(false);
+            }
+        };
+
+        let grad_ptr = self.ptrs.grad_buf;
+        let params_ptr = self.ptrs.params_ptr;
+        let n = self.total_params as i32;
+        let blocks = self.grad_norm_blocks as u32;
+
+        unsafe {
+            // grad += distill_weight * params_buf
+            self.stream
+                .launch_builder(&self.saxpy_f32_aux)
+                .arg(&grad_ptr)
+                .arg(&params_ptr)
+                .arg(&distill_weight)
+                .arg(&n)
+                .launch(LaunchConfig {
+                    grid_dim: (blocks, 1, 1),
+                    block_dim: (256, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .map_err(|e| MLError::ModelError(format!("distill saxpy (+params): {e}")))?;
+
+            // grad += (-distill_weight) * best_snapshot
+            let neg_weight = -distill_weight;
+            self.stream
+                .launch_builder(&self.saxpy_f32_aux)
+                .arg(&grad_ptr)
+                .arg(&best_ptr)
+                .arg(&neg_weight)
+                .arg(&n)
+                .launch(LaunchConfig {
+                    grid_dim: (blocks, 1, 1),
+                    block_dim: (256, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .map_err(|e| MLError::ModelError(format!("distill saxpy (-best): {e}")))?;
+        }
+
+        self.last_distill_active = true;
+        Ok(true)
+    }
+
+    // ── A3: LearningHealth signal writers / readers ───────────────────────
+
+    /// Write a scalar into the ISV pinned buffer at the given index.
+    /// No-op if the pinned pointer is null or index >= ISV_DIM.
+    pub fn write_isv_signal_at(&self, index: usize, value: f32) {
+        if self.isv_signals_pinned.is_null() { return; }
+        if index >= ISV_DIM { return; }
+        unsafe { *self.isv_signals_pinned.add(index) = value; }
+    }
+
+    /// Read atom utilization from q_readback_pinned[6] (one-step lag, written by
+    /// reduce_current_q_stats). Returns 0.0 if pointer is null.
+    pub fn read_atom_utilization(&self) -> f32 {
+        if self.q_readback_pinned.is_null() { return 0.0; }
+        unsafe { (*self.q_readback_pinned.add(6)).clamp(0.0, 1.0) }
+    }
+
+    /// Coarse Q-collapse detector derived from the single-sample Q readback at
+    /// `q_readback_pinned[7..7+total_actions]`.
+    ///
+    /// Returns a LARGE value (~100) when all Q values are (nearly) equal — i.e.
+    /// rank-1 / Q-collapse — and a value close to 1.0 when the Q spread is wide.
+    /// Intentionally coarse: a full SVD across a batch of Q samples would require
+    /// a separate larger buffer and is out of scope for the host-side health path.
+    ///
+    /// The downstream `NormalizedComponents::from_raw` thresholds are calibrated
+    /// for this behaviour (`spectral_gap_norm` drops toward 0 as this value rises).
+    pub fn compute_q_spectral_gap(&self) -> f32 {
+        if self.q_readback_pinned.is_null() { return 1.0; }
+        // Layout: q_readback_pinned[0..7] = stats, [7..7+total_actions] = per-action Q values.
+        let total_actions = self.total_actions() as usize;
+        if total_actions == 0 { return 1.0; }
+        unsafe {
+            let base = self.q_readback_pinned.add(7);
+            let mut max = f32::NEG_INFINITY;
+            let mut min = f32::INFINITY;
+            for i in 0..total_actions {
+                let v = *base.add(i);
+                if v.is_finite() {
+                    if v > max { max = v; }
+                    if v < min { min = v; }
+                }
+            }
+            let range = (max - min).abs();
+            // Range near zero means all actions have equal Q → rank-1-like collapse;
+            // return a large value to signal pathology via spectral_gap_norm ≈ 0.
+            if range < 1e-6 { 100.0 } else { 1.0 + range.recip() }
+        }
+    }
+
     /// Read eval_q_std_ema.
     pub fn eval_q_std_ema(&self) -> f32 { self.eval_q_std_ema }
 
@@ -10608,7 +10898,7 @@ impl GpuDqnTrainer {
     /// Reads: on_v_logits, on_b_logits, tg_v_logits, tg_b_logits,
     ///        on_next_v_logits, on_next_b_logits (Double DQN).
     /// Writes: per_sample_loss, td_errors, total_loss, save_current_lp, save_projected.
-    fn launch_c51_loss(&self) -> Result<(), MLError> {
+    fn launch_c51_loss(&mut self) -> Result<(), MLError> {
         let kernel = &self.c51_loss_kernel;
 
         let b = self.config.batch_size;
@@ -10674,6 +10964,14 @@ impl GpuDqnTrainer {
         let shmem_floats = na + na + max_branch * na + na + na + na + 8;
         let shmem_bytes = (shmem_floats * std::mem::size_of::<f32>()) as u32;
 
+        // B3/G4: mirror the kernel's tau_health_factor on host for logging.
+        let current_health = if self.isv_signals_pinned.is_null() {
+            0.5_f32
+        } else {
+            unsafe { (*self.isv_signals_pinned.add(LEARNING_HEALTH_INDEX)).clamp(0.0, 1.0) }
+        };
+        self.last_sarsa_tau_factor = 1.0 + 5.0 * (1.0 - current_health);
+
         unsafe {
             self.stream
                 .launch_builder(kernel)
@@ -10737,6 +11035,8 @@ impl GpuDqnTrainer {
                 .arg(&atom_positions_buf_ptr)
                 // ── Per-sample CVaR alpha from learned risk branch ──
                 .arg(&cvar_alpha_buf_ptr)
+                // ── B3/G4: ISV signals for health-scaled Expected SARSA temperature ──
+                .arg(&self.isv_signals_dev_ptr)
                 .launch(LaunchConfig {
                     grid_dim: (b as u32, 1, 1),
                     block_dim: (256, 1, 1),
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 01a423e26..c3bc6126f 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -732,6 +732,17 @@ pub struct GpuExperienceCollector {
     hindsight_relabel_kernel: CudaFunction,
     /// v8: TD(lambda) return computation kernel (loaded from nstep cubin).
     td_lambda_kernel: CudaFunction,
+
+    /// D4/N4: cached learning_health value from the trainer (updated once per epoch).
+    learning_health_cache: f32,
+    /// D4/N4: last effective cf_ratio computed from learning_health_cache.
+    last_cf_ratio_eff: f32,
+
+    /// D7/N7: Contrarian override flag cached from trainer. When non-zero, the
+    /// experience kernel negates Q values before Boltzmann softmax, converting
+    /// argmax-favoring sampling to argmin-favoring. Used briefly to escape Q-uniform
+    /// attractors when the policy is systematically anti-correlated with market.
+    contrarian_active_cache: u8,
 }
 
 impl Drop for GpuExperienceCollector {
@@ -1298,6 +1309,9 @@ impl GpuExperienceCollector {
             difficulty_scores_kernel,
             hindsight_relabel_kernel,
             td_lambda_kernel,
+            learning_health_cache: 1.0,  // D4/N4: assume healthy at construction
+            last_cf_ratio_eff: 0.5,       // D4/N4: standard cf_ratio at healthy state
+            contrarian_active_cache: 0,   // D7/N7: off by default
             reward_rank_kernel,
             reward_compute_abs_sharpe_kernel,
             bitonic_sort_step_kernel,
@@ -1340,6 +1354,34 @@ impl GpuExperienceCollector {
         self.readiness_dev_ptr = dev_ptr;
     }
 
+    /// D4/N4: propagate learning health from the trainer (called once per epoch).
+    /// The next call to collect_experiences_gpu will compute cf_ratio from this value.
+    pub fn set_learning_health(&mut self, health: f32) {
+        self.learning_health_cache = health.clamp(0.0, 1.0);
+    }
+
+    /// D4/N4: return the last effective cf_ratio (set during collect_experiences_gpu).
+    pub fn last_cf_ratio_eff(&self) -> f32 {
+        self.last_cf_ratio_eff
+    }
+
+    /// D7/N7: propagate contrarian flag from trainer (called once per epoch when
+    /// the state machine updates). When active, the action-selection kernel negates
+    /// Q values before Boltzmann softmax → argmin-favoring sampling.
+    pub fn set_contrarian_active(&mut self, active: bool) {
+        self.contrarian_active_cache = if active { 1 } else { 0 };
+    }
+
+    /// D7/N7: return the cached contrarian flag.
+    pub fn contrarian_active(&self) -> u8 {
+        self.contrarian_active_cache
+    }
+
+    /// D4/N4: compute cf_ratio from cached health.
+    fn current_cf_ratio(&self) -> f32 {
+        (0.5_f32 + 0.3_f32 * (1.0 - self.learning_health_cache)).clamp(0.0, 1.0)
+    }
+
     /// Set CVaR position scaling from IQN dual-head.
     /// The device pointer will be passed to the env_step kernel.
     /// Call with 0 to disable (NULL pointer = no scaling).
@@ -2348,6 +2390,7 @@ impl GpuExperienceCollector {
                     .arg(&(t as i32))                   // timestep for stateless Philox RNG
                     .arg(&self.per_sample_epsilon_ptr)   // IQL expectile gap epsilon (0=cosine schedule)
                     .arg(&self.isv_signals_dev_ptr)      // ISV signals for adaptive hold (0=NULL=static)
+                    .arg(&(self.contrarian_active_cache as i32))  // D7/N7: Q-negation flag
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_action_select t={t}: {e}"
@@ -2386,6 +2429,9 @@ impl GpuExperienceCollector {
 
             // ── 5. Environment step (reward v5: trade-aware hybrid) ──────
             // max_pos already defined above (action_select block)
+            // D4/N4: adaptive counterfactual ratio — 0.5 at healthy, 0.8 at collapse.
+            let cf_ratio = self.current_cf_ratio();
+            self.last_cf_ratio_eff = cf_ratio;
             let tx_cost = config.tx_cost_multiplier;
             let l_i32 = timesteps as i32;
             let total_actions_env = (self.branch_sizes[0] + self.branch_sizes[1]
@@ -2467,6 +2513,7 @@ impl GpuExperienceCollector {
                     .arg(&config.book_aggression_weight)    // micro-reward: book aggression weight
                     .arg(&config.hold_quality_weight)       // micro-reward: hold quality weight
                     .arg(&config.micro_reward_temp)         // micro-reward: tanh temperature
+                    .arg(&cf_ratio)                         // D4/N4: adaptive counterfactual flip ratio
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_env_step t={t}: {e}"
diff --git a/crates/ml/src/cuda_pipeline/learning_health.rs b/crates/ml/src/cuda_pipeline/learning_health.rs
new file mode 100644
index 000000000..67389d251
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/learning_health.rs
@@ -0,0 +1,186 @@
+//! LearningHealth signal — unified training health metric.
+//!
+//! Composes 7 normalized components into a scalar [0, 1] that senses training
+//! collapse and drives adaptive hyperparameters.
+
+/// Smoothstep function: smooth transition from 0 at `edge0` to 1 at `edge1`.
+/// Standard GLSL smoothstep formula.
+#[inline]
+pub fn smoothstep(edge0: f32, edge1: f32, x: f32) -> f32 {
+    if edge1 == edge0 { return if x >= edge0 { 1.0 } else { 0.0 }; }
+    let t = ((x - edge0) / (edge1 - edge0)).clamp(0.0, 1.0);
+    t * t * (3.0 - 2.0 * t)
+}
+
+/// Raw input signals for LearningHealth composition.
+#[derive(Debug, Clone, Copy, Default)]
+pub struct HealthComponents {
+    pub q_gap: f32,                // max - 2nd max of Q values across actions
+    pub q_var: f32,                // variance of Q across actions
+    pub atom_util: f32,            // C51 atom utilization fraction [0, 1]
+    pub grad_norm: f32,            // gradient L2 norm
+    pub ens_disagreement: f32,     // mean pairwise diff across ensemble heads
+    pub grad_consistency: f32,     // cosine similarity of successive gradients
+    pub spectral_gap: f32,         // sigma_1 / sigma_2 from SVD of Q matrix
+}
+
+/// Computed normalized signals in [0, 1].
+#[derive(Debug, Clone, Copy, Default)]
+pub struct NormalizedComponents {
+    pub q_gap_norm: f32,
+    pub q_var_norm: f32,
+    pub atom_util_norm: f32,
+    pub grad_stable: f32,
+    pub ens_agree: f32,
+    pub grad_consistency_norm: f32,
+    pub spectral_gap_norm: f32,
+}
+
+impl NormalizedComponents {
+    /// Normalize raw signals into [0, 1] via smoothstep curves. Thresholds encode
+    /// empirical collapse/health boundaries observed in production runs:
+    /// - q_gap: [0.01, 0.5] — below 0.01 the argmax is indistinguishable from noise; above 0.5 actions are well-separated
+    /// - q_var: [0.001, 0.1] — tight variance indicates Q-collapse to uniform
+    /// - atom_util: [0.2, 0.7] — C51 healthy when ≥ 70% of atoms carry >1% probability
+    /// - grad_norm: [10, 100] — inverted; higher grad_norm is less stable
+    /// - ens_disagreement: directly used; higher disagreement means less agreement on Q
+    /// - grad_consistency: [-0.2, 0.5] — cosine sim between successive gradient vectors
+    /// - spectral_gap: [2, 10] — inverted; sigma_1/sigma_2 above 10 means rank-1 collapse
+    pub fn from_raw(raw: &HealthComponents) -> Self {
+        Self {
+            q_gap_norm: smoothstep(0.01, 0.5, raw.q_gap),
+            q_var_norm: smoothstep(0.001, 0.1, raw.q_var),
+            atom_util_norm: smoothstep(0.2, 0.7, raw.atom_util),
+            grad_stable: 1.0 - smoothstep(10.0, 100.0, raw.grad_norm),
+            ens_agree: (1.0 - raw.ens_disagreement).clamp(0.0, 1.0),
+            grad_consistency_norm: smoothstep(-0.2, 0.5, raw.grad_consistency),
+            spectral_gap_norm: 1.0 - smoothstep(2.0, 10.0, raw.spectral_gap),
+        }
+    }
+
+    /// Weighted sum per spec Layer 1 composition.
+    pub fn compose(&self) -> f32 {
+        0.25 * self.q_gap_norm
+        + 0.15 * self.q_var_norm
+        + 0.15 * self.atom_util_norm
+        + 0.15 * self.grad_stable
+        + 0.10 * self.ens_agree
+        + 0.10 * self.grad_consistency_norm
+        + 0.10 * self.spectral_gap_norm
+    }
+}
+
+/// LearningHealth state: EMA-smoothed health value with bounds and warmup.
+#[derive(Debug, Clone)]
+pub struct LearningHealth {
+    /// Current EMA-smoothed health [0.2, 0.95].
+    pub value: f32,
+    /// Epoch counter — enables warmup mode for first N epochs.
+    pub epoch: u32,
+    /// Latest normalized components (for logging).
+    pub components: NormalizedComponents,
+    /// EMA smoothing factor (new sample weight).
+    ema_alpha: f32,
+    /// Warmup epochs — health clamped to 0.5 during warmup.
+    warmup_epochs: u32,
+}
+
+impl LearningHealth {
+    pub fn new() -> Self {
+        Self {
+            value: 0.5, // neutral start
+            epoch: 0,
+            components: NormalizedComponents::default(),
+            ema_alpha: 0.3,
+            warmup_epochs: 3,
+        }
+    }
+
+    /// Update health from new raw components. Returns updated health value.
+    pub fn update(&mut self, raw: &HealthComponents) -> f32 {
+        self.components = NormalizedComponents::from_raw(raw);
+        let new_sample = self.components.compose();
+
+        if self.epoch < self.warmup_epochs {
+            // Warmup: pin to 0.5, allow components to accumulate
+            self.value = 0.5;
+        } else {
+            // EMA update
+            let ema = (1.0 - self.ema_alpha) * self.value + self.ema_alpha * new_sample;
+            // Clamp to [0.2, 0.95] — hard floor prevents total adaptive-mechanism shutdown
+            // during transient collapse; hard ceiling keeps CQL/budget levers from pegging
+            // at extremes when health appears perfect.
+            self.value = ema.clamp(0.2, 0.95);
+        }
+
+        self.epoch += 1;
+        self.value
+    }
+}
+
+impl Default for LearningHealth {
+    fn default() -> Self { Self::new() }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_smoothstep_boundaries() {
+        assert_eq!(smoothstep(0.0, 1.0, -0.5), 0.0);
+        assert_eq!(smoothstep(0.0, 1.0, 1.5), 1.0);
+        assert!((smoothstep(0.0, 1.0, 0.5) - 0.5).abs() < 0.001);
+    }
+
+    #[test]
+    fn test_collapsed_inputs_give_low_health() {
+        let raw = HealthComponents {
+            q_gap: 0.01,          // collapsed
+            q_var: 0.001,         // collapsed
+            atom_util: 0.1,       // under-utilized
+            grad_norm: 50_000.0,  // exploding
+            ens_disagreement: 0.0,// everyone agrees (on uniform)
+            grad_consistency: -0.3, // gradients contradict
+            spectral_gap: 50.0,   // rank 1
+        };
+        let mut health = LearningHealth::new();
+        // Skip warmup
+        for _ in 0..5 { health.update(&raw); }
+        assert!(health.value <= 0.3, "Collapsed state should give health <= 0.3, got {}", health.value);
+    }
+
+    #[test]
+    fn test_healthy_inputs_give_high_health() {
+        let raw = HealthComponents {
+            q_gap: 1.0,            // strong differentiation
+            q_var: 0.5,            // good variance
+            atom_util: 0.9,        // atoms fully utilized
+            grad_norm: 5.0,        // stable
+            ens_disagreement: 0.1, // moderate diversity
+            grad_consistency: 0.9, // consistent direction
+            spectral_gap: 1.5,     // balanced rank
+        };
+        let mut health = LearningHealth::new();
+        for _ in 0..5 { health.update(&raw); }
+        assert!(health.value >= 0.7, "Healthy state should give health >= 0.7, got {}", health.value);
+    }
+
+    #[test]
+    fn test_warmup_clamps_to_neutral() {
+        let raw = HealthComponents {
+            q_gap: 0.01, q_var: 0.001, atom_util: 0.1,
+            grad_norm: 50_000.0, ens_disagreement: 0.0,
+            grad_consistency: -0.5, spectral_gap: 100.0,
+        };
+        let mut health = LearningHealth::new();
+        // First 3 updates are warmup — should stay at 0.5
+        for _ in 0..3 {
+            let v = health.update(&raw);
+            assert_eq!(v, 0.5, "Warmup should pin health to 0.5");
+        }
+        // 4th update starts EMA
+        health.update(&raw);
+        assert!(health.value < 0.5, "Post-warmup should start dropping with collapsed inputs");
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/meta_q_network.rs b/crates/ml/src/cuda_pipeline/meta_q_network.rs
new file mode 100644
index 000000000..754d4eacc
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/meta_q_network.rs
@@ -0,0 +1,234 @@
+//! N8: Meta-Q network — predicts collapse probability K epochs ahead.
+//!
+//! Small MLP (7→32→16→1 sigmoid) trained online via BCE loss on historical
+//! (health_components, future_collapse) pairs. Output is logged as a leading
+//! indicator — does NOT feed into LearningHealth until validated.
+
+use std::collections::VecDeque;
+
+pub const META_Q_INPUT_DIM: usize = 7;
+pub const META_Q_H1: usize = 32;
+pub const META_Q_H2: usize = 16;
+pub const META_Q_LOOKAHEAD: usize = 5;
+
+#[derive(Debug, Clone)]
+pub struct MetaQInput {
+    pub q_gap_ema: f32,
+    pub log_grad_norm: f32,
+    pub atom_util: f32,
+    pub ens_disagreement: f32,
+    pub regime_stability: f32,
+    pub spectral_gap_norm: f32,
+    pub grad_consistency: f32,
+}
+
+impl MetaQInput {
+    pub fn to_vec(&self) -> [f32; META_Q_INPUT_DIM] {
+        [
+            self.q_gap_ema,
+            self.log_grad_norm,
+            self.atom_util,
+            self.ens_disagreement,
+            self.regime_stability,
+            self.spectral_gap_norm,
+            self.grad_consistency,
+        ]
+    }
+}
+
+/// Small MLP on host (CPU). n_params ≈ 7×32 + 32×16 + 16×1 + biases = ~800.
+#[derive(Debug)]
+pub struct MetaQNetwork {
+    w1: Vec<f32>, b1: Vec<f32>,
+    w2: Vec<f32>, b2: Vec<f32>,
+    w3: Vec<f32>, b3: Vec<f32>,
+    m_w1: Vec<f32>, v_w1: Vec<f32>,
+    m_w2: Vec<f32>, v_w2: Vec<f32>,
+    m_w3: Vec<f32>, v_w3: Vec<f32>,
+    lr: f32,
+    step: u32,
+    buffer: VecDeque<(MetaQInput, f32)>,
+    buffer_cap: usize,
+}
+
+impl MetaQNetwork {
+    pub fn new() -> Self {
+        // Use rand with a fixed seed for determinism of initial weights.
+        // fastrand is already a dep in several places; fall back to a simple
+        // deterministic LCG if fastrand is unavailable.
+        let mut rng = SimpleLcg::new(0xDEADBEEF);
+        let xavier = |rng: &mut SimpleLcg, fan_in: usize, fan_out: usize| -> f32 {
+            let range = (6.0 / (fan_in + fan_out) as f32).sqrt();
+            (rng.next_f32() * 2.0 - 1.0) * range
+        };
+        let w1: Vec<f32> = (0..META_Q_INPUT_DIM * META_Q_H1)
+            .map(|_| xavier(&mut rng, META_Q_INPUT_DIM, META_Q_H1)).collect();
+        let w2: Vec<f32> = (0..META_Q_H1 * META_Q_H2)
+            .map(|_| xavier(&mut rng, META_Q_H1, META_Q_H2)).collect();
+        let w3: Vec<f32> = (0..META_Q_H2)
+            .map(|_| xavier(&mut rng, META_Q_H2, 1)).collect();
+        Self {
+            w1, b1: vec![0.0; META_Q_H1],
+            w2, b2: vec![0.0; META_Q_H2],
+            w3, b3: vec![0.0; 1],
+            m_w1: vec![0.0; META_Q_INPUT_DIM * META_Q_H1], v_w1: vec![0.0; META_Q_INPUT_DIM * META_Q_H1],
+            m_w2: vec![0.0; META_Q_H1 * META_Q_H2], v_w2: vec![0.0; META_Q_H1 * META_Q_H2],
+            m_w3: vec![0.0; META_Q_H2], v_w3: vec![0.0; META_Q_H2],
+            lr: 1e-3,
+            step: 0,
+            buffer: VecDeque::with_capacity(200),
+            buffer_cap: 200,
+        }
+    }
+
+    pub fn predict(&self, input: &MetaQInput) -> f32 {
+        let x = input.to_vec();
+        let mut h1 = vec![0.0f32; META_Q_H1];
+        for j in 0..META_Q_H1 {
+            let mut s = self.b1[j];
+            for i in 0..META_Q_INPUT_DIM { s += x[i] * self.w1[i * META_Q_H1 + j]; }
+            h1[j] = s.max(0.0);
+        }
+        let mut h2 = vec![0.0f32; META_Q_H2];
+        for j in 0..META_Q_H2 {
+            let mut s = self.b2[j];
+            for i in 0..META_Q_H1 { s += h1[i] * self.w2[i * META_Q_H2 + j]; }
+            h2[j] = s.max(0.0);
+        }
+        let mut logit = self.b3[0];
+        for i in 0..META_Q_H2 { logit += h2[i] * self.w3[i]; }
+        1.0 / (1.0 + (-logit).exp())
+    }
+
+    pub fn push_sample(&mut self, input: MetaQInput, label: f32) {
+        if self.buffer.len() >= self.buffer_cap { self.buffer.pop_front(); }
+        self.buffer.push_back((input, label));
+    }
+
+    pub fn train_step(&mut self) {
+        if self.buffer.len() < 32 { return; }
+        let mut rng = SimpleLcg::new(self.step as u64 ^ 0xAA55AA55);
+        for _ in 0..32 {
+            let idx = (rng.next_u32() as usize) % self.buffer.len();
+            let (input, label) = self.buffer[idx].clone();
+            self.train_one(&input, label);
+        }
+        self.step += 1;
+    }
+
+    fn train_one(&mut self, input: &MetaQInput, label: f32) {
+        let x = input.to_vec();
+        let mut h1_pre = vec![0.0f32; META_Q_H1];
+        let mut h1 = vec![0.0f32; META_Q_H1];
+        for j in 0..META_Q_H1 {
+            let mut s = self.b1[j];
+            for i in 0..META_Q_INPUT_DIM { s += x[i] * self.w1[i * META_Q_H1 + j]; }
+            h1_pre[j] = s; h1[j] = s.max(0.0);
+        }
+        let mut h2_pre = vec![0.0f32; META_Q_H2];
+        let mut h2 = vec![0.0f32; META_Q_H2];
+        for j in 0..META_Q_H2 {
+            let mut s = self.b2[j];
+            for i in 0..META_Q_H1 { s += h1[i] * self.w2[i * META_Q_H2 + j]; }
+            h2_pre[j] = s; h2[j] = s.max(0.0);
+        }
+        let mut logit = self.b3[0];
+        for i in 0..META_Q_H2 { logit += h2[i] * self.w3[i]; }
+        let pred = 1.0 / (1.0 + (-logit).exp());
+
+        let d_logit = pred - label;
+        let mut d_w3 = vec![0.0f32; META_Q_H2];
+        for i in 0..META_Q_H2 { d_w3[i] = d_logit * h2[i]; }
+        let d_b3 = d_logit;
+        let d_h2: Vec<f32> = (0..META_Q_H2).map(|i| d_logit * self.w3[i] * (if h2_pre[i] > 0.0 { 1.0 } else { 0.0 })).collect();
+        let mut d_w2 = vec![0.0f32; META_Q_H1 * META_Q_H2];
+        let d_b2 = d_h2.clone();
+        for i in 0..META_Q_H1 {
+            for j in 0..META_Q_H2 { d_w2[i * META_Q_H2 + j] = d_h2[j] * h1[i]; }
+        }
+        let d_h1: Vec<f32> = (0..META_Q_H1).map(|i| {
+            let mut s = 0.0f32;
+            for j in 0..META_Q_H2 { s += d_h2[j] * self.w2[i * META_Q_H2 + j]; }
+            s * (if h1_pre[i] > 0.0 { 1.0 } else { 0.0 })
+        }).collect();
+        let mut d_w1 = vec![0.0f32; META_Q_INPUT_DIM * META_Q_H1];
+        let d_b1 = d_h1.clone();
+        for i in 0..META_Q_INPUT_DIM {
+            for j in 0..META_Q_H1 { d_w1[i * META_Q_H1 + j] = d_h1[j] * x[i]; }
+        }
+        self.adam_update(d_w1, d_w2, d_w3, d_b1, d_b2, d_b3);
+    }
+
+    fn adam_update(&mut self, d_w1: Vec<f32>, d_w2: Vec<f32>, d_w3: Vec<f32>, d_b1: Vec<f32>, d_b2: Vec<f32>, d_b3: f32) {
+        let beta1 = 0.9f32; let beta2 = 0.999f32; let eps = 1e-8f32;
+        let t = (self.step + 1) as f32;
+        let bc1 = 1.0 - beta1.powf(t); let bc2 = 1.0 - beta2.powf(t);
+        let lr = self.lr;
+
+        let adam = |p: &mut [f32], m: &mut [f32], v: &mut [f32], g: &[f32]| {
+            for i in 0..p.len() {
+                m[i] = beta1 * m[i] + (1.0 - beta1) * g[i];
+                v[i] = beta2 * v[i] + (1.0 - beta2) * g[i] * g[i];
+                let mh = m[i] / bc1;
+                let vh = v[i] / bc2;
+                p[i] -= lr * mh / (vh.sqrt() + eps);
+            }
+        };
+        adam(&mut self.w1, &mut self.m_w1, &mut self.v_w1, &d_w1);
+        adam(&mut self.w2, &mut self.m_w2, &mut self.v_w2, &d_w2);
+        adam(&mut self.w3, &mut self.m_w3, &mut self.v_w3, &d_w3);
+        for i in 0..self.b1.len() { self.b1[i] -= lr * d_b1[i]; }
+        for i in 0..self.b2.len() { self.b2[i] -= lr * d_b2[i]; }
+        self.b3[0] -= lr * d_b3;
+    }
+}
+
+impl Default for MetaQNetwork {
+    fn default() -> Self { Self::new() }
+}
+
+/// Deterministic LCG for reproducible weight init and minibatch sampling.
+#[derive(Debug)]
+struct SimpleLcg { state: u64 }
+
+impl SimpleLcg {
+    fn new(seed: u64) -> Self { Self { state: seed.wrapping_add(1) } }
+    fn next_u32(&mut self) -> u32 {
+        self.state = self.state.wrapping_mul(6364136223846793005).wrapping_add(1442695040888963407);
+        (self.state >> 33) as u32
+    }
+    fn next_f32(&mut self) -> f32 {
+        (self.next_u32() as f32) / (u32::MAX as f32)
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn test_predict_returns_probability() {
+        let net = MetaQNetwork::new();
+        let input = MetaQInput {
+            q_gap_ema: 0.3, log_grad_norm: 2.0, atom_util: 0.5,
+            ens_disagreement: 0.1, regime_stability: 0.7,
+            spectral_gap_norm: 0.6, grad_consistency: 0.3,
+        };
+        let p = net.predict(&input);
+        assert!((0.0..=1.0).contains(&p), "probability out of range: {}", p);
+    }
+
+    #[test]
+    fn test_training_reduces_loss() {
+        let mut net = MetaQNetwork::new();
+        let collapse_input = MetaQInput {
+            q_gap_ema: 0.01, log_grad_norm: 10.0, atom_util: 0.1,
+            ens_disagreement: 0.0, regime_stability: 0.5,
+            spectral_gap_norm: 0.1, grad_consistency: -0.3,
+        };
+        for _ in 0..100 { net.push_sample(collapse_input.clone(), 1.0); }
+        for _ in 0..50 { net.train_step(); }
+        let p = net.predict(&collapse_input);
+        assert!(p > 0.5, "Trained net should predict >0.5 for collapse state, got {}", p);
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs
index c714a72f1..981ef4b26 100644
--- a/crates/ml/src/cuda_pipeline/mod.rs
+++ b/crates/ml/src/cuda_pipeline/mod.rs
@@ -37,6 +37,9 @@ pub mod gpu_iql_trainer;
 pub mod gpu_iqn_head;
 pub mod gpu_attention;
 pub mod decision_transformer;
+pub mod learning_health;
+pub mod q_snapshot;
+pub mod meta_q_network;
 // gpu_replay_buffer moved to ml-dqn crate
 
 /// Maximum bytes allowed for a single GPU upload (2 GB safety limit).
diff --git a/crates/ml/src/cuda_pipeline/q_snapshot.rs b/crates/ml/src/cuda_pipeline/q_snapshot.rs
new file mode 100644
index 000000000..a63d88441
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/q_snapshot.rs
@@ -0,0 +1,123 @@
+//! N1: Q-network weight snapshots ring buffer for temporal self-distillation.
+//!
+//! Keeps up to `MAX_SNAPSHOTS` historical weight checkpoints taken at high-health
+//! epochs. During collapse recovery, a gradient pull toward the best snapshot
+//! is added to the master gradient.
+
+use std::sync::Arc;
+use cudarc::driver::{CudaSlice, CudaStream, DevicePtr, DevicePtrMut};
+use crate::MLError;
+
+pub const MAX_SNAPSHOTS: usize = 5;
+pub const SNAPSHOT_HEALTH_THRESHOLD: f32 = 0.7;
+pub const DISTILL_HEALTH_THRESHOLD: f32 = 0.4;
+
+/// A single weight snapshot with metadata.
+pub struct QSnapshot {
+    pub weights: CudaSlice<f32>,
+    pub health: f32,
+    pub q_gap: f32,
+    pub epoch: u32,
+}
+
+impl std::fmt::Debug for QSnapshot {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        f.debug_struct("QSnapshot")
+            .field("health", &self.health)
+            .field("q_gap", &self.q_gap)
+            .field("epoch", &self.epoch)
+            .finish()
+    }
+}
+
+/// Ring buffer of best-health snapshots.
+pub struct SnapshotRing {
+    stream: Arc<CudaStream>,
+    pub(crate) snapshots: Vec<QSnapshot>,
+    param_count: usize,
+}
+
+impl std::fmt::Debug for SnapshotRing {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        f.debug_struct("SnapshotRing")
+            .field("param_count", &self.param_count)
+            .field("snapshot_count", &self.snapshots.len())
+            .finish()
+    }
+}
+
+impl SnapshotRing {
+    pub fn new(stream: Arc<CudaStream>, param_count: usize) -> Self {
+        Self {
+            stream,
+            snapshots: Vec::with_capacity(MAX_SNAPSHOTS),
+            param_count,
+        }
+    }
+
+    pub fn is_empty(&self) -> bool {
+        self.snapshots.is_empty()
+    }
+
+    pub fn param_count(&self) -> usize {
+        self.param_count
+    }
+
+    /// Copy current weights into a new snapshot if health >= threshold and
+    /// q_gap beats the worst stored snapshot (when at capacity).
+    pub fn maybe_snapshot(
+        &mut self,
+        params_src: &CudaSlice<f32>,
+        health: f32,
+        q_gap: f32,
+        epoch: u32,
+    ) -> Result<bool, MLError> {
+        if health < SNAPSHOT_HEALTH_THRESHOLD {
+            return Ok(false);
+        }
+
+        // At capacity, only replace if this q_gap beats the worst stored entry.
+        if self.snapshots.len() >= MAX_SNAPSHOTS {
+            let worst_q_gap = self.snapshots.iter()
+                .map(|s| s.q_gap)
+                .fold(f32::INFINITY, f32::min);
+            if q_gap <= worst_q_gap {
+                return Ok(false);
+            }
+        }
+
+        let n = self.param_count;
+        let num_bytes = n * std::mem::size_of::<f32>();
+        let mut new_weights = self.stream.alloc_zeros::<f32>(n)
+            .map_err(|e| MLError::ModelError(format!("snapshot alloc: {e}")))?;
+        {
+            let (src_ptr, _src_guard) = params_src.device_ptr(&self.stream);
+            let (dst_ptr, _dst_guard) = new_weights.device_ptr_mut(&self.stream);
+            #[allow(unsafe_code)]
+            unsafe {
+                cudarc::driver::result::memcpy_dtod_async(
+                    dst_ptr, src_ptr, num_bytes, self.stream.cu_stream(),
+                ).map_err(|e| MLError::ModelError(format!("snapshot dtod: {e}")))?;
+            }
+        }
+
+        if self.snapshots.len() >= MAX_SNAPSHOTS {
+            let worst_idx = self.snapshots.iter()
+                .enumerate()
+                .min_by(|(_, a), (_, b)| a.q_gap.partial_cmp(&b.q_gap).unwrap())
+                .map(|(i, _)| i)
+                .unwrap();
+            self.snapshots.swap_remove(worst_idx);
+        }
+
+        self.snapshots.push(QSnapshot { weights: new_weights, health, q_gap, epoch });
+        Ok(true)
+    }
+
+    /// Best (highest q_gap) snapshot raw device pointer, if any.
+    pub fn best_raw_ptr(&self) -> Option<u64> {
+        self.snapshots.iter()
+            .max_by(|a, b| a.q_gap.partial_cmp(&b.q_gap).unwrap())
+            .map(|s| s.weights.raw_ptr())
+    }
+}
diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs
index 0dcd32f7a..b84d36816 100644
--- a/crates/ml/src/trainers/dqn/fused_training.rs
+++ b/crates/ml/src/trainers/dqn/fused_training.rs
@@ -1157,23 +1157,26 @@ impl FusedTrainingCtx {
             ).map_err(|e| anyhow::anyhow!("HER in-place relabel kernel: {e}"))?;
         }
 
-        // NOTE: Per-component gradient budget clip REMOVED.
-        // The old code clipped C51+MSE gradient to a fraction of max_grad_norm,
-        // producing a CONSTANT gradient norm every step. Adam degenerates into
-        // SignSGD when it never sees true gradient magnitude variation.
-        // Safety is now handled by the single global clip inside Adam's kernel.
+        // NOTE: Per-component gradient budget clip REMOVED (old explicit clip → constant grad norm).
+        // B4/G5: Budgets are now applied as SCALE factors on individual SAXPY contributions,
+        // not as norm clips. This preserves gradient magnitude variation while steering
+        // component balance. Safety is still handled by the single global clip in Adam.
+
+        // B4/G5: Compute adaptive budgets once per step from ISV health/regime signals.
+        let (_c51_budget, iqn_budget, cql_budget, _ens_budget) = self.compute_adaptive_budgets();
 
         // DIAGNOSTIC: sync between each aux op to find hanging kernel
 
         // EMA target update — adaptive tau based on Q-divergence.
         {
             let dqn = agent.primary_dqn_mut();
-            let tau = compute_cosine_annealed_tau(
+            let tau_scheduled = compute_cosine_annealed_tau(
                 dqn.get_training_steps(),
                 dqn.config.tau,
                 dqn.config.tau_final,
                 dqn.config.tau_anneal_steps,
             );
+            let tau = self.trainer.apply_health_coupled_tau_floor(tau_scheduled);
             // Use cosine-annealed tau directly — fully deterministic.
             // Adaptive tau (from q_divergence) was removed because q_divergence
             // uses atomicAdd in the C51 loss kernel, introducing non-determinism
@@ -1424,16 +1427,17 @@ impl FusedTrainingCtx {
             if let Some(ref mut iqn) = self.gpu_iqn {
                 let d_h_s2_ptr = iqn.d_h_s2_raw_ptr();
                 self.trainer.apply_iqn_trunk_gradient(
-                    d_h_s2_ptr, &mut self.online_dueling,
+                    d_h_s2_ptr, &mut self.online_dueling, iqn_budget,
                 ).map_err(|e| anyhow::anyhow!("IQN trunk gradient: {e}"))?;
 
                 let dqn = agent.primary_dqn_mut();
-                let tau = compute_cosine_annealed_tau(
+                let tau_scheduled = compute_cosine_annealed_tau(
                     dqn.get_training_steps(),
                     dqn.config.tau,
                     dqn.config.tau_final,
                     dqn.config.tau_anneal_steps,
                 );
+                let tau = self.trainer.apply_health_coupled_tau_floor(tau_scheduled);
                 iqn.target_ema_update(tau as f32)
                     .map_err(|e| anyhow::anyhow!("IQN EMA update: {e}"))?;
 
@@ -1455,10 +1459,11 @@ impl FusedTrainingCtx {
 
 
         // CQL gradient + SAXPY (no per-component clip — Adam handles safety).
+        // B4/G5: cql_budget scales the SAXPY contribution (0.10×(1−regime)×health).
         if self.trainer.has_cql() {
             match self.trainer.apply_cql_gradient() {
                 Ok(true) => {
-                    self.trainer.apply_cql_saxpy()
+                    self.trainer.apply_cql_saxpy(cql_budget)
                         .map_err(|e| anyhow::anyhow!("CQL SAXPY: {e}"))?;
                 }
                 Ok(false) => {}
@@ -1829,23 +1834,25 @@ impl FusedTrainingCtx {
             iqn.increment_adam_step();
             // Update IQN tau via stable host address
             let dqn = agent.primary_dqn_mut();
-            let tau = compute_cosine_annealed_tau(
+            let tau_scheduled = compute_cosine_annealed_tau(
                 dqn.get_training_steps(),
                 dqn.config.tau,
                 dqn.config.tau_final,
                 dqn.config.tau_anneal_steps,
             );
+            let tau = self.trainer.apply_health_coupled_tau_floor(tau_scheduled);
             iqn.set_tau_host(tau as f32);
         }
         // Update DQN trainer tau via stable host address
         {
             let dqn = agent.primary_dqn_mut();
-            let tau = compute_cosine_annealed_tau(
+            let tau_scheduled = compute_cosine_annealed_tau(
                 dqn.get_training_steps(),
                 dqn.config.tau,
                 dqn.config.tau_final,
                 dqn.config.tau_anneal_steps,
             );
+            let tau = self.trainer.apply_health_coupled_tau_floor(tau_scheduled);
             self.trainer.set_tau_value(tau as f32);
         }
     }
@@ -2172,6 +2179,98 @@ impl FusedTrainingCtx {
         self.trainer.read_isv_regime()
     }
 
+    /// D8/N8: Forward (health, regime_stability) read through to caller.
+    pub(crate) fn read_isv_regime_via_trainer(&self) -> (f32, f32) {
+        self.trainer.read_isv_health_and_regime()
+    }
+
+    // ── A3: LearningHealth signal writers / readers ───────────────────────
+
+    /// Write a single f32 value to ISV signals buffer at given index (pinned HtoD).
+    pub(crate) fn write_isv_signal_at(&self, index: usize, value: f32) {
+        self.trainer.write_isv_signal_at(index, value);
+    }
+
+    /// Read atom utilization from the last Q-stats readback (pinned pointer offset 6).
+    pub(crate) fn read_atom_utilization(&self) -> f32 {
+        self.trainer.read_atom_utilization()
+    }
+
+    /// Last effective cql_alpha from the most recent apply_cql_gradient() call.
+    pub(crate) fn last_cql_alpha_eff(&self) -> f32 {
+        self.trainer.last_cql_alpha_eff
+    }
+
+    /// Last effective tau (cosine-annealed, potentially health-floored).
+    pub(crate) fn last_tau_eff(&self) -> f32 {
+        self.trainer.last_tau_eff
+    }
+
+    /// B3/G4: Last health-scaled Expected SARSA tau factor (1.0 at full health, 6.0 at collapse).
+    pub(crate) fn last_sarsa_tau_factor(&self) -> f32 {
+        self.trainer.last_sarsa_tau_factor
+    }
+
+    /// B4/G5: Compute adaptive gradient budgets from learning_health and regime_stability.
+    /// 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
+    ///
+    /// 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).
+    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);
+
+        self.trainer.last_iqn_budget_eff = iqn_budget;
+        self.trainer.last_cql_budget_eff = cql_budget;
+        self.trainer.last_c51_budget_eff = c51_budget;
+        self.trainer.last_ens_budget_eff = ens_budget;
+
+        (c51_budget, iqn_budget, cql_budget, ens_budget)
+    }
+
+    /// B4/G5: Last adaptive IQN gradient budget (0.10 + 0.30×health).
+    pub(crate) fn last_iqn_budget_eff(&self) -> f32 { self.trainer.last_iqn_budget_eff }
+    /// B4/G5: Last adaptive CQL gradient budget (0.10×(1−regime)×health).
+    pub(crate) fn last_cql_budget_eff(&self) -> f32 { self.trainer.last_cql_budget_eff }
+    /// B4/G5: Last adaptive C51 gradient budget (1−iqn−cql−ens).
+    pub(crate) fn last_c51_budget_eff(&self) -> f32 { self.trainer.last_c51_budget_eff }
+    /// B4/G5: Last ensemble gradient budget (constant 0.05).
+    pub(crate) fn last_ens_budget_eff(&self) -> f32 { self.trainer.last_ens_budget_eff }
+
+    /// Compute spectral gap on the last Q readback slice (coarse max/min ratio proxy).
+    pub(crate) fn compute_q_spectral_gap(&self) -> f32 {
+        self.trainer.compute_q_spectral_gap()
+    }
+
+    /// D1/N1: Conditionally snapshot current weights at epoch boundary.
+    /// Snapshots if health >= 0.7 and q_gap beats the worst stored entry.
+    pub(crate) fn maybe_snapshot_qnet(&mut self, health: f32, q_gap: f32, epoch: u32) -> anyhow::Result<bool> {
+        self.trainer.maybe_snapshot_params(health, q_gap, epoch)
+            .map_err(|e| anyhow::anyhow!("snapshot: {e}"))
+    }
+
+    /// D1/N1: Apply distillation gradient pull toward best snapshot.
+    /// No-op when health >= 0.4 or no snapshots exist yet.
+    pub(crate) fn apply_distillation(&mut self) -> anyhow::Result<bool> {
+        self.trainer.apply_distillation_gradient()
+            .map_err(|e| anyhow::anyhow!("distill: {e}"))
+    }
+
+    /// D1/N1: Whether distillation was active in the most recent epoch boundary.
+    pub(crate) fn last_distill_active(&self) -> bool {
+        self.trainer.last_distill_active
+    }
+
     pub(crate) fn read_popart_variance(&self) -> f32 {
         if !self.popart_enabled { return 0.0; }
         let mut var = [0.0_f32];
@@ -2666,6 +2765,17 @@ impl FusedTrainingCtx {
         self.trainer.set_adaptive_gamma(gamma);
     }
 
+    /// C4/P4: Apply temporal-coupled adjustment to the scheduled gamma.
+    /// Delegates to `GpuDqnTrainer::apply_adaptive_gamma`, which reads health
+    /// and regime_stability from ISV pinned memory, clamps to [0.9, 0.995],
+    /// and caches the result in `last_gamma_eff`.
+    pub(crate) fn apply_adaptive_gamma(&mut self, gamma_base: f32) -> f32 {
+        self.trainer.apply_adaptive_gamma(gamma_base)
+    }
+
+    /// C4/P4: Last effective gamma after temporal-coupled adjustment.
+    pub(crate) fn last_gamma_eff(&self) -> f32 { self.trainer.last_gamma_eff }
+
     /// G5: Set the EMA threshold for epistemic variance gating.
     pub(crate) fn set_var_ema(&self, val: f32) {
         self.trainer.set_var_ema(val);
diff --git a/crates/ml/src/trainers/dqn/smoke_tests/adaptive_learning.rs b/crates/ml/src/trainers/dqn/smoke_tests/adaptive_learning.rs
new file mode 100644
index 000000000..ed0261407
--- /dev/null
+++ b/crates/ml/src/trainers/dqn/smoke_tests/adaptive_learning.rs
@@ -0,0 +1,96 @@
+//! E1: Collapse-recovery smoke test for the LearningHealth adaptive system.
+//!
+//! Trains a short DQN run on fxcache data and asserts that the adaptive
+//! mechanisms (health-coupled CQL/tau/SARSA-temp/gradient-budget/gamma,
+//! PER diversity, snapshots/distillation, Q-gap barrier, plasticity injection,
+//! CF curriculum, IB, ensemble oracle, contrarian override, meta-Q) collectively
+//! prevent Q-value collapse. The specific thresholds encode the design goals:
+//!
+//!   - `q_gap_ema` must stay above 0.05 (non-trivial action differentiation)
+//!   - `learning_health.value` must stay above 0.3 (system not pinned in collapse mode)
+//!
+//! Baseline runs prior to this feature collapsed both signals to near-zero by
+//! epoch 8 — this test is a regression guard.
+//!
+//! Run: `FOXHUNT_TEST_DATA=test_data/futures-baseline \
+//!   cargo test -p ml --lib -- test_adaptive_learning_no_collapse --ignored --nocapture`
+
+use super::helpers::*;
+
+#[test]
+#[ignore] // GPU + fxcache data
+fn test_adaptive_learning_no_collapse() -> anyhow::Result<()> {
+    use crate::fxcache;
+
+    let cache_dir = feature_cache_dir();
+    assert!(cache_dir.exists(), "feature-cache not found at {:?}", cache_dir);
+
+    let entries: Vec<_> = std::fs::read_dir(&cache_dir)?
+        .filter_map(|e| e.ok())
+        .filter(|e| e.path().extension().and_then(|s| s.to_str()) == Some("fxcache"))
+        .collect();
+    assert!(!entries.is_empty(), "No .fxcache files in {:?}", cache_dir);
+
+    let fxcache_data = fxcache::load_fxcache(&entries[0].path())?;
+    let n = fxcache_data.bar_count.min(20_000);
+    let train_end = (n * 80) / 100;
+
+    let mut params = smoke_params();
+    params.epochs = 10;
+
+    let mut trainer = smoke_trainer_with(params)?;
+    let rt = tokio::runtime::Builder::new_current_thread()
+        .enable_all()
+        .build()?;
+
+    let features = &fxcache_data.features[..n];
+    let targets = &fxcache_data.targets[..n];
+    let ofi = &fxcache_data.ofi[..n.min(fxcache_data.ofi.len())];
+
+    let train_feat = &features[..train_end];
+    let train_targets = &targets[..train_end];
+    let val_feat = &features[train_end..n];
+    let val_targets = &targets[train_end..n];
+
+    rt.block_on(trainer.init_from_fxcache(features, targets, ofi))?;
+    trainer.set_training_range(0, train_end, train_end, n);
+    trainer.set_val_data_from_slices(val_feat, val_targets, train_end);
+    rt.block_on(trainer.reset_for_fold())?;
+
+    rt.block_on(trainer.train_fold_from_slices(
+        train_feat,
+        train_targets,
+        |_epoch, _bytes, _best| Ok("skip".to_owned()),
+    ))?;
+
+    // ── Post-training assertions ──────────────────────────────────────────
+    let final_q_gap = trainer.health_ema.q_gap_ema;
+    let final_health = trainer.learning_health.value;
+    let unhealthy_streak = trainer.unhealthy_epoch_count;
+
+    eprintln!(
+        "[ADAPTIVE] final q_gap_ema={:.4}, health={:.2}, unhealthy_streak={}, \
+         low_winrate={}, contrarian_active={}",
+        final_q_gap,
+        final_health,
+        unhealthy_streak,
+        trainer.low_winrate_count,
+        trainer.contrarian_active,
+    );
+
+    assert!(
+        final_q_gap > 0.05,
+        "Q-gap collapsed to {:.4} after 10 epochs — adaptive mechanisms failed \
+         to prevent collapse (threshold 0.05)",
+        final_q_gap
+    );
+
+    assert!(
+        final_health > 0.3,
+        "LearningHealth dropped to {:.2} — system is pinned in collapse mode \
+         (threshold 0.30)",
+        final_health
+    );
+
+    Ok(())
+}
diff --git a/crates/ml/src/trainers/dqn/smoke_tests/mod.rs b/crates/ml/src/trainers/dqn/smoke_tests/mod.rs
index db2b1b0ac..12a29580d 100644
--- a/crates/ml/src/trainers/dqn/smoke_tests/mod.rs
+++ b/crates/ml/src/trainers/dqn/smoke_tests/mod.rs
@@ -20,3 +20,5 @@ mod regression;
 pub mod reward_v8;
 #[cfg(test)]
 mod generalization;
+#[cfg(test)]
+mod adaptive_learning;
diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs
index cdaba041f..f37c91b6e 100644
--- a/crates/ml/src/trainers/dqn/trainer/constructor.rs
+++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs
@@ -723,6 +723,36 @@ impl DQNTrainer {
             enrichment: super::enrichment::EnrichmentState::new(
                 initial_epsilon,
             ),
+            learning_health: crate::cuda_pipeline::learning_health::LearningHealth::new(),
+            health_ema: Default::default(),
+            last_q_gap: None,
+            last_q_var: None,
+            last_grad_norm: None,
+            last_ens_disagreement: None,
+            last_atom_util: None,
+            last_cql_alpha_eff: None,
+            last_iqn_budget_eff: None,
+            last_cql_budget_eff: None,
+            last_c51_budget_eff: None,
+            last_tau_eff: None,
+            last_sarsa_tau_factor: None,
+            last_gamma_eff: None,
+            last_cf_ratio_eff: None,
+            last_distill_active: None,
+            last_barrier_loss: None,
+            last_plasticity_ready: None,
+            unhealthy_epoch_count: 0,
+            last_ib_penalty: None,
+            last_ensemble_collapse_score: None,
+            last_contrarian_active: None,
+            low_winrate_count: 0,
+            contrarian_active: false,
+            contrarian_remaining_epochs: 0,
+            last_epoch_win_rate: 0.5,
+            last_meta_q_pred: None,
+            meta_q: crate::cuda_pipeline::meta_q_network::MetaQNetwork::new(),
+            pending_meta_q_samples: std::collections::VecDeque::with_capacity(32),
+            health_history: std::collections::VecDeque::with_capacity(2 * crate::cuda_pipeline::meta_q_network::META_Q_LOOKAHEAD),
         })
     }
 
diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs
index e7442d17a..4f889ae72 100644
--- a/crates/ml/src/trainers/dqn/trainer/metrics.rs
+++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs
@@ -6,6 +6,43 @@ use super::DQNTrainer;
 use crate::TrainingMetrics;
 use super::super::config::DQNAgentType;
 
+/// Rolling EMA values for LearningHealth inputs.
+#[derive(Debug, Clone, Default)]
+pub(crate) struct HealthEmaTrackers {
+    pub q_gap_ema: f32,
+    pub q_var_ema: f32,
+    pub grad_norm_ema: f32,
+    prev_grad_norm: Option<f32>,
+    ema_alpha: f32, // initialized to 0.1 on first update
+}
+
+impl HealthEmaTrackers {
+    pub(crate) fn update(&mut self, q_gap: f32, q_var: f32, grad_norm: f32) {
+        if self.ema_alpha == 0.0 { self.ema_alpha = 0.1; }
+        self.q_gap_ema = (1.0 - self.ema_alpha) * self.q_gap_ema + self.ema_alpha * q_gap;
+        self.q_var_ema = (1.0 - self.ema_alpha) * self.q_var_ema + self.ema_alpha * q_var;
+        self.grad_norm_ema = (1.0 - self.ema_alpha) * self.grad_norm_ema + self.ema_alpha * grad_norm;
+    }
+
+    /// Simple scalar proxy for gradient direction consistency: 1 - |delta|/max.
+    /// Returns 0.0 on first call; +1.0 means no change, -1.0 means opposite sign.
+    /// This is a weaker proxy than full vector cosine similarity, but keeps A3
+    /// scope tight — can be upgraded in a follow-up if the Adam flat buffer
+    /// becomes available.
+    pub(crate) fn update_grad_consistency(&mut self, current_grad_norm: f32) -> f32 {
+        let cos = match self.prev_grad_norm {
+            None => 0.0,
+            Some(prev) => {
+                let denom = prev.abs().max(current_grad_norm.abs()).max(1e-6);
+                let delta = (current_grad_norm - prev).abs() / denom;
+                (1.0 - 2.0 * delta).clamp(-1.0, 1.0)
+            }
+        };
+        self.prev_grad_norm = Some(current_grad_norm);
+        cos
+    }
+}
+
 impl DQNTrainer {
     /// Full training loop with existing logic (Wave 12 Group 3)
     /// Calculate average metrics for an epoch
diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs
index 16c565a7b..1ababcb82 100644
--- a/crates/ml/src/trainers/dqn/trainer/mod.rs
+++ b/crates/ml/src/trainers/dqn/trainer/mod.rs
@@ -472,6 +472,72 @@ pub struct DQNTrainer {
 
     /// Self-improving enrichment state (persists across epochs).
     pub(crate) enrichment: enrichment::EnrichmentState,
+
+    // ── A3: LearningHealth (Layer 1) ─────────────────────────────────
+    /// Unified training health scalar [0.2, 0.95] — EMA-smoothed composition of
+    /// Q-gap, Q-variance, atom utilization, grad stability, ensemble agreement,
+    /// gradient consistency, and spectral gap. Broadcast to ISV[12] each epoch.
+    pub(crate) learning_health: crate::cuda_pipeline::learning_health::LearningHealth,
+    /// Rolling EMA trackers for LearningHealth input signals.
+    pub(crate) health_ema: metrics::HealthEmaTrackers,
+    /// Last computed Q-gap signal (mean of 4 per-branch EMA gaps). Cached for A4/B1.
+    pub(crate) last_q_gap: Option<f32>,
+    /// Last computed Q-variance proxy (epoch_q_max - epoch_q_min). Cached for A4.
+    pub(crate) last_q_var: Option<f32>,
+    /// Last grad_norm EMA fed into LearningHealth. Cached for A4/B4.
+    pub(crate) last_grad_norm: Option<f32>,
+    /// Last ensemble disagreement value (placeholder 0.1 until D6/N6). Cached for A4.
+    pub(crate) last_ens_disagreement: Option<f32>,
+    /// Last atom utilization from q_readback_pinned[6]. Cached for A4/G4.
+    pub(crate) last_atom_util: Option<f32>,
+
+    // ── A4: Extended HEALTH_DIAG placeholders ─────────────────────────────
+    // All fields start as None; B/C/D tasks populate them with real values.
+    /// B1/G2 — CQL alpha effective value
+    pub(crate) last_cql_alpha_eff: Option<f32>,
+    /// B4/G5 — per-component gradient budget allocations
+    pub(crate) last_iqn_budget_eff: Option<f32>,
+    pub(crate) last_cql_budget_eff: Option<f32>,
+    pub(crate) last_c51_budget_eff: Option<f32>,
+    /// B2/G3 — tau effective
+    pub(crate) last_tau_eff: Option<f32>,
+    /// B3/G4 — Expected SARSA temperature scaling factor
+    pub(crate) last_sarsa_tau_factor: Option<f32>,
+    /// C4/P4 — adaptive gamma
+    pub(crate) last_gamma_eff: Option<f32>,
+    /// D4/N4 — counterfactual ratio
+    pub(crate) last_cf_ratio_eff: Option<f32>,
+    /// D1/N1 — distillation active flag
+    pub(crate) last_distill_active: Option<bool>,
+    /// D2/N2 — barrier loss value
+    pub(crate) last_barrier_loss: Option<f32>,
+    /// D3/N3 — plasticity-ready flag (true = ready to trigger, false = in cooldown)
+    pub(crate) last_plasticity_ready: Option<bool>,
+    /// D3/N3: Consecutive epoch count with health < 0.3. When reaches 3, triggers plasticity injection.
+    pub(crate) unhealthy_epoch_count: u32,
+    /// D5/N5 — information bottleneck penalty
+    pub(crate) last_ib_penalty: Option<f32>,
+    /// D6/N6 — ensemble collapse score
+    pub(crate) last_ensemble_collapse_score: Option<f32>,
+    /// D7/N7 — contrarian override active
+    pub(crate) last_contrarian_active: Option<bool>,
+    /// D7/N7: Consecutive epochs with WinRate < 0.40. Triggers contrarian override at 5.
+    pub(crate) low_winrate_count: u32,
+    /// D7/N7: Whether contrarian override (argmin) is currently active.
+    pub(crate) contrarian_active: bool,
+    /// D7/N7: Epochs remaining in current contrarian window (2 epochs per activation).
+    pub(crate) contrarian_remaining_epochs: u32,
+    /// D7/N7: Most recent epoch win_rate (0..1) from financials. Carried into next process_epoch_boundary.
+    pub(crate) last_epoch_win_rate: f32,
+    /// D8/N8 — meta-Q collapse prediction
+    pub(crate) last_meta_q_pred: Option<f32>,
+    /// D8/N8: Meta-Q network predicting collapse probability K epochs ahead.
+    pub(crate) meta_q: crate::cuda_pipeline::meta_q_network::MetaQNetwork,
+    /// D8/N8: Pending inputs awaiting label resolution (K epochs delayed).
+    pub(crate) pending_meta_q_samples: std::collections::VecDeque<(crate::cuda_pipeline::meta_q_network::MetaQInput, u32)>,
+    /// D8/N8: Rolling history of (epoch, health) for meta-Q label resolution.
+    /// Capped at 2×META_Q_LOOKAHEAD.
+    pub(crate) health_history: std::collections::VecDeque<(u32, f32)>,
 }
 
 impl std::fmt::Debug for DQNTrainer {
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 95698c79e..67e306eb3 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -477,18 +477,8 @@ impl DQNTrainer {
                 }
             }
 
-            // Periodic shrink-and-perturb
-            let sp_interval = self.hyperparams.shrink_perturb_interval;
-            if sp_interval > 0 && epoch > 0 && epoch % sp_interval == 0 {
-                if let Some(ref mut fused) = self.fused_ctx {
-                    let alpha = self.hyperparams.shrink_perturb_alpha as f32;
-                    let sigma = self.hyperparams.shrink_perturb_sigma as f32;
-                    match fused.shrink_and_perturb(alpha, sigma) {
-                        Ok(()) => info!(epoch, alpha, sigma, "Periodic shrink-and-perturb applied"),
-                        Err(e) => warn!(epoch, "Shrink-and-perturb failed (non-fatal): {e}"),
-                    }
-                }
-            }
+            // D3/N3: Periodic shrink-and-perturb replaced by health-triggered path in
+            // process_epoch_boundary. No periodic trigger here.
 
             // ── Phase 3: Batched training from replay buffer ──
             let phase3_start = std::time::Instant::now();
@@ -630,9 +620,12 @@ impl DQNTrainer {
                 } else if util < 0.4 {
                     self.adaptive_gamma = (self.adaptive_gamma - 0.01).max(0.90);
                 }
-                // Wire adaptive gamma into GPU trainer for C51 Bellman projection
+                // C4/P4: Wire adaptive gamma into GPU trainer with temporal-coupled adjustment.
+                // apply_adaptive_gamma reads health+regime from ISV, adjusts, clamps [0.9, 0.995],
+                // and caches last_gamma_eff for HEALTH_DIAG.
                 if let Some(ref mut fused) = self.fused_ctx {
-                    fused.set_adaptive_gamma(self.adaptive_gamma);
+                    let gamma_scheduled = self.adaptive_gamma as f32;
+                    fused.apply_adaptive_gamma(gamma_scheduled);
                 }
                 info!("G4 gamma={:.3} util={:.2}", self.adaptive_gamma, util);
             }
@@ -723,8 +716,10 @@ impl DQNTrainer {
                     if let Some(eval_gamma) = result.gamma {
                         let blended = 0.9 * self.adaptive_gamma + 0.1 * eval_gamma;
                         self.adaptive_gamma = blended.clamp(0.85, 0.98);
+                        // C4/P4: apply temporal-coupled adjustment on top of EMA blend.
                         if let Some(ref mut fused) = self.fused_ctx {
-                            fused.set_adaptive_gamma(self.adaptive_gamma);
+                            let gamma_scheduled = self.adaptive_gamma as f32;
+                            fused.apply_adaptive_gamma(gamma_scheduled);
                         }
                     }
 
@@ -1783,6 +1778,338 @@ impl DQNTrainer {
                 })?;
         }
 
+        // ── LearningHealth computation (Layer 1) ──────────────────────────────
+        // Source signal adaptations vs. plan spec (A3 known deviations):
+        //   - q_gap: mean of per_branch_q_gap_ema (4 branches) via fused_ctx —
+        //     NOTE: per_branch_q_gap_ema() performs a synchronous stream flush + DtoH
+        //     (cuStreamSynchronize + cuMemcpyDtoH_v2). Acceptable at epoch boundary
+        //     because it runs once per epoch; DO NOT move into a per-step loop.
+        //   - q_var: epoch_q_max - epoch_q_min proxy instead of QValueStatsResult.q_variance
+        //     (q_variance is not separately tracked as a scalar field; range is a sufficient proxy).
+        //   - grad_consistency: scalar proxy (successive grad_norm delta) instead of full Adam
+        //     vector cosine similarity — per-component Adam buffers prevent flat-vector approach.
+        //   - spectral_gap: coarse max/min ratio on q_readback_pinned instead of full cuSOLVER SVD.
+        //   - ens_disagreement: 0.1 placeholder until D6/N6 wires real ensemble pairwise Q-gap.
+        {
+            use crate::cuda_pipeline::learning_health::HealthComponents;
+
+            // Per-branch Q-gap mean as q_gap signal (4 values, epoch-scoped DtoH).
+            let q_gap_signal: f32 = if let Some(ref fused) = self.fused_ctx {
+                let gaps = fused.per_branch_q_gap_ema();
+                ((gaps[0] + gaps[1] + gaps[2] + gaps[3]) * 0.25).max(0.0)
+            } else {
+                0.0
+            };
+
+            // q_var proxy: range (max - min) of per-step Q values this epoch. If no
+            // steps ran (epoch_q_min still INFINITY), fall back to 0.0 rather than
+            // propagating NaN/-inf into health normalization.
+            let q_var_signal = if self.epoch_q_min.is_finite() && self.epoch_q_max.is_finite() {
+                (self.epoch_q_max - self.epoch_q_min).max(0.0)
+            } else {
+                0.0
+            };
+
+            // Atom utilization from q_readback_pinned[6] (one-step lag).
+            let atom_util = self.fused_ctx.as_ref()
+                .map(|f| f.read_atom_utilization())
+                .unwrap_or(0.0);
+
+            // Spectral gap proxy from q_readback_pinned[7..7+total_actions].
+            let spectral_gap = self.fused_ctx.as_ref()
+                .map(|f| f.compute_q_spectral_gap())
+                .unwrap_or(1.0);
+
+            let grad_norm = avg_grad.abs();
+            let ens_disagreement = 0.1f32; // placeholder — D6/N6 task will replace with real ensemble pairwise Q-gap
+
+            // Update EMAs.
+            self.health_ema.update(q_gap_signal, q_var_signal, grad_norm);
+            let grad_consistency = self.health_ema.update_grad_consistency(grad_norm);
+
+            let raw = HealthComponents {
+                q_gap: self.health_ema.q_gap_ema,
+                q_var: self.health_ema.q_var_ema,
+                atom_util,
+                grad_norm: self.health_ema.grad_norm_ema,
+                ens_disagreement,
+                grad_consistency,
+                spectral_gap,
+            };
+            let health_value = self.learning_health.update(&raw);
+
+            // Cache for later tasks (A4 logging, B1, B4, G4 etc.)
+            self.last_q_gap = Some(raw.q_gap);
+            self.last_q_var = Some(raw.q_var);
+            self.last_grad_norm = Some(raw.grad_norm);
+            self.last_ens_disagreement = Some(raw.ens_disagreement);
+            self.last_atom_util = Some(raw.atom_util);
+
+            // Broadcast to GPU at ISV[LEARNING_HEALTH_INDEX=12].
+            if let Some(ref fused) = self.fused_ctx {
+                fused.write_isv_signal_at(
+                    crate::cuda_pipeline::gpu_dqn_trainer::LEARNING_HEALTH_INDEX,
+                    health_value,
+                );
+            }
+
+            // B1/G2: propagate last cql_alpha_eff for logging.
+            if let Some(ref fused) = self.fused_ctx {
+                self.last_cql_alpha_eff = Some(fused.last_cql_alpha_eff());
+            }
+
+            // B2/G3: propagate last tau_eff for logging.
+            if let Some(ref fused) = self.fused_ctx {
+                self.last_tau_eff = Some(fused.last_tau_eff());
+            }
+
+            // B3/G4: propagate SARSA tau factor for logging.
+            if let Some(ref fused) = self.fused_ctx {
+                self.last_sarsa_tau_factor = Some(fused.last_sarsa_tau_factor());
+            }
+
+            // B4/G5: propagate adaptive gradient budgets for logging.
+            if let Some(ref fused) = self.fused_ctx {
+                self.last_iqn_budget_eff = Some(fused.last_iqn_budget_eff());
+                self.last_cql_budget_eff = Some(fused.last_cql_budget_eff());
+                self.last_c51_budget_eff = Some(fused.last_c51_budget_eff());
+            }
+
+            // C4/P4: propagate adaptive gamma for logging.
+            if let Some(ref fused) = self.fused_ctx {
+                self.last_gamma_eff = Some(fused.last_gamma_eff());
+            }
+
+            // D1/N1: snapshot high-health checkpoints and pull toward best during collapse.
+            if let Some(ref mut fused) = self.fused_ctx {
+                let q_gap_for_snapshot = self.last_q_gap.unwrap_or(0.0);
+                let _ = fused.maybe_snapshot_qnet(health_value, q_gap_for_snapshot, epoch as u32);
+                if health_value < 0.4 {
+                    let _ = fused.apply_distillation();
+                }
+                self.last_distill_active = Some(fused.last_distill_active());
+            }
+
+            // D2/N2: Q-gap barrier constraint (scalar-only — informational loss).
+            // min_required = 0.05 × health; barrier_loss = 0.5 × max(0, min_required − q_gap)².
+            // Drives no gradient in this task; surfaces as HEALTH_DIAG `barrier=...` for visibility.
+            {
+                let q_gap_current = self.health_ema.q_gap_ema;
+                let min_required = 0.05_f32 * health_value;
+                let barrier = (min_required - q_gap_current).max(0.0);
+                self.last_barrier_loss = Some(0.5 * barrier * barrier);
+            }
+
+            // D3/N3: Health-triggered plasticity injection.
+            // Track consecutive unhealthy epochs; after 3, run shrink_perturb to break
+            // frozen representations and reset the counter.
+            {
+                let is_unhealthy = health_value < 0.3;
+                if is_unhealthy {
+                    self.unhealthy_epoch_count += 1;
+                } else {
+                    self.unhealthy_epoch_count = 0;
+                }
+
+                if self.unhealthy_epoch_count >= 3 {
+                    tracing::info!(
+                        "Plasticity injection triggered: health={:.2} for {} consecutive epochs",
+                        health_value, self.unhealthy_epoch_count
+                    );
+                    if let Some(ref mut fused) = self.fused_ctx {
+                        let alpha = self.hyperparams.shrink_perturb_alpha as f32;
+                        let sigma = self.hyperparams.shrink_perturb_sigma as f32;
+                        match fused.shrink_and_perturb(alpha, sigma) {
+                            Ok(()) => tracing::info!(epoch, alpha, sigma, "D3/N3 health-triggered shrink_and_perturb applied"),
+                            Err(e) => tracing::warn!(epoch, "D3/N3 shrink_and_perturb failed (non-fatal): {e}"),
+                        }
+                    }
+                    self.unhealthy_epoch_count = 0; // reset after trigger
+                    self.last_plasticity_ready = Some(false); // cooldown state
+                } else {
+                    self.last_plasticity_ready = Some(true);
+                }
+            }
+
+            // D5/N5: Information Bottleneck penalty (scalar-only — informational).
+            // IB penalizes lack of state-dependence in Q (low variance across batch).
+            // Minimum viable: host-side scalar computed from cached q_var_ema, surfaced
+            // via HEALTH_DIAG as `ib={..}`. Full backward-through-variance is deferred.
+            {
+                let ib_weight = 0.05_f32 * (1.0 - health_value).max(0.0);
+                let min_q_var = 0.01_f32;
+                let q_var_current = self.health_ema.q_var_ema;
+                let ib_penalty_raw = (min_q_var - q_var_current).max(0.0);
+                self.last_ib_penalty = Some(ib_weight * ib_penalty_raw);
+            }
+
+            // D6/N6: Ensemble as collapse oracle.
+            // High ensemble_collapse_score ⇒ ensembles agree (often on uniform during collapse).
+            // When score > 0.8 AND we haven't already triggered N3 recently, force a plasticity
+            // injection — this shortcut catches collapse even when the moving-average of
+            // individual health inputs hasn't accumulated to 3 unhealthy epochs yet.
+            // Uses last_ens_disagreement (Path C scalar proxy — populated by A3 readback).
+            {
+                use crate::cuda_pipeline::learning_health::smoothstep;
+                let mean_pairwise_gap = self.last_ens_disagreement.unwrap_or(0.1);
+                let ensemble_collapse_score = 1.0 - smoothstep(0.01, 0.1, mean_pairwise_gap);
+                self.last_ensemble_collapse_score = Some(ensemble_collapse_score);
+
+                if ensemble_collapse_score > 0.8 && self.unhealthy_epoch_count < 3 {
+                    tracing::info!(
+                        "Ensemble collapse detected (score={:.2}): forcing plasticity injection",
+                        ensemble_collapse_score
+                    );
+                    if let Some(ref mut fused) = self.fused_ctx {
+                        let alpha = self.hyperparams.shrink_perturb_alpha as f32;
+                        let sigma = self.hyperparams.shrink_perturb_sigma as f32;
+                        match fused.shrink_and_perturb(alpha, sigma) {
+                            Ok(()) => {
+                                self.unhealthy_epoch_count = 0;
+                                self.last_plasticity_ready = Some(false);
+                            }
+                            Err(e) => tracing::warn!("D6/N6 shrink_and_perturb failed: {e}"),
+                        }
+                    }
+                }
+            }
+
+            // D7/N7: Contrarian override state machine.
+            // When WinRate < 0.40 for 5 consecutive epochs AND health < 0.3, flip to argmin
+            // for 2 epochs to escape the mirror-image attractor. WinRate >= 0.45 resets.
+            // Note: last_epoch_win_rate is carried from previous epoch's financials computation
+            // (which runs after process_epoch_boundary). Epoch 0 defaults to 0.5 (no trigger).
+            {
+                let winrate = self.last_epoch_win_rate;
+
+                if winrate < 0.40 {
+                    self.low_winrate_count += 1;
+                } else if winrate >= 0.45 {
+                    self.contrarian_active = false;
+                    self.low_winrate_count = 0;
+                }
+
+                if self.low_winrate_count >= 5 && health_value < 0.3 && !self.contrarian_active {
+                    tracing::warn!(
+                        "CONTRARIAN OVERRIDE activated: WinRate={:.1}% for {} epochs, health={:.2}",
+                        winrate * 100.0, self.low_winrate_count, health_value
+                    );
+                    self.contrarian_active = true;
+                    self.contrarian_remaining_epochs = 2;
+                }
+
+                if self.contrarian_active {
+                    if self.contrarian_remaining_epochs == 0 {
+                        tracing::info!(
+                            "CONTRARIAN OVERRIDE deactivated: window expired"
+                        );
+                        self.contrarian_active = false;
+                    } else {
+                        self.contrarian_remaining_epochs -= 1;
+                    }
+                }
+
+                self.last_contrarian_active = Some(self.contrarian_active);
+
+                // D7/N7 Part B: forward the flag to the experience collector so the
+                // kernel negates Q values during Boltzmann sampling.
+                if let Some(ref mut collector) = self.gpu_experience_collector {
+                    collector.set_contrarian_active(self.contrarian_active);
+                }
+            }
+
+            // D8/N8: Meta-Q collapse predictor — leading indicator, NOT fed back into LearningHealth.
+            {
+                use crate::cuda_pipeline::meta_q_network::{MetaQInput, META_Q_LOOKAHEAD};
+
+                // Build input from current health components. regime_stability read via fused_ctx.
+                let (_, regime_stability) = if let Some(ref fused) = self.fused_ctx {
+                    fused.read_isv_regime_via_trainer()
+                } else {
+                    (0.5_f32, 0.5_f32)
+                };
+
+                let meta_q_input = MetaQInput {
+                    q_gap_ema: self.health_ema.q_gap_ema,
+                    log_grad_norm: self.health_ema.grad_norm_ema.max(1e-3).ln(),
+                    atom_util: self.last_atom_util.unwrap_or(0.0),
+                    ens_disagreement: self.last_ens_disagreement.unwrap_or(0.1),
+                    regime_stability,
+                    spectral_gap_norm: self.learning_health.components.spectral_gap_norm,
+                    grad_consistency: self.learning_health.components.grad_consistency_norm,
+                };
+
+                // Predict collapse probability
+                let meta_q_pred = self.meta_q.predict(&meta_q_input);
+                self.last_meta_q_pred = Some(meta_q_pred);
+
+                // Push current (input, epoch) into pending — label assigned K epochs later
+                let cur_epoch_u32 = self.current_epoch as u32;
+                self.pending_meta_q_samples.push_back((meta_q_input, cur_epoch_u32));
+
+                // Track health history for label resolution
+                if self.health_history.len() >= 2 * META_Q_LOOKAHEAD {
+                    self.health_history.pop_front();
+                }
+                self.health_history.push_back((cur_epoch_u32, health_value));
+
+                // Resolve labels for samples that are K epochs old
+                while let Some((pending_input, pending_epoch)) = self.pending_meta_q_samples.front().cloned() {
+                    if cur_epoch_u32 < pending_epoch + META_Q_LOOKAHEAD as u32 { break; }
+                    let collapsed = self.health_history.iter()
+                        .filter(|(ep, _)| *ep > pending_epoch && *ep <= pending_epoch + META_Q_LOOKAHEAD as u32)
+                        .any(|(_, h)| *h < 0.3);
+                    let label = if collapsed { 1.0 } else { 0.0 };
+                    self.meta_q.push_sample(pending_input, label);
+                    self.pending_meta_q_samples.pop_front();
+                }
+
+                // Train one step
+                self.meta_q.train_step();
+            }
+
+            // HEALTH_DIAG: components are [0, 1] normalized. effective = hyperparams after health-adaptation. novels = mechanism states.
+            tracing::info!(
+                "HEALTH_DIAG[{}]: health={:.2} components [q_gap={:.2} q_var={:.2} atoms={:.2} grad_stable={:.2} ens_agree={:.2} grad_cos={:.2} spectral={:.2}] effective [cql_alpha={:.4} iqn_budget={:.2} cql_budget={:.2} c51_budget={:.2} tau={:.5} sarsa_tau={:.2} gamma={:.3} cf_ratio={:.2}] novels [distill={} barrier={:.3} plasticity={} ib={:.3} ensemble_collapse={:.2} contrarian={} meta_q_pred={:.2}]",
+                epoch,
+                health_value,
+                self.learning_health.components.q_gap_norm,
+                self.learning_health.components.q_var_norm,
+                self.learning_health.components.atom_util_norm,
+                self.learning_health.components.grad_stable,
+                self.learning_health.components.ens_agree,
+                self.learning_health.components.grad_consistency_norm,
+                self.learning_health.components.spectral_gap_norm,
+                self.last_cql_alpha_eff.unwrap_or(0.0),
+                self.last_iqn_budget_eff.unwrap_or(0.40),
+                self.last_cql_budget_eff.unwrap_or(0.10),
+                self.last_c51_budget_eff.unwrap_or(0.45),
+                self.last_tau_eff.unwrap_or(0.005),
+                self.last_sarsa_tau_factor.unwrap_or(1.0),
+                self.last_gamma_eff.unwrap_or(0.99),
+                self.last_cf_ratio_eff.unwrap_or(0.5),
+                if self.last_distill_active.unwrap_or(false) { "on" } else { "off" },
+                self.last_barrier_loss.unwrap_or(0.0),
+                if self.last_plasticity_ready.unwrap_or(true) { "ready" } else { "cooldown" },
+                self.last_ib_penalty.unwrap_or(0.0),
+                self.last_ensemble_collapse_score.unwrap_or(0.0),
+                if self.last_contrarian_active.unwrap_or(false) { "on" } else { "off" },
+                self.last_meta_q_pred.unwrap_or(0.0),
+            );
+
+            // C1/P1: propagate health to GPU replay buffer for diversity-weighted priorities.
+            {
+                let mut agent = self.agent.write().await;
+                agent.primary_dqn_mut().memory.gpu.set_learning_health(health_value);
+            }
+
+            // D4/N4: propagate health to GPU experience collector for adaptive CF curriculum.
+            if let Some(ref mut collector) = self.gpu_experience_collector {
+                collector.set_learning_health(health_value);
+                self.last_cf_ratio_eff = Some(collector.last_cf_ratio_eff());
+            }
+        }
+
         // Feed per-step Q-value extremes into monitor (real min/max, not single-sample).
         if self.epoch_q_min.is_finite() && self.epoch_q_max.is_finite() {
             tracing::debug!(
@@ -2375,6 +2702,8 @@ impl DQNTrainer {
             }
 
             epoch_max_dd = financials.max_drawdown;
+            // D7/N7: carry win_rate into next epoch's process_epoch_boundary.
+            self.last_epoch_win_rate = financials.win_rate as f32;
             financials.sharpe
         };