From fc8dbb0a855c80aaffa9a131f771410d88e00c4f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 29 Apr 2026 18:29:20 +0200 Subject: [PATCH] feat(dqn): q_drift_rate_ema kernel + Rust wrapper (Plan C A.2 scaffolding) Single-block single-thread ISV producer mirrors h_s2_rms_ema / moe_lambda_eff pattern. Computes per-epoch drift_rate = |q_mean(t) - q_mean(t-1)| / max(|q_mean(t-1)|, ISV[Q_ABS_REF] + ISV[Q_DIR_ABS_REF], 1e-6) clipped to [0, 4] and writes to ISV[Q_DRIFT_RATE_INDEX=129]. Includes the full ISV-contract shift required for the kernel to load: - ISV_TOTAL_DIM 129 -> 130 - Q_DRIFT_RATE_INDEX = 129 (tail-appended after MOE_LAMBDA_EFF=128) - Cold-start ISV[129] = 0.0 in constructor (no-op dampening factor) - layout_fingerprint_seed entry Q_DRIFT_RATE=129 + ISV_TOTAL_DIM=130 - Cubin static, kernel field, kernel load, struct assignment Per feedback_no_partial_refactor.md the ISV slot + dim + fingerprint all migrate together (the wrapper references Q_DRIFT_RATE_INDEX so they cannot be split). Tau consumer + state reset registry + per-epoch producer launch land in the next commit. Audit doc dqn-wire-up-audit.md updated with the kernel + ISV slot description per Invariant 7. No callers in this commit; layout fingerprint shifts so existing checkpoints will fail-fast at load per feedback_no_legacy_aliases.md (expected for a real architecture change). Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 10 ++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 129 +++++++++++++++++- .../cuda_pipeline/q_drift_rate_ema_kernel.cu | 87 ++++++++++++ docs/dqn-wire-up-audit.md | 16 +++ 4 files changed, 240 insertions(+), 2 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index d61c58bce..2ea764085 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -201,6 +201,16 @@ fn main() { // Consumer (moe_load_balance_loss in moe_kernels.cu) reads slot 128 // at runtime — no DtoH per feedback_isv_for_adaptive_bounds.md. "moe_lambda_eff_kernel.cu", + // Plan C Phase 2 follow-up A.2 (2026-04-29): adaptive Polyak-tau + // dampening signal. Single-block single-thread cold-path producer + // mirroring moe_lambda_eff_kernel / kelly_cap_update precedent. Reads + // ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] and the host- + // passed q_mean_curr/prev scalars; writes + // ISV[Q_DRIFT_RATE_INDEX=129] = clip(|q_curr-q_prev|/denom, 0, 4). + // Consumer (tau_update_kernel.cu) multiplies tau_eff by + // 1/(1+ISV[129]) so tau ∈ [tau_base/5, tau_base] — monotone + // dampening under drift; healthy runs unaffected. + "q_drift_rate_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 d63752d6c..89582d6da 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -183,6 +183,17 @@ pub(crate) static AUX_HEADS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR" #[allow(dead_code)] pub(crate) static AUX_HEADS_LOSS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/aux_heads_loss_ema_kernel.cubin")); +/// Plan C Phase 2 follow-up A.2 (2026-04-29): q-drift rate ISV producer. +/// Single-thread single-block cold-path kernel mirroring +/// `moe_lambda_eff_kernel.cu` / `kelly_cap_update_kernel.cu`. Reads +/// ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] plus host-passed +/// `q_mean_curr` / `q_mean_prev` scalars; writes +/// ISV[Q_DRIFT_RATE_INDEX=129] = `clip(|Δq|/denom, 0, 4)`. Consumer: +/// `tau_update_kernel.cu` multiplies the cosine-scheduled `tau_base` by +/// `1/(1+ISV[129])` so the effective Polyak rate dampens monotonically as +/// q-drift rises. +static Q_DRIFT_RATE_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/q_drift_rate_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 @@ -441,7 +452,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -const ISV_TOTAL_DIM: usize = 129; +const ISV_TOTAL_DIM: usize = 130; /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -836,6 +847,36 @@ pub const MOE_GATE_ENTROPY_EMA_INDEX: usize = 126; /// flow through ISV; no DtoH on hot path. pub const MOE_LAMBDA_EFF_INDEX: usize = 128; +/// Plan C Phase 2 follow-up A.2 (2026-04-29): adaptive Polyak-tau drift signal. +/// +/// Per-epoch normalised q-drift rate ∈ [0, 4]. Producer: +/// `q_drift_rate_ema_update` GPU kernel (cold-path cadence, single-block +/// single-thread). Reads host-passed `q_mean_curr`, `q_mean_prev` plus +/// ISV[Q_ABS_REF_INDEX=16] and ISV[Q_DIR_ABS_REF_INDEX=21] to compute +/// `drift_rate = |q_curr − q_prev| / max(|q_prev|, |Q|_mag + |Q|_dir, 1e-6)` +/// clipped to `[0, 4]`. Consumer: `tau_update_kernel.cu` multiplies the +/// cosine-scheduled `tau_base` by `1 / (1 + ISV[129])` so the effective +/// Polyak rate decays from `tau_base` (healthy run, drift≈0) down to +/// `tau_base / 5` (peak drift, dampening saturated). +/// +/// Cold-start (constructor): 0.0 (no drift before any epochs have run). +/// FoldReset: 0.0 (drift is computed against `prev_epoch_q_mean` which +/// also resets to 0 at fold boundary, see A.1 in +/// `DQNTrainer::reset_for_fold`; the launch site gates on +/// `prev_epoch_q_mean.abs() > 1e-6` so the first epoch of a new fold +/// emits no drift). +/// +/// Tail-appended at index 129 (one past the previous `ISV_TOTAL_DIM=129` +/// boundary, raising it to 130) per `feedback_no_partial_refactor.md`. +/// Layout fingerprint changes — checkpoint-incompatible per +/// `feedback_no_legacy_aliases.md`. The `4.0` upper clip is an +/// architectural drift-saturation bound (matches the kill-criterion's +/// 3.0× ratio threshold in `training_loop.rs`); the `1e-6` denom floor +/// is a numerical-stability bound (Invariant 1 carve-out per +/// `feedback_isv_for_adaptive_bounds.md`). Per +/// `feedback_no_quickfixes.md` — neither is a tuned constant. +pub const Q_DRIFT_RATE_INDEX: usize = 129; + /// ISV slot [115] — low 32 bits of the u64 layout fingerprint (stored as raw f32 bits). /// /// Design note: this is NOT a version number. There is no ordered version space. @@ -951,7 +992,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { MOE_EXPERT_UTIL_EMA_BASE=118;MOE_EXPERT_UTIL_EMA_COUNT=8;\ MOE_GATE_ENTROPY_EMA=126;\ MOE_LAMBDA_EFF=128;\ - ISV_TOTAL_DIM=129;\ + Q_DRIFT_RATE=129;\ + ISV_TOTAL_DIM=130;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ @@ -3182,6 +3224,16 @@ pub struct GpuDqnTrainer { /// Loaded from `h_s2_rms_ema_kernel.cubin`. h_s2_rms_ema_kernel: CudaFunction, + // ── Plan C Phase 2 follow-up A.2: q-drift rate ISV producer ─────── + /// Single-thread single-block cold-path kernel. Reads ISV[Q_ABS_REF=16] + + /// ISV[Q_DIR_ABS_REF=21] plus host-passed `q_mean_curr` / `q_mean_prev` + /// scalars; writes `ISV[Q_DRIFT_RATE_INDEX=129] = clip(|Δq|/denom, 0, 4)` + /// at every epoch boundary. Consumer: `tau_update_kernel.cu` multiplies + /// the cosine-scheduled `tau_base` by `1/(1+ISV[129])` so tau ranges + /// `[tau_base/5, tau_base]` — monotone dampening under drift; healthy + /// runs unaffected. Loaded from `q_drift_rate_ema_kernel.cubin`. + q_drift_rate_ema_kernel: CudaFunction, + // ── Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMA producer ── /// 4-block kernel (one per off-median quantile τ ∈ {0.05, 0.25, 0.75, 0.95}) /// computing the mean |Q| over (B × TBA) at each fixed-τ position from @@ -8014,6 +8066,57 @@ impl GpuDqnTrainer { Ok(()) } + /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update` + /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129]. + /// + /// Reads ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] plus the + /// host-passed `q_mean_curr` / `q_mean_prev` scalars; writes + /// `clip(|q_curr − q_prev| / max(|q_prev|, ISV[16]+ISV[21], 1e-6), 0, 4)` + /// to ISV[129]. Consumer (`tau_update_kernel.cu`) multiplies the + /// cosine-scheduled `tau_base` by `1 / (1 + ISV[129])` so the effective + /// Polyak rate dampens monotonically as q-drift rises (range + /// `[tau_base/5, tau_base]`). + /// + /// Cold-path cadence — invoked once per epoch boundary AFTER the per- + /// step ISV producers (`launch_h_s2_rms_ema`, `launch_aux_heads_loss_ema`, + /// etc.) but BEFORE the next epoch's `launch_tau_update` reads the slot. + /// CPU-born inputs (`q_mean_curr`, `q_mean_prev`) are schedule values per + /// spec §4.C.6 — passed by value, not via mapped-pinned plumbing. + /// + /// Per `pearl_cold_path_no_exception_to_gpu_drives.md` — even a cold-path + /// scalar arithmetic stays on GPU when its inputs already live on GPU + /// (ISV[16], ISV[21] both produced by `q_stats_kernel.cu`). + pub fn launch_q_drift_rate_ema( + &self, + q_mean_curr: f32, + q_mean_prev: f32, + ) -> Result<(), MLError> { + // Same invariant rationale as `launch_h_s2_rms_ema`. Loud failure + // (debug_assert + kernel-launch error) preferred over silent skip. + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_q_drift_rate_ema: isv_signals_dev_ptr must be allocated by constructor"); + let isv_dev_ptr = self.isv_signals_dev_ptr; + let q_abs_ref_idx = Q_ABS_REF_INDEX as i32; + let q_dir_abs_ref_idx = Q_DIR_ABS_REF_INDEX as i32; + let q_drift_rate_idx = Q_DRIFT_RATE_INDEX as i32; + unsafe { + self.stream.launch_builder(&self.q_drift_rate_ema_kernel) + .arg(&isv_dev_ptr) + .arg(&q_abs_ref_idx) + .arg(&q_dir_abs_ref_idx) + .arg(&q_drift_rate_idx) + .arg(&q_mean_curr) + .arg(&q_mean_prev) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("q_drift_rate_ema_update: {e}")))?; + } + Ok(()) + } + /// Plan 4 Task 1B-iii: launch `vsn_mask_ema_update` (single-block, /// 256-thread shmem-reduction kernel; no atomicAdd). Reads /// `vsn_mask_buf [B, num_groups]` saved by the online-on-states VSN @@ -9574,6 +9677,16 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("h_s2_rms_ema_update load: {e}")))? }; + // Plan C Phase 2 follow-up A.2: load q_drift_rate_ema kernel + // (cold-path, per-epoch). Single-thread single-block ISV producer + // for ISV[Q_DRIFT_RATE_INDEX=129]; consumer is `tau_update_kernel`. + let q_drift_rate_ema_kernel = { + let module = stream.context().load_cubin(Q_DRIFT_RATE_EMA_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("q_drift_rate_ema cubin load: {e}")))?; + module.load_function("q_drift_rate_ema_update") + .map_err(|e| MLError::ModelError(format!("q_drift_rate_ema_update 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 @@ -11300,6 +11413,17 @@ impl GpuDqnTrainer { } *sig_ptr.add(MOE_GATE_ENTROPY_EMA_INDEX) = (MOE_EXPERT_UTIL_EMA_COUNT as f32).ln(); + /* Plan C Phase 2 follow-up A.2 (2026-04-29): + * cold-start ISV[Q_DRIFT_RATE_INDEX=129] at 0.0 so the + * very first `tau_update` consumer (which multiplies + * tau_base by 1/(1+ISV[129])) sees a no-op dampening + * factor of 1.0 before the producer kernel + * `q_drift_rate_ema_update` fires. FoldReset reapplies the + * same 0.0 (no drift at the start of a new fold; the + * first epoch's `prev_epoch_q_mean.abs() > 1e-6` gate in + * `training_loop.rs` prevents the producer launch until + * a meaningful prev value exists). */ + *sig_ptr.add(Q_DRIFT_RATE_INDEX) = 0.0_f32; // Layout fingerprint (ISV[58..60)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so @@ -12386,6 +12510,7 @@ impl GpuDqnTrainer { popart_count, reward_component_ema_kernel, h_s2_rms_ema_kernel, + q_drift_rate_ema_kernel, iqn_quantile_ema_kernel, q_mean_reduce_kernel, q_mean_subtract_kernel, diff --git a/crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu b/crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu new file mode 100644 index 000000000..61488d943 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu @@ -0,0 +1,87 @@ +/* q_drift_rate_ema_update — adaptive Polyak-tau dampening signal, GPU-driven. + * + * Plan C Phase 2 follow-up A.2 (researcher report 2026-04-29). Replaces the + * static `tau_base` Polyak EMA target sync with an ISV-driven adaptive + * `tau_eff` coupled to per-epoch q-drift rate. `tau_update_kernel.cu` reads + * ISV[Q_DRIFT_RATE_INDEX] (this kernel's output) at every epoch boundary and + * multiplies the cosine-scheduled tau by `1 / (1 + clip(drift, 0, 4))`, + * yielding a monotone dampening factor in `[1/5, 1]` — tau only ever + * decreases under drift; healthy runs are unaffected. + * + * Reads: ISV[Q_ABS_REF_INDEX] — magnitude-branch |Q| EMA (slot 16) + * ISV[Q_DIR_ABS_REF_INDEX] — direction-branch |Q| EMA (slot 21) + * q_mean_curr — host-passed scalar (epoch's q_mean) + * q_mean_prev — host-passed scalar (previous epoch's q_mean) + * Writes: ISV[Q_DRIFT_RATE_INDEX] — clipped drift rate ∈ [0, 4] (slot 129) + * + * Cold-path cadence — single-block, single-thread kernel (matches + * `tau_update_kernel.cu`, `kelly_cap_update_kernel.cu`, and + * `moe_lambda_eff_kernel.cu` shape). Launched once per epoch boundary AFTER + * the per-step ISV producers (`h_s2_rms_ema`, `aux_heads_loss_ema`, etc.) + * but BEFORE the next epoch's `tau_update` consumer reads ISV[129]. + * + * Formula: + * + * denom = max(|q_mean_prev|, ISV[Q_ABS_REF] + ISV[Q_DIR_ABS_REF], 1e-6) + * drift = |q_mean_curr - q_mean_prev| / denom + * ISV[129] = min(drift, 4.0) + * + * The `1e-6` in `denom` is a numerical-stability bound (Invariant 1 carve-out + * per `feedback_isv_for_adaptive_bounds.md`). The `4.0` upper clip is an + * architectural "drift starts here" bound matching the kill-criterion's + * 3.0× ratio threshold (in `training_loop.rs` Q-drift kill); a drift of 4 + * means the next-epoch tau_eff drops to `tau_base / 5` (5× slower target + * sync). `feedback_no_quickfixes.md` carve-out — both bounds are structural, + * not tuned. + * + * Per `pearl_cold_path_no_exception_to_gpu_drives.md`: even cold-path scalar + * arithmetic stays on GPU when its inputs already live on GPU — ISV[16,21] + * are produced by `q_stats_kernel.cu` and consumed here without DtoH. The + * two host-passed q-mean scalars are CPU-born monitor outputs (schedule + * inputs per spec §4.C.6), legal as kernel by-value arguments. + * + * No atomicAdd (per `feedback_no_atomicadd.md`); a single-thread reads three + * slots and writes one. The `__threadfence_system()` matches + * `moe_lambda_eff_update`'s emission so the mapped-pinned host_ptr sees the + * fresh value before the next `launch_tau_update` reads ISV[129]. + * + * Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic + * drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold- + * start) + reset (0.0 fold boundary) + observability (HEALTH_DIAG via slot + * 129 mirror). + */ + +extern "C" __global__ void q_drift_rate_ema_update( + float* __restrict__ isv, /* ISV bus [ISV_TOTAL_DIM] */ + int isv_q_abs_ref_idx, /* Q_ABS_REF_INDEX = 16 */ + int isv_q_dir_abs_ref_idx, /* Q_DIR_ABS_REF_INDEX = 21 */ + int isv_q_drift_rate_idx, /* Q_DRIFT_RATE_INDEX = 129 */ + float q_mean_curr, /* current epoch's q_mean (host-passed) */ + float q_mean_prev /* previous epoch's q_mean (host-passed) */ +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float q_abs_ref = isv[isv_q_abs_ref_idx]; + const float q_dir_abs_ref = isv[isv_q_dir_abs_ref_idx]; + + /* Denominator picks the largest of: previous q_mean magnitude (so drift + * is meaningfully scale-invariant when q has grown), the two per-branch + * |Q| EMAs (so the ratio uses the policy's currently-observed Q-scale, + * matching the kill-criterion's adaptive floor), and a 1e-6 numerical + * floor (Invariant 1 carve-out — guards against the cold-start case + * where every term is 0 before any epoch has produced statistics). */ + const float prev_abs = fabsf(q_mean_prev); + const float ref_sum = q_abs_ref + q_dir_abs_ref; + const float denom = fmaxf(fmaxf(prev_abs, ref_sum), 1e-6f); + const float diff_abs = fabsf(q_mean_curr - q_mean_prev); + const float drift_raw = diff_abs / denom; + + /* Clip to [0, 4]. Lower bound is structurally guaranteed (fabsf output + * is non-negative, division by positive denom preserves sign). Upper + * bound 4 matches the kill-criterion's 3.0× ratio threshold; beyond 4 + * the dampening factor 1/(1+drift) saturates at 0.2 = tau_base/5. */ + const float drift_clipped = fminf(drift_raw, 4.0f); + + isv[isv_q_drift_rate_idx] = drift_clipped; + __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 d96b67616..a5356460f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2174,3 +2174,19 @@ Assertions: Plan C T11 follow-up A.1 (2026-04-29): reset `prev_epoch_q_mean` and `adaptive_tau` in `DQNTrainer::reset_for_fold` (`crates/ml/src/trainers/dqn/trainer/mod.rs`). Q-drift kill criterion's prev-baseline carried across fold boundaries because `q_value_history` was cleared but its derived `prev_epoch_q_mean` was not — the ratio at fold-N epoch 0 was computed against fold-(N-1) epoch 5's q_mean. Companion `adaptive_tau` modulation state also reset to `hyperparams.tau` baseline. Independent bug fix; applies regardless of Plan C Thompson outcome. Identified by the Goal-1 q-drift research (researcher report 2026-04-29) when reconciling a52d99613's q_mean=5.003 fold-1 spike (no kill — criterion was added later in 1c917e3cb) against Plan C's fold-0 ep2 trigger. + +Plan C Phase 2 follow-up A.2 (2026-04-29): adaptive Polyak-tau coupled to per-epoch q-drift rate. Two-commit landing per `feedback_no_partial_refactor.md` — kernel + Rust wrapper first (additive, no callers), then ISV slot + reset registry + producer launch + tau_eff consumer wired together. + +* **New CUDA kernel** `crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu` (single-block single-thread cold-path producer mirroring `moe_lambda_eff_kernel.cu` / `kelly_cap_update_kernel.cu` shape). Reads ISV[Q_ABS_REF_INDEX=16] + ISV[Q_DIR_ABS_REF_INDEX=21] plus host-passed `q_mean_curr` / `q_mean_prev` scalars; writes `clip(|q_curr − q_prev| / max(|q_prev|, |Q|_mag + |Q|_dir, 1e-6), 0, 4)` to ISV[Q_DRIFT_RATE_INDEX=129]. Registered in `crates/ml/build.rs` kernel list (count grew by 1). + +* **New ISV slot** `Q_DRIFT_RATE_INDEX = 129` (tail-appended after `MOE_LAMBDA_EFF_INDEX = 128`, raising `ISV_TOTAL_DIM` 129→130). Cold-start 0.0 (constructor `*sig_ptr.add(Q_DRIFT_RATE_INDEX) = 0.0_f32`); FoldReset 0.0 via new `isv_q_drift_rate` registry entry in `state_reset_registry.rs` + dispatch arm in `training_loop.rs::reset_named_state`. Layout fingerprint shifts (one slot added + ISV_TOTAL_DIM bump) — checkpoint-incompatible per `feedback_no_legacy_aliases.md`, expected for a real architecture change. + +* **Rust orchestrator wrapper** `GpuDqnTrainer::launch_q_drift_rate_ema(q_mean_curr: f32, q_mean_prev: f32)` in `gpu_dqn_trainer.rs` — same pattern as `launch_h_s2_rms_ema` (debug_assert non-zero `isv_signals_dev_ptr`, single-thread launch config, kernel-launch error mapped to `MLError::ModelError`). + +* **Per-epoch producer launch** in `training_loop.rs` directly between the Q-drift kill criterion and the `prev_epoch_q_mean` update (so both `q_mean_curr` and `q_mean_prev` are in scope, AND the launch is gated on the same `prev_epoch_q_mean.abs() > 1e-6` cold-start guard the kill check uses). Fires once per epoch boundary; the next epoch's `launch_tau_update` (top of loop iteration) reads the freshly-written ISV[129]. + +* **Tau consumer wire-up** in `tau_update_kernel.cu` — after the existing cosine-schedule + health-coupled-floor clamp, the kernel multiplies `tau_eff` by `1 / (1 + clip(ISV[129], 0, 4))` so the effective Polyak rate dampens monotonically as q-drift rises (range `[tau_base/5, tau_base]`). Defensive `fminf(..., 4.0f)` in the kernel guards against producer-side bugs even though the producer already clips to `[0, 4]` (Invariant 1 carve-out, belt-and-braces). + +Predicted impact: Plan C fold 0 ep2 with `q_mean(t)=0.82, q_mean(t-1)=-0.018, ISV[16]+ISV[21]≈0.1` → drift_rate ≈ |0.84|/max(0.018, 0.1, ε) = 8.4 → clipped to 4 → tau_eff = tau_base × 0.2 (5× slower target sync, dampens Q-target optimism through inflation spike). a52d99613 fold 0 with q_mean staying ~0.21 → drift_rate ≈ 0 → tau_eff = tau_base (unchanged, healthy run unaffected). + +Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold-start) + reset (0.0 fold boundary) + observability. Per `pearl_cold_path_no_exception_to_gpu_drives.md` — even cold-path scalar arithmetic stays on GPU when its reduction inputs already live on GPU (ISV[16,21] both produced by `q_stats_kernel.cu`). Per `feedback_adaptive_not_tuned.md` — tau decay is ISV-driven, not a constant. Per `feedback_isv_for_adaptive_bounds.md` — the `4.0` upper clip and `1e-6` denom floor are structural bounds (drift saturation matching the kill-criterion's 3.0× ratio threshold; numerical-stability floor) carried in ISV slot semantics.