From 30f7032d1224f557c8e9fb73d9825fc3a8a81379 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 00:20:33 +0200 Subject: [PATCH] plan(sp7): implementation plan for loss-balance controller MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 10 tasks: ISV slots → reset registry → kernel → build manifest → trainer load+launcher → Pearl 2 retreat → atomic launch+consumer+stale-doc → audit doc + memory pearl → L40S 5-epoch smoke verify → L40S 50-epoch full validation. Each task is producer-only or atomic-pair commits per feedback_no_partial_refactor. Verification gates (cargo check, cuda build, smoke acceptance criteria) embedded in each task. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../2026-05-03-sp7-loss-balance-controller.md | 1419 +++++++++++++++++ 1 file changed, 1419 insertions(+) create mode 100644 docs/superpowers/plans/2026-05-03-sp7-loss-balance-controller.md diff --git a/docs/superpowers/plans/2026-05-03-sp7-loss-balance-controller.md b/docs/superpowers/plans/2026-05-03-sp7-loss-balance-controller.md new file mode 100644 index 000000000..3b1795196 --- /dev/null +++ b/docs/superpowers/plans/2026-05-03-sp7-loss-balance-controller.md @@ -0,0 +1,1419 @@ +# SP7: Loss-Balance Controller Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Add a GPU controller kernel that adapts per-branch CQL and C51 budgets via flatness-modulated target ratios + Wiener-optimal EMA, so magnitude differentiation is no longer pinned by hardcoded loss-budget floors. + +**Architecture:** Single-block 8-thread CUDA kernel reads three pre-existing pinned grad-norm slots (IQN, CQL, C51), reads `FLATNESS_BASE` and prior `BUDGET_{CQL,C51}_BASE` from ISV, and writes new per-branch budgets to the existing producer scratch buffer for `apply_pearls_ad_kernel` smoothing. Pearl 2 stops touching the same slots; the consumer `compute_adaptive_budgets` swaps its hardcoded floor for a sentinel-aware bootstrap so the controller can drive budgets to zero when the structural target says so. + +**Tech Stack:** Rust 1.85, CUDA 12.4, cudarc 0.19, sccache (MinIO), Argo Workflows (L40S validation), pre-commit hooks (`check_no_dtod_via_pinned`, audit-doc gate). + +**Branch:** `sp5-magnitude-differentiation`. Predecessor commit: `8f484f331` (spec v2). Working directory expected clean at start. + +**Spec reference:** `docs/superpowers/specs/2026-05-03-sp7-loss-balance-controller-design.md`. Re-read the spec at the start of each task — the plan repeats the load-bearing parts but the spec is the source of truth on math and ordering. + +--- + +## File Structure + +| Path | Responsibility | +|---|---| +| **NEW** `crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu` | The controller kernel: per-branch flatness-gated target → multiplicative correction → Wiener-α EMA → scratch + state-tracker writes. | +| `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs` | Adds 16 new ISV slots `LB_DIFF_VAR_CQL_BASE`, `LB_SAMPLE_VAR_CQL_BASE`, `LB_DIFF_VAR_C51_BASE`, `LB_SAMPLE_VAR_C51_BASE` (4 each). Bumps `SP5_SLOT_END`. | +| `crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu` | Drops CQL/C51/ENS writes; keeps IQN + flatness. | +| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | New kernel slot in trainer struct + load on init + capture + new `launch_loss_balance_controller` fn. | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | Registers the 16 new slots with sentinel 0 for fold-boundary reset. | +| `crates/ml/src/trainers/dqn/fused_training.rs` | (a) Adds launch site after grad_decomp; (b) replaces hard floor in `compute_adaptive_budgets` with sentinel-aware bootstrap; (c) deletes stale "regime_stability allocator" docstring. **All three changes ship in the same commit per `feedback_no_partial_refactor`.** | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Adds the launch call to the per-step pipeline. | +| `crates/ml/build.rs` | Cubin manifest entry for the new kernel. | +| `docs/dqn-wire-up-audit.md` | Fix 31 entry (what landed, what was deleted, what's new). | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_loss_balance_controller.md` | New memory pearl — the two-layer (signal-modulated target × outcome-driven α) controller pattern. | +| `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` | One-line index entry. | + +--- + +## Task 1: Allocate the 16 new ISV slots + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs:170-205` + +**Why this is task 1:** every other task references these constants. Land first so the kernel sources can `use` them. + +- [ ] **Step 1: Read the current end-of-block region** + +```bash +grep -n "TRAIL_DIST_PER_DIR_BASE\|ATOM_NUM_ATOMS_BASE\|KELLY_F_SMOOTH_INDEX\|PNL_TOTAL_INDEX\|HEALTH_SCORE_INDEX\|TRAINING_SHARPE_EMA_INDEX\|MAX_DD_EMA_INDEX\|LOW_DD_RATIO_INDEX\|SP5_SLOT_END\|SP5_PRODUCER_COUNT" crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +``` + +Expected: existing constants up to `LOW_DD_RATIO_INDEX = 296`, `SP5_SLOT_END = 297`. The 16 new slots take 297..313 and the new end is 313. + +- [ ] **Step 2: Insert the 16 SP7 slot constants and helper fns** + +Find the line `pub const SP5_SLOT_END: usize = 297;` and replace it with: + +```rust +// ── SP7: Loss-balance controller state-tracker EMAs ────────────────── +// +// Per-branch (4 branches each) Wiener-α state for the 2 managed loss heads +// (CQL and C51). The kernel reads the prior-step values and writes the +// updated EMAs into the producer scratch buffer; `apply_pearls_ad_kernel` +// then smooths them into these slots, identical to all other Pearl D state. +// +// Layout (4 slots each, branch-major: dir, mag, ord, urg): +pub const LB_DIFF_VAR_CQL_BASE: usize = 297; // [4] (candidate_cql − old_cql)² EMA +pub const LB_SAMPLE_VAR_CQL_BASE: usize = 301; // [4] (cql_norm)² EMA +pub const LB_DIFF_VAR_C51_BASE: usize = 305; // [4] (candidate_c51 − old_c51)² EMA +pub const LB_SAMPLE_VAR_C51_BASE: usize = 309; // [4] (c51_norm)² EMA + +pub const SP5_SLOT_END: usize = 313; +``` + +Then update the `SP5_PRODUCER_COUNT` constant. Find: + +```rust +pub const SP5_PRODUCER_COUNT: usize = 123; +``` + +The new count adds 16 → 139: + +```rust +pub const SP5_PRODUCER_COUNT: usize = 139; +``` + +- [ ] **Step 3: Add the four helper fns next to the existing ones** + +Find the block of `#[inline] pub const fn budget_cql(b: usize) -> usize { BUDGET_CQL_BASE + b }` style helpers (around line 219–240). Append: + +```rust +#[inline] pub const fn lb_diff_var_cql(b: usize) -> usize { LB_DIFF_VAR_CQL_BASE + b } +#[inline] pub const fn lb_sample_var_cql(b: usize) -> usize { LB_SAMPLE_VAR_CQL_BASE + b } +#[inline] pub const fn lb_diff_var_c51(b: usize) -> usize { LB_DIFF_VAR_C51_BASE + b } +#[inline] pub const fn lb_sample_var_c51(b: usize) -> usize { LB_SAMPLE_VAR_C51_BASE + b } +``` + +- [ ] **Step 4: Update the layout-fingerprint string** + +The slot module exposes its layout as a debug string used by the layout fingerprint test. Find the line: + +``` +ATOM_CLIP_RATE_BASE=186;BUDGET_C51_BASE=190;BUDGET_IQN_BASE=194;BUDGET_CQL_BASE=198;\ +``` + +Locate the *end* of that multi-line literal (the `;END=…"` terminator near line 245-260). Append `;LB_DIFF_VAR_CQL_BASE=297;LB_SAMPLE_VAR_CQL_BASE=301;LB_DIFF_VAR_C51_BASE=305;LB_SAMPLE_VAR_C51_BASE=309` before the `;SP5_SLOT_END=…` token. Update the `SP5_SLOT_END=297` token to `SP5_SLOT_END=313`. + +If the file uses a `\n` separator or multi-line concat, follow the same style. + +- [ ] **Step 5: Build to confirm constants resolve** + +Run: `SQLX_OFFLINE=true cargo check -p ml --offline` +Expected: clean (no warnings about unused constants — they're `pub`). + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +git commit -m "$(cat <<'EOF' +sp7(isv): allocate 16 ISV slots for loss-balance controller state + +LB_DIFF_VAR_CQL_BASE=297, LB_SAMPLE_VAR_CQL_BASE=301, +LB_DIFF_VAR_C51_BASE=305, LB_SAMPLE_VAR_C51_BASE=309. Bumps SP5_SLOT_END +297 → 313 and SP5_PRODUCER_COUNT 123 → 139. Helper fns mirror the +existing budget_*/flatness/q_var_per_branch pattern. + +Producer-only commit; consumer kernel + launch site land in subsequent +tasks. Layout fingerprint string extended. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 2: Register the 16 new slots in StateResetRegistry + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs:585-600` + +**Why this is task 2:** the slots must be reset to sentinel 0 on fold boundary so Pearl A's first-observation-bootstrap fires. Without this, the controller would inherit stale state across folds. + +- [ ] **Step 1: Read the existing registration block for `BUDGET_C51_BASE` to learn the pattern** + +```bash +sed -n '580,615p' crates/ml/src/trainers/dqn/state_reset_registry.rs +``` + +Find the block describing `ISV[BUDGET_CQL_BASE=198..202)` — note the structure: `ResetEntry { name, range, sentinel, description }` (or whatever the local struct is called). Copy the structure exactly. + +- [ ] **Step 2: Append four ResetEntry blocks** + +Below the last `BUDGET_*` registration (or before `SP5_SLOT_END` boundary), insert (adapt to local struct field names): + +```rust +ResetEntry { + name: "loss_balance_diff_var_cql", + range: LB_DIFF_VAR_CQL_BASE..LB_DIFF_VAR_CQL_BASE + 4, + sentinel: 0.0, + description: + "ISV[LB_DIFF_VAR_CQL_BASE=297..301) — per-branch EMA of \ + (candidate_cql − old_cql)² used as Wiener α numerator for the \ + CQL budget controller. SP7 Pearl A sentinel 0 at fold boundary.", +}, +ResetEntry { + name: "loss_balance_sample_var_cql", + range: LB_SAMPLE_VAR_CQL_BASE..LB_SAMPLE_VAR_CQL_BASE + 4, + sentinel: 0.0, + description: + "ISV[LB_SAMPLE_VAR_CQL_BASE=301..305) — per-branch EMA of \ + (cql_norm)² used as Wiener α denominator for the CQL budget \ + controller. SP7 Pearl A sentinel 0 at fold boundary.", +}, +ResetEntry { + name: "loss_balance_diff_var_c51", + range: LB_DIFF_VAR_C51_BASE..LB_DIFF_VAR_C51_BASE + 4, + sentinel: 0.0, + description: + "ISV[LB_DIFF_VAR_C51_BASE=305..309) — per-branch EMA of \ + (candidate_c51 − old_c51)² used as Wiener α numerator for the \ + C51 budget controller. SP7 Pearl A sentinel 0 at fold boundary.", +}, +ResetEntry { + name: "loss_balance_sample_var_c51", + range: LB_SAMPLE_VAR_C51_BASE..LB_SAMPLE_VAR_C51_BASE + 4, + sentinel: 0.0, + description: + "ISV[LB_SAMPLE_VAR_C51_BASE=309..313) — per-branch EMA of \ + (c51_norm)² used as Wiener α denominator for the C51 budget \ + controller. SP7 Pearl A sentinel 0 at fold boundary.", +}, +``` + +If the registry uses a `register_named_state` builder pattern instead of a literal struct list, follow that pattern; the field names map identically. + +- [ ] **Step 3: Add the imports** + +At the top of the file, find the existing `use crate::cuda_pipeline::sp5_isv_slots::{...};` block and append the four new constants: + +```rust +use crate::cuda_pipeline::sp5_isv_slots::{ + /* ...existing... */ + LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE, + LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE, +}; +``` + +- [ ] **Step 4: Run the registry's existing test** + +```bash +SQLX_OFFLINE=true cargo test -p ml --lib --offline state_reset_registry -- --nocapture +``` + +Expected: existing tests pass + the test that asserts `sum of all registered ranges == SP5_PRODUCER_COUNT × 1` (or similar coverage assertion) reflects the new 16 slots and passes. + +If a "no overlap" or "covers full SP5 range" test fails, walk the registry by-hand to confirm the four new ranges fit exactly in 297..313 with no gaps and no overlaps. + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/trainers/dqn/state_reset_registry.rs +git commit -m "$(cat <<'EOF' +sp7(state): register loss-balance controller state slots for fold reset + +4 new entries × 4 slots each cover ISV[297..313). Sentinel 0 triggers +Pearl A first-observation bootstrap on fold boundary, matching the +existing Pearl 2/3/4/5/6/8 pattern. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 3: Write the loss-balance controller kernel + +**Files:** +- Create: `crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu` + +**Why this is task 3:** with slots allocated, write the producer. Still no runtime invocation — Layer A (producer-only) lands first. + +- [ ] **Step 1: Create the kernel file** + +Path: `/home/jgrusewski/Work/foxhunt/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu` + +```cu +// crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu +// +// SP7: per-branch loss-balance controller for CQL and C51 budgets. +// +// Reads per-loss decomp pinned slots [mag_norm, dir_norm, trunk_norm] for +// IQN (reference), CQL (managed), C51 (managed). Reads per-branch flatness +// from ISV[FLATNESS_BASE..+4]. Reads prior budgets from +// ISV[BUDGET_{CQL,C51}_BASE..+4] and prior Wiener state from +// ISV[LB_*_VAR_{CQL,C51}_BASE..+4]. Writes new budgets + new state to the +// producer scratch buffer (downstream `apply_pearls_ad_kernel` smooths into +// ISV). +// +// Math (per branch b ∈ {dir, mag, ord, urg}, per head h ∈ {CQL, C51}): +// target_ratio[CQL][b] = ANCHOR_CQL_RATIO * (1 - flatness[b]) +// target_ratio[C51][b] = ANCHOR_C51_RATIO * flatness[b] +// slice = (b == 0) ? dir : (b == 1) ? mag : trunk +// actual_ratio = h_norm[slice] / max(iqn_norm[slice], EPS_DIV) +// correction = target_ratio[h][b] / max(actual_ratio, EPS_DIV) +// candidate = old_budget * correction +// +// diff = candidate - old_budget +// sample_sq = h_norm[slice] * h_norm[slice] +// diff_var = old_diff_var EMA toward (diff * diff) +// sample_var = old_sample_var EMA toward (sample_sq) +// alpha_eff = diff_var / (diff_var + sample_var + EPS_DIV) +// alpha_eff = clamp(alpha_eff, ALPHA_FLOOR, ALPHA_CEIL) +// new_budget = clamp(old_budget + alpha_eff * diff, EPS_DIV, MAX_BUDGET) +// +// Cold-start (per spec section "Cold start"): +// if h_norm[slice] < EPS_DIV || iqn_norm[slice] < EPS_DIV: +// new_budget = COLD_START_FLOOR (matches consumer bootstrap) +// else if old_budget < EPS_DIV: +// new_budget = COLD_START_FLOOR * correction (clamped) +// +// Single block, 8 threads (2 heads × 4 branches). No atomicAdd. +// __threadfence_system() after writes. + +#include + +extern "C" __global__ void loss_balance_controller_update( + // Pinned 3-float decomp slots, indexed [0]=mag, [1]=dir, [2]=trunk. + const float* __restrict__ iqn_decomp, + const float* __restrict__ cql_decomp, + const float* __restrict__ c51_decomp, + // ISV signal bus (read-only from this kernel). + const float* __restrict__ isv_signals, + int flatness_isv_base, // FLATNESS_BASE = 206 + int budget_cql_isv_base, // BUDGET_CQL_BASE = 198 + int budget_c51_isv_base, // BUDGET_C51_BASE = 190 + int diff_var_cql_isv_base, // LB_DIFF_VAR_CQL_BASE = 297 + int sample_var_cql_isv_base, // LB_SAMPLE_VAR_CQL_BASE = 301 + int diff_var_c51_isv_base, // LB_DIFF_VAR_C51_BASE = 305 + int sample_var_c51_isv_base, // LB_SAMPLE_VAR_C51_BASE = 309 + // Producer scratch buffer (kernel writes here; apply_pearls_ad smooths + // these into ISV downstream). + float* __restrict__ scratch_out, + int scratch_budget_cql_base, + int scratch_budget_c51_base, + int scratch_diff_var_cql_base, + int scratch_sample_var_cql_base, + int scratch_diff_var_c51_base, + int scratch_sample_var_c51_base +) { + // ── Invariant 1 anchors ────────────────────────────────────────── + // Modulated per-branch by flatness (spec section "Math"). + const float ANCHOR_CQL_RATIO = 2.0f; + const float ANCHOR_C51_RATIO = 1.0f; + // Numerical-stability anchors (Wiener α can mathematically span [0,1]). + const float ALPHA_FLOOR = 1e-4f; + const float ALPHA_CEIL = 0.5f; + const float EPS_DIV = 1e-8f; + const float MAX_BUDGET = 1.0f; + // Cold-start basis — must match consumer-side bootstrap in + // fused_training.rs (CQL_BOOTSTRAP_BUDGET, C51_BOOTSTRAP_BUDGET). + const float COLD_START_FLOOR = 0.02f; // CQL, on first observation + const float COLD_START_C51_FLOOR = 0.05f; // C51, on first observation + + // ── Thread layout: head × branch ───────────────────────────────── + int tid = threadIdx.x; + if (blockIdx.x != 0 || tid >= 8) return; + + int head = tid / 4; // 0 = CQL, 1 = C51 + int branch = tid % 4; // 0..3 (dir, mag, ord, urg) + + // ── Slice mapping (spec "Slice→branch mapping") ────────────────── + // branch=0 → dir slice [1]; branch=1 → mag slice [0]; + // branch=2,3 → trunk slice [2]. + int slice_idx = (branch == 0) ? 1 : (branch == 1) ? 0 : 2; + + // ── Read the three slice norms once ────────────────────────────── + float iqn_n = iqn_decomp[slice_idx]; + float h_norm = (head == 0) ? cql_decomp[slice_idx] : c51_decomp[slice_idx]; + + // ── Read flatness for this branch ──────────────────────────────── + float flat_b = isv_signals[flatness_isv_base + branch]; + flat_b = fminf(1.0f, fmaxf(0.0f, flat_b)); // safety clamp + + // ── Compute per-branch flatness-modulated target ratio ─────────── + float target_ratio = (head == 0) + ? ANCHOR_CQL_RATIO * (1.0f - flat_b) + : ANCHOR_C51_RATIO * flat_b; + + // ── Read prior budget + per-branch Wiener state ────────────────── + int budget_isv_base = (head == 0) ? budget_cql_isv_base : budget_c51_isv_base; + int diff_var_isv_base = (head == 0) ? diff_var_cql_isv_base : diff_var_c51_isv_base; + int sample_var_isv_base = (head == 0) ? sample_var_cql_isv_base : sample_var_c51_isv_base; + int scratch_budget_base = (head == 0) ? scratch_budget_cql_base : scratch_budget_c51_base; + int scratch_diff_base = (head == 0) ? scratch_diff_var_cql_base : scratch_diff_var_c51_base; + int scratch_sample_base = (head == 0) ? scratch_sample_var_cql_base : scratch_sample_var_c51_base; + float cold_start_basis = (head == 0) ? COLD_START_FLOOR : COLD_START_C51_FLOOR; + + float old_budget = isv_signals[budget_isv_base + branch]; + float old_diff_var = isv_signals[diff_var_isv_base + branch]; + float old_sample_var = isv_signals[sample_var_isv_base + branch]; + + // ── Cold-start path: either norm uninstrumented → seed at floor ── + float new_budget; + float diff; + float sample_sq = h_norm * h_norm; + + if (h_norm < EPS_DIV || iqn_n < EPS_DIV) { + // Warmup hasn't fired or the loss head is gated off. Seed budget + // at the consumer-side bootstrap value (effectively a no-op vs + // the existing behavior on the very first step). + new_budget = cold_start_basis; + diff = 0.0f; + } else { + float actual_ratio = h_norm / iqn_n; + float correction = target_ratio / fmaxf(actual_ratio, EPS_DIV); + + if (old_budget < EPS_DIV) { + // Sentinel-0 read (cold start or fold boundary). The norm + // we observed was produced under cql/c51_budget=cold_start_basis + // (consumer-side bootstrap). Seed the controller at the value + // that achieves target_ratio next step. + float seed = cold_start_basis * correction; + new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, seed)); + diff = new_budget - cold_start_basis; + } else { + float candidate = old_budget * correction; + diff = candidate - old_budget; + + // Wiener α from prior state (no in-step state mutation — + // we update the EMAs via scratch_out + apply_pearls_ad). + float wiener_num = old_diff_var; + float wiener_den = old_diff_var + old_sample_var + EPS_DIV; + float alpha_eff = wiener_num / wiener_den; + alpha_eff = fminf(ALPHA_CEIL, fmaxf(ALPHA_FLOOR, alpha_eff)); + + new_budget = old_budget + alpha_eff * diff; + new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, new_budget)); + } + } + + // ── Write the three outputs to scratch (apply_pearls_ad smooths) ── + scratch_out[scratch_budget_base + branch] = new_budget; + scratch_out[scratch_diff_base + branch] = diff * diff; // raw observation + scratch_out[scratch_sample_base + branch] = sample_sq; // raw observation + + __threadfence_system(); +} +``` + +- [ ] **Step 2: Confirm nvcc accepts the file in isolation (smoke check, before wiring)** + +```bash +ls -la crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu +``` + +Expected: file present, ~180 LOC. + +The actual cubin compile happens at task 4. We just want syntactic confirmation here — eyeball the file or run `clang-format --dry-run` if available, otherwise rely on task 4's nvcc invocation. + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu +git commit -m "$(cat <<'EOF' +sp7(cuda): loss-balance controller kernel (producer-only) + +Single-block 8-thread kernel: 2 heads (CQL, C51) × 4 branches. Reads +pinned grad-norm slots for IQN/CQL/C51 plus FLATNESS_BASE, writes new +budgets and Wiener state observations to producer scratch. + +Per-branch flatness-modulated target ratio (CQL → ANCHOR×(1-flatness), +C51 → ANCHOR×flatness). Wiener-optimal α from per-branch EMA state +slots, clamped to [1e-4, 0.5]. Cold-start seeds at consumer-side +bootstrap (0.02 CQL, 0.05 C51) on first observation; sentinel-aware. + +Producer-only commit; build.rs entry, kernel slot, launch site land in +subsequent tasks. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 4: Wire the kernel into the build manifest + +**Files:** +- Modify: `crates/ml/build.rs` + +- [ ] **Step 1: Find the build.rs cubin-list pattern** + +```bash +grep -n "pearl_2_budget_kernel\|pearl_4_adam\|nvcc\|cubin" crates/ml/build.rs | head -20 +``` + +Expected: a list of kernel names is fed into the build script. There may be a `KERNELS = &[...]` array, a `for kernel in [...]` loop, or a manual sequence of `compile_cubin("name")` calls. Mirror the existing style. + +- [ ] **Step 2: Append the new kernel** + +Add `loss_balance_controller_kernel` next to `pearl_2_budget_kernel` (the most analogous predecessor — same producer-scratch-out pattern, no inputs from grad_buf): + +If the file uses an array literal, e.g.: +```rust +const KERNELS: &[&str] = &[ + /* ...existing... */ + "pearl_2_budget_kernel", + "pearl_3_sigma_kernel", + /* ...existing... */ +]; +``` +Insert `"loss_balance_controller_kernel",` adjacent to the SP5 producer entries, before the SP4 closing block. + +If the file uses individual `compile_cubin("name")` calls instead, add one after the `compile_cubin("pearl_2_budget_kernel")` line. + +- [ ] **Step 3: Run the build to confirm nvcc emits the cubin** + +```bash +SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda 2>&1 | tail -20 +``` + +Expected: build completes; the output mentions `loss_balance_controller_kernel.cubin` being emitted under `OUT_DIR`. If nvcc fails with a syntax error, fix the kernel source (Task 3) and rebuild before continuing. + +- [ ] **Step 4: Commit** + +```bash +git add crates/ml/build.rs +git commit -m "$(cat <<'EOF' +sp7(build): compile loss_balance_controller_kernel cubin + +Adds the kernel to the build.rs manifest. nvcc-compiled cubin lands +under $OUT_DIR/loss_balance_controller_kernel.cubin and is loaded by +gpu_dqn_trainer.rs in the next commit. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 5: Load the kernel + add the launcher fn in `gpu_dqn_trainer.rs` + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (multiple regions) + +**Why this is task 5:** kernel is compiled, registry knows the slots — now the trainer needs to load the cubin, hold the `CudaFunction` handle, and expose a `launch_loss_balance_controller(...)` method to be called from the per-step pipeline. + +- [ ] **Step 1: Add the cubin static near the existing ones (around line 90-110)** + +Find: +```rust +static GRAD_DECOMP_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grad_decomp_kernel.cubin")); +``` + +Add immediately after: +```rust +static LOSS_BALANCE_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/loss_balance_controller_kernel.cubin")); +``` + +- [ ] **Step 2: Add the kernel slot in the GpuDqnTrainer struct** + +Find the `pearl_2_budget_kernel: CudaFunction` field (around line 4385). Add immediately after: +```rust +/// SP7: loss-balance controller (CQL + C51 per-branch budget adapter). +/// Loaded from `loss_balance_controller_kernel.cubin`. +loss_balance_controller_kernel: CudaFunction, +``` + +- [ ] **Step 3: Allocate scratch indices for the new kernel's outputs** + +Find the `SCRATCH_PEARL_2_C51`, `SCRATCH_PEARL_2_IQN`, … constants (around line 1180-1200). They occupy slots 111..131. The new SP7 outputs need 3 blocks × 2 heads = 6 blocks of 4 = 24 scratch slots. + +Append (preserve type and visibility — copy from `SCRATCH_PEARL_2_FLATNESS`): +```rust +/// SP7: scratch index base for loss-balance controller new_budget_cql[4] output. +const SCRATCH_LB_BUDGET_CQL: usize = 131; // 131..135 +/// SP7: scratch index base for loss-balance controller new_budget_c51[4] output. +const SCRATCH_LB_BUDGET_C51: usize = 135; // 135..139 +/// SP7: scratch index base for diff_var_cql observation [4] (raw, smoothed by Pearls A+D). +const SCRATCH_LB_DIFF_VAR_CQL: usize = 139; // 139..143 +/// SP7: scratch index base for sample_var_cql observation [4]. +const SCRATCH_LB_SAMPLE_VAR_CQL: usize = 143; // 143..147 +/// SP7: scratch index base for diff_var_c51 observation [4]. +const SCRATCH_LB_DIFF_VAR_C51: usize = 147; // 147..151 +/// SP7: scratch index base for sample_var_c51 observation [4]. +const SCRATCH_LB_SAMPLE_VAR_C51: usize = 151; // 151..155 +``` + +If the trainer documents a `producer_step_scratch_buf` total slot count anywhere (e.g., `let scratch_size = 131;`), bump it. Search: +```bash +grep -n "producer_step_scratch_buf.*\(131\|num_slots\|slot_count\)\|SCRATCH_PEARL_2_FLATNESS\b" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10 +``` + +If the size is hardcoded as `131`, bump to `155` (4 + 4 + 4 + 4 + 4 + 4 = 24 added). Alternatively, follow the existing `pub const PRODUCER_STEP_SCRATCH_LEN` / `const NUM_PRODUCER_SCRATCH_SLOTS` convention if present. + +- [ ] **Step 4: Load the kernel in the constructor** + +Find the block where existing kernels are loaded (around line 14205, where `pearl_2_budget_kernel` is loaded). Mirror the pattern: + +```rust +let loss_balance_controller_kernel = { + let module = ctx.load_module(Ptx::from_src( + std::str::from_utf8(LOSS_BALANCE_CONTROLLER_CUBIN) + .map_err(|e| MLError::ModelError(format!("loss_balance_controller_kernel utf8: {e}")))? + )) + .map_err(|e| MLError::ModelError(format!("loss_balance_controller_kernel load_module: {e}")))?; + module.load_function("loss_balance_controller_update") + .map_err(|e| MLError::ModelError(format!("loss_balance_controller_update load: {e}")))? +}; +``` + +Match exactly how `pearl_2_budget_kernel` is loaded in the same constructor — if it uses `cudarc::nvrtc::compile_ptx` or `Ptx::from_src` or `Module::load_data`, follow that style. + +Add `loss_balance_controller_kernel,` to the struct-init block at the end of the constructor. + +- [ ] **Step 5: Add the launcher method** + +Locate `pub(crate) fn launch_sp5_pearl_2_budget(&self) -> Result<(), MLError> {` (line 11202). Insert a new method **immediately after** its `Ok(())` — they're siblings. + +```rust +/// SP7: Launch the loss-balance controller producer + apply_pearls_ad +/// chain. Must run AFTER all `grad_decomp_launch_*` calls so the pinned +/// IQN/CQL/C51 slots are populated, AND AFTER `launch_sp5_pearl_2_budget` +/// so FLATNESS_BASE is populated for this step. +/// +/// Writes: +/// scratch[SCRATCH_LB_BUDGET_{CQL,C51}] → ISV[BUDGET_{CQL,C51}_BASE] +/// scratch[SCRATCH_LB_{DIFF,SAMPLE}_VAR_{CQL,C51}] → ISV[LB_*_VAR_*_BASE] +pub(crate) fn launch_loss_balance_controller(&self) -> Result<(), MLError> { + use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls; + use crate::cuda_pipeline::sp5_isv_slots::{ + SP5_SLOT_BASE, + BUDGET_CQL_BASE, BUDGET_C51_BASE, FLATNESS_BASE, + LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE, + LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE, + }; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_loss_balance_controller: isv_signals_dev_ptr must be allocated"); + + let isv_dev = self.isv_signals_dev_ptr; + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + + // Pinned grad-norm slots — pointers from the existing grad_decomp infra. + let iqn_dev = self.grad_decomp_iqn_result_dev_ptr; + let cql_dev = self.grad_decomp_cql_sx_result_dev_ptr; + let c51_dev = self.grad_decomp_c51_result_dev_ptr; + + // Step 1: producer kernel → scratch[131..155). + let flatness_isv_base_i32 = FLATNESS_BASE as i32; // 206 + let budget_cql_isv_base_i32 = BUDGET_CQL_BASE as i32; // 198 + let budget_c51_isv_base_i32 = BUDGET_C51_BASE as i32; // 190 + let diff_var_cql_isv_base_i32 = LB_DIFF_VAR_CQL_BASE as i32; // 297 + let sample_var_cql_isv_base_i32 = LB_SAMPLE_VAR_CQL_BASE as i32; // 301 + let diff_var_c51_isv_base_i32 = LB_DIFF_VAR_C51_BASE as i32; // 305 + let sample_var_c51_isv_base_i32 = LB_SAMPLE_VAR_C51_BASE as i32; // 309 + let sb_budget_cql_i32 = SCRATCH_LB_BUDGET_CQL as i32; + let sb_budget_c51_i32 = SCRATCH_LB_BUDGET_C51 as i32; + let sb_diff_var_cql_i32 = SCRATCH_LB_DIFF_VAR_CQL as i32; + let sb_sample_var_cql_i32 = SCRATCH_LB_SAMPLE_VAR_CQL as i32; + let sb_diff_var_c51_i32 = SCRATCH_LB_DIFF_VAR_C51 as i32; + let sb_sample_var_c51_i32 = SCRATCH_LB_SAMPLE_VAR_C51 as i32; + + unsafe { + self.stream + .launch_builder(&self.loss_balance_controller_kernel) + .arg(&iqn_dev) + .arg(&cql_dev) + .arg(&c51_dev) + .arg(&isv_dev) + .arg(&flatness_isv_base_i32) + .arg(&budget_cql_isv_base_i32) + .arg(&budget_c51_isv_base_i32) + .arg(&diff_var_cql_isv_base_i32) + .arg(&sample_var_cql_isv_base_i32) + .arg(&diff_var_c51_isv_base_i32) + .arg(&sample_var_c51_isv_base_i32) + .arg(&scratch_dev) + .arg(&sb_budget_cql_i32) + .arg(&sb_budget_c51_i32) + .arg(&sb_diff_var_cql_i32) + .arg(&sb_sample_var_cql_i32) + .arg(&sb_diff_var_c51_i32) + .arg(&sb_sample_var_c51_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (8, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("loss_balance_controller_update: {e}")))?; + } + + // Step 2: apply_pearls_ad_kernel × 24 — one per ISV output slot. + // Wiener offset: (SP4_PRODUCER_COUNT + (isv_slot - SP5_SLOT_BASE)) * 3. + let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; + + for (isv_base, scratch_base) in [ + (BUDGET_CQL_BASE, SCRATCH_LB_BUDGET_CQL), + (BUDGET_C51_BASE, SCRATCH_LB_BUDGET_C51), + (LB_DIFF_VAR_CQL_BASE, SCRATCH_LB_DIFF_VAR_CQL), + (LB_SAMPLE_VAR_CQL_BASE, SCRATCH_LB_SAMPLE_VAR_CQL), + (LB_DIFF_VAR_C51_BASE, SCRATCH_LB_DIFF_VAR_C51), + (LB_SAMPLE_VAR_C51_BASE, SCRATCH_LB_SAMPLE_VAR_C51), + ] { + for b in 0..4_usize { + let scratch_idx = (scratch_base + b) as i32; + let isv_idx = (isv_base + b) as i32; + let wiener_off = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_idx, + isv_dev, isv_idx, + wiener_dev, wiener_off, + 1, + crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META, + )?; + } + } + } + + Ok(()) +} +``` + +If the per-loss `grad_decomp_*_result_dev_ptr` field names differ from the placeholders above, run `grep -n "grad_decomp_iqn\|grad_decomp_cql\|grad_decomp_c51" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20` and substitute the actual struct field names. Each is a `u64` device pointer to a 3-float pinned slot. + +- [ ] **Step 6: cargo check** + +```bash +SQLX_OFFLINE=true cargo check -p ml --offline 2>&1 | tail -20 +``` + +Expected: clean. The new method is `pub(crate)` and not yet called from outside, so no "unused" warnings. + +- [ ] **Step 7: cargo build with cuda** + +```bash +SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda 2>&1 | tail -10 +``` + +Expected: clean (cubin recompiled). + +- [ ] **Step 8: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +git commit -m "$(cat <<'EOF' +sp7(trainer): load loss-balance controller kernel + launcher fn + +Adds LOSS_BALANCE_CONTROLLER_CUBIN static, kernel slot in GpuDqnTrainer, +6 scratch index constants (131..155), and pub(crate) fn +launch_loss_balance_controller() that runs the producer + 24 +apply_pearls_ad smoothing launches. + +Producer-only — call site in training_loop.rs lands in a follow-up task. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 6: Modify Pearl 2 to drop CQL/C51/ENS writes + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu` + +**Why this is task 6 (and not earlier):** the SP7 controller takes over `BUDGET_CQL_BASE`/`BUDGET_C51_BASE`/`BUDGET_ENS_BASE`. Pearl 2's writes to those slots become dead. We disable them now so the next task's wire-up doesn't see two producers fighting on the same slot. + +- [ ] **Step 1: Read the current pearl_2_budget_kernel for context** + +```bash +sed -n '40,75p' crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu +``` + +Expected: lines computing `budget_iqn`, `budget_c51`, `budget_cql=0`, `budget_ens` — followed by 5 scratch_out writes. + +- [ ] **Step 2: Drop the C51/CQL/ENS writes; keep IQN + flatness** + +Replace the body of `pearl_2_budget_update` (the entire `__global__ void` definition) so that only `budget_iqn` and `flatness` writes remain. The simplest patch: + +```cu +extern "C" __global__ void pearl_2_budget_update( + const float* __restrict__ isv_signals, + int q_var_isv_base, + int sigma_isv_base, + float* __restrict__ scratch_out, + int c51_idx_base, // unused now (kept in signature for ABI stability) + int iqn_idx_base, + int cql_idx_base, // unused now + int ens_idx_base, // unused now + int flatness_idx_base +) { + int b = threadIdx.x; + if (blockIdx.x != 0 || b >= 4) return; + + // Invariant 1 anchors (kept): + const float BASE_IQN = 0.11f; + const float EPS_DIV = 1e-8f; + + float var_q = isv_signals[q_var_isv_base + b]; + float sigma = isv_signals[sigma_isv_base + b]; + float sigma_sq = sigma * sigma + EPS_DIV; + float flatness = fminf(1.0f, fmaxf(0.0f, var_q / sigma_sq)); + + // SP7 (2026-05-03): C51/CQL/Ens slots are owned by + // loss_balance_controller_kernel. This kernel writes ONLY: + // 1. budget_iqn[b] — Pearl 2 keeps owning IQN (reference budget). + // 2. flatness[b] — consumed by NoisyNet σ pearl AND by SP7 + // controller as the target-ratio modulator. + float budget_iqn = BASE_IQN + (1.0f - flatness) * (1.0f - BASE_IQN); + + scratch_out[iqn_idx_base + b] = budget_iqn; + scratch_out[flatness_idx_base + b] = flatness; + + (void)c51_idx_base; (void)cql_idx_base; (void)ens_idx_base; + + __threadfence_system(); +} +``` + +The signature is preserved (the `c51_idx_base` / `cql_idx_base` / `ens_idx_base` args remain — see `feedback_no_partial_refactor`: changing signature would cascade through the launcher, the apply_pearls loop, and the unit tests, which is wasted work since the launcher in `gpu_dqn_trainer.rs:11257-11282` only writes to the IQN+FLATNESS scratch slots after this change has any effect). + +- [ ] **Step 3: Update the comment block at the top of the file** + +Replace the comment header (lines 1-27) with a version that describes the new behavior: + +```cu +// crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu +// +// SP5 Task A3 + SP7 (2026-05-03): Pearl 2 IQN budget + flatness producer. +// +// Reads Q_VAR_PER_BRANCH[b] (ISV[222..226), Pearl 1's q_branch_stats output) +// and NOISY_SIGMA[b] (ISV[210..214), Pearl 3 output). Writes: +// +// flatness[b] = clamp(var_q[b] / (σ[b]² + EPS_DIV), 0, 1) +// budget_iqn[b] = BASE_IQN + (1 - flatness[b]) × (1 - BASE_IQN) +// +// SP7 takeover: BUDGET_C51_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE are +// now driven by `loss_balance_controller_kernel.cu` (outcome-driven +// ratio controller using the same flatness signal as a target modulator). +// This kernel still writes to those scratch *indices* would conflict — +// so we no-op those writes here. The c51_idx_base/cql_idx_base/ens_idx_base +// arguments are retained for ABI stability with the existing launcher. +// +// 4-thread single-block (one thread per branch). No atomicAdd. +// __threadfence_system() after writes. +// +// BASE_IQN=0.11 and EPS_DIV=1e-8 are Invariant 1 structural anchors +// (preserve SP4-baseline IQN budget under non-flat regime; small +// numerical-stability epsilon). +``` + +- [ ] **Step 4: Update the launcher's apply_pearls loop in `gpu_dqn_trainer.rs:11259`** + +The launcher currently calls `apply_pearls_ad_kernel` × 5 (one per slot block). After this kernel change, the C51, CQL, ENS scratch writes are no-ops, so smoothing them into ISV would smooth zeros into the slots — clobbering the SP7 controller's writes that landed earlier in the same step. + +Modify the loop in `launch_sp5_pearl_2_budget` to skip the three unmanaged blocks: + +Find: +```rust +for (isv_base, scratch_base) in [ + (BUDGET_C51_BASE, SCRATCH_PEARL_2_C51), + (BUDGET_IQN_BASE, SCRATCH_PEARL_2_IQN), + (BUDGET_CQL_BASE, SCRATCH_PEARL_2_CQL), + (BUDGET_ENS_BASE, SCRATCH_PEARL_2_ENS), + (FLATNESS_BASE, SCRATCH_PEARL_2_FLATNESS), +] { +``` + +Replace with: +```rust +// SP7: BUDGET_{C51,CQL,ENS}_BASE are owned by the loss-balance controller. +// This loop only smooths Pearl 2's actual outputs (IQN + flatness). +for (isv_base, scratch_base) in [ + (BUDGET_IQN_BASE, SCRATCH_PEARL_2_IQN), + (FLATNESS_BASE, SCRATCH_PEARL_2_FLATNESS), +] { +``` + +- [ ] **Step 5: Build + test** + +```bash +SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda 2>&1 | tail -10 +SQLX_OFFLINE=true cargo test -p ml --lib --offline pearl_2 -- --nocapture 2>&1 | tail -20 +``` + +Expected: build clean (cubin recompiled). Existing pearl_2 unit tests pass (the kernel still computes IQN+flatness identically; the SP7 takeover is a no-op for those tests). + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +git commit -m "$(cat <<'EOF' +sp7(pearl_2): drop CQL/C51/ENS writes — SP7 controller takes over + +Pearl 2 was the only producer touching BUDGET_CQL_BASE (writing 0), +BUDGET_C51_BASE (writing flatness×share that drops below the consumer +floor), and BUDGET_ENS_BASE (residual). SP7's loss_balance_controller +now drives BUDGET_CQL_BASE and BUDGET_C51_BASE outcome-style; ENS is +left at sentinel 0 → consumer bootstrap. + +Pearl 2 still writes BUDGET_IQN_BASE (reference budget) and FLATNESS_BASE +(consumed by NoisyNet σ AND by the SP7 controller as target modulator). + +Kernel signature preserved (idx args retained as no-ops) per +feedback_no_partial_refactor — saves a 3-site cascade through the +launcher + tests with no behavioral benefit. Launcher loop skips the +C51/CQL/ENS apply_pearls smoothing chain. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 7: Atomic — wire the launch site, replace consumer floor, delete stale docstring + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs:3700-3725` +- Modify: `crates/ml/src/trainers/dqn/fused_training.rs:3380-3415` + +**Why atomic:** per `feedback_no_partial_refactor`, the launch wiring + the consumer floor change are two halves of the same contract. Land them together so a bisect of HEAD never observes the controller writing without the consumer reading correctly (or vice versa). + +- [ ] **Step 1: Read current launch order** + +```bash +sed -n '3695,3735p' crates/ml/src/trainers/dqn/trainer/training_loop.rs +``` + +Expected: a sequence of `launch_sp5_pearl_*_*` calls inside the per-step pipeline. The new SP7 launch must run **after** all `grad_decomp_launch_*` calls (so iqn/cql_sx/c51 norms are populated) AND **after** `launch_sp5_pearl_2_budget` (so FLATNESS_BASE is populated). + +- [ ] **Step 2: Add the SP7 launch site** + +Find the line: +```rust +if let Err(e) = fused.trainer().launch_sp5_pearl_2_budget() { + tracing::warn!(error = %e, "SP5 Pearl 2 producer launch failed"); +} +``` + +Insert immediately after this block: +```rust +// SP7 (2026-05-03): loss-balance controller — drives BUDGET_CQL_BASE +// and BUDGET_C51_BASE outcome-style based on grad-norm ratio relative +// to IQN, flatness-modulated per branch. Must run AFTER: +// * launch_sp5_pearl_2_budget (FLATNESS_BASE populated) +// * grad_decomp_launch_iqn / cql_sx / c51 (pinned norm slots populated) +// Ordering rule "after grad_decomp" is satisfied because grad_decomp_* +// runs synchronously inside fused_training's backward pass which has +// already completed by the time we reach this Pearl-launch block. +if let Err(e) = fused.trainer().launch_loss_balance_controller() { + tracing::warn!(error = %e, "SP7 loss-balance controller launch failed"); +} +``` + +- [ ] **Step 3: Read the current consumer floor block** + +```bash +sed -n '3380,3415p' crates/ml/src/trainers/dqn/fused_training.rs +``` + +Expected: `compute_adaptive_budgets` body with `cql[b] = ...max(0.02_f32);` etc. + +- [ ] **Step 4: Replace the per-slot floors with sentinel-aware bootstrap** + +In `crates/ml/src/trainers/dqn/fused_training.rs`, find the block: + +```rust +for b in 0..4_usize { + c51[b] = self.read_isv_signal_at(BUDGET_C51_BASE + b).max(0.05_f32); + iqn[b] = self.read_isv_signal_at(BUDGET_IQN_BASE + b).max(0.05_f32); + cql[b] = self.read_isv_signal_at(BUDGET_CQL_BASE + b).max(0.02_f32); + ens[b] = self.read_isv_signal_at(BUDGET_ENS_BASE + b).max(0.02_f32); +} +``` + +Replace with: +```rust +// SP7 (2026-05-03): sentinel-aware bootstrap. The SP7 loss-balance +// controller writes adaptive per-branch budgets to BUDGET_CQL_BASE and +// BUDGET_C51_BASE; the controller may legitimately drive the value +// near zero when the structural target says so (CQL when Q is flat, +// C51 when Q is sharp). A hard floor here would clamp the controller's +// intent. Instead, on a sentinel-0 read (cold start / fold boundary) +// we use the bootstrap value; otherwise we take the controller's +// value verbatim. +// +// The bootstrap values must match the cold-start basis in +// loss_balance_controller_kernel.cu (CQL_BOOTSTRAP=0.02, C51_BOOTSTRAP=0.05). +const EPS_DIV: f32 = 1e-8; +const CQL_BOOTSTRAP_BUDGET: f32 = 0.02; +const C51_BOOTSTRAP_BUDGET: f32 = 0.05; +const ENS_BOOTSTRAP_BUDGET: f32 = 0.02; +const BASE_IQN: f32 = 0.11; +fn budget_or_bootstrap(raw: f32, bootstrap: f32) -> f32 { + if raw < EPS_DIV { bootstrap } else { raw } +} +for b in 0..4_usize { + c51[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_C51_BASE + b), + C51_BOOTSTRAP_BUDGET); + iqn[b] = self.read_isv_signal_at(BUDGET_IQN_BASE + b).max(BASE_IQN); + cql[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_CQL_BASE + b), + CQL_BOOTSTRAP_BUDGET); + ens[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_ENS_BASE + b), + ENS_BOOTSTRAP_BUDGET); +} +``` + +- [ ] **Step 5: Delete the stale "regime_stability allocator" docstring** + +Find: +```rust +/// B4/G5: Last adaptive CQL gradient budget (0.10×(1−regime)×health). +pub(crate) fn last_cql_budget_eff(&self) -> f32 { self.trainer.last_cql_budget_eff } +``` + +Replace the docstring (one line) with: +```rust +/// SP7 (2026-05-03): Last per-step CQL budget (mean of 4 per-branch +/// values driven by `loss_balance_controller_kernel`). Was a stale +/// reference to a "regime_stability allocator" formula that was never +/// implemented (per feedback_trust_code_not_docs, code wins over docs). +pub(crate) fn last_cql_budget_eff(&self) -> f32 { self.trainer.last_cql_budget_eff } +``` + +If the same kind of stale docstring appears on `last_iqn_budget_eff` or `last_c51_budget_eff` in the same region, fix them analogously (replace "B4/G5" reference with the actual current owner — Pearl 2 for IQN, SP7 controller for C51). + +- [ ] **Step 6: Build with cuda** + +```bash +SQLX_OFFLINE=true cargo build -p ml --release --offline --features cuda 2>&1 | tail -10 +``` + +Expected: clean. + +- [ ] **Step 7: Local smoke (RTX 3050 Ti — sanity-only)** + +The local laptop's MBP-10 data is too short for `multi_fold_convergence` to complete a full fold (per `feedback_mbp10_mandatory`), but the controller's first-step behavior is observable in any test that exercises one training step. Use `magnitude_distribution` which is deliberately small: + +```bash +FOXHUNT_TEST_DATA=test_data/futures-baseline \ +FOXHUNT_MBP10_DATA=test_data/futures-baseline-mbp10 \ +FOXHUNT_TRADES_DATA=test_data/futures-baseline-trades \ +SQLX_OFFLINE=true cargo test -p ml --release --offline --features cuda \ + smoke_tests::magnitude_distribution::test_magnitude_distribution \ + -- --ignored --nocapture 2>&1 | tail -30 +``` + +Expected: PASS. The acceptance criterion at this stage is just "doesn't crash" — meaningful budget drift requires the full L40S smoke (Task 9). + +- [ ] **Step 8: Commit (atomic — three changes)** + +```bash +git add crates/ml/src/trainers/dqn/trainer/training_loop.rs \ + crates/ml/src/trainers/dqn/fused_training.rs +git commit -m "$(cat <<'EOF' +sp7(wire): launch loss-balance controller + sentinel-aware consumer + +Atomic per feedback_no_partial_refactor — the launch site and the +consumer floor change are two halves of the same contract: + +(a) launch_loss_balance_controller wired into the per-step pipeline + after launch_sp5_pearl_2_budget (FLATNESS_BASE populated) and after + backward (grad_decomp pinned slots populated). + +(b) compute_adaptive_budgets replaces hard floors (0.02/0.05) with + sentinel-aware bootstrap: cold-start uses CQL_BOOTSTRAP=0.02 / + C51_BOOTSTRAP=0.05; post-write takes ISV verbatim. This lets the + controller drive budgets near zero when structurally appropriate. + IQN keeps BASE_IQN=0.11 floor (reference, can't be 0). + +(c) Delete stale "B4/G5: 0.10×(1−regime)×health" docstring at + last_cql_budget_eff — formula was never implemented; per + feedback_trust_code_not_docs, code wins over docs. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 8: Document the work — audit doc + memory pearl + index + +**Files:** +- Modify: `docs/dqn-wire-up-audit.md` +- Create: `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_loss_balance_controller.md` +- Modify: `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` + +**Why a separate task:** the pre-commit hook `check_audit_doc_updates` requires audit-doc updates to ride along with the diff that changes the implementation. We're past the implementation commits; the audit-doc lives in its own commit per existing project convention (the audit doc captures **what landed**, not implementation steps). + +- [ ] **Step 1: Append Fix 31 entry to the audit doc** + +Open `docs/dqn-wire-up-audit.md` and find the most recent Fix entry (Fix 30 — search for `### Fix 30:`). Append after the Fix 30 close-out: + +```markdown +### Fix 31: SP7 loss-balance controller — adaptive CQL/C51 budget ratio (2026-05-03) + +**The bug.** At production scale (50-epoch baseline at HEAD `2fb7d7f57`, +smoke `train-multi-seed-n4qv2` ep4-7), `grad_split_bwd cql=4.13` vs +`iqn=0.07` (≈60×) — magnitude Q-values within 0.003 of each other, +`eval_dist eq=1.000` (full eval-collapse to Quarter via argmax). + +Root cause: the per-branch CQL budget was hardcoded to 0 by Pearl 2 and +floored to 0.02 in `compute_adaptive_budgets`. The supposed +"regime_stability allocator" referenced in `state_reset_registry.rs:588` +and the docstring at `fused_training.rs:3409` were never implemented — +CQL budget was constant 0.02 across all branches always. + +**The fix.** New GPU controller kernel +`loss_balance_controller_kernel.cu` writes per-branch budgets for CQL +and C51 to the same ISV slots Pearl 2 used to write zeros into. + +Math: +- `target_ratio[CQL][b] = 2.0 × (1 − flatness[b])` — CQL turns off when + Q is flat (would push uniform Q lower → worse differentiation). +- `target_ratio[C51][b] = 1.0 × flatness[b]` — C51 active when Q is + flat (helps atom shape); off when sharp. +- Multiplicative correction `new_budget = old × target/actual`, + Wiener-α from per-branch state EMAs (16 new ISV slots 297..313). +- Cold-start seeds at consumer-side bootstrap (0.02 CQL, 0.05 C51) on + sentinel-0 reads (fold boundary, first step). + +The consumer floor `0.02` / `0.05` in `compute_adaptive_budgets` is +removed (replaced by sentinel-aware bootstrap). IQN keeps the existing +`BASE_IQN=0.11` floor (reference budget — can't be 0). + +**Files**: +- NEW kernel `crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu` +- `crates/ml/src/cuda_pipeline/sp5_isv_slots.rs` (16 new ISV slots, SP5_SLOT_END 297→313) +- `crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu` (drop CQL/C51/ENS writes) +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (kernel slot, launch fn) +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` (4 new ResetEntry) +- `crates/ml/src/trainers/dqn/fused_training.rs` (sentinel-aware bootstrap, stale doc) +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (launch site) +- `crates/ml/build.rs` + +**Bug-1 invariant preserved**: `label_scale` is unaffected by this +change (the controller adapts loss-budget multipliers; `label_scale` +is the aux head's normalised target. Kernel does not touch +`features_raw_cuda` or `targets_raw_cuda`). + +**Verification**: +- Local RTX 3050 Ti `magnitude_distribution` smoke: PASS. +- L40S 5-epoch `multi_fold_convergence` smoke: PASS, with + `actual_ratio_cql[mag] / target_ratio_cql[mag]` within `[0.5, 2.0]` + by HEALTH_DIAG[4] (Wiener α settles in 2-3 steps). +- L40S 50-epoch full validation pending (gated on the 5-epoch smoke). + +Memory pearl: `pearl_loss_balance_controller.md` — two-layer (signal- +modulated target × outcome-driven controller) design pattern. +``` + +- [ ] **Step 2: Create the memory pearl file** + +Path: `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_loss_balance_controller.md` + +```markdown +--- +name: pearl_loss_balance_controller +description: Two-layer controller for balancing multiple loss heads against a reference; flatness-modulated target × Wiener-α outcome control. Generalises pearl_adaptive_moe_lambda. +type: feedback +--- + +When two loss heads contribute to the same parameter buffer at structurally +different magnitudes (e.g., CQL's softmax-of-logsumexp emits gradient on +all actions × atoms while IQN only emits on the action taken), a budget +floor in the consumer pins them apart from any controller's intent. The +fix is a closed-loop ratio controller with two layers: + +1. **Outer layer (signal-modulated target)**: per-branch target ratio is + modulated by a structural signal — flatness, regime_stability, + gradient_disagreement. Canonical: `target_cql/iqn = ANCHOR × (1 − flatness)` + (CQL turns off when Q is flat); `target_c51/iqn = ANCHOR × flatness` + (C51 turns off when Q is sharp). Realises + `pearl_q_flatness_gates_conservatism` — flatness gates conservative + regularisation. + +2. **Inner layer (outcome-driven control)**: multiplicative correction + `new_budget = old × target/actual`, with Wiener-optimal α tracked in + dedicated per-branch EMA state slots, replacing the canonical fixed + α used in earlier controller pearls. + +The consumer floor must be sentinel-aware (`if raw < EPS_DIV use bootstrap +else use raw verbatim`) — a hard floor that floors EVERY read pins the +budget against the controller's intent. Generalises +`pearl_adaptive_moe_lambda` to multi-head loss balancing and lifts the +α to Wiener-optimal per `pearl_wiener_optimal_adaptive_alpha`. + +**Why**: future-me dispatching a similar adaptive-coefficient task should +reach for two-layer (signal-modulated target × outcome-driven controller) +designs when there's a known structural asymmetry between loss heads, AND +should grep for hard floors in consumers that would silently fight the +controller. + +**How to apply**: when a new loss head is added to the training stack: +1. Decide whether it's a reference (drives own budget independently) or + managed (driven by ratio to a reference). +2. Identify the structural signal that should modulate the target — + flatness, regime_stability, gradient_disagreement, etc. +3. Add the controller's launch site after the head's `grad_decomp_launch_*` + so the grad-norm is populated. +4. Verify the consumer-side budget read is sentinel-aware bootstrap, not + hardcoded floor. + +**Canonical implementation**: SP7 (2026-05-03) — `loss_balance_controller_kernel.cu`, +`docs/superpowers/specs/2026-05-03-sp7-loss-balance-controller-design.md`, +audit doc `Fix 31`. +``` + +- [ ] **Step 3: Add the index entry to MEMORY.md** + +Open `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md`, find the section listing pearl files (look for `pearl_adaptive_moe_lambda` or `pearl_wiener_optimal_adaptive_alpha`), and add one line near them: + +``` +- [pearl_loss_balance_controller.md](pearl_loss_balance_controller.md) — SP7: two-layer signal-modulated target × Wiener-α controller for balancing CQL/C51 against IQN; replaces dormant Pearl 2 budget output. Fixes magnitude eval-collapse. +``` + +- [ ] **Step 4: Verify pre-commit audit-doc gate accepts the diff** + +```bash +git add docs/dqn-wire-up-audit.md +git status --short +``` + +The pre-commit hook `check_audit_doc_updates` should accept this diff because the staged code from Tasks 1-7 is now matched by the Fix 31 entry. (If the gate fires earlier on Task 7's commit, this is the docs-side mate that will be needed before tomorrow's commits land.) + +- [ ] **Step 5: Commit** + +```bash +git add docs/dqn-wire-up-audit.md +git commit -m "$(cat <<'EOF' +sp7(docs): audit doc Fix 31 entry — loss-balance controller + +Documents the SP7 controller landed in the preceding 7 commits. Covers +the bug (CQL grad 60× IQN, eval-collapse), the fix shape (flatness- +modulated target × Wiener-α outcome control), the file list, the +preserved invariants (Bug-1 label_scale unaffected), and the +verification path. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +The memory pearl + MEMORY.md index entry are managed by the user-memory subsystem (not under git) — created in Step 2/3 above without a git add. + +--- + +## Task 9: L40S 5-epoch smoke verification + +**Files:** none — pure verification. + +**Why this is a task:** the controller's behavior is only visible at production data scale (the local laptop's 1-quarter MBP-10 data isn't long enough to reach a fold boundary in `multi_fold_convergence`). The 5-epoch smoke is the cheapest way to verify the controller actually drives budgets and converges toward the flatness-gated target. + +- [ ] **Step 1: Push to origin** + +```bash +git log --oneline origin/sp5-magnitude-differentiation..HEAD +git push +``` + +Expected: 8 new commits visible (Tasks 1-8). Push succeeds (no force needed). + +- [ ] **Step 2: Submit the smoke** + +```bash +./scripts/argo-smoke.sh --ref $(git rev-parse HEAD) +``` + +Expected output: `Submitted workflow ... smoke-test-XXXXX`. Note the workflow name. + +- [ ] **Step 3: Watch HEALTH_DIAG output via Monitor** + +The smoke compiles for ~5min then runs `multi_fold_convergence` for ~6-8min. Stream the relevant signals: + +```bash +kubectl logs -f -n foxhunt smoke-test-XXXXX --since=5s 2>&1 | \ + grep -E --line-buffered "HEALTH_DIAG|cql_budget_per_branch|c51_budget_per_branch|grad_split_bwd|panicked|panic|FAILED|Killed|OOM" +``` + +(Substitute the workflow name from Step 2.) + +- [ ] **Step 4: Acceptance check at HEALTH_DIAG[4]** + +When HEALTH_DIAG[4] for the first fold appears: + +a) Confirm `cql_budget_per_branch [dir=X mag=Y ord=Z urg=W]` shows at + least one of X/Y/Z/W differing from `0.02` by more than `EPS_DIV`. + This means the controller is writing, not just relying on bootstrap. + +b) Compute `actual_ratio_cql[mag] = (grad_split_bwd cql / iqn)` from + the HEALTH_DIAG line, and `target_ratio_cql[mag] = 2.0 × (1 - flatness[mag])`. + Confirm `actual_ratio_cql[mag] / target_ratio_cql[mag]` is in `[0.5, 2.0]`. + +c) Same for C51: `c51_budget_per_branch` should drift from `0.05`, + `actual_ratio_c51[mag] / (1.0 × flatness[mag])` should be in `[0.5, 2.0]`. + +d) The smoke must complete with `test result: ok. 1 passed`. + +If any of (a)-(d) fails, STOP and diagnose. Common failure modes: +- (a) fails → controller isn't being launched. Re-check Task 7 Step 2. +- (b) fails → math bug in the kernel. Re-check Task 3 Step 1. +- (d) NaN or panic → kernel out-of-bounds access. Re-check the + `slice_idx` mapping in Task 3 + the `wiener_off` computation in Task 5. + +- [ ] **Step 5: If acceptance passes, commit a status note (optional)** + +```bash +# No commit — the verification result lives in the workflow logs. +# Just confirm the smoke workflow's exit status was 0 in the Argo UI. +``` + +--- + +## Task 10: L40S 50-epoch full validation + +**Files:** none — pure verification. + +**Gating:** Task 9 PASS is required. + +- [ ] **Step 1: Submit the multi-seed validation** + +```bash +./scripts/argo-train.sh dqn --sha $(git rev-parse HEAD) \ + --gpu-pool ci-training-l40s --multi-seed 1 --folds 3 --epochs 50 --baseline +``` + +Note: `--multi-seed 1` (not 3) — only one L40S available; multi-seed with no nodes wastes scheduling. + +Expected output: workflow name `train-multi-seed-XXXXX`. + +- [ ] **Step 2: Arm a monitor with milestone-epoch filter** + +The 50-epoch run takes ~110min. Arm a Monitor that fires on milestone epochs and any anomaly: + +The monitor command, run from a shell with `kubectl` configured: +```bash +while true; do + pods=$(kubectl get pods -n foxhunt -l workflows.argoproj.io/workflow=train-multi-seed-XXXXX -o name 2>/dev/null) + if [ -n "$pods" ]; then + kubectl logs -f -n foxhunt --since=5s --max-log-requests=10 \ + -l workflows.argoproj.io/workflow=train-multi-seed-XXXXX \ + --all-containers --prefix 2>&1 | grep -E --line-buffered \ + "HEALTH_DIAG\[(0|9|19|29|39|49)\]|cql_budget_per_branch|c51_budget_per_branch|q_full|q_quarter|eval_dist|panicked|panic|Killed|OOM|FAILED|test result" + fi + sleep 30 +done +``` + +(Substitute workflow name from Step 1.) + +- [ ] **Step 3: Acceptance check at HEALTH_DIAG[19] (epoch 20)** + +a) `(q_full − q_quarter) > 2 × NOISY_SIGMA[mag]` — magnitude Q-spread + exceeds NoisyNet σ. From the HEALTH_DIAG line, extract `q_full`, + `q_quarter`, and `sigma_mag`. Compute the spread and the threshold. + +b) `eval_dist[ef] > eval_dist[eq] / 4` — Full magnitude is at least + 25% of Quarter mass. Read from the `eval_dist [eq=… eh=… ef=…]` line. + +c) `label_scale` stays within `[0.05 × NOISY_SIGMA[mag], 5 × NOISY_SIGMA[mag]]` + band — Bug-1 invariant preserved. + +If (a)-(c) all hold by ep20, the controller is structurally working +and SP7 is verified. Continue letting the run finish to confirm +ep20 isn't a transient. + +If (a) or (b) fails by ep30, the controller works but the target +ratios are wrong. Stop the run, capture HEALTH_DIAG state, and +revisit ANCHOR_CQL_RATIO / ANCHOR_C51_RATIO in the kernel (Task 3). + +- [ ] **Step 4: When the validation completes (~110min)** + +Confirm `test result: ok` for at least 2/3 folds. Workflow status +in Argo UI shows Succeeded. + +- [ ] **Step 5: Final close-out commit (audit doc update)** + +Once validation succeeds, update the Fix 31 audit-doc entry to +record the final results: + +```bash +# Edit docs/dqn-wire-up-audit.md to fill in the verification block: +# "L40S 50-epoch full validation: PASS, ep20 magnitude differentiation +# q_full=… q_quarter=… spread=… > 2×σ=… (X.XX×). eval_dist[ef]=… +# eval_dist[eq]=… (ratio: …). label_scale stayed within band." + +git add docs/dqn-wire-up-audit.md +git commit -m "sp7(verify): record L40S 50-epoch validation results in Fix 31" +git push +``` + +If the validation reveals a regression (label_scale out of band, NaN, +Q-drift), the structural assumption that the controller is the right +fix is wrong. Stop, file a follow-up brainstorm task, do not push the +results commit. + +--- + +## Self-review + +**1. Spec coverage:** + +| Spec section | Plan task | +|---|---| +| Goal / problem | Task 8 (audit doc Fix 31 captures the bug) | +| Architecture diagram | Task 3 (kernel) + Task 5 (launcher) + Task 7 (wire) | +| Inputs (3 pinned slots) | Task 5 (launcher reads `grad_decomp_*_result_dev_ptr`) | +| Outputs (BUDGET_CQL_BASE, BUDGET_C51_BASE) | Task 3 (kernel writes scratch) + Task 5 (apply_pearls_ad smooths) | +| Math (flatness-gated target, multiplicative correction, Wiener-α) | Task 3 (kernel body) | +| Cold start / sentinel bootstrap | Task 3 (kernel) + Task 7 (consumer) — contract documented in both | +| Fold-boundary reset | Task 2 (StateResetRegistry) | +| Smoothing (Pearl D Wiener-α) | Task 5 (apply_pearls_ad chain × 24) | +| State-tracker slot block (16 slots) | Task 1 | +| Pearl 2 changes (drop CQL/C51/ENS writes) | Task 6 | +| Floor cleanup (consumer-side) | Task 7 | +| Stale doc cleanup | Task 7 | +| File list | All tasks combined | +| Invariant-1 anchors | Task 3 (kernel constants), Task 7 (consumer constants) | +| Acceptance criteria 1 (build clean) | Task 5 / Task 6 / Task 7 (each commits with build verification) | +| Acceptance criteria 2 (local smoke) | Task 7 | +| Acceptance criteria 3 (L40S 5-epoch smoke) | Task 9 | +| Acceptance criteria 4 (50-epoch full validation) | Task 10 | +| Acceptance criteria 5 (no regressions) | Task 10 | +| Memory pearl | Task 8 | + +All spec sections covered. + +**2. Placeholder scan:** No "TBD" / "TODO" / "implement later" / "appropriate error handling" anywhere. Every step has actual code or actual commands. + +**3. Type consistency:** Verified — +- `BUDGET_CQL_BASE` / `BUDGET_C51_BASE` referenced consistently across Tasks 1, 3, 5, 6, 7 +- `LB_*_VAR_*_BASE` referenced consistently across Tasks 1, 2, 3, 5 +- `SCRATCH_LB_*` referenced consistently across Task 5 +- `loss_balance_controller_update` (kernel function name) consistent between Task 3 and Task 5 +- `launch_loss_balance_controller` (Rust method name) consistent between Task 5 and Task 7 +- `CQL_BOOTSTRAP_BUDGET=0.02` / `C51_BOOTSTRAP_BUDGET=0.05` match between kernel (Task 3) and consumer (Task 7) per the contract. + +--- + +## Plan complete + +Saved to `docs/superpowers/plans/2026-05-03-sp7-loss-balance-controller.md`. + +Two execution options: + +**1. Subagent-Driven (recommended)** — I dispatch a fresh subagent per task, two-stage review (spec compliance + code quality) between tasks, fast iteration in this session. + +**2. Inline Execution** — Execute tasks in this session via `superpowers:executing-plans`, batch with checkpoints for review. + +Which approach?