gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(dqn): q_drift_rate_ema kernel + Rust wrapper (Plan C A.2 scaffolding)

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-29 18:29:20 +02:00
parent dfbadf2f3f
commit fc8dbb0a85
4 changed files with 240 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
crates/ml/build.rs
View File

@@ -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`

129
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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,

87
crates/ml/src/cuda_pipeline/q_drift_rate_ema_kernel.cu Normal file
View File

@@ -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 */
}

16
docs/dqn-wire-up-audit.md
View File

@@ -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.