From d8666a232ff998b7e14fe060055b29785d33cc6f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 30 Apr 2026 21:36:11 +0200 Subject: [PATCH] =?UTF-8?q?docs(sp4):=20fix=20all=2011=20review=20findings?= =?UTF-8?q?=20=E2=80=94=20spec=20is=20now=20source=20of=20truth?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Resolved every issue from the critical self-review: 1. Param-group count: 7 → 8 throughout. Slot total: 36 → 40 (8 groups × 3 per-group families = 24 + 16 single = 40). IQL high-tau and low-tau are 2 distinct param groups (separate buffers + Adam states), not one. The "(Resolved during implementation)" hand-wave deleted. 2. Reset semantics section rewritten — no longer references Xavier- derived bootstraps. Sentinel 0 per Pearl A; Wiener-state triples reset to 0 alongside ISV slots (160 reset entries total). 3-4. Weight decay and L1 lambda hardcoded α/bootstrap removed. Both now follow universal Pearls A+D contract (sentinel-detect + Wiener adaptive). For L1, the natural curriculum (λ ramps as gradient differentiates) emerges from the signal itself, no "Bootstrap = 0.0" constant needed. 5. ε naming collision resolved: ε_div = 1e-8 (Pearl D division-safety), ε_clamp_floor = 1.0 (Pearl A consumer cold-start floor). Distinct names, distinct purposes. 6. Per-signal kernel signature updated: removed stale `ema_alpha` arg, added `wiener_state` pointer + `wiener_state_offset` + uniform `alpha_meta` (structural meta-EMA constant). 7. Pearl C engagement counter race fixed. Replaced `clamp_engage_buf [group] += 1` (race) with proper register-then-tree-reduce pattern mirroring `dqn_grad_norm_kernel`. Per-thread register counter → block-shared tree-reduce → single block-leader writes to `clamp_engage_per_block_buf[engage_buf_offset + blockIdx.x]`. Host sums across blocks. No atomicAdd, consistent with feedback. 8. Pearl D state buffer allocation spelled out: `wiener_state_buf: MappedF32Buffer` of size 141 floats (47 slots × 3), per `feedback_no_htod_htoh_only_mapped_pinned`. Reset-registry entries: 40 + 40×3 (new bound slots + Wiener triples) + 7×3 (retrofit existing producers' Wiener states) = 181 reset entries. 9. `grad_norm_slow_ema` retirement spelled out in Layer C: removed entirely (sole consumer Mech 6 migrates to ISV[GRAD_CLIP_BOUND]). Also documented Q_ABS_REF=16 and H_S2_RMS_EMA=96 transition to "monitoring-only ISV" — producers stay running, only the orphaned Mech 1/2/5/6/9/10 consumer reads removed. 10. Histogram bin range fixed: linear-spaced bins from 0 to step_max (avoids log(0) singularity from earlier log-spaced design). Linear is also better for p99: top-of-distribution gets ~0.4% resolution per bin. Degenerate "step_max == 0" branch handles all-zero signals gracefully (skip ISV update, leave previous bound). 11. Pearl D-subsumes-Pearl-A claim CORRECTED: mathematically wrong. At t=0, Pearl D's formula yields `x_mean[0] = 0`, not `x[0]`. Pearl A's first-observation replacement requires an explicit sentinel-detection branch in the producer. Both pearls are necessary and complementary; they are not hierarchical. Slot count math now consistent: 1 target_q + 4 atom_pos + 3×8 per-group + 1 grad_clip + 1 h_s2 + 8 wd_rate + 1 l1_lambda = 40. Producer count: 15 fused (1 + 4 + 8 + 1 + 1). Co-Authored-By: Claude Opus 4.7 (1M context) --- ...-signal-driven-magnitude-control-design.md | 224 +++++++++++------- 1 file changed, 136 insertions(+), 88 deletions(-) diff --git a/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md b/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md index 32cad89f8..de1e9c53a 100644 --- a/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md +++ b/docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md @@ -32,7 +32,7 @@ This revision (post-self-review) addresses 8 critical issues from the first draf **In scope (resolved via Pearl D — Wiener-optimal adaptive α):** -- **EMA rates (α values in producer kernels)** — derived per-step from observed signal variance via Wiener-filter optimal smoothing (see "Pearl D: Wiener-optimal adaptive α" below). The recursion problem dissolves because the optimal-smoothing formula `α* = diff_var / (diff_var + sample_var)` collapses to first-observation (α=1) on step 0 and to anchored-anchor (α=0) on stationary signals. No tuning constants. This applies to the 36 new SP4 producers AND retrofits to the ~7 existing pre-SP4 ISV producers (h_s2_rms_ema, grad_norm_slow_ema, aux_heads_loss_ema, moe_gate_entropy_ema, vsn_mask_ema, iqn_quantile_ema, reward_component_ema) — single coordinated migration per `feedback_no_partial_refactor`. +- **EMA rates (α values in producer kernels)** — derived per-step from observed signal variance via Wiener-filter optimal smoothing (see "Pearl D: Wiener-optimal adaptive α" below). The recursion problem dissolves via a single uniform structural meta-α (= 1/N_step_history). No per-signal tuning constants. This applies to the 40 new SP4 ISV slots (15 producer kernels via Pearl B fusion) AND retrofits to the ~7 existing pre-SP4 ISV producers (h_s2_rms_ema, grad_norm_slow_ema, aux_heads_loss_ema, moe_gate_entropy_ema, vsn_mask_ema, iqn_quantile_ema, reward_component_ema) — single coordinated migration per `feedback_no_partial_refactor`. Note: `grad_norm_slow_ema` is retired entirely in Layer C since its sole consumer (Mech 6's `update_adaptive_clip` upper-bound) migrates to ISV[GRAD_CLIP_BOUND]. **Out of scope (categories D + E from the brainstorm):** mathematically derived theoretical constants (Adam β₁=0.9, β₂=0.999; Adam ε=1e-8 numerical denominator; Xavier init formula `√(2/K)`; attention 1/√d), and architectural shape constants (hidden_dim=256, num_atoms=51, num_quantiles=5, layer count, batch size). @@ -52,27 +52,34 @@ Consumer (clamp or diagnostic): **Why p99 over max-abs:** the EMA of `max(|v|)` is outlier-driven by construction — a single bad sample inflates the EMA and the bound chases the inflation. p99 is a robust statistic: a single outlier (in the top 1%) does not move the 99th percentile. Clamping at p99 means clipping the worst 1% of magnitudes to the top-of-typical, which is a precise statistical statement. -**Algorithm: dynamic-range histogram (NOT P²).** +**Algorithm: dynamic-range linear histogram (NOT P², NOT log-spaced).** -The first-draft spec proposed P² (Jain-Chlamtac), but P² is a sequential algorithm and does not parallelize cleanly across GPU threads. This revision uses a three-pass dynamic-range histogram in a single block: +The first-draft spec proposed P² (Jain-Chlamtac), but P² is a sequential algorithm and does not parallelize cleanly across GPU threads. An earlier revision proposed log-spaced bins from 0 to step_max, but `log(0) = -∞` creates a singularity. This revision uses **linear-spaced** bins, which avoid the log-singularity and give the right resolution where p99 lives (top of distribution): ``` Pass 1 (block-wide max-reduce): step_max = max(|v|) over the buffer + if step_max == 0: ISV[BOUND] is left unchanged this step (degenerate signal — skip update) Pass 2 (block-wide bin): - each thread reads strided slice, increments shared-memory bin counters - bins are LOG-spaced from 0 to step_max (256 bins; bin_i covers - [step_max × 2^(i/256 - 1), step_max × 2^((i+1)/256 - 1)] roughly) + each thread reads strided slice + bin_idx = floor(|v| / step_max × 256) (clamped to [0, 255]) + bins are LINEAR-spaced over [0, step_max], 256 bins of width step_max/256 + increments use per-thread register counter → block-shared tree-reduce + (no atomicAdd; per `feedback_no_atomicadd`) Pass 3 (block-wide cumulative-from-top): - thread 0 walks bins from highest to lowest, accumulates count - when cumulative ≥ N/100, stop — that bin's lower edge is p99 + thread 0 walks bins from index 255 down to 0, accumulates count + when cumulative ≥ N/100, stop — that bin's upper edge is p99 ``` -All three passes use `__syncthreads()` between them, no scratch buffer beyond shared memory (256 ints + 1 float). Single-block, 256-thread, deterministic, no atomicAdd (per `feedback_no_atomicadd`). Cost: ~5-10 μs per launch on L40S for buffers up to 64k elements. +Linear bins are the right choice for p99 because we care about resolution AT the top of the distribution. With 256 linear bins from 0 to step_max, the top bin width is `step_max/256 ≈ 0.4%` of step_max — sufficient for 1% quantile estimation with safety margin. -The 256-bin choice and log-spacing parameters are **derived from the resolution required for a 1% quantile** (need ≥100 bins for 1% precision; 256 gives ~0.4% precision with safety margin). This is a numerical-precision choice, not a tuning knob — same theoretical-constant category as floating-point precision. +All three passes use `__syncthreads()` between them, no scratch buffer beyond shared memory (256 ints + 1 float). Single-block, 256-thread, deterministic, no atomicAdd. Cost: ~5-10 μs per launch on L40S for buffers up to 64k elements. + +The 256-bin choice is **derived from the resolution required for a 1% quantile** (need ≥100 bins for 1% precision; 256 gives ~0.4% precision with safety margin). This is a numerical-precision choice, not a tuning knob — same theoretical-constant category as floating-point precision. + +The "if step_max == 0 skip update" branch handles the degenerate case where the signal is identically zero (e.g., grads after a NaN-sanitization). The producer leaves the ISV slot unchanged, so the consumer reads the previous bound. No NaN propagation through the histogram math. ### Pre-clamp vs post-clamp producer reads (the feedback-loop fix) @@ -116,22 +123,38 @@ The producer kernel reads PRE-clamp buffer for histogram-p99. Pass 1 computes `s **In-kernel clamps (Mech 6, 9 — weights and Adam m/v) — diagnostic INSIDE the kernel + Pearl C self-correction:** The producer for these signals reads POST-clamp distribution (since the kernel mutates in-place and there's no separate pre-clamp buffer). Post-clamp `|v| ≤ bound` by construction, so a producer-side magnitude comparison NEVER fires. -The diagnostic must therefore live inside the Adam (or weight-update) kernel itself, immediately at the clamp step. For each Adam kernel, two per-thread atomic-or-style flag/counter operations: +The diagnostic must therefore live inside the Adam (or weight-update) kernel itself, immediately at the clamp step. For each Adam kernel, the engagement count uses the existing block-wide reduction pattern from `dqn_grad_norm_kernel` (per-thread register counter → block-shared tree-reduce → single block-leader writes the per-block count to a per-(block, group) output array). ```cuda /* Mech 9 clamp + SP4 diagnostic + Pearl C engagement counter */ +__shared__ int s_block_engage[BLOCK_REDUCE_WIDTH]; // shared-mem reduction scratch +int local_engage = 0; // per-thread register counter + +// ... existing Adam compute → p_val ready for clamp/writeback ... if (weight_clamp_max_abs > 0.0f) { if (fabsf(p_val) > weight_clamp_max_abs) { - nan_flags_buf[diag_slot] = 1; // sticky flag — clamp engaging - clamp_engage_buf[group] += 1; // per-group engagement count (per-thread shared-mem then block-reduce) + nan_flags_buf[diag_slot] = 1; // sticky flag — race-free idempotent write + local_engage = 1; // per-thread register, no contention } p_val = fminf(fmaxf(p_val, -weight_clamp_max_abs), weight_clamp_max_abs); } + +/* Block-wide tree-reduce (no atomicAdd, mirrors dqn_grad_norm_kernel pattern) */ +s_block_engage[threadIdx.x] = local_engage; +__syncthreads(); +for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (threadIdx.x < s) s_block_engage[threadIdx.x] += s_block_engage[threadIdx.x + s]; + __syncthreads(); +} +if (threadIdx.x == 0) { + /* One write per block per Adam launch — block_idx selects the slot */ + clamp_engage_per_block_buf[engage_buf_offset + blockIdx.x] = s_block_engage[0]; +} ``` -The flag write is per-thread but idempotent (all threads writing 1 to the same address — race-free per the existing `nan_flags_buf` convention used at `dqn_utility_kernels.cu:1808`). The engagement count uses block-local shared-memory counters reduced once per block (no atomicAdd, consistent with `feedback_no_atomicadd`). Cost: two comparisons + two stores per affected thread per Adam launch. +The `nan_flags_buf` write is per-thread but idempotent (all threads writing 1 — race-free per existing convention at `dqn_utility_kernels.cu:1808`). The engagement count uses register-then-tree-reduce (no atomicAdd, consistent with `feedback_no_atomicadd`). Each block writes its local count to a separate slot of `clamp_engage_per_block_buf`, then the host sums across blocks (cheap reduction over `n_blocks` ≤ ~256 floats). Cost: two comparisons + one register update per affected thread + one block-wide reduction per Adam launch. -Each in-kernel-clamp Adam kernel takes a `diag_slot: i32` and `engage_buf_offset: i32` alongside `weight_clamp_max_abs: f32`. Launch site passes the appropriate offsets for that param-group. +Each in-kernel-clamp Adam kernel takes a `diag_slot: i32` and `engage_buf_offset: i32` alongside `weight_clamp_max_abs: f32`. Launch site passes the appropriate offsets for that param-group. `clamp_engage_per_block_buf` is a mapped-pinned `[N_GROUPS × MAX_BLOCKS]` int buffer (≤ 8 × 256 = 2048 ints ≈ 8 KiB). ### Pearl C — engagement-rate self-correction (post-clamp feedback-loop fix) @@ -163,24 +186,24 @@ This drops the `× 10` and `× 1e3` headroom patterns entirely. All three diagno ## Migration scope -### ISV slot allocation (36 new slots) +### ISV slot allocation (40 new slots) | Slot family | Count | Tracks | Consumed by | |---|---|---|---| | `TARGET_Q_BOUND` | 1 | p99(\|target_q\|) | Mech 1 clamp + slot 46 diagnostic | | `ATOM_POS_BOUND[branch]` | 4 (per branch) | p99(\|atom_positions[branch]\|) | Mech 2 clamps × 3 sites + slot 47 diagnostic | -| `WEIGHT_BOUND[group]` | 7 (per param-group) | p99(\|params\|) per group | Mech 9 clamps × 5 Adam kernels + slots 44/45/48 diagnostics | -| `ADAM_M_BOUND[group]` | 7 (per param-group) | p99(\|adam_m\|) per group | slots 36-39 diagnostics | -| `ADAM_V_BOUND[group]` | 7 (per param-group) | p99(\|adam_v\|) per group | slots 40-43 diagnostics | +| `WEIGHT_BOUND[group]` | 8 (per param-group) | p99(\|params\|) per group | Mech 9 clamps × 5 Adam kernels + slots 44/45/48 diagnostics | +| `ADAM_M_BOUND[group]` | 8 (per param-group) | p99(\|adam_m\|) per group | slots 36-39 diagnostics | +| `ADAM_V_BOUND[group]` | 8 (per param-group) | p99(\|adam_v\|) per group | slots 40-43 diagnostics | | `GRAD_CLIP_BOUND` | 1 | p99(grad_norm) | Mech 6 + adaptive_clip | | `H_S2_BOUND` | 1 | p99(\|h_s2\|) | Mech 10 clamp + slot 49 diagnostic | -| `WD_RATE[group]` | 7 (per param-group) | EMA of \|w·g\|/\|\|w\|\|² per group | Adam kernel `weight_decay` arg | +| `WD_RATE[group]` | 8 (per param-group) | EMA of \|w·g\|/\|\|w\|\|² per group | Adam kernel `weight_decay` arg | | `L1_LAMBDA[group]` | 1 (only trunk has L1) | (mean(\|g\|)/mean(\|w\|)) × entropy_deficit | Adam kernel `l1_lambda` arg | -| **Total** | **36** | | | +| **Total** | **40** | | | ### Param-group definitions (for the `[group]` slots) -7 distinct param groups corresponding to the 5 Adam kernels' param buffers. Each producer reads its group's `(buf_ptr, offset, length)` and computes p99 over that slice: +**8 distinct param groups** corresponding to the live training-path Adam kernels' param buffers (DT excluded — offline-RL, outside SP4 scope). Each producer reads its group's `(buf_ptr, offset, length)` and computes p99 over that slice: | Group index | Adam kernel | Param-buffer slice | |---|---|---| @@ -188,42 +211,48 @@ This drops the `× 10` and `× 1e3` headroom patterns entirely. All three diagno | 1 | `dqn_adam_update_kernel` | value-head slice (within main params, contiguous) | | 2 | `dqn_adam_update_kernel` | branch-heads slice (4 branches concatenated) | | 3 | `iqn_adam_kernel` | IQN online_params (separate buffer entirely) | -| 4 | `iql_adam_kernel` (high-tau) | IQL value-head params | -| 5 | `iql_adam_kernel` (low-tau) | IQL low-tau params | +| 4 | `iql_adam_kernel` (high-tau instance) | IQL value-head params (`gpu_iql.params`) | +| 5 | `iql_adam_kernel` (low-tau instance) | IQL low-tau params (`gpu_iql_low.params`) | | 6 | `attn_adam_kernel` | attention/TLOB params (shared cubin) | | 7 | `curiosity_adam_step` | curiosity-module params | -(Note: 8 Adam-launch entry points map to 7 param groups because IQL high-tau and low-tau may share a cubin. Resolved during implementation — minor clarification.) +IQL high-tau and low-tau share the `iql_adam_kernel` cubin but have separate param buffers and Adam states (one `GpuIql` instance per τ value). They are 2 distinct param groups (4, 5), 8 groups total. The trunk/value-head/branch split for the main DQN params buffer requires the producer to know offsets. These offsets are already known to the Mech 9 launch sites (since they pass `total_params` etc.). Producer kernels take `(offset, length)` as launch args, same pattern. ### EMA rates — Pearl D: Wiener-optimal adaptive α -**No hardcoded α values across all 43 producer kernels (36 new + 7 retrofit).** Each producer derives its own α per step from observed signal variances via Wiener-filter optimal smoothing. +**No hardcoded α values across all 47 producer kernels (40 new + 7 retrofit).** Each producer derives its own α per step from observed signal variances via Wiener-filter optimal smoothing. For a signal `x[t]` (the producer's per-step observation, e.g., step_p99): ``` -sample_var[t] = EMA( (x[t] - x_mean[t-1])² ) // signal variance around mean -diff_var[t] = EMA( (x[t] - x[t-1])² ) // step-to-step variance -α*[t] = diff_var[t] / (diff_var[t] + sample_var[t] + ε_num) +sample_var[t] = (1-α_meta) · sample_var[t-1] + α_meta · (x[t] - x_mean[t-1])² +diff_var[t] = (1-α_meta) · diff_var[t-1] + α_meta · (x[t] - x_lag[t])² +α*[t] = diff_var[t] / (diff_var[t] + sample_var[t] + ε_div) x_mean[t] = (1 - α*[t]) · x_mean[t-1] + α*[t] · x[t] +x_lag[t+1] = x[t] ``` -`ε_num = 1e-8` is numerical safety against 0/0 (Adam-ε category, not magnitude tuning). +`ε_div = 1e-8` is **division-safety** (Adam-ε numerical category — protects 0/0 in α* formula, not a magnitude estimate). -**Why this is theoretically optimal:** for a signal modeled as `x[t] = μ[t] + noise[t]` where `μ[t]` is the latent value we want to track and `noise[t]` is observation noise, the optimal estimator weight under Wiener filter analysis is exactly `α* = process_variance / (process_variance + noise_variance)`. `diff_var` estimates process variance (how much μ moves between steps); `sample_var` estimates noise variance. +`α_meta` is the META-EMA rate for tracking the variances themselves. To avoid the recursion problem (what controls α_meta?), it is fixed at `α_meta = 1 / N_step_history` where N_step_history is structurally determined by the smoke's per-fold step count (typically ~1000 → α_meta = 1e-3). Documented as a structural-design parameter (analogous to "buffer size" — derived from training cardinality, not magnitude tuning). One value across all producers (uniform meta-smoothing). + +**Why this is theoretically optimal:** for a signal modeled as `x[t] = μ[t] + noise[t]` where `μ[t]` is the latent value we want to track and `noise[t]` is observation noise, the optimal Wiener-filter estimator weight is exactly `α* = process_variance / (process_variance + noise_variance)`. `diff_var` estimates process variance (how much μ moves between steps); `sample_var` estimates noise variance. **Behavior:** -- **Step 0 (cold start):** sample_var = 0 (no history), diff_var = 0 (no prior step), α* = 0/(0+ε) ≈ 0. But Pearl A (first-observation) overrides at step 0: producer detects sentinel `x_mean = 0`, replaces directly with x[0]. From step 1, α* derives normally. +- **Step 0 (cold start):** Pearl A (sentinel-detection branch in the producer) overrides Pearl D entirely. Producer sees `x_mean[t-1] = 0` (sentinel) and writes `x_mean[0] = x[0]` directly, bypassing the Wiener formula. Wiener state buffers initialize at `(sample_var=0, diff_var=0, x_lag=x[0])` for the next step. +- **Step 1+:** Pearl D applies normally. α* derives from variance ratio. - **Stationary signal:** μ is constant. diff_var → 0 (no step-to-step movement). sample_var > 0 (noise around the constant). α* → 0. Mean stays anchored. Correct behavior. - **Non-stationary signal:** μ is moving. diff_var dominates. α* → 1. Mean tracks new observations. Correct behavior. - **Mid-regime:** α* in (0, 1), exact value reflects signal-to-noise ratio of process change. -**Producer state per signal:** 3 floats (`x_mean`, `sample_var`, `diff_var`) + 1 float for `x[t-1]` lag = 4 floats. Plus the producer's own working state (e.g., histogram bins). Negligible. +**Producer state per ISV bound slot:** 3 floats (`sample_var`, `diff_var`, `x_lag`) stored in a dedicated mapped-pinned buffer `wiener_state_buf: MappedF32Buffer` of total size `40 × 3 = 120 floats` (480 bytes). Per `feedback_no_htod_htoh_only_mapped_pinned`. Plus the producer's own working state (e.g., histogram bins in shared memory). -**Subsumes Pearl A:** the first-observation pattern is the t=0 limit of Pearl D. Pearl A is the t=0 special case; Pearl D is the general formula. Both are documented but Pearl D's implementation handles both. +The current EMA value (`x_mean`) IS the ISV bound slot itself — no separate storage. Producer reads `isv[BOUND_X]` (= old x_mean), computes new x_mean using Pearl D formula + state buffer, writes back to `isv[BOUND_X]`. -**Subsumes the per-signal time-window argument from the first-draft spec.** No more "α=0.05 for fast signals, α=0.001 for slow signals" choice — α adapts per-signal per-step automatically. The half-life table from the first draft is replaced by the formula. +**Pearl D does NOT subsume Pearl A.** Pearl A's first-observation replacement requires an explicit sentinel-detection branch in the producer (`if isv[BOUND_X] == 0.0 then write step_observation directly, bypassing Wiener`). At step 0 Pearl D's formula gives `x_mean[0] = (1-0)·0 + 0·x[0] = 0` (NOT `x[0]`) — wrong. Both pearls are necessary; Pearl A handles t=0, Pearl D handles t≥1. The earlier claim that D subsumes A was mathematically incorrect and is corrected here. + +**Subsumes the per-signal time-window argument from the first-draft spec.** No more "α=0.05 for fast signals, α=0.001 for slow signals" per-producer choice — α* adapts per-signal per-step automatically. ### Cold-start — Pearl A: First-observation bootstrap @@ -231,26 +260,32 @@ x_mean[t] = (1 - α*[t]) · x_mean[t-1] + α*[t] · x[t] ``` At fold-boundary reset: isv[BOUND_X] = 0 (sentinel "uninitialized") -Producer step 0 sees 0: isv[BOUND_X] = step_observation (full replacement, no EMA blend) -Producer step 1+: α*-blended EMA per Pearl D -Consumer cold-start: bound = isv[BOUND_X].max(ε_num) (ε_num = 1.0, numerical safety) + wiener_state[BOUND_X] = (0, 0, 0) (sample_var, diff_var, x_lag all 0) +Producer step 0 detects sentinel (isv[BOUND_X] == 0): + isv[BOUND_X] = step_observation (full replacement, NO Wiener math) + wiener_state[BOUND_X].x_lag = step_observation +Producer step 1+: Pearl D Wiener-optimal blend (sentinel no longer 0) +Consumer cold-start: bound = isv[BOUND_X].max(ε_clamp_floor) (ε_clamp_floor = 1.0) ``` -`ε_num = 1.0` is purely numerical (prevents `bound = 0` from collapsing the clamp at step 0 before producer runs). Same Adam-ε category — protects arithmetic, not a magnitude estimate. - -**The bootstrap section in the first draft (Xavier formulas, `2.33`, `√2`, `√(2/K_in)`) is replaced entirely by this 4-line pattern.** The pattern applies uniformly to all 43 producers (36 new + 7 retrofit) — no per-bound formula table needed, no `K_in[group]` lookup, no `LR_init` read at construction. +**Two epsilons (distinct names, distinct purposes):** +- `ε_clamp_floor = 1.0` — consumer-side numerical safety floor. Prevents `bound = 0` from collapsing the clamp at step 0 before any producer runs. Same Adam-ε category — protects arithmetic, not a magnitude estimate. +- `ε_div = 1e-8` — producer-side division safety in Pearl D's α* formula. Prevents 0/0. Different value, different role; named separately to avoid confusion. **Why this is a pearl:** - Eliminates every magnitude-derived constant from cold-start. - Adapts to actual signal distribution at fold start. F1's natural p99 may differ from F0's; the bootstrap automatically tracks. - Pure form of "the bound lives in ISV": even cold-start values come from observation, not theory. -- Subsumed by Pearl D as the t=0 edge case (when sample_var = diff_var = 0, the formula degenerates to "use the observation directly"). + +**Pearl A and Pearl D are complementary, NOT hierarchical.** Pearl A handles t=0 (sentinel-detection branch); Pearl D handles t≥1 (Wiener formula). The earlier claim that D subsumes A was mathematically incorrect — Pearl D's formula at t=0 yields `x_mean[0] = 0`, not `x[0]`. Both mechanisms are explicit in producer code. ### Reset semantics -Per `StateResetRegistry`: at fold boundary, all 28 ISV bound slots reset to their bootstrap values (computed once at trainer construction from K_in dimensions and LR). Producer kernels start fresh, EMA tracks from bootstrap. This automatically handles the cross-fold anchor staleness without needing Mech 8. +Per `StateResetRegistry`: at fold boundary, all 40 ISV bound slots AND their `wiener_state` triples (sample_var, diff_var, x_lag) reset to **0** (Pearl A sentinel — no theoretical-formula bootstrap). The producer's first step after reset sees the sentinel and replaces with `step_observation` (Pearl A). Subsequent steps use Pearl D's adaptive Wiener formula. -Reset entries follow existing pattern in `state_reset_registry.rs` (one entry per slot with cold-start value and description). +Reset entries follow existing pattern in `state_reset_registry.rs`: one entry per ISV slot for `isv[BOUND_X] = 0` plus three entries per slot for `wiener_state[BOUND_X].{sample_var, diff_var, x_lag} = 0`. Total: `40 + 40×3 = 160` reset entries (compact tabular addition; existing infrastructure handles iteration). + +This automatically handles the cross-fold anchor-staleness pathology without needing Mech 8. ### Weight decay producer signal @@ -267,9 +302,9 @@ Producer kernel `wd_rate_update[group]`: - Pass 2 (block-wide reduce): compute `w_norm_sq = Σ_i w[i]²` over the slice - Final: `step_wd = |w_dot_g| / max(w_norm_sq, bootstrap²)` ; EMA into `ISV[WD_RATE[group]]` -Adam kernels read `weight_decay = ISV[WD_RATE[group]]` from the launch site (no kernel-internal change beyond consuming the launch-site-supplied value, replacing the existing scalar field). Per-group → 7 producers, 7 ISV slots. +Adam kernels read `weight_decay = ISV[WD_RATE[group]]` from the launch site (no kernel-internal change beyond consuming the launch-site-supplied value, replacing the existing scalar field). Per-group → 8 producers, 8 ISV slots. -EMA half-life: α=0.005 (~140 steps, matches weight evolution time-scale). Bootstrap: `1.0` (cold-start neutral; first-step EMA pulls toward observed signal). +**EMA rate and bootstrap follow the universal Pearls A+D contract** — no hardcoded α (Pearl D Wiener-optimal adaptive α from variance ratio), no theoretical bootstrap (Pearl A sentinel-detection at t=0). The producer's per-step output `step_wd = |w·g|/||w||²` flows through the same Wiener-EMA path as every other producer. This is the same equilibrium-derivation principle the Adam paper uses for the β values — derived from a continuous-time stability analysis. Theoretically grounded, not tuned. @@ -303,15 +338,13 @@ Producer kernel `l1_lambda_update[group]`: Adam kernel reads `l1_lambda = ISV[L1_LAMBDA[group]]` at the launch site, passes through to the existing L1 proximal step. -EMA half-life: α=0.05 (~14 steps; L1 strength can react quickly to gradient-direction shifts since differentiation is the signal we're tracking). - -Bootstrap: `0.0` (cold-start no L1 — no pruning until network has differentiated). The fold-boundary reset to 0.0 means each new fold starts L1-free and the producer ramps up λ as gradient differentiates within the new fold's data regime. **This is exactly the right curriculum** — you don't want to prune during the first ~tens of steps when the network is learning the new fold's structure. +**EMA rate and bootstrap follow the universal Pearls A+D contract** — no hardcoded α, no fixed bootstrap. Pearl A sentinel-detection at t=0; Pearl D Wiener-optimal adaptive α at t≥1. The natural curriculum behavior emerges automatically: at fold boundary, the producer's first observation `step_λ = (mean|g|/mean|w|) × D` is small because `D ≈ 0` (gradient still uniform across features pre-differentiation), so the bound starts small and ramps as `D` grows. No `Bootstrap = 0.0` constant needed — the curriculum is in the signal itself. **Pearl name (post-validation):** `pearl_signal_driven_regularisation_strength` — extends `pearl_adaptive_moe_lambda` to a class of regularization-strength controllers driven by entropy-deficit of the regularized quantity. Captured in memory once SP4 ships and validates. ## Producer architecture — Pearl B: Fused per-param-group statistics oracle -**Producer count: 14 kernels (NOT 36).** The first-draft "36 separate producer kernels" wasted memory bandwidth — the 4 per-group signals (WEIGHT_BOUND, ADAM_M_BOUND, ADAM_V_BOUND, WD_RATE) all read the same 4 buffers (`params, grads, adam_m, adam_v`) for the same group. Fusing them into one kernel per group reads each buffer once and computes all 4 stats in a single multi-pass operation. L1_LAMBDA (trunk only) extends the trunk's fused producer with one more pass for the gradient-direction entropy. +**Producer count: 15 kernels (NOT 40).** The naive "40 separate producer kernels" would waste memory bandwidth — the 4 per-group signals (WEIGHT_BOUND, ADAM_M_BOUND, ADAM_V_BOUND, WD_RATE) all read the same 4 buffers (`params, grads, adam_m, adam_v`) for the same group. Fusing them into one kernel per group reads each buffer once and computes all 4 stats in a single multi-pass operation. L1_LAMBDA (trunk only) extends the trunk's fused producer with one more pass for the gradient-direction entropy. ### Final producer kernel inventory @@ -319,22 +352,23 @@ Bootstrap: `0.0` (cold-start no L1 — no pruning until network has differentiat |---|---|---|---| | `target_q_p99_update` | denoise_target_q_buf | TARGET_Q_BOUND | 1 | | `atom_pos_p99_update[branch]` | atom_positions[branch] slice | ATOM_POS_BOUND[branch] | 4 | -| `param_group_stats_oracle[group]` | params, grads, adam_m, adam_v (all per-group slice) | WEIGHT_BOUND[group], ADAM_M_BOUND[group], ADAM_V_BOUND[group], WD_RATE[group] | 7 | -| `l1_lambda_update` (extends trunk's oracle with gradient-entropy pass) | grads (trunk slice) | L1_LAMBDA[trunk] | (folded into group-0 oracle) | +| `param_group_stats_oracle[group]` | params, grads, adam_m, adam_v (all per-group slice) | WEIGHT_BOUND[group], ADAM_M_BOUND[group], ADAM_V_BOUND[group], WD_RATE[group] (+ L1_LAMBDA for group 0 only) | 8 | | `grad_norm_p99_update` | grad_norm_buf | GRAD_CLIP_BOUND | 1 | | `h_s2_p99_update` | save_h_s2 | H_S2_BOUND | 1 | -| **Total** | | **36 ISV slots** | **14 kernels** | +| **Total** | | **40 ISV slots** | **15 kernels** | The fused per-group oracle uses 4 sequential passes inside one kernel (sharing block sync barriers): -- Pass A: max-reduce + histogram-p99 of `|params|` → WEIGHT_BOUND[group] -- Pass B: max-reduce + histogram-p99 of `|adam_m|` → ADAM_M_BOUND[group] -- Pass C: max-reduce + histogram-p99 of `|adam_v|` → ADAM_V_BOUND[group] -- Pass D: dot-product `Σ w·g`, sum-of-squares `Σ w²`, plus `mean|g|`, `mean|w|` (single block-wide reduce) → WD_RATE[group] = |w·g|/||w||² -- (Group 0 only) Pass E: per-feature gradient L2 norm `||g_w[:,i]||`, normalize, compute entropy H, deficit D → L1_LAMBDA[trunk] +- Pass A: max-reduce + linear-histogram-p99 of `|params|` → WEIGHT_BOUND[group] +- Pass B: max-reduce + linear-histogram-p99 of `|adam_m|` → ADAM_M_BOUND[group] +- Pass C: max-reduce + linear-histogram-p99 of `|adam_v|` → ADAM_V_BOUND[group] +- Pass D: 4-way parallel block-wide reduce computing `Σ w·g`, `Σ w²`, `Σ |g|`, `Σ |w|` (each thread maintains 4 register accumulators; tree-reduce per accumulator using shared memory) → WD_RATE[group] = |w·g|/||w||²; the `mean|g|`, `mean|w|` are also available for L1's λ scale. +- (Group 0 only) Pass E: per-feature gradient L2 norm `||g_w[:,i]||` for i ∈ [0, K); single-thread normalize, entropy H, deficit D; L1_LAMBDA[trunk] = (mean|g|/mean|w|) × D. -Each pass uses the existing block-shared-memory infrastructure. Memory bandwidth: each buffer read once (4× reduction from non-fused). Per-step cost per oracle: ~30-50 μs (proportional to sum of buffer sizes). +All passes apply Pearls A+D to update each ISV bound slot via Wiener-optimal EMA with sentinel-detection (no hardcoded α, no theoretical bootstrap). -**Total per-step launches: 14** (vs 36 non-fused). Latency overhead ~0.3% of step time. +Each pass uses block-shared-memory infrastructure. Memory bandwidth: each buffer read once (4× reduction from non-fused). Per-step cost per oracle: ~30-50 μs (proportional to sum of buffer sizes). + +**Total per-step launches: 15** (vs 40 non-fused). Latency overhead ~0.3% of step time. **Why this is a pearl, not just an optimization:** the architectural insight is that per-param-group bounds are a *unit* — they all derive from the same buffers, serve the same Adam kernel, and must be reset together at fold boundaries. Fusing at the producer level mirrors the conceptual grouping. Future per-group signals (spectral-norm tracking, gradient-direction stability, etc.) extend the existing oracle by adding a pass — no new kernel infrastructure required. @@ -342,20 +376,22 @@ Each pass uses the existing block-shared-memory infrastructure. Memory bandwidth ### Per-signal kernels -The 14 producers, mirroring the existing `h_s2_rms_ema_update` shape: +The 15 producers, mirroring the existing `h_s2_rms_ema_update` shape but extended with Pearls A+D (sentinel-detection + Wiener-state pointers; NO `ema_alpha` arg). -Magnitude-bound producer signature (`target_q`, `atom_pos`, `grad_norm`, `h_s2`): +Single-output producer signature (`target_q`, `atom_pos`, `grad_norm`, `h_s2`): ```cuda extern "C" __global__ void _p99_update( - const float* __restrict__ buf, - int offset, - int length, - float* __restrict__ isv, - int isv_bound_index, - float ema_alpha, - int* __restrict__ nan_flags_buf, - int diag_slot + const float* __restrict__ buf, /* signal buffer */ + int offset, /* slice offset (bytes/4) */ + int length, /* slice length */ + float* __restrict__ isv, /* ISV bus */ + int isv_bound_index, /* this signal's ISV slot */ + float* __restrict__ wiener_state, /* mapped-pinned [N_BOUNDS × 3] */ + int wiener_state_offset, /* this signal's offset into wiener_state (= isv_bound_index × 3) */ + float alpha_meta, /* Pearl D meta-α for variance EMAs (single uniform value, structurally derived) */ + int* __restrict__ nan_flags_buf, /* sticky-flag diagnostic */ + int diag_slot /* this signal's diagnostic slot */ ); ``` @@ -391,17 +427,19 @@ Per `feedback_no_partial_refactor`, the contract change is one atomic commit. Th ### Layer A — additive infrastructure (no behavior change) -- Add 36 ISV slot indices to the bus (extend `ISV_TOTAL_DIM` as needed) -- Implement 14 producer kernels per Pearl B (1 target_q + 4 atom_pos + 7 fused-per-group oracle + 1 grad_norm + 1 h_s2), each with unit tests for every output slot it writes +- Add 40 ISV slot indices to the bus (extend `ISV_TOTAL_DIM` as needed) +- Allocate `wiener_state_buf: MappedF32Buffer` of size `40 × 3 = 120 floats` for Pearl D's per-slot (sample_var, diff_var, x_lag) state (per `feedback_no_htod_htoh_only_mapped_pinned`) +- Allocate `clamp_engage_per_block_buf: MappedI32Buffer` of size `8 × MAX_BLOCKS_PER_ADAM` for Pearl C's per-(block,group) engagement counters (≤ 2048 ints) +- Implement 15 producer kernels per Pearl B (1 target_q + 4 atom_pos + 8 fused-per-group oracle + 1 grad_norm + 1 h_s2), each with unit tests for every output slot it writes - Implement Pearl D (Wiener-optimal adaptive α) inside each producer's EMA-update step -- Implement Pearl A (sentinel = first-observation replacement) at the EMA step -- Implement Pearl C (engagement-rate self-correction) for in-kernel-clamp signals: per-Adam-kernel block-shared-memory engagement counter + host-side rate-deficit EMA + force-bump trigger -- Wire 14 launch sites (4 in captured-graph for buffer-based bounds; 10 in cold-path `training_loop.rs` for weight/Adam/regularization/grad-norm/h_s2) -- Retrofit Pearls A and D into the 7 existing pre-SP4 ISV producers (h_s2_rms_ema, grad_norm_slow_ema, aux_heads_loss_ema, moe_gate_entropy_ema, vsn_mask_ema, iqn_quantile_ema, reward_component_ema) — replace their hardcoded α with adaptive α; sentinel-bootstrap at fold boundary -- Add 36 `StateResetRegistry` entries (sentinel `0.0` per Pearl A — no Xavier formulas) +- Implement Pearl A (sentinel = first-observation replacement) at the EMA step (explicit branch — Pearl A is NOT subsumed by Pearl D) +- Implement Pearl C (engagement-rate self-correction) for in-kernel-clamp signals: per-Adam-kernel register-then-tree-reduce engagement counter + host-side rate-deficit EMA + force-bump trigger +- Wire 15 launch sites (4 in captured-graph for buffer-based bounds: target_q, 4 atom_pos; 11 in cold-path `training_loop.rs` for 8 fused oracles + grad-norm + h_s2) +- Retrofit Pearls A and D into the 7 existing pre-SP4 ISV producers (h_s2_rms_ema, grad_norm_slow_ema, aux_heads_loss_ema, moe_gate_entropy_ema, vsn_mask_ema, iqn_quantile_ema, reward_component_ema) — replace their hardcoded α with adaptive α; sentinel-bootstrap at fold boundary; allocate Wiener state for each in `wiener_state_buf` (extending its size to `(40 + 7) × 3 = 141 floats`) +- Add 40 `StateResetRegistry` entries for the new ISV bound slots (each resets to sentinel `0.0`) + 40 entries for the corresponding Wiener-state triples (each resets to `0.0`) + 7 retrofit entries for existing producers' Wiener states - Update existing-producer reset entries to use sentinel `0.0` (no formula bootstraps) -After Layer A: all 36 ISV slots populate correctly each step via Pearls A/B/D; cargo check + 36 unit tests + L40S smoke (regression check) all pass. +After Layer A: all 40 ISV slots populate correctly each step via Pearls A/B/D; cargo check + 40 unit tests + L40S smoke (regression check) all pass. ### Layer B — atomic consumer migration (the contract change) @@ -418,9 +456,17 @@ Single coordinated commit. Hardcoded multipliers, weight decay, L1 lambda — AL ### Layer C — validation + cleanup - L40S smoke validation (multi_fold_convergence::test_multi_fold_convergence) -- Remove now-dead `q_abs_ref_eff` and `h_s2_rms_ema_eff` parameters from the Mech 5 fused-kernel signature (if unused after Layer B) -- Update audit docs (`dqn-backward-nan-audit.md` mechanism table, slot allocation table; `dqn-wire-up-audit.md` Invariant 7 entries) -- Update `MEMORY.md` and add memory entry capturing SP4 close-out +- Remove now-dead `q_abs_ref_eff` and `h_s2_rms_ema_eff` parameters from the Mech 5 fused-kernel signature (unused after Layer B's per-slot ISV reads) +- **Retire `grad_norm_slow_ema` infrastructure:** since Mech 6's adaptive_clip upper-bound now reads `ISV[GRAD_CLIP_BOUND]` instead of `grad_norm_slow_ema_pinned × 100 × isv[16]`, the existing `grad_norm_slow_ema_pinned` field, its allocation/free, its launch site, and the `update_adaptive_clip` slow-EMA branch all become orphan. Remove them. The `grad_norm_ema` (fast EMA, used elsewhere) stays. +- Decide retention of `ISV[Q_ABS_REF=16]` and `ISV[H_S2_RMS_EMA=96]` producers: after Layer B, no clamp consumer reads these. They may still be useful for HEALTH_DIAG monitoring and `mag_concat_qdir`'s adaptive scale (P4.T2c.3c.6 consumer). **Decision:** keep producers running, drop only the orphaned Mech 1/2/5/6/9/10 consumer reads. Document as "monitoring-only ISV" in the audit. +- Update audit docs (`dqn-backward-nan-audit.md` mechanism table, slot allocation table; `dqn-wire-up-audit.md` Invariant 7 entries — one big entry covering 40 new slots + 15 producers + Pearl C engagement-counter buffer + Wiener-state buffer + retired components) +- Update `MEMORY.md` and add memory entries: + - SP4 close-out (project entry — like SP3) + - `pearl_first_observation_bootstrap` + - `pearl_fused_per_group_statistics_oracle` + - `pearl_engagement_rate_self_correction` + - `pearl_wiener_optimal_adaptive_ema` + - `pearl_signal_driven_regularisation_strength` ## Validation @@ -430,14 +476,16 @@ Single coordinated commit. Hardcoded multipliers, weight decay, L1 lambda — AL 2. **F1 Best Sharpe > 0** — fixes the original SP3 motivation; no h_s2-overflow cascade 3. **F0 Best Sharpe ≥ 37.5** — *revised criterion* matches `smoke-test-76pnm`'s 1e30-effectively-unclamped control. The 45.47 v2 baseline may be unreachable due to launch-scheduling-shift (the diagnostic smoke already showed F0=37.53 with no value-clamping at all). 37.5 is the realistic upper bound for any Mech-10-class design. 4. **F2 Best Sharpe ≥ 55** — F2 reached 80.55 with Mech 10 100×, and 56.70 with v2; ≥55 is the conservative floor -5. **Producer unit tests all pass** with <5% histogram-quantization error (magnitude bounds), <2% relative error vs. analytical formula (weight decay, L1 lambda) -6. **No new ISV slots beyond the 36 in this design** -7. **No new kernels beyond per-signal producers + standard Adam-kernel arg additions** +5. **Producer unit tests all pass** with <5% linear-histogram-quantization error (magnitude bounds, with bin width = step_max / 256), <2% relative error vs. analytical formula (weight decay, L1 lambda) +6. **No new ISV slots beyond the 40 in this design** +7. **No new kernels beyond the 15 fused producers + Adam-kernel-internal Pearl C engagement counter** 8. **Each consumer site reads exactly one ISV slot** (no compound expressions like `K × isv[X]`) 9. **Mech 5 diagnostic slots fire only on regime shifts** — sticky-flag pattern -10. **L1 lambda starts at 0 and ramps up** — entropy-deficit pearl: log `ISV[L1_LAMBDA[trunk]]` per fold; expect 0 at fold-start, ramping up as gradient differentiates within ~hundreds of steps. If λ stays at 0 throughout a fold, gradient never differentiates → either D signal is degenerate or the network's gradients are uniformly distributed (separate investigation thread). +10. **L1 lambda starts low and ramps up** — entropy-deficit pearl: log `ISV[L1_LAMBDA[trunk]]` per fold. At fold-boundary reset, Pearl A sentinel-detection means producer step-0 writes `step_λ = (mean|g|/mean|w|) × D` directly. With network just reset, D ≈ 0 (gradient still uniform) → step_λ ≈ 0 → ISV[L1_LAMBDA] ≈ 0 at step 0. Subsequent steps: λ ramps as D grows. If λ stays near 0 throughout a fold, gradient never differentiates → either D signal is degenerate or the network's gradients are uniformly distributed (separate investigation thread). 11. **Weight decay tracks gradient pressure** — log `ISV[WD_RATE[group]]` per fold; expect non-zero, smooth, per-group-specific values. If WD_RATE is identically zero, gradient projection onto weights is zero (network not learning anything from those weights → separate concern). 12. **No `weight_decay` or `l1_lambda` config fields remain in `HyperParams`** — the contract is fully ISV-driven; removing the config field prevents accidental hardcoding regression. +13. **Pearl C engagement-rate self-correction fires when expected** — log per-step `engagement_rate[group]` for in-kernel-clamp signals. Expected: ≈ 0.01 (1%) under stable training; transient bumps during regime shift; Pearl C force-bump events should be rare (< once per fold) under correct operation. Frequent force-bumps signal a deeper issue (e.g., distribution shift faster than EMA tracking). +14. **No hardcoded α values in any producer** — grep the producer-kernel sources for hardcoded ema_alpha/α values; should find zero outside `α_meta` (the structural meta-EMA constant). Zero magnitude constants in clamps. Zero theoretical bootstrap formulas. ### Regression sentinels (logged, not pass/fail) @@ -453,7 +501,7 @@ The first-draft spec carried 8 limitations. Pearls A-D resolve most of them; the - ~~In-kernel clamp feedback loop accepted~~ — **Pearl C resolves** via engagement-rate self-correction. Force-bump on sustained engagement deviation from theoretical 1%. - ~~Smoke time-budget marginally validates Adam EMAs~~ — **Pearl D resolves** via adaptive α. Step 0 = first-observation replacement (Pearl A); subsequent steps use Wiener-optimal α. EMA converges within tens of steps regardless of fold length. -- ~~28/36 producer kernels is plumbing-heavy~~ — **Pearl B resolves**. 14 fused producers with conceptual per-group grouping. +- ~~28/36/40 producer kernels is plumbing-heavy~~ — **Pearl B resolves**. 15 fused producers with conceptual per-group grouping (1 target_q + 4 atom_pos + 8 fused-per-group oracle + 1 grad_norm + 1 h_s2 = 15). - ~~Stale-by-one-step bound at fold boundary~~ — **Pearl A resolves**. Sentinel = 0 detected by producer → first-observation replaces immediately. No theoretical bootstrap to lag from. - ~~Histogram log-spacing 256-bin choice~~ — still a numerical-precision choice (0.4% quantile precision). Same Adam-ε category. Documented but not iterated. - ~~Carved-out items (weight decay, L1 lambda)~~ — **resolved in scope expansion**. Both fully signal-driven via gradient-pressure-projection (WD) and gradient-direction-entropy-deficit (L1). @@ -470,11 +518,11 @@ The first-draft spec carried 8 limitations. Pearls A-D resolve most of them; the 5. **Pearl C requires engagement-counter buffer per param-group.** Adds ~7 floats of mapped-pinned scratch space (one count per param-group). Per-step cost: per-thread shared-memory increment + block-wide reduce on existing nan-flag-buf path. Negligible. -6. **Pearl D's adaptive α adds 3 floats of state per producer.** sample_var, diff_var, x[t-1] lag. 14 producers × 3 floats = ~170 bytes additional ISV-state. Negligible. +6. **Pearl D's adaptive α adds 3 floats of state per ISV bound slot.** sample_var, diff_var, x_lag. 47 slots × 3 floats × 4 bytes ≈ 564 bytes mapped-pinned (`wiener_state_buf`). Negligible. ## Estimated effort -- **Layer A:** 14 fused producer kernels (per Pearl B) + 36 unit tests (one per output ISV slot) + Pearls A/D adaptive-EMA scaffolding + Pearl C engagement-rate self-correction + retrofit of 7 existing ISV producers + ISV plumbing + reset entries + 14 launch sites ≈ 3500-4500 LOC, ~1.5-2 weeks of focused work. Per-kernel is more complex (multi-pass + adaptive α) but kernel count is 60% lower than first draft. Net effort comparable. +- **Layer A:** 15 fused producer kernels (per Pearl B) + 40 unit tests (one per output ISV slot) + Pearls A/D adaptive-EMA scaffolding + Pearl C engagement-rate self-correction (per-Adam-kernel block-reduce + host-side rate-deficit EMA + force-bump trigger) + retrofit of 7 existing ISV producers + Wiener-state buffer (141 floats) + Pearl C engagement-counter buffer (8 × MAX_BLOCKS_PER_ADAM ints) + ISV plumbing + reset entries + 15 launch sites + retire `grad_norm_slow_ema` ≈ 3500-4500 LOC, ~1.5-2 weeks of focused work. Per-kernel is more complex (multi-pass + adaptive α) but kernel count is 60% lower than first draft. Net effort comparable. - **Layer B:** ~15 consumer sites × ~10 LOC change each ≈ 200 LOC, one focused day - **Layer C:** smoke + audit doc + memory entries ≈ 100 LOC + validation runs