gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

chore(sp4): delete compute_adaptive_tau orphan helper per feedback_wire_everything_up

Browse Source 
`compute_adaptive_tau` (gpu_dqn_trainer.rs:7045) had zero production
callers — its q_div_ema field's doc-comment explicitly flagged it
"NOTE: only consumed by the orphan helper compute_adaptive_tau (zero..)".
Pre-SP3 attempt at adaptive-tau control, superseded by the SP3/SP4 ISV-
driven tau controller (`tau_kernel.cu` + ISV[TAU_INDEX]).

Removed:
  - compute_adaptive_tau method (CPU EMA + ratio-scaled tau formula)
  - q_divergence_readback method (only caller was compute_adaptive_tau)
  - q_div_ema field + initialiser + fold-reset assignment
  - q_divergence_pinned + q_divergence_dev_ptr struct fields, alloc, free
  - q_divergence_dev_ptr memsets at C51 launch sites (2)
  - q_divergence_dev_ptr arg in launch_c51_loss kernel call
  - c51_loss_batched kernel parameter `q_divergence` (never written by
    kernel body — comment "removed from hot path — zero atomicAdd"
    confirmed buffer was inert; deleting parameter keeps kernel ABI clean
    per feedback_no_partial_refactor)
  - Stale doc-comments referencing q_divergence in c51 reduction kernel
  - readback_pinned [12]=q_divergence layout doc (slot was never read)

cargo check clean. SP4 + state_reset_registry lib tests pass (11/11).
16/16 SP4 producer GPU tests pass on RTX 3050 Ti. No behavior change —
dead code only.

Refs: feedback_no_cpu_compute_strict sweep audit (commit 6a6b58aec)
flagged this as orphan; feedback_wire_everything_up mandates deletion.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 16:05:39 +02:00
parent 6a6b58aec5
commit a385c1d2be
3 changed files with 8 additions and 84 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
crates/ml/src/cuda_pipeline/c51_loss_kernel.cu
View File

@@ -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.
* ══════════════════════════════════════════════════════════════════════ */

79
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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<f32>, // [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::<f32>(), 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 ──

6
docs/dqn-wire-up-audit.md
View File

@@ -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.