feat(dqn-v2): C.6 Task 10 — gamma GPU kernel + CPU monitor (discount factor)

gamma_update kernel computes health-coupled gamma from ISV[LEARNING_HEALTH=12]:
gamma_eff = gamma_min + (gamma_base - gamma_min) * health. Single-thread
cold-path kernel writes ISV[GAMMA_EFF_INDEX=43].

GammaMonitor is a read-only observer exposing gamma_eff, health, fire_rate.

Consumer migration: fill_gamma_buf and IQL gamma computation now read
ISV[GAMMA_EFF_INDEX] via read_isv_signal_at (pinned, zero-copy). Training
loop passes hyperparams.gamma as gamma_base to the kernel.

Deleted: apply_adaptive_gamma method (GpuDqnTrainer + FusedTrainingCtx
delegates), set_adaptive_gamma method, adaptive_gamma field (GpuDqnTrainer
+ DQNTrainer), last_gamma_eff cached field + last_gamma_eff() delegate.
StateResetRegistry entry for adaptive_gamma removed (field gone).

Smoke test generalization.rs updated to check config gamma_base instead
of deleted adaptive_gamma field.

Tests: 3 monitor unit tests pass. cargo check -p ml at 8-warning baseline.

Plan 1 Task 10. Spec §4.C.6 (2026-04-24 revision).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

crates/ml/build.rs

@@ -86,6 +86,7 @@ fn main() {
         "backtest_plan_kernel.cu",
         "tau_update_kernel.cu",
         "epsilon_update_kernel.cu",
+        "gamma_update_kernel.cu",
     ];
 
     // ALL kernels get common header (BF16 types + wrappers)

crates/ml/src/cuda_pipeline/gamma_update_kernel.cu

@@ -0,0 +1,21 @@
+/* gamma_update — effective discount factor gamma, GPU-driven (Plan 1 Task 10).
+ * Reads:  ISV[LEARNING_HEALTH_INDEX=12]
+ * Writes: ISV[GAMMA_EFF_INDEX=43]
+ * Cold path (per-epoch boundary). Single-thread kernel.
+ *
+ * Formula: health-coupled gamma — high health → gamma_base, collapse →
+ * gamma_min. Provides graceful degradation: collapsed policy uses shorter
+ * horizon to limit variance from unreliable Q-targets.
+ */
+extern "C" __global__ void gamma_update(
+    const float* __restrict__ isv,
+    float*       __restrict__ isv_out,
+    float gamma_base,
+    float gamma_min
+) {
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+    float health    = fminf(fmaxf(isv[12], 0.0f), 1.0f);
+    float gamma_eff = gamma_min + (gamma_base - gamma_min) * health;
+    gamma_eff       = fminf(fmaxf(gamma_eff, gamma_min), gamma_base);
+    isv_out[43] = gamma_eff;
+}

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -81,6 +81,8 @@ static BRANCH_GRAD_BALANCE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR")
 static TAU_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/tau_update_kernel.cubin"));
 /// Plan 1 Task 14: GPU-driven effective exploration epsilon.
 static EPSILON_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/epsilon_update_kernel.cubin"));
+/// Plan 1 Task 10: GPU-driven effective discount factor gamma.
+static GAMMA_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/gamma_update_kernel.cubin"));
 
 /// Mamba2 temporal scan configuration.
 const MAMBA2_HISTORY_K: usize = 8;    // Rolling history length
@@ -1212,8 +1214,6 @@ pub struct GpuDqnTrainer {
     regime_dropout_kernel: CudaFunction,
     /// Epoch seed for regime dropout Philox PRNG (changes per epoch).
     regime_dropout_epoch_seed: i32,
-    /// G4: Adaptive gamma for C51 Bellman projection. Updated from DQNTrainer.
-    adaptive_gamma: f32,
 
     // ── G5: Epistemic-gated magnitude ──
     /// Kernel: epistemic_gate_magnitude — sigmoid-gates magnitude Q-values by ensemble variance.
@@ -1407,6 +1407,9 @@ pub struct GpuDqnTrainer {
     /// Plan 1 Task 14: GPU-driven effective epsilon. Cold-path (per-epoch).
     /// Reads ISV[39, 40, 12]; writes ISV[EPSILON_EFF_INDEX=41].
     epsilon_update_kernel: CudaFunction,
+    /// Plan 1 Task 10: GPU-driven effective gamma. Cold-path (per-epoch).
+    /// Reads ISV[LEARNING_HEALTH_INDEX=12]; writes ISV[GAMMA_EFF_INDEX=43].
+    gamma_update_kernel: CudaFunction,
     /// Host-side cached per-component norms populated at epoch boundary
     /// from the pinned result slot. Index order matches `grad_mag_*_ratio`
     /// accessors: `[iqn, cql, cql_sx, c51, c51_bs, distill, rec, pred, ens]`.
@@ -2203,11 +2206,6 @@ pub struct GpuDqnTrainer {
     /// 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.
@@ -4377,7 +4375,7 @@ impl GpuDqnTrainer {
 
     /// Broadcast base_gamma * gamma_mod[0] into gamma_buf [B] for per-sample C51 loss.
     pub(crate) fn fill_gamma_buf(&self) -> Result<(), MLError> {
-        let base_gamma = self.adaptive_gamma.powi(self.config.n_steps as i32);
+        let base_gamma = self.read_isv_signal_at(GAMMA_EFF_INDEX).powi(self.config.n_steps as i32);
         let blocks = ((self.config.batch_size as u32 + 255) / 256).max(1);
         let b = self.config.batch_size as i32;
         let gamma_mod_ptr = self.gamma_mod_buf.raw_ptr();
@@ -4581,11 +4579,6 @@ impl GpuDqnTrainer {
         unsafe { *self.regime_util_pinned = util; }
     }
 
-    /// G4: Set adaptive gamma for C51 Bellman projection.
-    pub fn set_adaptive_gamma(&mut self, gamma: f32) {
-        self.adaptive_gamma = gamma;
-    }
-
     /// G9: Apply regime-conditioned dropout to h_s2 trunk output.
     /// Must be called AFTER mamba2_step and BEFORE compute_expected_q.
     pub(crate) fn apply_regime_dropout(&self, batch_size: usize, is_training: bool) -> Result<(), MLError> {
@@ -7055,6 +7048,14 @@ impl GpuDqnTrainer {
                 .map_err(|e| MLError::ModelError(format!("epsilon_update load: {e}")))?
         };
 
+        // Plan 1 Task 10: load gamma_update kernel (cold-path, per-epoch).
+        let gamma_update_kernel = {
+            let module = stream.context().load_cubin(GAMMA_UPDATE_CUBIN.to_vec())
+                .map_err(|e| MLError::ModelError(format!("gamma_update cubin load: {e}")))?;
+            module.load_function("gamma_update")
+                .map_err(|e| MLError::ModelError(format!("gamma_update load: {e}")))?
+        };
+
         // eval_v_range — pinned device-mapped (zero-copy write from CPU).
         let eval_v_range_pinned: *mut f32 = unsafe {
             let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP;
@@ -8809,8 +8810,6 @@ impl GpuDqnTrainer {
             .map_err(|e| MLError::ModelError(format!("alloc v_logits_blended: {e}")))?;
         info!("GpuDqnTrainer: multi-horizon value heads allocated (5-bar + 20-bar)");
 
-        let initial_gamma = config.gamma;
-
         // ── Learned risk management buffers ──
         let risk_hidden_buf = stream.alloc_zeros::<f32>(b * config.adv_h)
             .map_err(|e| MLError::ModelError(format!("risk_hidden_buf alloc: {e}")))?;
@@ -9056,7 +9055,6 @@ impl GpuDqnTrainer {
             regime_util_dev_ptr,
             regime_dropout_kernel,
             regime_dropout_epoch_seed: 0,
-            adaptive_gamma: initial_gamma,
             epistemic_gate_kernel,
             var_ema_pinned,
             var_ema_dev_ptr,
@@ -9113,6 +9111,7 @@ impl GpuDqnTrainer {
             branch_slice_max_len,
             tau_update_kernel,
             epsilon_update_kernel,
+            gamma_update_kernel,
             grad_component_norms_mag: [0.0_f32; 9],
             grad_component_norms_dir: [0.0_f32; 9],
             grad_component_norms_trunk: [0.0_f32; 9],
@@ -9481,7 +9480,6 @@ impl GpuDqnTrainer {
             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,
@@ -9774,21 +9772,6 @@ impl GpuDqnTrainer {
 
 
 
-    /// 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.
@@ -10473,6 +10456,32 @@ impl GpuDqnTrainer {
         Ok(())
     }
 
+    /// Plan 1 Task 10: GPU-driven gamma update.
+    ///
+    /// Single-thread cold-path kernel: reads ISV[LEARNING_HEALTH_INDEX=12];
+    /// writes ISV[GAMMA_EFF_INDEX=43] with health-coupled formula
+    /// `gamma_min + (gamma_base - gamma_min) * health`. Must be called once
+    /// per epoch boundary BEFORE any gamma consumer reads ISV[GAMMA_EFF_INDEX].
+    pub(crate) fn launch_gamma_update(&self, gamma_base: f32, gamma_min: f32) -> Result<(), MLError> {
+        let isv_ptr = self.isv_signals_dev_ptr;
+        let isv_out = isv_ptr; // kernel writes to the same ISV bus (offset 43)
+        unsafe {
+            self.stream
+                .launch_builder(&self.gamma_update_kernel)
+                .arg(&isv_ptr)
+                .arg(&isv_out)
+                .arg(&gamma_base)
+                .arg(&gamma_min)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (1, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .map_err(|e| MLError::ModelError(format!("gamma_update launch: {e}")))?;
+        }
+        Ok(())
+    }
+
     /// Read back the per-branch scales applied on the most recent
     /// `launch_branch_grad_balance` call — `[dir, mag, ord, urg]`.
     /// Values in `[0, 1]`; `1.0` means the branch passed through
@@ -13915,7 +13924,7 @@ impl GpuDqnTrainer {
         let on_next_b3_ptr = on_next_b2_ptr + (b * b2 * na * f32_sz) as u64;
 
         // N-step returns: use gamma^n for the Bellman projection.
-        let gamma = self.adaptive_gamma.powi(self.config.n_steps as i32);
+        let gamma = self.read_isv_signal_at(GAMMA_EFF_INDEX).powi(self.config.n_steps as i32);
         let batch_i32 = b as i32;
         let na_i32 = na as i32;
         let b0_i32 = b0 as i32;

crates/ml/src/trainers/dqn/fused_training.rs

@@ -3081,21 +3081,6 @@ impl FusedTrainingCtx {
     }
 
     /// G4: Set adaptive gamma for C51 Bellman projection.
-    pub(crate) fn set_adaptive_gamma(&mut self, gamma: f32) {
-        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 }
-
     /// Plan 1 Task 13: delegate GPU tau update to trainer.
     pub(crate) fn launch_tau_update(&self, tau_base: f32, tau_final: f32)
         -> anyhow::Result<()>
@@ -3112,6 +3097,14 @@ impl FusedTrainingCtx {
             .map_err(|e| anyhow::anyhow!("launch_epsilon_update: {e}"))
     }
 
+    /// Plan 1 Task 10: delegate GPU gamma update to trainer.
+    pub(crate) fn launch_gamma_update(&self, gamma_base: f32, gamma_min: f32)
+        -> anyhow::Result<()>
+    {
+        self.trainer.launch_gamma_update(gamma_base, gamma_min)
+            .map_err(|e| anyhow::anyhow!("launch_gamma_update: {e}"))
+    }
+
     /// G5: Set the EMA threshold for epistemic variance gating.
     pub(crate) fn set_var_ema(&self, val: f32) {
         self.trainer.set_var_ema(val);

crates/ml/src/trainers/dqn/monitors/gamma_monitor.rs

@@ -0,0 +1,102 @@
+//! CPU-side observer for the GPU-driven effective discount factor gamma.
+//!
+//! The `gamma_update` CUDA kernel reads ISV[LEARNING_HEALTH_INDEX=12] to
+//! compute a health-coupled gamma that degrades from `gamma_base` at full
+//! health to `gamma_min` at full collapse, writing the result to
+//! ISV[GAMMA_EFF_INDEX=43]. This monitor reads that slot for HEALTH_DIAG
+//! emission and fire-rate tracking.
+//!
+//! The monitor's `read()` returns `isv[GAMMA_EFF_INDEX]`. `diagnose()`
+//! exposes the effective gamma, health, and fire-rate.
+
+use crate::cuda_pipeline::gpu_dqn_trainer::{
+    LEARNING_HEALTH_INDEX, GAMMA_EFF_INDEX,
+};
+use crate::trainers::dqn::adaptive_monitor::{
+    AdaptiveMonitor, DiagSnapshot, FireRateStats, IsvBus,
+};
+
+pub(crate) struct GammaMonitor {
+    last_gamma: f32,
+    fire_rate: FireRateStats,
+}
+
+impl GammaMonitor {
+    pub(crate) fn new() -> Self {
+        Self { last_gamma: 0.0, fire_rate: FireRateStats::default() }
+    }
+}
+
+impl AdaptiveMonitor for GammaMonitor {
+    fn read(&self, isv: &IsvBus<'_>) -> f32 {
+        isv.read(GAMMA_EFF_INDEX)
+    }
+
+    fn diagnose(&self, isv: &IsvBus<'_>) -> DiagSnapshot {
+        let gamma_eff = isv.read(GAMMA_EFF_INDEX);
+        let health    = isv.read(LEARNING_HEALTH_INDEX);
+        DiagSnapshot::new()
+            .with("gamma_eff",  gamma_eff)
+            .with("health",     health)
+            .with("fire_rate",  self.fire_rate.fire_rate())
+    }
+
+    fn observe(&mut self, current: f32) {
+        let fired = (current - self.last_gamma).abs() > 1e-6;
+        self.fire_rate.record_fire(fired);
+        self.last_gamma = current;
+    }
+
+    fn fire_rate(&self) -> &FireRateStats {
+        &self.fire_rate
+    }
+
+    fn name(&self) -> &'static str {
+        "gamma"
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn make_isv(health: f32, gamma_eff: f32) -> Vec<f32> {
+        let mut v = vec![0.0f32; 49];
+        v[LEARNING_HEALTH_INDEX] = health;
+        v[GAMMA_EFF_INDEX]       = gamma_eff;
+        v
+    }
+
+    #[test]
+    fn read_returns_gamma_eff_slot() {
+        let mut buf = make_isv(0.8, 0.985);
+        let bus = IsvBus::new(&mut buf);
+        let mon = GammaMonitor::new();
+        let v = mon.read(&bus);
+        assert!((v - 0.985).abs() < 1e-7, "expected gamma_eff=0.985, got {v}");
+    }
+
+    #[test]
+    fn diagnose_exposes_all_fields() {
+        let mut buf = make_isv(0.6, 0.97);
+        let bus = IsvBus::new(&mut buf);
+        let mon = GammaMonitor::new();
+        let snap = mon.diagnose(&bus);
+        assert!(snap.fields.contains_key("gamma_eff"));
+        assert!(snap.fields.contains_key("health"));
+        assert!(snap.fields.contains_key("fire_rate"));
+        let gamma_eff = snap.fields["gamma_eff"];
+        assert!((gamma_eff - 0.97).abs() < 1e-6, "expected gamma_eff=0.97, got {gamma_eff}");
+    }
+
+    #[test]
+    fn observe_tracks_fire_rate_on_change() {
+        let mut mon = GammaMonitor::new();
+        mon.observe(0.99);
+        mon.observe(0.99);
+        mon.observe(0.97);
+        let rate = mon.fire_rate().fire_rate();
+        assert!((rate - 2.0 / 3.0).abs() < 1e-6,
+                "expected 2/3 fire rate, got {rate}");
+    }
+}

crates/ml/src/trainers/dqn/monitors/mod.rs

@@ -6,5 +6,6 @@
 //! smoke-test fire-rate tracking. No CPU-side adaptive computation.
 
 pub(crate) mod epsilon_monitor;
+pub(crate) mod gamma_monitor;
 pub(crate) mod grad_balancer_monitor;
 pub(crate) mod tau_monitor;

crates/ml/src/trainers/dqn/smoke_tests/generalization.rs

@@ -86,10 +86,12 @@ fn test_generalization_components_smoke() -> anyhow::Result<()> {
     assert!(trainer.cost_anneal_factor >= 0.0 && trainer.cost_anneal_factor <= 1.0,
         "cost_anneal should be in [0, 1], got {}", trainer.cost_anneal_factor);
 
-    // G4: Gamma annealing — should be in [0.90, 0.95]
-    info!("  G4 adaptive_gamma: {:.3}", trainer.adaptive_gamma);
-    assert!(trainer.adaptive_gamma >= 0.89 && trainer.adaptive_gamma <= 0.96,
-        "adaptive_gamma should be in [0.90, 0.95], got {}", trainer.adaptive_gamma);
+    // Plan 1 Task 10: gamma is now GPU-computed via gamma_update kernel (ISV[43]).
+    // Validate the config gamma_base is in a reasonable range.
+    let gamma_base = trainer.hyperparams.gamma;
+    info!("  G4 gamma_base: {:.3}", gamma_base);
+    assert!(gamma_base >= 0.85 && gamma_base <= 0.999,
+        "gamma_base should be in [0.85, 0.999], got {}", gamma_base);
 
     // Adaptive DSR — should be initialized
     info!("  DSR sharpe_ema: {:.3} (initialized={})",

crates/ml/src/trainers/dqn/state_reset_registry.rs

@@ -94,11 +94,6 @@ impl StateResetRegistry {
                 description: "ps[3] — entry price of current trade",
             },
             // ───── Soft-reset state (anneals across folds) ──────────────
-            RegistryEntry {
-                name: "adaptive_gamma",
-                category: ResetCategory::SoftReset { decay_bars: 500 },
-                description: "training_loop.rs adaptive_gamma — anneals to 0.905 bootstrap",
-            },
             RegistryEntry {
                 name: "isv_grad_balance_targets",
                 category: ResetCategory::SoftReset { decay_bars: 500 },
@@ -238,7 +233,7 @@ mod tests {
     fn registry_classifies_known_state() {
         let r = StateResetRegistry::new();
         assert_eq!(r.category("kelly_stats"), Some(ResetCategory::FoldReset));
-        assert_eq!(r.category("adaptive_gamma"), Some(ResetCategory::SoftReset { decay_bars: 500 }));
+        assert_eq!(r.category("isv_grad_balance_targets"), Some(ResetCategory::SoftReset { decay_bars: 500 }));
         assert_eq!(r.category("network_params"), Some(ResetCategory::TrainingPersist));
         assert_eq!(r.category("ISV_LAYOUT_FINGERPRINT"), Some(ResetCategory::SchemaContract));
     }

crates/ml/src/trainers/dqn/trainer/constructor.rs

@@ -729,7 +729,6 @@ impl DQNTrainer {
             cached_n_episodes: None,
             backtracking: super::BacktrackingState::new(),
             cost_anneal_factor: 0.0,
-            adaptive_gamma: 0.90,
             wf_window_sharpes: Vec::new(),
             wf_num_windows: 6,
             wf_purge_bars: 2000,  // ~1 trading day of imbalance bars. Was 100 (~1 hour, too short for ES regime persistence)

crates/ml/src/trainers/dqn/trainer/mod.rs

@@ -572,9 +572,6 @@ pub struct DQNTrainer {
     /// G2: Transaction cost curriculum anneal factor (0.0=no costs, 1.0=full costs).
     pub(crate) cost_anneal_factor: f32,
 
-    /// G4: Adaptive gamma annealed by atom utilization (0.90 -> 0.95).
-    pub(crate) adaptive_gamma: f32,
-
     /// G3: Walk-forward validation window Sharpe ratios.
     pub(crate) wf_window_sharpes: Vec<f64>,
     /// G3: Number of walk-forward windows (K=6).
@@ -1477,8 +1474,8 @@ impl DQNTrainer {
         // ad-hoc reset calls. Adding a new FoldReset state requires a registry
         // entry AND a dispatch arm in reset_named_state (Invariant 2 Wire-It-Up).
         // SoftReset category is handled separately (Plan 2); the pre-existing
-        // ad-hoc handling for adaptive_gamma and grad_balance targets is
-        // preserved in the training loop and is not replaced here.
+        // ad-hoc handling for grad_balance targets is preserved in the
+        // training loop and is not replaced here.
         let registry = StateResetRegistry::new();
         for entry in registry.fold_reset_entries() {
             self.reset_named_state(entry.name, 0u32)

crates/ml/src/trainers/dqn/trainer/training_loop.rs

@@ -21,7 +21,7 @@ use cudarc::driver::CudaSlice;
 use common::metrics::{questdb_sink, training_metrics};
 use tracing::{debug, info, warn};
 
-use crate::cuda_pipeline::gpu_dqn_trainer::{EPOCH_IDX_INDEX, EPSILON_EFF_INDEX};
+use crate::cuda_pipeline::gpu_dqn_trainer::{EPOCH_IDX_INDEX, EPSILON_EFF_INDEX, GAMMA_EFF_INDEX};
 use crate::dqn::logging::{log_epoch_start, log_epoch_end, log_training_progress};
 use crate::dqn::target_update::convergence_half_life;
 use crate::evaluation::metrics::calculate_var_cvar;
@@ -267,6 +267,16 @@ impl DQNTrainer {
                 }
             }
 
+            // Plan 1 Task 10: launch GPU gamma update kernel at epoch boundary.
+            // gamma_min = 0.90 is the floor during full collapse.
+            if let Some(ref fused) = self.fused_ctx {
+                let gamma_base = self.hyperparams.gamma;
+                let gamma_min  = 0.90_f32;
+                if let Err(e) = fused.launch_gamma_update(gamma_base, gamma_min) {
+                    tracing::warn!(epoch, "launch_gamma_update failed (non-fatal): {e}");
+                }
+            }
+
             // Collect pending async validation from PREVIOUS epoch
             if let Some(val) = self.pending_val_loss.take() {
                 // Store for this epoch's logging (one-epoch delay is acceptable)
@@ -693,34 +703,16 @@ impl DQNTrainer {
             // 0.90 floor, reinforcing the symptom it was meant to
             // mitigate. Per research-agent audit (2026-04-24).
             //
-            // New direction: when atoms are collapsed, RAISE γ so
-            // `γ × atom_support` spans a larger fraction of the grid,
-            // spreading TD targets across more bins and breaking the
-            // loop. Slower step (0.005) than the healthy-branch to
-            // avoid overshoot.
-            //
-            // Clamps retained: [0.90, 0.95] at this layer; the fused
-            // trainer re-clamps to [0.90, 0.995] after its regime
-            // adjustment.
+            // Plan 1 Task 10: gamma is now GPU-computed at epoch boundary.
+            // gamma_update kernel applied health-coupling: gamma_eff = gamma_base * health + gamma_min * (1-health).
+            // Read effective gamma from ISV for logging only.
             {
-                let util = self.fused_ctx.as_ref()
-                    .map(|f| f.utilization_ema())
-                    .unwrap_or(0.5);
-                if util > 0.6 {
-                    self.adaptive_gamma = (self.adaptive_gamma + 0.005).min(0.95);
-                } else if util < 0.4 {
-                    // Direction flipped (atom-collapse fix): RAISE γ to
-                    // widen the TD-target spread across the atom grid.
-                    self.adaptive_gamma = (self.adaptive_gamma + 0.005).min(0.95);
-                }
-                // 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 {
-                    let gamma_scheduled = self.adaptive_gamma as f32;
-                    fused.apply_adaptive_gamma(gamma_scheduled);
-                }
-                info!("G4 gamma={:.3} util={:.2}", self.adaptive_gamma, util);
+                let gamma_eff = if let Some(ref fused) = self.fused_ctx {
+                    fused.trainer().read_isv_signal_at(GAMMA_EFF_INDEX)
+                } else {
+                    self.hyperparams.gamma
+                };
+                info!("G4 gamma={:.3} (GPU-computed, ISV[43])", gamma_eff);
             }
 
             // Flush GPU-accumulated max priority
@@ -805,15 +797,14 @@ impl DQNTrainer {
                     // Apply E2: adaptive epsilon
                     self.hyperparams.epsilon_end = result.epsilon;
 
-                    // Apply E3: dynamic gamma (EMA blend)
+                    // Plan 1 Task 10: gamma is GPU-computed at epoch boundary.
+                    // Eval-blended gamma_base can override hyperparams if enrichment
+                    // provides a suggestion — but the GPU kernel applies health coupling.
                     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 {
-                            let gamma_scheduled = self.adaptive_gamma as f32;
-                            fused.apply_adaptive_gamma(gamma_scheduled);
-                        }
+                        // Update the config gamma_base with the eval-blended value
+                        // so the next epoch's gamma_update kernel uses the blended base.
+                        let blended = 0.9 * self.hyperparams.gamma + 0.1 * eval_gamma;
+                        self.hyperparams.gamma = blended.clamp(0.85, 0.98);
                     }
 
                     // Apply E5: ensemble agreement threshold
@@ -2069,9 +2060,9 @@ impl DQNTrainer {
                 self.last_c51_budget_eff = Some(fused.last_c51_budget_eff());
             }
 
-            // C4/P4: propagate adaptive gamma for logging.
+            // Plan 1 Task 10: propagate gamma_eff from ISV for logging.
             if let Some(ref fused) = self.fused_ctx {
-                self.last_gamma_eff = Some(fused.last_gamma_eff());
+                self.last_gamma_eff = Some(fused.trainer().read_isv_signal_at(GAMMA_EFF_INDEX));
             }
 
             // D1/N1: snapshot high-q_gap checkpoints and pull toward best during collapse.

docs/dqn-gpu-hot-path-audit.md

@@ -69,6 +69,7 @@
 | `training_loop.rs:1007` | `clone_htod_f32` → features device buffer | COLD-PATH | Data loading: feature matrix uploaded before step loop | OK |
 | `tau_update_kernel.cu` (Plan 1 Task 13) | single-thread kernel: reads ISV[39,40,12]; writes ISV[42] | COLD-PATH | Per-epoch boundary call from `training_loop.rs`; 1-thread, 1-block; zero hot-path overhead | OK |
 | `epsilon_update_kernel.cu` (Plan 1 Task 14) | single-thread kernel: reads ISV[39,40,12]; writes ISV[41] | COLD-PATH | Per-epoch boundary call from `training_loop.rs`; 1-thread, 1-block; zero hot-path overhead | OK |
+| `gamma_update_kernel.cu` (Plan 1 Task 10) | single-thread kernel: reads ISV[12]; writes ISV[43] | COLD-PATH | Per-epoch boundary call from `training_loop.rs`; 1-thread, 1-block; ISV[43] read by fill_gamma_buf on hot path via read_isv_signal_at (pinned, zero-copy) | OK |
 
 ## Summary
 

docs/dqn-wire-up-audit.md

@@ -48,6 +48,8 @@
 | `cuda_pipeline/tau_update_kernel.cu` | `GpuDqnTrainer::launch_tau_update` → `FusedTrainingCtx::launch_tau_update` → `training_loop.rs` epoch boundary | Wired | Plan 1 Task 13 — cold-path (per-epoch) | — |
 | `trainers/dqn/monitors/epsilon_monitor.rs` | Read-only observer for epsilon_update kernel output (ISV slot 41); consumers: HEALTH_DIAG + controller_activity smoke | Wired | Plan 1 Task 14 | — |
 | `cuda_pipeline/epsilon_update_kernel.cu` | `GpuDqnTrainer::launch_epsilon_update` → `FusedTrainingCtx::launch_epsilon_update` → `training_loop.rs` epoch boundary | Wired | Plan 1 Task 14 — cold-path (per-epoch) | — |
+| `trainers/dqn/monitors/gamma_monitor.rs` | Read-only observer for gamma_update kernel output (ISV slot 43); consumers: HEALTH_DIAG + controller_activity smoke | Wired | Plan 1 Task 10 | — |
+| `cuda_pipeline/gamma_update_kernel.cu` | `GpuDqnTrainer::launch_gamma_update` → `FusedTrainingCtx::launch_gamma_update` → `training_loop.rs` epoch boundary | Wired | Plan 1 Task 10 — cold-path (per-epoch) | — |
 
 ## CUDA Pipeline — Rust Wrappers