diff --git a/crates/ml/build.rs b/crates/ml/build.rs index bffc32839..5137bb03f 100644 --- a/crates/ml/build.rs +++ b/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) diff --git a/crates/ml/src/cuda_pipeline/gamma_update_kernel.cu b/crates/ml/src/cuda_pipeline/gamma_update_kernel.cu new file mode 100644 index 000000000..87253fea9 --- /dev/null +++ b/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; +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 9976a3101..19ba5878e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/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::(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; diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 360631a3c..1eba72ae7 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/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); diff --git a/crates/ml/src/trainers/dqn/monitors/gamma_monitor.rs b/crates/ml/src/trainers/dqn/monitors/gamma_monitor.rs new file mode 100644 index 000000000..273bf9b84 --- /dev/null +++ b/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 { + 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}"); + } +} diff --git a/crates/ml/src/trainers/dqn/monitors/mod.rs b/crates/ml/src/trainers/dqn/monitors/mod.rs index 140333d87..79fc05f2d 100644 --- a/crates/ml/src/trainers/dqn/monitors/mod.rs +++ b/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; diff --git a/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs b/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs index 12ab4018f..731246b9c 100644 --- a/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs +++ b/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={})", diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 6de186b28..5680be729 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/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)); } diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index a19f5e0c7..204bc9f90 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/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) diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 4ffa766d4..5504ba7a0 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/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, /// 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) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 377ce4c7b..99dc486c1 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/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. diff --git a/docs/dqn-gpu-hot-path-audit.md b/docs/dqn-gpu-hot-path-audit.md index a732739ec..e59af6922 100644 --- a/docs/dqn-gpu-hot-path-audit.md +++ b/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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index ec0b60eb1..65f64d363 100644 --- a/docs/dqn-wire-up-audit.md +++ b/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