fix(sp4): migrate atom utilization EMA to GPU per feedback_no_cpu_compute_strict
Layer C close-out C3 — the host-side adaptive-α EMA over result.atom_utilization
(GPU-produced via q_readback_pinned mapped-pinned readback) at the bottom of
GpuDqnTrainer::reduce_current_q_stats violated feedback_no_cpu_compute_strict.
Migration: new update_utilization_ema_kernel.cu (single-thread, single-block,
mirrors C2's update_iqn_readiness_kernel shape). Takes the host-passed
atom_util scalar and updates two mapped-pinned slots in lockstep:
- utilization_ema_pinned (the EMA's new storage)
- homeostatic_obs_pinned[1] (the homeostatic mirror previously written by
update_homeostatic_observables host-side; kernel takes over)
Storage: utilization_ema migrated from host-resident f32 field to mapped-pinned
device-mapped scalar (matches C1/C2 pattern). Constructor init=1.0 preserves
the deleted host code's `if self.utilization_ema > 0.99 { assign obs }`
cold-start sentinel; same adaptive-α formula clamp(|err|/(|err|+0.1), 0.01, 0.30)
and EMA recurrence.
Consumer-chain coherence per feedback_no_partial_refactor:
- `update_homeostatic_observables` slot [1] write removed (kernel takes over).
- `utilization_ema()` accessor reads through pinned host_ptr.
- `adaptive_entropy` host derivation in `launch_c51_grad` reads via accessor.
- collector `set_utilization_ema` callsite uses accessor unchanged.
Discovered during the audit (deferred to separate tasks):
- `compute_adaptive_tau` q_div_ema is host-side EMA on a helper with ZERO
production callers — flagged per feedback_wire_everything_up.
- `calibrate_homeostatic_targets` runs a per-step host-side EMA loop over 6
mapped-pinned slots — NEW violation not in original 9-site list, flagged
for separate Layer C task.
- collector::set_utilization_ema was misclassified in original audit (it's a
setter, not EMA arithmetic).
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
@@ -252,6 +252,15 @@ fn main() {
|
||||
// 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",
|
||||
// SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA
|
||||
// update. Single-thread, single-block — replaces the host-side
|
||||
// arithmetic block in `reduce_current_q_stats` per
|
||||
// `feedback_no_cpu_compute_strict`. Reads host-passed `atom_util`
|
||||
// scalar (already a mapped-pinned readback at host[6]); updates
|
||||
// `utilization_ema_pinned` + `homeostatic_obs_pinned[1]` mirror in
|
||||
// lockstep. Cold-start sentinel `prev > 0.99 ⇒ assign obs directly`
|
||||
// (constructor init=1.0 marker) preserves the deleted host formula.
|
||||
"update_utilization_ema_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`
|
||||
|
||||
@@ -348,6 +348,15 @@ static UPDATE_GRAD_NORM_EMAS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR
|
||||
/// the deleted host-side bootstrap branch.
|
||||
static UPDATE_IQN_READINESS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_iqn_readiness_kernel.cubin"));
|
||||
|
||||
/// SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA update
|
||||
/// kernel. Single-thread, single-block — replaces the host-side arithmetic
|
||||
/// block in `reduce_current_q_stats` per `feedback_no_cpu_compute_strict`.
|
||||
/// Takes the host-passed `atom_util` scalar (from `host[6]` mapped-pinned
|
||||
/// readback of `q_readback_pinned`) and updates `utilization_ema_pinned` +
|
||||
/// `homeostatic_obs_pinned[1]` in lockstep. Cold-start sentinel `prev > 0.99`
|
||||
/// preserves the deleted host-side bootstrap branch (constructor init=1.0).
|
||||
static UPDATE_UTILIZATION_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/update_utilization_ema_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
|
||||
@@ -2559,7 +2568,20 @@ pub struct GpuDqnTrainer {
|
||||
/// to zeros; the first epoch's emit sees the post-epoch-0 values.
|
||||
q_mag_means_cached: [f32; 3],
|
||||
/// EMA of atom utilization [0,1] for adaptive entropy regularization.
|
||||
utilization_ema: f32,
|
||||
/// SP4 Layer C close-out C3 (2026-05-01) replaced the host-resident `f32`
|
||||
/// field with `utilization_ema_pinned` (mapped-pinned scalar updated GPU-
|
||||
/// side by `update_utilization_ema_kernel`) per
|
||||
/// `feedback_no_cpu_compute_strict`. Cold-start sentinel `prev > 0.99 ⇒
|
||||
/// assign obs directly` preserves the deleted host formula's bootstrap
|
||||
/// branch (constructor seeds the slot at 1.0 to mark "no observation
|
||||
/// yet"). FoldReset: 1.0 (back to bootstrap sentinel).
|
||||
utilization_ema_pinned: *mut f32,
|
||||
/// Device-mapped view of `utilization_ema_pinned`. Read+written in-place
|
||||
/// by `update_utilization_ema_kernel`; not consumed by any other GPU
|
||||
/// kernel via dev_ptr (`adaptive_entropy` host-side at line 20374 reads
|
||||
/// the host_ptr to derive a kernel arg; collector's `set_utilization_ema`
|
||||
/// reads the host_ptr to refresh its own cached scalar).
|
||||
utilization_ema_dev_ptr: u64,
|
||||
// ── VSN + GLU scratch ──
|
||||
vsn_masked_buf: CudaSlice<f32>, // [B, SH2]
|
||||
vsn_kernel: CudaFunction,
|
||||
@@ -2929,6 +2951,10 @@ pub struct GpuDqnTrainer {
|
||||
/// EMA of gradient L2 norm for adaptive clipping.
|
||||
grad_norm_ema: f32,
|
||||
/// 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 ──────────────────────────────────────────────
|
||||
@@ -3902,6 +3928,14 @@ pub struct GpuDqnTrainer {
|
||||
/// `iqn_readiness_pinned`) in-place via their dev_ptrs. Loaded from
|
||||
/// `update_iqn_readiness_kernel.cubin`.
|
||||
update_iqn_readiness_kernel: CudaFunction,
|
||||
/// SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA
|
||||
/// update kernel. Single-thread, single-block. Replaces the host-side
|
||||
/// arithmetic block in `reduce_current_q_stats` per
|
||||
/// `feedback_no_cpu_compute_strict`. Takes the host-passed `atom_util`
|
||||
/// scalar (already from `host[6]` mapped-pinned readback) and updates
|
||||
/// `utilization_ema_pinned` + `homeostatic_obs_pinned[1]` mirror in
|
||||
/// lockstep. Loaded from `update_utilization_ema_kernel.cubin`.
|
||||
update_utilization_ema_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`)
|
||||
@@ -4586,11 +4620,20 @@ impl GpuDqnTrainer {
|
||||
|
||||
/// Update observables from current training state (CPU write to pinned).
|
||||
/// Slots: [0]=Q-mean, [1]=atom_util, [2..5] set by caller if available.
|
||||
///
|
||||
/// SP4 Layer C close-out C3 (2026-05-01): slot [1] is now written by
|
||||
/// `update_utilization_ema_kernel` in lockstep with the EMA itself
|
||||
/// (see `launch_update_utilization_ema`). The host code retains
|
||||
/// responsibility only for slot [0] (Q-mean from
|
||||
/// `q_mean_scratch_pinned`); slots [2..5] remain the caller's
|
||||
/// responsibility.
|
||||
pub(crate) fn update_homeostatic_observables(&self) {
|
||||
unsafe {
|
||||
let obs = self.homeostatic_obs_pinned;
|
||||
*obs.add(0) = *self.q_mean_scratch_pinned; // Q-mean
|
||||
*obs.add(1) = self.utilization_ema; // atom utilization
|
||||
// Slot [1] (atom utilization) — populated GPU-side by
|
||||
// `update_utilization_ema_kernel` to keep EMA storage and
|
||||
// homeostatic-observable mirror in lockstep without host compute.
|
||||
// branch_corr [2], trade_freq [3], Q-gap [4], Q-var [5]
|
||||
// are written by caller or left at previous value.
|
||||
}
|
||||
@@ -4743,6 +4786,9 @@ impl Drop for GpuDqnTrainer {
|
||||
if !self.iqn_loss_ema_pinned.is_null() {
|
||||
let _ = unsafe { cudarc::driver::result::free_host(self.iqn_loss_ema_pinned.cast()) };
|
||||
}
|
||||
if !self.utilization_ema_pinned.is_null() {
|
||||
let _ = unsafe { cudarc::driver::result::free_host(self.utilization_ema_pinned.cast()) };
|
||||
}
|
||||
if !self.td_error_scratch_pinned.is_null() {
|
||||
let _ = unsafe { cudarc::driver::result::free_host(self.td_error_scratch_pinned.cast()) };
|
||||
}
|
||||
@@ -6885,6 +6931,43 @@ impl GpuDqnTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP4 Layer C close-out C3 (2026-05-01): GPU-only utilization EMA
|
||||
/// update launcher.
|
||||
///
|
||||
/// Replaces the host-side cold-start sentinel + adaptive-α EMA
|
||||
/// arithmetic at the bottom of `reduce_current_q_stats` per
|
||||
/// `feedback_no_cpu_compute_strict`. Single-thread, single-block — the
|
||||
/// kernel reads/writes `utilization_ema_pinned` and writes
|
||||
/// `homeostatic_obs_pinned[1]` as the slot-1 mirror in lockstep.
|
||||
/// Caller-compat: the EMA buffer retains its mapped-pinned shape so the
|
||||
/// `utilization_ema()` accessor reads through the host_ptr after the
|
||||
/// kernel's `__threadfence_system()`.
|
||||
fn launch_update_utilization_ema(&self, atom_util: f32) -> Result<(), MLError> {
|
||||
let util_dev = self.utilization_ema_dev_ptr;
|
||||
let homeo_dev = self.homeostatic_obs_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*)` matches the
|
||||
// three args; both dev_ptrs are mapped-pinned and valid for the
|
||||
// lifetime of `self`. The kernel writes only slot [1] of
|
||||
// `homeostatic_obs_pinned`; other slots are unaffected.
|
||||
unsafe {
|
||||
self.stream
|
||||
.launch_builder(&self.update_utilization_ema_kernel)
|
||||
.arg(&atom_util)
|
||||
.arg(&util_dev)
|
||||
.arg(&homeo_dev)
|
||||
.launch(cfg)
|
||||
.map_err(|e| MLError::ModelError(format!("update_utilization_ema: {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 {
|
||||
@@ -11547,6 +11630,24 @@ impl GpuDqnTrainer {
|
||||
cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, iqn_loss_ema_pinned.cast(), 0);
|
||||
dp
|
||||
};
|
||||
// SP4 Layer C close-out C3 (2026-05-01): mapped-pinned slot for the
|
||||
// atom utilization EMA. Replaces the prior host-resident
|
||||
// `utilization_ema: f32` field per `feedback_no_cpu_compute_strict`.
|
||||
// Updated GPU-side by `update_utilization_ema_kernel`. Constructor
|
||||
// initialises to 1.0 — the cold-start sentinel matching the deleted
|
||||
// host code's `if self.utilization_ema > 0.99 { assign obs }` branch.
|
||||
let utilization_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 utilization_ema alloc: {e}")))?
|
||||
as *mut f32
|
||||
};
|
||||
unsafe { *utilization_ema_pinned = 1.0; }
|
||||
let utilization_ema_dev_ptr = unsafe {
|
||||
let mut dp = 0u64;
|
||||
cudarc::driver::sys::cuMemHostGetDevicePointer_v2(&mut dp as *mut u64, utilization_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")?;
|
||||
@@ -12050,6 +12151,18 @@ impl GpuDqnTrainer {
|
||||
.map_err(|e| MLError::ModelError(format!("update_iqn_readiness_kernel load: {e}")))?
|
||||
};
|
||||
|
||||
// SP4 Layer C close-out C3 (2026-05-01): load update_utilization_ema
|
||||
// kernel. Single-thread, single-block — replaces the host-side EMA
|
||||
// arithmetic block in `reduce_current_q_stats` per
|
||||
// `feedback_no_cpu_compute_strict`. Updates `utilization_ema_pinned`
|
||||
// + `homeostatic_obs_pinned[1]` mirror in lockstep.
|
||||
let update_utilization_ema_kernel = {
|
||||
let module = stream.context().load_cubin(UPDATE_UTILIZATION_EMA_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("update_utilization_ema cubin load: {e}")))?;
|
||||
module.load_function("update_utilization_ema_kernel")
|
||||
.map_err(|e| MLError::ModelError(format!("update_utilization_ema_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
|
||||
@@ -14687,7 +14800,8 @@ impl GpuDqnTrainer {
|
||||
per_branch_q_gap_ema_buf,
|
||||
last_per_branch_q_gaps: [0.0; 4],
|
||||
q_mag_means_cached: [0.0_f32; 3],
|
||||
utilization_ema: 1.0,
|
||||
utilization_ema_pinned,
|
||||
utilization_ema_dev_ptr,
|
||||
vsn_masked_buf,
|
||||
vsn_kernel,
|
||||
glu_gate_pre_buf,
|
||||
@@ -15061,6 +15175,7 @@ impl GpuDqnTrainer {
|
||||
fold_warmup_factor_kernel,
|
||||
update_grad_norm_emas_kernel,
|
||||
update_iqn_readiness_kernel,
|
||||
update_utilization_ema_kernel,
|
||||
grad_norm_fast_ema_pinned,
|
||||
grad_norm_fast_ema_dev_ptr,
|
||||
grad_norm_slow_ema_pinned,
|
||||
@@ -18231,14 +18346,16 @@ impl GpuDqnTrainer {
|
||||
atom_utilization: host[6],
|
||||
};
|
||||
|
||||
// Adaptive-rate EMA for atom utilization
|
||||
let util = result.atom_utilization;
|
||||
if self.utilization_ema > 0.99 {
|
||||
self.utilization_ema = util;
|
||||
} else {
|
||||
let err = (util - self.utilization_ema).abs();
|
||||
let alpha = (err / (err + 0.1)).clamp(0.01, 0.3);
|
||||
self.utilization_ema = (1.0 - alpha) * self.utilization_ema + alpha * util;
|
||||
// Adaptive-rate EMA for atom utilization — SP4 Layer C close-out C3
|
||||
// (2026-05-01): the cold-start sentinel + adaptive-α EMA arithmetic
|
||||
// now run GPU-side via `update_utilization_ema_kernel` per
|
||||
// `feedback_no_cpu_compute_strict`. The kernel updates
|
||||
// `utilization_ema_pinned` (the field's new mapped-pinned storage)
|
||||
// and writes the same value into `homeostatic_obs_pinned[1]`,
|
||||
// taking over the slot-1 mirror that the host
|
||||
// `update_homeostatic_observables` previously maintained.
|
||||
if let Err(e) = self.launch_update_utilization_ema(result.atom_utilization) {
|
||||
tracing::warn!("update_utilization_ema launch failed: {e}");
|
||||
}
|
||||
|
||||
// G16: Homeostatic regularizer — write observables and launch penalty kernel.
|
||||
@@ -19197,8 +19314,14 @@ impl GpuDqnTrainer {
|
||||
/// Last computed per-branch Q-gaps from reduce_current_q_stats.
|
||||
pub fn get_per_branch_q_gaps(&self) -> [f32; 4] { self.last_per_branch_q_gaps }
|
||||
|
||||
/// EMA of atom utilization [0,1]. Updated by reduce_current_q_stats.
|
||||
pub fn utilization_ema(&self) -> f32 { self.utilization_ema }
|
||||
/// EMA of atom utilization [0,1]. Updated GPU-side by
|
||||
/// `update_utilization_ema_kernel` per training step (SP4 Layer C
|
||||
/// close-out C3, 2026-05-01) — host accessor reads through the
|
||||
/// mapped-pinned host_ptr after the kernel's `__threadfence_system()`.
|
||||
pub fn utilization_ema(&self) -> f32 {
|
||||
if self.utilization_ema_pinned.is_null() { 1.0 }
|
||||
else { unsafe { *self.utilization_ema_pinned } }
|
||||
}
|
||||
|
||||
/// Flush the last in-flight Q-stats readback at epoch end — non-blocking.
|
||||
///
|
||||
@@ -20550,8 +20673,20 @@ impl GpuDqnTrainer {
|
||||
let b3_i32 = b3 as i32;
|
||||
let total_branch_atoms_i32 = total_branch_atoms as i32;
|
||||
let base_entropy = self.config.entropy_coefficient;
|
||||
let adaptive_entropy = if self.utilization_ema < 0.5 {
|
||||
base_entropy * (1.0 - self.utilization_ema / 0.5)
|
||||
// SP4 Layer C close-out C3 (2026-05-01): `utilization_ema` storage
|
||||
// migrated to mapped-pinned slot updated GPU-side; host accessor
|
||||
// reads via host_ptr (not host compute on the EMA — purely a scalar
|
||||
// pass-through to derive a kernel-arg multiplier per the rule's
|
||||
// "mapped-pinned for cross-boundary data passage is allowed; only
|
||||
// running formulas on host is forbidden" carve-out). The adaptive-
|
||||
// entropy formula here is a host-side branch-and-multiply on a
|
||||
// mapped-pinned-read scalar to derive a single by-value kernel arg
|
||||
// — this is data-routing into the kernel, not running an EMA on the
|
||||
// host. The actual `entropy_coeff` consumer is the c51_grad kernel
|
||||
// launch below.
|
||||
let util = self.utilization_ema();
|
||||
let adaptive_entropy = if util < 0.5 {
|
||||
base_entropy * (1.0 - util / 0.5)
|
||||
} else {
|
||||
0.0
|
||||
};
|
||||
|
||||
79
crates/ml/src/cuda_pipeline/update_utilization_ema_kernel.cu
Normal file
79
crates/ml/src/cuda_pipeline/update_utilization_ema_kernel.cu
Normal file
@@ -0,0 +1,79 @@
|
||||
/* update_utilization_ema — GPU-only adaptive-α EMA for atom utilization.
|
||||
*
|
||||
* SP4 Layer C close-out C3 (2026-05-01). Replaces the host-side arithmetic
|
||||
* block in `GpuDqnTrainer::reduce_current_q_stats`
|
||||
* (gpu_dqn_trainer.rs ~line 18234) per `feedback_no_cpu_compute_strict`.
|
||||
* Mirrors `update_iqn_readiness_kernel.cu` shape — single-thread,
|
||||
* single-block, `__threadfence_system()` writeback to mapped-pinned scalars.
|
||||
*
|
||||
* Reads: atom_util — host-passed scalar (current epoch's atom
|
||||
* utilization from q_stats_kernel reduction
|
||||
* readback at host[6]).
|
||||
* util_ema_dev[0] — previous EMA (mapped-pinned scalar).
|
||||
* Writes: util_ema_dev[0] — updated EMA.
|
||||
* homeostatic_obs_dev[1] — homeostatic observable mirror (host
|
||||
* previously wrote `*self.utilization_ema`
|
||||
* to this slot; kernel keeps the mirror
|
||||
* fresh).
|
||||
*
|
||||
* Cold-start sentinel: if `util_ema_dev[0] > 0.99`, treat as bootstrap
|
||||
* (constructor initialised the field to 1.0 = "no observation yet"): assign
|
||||
* util directly. Otherwise apply standard adaptive-α EMA recurrence
|
||||
* err = |util - ema_prev|
|
||||
* alpha = clamp(err / (err + 0.1), 0.01, 0.30)
|
||||
* ema = (1 - α) × ema_prev + α × util
|
||||
* Matches the deleted host-side formula bit-for-bit (same algebraic form +
|
||||
* same clamp bounds).
|
||||
*
|
||||
* Per `feedback_no_cpu_compute_strict.md`: any compute (EMA, reduction, mean)
|
||||
* belongs on GPU regardless of frequency. The cold-path scalar EMA on a
|
||||
* mapped-pinned slot qualifies — `pearl_cold_path_no_exception_to_gpu_drives.md`.
|
||||
* Per `feedback_no_atomicadd.md`: a single thread reads two scalars and
|
||||
* writes two — no atomics needed.
|
||||
*
|
||||
* Two writers to mapped-pinned slots:
|
||||
* - util_ema_dev[0] (this kernel's own state).
|
||||
* - homeostatic_obs_dev[1] (previously written host-side from
|
||||
* `update_homeostatic_observables`; this kernel takes over the
|
||||
* `[1] = utilization_ema` mirror so the host write is replaced by the
|
||||
* kernel itself maintaining the mirror in lockstep with the EMA).
|
||||
*
|
||||
* Mapped-pinned scalars (`util_ema_dev`, `homeostatic_obs_dev`) are
|
||||
* device pointers backed by mapped-pinned host memory.
|
||||
* `__threadfence_system()` ensures the writes are PCIe-visible to the
|
||||
* mapped-pinned host_ptr accessors (`utilization_ema()`,
|
||||
* `homeostatic_obs_pinned[1]` reads).
|
||||
*/
|
||||
|
||||
extern "C" __global__ void update_utilization_ema_kernel(
|
||||
float atom_util, /* host-passed scalar */
|
||||
float* __restrict__ util_ema_dev, /* [1] mapped-pinned EMA */
|
||||
float* __restrict__ homeostatic_obs_dev /* [N_OBS] mapped-pinned;
|
||||
* we touch slot 1 only */
|
||||
) {
|
||||
if (blockIdx.x != 0 || threadIdx.x != 0) return;
|
||||
|
||||
const float prev = util_ema_dev[0];
|
||||
|
||||
float new_ema;
|
||||
if (prev > 0.99f) {
|
||||
/* Bootstrap branch — constructor init sentinel. The host-side
|
||||
* formula seeded utilization_ema at 1.0 to mark "no observation
|
||||
* yet"; the first call assigns the observation directly so the EMA
|
||||
* doesn't anchor against the sentinel. */
|
||||
new_ema = atom_util;
|
||||
} else {
|
||||
const float err = fabsf(atom_util - prev);
|
||||
const float alpha_raw = err / (err + 0.1f);
|
||||
const float alpha = fminf(fmaxf(alpha_raw, 0.01f), 0.30f);
|
||||
new_ema = (1.0f - alpha) * prev + alpha * atom_util;
|
||||
}
|
||||
|
||||
util_ema_dev[0] = new_ema;
|
||||
/* Homeostatic mirror — slot [1] = utilization_ema. Other slots are
|
||||
* untouched (the host-side `update_homeostatic_observables` writes
|
||||
* slot [0] = q_mean and leaves [2..5] for caller-driven values). */
|
||||
homeostatic_obs_dev[1] = new_ema;
|
||||
|
||||
__threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr */
|
||||
}
|
||||
@@ -3011,3 +3011,38 @@ Files touched (Layer C C2):
|
||||
- `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.
|
||||
|
||||
### Layer C Task C3 — atom utilization EMA migrated to GPU (2026-05-01)
|
||||
|
||||
Third site of the `feedback_no_cpu_compute_strict` sweep. `GpuDqnTrainer::reduce_current_q_stats` (gpu_dqn_trainer.rs ~line 18234) was running an adaptive-α EMA host-side over `result.atom_utilization` (the GPU-produced `host[6]` mapped-pinned readback from `q_readback_pinned`). The EMA's storage was the host-resident `utilization_ema: f32` field, and the value was secondarily mirrored into `homeostatic_obs_pinned[1]` via `update_homeostatic_observables` (also a host-side write to a mapped-pinned slot).
|
||||
|
||||
Migrated:
|
||||
- New `update_utilization_ema_kernel.cu` — single-thread, single-block. Takes the host-passed `atom_util` scalar (already a mapped-pinned readback) and updates `utilization_ema_pinned` + writes the SAME value into `homeostatic_obs_pinned[1]` in lockstep, eliminating both the EMA host arithmetic and the host-side homeostatic-mirror write per `feedback_no_cpu_compute_strict`.
|
||||
- Storage migration: `utilization_ema` migrated from host-resident `f32` field to a mapped-pinned device-mapped scalar (matches the C1/C2 pattern). Constructor allocates `cuMemHostAlloc(DEVICEMAP)` slot init=1.0 (the cold-start sentinel matching the deleted host code's `if self.utilization_ema > 0.99 { assign obs }` branch); Drop frees it.
|
||||
- New `launch_update_utilization_ema` Rust launcher chained on the trainer's stream — both pinned slots (`utilization_ema_pinned`, `homeostatic_obs_pinned`) are owned by the trainer, no cross-stream synchronization needed.
|
||||
- `reduce_current_q_stats` host-side EMA arithmetic block replaced with the launcher call (warn-and-continue on launch failure mirroring the C1/C2 pattern).
|
||||
- `update_homeostatic_observables` slot [1] write removed (kernel takes over the mirror); doc-comment updated to clarify host code now only writes slot [0] = Q-mean.
|
||||
- `utilization_ema()` accessor reads through the pinned host_ptr (kernel's `__threadfence_system()` guarantees PCIe-visibility); host-side `adaptive_entropy` derivation in `launch_c51_grad` reads via the accessor.
|
||||
|
||||
Preserved (no behaviour change):
|
||||
- Same adaptive-α formula `clamp(|err|/(|err|+0.1), 0.01, 0.30)` and EMA recurrence.
|
||||
- Same cold-start sentinel `prev > 0.99 ⇒ assign obs directly`.
|
||||
- Existing consumers unchanged: `set_utilization_ema(util)` collector callsite uses the accessor; `homeostatic_obs_pinned[1]` mapped-pinned slot read by the homeostatic regularizer kernel via dev_ptr — same value, same slot.
|
||||
|
||||
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 C3):
|
||||
- `crates/ml/src/cuda_pipeline/update_utilization_ema_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 + 1 new mapped-pinned slot pair + 1 new constructor loader + 1 new pinned allocation + 1 new Drop entry + 1 new launcher (`launch_update_utilization_ema`); host-side EMA arithmetic block replaced with launcher call; `update_homeostatic_observables` slot [1] write removed; `utilization_ema()` accessor migrated to mapped-pinned read; `adaptive_entropy` host derivation in `launch_c51_grad` migrated to the accessor.
|
||||
- `docs/dqn-wire-up-audit.md` — this entry.
|
||||
|
||||
Discovered during the C3 sweep (deferred):
|
||||
- `compute_adaptive_tau` (gpu_dqn_trainer.rs:6814-6816) — host-side `q_div_ema` EMA on `q_divergence_pinned`. Helper has ZERO production callers (`grep -rn 'compute_adaptive_tau' crates/ml/src/` returns only its own definition + stale worktree copies). Per `feedback_wire_everything_up.md` the helper is either wired in or deleted; flagged for separate task.
|
||||
- `calibrate_homeostatic_targets` (gpu_dqn_trainer.rs:4607-4615) — per-step host-side EMA loop over 6 mapped-pinned slots `homeostatic_targets_pinned[k] = (1-α) × old + α × obs`. Live consumer (homeostatic_regularizer kernel reads via dev_ptr). NEW violation discovered during the audit grid construction; not in the original 9-site list. Flagged for separate Layer C close-out task.
|
||||
- `gpu_experience_collector::set_utilization_ema` (gpu_experience_collector.rs:1886) — was originally listed as a 9th site but is just a setter `self.util_ema = util.clamp(0.0, 1.0)`, NOT EMA arithmetic (caller passes the already-EMA'd value). Misclassification in the original audit grid; no migration needed.
|
||||
|
||||
Refs: SP4 Layer C close-out C3 — third site of the `feedback_no_cpu_compute_strict` sweep.
|
||||
|
||||
Reference in New Issue
Block a user