diff --git a/crates/ml/build.rs b/crates/ml/build.rs index d46e08a8d..fd0ff1299 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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` diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index c38350637..c390dc506 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, // [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::(), 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 }; diff --git a/crates/ml/src/cuda_pipeline/update_utilization_ema_kernel.cu b/crates/ml/src/cuda_pipeline/update_utilization_ema_kernel.cu new file mode 100644 index 000000000..bcb7e2c72 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/update_utilization_ema_kernel.cu @@ -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 */ +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 680b4717d..ecf45c418 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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.