From 552d91bf45dc11eed4da1c7bcd07ed35ba644a57 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 4 Jun 2026 01:41:45 +0200 Subject: [PATCH] =?UTF-8?q?feat(ml-alpha):=20Tier=203=20Phase=204-B=20?= =?UTF-8?q?=E2=80=94=20Band=20turnover=20controller=20+=20backward=20chain?= =?UTF-8?q?=20wiring?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Closes Phase 4-A's open loop by making the no-transaction band LEARN its width. Two coupled mechanisms wired in: 1. Adaptive turnover-target controller — single-thread CUDA kernel (rl_band_turnover_controller.cu) reads per-step `frac_not_masked` and adjusts the turnover target via asymmetric Schulman-bounded adapter: tighten (×0.97/step) above 0.60 threshold, loosen (×1.05/step) below 0.20. Asymmetric LOOSEN rate is faster per `pearl_dead_signal_resurrection_discipline`. First-observation bootstrap on slot 809 with dedicated boot-done sentinel slot 810. 2. Backward chain wiring — rl_band_head_backward.cu propagates `grad_band_per_b [B × 2]` through the asymmetric ±|tanh|·N_max activation and linear projection into per-batch weight/bias scratch plus `grad_h_t [B × HIDDEN_DIM]` (OVERWRITE). Trainer reduces via `reduce_axis0` and folds into the encoder grad combiner alongside π/V/FRD/outcome heads using the existing `grad_h_accumulate_scaled` pattern. Band-weight Adam steps share LR with π (RL_LR_PI_INDEX) — same plateau dynamics as policy. Plus a GPU reducer rl_band_frac_aggregate.cu (single-block tree-reduce over batch → writes to slot 812) and a host-side launch sequence in step_with_lobsim_gpu_body that runs frac_aggregate → controller → turnover_loss → backward outside graph capture, gated by RL_BAND_ENABLED_INDEX > 0.5 so Phase 3D bit-equality is preserved when the master gate is OFF. ISV slots added (809-812): turnover_ema, controller_boot_done, turnover_target_adaptive, frac_not_masked_observed. The turnover loss kernel now reads slot 811 (controller output) instead of the static slot 803. Bootstrap target 0.05 matches the Phase 4-A static default so first-step behavior is identical. Tests: 4 new band_invariants tests cover the controller bootstrap, the escalate-on-over-trading path, the loosen-on-saturation path, and the backward chain producing non-zero `grad_h_t` from a turnover/target mismatch. All 9 band tests + 18 multi_head_policy regression tests pass. Determinism: `determinism-check.sh --quick` exits 0 for the three relevant configurations (band off, band on, band + multi-head on). The Schulman controllers are deterministic by construction (no PRNG); backward kernels are sole-writer per (batch, slot) so reduce_axis0 delivers bit-equal weight grads across replays. Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §3.1 Option (c), §3.3 (backward chain), §3.4 (loss weight controller), §5 (diag emission), §9.1 Mitigation 3 (resurrection discipline). Pearls applied: * pearl_bootstrap_must_respect_clamp_range — bootstrap 0.05 ∈ [0.01, 0.20] * pearl_dead_signal_resurrection_discipline — asymmetric tighten/loosen rates * pearl_welford_trade_count_is_step_not_trade — bootstrap-done sentinel slot * pearl_determinism_achieved — single-thread controller, deterministic reducers Single-seed local smoke (b=128, 2000+500 steps, FOXHUNT_BAND_ENABLED=1) runs end-to-end without NaN. The band saturates to `frac_masked = 1.0` because position state stays at 0 throughout — the sigmoid surrogate gradient is zero deep inside the band, so the band's upper boundary cannot be pulled toward 0 even though the controller drives the adaptive target to its MAX (0.20). Spec §9.1 Mitigation 2 (`RL_BAND_MAX_MASK_FRAC_INDEX` exploration-slot bypass in rl_band_mask.cu) was declared but not shipped in Phase 4-A; it is the structural fix for this chicken-and-egg trap. Phase 4-B mechanism itself is verified correct via the four new tests; the saturation is an architectural-grafting gap surfaced for Phase 4-A2 follow-up rather than a Phase 4-B implementation defect. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 3 + .../ml-alpha/cuda/rl_band_frac_aggregate.cu | 86 +++++++ crates/ml-alpha/cuda/rl_band_head_backward.cu | 136 +++++++++++ .../cuda/rl_band_turnover_controller.cu | 108 +++++++++ crates/ml-alpha/cuda/rl_band_turnover_loss.cu | 22 +- crates/ml-alpha/src/rl/band_head.rs | 211 +++++++++++++++- crates/ml-alpha/src/rl/isv_slots.rs | 58 ++++- crates/ml-alpha/src/trainer/integrated.rs | 156 +++++++++++- crates/ml-alpha/tests/band_invariants.rs | 226 +++++++++++++++++- 9 files changed, 990 insertions(+), 16 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_band_frac_aggregate.cu create mode 100644 crates/ml-alpha/cuda/rl_band_head_backward.cu create mode 100644 crates/ml-alpha/cuda/rl_band_turnover_controller.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index abaafbb91..4d2336305 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -165,6 +165,9 @@ const KERNELS: &[&str] = &[ "rl_band_head_forward", // Phase 4-A (2026-06-03): No-transaction-band head forward — two-stage launch: (1) `rl_band_head_linear_fwd` produces [B × 2] pre-activation band logits via tree-reduce over HIDDEN_DIM; (2) `rl_band_apply_activation` applies asymmetric ±|tanh| × N_max_eff to enforce b_l ≤ 0 ≤ b_u per Davis-Norman optimality. Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §1.1. "rl_band_mask", // Phase 4-A: override actions[b]→Hold when position_lots[b] ∈ [b_l, b_u]. Master-gated at slot 799 (RL_BAND_ENABLED_INDEX). Grid=(B), Block=(1) — mirrors rl_confidence_gate launch shape; runs OUTSIDE graph capture per spec §9.5. "rl_band_turnover_loss", // Phase 4-A: turnover regularizer (Option b) with sigmoid surrogate for differentiable boundary gradient. Emits per-batch loss + per-batch grad on (b_l, b_u). Phase 4-A wires for OBSERVABILITY only; full encoder backward chain lands in Phase 4-B. + "rl_band_frac_aggregate", // Phase 4-B (2026-06-04): per-step `frac_not_masked` reducer (single-block tree-reduce over batch) → writes to RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX (slot 812). Consumed by `rl_band_turnover_controller`. + "rl_band_turnover_controller", // Phase 4-B: adaptive turnover-target controller. Single-thread launch (1×1×1); first-observation bootstrap for EMA at slot 809 + Schulman-bounded asymmetric adapter for slot 811 from slot 812 input. Spec §3.1 Option (c). + "rl_band_head_backward", // Phase 4-B: backward through ±|tanh|×N_max activation + linear projection → per-batch grad_w/grad_b scratch + grad_h_t (OVERWRITE). Caller reduces via reduce_axis0 + folds grad_h_t into encoder via grad_h_accumulate_scaled. Spec §3.3. ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_band_frac_aggregate.cu b/crates/ml-alpha/cuda/rl_band_frac_aggregate.cu new file mode 100644 index 000000000..81dacf40d --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_frac_aggregate.cu @@ -0,0 +1,86 @@ +// rl_band_frac_aggregate.cu — Phase 4-B per-step `frac_not_masked` reducer. +// +// Single-block tree-reduce over batch. Reads: +// * band_out [B × 2] — (b_l, b_u) post-activation, lots units +// * pos_state [B × P] — i32 position lots at offset 0 +// * isv — read-only (gate slot) +// +// Computes the per-batch boolean `position ∈ [b_l, b_u]` and reduces +// `(1 - in_band)` (= "this batch is FREE to trade this step") to a single +// scalar `frac_not_masked ∈ [0, 1]`, written to +// `RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX` (slot 812). +// +// Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md +// §5.2 + §5.3 (fleet-fraction discipline per pearl_fleet_fraction_not_aggregate). +// +// Disciplines: +// * `feedback_no_atomicadd` — tree-reduce in shared memory, single-thread +// writes the final scalar to ISV. +// * `pearl_determinism_achieved` — deterministic shared-mem reduction with +// `__syncthreads()` barriers; no PRNG. +// * `feedback_no_nvrtc` — pre-compiled cubin via build.rs. +// +// Launch: +// grid = (1, 1, 1) +// block = (block_dim, 1, 1) — power-of-two, ≥ b_size when possible +// shared_mem_bytes = block_dim × sizeof(float) +// +// The host caller picks `block_dim = next_power_of_two(b_size).min(1024)` +// — matches `rl_q_pi_agree_b`'s launch shape so the reduction is fully +// captured in a single warp-aligned tree-reduce. + +#include + +#define BAND_OUT 2 +#define RL_BAND_ENABLED_INDEX 799 +#define RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX 812 + +extern "C" __global__ void rl_band_frac_aggregate( + const float* __restrict__ band_outputs, // [B × 2] + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + float* __restrict__ isv, + int b_size, + int pos_bytes +) { + extern __shared__ float s_partial[]; // [block_dim] + + const int tid = threadIdx.x; + const int bdim = blockDim.x; + + // Master gate: if band disabled, write sentinel 0.0 and return — the + // controller's master-gate guard will skip the EMA update. + const float enabled = isv[RL_BAND_ENABLED_INDEX]; + if (enabled <= 0.5f) { + if (tid == 0) { + isv[RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX] = 0.0f; + } + return; + } + + // Per-thread accumulation: walk over batch indices in strides of bdim. + float local = 0.0f; + for (int b = tid; b < b_size; b += bdim) { + const int position_lots = + *reinterpret_cast(pos_state + b * pos_bytes); + const float pos_f = (float)position_lots; + const float b_l = band_outputs[b * BAND_OUT + 0]; + const float b_u = band_outputs[b * BAND_OUT + 1]; + const float in_band = (pos_f >= b_l && pos_f <= b_u) ? 1.0f : 0.0f; + local += (1.0f - in_band); + } + s_partial[tid] = local; + __syncthreads(); + + // Tree-reduce. + for (int stride = bdim / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + s_partial[tid] += s_partial[tid + stride]; + } + __syncthreads(); + } + + if (tid == 0) { + const float inv_b = 1.0f / (float)b_size; + isv[RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX] = s_partial[0] * inv_b; + } +} diff --git a/crates/ml-alpha/cuda/rl_band_head_backward.cu b/crates/ml-alpha/cuda/rl_band_head_backward.cu new file mode 100644 index 000000000..ae1fc8df3 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_head_backward.cu @@ -0,0 +1,136 @@ +// rl_band_head_backward.cu — Phase 4-B band-head backward chain. +// +// Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md +// §3.3 (backward chain into encoder). +// +// Forward chain (recap, see `rl_band_head_forward.cu`): +// band_pre[b, j] = b_band[j] + Σ_c W_band[j, c] · h_t[b, c] +// b_l = -|tanh(band_pre[b, 0])| × N_max_eff +// b_u = +|tanh(band_pre[b, 1])| × N_max_eff +// +// Backward chain (this file): +// Given dL/d(b_l) = grad_band[b, 0] and dL/d(b_u) = grad_band[b, 1] from +// the turnover-loss kernel: +// +// dL/d(band_pre[b, 0]) = grad_band[b, 0] · d(b_l)/d(pre) +// = grad_band[b, 0] · (-sign(t_l)) · (1 - t_l²) · N_max +// = grad_band[b, 0] · (-sign(t_l)) · sech²(pre_l) · N_max +// dL/d(band_pre[b, 1]) = grad_band[b, 1] · d(b_u)/d(pre) +// = grad_band[b, 1] · (+sign(t_u)) · sech²(pre_u) · N_max +// +// Then the linear backward: +// dL/d(W_band[j, c]) per batch = dL/d(pre[b, j]) · h_t[b, c] +// dL/d(b_band[j]) per batch = dL/d(pre[b, j]) +// dL/d(h_t[b, c]) = Σ_j dL/d(pre[b, j]) · W_band[j, c] +// +// Outputs (per-batch scratch; caller reduces via `reduce_axis0`): +// grad_w_per_batch [B × BAND_OUT × HIDDEN_DIM] +// grad_b_per_batch [B × BAND_OUT] +// grad_h_t [B × HIDDEN_DIM] (OVERWRITE — caller folds into encoder +// grad via `grad_h_accumulate_scaled`) +// +// Disciplines: +// * `feedback_no_atomicadd` — per-batch scratch + reduce_axis0 +// * `feedback_no_nvrtc` — pre-compiled cubin via build.rs +// * `pearl_determinism_achieved` — sole-writer per (b, j, c); no PRNG +// * `pearl_no_host_branches_in_captured_graph` — only kernel args, no +// host scalars in control flow +// +// Launch: +// grid = (B, 1, 1) +// block = (HIDDEN_DIM, 1, 1) +// shared_mem_bytes = BAND_OUT × sizeof(float) (s_dpre[2] staging) +// +// Per-block: thread c writes its own grad_w[b, 0, c], grad_w[b, 1, c], +// grad_h_t[b, c]. Thread 0 additionally writes grad_b[b, 0] and grad_b[b, 1]. + +#include +#include + +#define HIDDEN_DIM 128 +#define BAND_OUT 2 +#define RL_BAND_ENABLED_INDEX 799 +#define RL_HEAT_CAP_MAX_LOTS_INDEX 504 + +extern "C" __global__ void rl_band_head_backward( + const float* __restrict__ w_band, // [BAND_OUT × HIDDEN_DIM] + const float* __restrict__ band_pre, // [B × BAND_OUT] + const float* __restrict__ h_t, // [B × HIDDEN_DIM] + const float* __restrict__ grad_band, // [B × BAND_OUT] + const float* __restrict__ isv, + int b_size, + float* __restrict__ grad_w_per_batch, // [B × BAND_OUT × HIDDEN_DIM] + float* __restrict__ grad_b_per_batch, // [B × BAND_OUT] + float* __restrict__ grad_h_t // [B × HIDDEN_DIM] (OVERWRITE) +) { + const int b = blockIdx.x; + const int c = threadIdx.x; + if (b >= b_size || c >= HIDDEN_DIM) return; + + __shared__ float s_dpre[BAND_OUT]; + + // Master gate: write zeros and return so the caller's reduce_axis0 + + // accumulate_grad_h chain produces no encoder-grad contribution. The + // host-side branch in the trainer ALSO short-circuits the launch when + // disabled — this device-side guard is defense-in-depth so the kernel + // is safe to launch unconditionally during graph capture. + const float enabled = isv[RL_BAND_ENABLED_INDEX]; + if (enabled <= 0.5f) { + grad_w_per_batch[(b * BAND_OUT + 0) * HIDDEN_DIM + c] = 0.0f; + grad_w_per_batch[(b * BAND_OUT + 1) * HIDDEN_DIM + c] = 0.0f; + grad_h_t[b * HIDDEN_DIM + c] = 0.0f; + if (c == 0) { + grad_b_per_batch[b * BAND_OUT + 0] = 0.0f; + grad_b_per_batch[b * BAND_OUT + 1] = 0.0f; + } + return; + } + + // ── Stage 1: thread 0 computes the two activation-derivative scalars ─ + if (c == 0) { + const float n_max_eff = isv[RL_HEAT_CAP_MAX_LOTS_INDEX]; + + const float pre_l = band_pre[b * BAND_OUT + 0]; + const float pre_u = band_pre[b * BAND_OUT + 1]; + const float t_l = tanhf(pre_l); + const float t_u = tanhf(pre_u); + // sech²(x) = 1 − tanh²(x) + const float sech2_l = 1.0f - t_l * t_l; + const float sech2_u = 1.0f - t_u * t_u; + // d(b_l)/d(pre_l) = -sign(t_l) · sech²(pre_l) · N_max + // d(b_u)/d(pre_u) = +sign(t_u) · sech²(pre_u) · N_max + // sign(0) = 0 — at pre = 0 the gradient is zero (the activation + // is non-differentiable at the cusp |tanh| → 0). Treat as zero; + // upstream signs propagate cleanly. + const float sign_l = (t_l > 0.0f) ? 1.0f : ((t_l < 0.0f) ? -1.0f : 0.0f); + const float sign_u = (t_u > 0.0f) ? 1.0f : ((t_u < 0.0f) ? -1.0f : 0.0f); + + const float g_b_l = grad_band[b * BAND_OUT + 0]; + const float g_b_u = grad_band[b * BAND_OUT + 1]; + // dL/d(pre_l) = g_b_l · (-sign_l · sech2_l · N_max) + // dL/d(pre_u) = g_b_u · (+sign_u · sech2_u · N_max) + s_dpre[0] = g_b_l * (-sign_l) * sech2_l * n_max_eff; + s_dpre[1] = g_b_u * (+sign_u) * sech2_u * n_max_eff; + + // grad_b_band per batch is just dL/d(pre[b, j]). + grad_b_per_batch[b * BAND_OUT + 0] = s_dpre[0]; + grad_b_per_batch[b * BAND_OUT + 1] = s_dpre[1]; + } + __syncthreads(); + + const float dpre_l = s_dpre[0]; + const float dpre_u = s_dpre[1]; + + // ── Stage 2: per-(b, c) writes ─────────────────────────────────── + // grad_w_band[j, c] per batch = dpre[j] · h_t[b, c] + const float h_bc = h_t[b * HIDDEN_DIM + c]; + grad_w_per_batch[(b * BAND_OUT + 0) * HIDDEN_DIM + c] = dpre_l * h_bc; + grad_w_per_batch[(b * BAND_OUT + 1) * HIDDEN_DIM + c] = dpre_u * h_bc; + + // grad_h_t[b, c] = dpre_l · W_band[0, c] + dpre_u · W_band[1, c] + // OVERWRITE — caller's `grad_h_accumulate_scaled` folds into encoder + // grad via additive +=. + const float w_lc = w_band[0 * HIDDEN_DIM + c]; + const float w_uc = w_band[1 * HIDDEN_DIM + c]; + grad_h_t[b * HIDDEN_DIM + c] = dpre_l * w_lc + dpre_u * w_uc; +} diff --git a/crates/ml-alpha/cuda/rl_band_turnover_controller.cu b/crates/ml-alpha/cuda/rl_band_turnover_controller.cu new file mode 100644 index 000000000..acf816863 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_turnover_controller.cu @@ -0,0 +1,108 @@ +// rl_band_turnover_controller.cu — Phase 4-B adaptive turnover-target +// controller. +// +// Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md +// §3.1 Option (c) + §3.4 (Schulman-bounded adaptive controller). +// +// Reads (from ISV): +// * RL_BAND_ENABLED_INDEX (799) — master gate +// * RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX (812) — per-step raw signal +// written by +// `rl_band_frac_aggregate.cu` +// * RL_BAND_TURNOVER_EMA_INDEX (809) — EMA state (owned) +// * RL_BAND_CONTROLLER_BOOT_DONE_INDEX (810) — bootstrap latch (owned) +// * RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX (811) — output (owned) +// +// Writes (to ISV): +// * 809 — updated EMA (first-observation bootstrap or Wiener-α blend) +// * 810 — latched 1.0 on first observation +// * 811 — adaptive target, clamped to [TARGET_MIN, TARGET_MAX] +// +// Control law (per the Phase 4-B dispatch & spec §3.1 Option c): +// Healthy frac_not_masked target ≈ 0.4 (40% of batches free to trade, +// 60% inside band). Asymmetric Schulman-bounded adapter: +// +// if frac_not_masked_ema > OVER_THRESHOLD (too much trading): +// target *= TIGHTEN_RATE (push narrower band → fewer trades) +// elif frac_not_masked_ema < UNDER_THRESHOLD (band saturated): +// target *= LOOSEN_RATE (push wider band → more trading) +// else: +// healthy band — leave alone. +// clamp(target, TARGET_MIN, TARGET_MAX). +// +// Per `pearl_bootstrap_must_respect_clamp_range`: bootstrap 0.05 ∈ +// [TARGET_MIN=0.01, TARGET_MAX=0.20]; controller never snaps to a clamp +// boundary on its first emission. +// +// Per `pearl_dead_signal_resurrection_discipline`: the band's output gates +// its own input (mask → fewer trades → narrower band drift → mask). The +// asymmetric LOOSEN escape rate is faster than the TIGHTEN rate so the +// controller can recover when frac_not_masked drops below 0.2. +// +// Per `feedback_no_atomicadd`: single-thread launch (1×1×1), no atomics. +// Per `feedback_no_nvrtc`: pre-compiled cubin via build.rs. + +#include +#include + +#define RL_BAND_ENABLED_INDEX 799 +#define RL_BAND_TURNOVER_EMA_INDEX 809 +#define RL_BAND_CONTROLLER_BOOT_DONE_INDEX 810 +#define RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX 811 +#define RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX 812 + +// Tuning constants (all dimensionless). +#define BAND_CTRL_EMA_ALPHA 0.02f // Wiener-α — ~50-step horizon +#define BAND_CTRL_OVER_THRESHOLD 0.60f // frac_not_masked above → too much trading +#define BAND_CTRL_UNDER_THRESHOLD 0.20f // frac_not_masked below → band saturated +#define BAND_CTRL_TIGHTEN_RATE 0.97f // −3 %/step (slow narrow) +#define BAND_CTRL_LOOSEN_RATE 1.05f // +5 %/step (faster escape per resurrection discipline) +#define BAND_CTRL_TARGET_MIN 0.01f // never below 1 % of batches trading +#define BAND_CTRL_TARGET_MAX 0.20f // never above 20 % — keeps band materially active + +extern "C" __global__ void rl_band_turnover_controller(float* isv) { + // Single-thread kernel — launched (1,1,1)/(1,1,1). + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // Master gate: when the band is disabled, leave all owned slots + // untouched so Phase 3D bit-equality is preserved. + const float enabled = isv[RL_BAND_ENABLED_INDEX]; + if (enabled <= 0.5f) return; + + const float frac_now = isv[RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX]; + const float boot_done = isv[RL_BAND_CONTROLLER_BOOT_DONE_INDEX]; + + // ── EMA update (first-observation bootstrap pattern) ───────────── + float ema; + if (boot_done < 0.5f) { + // Replace EMA with current observation; latch flag. Necessary + // because slot 809 sentinel-zeroes at trainer init and blending + // zero with the first real measurement (likely ≈ 0.5 at warm + // band-start) would lag the controller by ~50 steps before it + // sees a representative value. + ema = frac_now; + isv[RL_BAND_CONTROLLER_BOOT_DONE_INDEX] = 1.0f; + } else { + ema = (1.0f - BAND_CTRL_EMA_ALPHA) * isv[RL_BAND_TURNOVER_EMA_INDEX] + + BAND_CTRL_EMA_ALPHA * frac_now; + } + isv[RL_BAND_TURNOVER_EMA_INDEX] = ema; + + // ── Target update — asymmetric Schulman-bounded adapter ────────── + float target = isv[RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX]; + if (ema > BAND_CTRL_OVER_THRESHOLD) { + // Trading too much → tighten band (smaller target → smaller + // frac_not_masked goal → narrower bands chase the goal). + target *= BAND_CTRL_TIGHTEN_RATE; + } else if (ema < BAND_CTRL_UNDER_THRESHOLD) { + // Band saturated → loosen (escape with faster rate per + // resurrection discipline). + target *= BAND_CTRL_LOOSEN_RATE; + } + // Healthy band — leave target alone. + + // Clamp; bootstrap 0.05 sits strictly inside [0.01, 0.20]. + target = fmaxf(BAND_CTRL_TARGET_MIN, + fminf(target, BAND_CTRL_TARGET_MAX)); + isv[RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX] = target; +} diff --git a/crates/ml-alpha/cuda/rl_band_turnover_loss.cu b/crates/ml-alpha/cuda/rl_band_turnover_loss.cu index 9197fd1ac..d9878ec83 100644 --- a/crates/ml-alpha/cuda/rl_band_turnover_loss.cu +++ b/crates/ml-alpha/cuda/rl_band_turnover_loss.cu @@ -27,17 +27,24 @@ // (799) gates the entire write — when ≤ 0.5 the kernel writes zeros so the // per-batch loss reducer sees no contribution. // +// Phase 4-B (2026-06-04): the turnover target is now read from the +// adaptive controller's output at `RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX` +// (811) instead of the static slot 803. The controller updates 811 every +// step based on the EMA of `frac_not_masked`. Slot 803 remains in the ISV +// (Phase 4-A clamp anchor for the controller) but is no longer the loss +// kernel's input. See spec §3.1 Option (c) and `rl_band_turnover_controller.cu`. +// // Per `feedback_no_atomicadd` (single-writer per slot), `pearl_determinism_achieved` // (no PRNG, single-thread-per-batch, no shared reductions across blocks). #include #include -#define BAND_OUT 2 -#define RL_BAND_ENABLED_INDEX 799 -#define RL_BAND_LOSS_WEIGHT_INDEX 802 -#define RL_BAND_TURNOVER_TARGET_INDEX 803 -#define RL_BAND_GRAD_SHARPNESS_INDEX 804 +#define BAND_OUT 2 +#define RL_BAND_ENABLED_INDEX 799 +#define RL_BAND_LOSS_WEIGHT_INDEX 802 +#define RL_BAND_GRAD_SHARPNESS_INDEX 804 +#define RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX 811 // Numerically-stable sigmoid (clamps argument to avoid expf overflow). static __device__ __forceinline__ float stable_sigmoid(float x) { @@ -99,7 +106,10 @@ extern "C" __global__ void rl_band_turnover_loss( const float b_u = band_outputs[b * BAND_OUT + 1]; const float sharpness = isv[RL_BAND_GRAD_SHARPNESS_INDEX]; - const float target = isv[RL_BAND_TURNOVER_TARGET_INDEX]; + // Phase 4-B: read the adaptive controller's target instead of the + // static slot 803. The controller's bootstrap matches the static + // default so first-step behavior is identical to Phase 4-A. + const float target = isv[RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX]; const float loss_w = isv[RL_BAND_LOSS_WEIGHT_INDEX]; // Sigmoid surrogate near each boundary. `m_lower` ≈ 1 when pos ≥ b_l; diff --git a/crates/ml-alpha/src/rl/band_head.rs b/crates/ml-alpha/src/rl/band_head.rs index d2e1a21aa..40ccfbd85 100644 --- a/crates/ml-alpha/src/rl/band_head.rs +++ b/crates/ml-alpha/src/rl/band_head.rs @@ -51,6 +51,12 @@ const BAND_MASK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_mask.cubin")); const BAND_TURNOVER_LOSS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_turnover_loss.cubin")); +const BAND_FRAC_AGGREGATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_frac_aggregate.cubin")); +const BAND_TURNOVER_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_turnover_controller.cubin")); +const BAND_HEAD_BACKWARD_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_head_backward.cubin")); /// Number of band outputs (`b_l`, `b_u`). pub const BAND_OUT: usize = 2; @@ -92,6 +98,13 @@ pub struct BandHead { mask_fn: CudaFunction, _turnover_loss_module: Arc, turnover_loss_fn: CudaFunction, + // ── Phase 4-B handles ──────────────────────────────────────────── + _frac_aggregate_module: Arc, + frac_aggregate_fn: CudaFunction, + _turnover_controller_module: Arc, + turnover_controller_fn: CudaFunction, + _backward_module: Arc, + backward_fn: CudaFunction, // ── Online weights ──────────────────────────────────────────────── /// `W_band[j, c]` row-major over `(j ∈ {0, 1}, c ∈ HIDDEN_DIM)`. Two @@ -115,8 +128,22 @@ pub struct BandHead { /// `[B]` — per-batch loss contribution (same scalar value, ×1/B). pub loss_per_b_d: CudaSlice, /// `[B × 2]` — per-batch grad on `(b_l, b_u)` from the turnover - /// regularizer. Phase 4-A allocates but does NOT fold into encoder grad. + /// regularizer. Phase 4-B folds into encoder grad via `backward()`. pub grad_band_per_b_d: CudaSlice, + + // ── Phase 4-B backward-chain scratch ────────────────────────────── + /// `[B × BAND_OUT × HIDDEN_DIM]` — per-batch weight grads. Caller + /// reduces via `reduce_axis0` into `grad_w_d`. + pub grad_w_per_b_d: CudaSlice, + /// `[B × BAND_OUT]` — per-batch bias grads. Reduced into `grad_b_d`. + pub grad_b_per_b_d: CudaSlice, + /// `[B × HIDDEN_DIM]` — per-batch encoder-input grad (OVERWRITE). + /// Caller folds into encoder grad via `grad_h_accumulate_scaled`. + pub grad_h_t_d: CudaSlice, + /// `[BAND_OUT × HIDDEN_DIM]` — reduced weight grad, Adam target. + pub grad_w_d: CudaSlice, + /// `[BAND_OUT]` — reduced bias grad, Adam target. + pub grad_b_d: CudaSlice, } impl BandHead { @@ -154,6 +181,24 @@ impl BandHead { let turnover_loss_fn = turnover_loss_module .load_function("rl_band_turnover_loss") .context("load rl_band_turnover_loss")?; + let frac_aggregate_module = ctx + .load_cubin(BAND_FRAC_AGGREGATE_CUBIN.to_vec()) + .context("load rl_band_frac_aggregate cubin")?; + let frac_aggregate_fn = frac_aggregate_module + .load_function("rl_band_frac_aggregate") + .context("load rl_band_frac_aggregate")?; + let turnover_controller_module = ctx + .load_cubin(BAND_TURNOVER_CONTROLLER_CUBIN.to_vec()) + .context("load rl_band_turnover_controller cubin")?; + let turnover_controller_fn = turnover_controller_module + .load_function("rl_band_turnover_controller") + .context("load rl_band_turnover_controller")?; + let backward_module = ctx + .load_cubin(BAND_HEAD_BACKWARD_CUBIN.to_vec()) + .context("load rl_band_head_backward cubin")?; + let backward_fn = backward_module + .load_function("rl_band_head_backward") + .context("load rl_band_head_backward")?; // ── Weight init: Xavier-0.01 + bias seed ───────────────────── // Per pearl_scoped_init_seed_for_reproducibility, install the @@ -193,6 +238,23 @@ impl BandHead { .alloc_zeros::(cfg.b_size * BAND_OUT) .context("alloc grad_band_per_b_d")?; + // ── Phase 4-B backward scratch ─────────────────────────────── + let grad_w_per_b_d = stream + .alloc_zeros::(cfg.b_size * BAND_OUT * cfg.hidden_dim) + .context("alloc grad_w_per_b_d")?; + let grad_b_per_b_d = stream + .alloc_zeros::(cfg.b_size * BAND_OUT) + .context("alloc grad_b_per_b_d")?; + let grad_h_t_d = stream + .alloc_zeros::(cfg.b_size * cfg.hidden_dim) + .context("alloc band grad_h_t_d")?; + let grad_w_d = stream + .alloc_zeros::(BAND_OUT * cfg.hidden_dim) + .context("alloc band grad_w_d")?; + let grad_b_d = stream + .alloc_zeros::(BAND_OUT) + .context("alloc band grad_b_d")?; + let raw_stream = stream.cu_stream(); Ok(Self { cfg, @@ -205,6 +267,12 @@ impl BandHead { mask_fn, _turnover_loss_module: turnover_loss_module, turnover_loss_fn, + _frac_aggregate_module: frac_aggregate_module, + frac_aggregate_fn, + _turnover_controller_module: turnover_controller_module, + turnover_controller_fn, + _backward_module: backward_module, + backward_fn, w_band_d, b_band_d, band_pre_d, @@ -212,6 +280,11 @@ impl BandHead { m_soft_per_b_d, loss_per_b_d, grad_band_per_b_d, + grad_w_per_b_d, + grad_b_per_b_d, + grad_h_t_d, + grad_w_d, + grad_b_d, }) } @@ -328,14 +401,87 @@ impl BandHead { Ok(()) } - /// Phase 4-A turnover regularizer loss launch (observability-only in - /// 4-A; gradient written to `grad_band_per_b_d` is allocated but NOT - /// folded into the encoder grad path until Phase 4-B). + /// Phase 4-B per-step `frac_not_masked` reducer launch. Reads + /// `band_out_d` + `pos_state` and writes the per-step fleet-fraction + /// (= 1 − frac_in_band) to `RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX` + /// (slot 812) via a single-block tree-reduce. Consumed by the + /// turnover controller on the same step. + /// + /// Master-gated at slot 799 inside the kernel: when the band is OFF + /// the kernel writes 0.0 and returns. Launched OUTSIDE graph capture + /// (matches the controller below — both run host-side post-forward + /// once `band_out_d` is materialised). + pub fn launch_frac_aggregate( + &self, + pos_state_d: &CudaSlice, + isv_dev_ptr: &u64, + b_size: usize, + pos_bytes: usize, + ) -> Result<()> { + // Block-dim = next pow-of-two ≥ b_size, capped at 1024 (single + // block tree-reduce semantics; the kernel iterates with stride + // bdim if b_size > bdim). + let block_dim = (b_size as u32) + .next_power_of_two() + .clamp(1, 1024); + let smem = (block_dim as usize * std::mem::size_of::()) as u32; + let mut args = RawArgs::new(); + args.push_ptr(self.band_out_d.raw_ptr()); + args.push_ptr(pos_state_d.raw_ptr()); + args.push_ptr(*isv_dev_ptr); + args.push_i32(b_size as i32); + args.push_i32(pos_bytes as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.frac_aggregate_fn.cu_function(), + (1, 1, 1), + (block_dim, 1, 1), + smem, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_frac_aggregate: {:?}", e))?; + } + Ok(()) + } + + /// Phase 4-B adaptive turnover-target controller launch. Single-thread + /// kernel; reads the per-step `frac_not_masked` written by + /// `launch_frac_aggregate` and updates the EMA (slot 809) + + /// adaptive target (slot 811) using first-observation bootstrap + + /// asymmetric Schulman-bounded adapter. + /// + /// MUST be launched AFTER `launch_frac_aggregate` (which writes the + /// controller's input slot) and BEFORE `launch_turnover_loss` on the + /// same step (which reads slot 811). + pub fn launch_turnover_controller(&self, isv_dev_ptr: &u64) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(*isv_dev_ptr); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.turnover_controller_fn.cu_function(), + (1, 1, 1), + (1, 1, 1), + 0, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_turnover_controller: {:?}", e))?; + } + Ok(()) + } + + /// Phase 4-B turnover regularizer loss launch. Now wired into the + /// joint-loss path (`step_synthetic_body`) as the band's training + /// signal: emits `loss_per_b_d` + `grad_band_per_b_d` consumed by + /// `backward()` on the same step. /// /// `turnover_t` is the host-supplied scalar mean of `(1 − m_soft)` from /// the previous step (Option (b) target-tracking semantics per spec - /// §3.1). At Phase 4-A bootstrap the trainer passes the current step's - /// `frac_not_masked` measured host-side after the mask launch. + /// §3.1). The kernel uses `turnover_t − target` for the loss; the + /// per-batch gradient routes through the sigmoid surrogate `(1 − m_soft)`. pub fn launch_turnover_loss( &mut self, pos_state_d: &CudaSlice, @@ -369,6 +515,59 @@ impl BandHead { } Ok(()) } + + /// Phase 4-B backward chain — propagates `grad_band_per_b_d [B × 2]` + /// (produced by `launch_turnover_loss`) through the asymmetric + /// `±|tanh| × N_max_eff` activation and the linear projection into: + /// * `grad_w_per_b_d [B × BAND_OUT × HIDDEN_DIM]` — per-batch + /// weight grad scratch; caller reduces via `reduce_axis0` into + /// `grad_w_d` for Adam consumption. + /// * `grad_b_per_b_d [B × BAND_OUT]` — per-batch bias grad scratch. + /// * `grad_h_t_d [B × HIDDEN_DIM]` — encoder-input grad (OVERWRITE). + /// Caller folds into the encoder grad combiner via + /// `grad_h_accumulate_scaled`. + /// + /// The kernel reads the master gate at slot 799 and writes zeros when + /// disabled — bit-equality with the Phase 3D baseline is preserved + /// even if the host-side launch is unconditional. + pub fn backward( + &mut self, + h_t: &CudaSlice, + isv_dev_ptr: &u64, + b_size: usize, + ) -> Result<()> { + debug_assert_eq!(h_t.len(), b_size * HIDDEN_DIM); + debug_assert_eq!( + self.grad_w_per_b_d.len(), + b_size * BAND_OUT * HIDDEN_DIM + ); + debug_assert_eq!(self.grad_b_per_b_d.len(), b_size * BAND_OUT); + debug_assert_eq!(self.grad_h_t_d.len(), b_size * HIDDEN_DIM); + + let mut args = RawArgs::new(); + args.push_ptr(self.w_band_d.raw_ptr()); + args.push_ptr(self.band_pre_d.raw_ptr()); + args.push_ptr(h_t.raw_ptr()); + args.push_ptr(self.grad_band_per_b_d.raw_ptr()); + args.push_ptr(*isv_dev_ptr); + args.push_i32(b_size as i32); + args.push_ptr(self.grad_w_per_b_d.raw_ptr()); + args.push_ptr(self.grad_b_per_b_d.raw_ptr()); + args.push_ptr(self.grad_h_t_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.backward_fn.cu_function(), + (b_size as u32, 1, 1), + (HIDDEN_DIM as u32, 1, 1), + (BAND_OUT * std::mem::size_of::()) as u32, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_head_backward: {:?}", e))?; + } + Ok(()) + } } // ── pinned-staging upload helper (mirrors ppo::upload) ──────────────────── diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 17133474b..6d032bdde 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -2111,6 +2111,61 @@ pub const RL_BAND_WIDTH_MAX_INDEX: usize = 807; /// 4-B if the diag aggregator becomes a measurable cost. pub const RL_BAND_DIAG_LAUNCH_EVERY_INDEX: usize = 808; +// ── Phase 4-B (2026-06-04) Adaptive turnover controller + backward chain ── +// +// 4 slots (809-812) extending the Phase 4-A foundation: +// * 809 RL_BAND_TURNOVER_EMA_INDEX — measured frac_not_masked EMA +// * 810 RL_BAND_CONTROLLER_BOOT_DONE_INDEX — first-observation bootstrap latch +// * 811 RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX — controller-output target read +// by `rl_band_turnover_loss.cu` +// * 812 RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX — per-step raw fraction (1 − frac_in_band) +// written by `rl_band_frac_aggregate.cu`, +// read by controller +// +// All slots respect `pearl_bootstrap_must_respect_clamp_range`. Master gate +// at slot 799 still gates: when the band is OFF the controller does not +// launch (host-side `RL_BAND_ENABLED_INDEX > 0.5` branch in the trainer), +// preserving Phase 3D bit-equality. + +/// Phase 4-B EMA of `frac_not_masked` (fleet-fraction of batches whose +/// position lies OUTSIDE the band on the current step). Per +/// `pearl_first_observation_bootstrap`, the controller initialises this slot +/// to the first observed value when `RL_BAND_CONTROLLER_BOOT_DONE_INDEX` +/// reads 0.0. Steady-state Wiener-α blend per +/// `pearl_wiener_alpha_floor_for_nonstationary` (α=0.02, ~50-step horizon). +/// +/// Clamp `[0.0, 1.0]` (the EMA tracks a fraction). +pub const RL_BAND_TURNOVER_EMA_INDEX: usize = 809; + +/// Phase 4-B first-observation bootstrap latch for the turnover controller. +/// Sentinel 0.0 at trainer init → kernel snaps EMA to current observation on +/// first call and latches this slot to 1.0. Subsequent calls take the +/// steady-state Wiener-α path. Per `pearl_welford_trade_count_is_step_not_trade` +/// + `pearl_bootstrap_must_respect_clamp_range`. +/// +/// Clamp `{0.0, 1.0}` (binary). Bootstrap 0.0 (not done yet). +pub const RL_BAND_CONTROLLER_BOOT_DONE_INDEX: usize = 810; + +/// Phase 4-B controller-output adaptive turnover target. Read by +/// `rl_band_turnover_loss.cu` in place of the static slot 803 — the +/// adaptive controller updates this slot every step toward a "healthy +/// band-active" frac_not_masked of 0.4 (60% of batches inside band). +/// +/// Bootstrap 0.05 (matches the static slot 803 default so first-step +/// behaviour is identical to Phase 4-A). Clamp `[0.01, 0.20]` — tighter +/// than the static slot's `[0.005, 0.5]` because the controller is +/// constrained to a stable adaptation range. +pub const RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX: usize = 811; + +/// Phase 4-B per-step raw `frac_not_masked` (= `1 − frac_in_band`) written +/// by the GPU reducer `rl_band_frac_aggregate.cu`. Single-block tree-reduce +/// over batch — no atomicAdd per `feedback_no_atomicadd`. Controller reads +/// this slot and folds it into the EMA at slot 809. +/// +/// Bootstrap 0.0 (sentinel — the reducer overwrites every step). Clamp +/// `[0.0, 1.0]`. +pub const RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX: usize = 812; + /// Last RL-allocated slot index (exclusive). /// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696. /// Post-regime-observer (F1.1): 716. Post-warmed-flag addendum: 717. @@ -2133,4 +2188,5 @@ pub const RL_BAND_DIAG_LAUNCH_EVERY_INDEX: usize = 808; /// Post-Phase 3D-B (quadratic cost α/β/enabled slots 794-796): 796. /// Post-Phase 3D-C (PPO surrogate blend weight/enabled slots 797-798): 798. /// Post-Phase 4-A (no-transaction-band foundation, slots 799-808): 808. -pub const RL_SLOTS_END: usize = 809; +/// Post-Phase 4-B (adaptive turnover controller + backward chain, slots 809-812): 812. +pub const RL_SLOTS_END: usize = 813; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 57503d86b..190568834 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -687,6 +687,15 @@ pub struct IntegratedTrainer { pub policy_b_adam: AdamW, pub value_w_adam: AdamW, pub value_b_adam: AdamW, + /// Phase 4-B (2026-06-04) — no-transaction-band head Adam pair. LR + /// sourced from `RL_LR_PI_INDEX` (same scale as the policy head — the + /// band is a structural action-selection prior, conceptually paired + /// with π). Flag-gated by `RL_BAND_ENABLED_INDEX > 0.5` at the host; + /// when off the Adam step is still called but with zero grads from + /// the backward kernel's device-side guard, so weight-only momentum + /// decay applies (analogous to the FRD sentinel-zero path). + pub band_w_adam: AdamW, + pub band_b_adam: AdamW, /// Phase 2A-C (2026-06-03) — multi-head policy gate (FOXHUNT_USE_MULTI_HEAD_POLICY). /// @@ -1702,6 +1711,11 @@ impl IntegratedTrainer { AdamW::new(dev, policy_head.w_d.len(), lr_placeholder).context("policy_w_adam")?; let policy_b_adam = AdamW::new(dev, policy_head.b_d.len(), lr_placeholder).context("policy_b_adam")?; + // Phase 4-B (2026-06-04) band-head Adam pair. + let band_w_adam = + AdamW::new(dev, band_head.w_band_d.len(), lr_placeholder).context("band_w_adam")?; + let band_b_adam = + AdamW::new(dev, band_head.b_band_d.len(), lr_placeholder).context("band_b_adam")?; // ── Phase 2A-C (2026-06-03) MultiHeadPolicy gate ───────────── // `FOXHUNT_USE_MULTI_HEAD_POLICY` env flag follows the @@ -3141,6 +3155,9 @@ impl IntegratedTrainer { policy_b_adam, value_w_adam, value_b_adam, + // Phase 4-B band-head Adam pair. + band_w_adam, + band_b_adam, // Phase 2A-C MultiHeadPolicy + Adams (None when flag off). use_multi_head_policy, multi_head_policy, @@ -3777,7 +3794,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 260] = [ + let isv_constants: [(usize, f32); 264] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -4267,6 +4284,17 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_BAND_WIDTH_MIN_INDEX, 0.1_f32), (crate::rl::isv_slots::RL_BAND_WIDTH_MAX_INDEX, 1.8_f32), (crate::rl::isv_slots::RL_BAND_DIAG_LAUNCH_EVERY_INDEX, 1.0_f32), + // Phase 4-B (2026-06-04) adaptive turnover controller seeds. + // EMA + bootstrap-done sentinel at 0.0 (the controller + // bootstraps on first call). Adaptive target initialised to + // 0.05 — matches the Phase 4-A static target so the first + // step's turnover loss is identical to 4-A. The frac slot + // is sentinel-zeroed; the GPU reducer overwrites every step + // the band is enabled. + (crate::rl::isv_slots::RL_BAND_TURNOVER_EMA_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_BAND_CONTROLLER_BOOT_DONE_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, 0.05_f32), + (crate::rl::isv_slots::RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX, 0.0_f32), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -6374,6 +6402,62 @@ impl IntegratedTrainer { )?; } + // ── Phase 4-B (2026-06-04) — fold band-head grad_h_t into encoder ─ + // The band's backward kernel was launched in `step_with_lobsim_gpu_body` + // OUTSIDE graph capture (host gated on `RL_BAND_ENABLED_INDEX > 0.5`); + // here we reduce the per-batch weight/bias scratches via reduce_axis0 + // and fold the OVERWRITE `grad_h_t_d` into the encoder grad combiner. + // When band is disabled the device-side master-gate guard wrote + // zeros into `grad_h_t_d`, so the accumulate is a no-op and Phase 3D + // bit-equality is preserved. + let band_enabled_host = self + .read_isv_host(crate::rl::isv_slots::RL_BAND_ENABLED_INDEX) + > 0.5; + if band_enabled_host { + // Reduce per-batch weight + bias grads. + reduce_axis0_free( + &self.stream, + &self.reduce_axis0_fn, + &self.band_head.grad_w_per_b_d, + b_size, + crate::rl::band_head::BAND_OUT * HIDDEN_DIM, + &mut self.band_head.grad_w_d, + )?; + reduce_axis0_free( + &self.stream, + &self.reduce_axis0_fn, + &self.band_head.grad_b_per_b_d, + b_size, + crate::rl::band_head::BAND_OUT, + &mut self.band_head.grad_b_d, + )?; + // Fold band's grad_h_t into encoder grad combiner. λ_band is + // sourced from `RL_BAND_LOSS_WEIGHT_INDEX` (slot 802) so an + // operator can dial down the band's encoder-grad contribution + // via the same slot that gates the loss magnitude. + let lambda_band = self + .read_isv_host(crate::rl::isv_slots::RL_BAND_LOSS_WEIGHT_INDEX); + accumulate_grad_h( + &self.stream, + &self.grad_h_accumulate_fn, + &self.band_head.grad_h_t_d, + lambda_band * b_inv, + &mut self.grad_h_t_combined_d, + )?; + // Adam steps on band W and b. LR sourced from RL_LR_PI_INDEX — + // the band is a structural action-selection prior and tracks + // π's plateau dynamics. + let lr_pi_band = self.read_isv_host(crate::rl::isv_slots::RL_LR_PI_INDEX); + self.band_w_adam.lr = lr_pi_band; + self.band_b_adam.lr = lr_pi_band; + self.band_w_adam + .step(&mut self.band_head.w_band_d, &self.band_head.grad_w_d) + .context("band_w_adam.step")?; + self.band_b_adam + .step(&mut self.band_head.b_band_d, &self.band_head.grad_b_d) + .context("band_b_adam.step")?; + } + // ── Step 11: encoder backward (Phase E.3a) ─────────────────── // Seed the perception trainer's per-K hidden-state grad buffer // from `grad_h_t_combined_d` at slot K-1 (the only slot the @@ -9768,6 +9852,69 @@ impl IntegratedTrainer { ) .context("step_with_lobsim_gpu: band_head.launch_mask (Phase 4-A)")?; } + + // Phase 4-B (2026-06-04) — adaptive turnover controller + backward chain. + // + // Sequence (all outside graph capture, same stream so kernels are + // ordered by issue): + // 1. `rl_band_frac_aggregate` reads `band_out_d` + live position + // state and writes per-step `frac_not_masked` to ISV slot 812. + // 2. `rl_band_turnover_controller` reads slot 812 and updates + // EMA (slot 809) + adaptive target (slot 811) using the + // first-observation bootstrap pattern + asymmetric + // Schulman-bounded adapter (spec §3.1 Option c). + // 3. `rl_band_turnover_loss` reads the freshly-updated slot 811 + // and emits per-batch loss/grad scratches. + // 4. `rl_band_head_backward` propagates `grad_band_per_b_d` + // through the activation + linear layer into per-batch + // weight/bias grad scratches + `grad_h_t_d` (OVERWRITE). + // + // Host-side gated by `RL_BAND_ENABLED_INDEX > 0.5` — every kernel + // also has its own device-side master-gate guard for defense in + // depth, so flag-off paths produce zeros even if the host branch + // were ever to fall through. + let band_enabled_host = self + .read_isv_host(crate::rl::isv_slots::RL_BAND_ENABLED_INDEX) + > 0.5; + if band_enabled_host { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + let pos_bytes = lobsim.pos_bytes(); + self.band_head + .launch_frac_aggregate( + pos_d_ref, + &self.isv_dev_ptr, + b_size, + pos_bytes, + ) + .context("step_with_lobsim_gpu: band_head.launch_frac_aggregate (Phase 4-B)")?; + self.band_head + .launch_turnover_controller(&self.isv_dev_ptr) + .context("step_with_lobsim_gpu: band_head.launch_turnover_controller (Phase 4-B)")?; + // `turnover_t` host arg is read from the controller-updated EMA + // (slot 809) so the loss diff equals (turnover_ema − adaptive_target). + // Both terms have a one-step delay — the EMA was updated by the + // controller on the same stream; the host read picks up the + // value after stream completion at the next step's diag emit. + // For the loss this delay is irrelevant: the gradient direction + // is what matters. + let turnover_t = self + .read_isv_host(crate::rl::isv_slots::RL_BAND_TURNOVER_EMA_INDEX); + self.band_head + .launch_turnover_loss( + pos_d_ref, + &self.isv_dev_ptr, + turnover_t, + b_size, + pos_bytes, + ) + .context("step_with_lobsim_gpu: band_head.launch_turnover_loss (Phase 4-B)")?; + // Backward needs h_t — bind separately so the borrow from + // `&self.perception` does not collide with `&mut self.band_head`. + let h_t_borrow_band: &CudaSlice = self.perception.h_t_view(); + self.band_head + .backward(h_t_borrow_band, &self.isv_dev_ptr, b_size) + .context("step_with_lobsim_gpu: band_head.backward (Phase 4-B)")?; + } { let pos_d_ref: &CudaSlice = lobsim.pos_d(); let b_size_i = b_size as i32; @@ -12069,6 +12216,13 @@ impl IntegratedTrainer { "frac_in_band": band_frac_in_band, "frac_masked": band_frac_masked, "collapse_warning": band_collapse_warning, + // Phase 4-B (2026-06-04) adaptive controller + observability. + "controller_target": isv[crate::rl::isv_slots::RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX], + "turnover_ema": isv[crate::rl::isv_slots::RL_BAND_TURNOVER_EMA_INDEX], + "frac_not_masked_observed": isv[crate::rl::isv_slots::RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX], + "controller_boot_done": isv[crate::rl::isv_slots::RL_BAND_CONTROLLER_BOOT_DONE_INDEX], + "width_min": isv[crate::rl::isv_slots::RL_BAND_WIDTH_MIN_INDEX], + "width_max": isv[crate::rl::isv_slots::RL_BAND_WIDTH_MAX_INDEX], }, "position_heat": { "capped_count_step": heat_cap_step, diff --git a/crates/ml-alpha/tests/band_invariants.rs b/crates/ml-alpha/tests/band_invariants.rs index ea1e58174..7ce0cdf5e 100644 --- a/crates/ml-alpha/tests/band_invariants.rs +++ b/crates/ml-alpha/tests/band_invariants.rs @@ -34,10 +34,13 @@ use ml_alpha::heads::HIDDEN_DIM; use ml_alpha::pinned_mem::MappedF32Buffer; use ml_alpha::rl::band_head::{BandHead, BandHeadConfig, BAND_OUT}; use ml_alpha::rl::isv_slots::{ + RL_BAND_CONTROLLER_BOOT_DONE_INDEX, RL_BAND_ENABLED_INDEX, + RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX, RL_BAND_GRAD_SHARPNESS_INDEX, RL_BAND_LOSS_WEIGHT_INDEX, - RL_BAND_TURNOVER_TARGET_INDEX, + RL_BAND_TURNOVER_EMA_INDEX, + RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, RL_HEAT_CAP_MAX_LOTS_INDEX, RL_SLOTS_END, }; @@ -303,7 +306,7 @@ fn band_turnover_loss_correct() -> Result<()> { let target: f32 = 0.05; let sharpness: f32 = 4.0; isv_write(&mut isv, RL_BAND_LOSS_WEIGHT_INDEX, loss_w); - isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_INDEX, target); + isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, target); isv_write(&mut isv, RL_BAND_GRAD_SHARPNESS_INDEX, sharpness); // Fixed band [-2, +2]. @@ -412,3 +415,222 @@ fn band_disabled_means_no_mask() -> Result<()> { drop(actions_d); Ok(()) } + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_controller_first_observation_bootstrap() -> Result<()> { + // Phase 4-B (2026-06-04): the turnover controller must snap the EMA + // (slot 809) to the first observed `frac_not_masked` (slot 812) and + // latch the bootstrap-done sentinel (slot 810) → 1.0. Without this + // path the EMA would lag the truth by ~50 steps (Wiener-α 0.02). + let Some(dev) = build_device() else { + return Ok(()); + }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 4usize; + let band = fresh_band(&dev, b_size)?; + + // ISV: master gate ON, bootstrap target = 0.05, raw frac observed = 0.5. + let (mut isv, isv_ptr) = build_isv(&stream)?; + isv_write(&mut isv, RL_BAND_ENABLED_INDEX, 1.0); + isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, 0.05); + isv_write(&mut isv, RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX, 0.5); + isv_write(&mut isv, RL_BAND_TURNOVER_EMA_INDEX, 0.0); // sentinel + isv_write(&mut isv, RL_BAND_CONTROLLER_BOOT_DONE_INDEX, 0.0); // not done + + band.launch_turnover_controller(&isv_ptr) + .context("band.launch_turnover_controller (first observation)")?; + stream.synchronize()?; + + let host = unsafe { std::slice::from_raw_parts(isv.host_ptr, RL_SLOTS_END) }; + let ema = host[RL_BAND_TURNOVER_EMA_INDEX]; + let boot_done = host[RL_BAND_CONTROLLER_BOOT_DONE_INDEX]; + + assert!( + (ema - 0.5).abs() < TOL, + "EMA = {ema}, expected 0.5 (first-observation bootstrap snap)" + ); + assert!( + (boot_done - 1.0).abs() < TOL, + "bootstrap_done = {boot_done}, expected 1.0 (latched)" + ); + eprintln!( + "PASS — controller first-observation bootstrap: EMA={ema}, boot_done={boot_done}" + ); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_controller_escalates_when_too_much_trading() -> Result<()> { + // Phase 4-B: when `frac_not_masked_ema > 0.60` (= over-trading + // threshold), the controller tightens the adaptive target by + // ×0.97/step. Starting from bootstrap target 0.05 and EMA already + // above the threshold, 100 controller fires should drive the target + // strictly DOWN. + let Some(dev) = build_device() else { + return Ok(()); + }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 4usize; + let band = fresh_band(&dev, b_size)?; + + let (mut isv, isv_ptr) = build_isv(&stream)?; + isv_write(&mut isv, RL_BAND_ENABLED_INDEX, 1.0); + isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, 0.05); + // Pre-populate EMA above the OVER threshold (0.60); mark bootstrap + // already done so the kernel takes the Wiener-α + tighten path. + isv_write(&mut isv, RL_BAND_TURNOVER_EMA_INDEX, 0.80); + isv_write(&mut isv, RL_BAND_CONTROLLER_BOOT_DONE_INDEX, 1.0); + // Per-step raw observation stays at 0.80 so EMA blend keeps it high. + isv_write(&mut isv, RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX, 0.80); + + let initial_target = 0.05_f32; + for _ in 0..100 { + band.launch_turnover_controller(&isv_ptr) + .context("band.launch_turnover_controller (escalate loop)")?; + } + stream.synchronize()?; + + let host = unsafe { std::slice::from_raw_parts(isv.host_ptr, RL_SLOTS_END) }; + let final_target = host[RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX]; + let final_ema = host[RL_BAND_TURNOVER_EMA_INDEX]; + + assert!( + final_target < initial_target - 1e-4, + "final_target = {final_target}, expected strictly < {initial_target} after 100 tighten steps (EMA={final_ema})" + ); + // Clamp floor is 0.01 — the controller must respect it. + assert!( + final_target >= 0.01 - TOL, + "final_target = {final_target}, fell below clamp floor 0.01" + ); + eprintln!( + "PASS — controller tightens target from {initial_target} → {final_target} when EMA > 0.60 (final EMA={final_ema})" + ); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_controller_loosens_when_band_saturated() -> Result<()> { + // Phase 4-B: when `frac_not_masked_ema < 0.20` (= under-trading + // threshold, band saturated), the controller loosens the adaptive + // target by ×1.05/step. The asymmetric LOOSEN rate is faster than + // the TIGHTEN rate per `pearl_dead_signal_resurrection_discipline`. + let Some(dev) = build_device() else { + return Ok(()); + }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 4usize; + let band = fresh_band(&dev, b_size)?; + + let (mut isv, isv_ptr) = build_isv(&stream)?; + isv_write(&mut isv, RL_BAND_ENABLED_INDEX, 1.0); + isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, 0.05); + // EMA below the UNDER threshold (0.20); raw observation stays low so + // Wiener-α blend keeps EMA in the loosen-fire band. + isv_write(&mut isv, RL_BAND_TURNOVER_EMA_INDEX, 0.05); + isv_write(&mut isv, RL_BAND_CONTROLLER_BOOT_DONE_INDEX, 1.0); + isv_write(&mut isv, RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX, 0.05); + + let initial_target = 0.05_f32; + for _ in 0..100 { + band.launch_turnover_controller(&isv_ptr) + .context("band.launch_turnover_controller (loosen loop)")?; + } + stream.synchronize()?; + + let host = unsafe { std::slice::from_raw_parts(isv.host_ptr, RL_SLOTS_END) }; + let final_target = host[RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX]; + let final_ema = host[RL_BAND_TURNOVER_EMA_INDEX]; + + assert!( + final_target > initial_target + 1e-4, + "final_target = {final_target}, expected strictly > {initial_target} after 100 loosen steps (EMA={final_ema})" + ); + // Clamp ceiling is 0.20 — the controller must respect it. + assert!( + final_target <= 0.20 + TOL, + "final_target = {final_target}, exceeded clamp ceiling 0.20" + ); + eprintln!( + "PASS — controller loosens target from {initial_target} → {final_target} when EMA < 0.20 (final EMA={final_ema})" + ); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_backward_chain_propagates_to_encoder() -> Result<()> { + // Phase 4-B: with the band enabled and a non-zero turnover/target + // mismatch, the backward kernel must produce a non-zero `grad_h_t_d` + // so the encoder grad combiner has something to fold in. Validates + // that the per-batch grad on (b_l, b_u) is correctly chained through + // the `±|tanh|` activation + linear projection into the encoder's + // hidden-state gradient. + let Some(dev) = build_device() else { + return Ok(()); + }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 4usize; + let mut band = fresh_band(&dev, b_size)?; + + let (mut isv, isv_ptr) = build_isv(&stream)?; + isv_write(&mut isv, RL_BAND_ENABLED_INDEX, 1.0); + isv_write(&mut isv, RL_HEAT_CAP_MAX_LOTS_INDEX, 8.0); + isv_write(&mut isv, RL_BAND_LOSS_WEIGHT_INDEX, 0.1); // larger λ → larger grad + isv_write(&mut isv, RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX, 0.05); + isv_write(&mut isv, RL_BAND_GRAD_SHARPNESS_INDEX, 4.0); + + // Random h_t and run the forward to populate band_pre + band_out. + use rand::{Rng, SeedableRng}; + use rand_chacha::ChaCha8Rng; + let mut r = ChaCha8Rng::seed_from_u64(0xBA_C4_BAC4); + let h_t: Vec = (0..b_size * HIDDEN_DIM) + .map(|_| r.gen_range(-1.0..1.0)) + .collect(); + let h_t_d = upload_f32(&stream, &h_t)?; + band.forward(&h_t_d, &isv_ptr, b_size) + .context("band.forward (backward test)")?; + + // Stage a position state that places batches at varied positions + // so the sigmoid surrogates produce non-zero gradient at the + // boundaries (avoids the deep-interior saturation where surrogate + // derivatives are vanishingly small). + let positions: Vec = vec![3, 4, 5, 6]; + let pos_d = upload_pos(&stream, &positions)?; + + // turnover_t = 0.40 → diff = 0.35, large positive (over-trading) → grad_band ≠ 0. + let turnover_t: f32 = 0.40; + band.launch_turnover_loss(&pos_d, &isv_ptr, turnover_t, b_size, POS_BYTES) + .context("band.launch_turnover_loss (backward test)")?; + + // Backward — h_t is the input. + band.backward(&h_t_d, &isv_ptr, b_size) + .context("band.backward")?; + + let grad_h_t = read_slice_d_pub(&stream, &band.grad_h_t_d, b_size * HIDDEN_DIM)?; + let grad_b = read_slice_d_pub(&stream, &band.grad_b_per_b_d, b_size * BAND_OUT)?; + + // At least ONE element of grad_h_t must be non-trivially non-zero — + // signals the chain `grad_band → (sech² · sign) → W_band → h_t` is + // wired. Tolerance is loose since the magnitudes are small (≤ 1e-5) + // by construction (small λ, sigmoid surrogate at moderate boundary + // distance). + let max_abs = grad_h_t.iter().fold(0.0_f32, |acc, &x| acc.max(x.abs())); + assert!( + max_abs > 1e-9, + "max|grad_h_t| = {max_abs}, expected > 1e-9 (backward chain dead?)" + ); + + let max_abs_b = grad_b.iter().fold(0.0_f32, |acc, &x| acc.max(x.abs())); + assert!( + max_abs_b > 1e-9, + "max|grad_b| = {max_abs_b}, expected > 1e-9 (bias grad dead?)" + ); + eprintln!( + "PASS — band backward chain produces max|grad_h_t| = {max_abs}, max|grad_b| = {max_abs_b}" + ); + Ok(()) +}