diff --git a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu index 8db66c077..f54bdffc6 100644 --- a/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu +++ b/crates/ml/src/cuda_pipeline/c51_loss_kernel.cu @@ -571,9 +571,6 @@ extern "C" __global__ void c51_loss_batched( /* ── Spectral decoupling: L2 penalty on Q-value logit magnitudes ── */ float spectral_decoupling_lambda, /* Pezeshki et al. 2021 (0.0=disabled, 0.01=default) */ - /* ── Online-target Q-divergence for adaptive tau ── */ - float* __restrict__ q_divergence, /* [1] output: sum of (E[Q_online]-E[Q_target])^2, written by reduction kernel (no atomicAdd) */ - /* ── CVaR alpha from IQN readiness ── */ const float* __restrict__ iqn_readiness_ptr, /* [1] pinned device-mapped for CVaR alpha */ @@ -1135,8 +1132,7 @@ extern "C" __global__ void c51_loss_batched( float weighted_loss = clamped_ce * is_weight; per_sample_loss[sample_id] = (weighted_loss); td_errors[sample_id] = clamped_ce; - /* total_loss reduced by separate deterministic kernel — no atomicAdd. - * q_divergence removed from hot path — zero atomicAdd. */ + /* total_loss reduced by separate deterministic kernel — no atomicAdd. */ } } @@ -1145,7 +1141,6 @@ extern "C" __global__ void c51_loss_batched( * * Sums per_sample_loss[0..batch_size-1] / batch_size into total_loss[0]. * Single-block sequential reduction — fully deterministic (fixed summation order). - * Also sums q_divergence contributions if q_div_per_sample is non-null. * * Launch: grid=(1), block=(1). Trivially deterministic. * ══════════════════════════════════════════════════════════════════════ */ diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 2b0f648e1..67223df24 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2664,9 +2664,6 @@ pub struct GpuDqnTrainer { /// MSE loss accumulator [1] — pinned device-mapped. GPU writes via dev_ptr, CPU reads via host ptr. mse_loss_pinned: *mut f32, mse_loss_dev_ptr: u64, - /// [1] Mean squared Q-divergence — pinned device-mapped. GPU writes via dev_ptr, CPU reads via host ptr. - q_divergence_pinned: *mut f32, - q_divergence_dev_ptr: u64, // ── Forward-only Q-value output ───────────────────────────────── q_out_buf: CudaSlice, // [B, TOTAL_ACTIONS(11)] @@ -2984,12 +2981,6 @@ pub struct GpuDqnTrainer { /// Device-mapped view of `outlier_diag_pinned`. Written by /// `update_adaptive_clip_kernel`; read host-side after launch. outlier_diag_dev_ptr: u64, - /// EMA of Q-divergence for adaptive tau computation. - /// NOTE: only consumed by the orphan helper `compute_adaptive_tau` (zero - /// production callers as of 2026-05-01); per `feedback_wire_everything_up.md` - /// either wire the helper into the production path or delete it. Tracked - /// as a separate item from the SP4 Layer C close-out sweep. - q_div_ema: f32, // ── Training state ────────────────────────────────────────────── pub(crate) adam_step: i32, @@ -3042,8 +3033,7 @@ pub struct GpuDqnTrainer { /// [8]=atom_entropy, [9]=atom_utilization (every 50 steps, 7 q_stats) /// [10]=causal_mean_sens (every N steps) /// [11]=ensemble_diversity_loss (epoch boundary) - /// [12]=q_divergence (per-step, online vs target MSE) - /// [13..16]=reserved (future use) + /// [12..16]=reserved (future use) /// Pinned memory enables true async cuMemcpyDtoHAsync without CPU blocking. readback_pinned: *mut f32, /// Whether there's an in-flight readback to collect @@ -4515,14 +4505,6 @@ impl GpuDqnTrainer { // of the model architecture and loss function, not the data window. // Resetting it would leave the first fold-2 epochs unprotected. - // Reset Q-divergence EMA — fold 2's divergence baseline differs from fold 1. - self.q_div_ema = 0.0; - unsafe { - cudarc::driver::sys::cuMemsetD32Async( - self.q_divergence_dev_ptr, 0, 1, self.stream.cu_stream(), - ); - } - tracing::info!("Adam optimizer + PopArt + grad_clip state reset for new fold"); Ok(()) } @@ -4816,9 +4798,6 @@ impl Drop for GpuDqnTrainer { if !self.per_branch_q_stats_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.per_branch_q_stats_pinned.cast()) }; } - if !self.q_divergence_pinned.is_null() { - let _ = unsafe { cudarc::driver::result::free_host(self.q_divergence_pinned.cast()) }; - } if !self.iqn_readiness_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.iqn_readiness_pinned.cast()) }; } @@ -7022,13 +7001,6 @@ impl GpuDqnTrainer { self.readback_pinned } - /// Read Q-divergence from pinned readback buffer (offset [12]). - /// Returns the mean squared Q-divergence between online and target networks. - /// Pinned device-mapped: reads directly from host pointer (no DtoH copy). - pub(crate) fn q_divergence_readback(&self) -> f32 { - unsafe { *self.q_divergence_pinned } - } - /// Current IQN readiness scalar [0, 1]. /// 0 = uncertain/exploring, 1 = converged/exploiting. /// Used by the experience collector's quantile_q_select kernel. @@ -7039,27 +7011,6 @@ impl GpuDqnTrainer { /// Device pointer to IQN readiness [1] pinned — for plan enforcement gating in env_step. pub fn iqn_readiness_dev_ptr(&self) -> u64 { self.iqn_readiness_dev_ptr } - /// Compute adaptive tau from Q-divergence signal. - /// When divergence is high (networks drifted apart), increase tau to close the gap faster. - /// Uses EMA-smoothed divergence as the baseline. - pub fn compute_adaptive_tau(&mut self, base_tau: f32) -> f32 { - let raw_div = unsafe { *self.q_divergence_pinned }; - // EMA smooth the divergence signal - const DIV_BETA: f32 = 0.95; - if self.q_div_ema <= 0.0 { - self.q_div_ema = raw_div.max(1e-10); - } else { - self.q_div_ema = DIV_BETA * self.q_div_ema + (1.0 - DIV_BETA) * raw_div; - } - // Scale tau: when divergence is high, increase tau to close the gap faster - // baseline_div is the "normal" divergence level — tau stays at base_tau - let baseline_div = self.q_div_ema; // EMA IS the baseline - let div_ratio = raw_div / baseline_div.max(1e-10); - // tau range: [base_tau * 0.5, base_tau * 10] - let scaled_tau = base_tau * div_ratio.sqrt().clamp(0.5, 10.0); - scaled_tau - } - /// Reference to saved h_s2 activations from the last forward pass. /// /// Shape: `[B, SHARED_H2]` — the shared trunk output for online states. @@ -11621,19 +11572,6 @@ impl GpuDqnTrainer { cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, mse_loss_pinned.cast(), 0); dp }; - // q_divergence — pinned device-mapped (zero-copy readback). - let q_divergence_pinned: *mut f32 = unsafe { - let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; - cudarc::driver::result::malloc_host(std::mem::size_of::(), flags) - .map_err(|e| MLError::ModelError(format!("pinned q_divergence alloc: {e}")))? - as *mut f32 - }; - unsafe { *q_divergence_pinned = 0.0; } - let q_divergence_dev_ptr = unsafe { - let mut dp = 0u64; - cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, q_divergence_pinned.cast(), 0); - dp - }; // iqn_readiness — pinned device-mapped (CVaR alpha for c51_loss_kernel). let iqn_readiness_pinned: *mut f32 = unsafe { let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP; @@ -14930,8 +14868,6 @@ impl GpuDqnTrainer { total_loss_dev_ptr, mse_loss_pinned, mse_loss_dev_ptr, - q_divergence_pinned, - q_divergence_dev_ptr, q_out_buf, eval_td_snapshot, eval_loss_snapshot, @@ -15018,7 +14954,6 @@ impl GpuDqnTrainer { grad_norm_ema_dev_ptr, outlier_diag_pinned, outlier_diag_dev_ptr, - q_div_ema: 0.0, adam_step: 0, total_params, non_isv_params, @@ -19573,11 +19508,6 @@ impl GpuDqnTrainer { self.launch_mse_grad_to_scratch()?; // C51 path → main buffers (already zeroed above) - unsafe { - cudarc::driver::sys::cuMemsetD32Async( - self.q_divergence_dev_ptr, 0, 1, self.stream.cu_stream(), - ); - } self.fill_gamma_buf()?; self.launch_c51_loss()?; self.launch_loss_reduce(self.total_loss_dev_ptr)?; @@ -19733,11 +19663,6 @@ impl GpuDqnTrainer { self.launch_mse_grad_to_scratch()?; // C51 path - unsafe { - cudarc::driver::sys::cuMemsetD32Async( - self.q_divergence_dev_ptr, 0, 1, self.stream.cu_stream(), - ); - } self.fill_gamma_buf()?; self.launch_c51_loss()?; self.launch_loss_reduce(self.total_loss_dev_ptr)?; @@ -20661,8 +20586,6 @@ impl GpuDqnTrainer { .arg(&self.ensemble_disagreement_weight) // ── Spectral decoupling (1) ── .arg(&spectral_lambda_f32) - // ── Q-divergence accumulator (1) ── - .arg(&self.q_divergence_dev_ptr) // ── CVaR alpha from IQN readiness ── .arg(&self.iqn_readiness_dev_ptr) // ── Adam step counter for stochastic Expected SARSA ── diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 594b518d3..c5270030e 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3135,3 +3135,9 @@ Sweep verification: - `cargo test -p ml --test sp4_producer_unit_tests --offline --release -- --ignored` 16/16 SP4 GPU tests pass on RTX 3050 Ti at each commit. Refs: SP4 Layer C close-out comprehensive sweep — `feedback_no_cpu_compute_strict.md` (saved 2026-05-01) zero-tolerance enforcement in SP4-touched code paths going forward. + +#### Sweep close-out (2026-05-01 follow-up) + +Site #2 (`compute_adaptive_tau` orphan) deleted per `feedback_wire_everything_up`. The pre-SP3 helper had zero production callers; the canonical adaptive-tau path is the SP3/SP4 ISV-driven `tau_kernel.cu` + `ISV[TAU_INDEX]` controller. Removed: `compute_adaptive_tau` method, `q_divergence_readback` accessor (only called by the deleted helper), `q_div_ema` field + initialiser + fold reset, `q_divergence_pinned` + `q_divergence_dev_ptr` (allocation, free, struct fields, two memsets at C51 launch sites, kernel arg in `launch_c51_loss`), and the inert `q_divergence` parameter from `c51_loss_batched` (the kernel never wrote to it; the comment "removed from hot path — zero atomicAdd" had been confirming this for releases). Behavior preservation is trivial — dead code only. + +Refs: SP4 Layer C close-out follow-up — site #2 of the sweep audit grid resolved.