From 605a8f5268c26c4497b6794cfbe3ee3d81cbd72f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 14:49:47 +0200 Subject: [PATCH] fix(sp4): migrate IQN readiness gauge update to GPU per feedback_no_cpu_compute_strict MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Layer C close-out C2 — the host-side EMA arithmetic block in GpuDqnTrainer::update_iqn_readiness violated feedback_no_cpu_compute_strict (any compute — EMA, reduction, mean — belongs on GPU regardless of frequency). The cold-start sentinel + adaptive-α EMA + improvement-gauge formula now run GPU-side via update_iqn_readiness_kernel (single-thread, single-block, mirrors update_grad_norm_emas_kernel shape). The kernel takes the host-passed iqn_loss scalar (from gpu_iqn.read_total_loss() mapped-pinned readback) and updates three coupled mapped-pinned scalars (iqn_loss_initial_pinned, iqn_loss_ema_pinned, iqn_readiness_pinned) in-place. __threadfence_system() guarantees PCIe-visibility to mapped-pinned host_ptr accessors and to other-stream kernel reads via dev_ptrs. Storage: iqn_loss_initial and iqn_loss_ema migrated from host-resident f32 fields to mapped-pinned device-mapped scalars (matches grad_norm_fast/slow_ema pattern from C1 redesigned). New accessors iqn_loss_ema_value() and iqn_loss_initial_value() read through the host_ptrs. iqn_readiness pinned slot remains the same — c51_loss_kernel CVaR α consumer dev_ptr unchanged. Cold-start sentinel (prev_initial < 1e-12 ⇒ assign loss directly + readiness=0) preserves the deleted host-side bootstrap branch exactly. Same adaptive-α formula clamp(|err|/(|err|+0.01), 0.01, 0.30) and improvement gauge clamped to [0,1]. state_reset_registry entries updated to reflect the new mapped-pinned storage and producer relocation; reset_iqn_readiness_state writes 0.0 through the pinned slots so the next kernel launch re-enters the bootstrap branch. Verification: SP4 lib tests + 16 SP4 GPU producer unit tests pass on RTX 3050 Ti. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 15 ++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 213 ++++++++++++++++-- .../update_iqn_readiness_kernel.cu | 84 +++++++ .../src/trainers/dqn/state_reset_registry.rs | 6 +- docs/dqn-wire-up-audit.md | 44 ++++ 5 files changed, 342 insertions(+), 20 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 1657a029d..d46e08a8d 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -237,6 +237,21 @@ fn main() { // Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly) preserves // the deleted host-side formula's semantics exactly. "update_grad_norm_emas_kernel.cu", + // SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness + // gauge update. Single-thread single-block — replaces the host-side + // arithmetic block in `update_iqn_readiness` per + // `feedback_no_cpu_compute_strict`. Reads host-passed `iqn_loss` + // scalar (already from `read_total_loss` mapped-pinned readback), + // updates the three coupled scalars (loss_initial anchor, loss_ema + // streaming smoothed loss, readiness gauge ∈ [0,1]) all backed by + // mapped-pinned device-mapped storage. Cold-start sentinel + // (`loss_initial < 1e-12` ⇒ bootstrap) preserves the deleted host- + // side formula's semantics exactly. Same adaptive-α formula + // (clamp(|err|/(|err|+0.01), 0.01, 0.30)) and improvement gauge + // ((initial - ema) / max(initial, 1e-6)) clamped to [0, 1]. + // Consumer (c51_loss_kernel CVaR α via `iqn_readiness_dev_ptr`) is + // unchanged — same dev_ptr, kernel writes the gauge through it. + "update_iqn_readiness_kernel.cu", // HEALTH_DIAG GPU port — Phase 2A landing (2026-04-28). Single-block // single-thread `health_diag_isv_mirror` kernel copies a curated set // of ISV signal-bus slots into the mapped-pinned `HealthDiagSnapshot` diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 1791fe0ac..c38350637 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -338,6 +338,16 @@ static FOLD_WARMUP_FACTOR_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), /// mapped-pinned EMA scalars consumed by `fold_warmup_factor_kernel`. static UPDATE_GRAD_NORM_EMAS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_grad_norm_emas_kernel.cubin")); +/// SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness gauge update +/// kernel. Single-thread, single-block — replaces the host-side arithmetic +/// block in `update_iqn_readiness` per `feedback_no_cpu_compute_strict`. +/// Takes the host-passed `iqn_loss` scalar (already from `read_total_loss` +/// mapped-pinned readback) and updates three coupled mapped-pinned scalars +/// (`iqn_loss_initial_pinned`, `iqn_loss_ema_pinned`, `iqn_readiness_pinned`) +/// in-place via their dev_ptrs. Cold-start sentinel `prev < 1e-12` preserves +/// the deleted host-side bootstrap branch. +static UPDATE_IQN_READINESS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_iqn_readiness_kernel.cubin")); + /// Plan 4 Task 1B-ii: per-group VSN MLP hidden dim. `Linear_1[g]` is /// `[VSN_HIDDEN_DIM, group_dim_g]`, `Linear_2[g]` is `[1, VSN_HIDDEN_DIM]`. /// 16 chosen as a tractable expansion that fits comfortably against the @@ -2861,12 +2871,43 @@ pub struct GpuDqnTrainer { pub(crate) eval_q_std_ema: [f32; 4], eval_ema_initialized: [bool; 4], /// Adaptive IQN lambda readiness: 0=uncertain (suppress gradient), 1=converged (full weight). + /// Cached host shadow for compatibility with non-host-pinned readers; the + /// authoritative value lives in `iqn_readiness_pinned` (see below). After + /// SP4 Layer C close-out C2 (2026-05-01) `update_iqn_readiness_kernel` + /// writes the pinned slot directly; this field is kept in sync via the + /// `iqn_readiness()` accessor reading through the pinned host_ptr. iqn_readiness: f32, /// IQN readiness — pinned device-mapped for CUDA graph. GPU reads via dev_ptr. + /// SP4 Layer C close-out C2 (2026-05-01) made this the authoritative + /// storage written by `update_iqn_readiness_kernel`; previously the host + /// `iqn_readiness` field was the source of truth and the pinned slot a + /// mirror per `feedback_no_cpu_compute_strict`. iqn_readiness_pinned: *mut f32, iqn_readiness_dev_ptr: u64, - iqn_loss_ema: f32, - iqn_loss_initial: f32, + /// Streaming smoothed IQN loss — mapped-pinned scalar. SP4 Layer C + /// close-out C2 (2026-05-01) replaced the host `iqn_loss_ema: f32` field + /// with this mapped-pinned slot updated GPU-side by + /// `update_iqn_readiness_kernel` per `feedback_no_cpu_compute_strict`. + /// Cold-start sentinel `prev < 1e-12 ⇒ assign loss directly` preserves + /// the deleted host-side formula's bootstrap branch exactly. FoldReset: + /// 0.0 (the new fold's first call recaptures the bootstrap anchor). + iqn_loss_ema_pinned: *mut f32, + /// Device-mapped view of `iqn_loss_ema_pinned`. Read+written in-place by + /// `update_iqn_readiness_kernel`; not consumed by any other GPU kernel + /// (the gauge produced from it lives in `iqn_readiness_pinned`). + iqn_loss_ema_dev_ptr: u64, + /// Fold-anchor for the IQN improvement gauge — mapped-pinned scalar. + /// SP4 Layer C close-out C2 (2026-05-01) replaced the host + /// `iqn_loss_initial: f32` field with this mapped-pinned slot updated + /// GPU-side by `update_iqn_readiness_kernel`. Captured once per fold on + /// the first call where `prev < 1e-12` and never overwritten afterward + /// (the kernel's bootstrap branch is the only writer of this field after + /// fold reset). FoldReset: 0.0 so the next call rebootstraps the anchor. + iqn_loss_initial_pinned: *mut f32, + /// Device-mapped view of `iqn_loss_initial_pinned`. Read+conditionally- + /// written in-place by `update_iqn_readiness_kernel`; not consumed by any + /// other GPU kernel. + iqn_loss_initial_dev_ptr: u64, /// Learning rate — pinned device-mapped host memory. /// CUDA graph captures the pointer (not the value). CPU writes new LR each epoch, @@ -3851,6 +3892,16 @@ pub struct GpuDqnTrainer { /// in-place via their mapped-pinned dev_ptrs. Loaded from /// `update_grad_norm_emas_kernel.cubin`. update_grad_norm_emas_kernel: CudaFunction, + /// SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness gauge + /// update kernel. Single-thread, single-block. Replaces the host-side + /// arithmetic block in `update_iqn_readiness` per + /// `feedback_no_cpu_compute_strict`. Takes the host-passed `iqn_loss` + /// scalar (already from `read_total_loss` mapped-pinned readback) and + /// updates the three coupled mapped-pinned scalars + /// (`iqn_loss_initial_pinned`, `iqn_loss_ema_pinned`, + /// `iqn_readiness_pinned`) in-place via their dev_ptrs. Loaded from + /// `update_iqn_readiness_kernel.cubin`. + update_iqn_readiness_kernel: CudaFunction, /// Mapped-pinned host pointer for the fast (α=0.1, ~10-step horizon) /// grad-norm EMA. Updated GPU-side by `update_grad_norm_emas_kernel` /// (chained on the producer's stream after `launch_grad_norm_finalize`) @@ -4686,6 +4737,12 @@ impl Drop for GpuDqnTrainer { if !self.iqn_readiness_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.iqn_readiness_pinned.cast()) }; } + if !self.iqn_loss_initial_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.iqn_loss_initial_pinned.cast()) }; + } + if !self.iqn_loss_ema_pinned.is_null() { + let _ = unsafe { cudarc::driver::result::free_host(self.iqn_loss_ema_pinned.cast()) }; + } if !self.td_error_scratch_pinned.is_null() { let _ = unsafe { cudarc::driver::result::free_host(self.td_error_scratch_pinned.cast()) }; } @@ -6763,23 +6820,71 @@ impl GpuDqnTrainer { /// Update IQN readiness from the current IQN loss. /// Readiness ramps 0→1 as loss decreases from initial value. /// Uses adaptive-rate EMA (same pattern as eval_v_range). + /// + /// SP4 Layer C close-out C2 (2026-05-01): the EMA arithmetic + + /// improvement-gauge formula now run GPU-side via + /// `update_iqn_readiness_kernel` per `feedback_no_cpu_compute_strict`. + /// The host signature is preserved (callers pass the loss scalar from + /// `gpu_iqn.read_total_loss()` mapped-pinned readback); the kernel is + /// launched on the trainer's stream and updates the three coupled + /// mapped-pinned scalars (`iqn_loss_initial_pinned`, `iqn_loss_ema_pinned`, + /// `iqn_readiness_pinned`) in-place via their dev_ptrs. The host + /// `iqn_readiness` shadow field is refreshed from the pinned slot after + /// launch so existing host-side accessors observe the same value (the + /// kernel's `__threadfence_system()` ensures PCIe-visibility before the + /// host_ptr deref). pub fn update_iqn_readiness(&mut self, iqn_loss: f32) { - if self.iqn_loss_initial < 1e-12 { - self.iqn_loss_initial = iqn_loss; - self.iqn_loss_ema = iqn_loss; - self.iqn_readiness = 0.0; - } else { - let err = (iqn_loss - self.iqn_loss_ema).abs(); - let alpha = (err / (err + 0.01)).clamp(0.01, 0.3); - self.iqn_loss_ema = (1.0 - alpha) * self.iqn_loss_ema + alpha * iqn_loss; - let improvement = (self.iqn_loss_initial - self.iqn_loss_ema) / self.iqn_loss_initial.max(1e-6); - self.iqn_readiness = improvement.clamp(0.0, 1.0); + if let Err(e) = self.launch_update_iqn_readiness(iqn_loss) { + tracing::warn!("update_iqn_readiness launch failed: {e}"); + return; } + // Refresh host shadow from the pinned slot — kernel's + // __threadfence_system() makes the write PCIe-visible to host_ptr + // dereferences on this stream. unsafe { - *self.iqn_readiness_pinned = self.iqn_readiness; + self.iqn_readiness = *self.iqn_readiness_pinned; } } + /// SP4 Layer C close-out C2 (2026-05-01): GPU-only IQN readiness gauge + /// update launcher. + /// + /// Replaces the host-side cold-start sentinel + adaptive-α EMA + improvement- + /// gauge arithmetic in `update_iqn_readiness` per + /// `feedback_no_cpu_compute_strict`. Single-thread, single-block — the + /// kernel reads/writes three mapped-pinned scalars in-place. Caller-compat: + /// the EMA buffers retain their mapped-pinned shape so c51_loss_kernel + /// continues to read `iqn_readiness_pinned` via dev_ptr (CVaR α) unchanged, + /// and the host-side accessors (`iqn_readiness()`, `iqn_loss_ema_value()`, + /// `iqn_loss_initial_value()`) keep observing the same memory after the + /// kernel's `__threadfence_system()`. + fn launch_update_iqn_readiness(&self, iqn_loss: f32) -> Result<(), MLError> { + let initial_dev = self.iqn_loss_initial_dev_ptr; + let ema_dev = self.iqn_loss_ema_dev_ptr; + let readiness_dev = self.iqn_readiness_dev_ptr; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + + // Safety: kernel signature `(float, float*, float*, float*)` matches + // the four args below; the three dev_ptrs are mapped-pinned and valid + // for the lifetime of `self`. + unsafe { + self.stream + .launch_builder(&self.update_iqn_readiness_kernel) + .arg(&iqn_loss) + .arg(&initial_dev) + .arg(&ema_dev) + .arg(&readiness_dev) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("update_iqn_readiness: {e}")))?; + } + Ok(()) + } + /// Raw pointer to the pinned host readback buffer [16 × f32]. /// Used by `FusedTrainingCtx` for async diversity loss readback at offset 9. pub fn readback_pinned_ptr(&self) -> *mut f32 { @@ -11411,6 +11516,37 @@ impl GpuDqnTrainer { cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, iqn_readiness_pinned.cast(), 0); dp }; + // SP4 Layer C close-out C2 (2026-05-01): mapped-pinned slots for the + // IQN readiness gauge state (`iqn_loss_initial`, `iqn_loss_ema`). + // Replace the prior host-resident `f32` fields per + // `feedback_no_cpu_compute_strict`. Updated GPU-side by + // `update_iqn_readiness_kernel` alongside `iqn_readiness_pinned`; + // accessed host-side via the kept `iqn_loss_initial_value()` / + // `iqn_loss_ema_value()` accessors that read the pinned slot directly. + let iqn_loss_initial_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 iqn_loss_initial alloc: {e}")))? + as *mut f32 + }; + unsafe { *iqn_loss_initial_pinned = 0.0; } + let iqn_loss_initial_dev_ptr = unsafe { + let mut dp = 0u64; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, iqn_loss_initial_pinned.cast(), 0); + dp + }; + let iqn_loss_ema_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 iqn_loss_ema alloc: {e}")))? + as *mut f32 + }; + unsafe { *iqn_loss_ema_pinned = 0.0; } + let iqn_loss_ema_dev_ptr = unsafe { + let mut dp = 0u64; + cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, iqn_loss_ema_pinned.cast(), 0); + dp + }; let total_actions = config.branch_0_size + config.branch_1_size + config.branch_2_size + config.branch_3_size; let q_out_buf = alloc_f32(&stream, b * total_actions, "q_out")?; let eval_td_snapshot = alloc_f32(&stream, b, "eval_td_snapshot")?; @@ -11901,6 +12037,19 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("update_grad_norm_emas_kernel load: {e}")))? }; + // SP4 Layer C close-out C2 (2026-05-01): load update_iqn_readiness + // kernel. Single-thread, single-block — replaces the host-side EMA + // arithmetic block in `update_iqn_readiness` per + // `feedback_no_cpu_compute_strict`. Updates the three coupled + // mapped-pinned scalars (`iqn_loss_initial_pinned`, + // `iqn_loss_ema_pinned`, `iqn_readiness_pinned`) in-place. + let update_iqn_readiness_kernel = { + let module = stream.context().load_cubin(UPDATE_IQN_READINESS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("update_iqn_readiness cubin load: {e}")))?; + module.load_function("update_iqn_readiness_kernel") + .map_err(|e| MLError::ModelError(format!("update_iqn_readiness_kernel load: {e}")))? + }; + // Plan 4 Task 3 (E.3): load iqn_quantile_ema kernel (cold-path, per-step). // 4-block kernel — one block per off-median fixed quantile in // FIXED_TAUS = [0.05, 0.25, 0.50, 0.75, 0.95]. Producer for @@ -14648,8 +14797,10 @@ impl GpuDqnTrainer { iqn_readiness: 0.0, iqn_readiness_pinned, iqn_readiness_dev_ptr, - iqn_loss_ema: 0.0, - iqn_loss_initial: 0.0, + iqn_loss_ema_pinned, + iqn_loss_ema_dev_ptr, + iqn_loss_initial_pinned, + iqn_loss_initial_dev_ptr, lr_pinned, lr_dev_ptr, aux_lr_pinned, @@ -14909,6 +15060,7 @@ impl GpuDqnTrainer { q_drift_rate_ema_kernel, fold_warmup_factor_kernel, update_grad_norm_emas_kernel, + update_iqn_readiness_kernel, grad_norm_fast_ema_pinned, grad_norm_fast_ema_dev_ptr, grad_norm_slow_ema_pinned, @@ -15940,15 +16092,42 @@ impl GpuDqnTrainer { /// mapped pinned slot so any in-flight kernel reads the reset value /// without an HtoD copy. `update_iqn_readiness` already handles the /// `iqn_loss_initial < 1e-12` case as the bootstrap branch. + /// + /// SP4 Layer C close-out C2 (2026-05-01): `iqn_loss_initial` and + /// `iqn_loss_ema` storage migrated to mapped-pinned slots updated by + /// `update_iqn_readiness_kernel`; the fold-boundary reset writes 0.0 + /// through their host_ptrs so the next kernel launch (from the new + /// fold's first `update_iqn_readiness` call) sees `prev_initial < 1e-12` + /// and re-enters the bootstrap branch. pub fn reset_iqn_readiness_state(&mut self) { - self.iqn_loss_initial = 0.0; - self.iqn_loss_ema = 0.0; self.iqn_readiness = 0.0; if !self.iqn_readiness_pinned.is_null() { // Pinned device-mapped: writing the host slot propagates // immediately to GPU via the mapping; no HtoD copy issued. unsafe { *self.iqn_readiness_pinned = 0.0; } } + if !self.iqn_loss_initial_pinned.is_null() { + unsafe { *self.iqn_loss_initial_pinned = 0.0; } + } + if !self.iqn_loss_ema_pinned.is_null() { + unsafe { *self.iqn_loss_ema_pinned = 0.0; } + } + } + + /// Read the streaming smoothed IQN loss from the mapped-pinned slot. + /// SP4 Layer C close-out C2 (2026-05-01) accessor — the pinned slot is + /// the authoritative storage; the kernel's `__threadfence_system()` + /// guarantees PCIe-visibility before this read sees a fresh value. + pub fn iqn_loss_ema_value(&self) -> f32 { + if self.iqn_loss_ema_pinned.is_null() { 0.0 } + else { unsafe { *self.iqn_loss_ema_pinned } } + } + + /// Read the IQN loss fold-anchor from the mapped-pinned slot. Same + /// rationale as `iqn_loss_ema_value()`. + pub fn iqn_loss_initial_value(&self) -> f32 { + if self.iqn_loss_initial_pinned.is_null() { 0.0 } + else { unsafe { *self.iqn_loss_initial_pinned } } } /// Set per_branch_q_gap_ema — HtoD for trajectory backtracking restore. diff --git a/crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu b/crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu new file mode 100644 index 000000000..741680362 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu @@ -0,0 +1,84 @@ +/* update_iqn_readiness — GPU-only IQN readiness gauge update. + * + * SP4 Layer C close-out C2 (2026-05-01). Replaces the host-side arithmetic + * block in `GpuDqnTrainer::update_iqn_readiness` + * (gpu_dqn_trainer.rs ~line 6766) per `feedback_no_cpu_compute_strict`. Mirrors + * `update_grad_norm_emas_kernel.cu` shape — single-thread, single-block, + * `__threadfence_system()` writeback to mapped-pinned scalars consumed by + * downstream GPU kernels (c51_loss CVaR α / IQN gradient gating). + * + * Reads: iqn_loss — host-passed scalar (current epoch's IQN total + * loss; `read_total_loss` already populated this + * from the IQN's mapped-pinned readback before + * this launch). + * loss_initial_dev[0] — previous fold-anchor (mapped-pinned scalar). + * loss_ema_dev[0] — previous streaming smoothed loss + * (mapped-pinned scalar). + * readiness_dev[0] — previous readiness gauge (mapped-pinned + * scalar; same slot c51_loss_kernel reads as + * CVaR α via dev_ptr). + * Writes: loss_initial_dev[0] — captured on first call (where prev < 1e-12), + * never overwritten after. + * loss_ema_dev[0] — adaptive-α EMA update. + * readiness_dev[0] — `(initial - ema) / max(initial, 1e-6)`, + * clamped to [0, 1]. + * + * Cold-start sentinel: if `loss_initial_dev[0] < 1e-12`, treat as bootstrap + * (first call after fold-boundary reset OR the very first call): assign + * initial=loss, ema=loss, readiness=0. Otherwise apply standard adaptive-α + * EMA recurrence + * err = |loss - ema_prev| + * alpha = clamp(err / (err + 0.01), 0.01, 0.30) + * ema = (1 - α) × ema_prev + α × loss + * readiness = clamp((initial - ema) / max(initial, 1e-6), 0, 1) + * Matches the deleted host-side formula bit-for-bit (same algebraic form + + * same clamp bounds; IEEE-754 fp32 commutativity preserved). + * + * Per `feedback_no_cpu_compute_strict.md`: any compute (EMA, reduction, mean) + * belongs on GPU regardless of frequency, regardless of whether it's hot-path + * or cold-path. The cold-path scalar arithmetic on three mapped-pinned + * scalars qualifies — pearl `pearl_cold_path_no_exception_to_gpu_drives.md`. + * Per `feedback_no_atomicadd.md`: a single thread reads four scalars and + * writes three — no atomics needed. + * + * Mapped-pinned scalars (`loss_initial_dev`, `loss_ema_dev`, `readiness_dev`) + * are 1-element device pointers backed by mapped-pinned host memory. + * `__threadfence_system()` ensures the writes are PCIe-visible to the + * mapped-pinned host_ptr accessors (`iqn_readiness()`, `iqn_loss_ema_value()`, + * `iqn_loss_initial_value()`) plus other-stream kernel reads via the same + * dev_ptrs. + */ + +extern "C" __global__ void update_iqn_readiness_kernel( + float iqn_loss, /* host-passed scalar */ + float* __restrict__ loss_initial_dev, /* [1] mapped-pinned anchor */ + float* __restrict__ loss_ema_dev, /* [1] mapped-pinned smoothed loss */ + float* __restrict__ readiness_dev /* [1] mapped-pinned gauge ∈ [0,1] */ +) { + if (blockIdx.x != 0 || threadIdx.x != 0) return; + + const float prev_initial = loss_initial_dev[0]; + + if (prev_initial < 1e-12f) { + /* Bootstrap branch — first call OR post-fold-reset. Capture anchor, + * seed EMA at the same value, gauge starts at 0 (readiness ramps as + * EMA drops below initial). */ + loss_initial_dev[0] = iqn_loss; + loss_ema_dev[0] = iqn_loss; + readiness_dev[0] = 0.0f; + } else { + const float prev_ema = loss_ema_dev[0]; + const float err = fabsf(iqn_loss - prev_ema); + const float alpha_raw = err / (err + 0.01f); + const float alpha = fminf(fmaxf(alpha_raw, 0.01f), 0.30f); + const float new_ema = (1.0f - alpha) * prev_ema + alpha * iqn_loss; + loss_ema_dev[0] = new_ema; + + const float denom = fmaxf(prev_initial, 1e-6f); + const float improve = (prev_initial - new_ema) / denom; + const float readiness = fminf(fmaxf(improve, 0.0f), 1.0f); + readiness_dev[0] = readiness; + } + + __threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr */ +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 714d82190..1ccb58791 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -91,17 +91,17 @@ impl StateResetRegistry { RegistryEntry { name: "isv_iqn_loss_initial", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.iqn_loss_initial — fold-anchor for IQN improvement gauge; producer is `update_iqn_readiness` (gpu_dqn_trainer.rs:5804); never reset before Fix 3 of 2026-04-28 fold-boundary audit", + description: "GpuDqnTrainer.iqn_loss_initial_pinned — fold-anchor for IQN improvement gauge stored in a mapped-pinned scalar; producer is `update_iqn_readiness_kernel` (SP4 Layer C close-out C2, 2026-05-01 — migrated from host-side `f32` field per `feedback_no_cpu_compute_strict`); cold-start sentinel `prev < 1e-12 ⇒ bootstrap` keeps the deleted host formula's semantics. Reset to 0.0 at fold boundary so the new fold's first kernel launch re-enters the bootstrap branch.", }, RegistryEntry { name: "isv_iqn_loss_ema", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.iqn_loss_ema — streaming smoothed IQN loss; producer is `update_iqn_readiness`; coupled to isv_iqn_loss_initial in the readiness fraction", + description: "GpuDqnTrainer.iqn_loss_ema_pinned — streaming smoothed IQN loss in a mapped-pinned scalar; producer is `update_iqn_readiness_kernel` (SP4 Layer C close-out C2, 2026-05-01 — migrated from host-side `f32` field per `feedback_no_cpu_compute_strict`); coupled to isv_iqn_loss_initial in the readiness fraction `(initial - ema) / max(initial, 1e-6)`. Reset to 0.0 at fold boundary alongside the anchor so the new fold's first kernel launch reseeds via bootstrap.", }, RegistryEntry { name: "isv_iqn_readiness", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.iqn_readiness scalar + pinned device-mapped slot (read by c51_loss_kernel as CVaR α and by IQN gradient-weight gating); reset to 0.0 at fold boundary so the new fold's first `update_iqn_readiness` recaptures the bootstrap anchor", + description: "GpuDqnTrainer.iqn_readiness_pinned (mapped-pinned device slot read by c51_loss_kernel as CVaR α and by IQN gradient-weight gating); produced GPU-side by `update_iqn_readiness_kernel` per training step (SP4 Layer C close-out C2, 2026-05-01) — replaces the prior host-side `(1-α) × prev + α × loss` arithmetic in `update_iqn_readiness` per `feedback_no_cpu_compute_strict`. Reset to 0.0 at fold boundary so the new fold's first `update_iqn_readiness` recaptures the bootstrap anchor.", }, // ───── Window-reset state ─────────────────────────────────── RegistryEntry { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 663b9926c..680b4717d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2967,3 +2967,47 @@ Files touched (Layer C C1 redesigned): No behavior change — pure architectural fix. The L40S smoke `smoke-test-tkkx6` (against `45da3eac6`) remains valid as the baseline; this commit doesn't need a new smoke since the only change is "where is the arithmetic computed", not "what is computed". Refs: SP4 Layer C C1 redesigned (was: retire). Original plan line ~2049 + Task C1 ~lines 2184-2221. follow-up: symbolic doc-anchors + dedicated q_dir_grad sub-buffer table + p99 helper extraction + +### Layer C Task C2 — IQN readiness gauge update migrated to GPU (2026-05-01) + +C1 redesigned closed out one site of the codebase-wide `feedback_no_cpu_compute_strict` sweep. C2 closes the next site: `GpuDqnTrainer::update_iqn_readiness` (gpu_dqn_trainer.rs ~line 6766) was running the cold-start sentinel + adaptive-α EMA + improvement-gauge formula host-side: + +``` +if self.iqn_loss_initial < 1e-12 { /* bootstrap */ } +else { + let err = (iqn_loss - self.iqn_loss_ema).abs(); + let alpha = (err / (err + 0.01)).clamp(0.01, 0.3); + self.iqn_loss_ema = (1.0 - alpha) * self.iqn_loss_ema + alpha * iqn_loss; + let improvement = (self.iqn_loss_initial - self.iqn_loss_ema) / self.iqn_loss_initial.max(1e-6); + self.iqn_readiness = improvement.clamp(0.0, 1.0); +} +``` + +Three coupled scalars on `GpuDqnTrainer`: `iqn_loss_initial` (fold-anchor), `iqn_loss_ema` (streaming smoothed loss), `iqn_readiness` (gauge ∈ [0,1] consumed by `c51_loss_kernel` as CVaR α via `iqn_readiness_dev_ptr`). Producer cadence: per training step from `fused_training.rs:1769` after `gpu_iqn.read_total_loss()` provides the loss via mapped-pinned readback. + +Migrated: +- New `update_iqn_readiness_kernel.cu` — single-thread, single-block. Takes the host-passed `iqn_loss` scalar (already a mapped-pinned readback from the IQN's `total_loss_pinned`) and updates three coupled mapped-pinned scalars (`iqn_loss_initial_pinned`, `iqn_loss_ema_pinned`, `iqn_readiness_pinned`) in-place via their dev_ptrs. Mirrors `update_grad_norm_emas_kernel` shape (cold-path scalar arithmetic on mapped-pinned slots). +- Storage migration: `iqn_loss_initial` and `iqn_loss_ema` migrated from host-resident `f32` fields to mapped-pinned device-mapped scalars. Constructor allocates two new `cuMemHostAlloc(DEVICEMAP)` slots; Drop frees them. The host shadow `iqn_readiness: f32` is kept as a cached value refreshed from the pinned slot after each kernel launch (preserves the `iqn_readiness()` accessor signature for non-host-pinned callers). +- New accessors `iqn_loss_ema_value()` and `iqn_loss_initial_value()` read through the host_ptrs — `__threadfence_system()` in the kernel guarantees PCIe-visibility before the host_ptr deref returns. +- `launch_update_iqn_readiness` Rust launcher chained on the trainer's stream — the kernel's three pinned slots are owned by the trainer, no cross-stream synchronization needed. +- `update_iqn_readiness` host-side body replaced with the launcher call + post-launch host shadow refresh (warn-and-continue on launch failure mirroring the C1 pattern). +- `reset_iqn_readiness_state` writes 0.0 through all three pinned slots so the next kernel launch re-enters the bootstrap branch (same fold-boundary semantic as before). + +Preserved (no behaviour change): +- Same adaptive-α formula `clamp(|err|/(|err|+0.01), 0.01, 0.30)` and improvement gauge `(initial - ema) / max(initial, 1e-6)` clamped to [0, 1]. +- Same cold-start sentinel `prev_initial < 1e-12 ⇒ bootstrap (initial=loss, ema=loss, readiness=0)`. +- `iqn_readiness_dev_ptr` consumer chain (`c51_loss_kernel` CVaR α + IQN gradient-weight gating + `gpu_experience_collector::set_iqn_readiness` quantile-blend selection) — unchanged. + +Verification: +- `cargo check -p ml --offline`: clean, same pre-existing warnings. +- `cargo test -p ml --lib --offline -- sp4 state_reset_registry`: 11 passed. +- `cargo test -p ml --test sp4_producer_unit_tests --offline --release -- --ignored`: 16/16 SP4 GPU tests pass on RTX 3050 Ti. + +Files touched (Layer C C2): +- `crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu` — new kernel. +- `crates/ml/build.rs` — 1 new kernel registration. +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 1 new CUBIN static + 1 new struct field + 2 new mapped-pinned slot pairs + 1 new constructor loader + 2 new pinned allocations + 2 new Drop entries + 1 new launcher (`launch_update_iqn_readiness`); host-side EMA arithmetic block replaced with launcher call + host shadow refresh; field doc-comments updated; `reset_iqn_readiness_state` rewires through pinned slots; 2 new accessors. +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — three FoldReset entries updated to reflect new mapped-pinned storage and producer relocation. +- `docs/dqn-wire-up-audit.md` — this entry. + +Refs: SP4 Layer C close-out C2 — second site of the `feedback_no_cpu_compute_strict` sweep. Audit grid (8 sites total) + group sequencing in `feedback_no_cpu_compute_strict.md` continuation.