fix(sp4): migrate IQN readiness gauge update to GPU per feedback_no_cpu_compute_strict
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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`
|
||||
|
||||
@@ -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::<f32>(), 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::<f32>(), 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.
|
||||
|
||||
84
crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu
Normal file
84
crates/ml/src/cuda_pipeline/update_iqn_readiness_kernel.cu
Normal file
@@ -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 */
|
||||
}
|
||||
@@ -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 {
|
||||
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user