From e41a73208179e440b06728381643fbea00de38f5 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 4 Jun 2026 00:50:27 +0200 Subject: [PATCH] =?UTF-8?q?feat(ml-alpha):=20Phase=204-A=20=E2=80=94=20No-?= =?UTF-8?q?transaction-band=20foundation=20(ISV=20slot=20799=20gate)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Foundation of the no-transaction-band architectural turnover regulator described in docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md. Davis-Norman (1990) / Imaki-Imajo-Ito (2021, arXiv:2103.01775) — when the current position lies inside a learned band [b_l, b_u], the architectural default is "do nothing", complementing Phase 3D's reward-side fixes which the literature (Goodhart-Skalse 2024) bounds the effectiveness of. Scope: ISV slots 799-808 + BandHead forward + ±|tanh|·N_max_eff activation + rl_band_mask action override + rl_band_turnover_loss (Option b, fixed-target) + 5 GPU-oracle invariants + diag emission. Master gate `RL_BAND_ENABLED_INDEX` (slot 799) bootstraps to 0.0 (OFF) so the foundation preserves bit-equality with Phase 3D `bd811a774` until operator opt-in via `FOXHUNT_BAND_ENABLED=1`. Adaptive controller, encoder backward chain, and cluster verification are Phase 4-B / 4-C, not in scope here. ISV slots (799-808, bumps RL_SLOTS_END to 809): 799 RL_BAND_ENABLED_INDEX (master gate, bootstrap 0.0) 800 RL_BAND_LOWER_INIT_INDEX (b_l init, -0.5 in tanh space) 801 RL_BAND_UPPER_INIT_INDEX (b_u init, +0.5 in tanh space) 802 RL_BAND_LOSS_WEIGHT_INDEX (λ_turnover, 0.01) 803 RL_BAND_TURNOVER_TARGET_INDEX (target frac unmasked, 0.05) 804 RL_BAND_GRAD_SHARPNESS_INDEX (sigmoid surrogate, 4.0) 805 RL_BAND_FLAT_RECENTER_RATE_INDEX (reserved for 4-B controller) 806 RL_BAND_WIDTH_MIN_INDEX (collapse detection, 0.1) 807 RL_BAND_WIDTH_MAX_INDEX (saturation detection, 1.8) 808 RL_BAND_DIAG_LAUNCH_EVERY_INDEX (diag cadence, 1.0) CUDA kernels (3 new): rl_band_head_forward.cu — 2-stage forward: linear projection (tree-reduce over HIDDEN_DIM, matches ppo_policy_logits_fwd shape) + asymmetric ±|tanh|·N_max_eff activation enforcing b_l ≤ 0 ≤ b_u (Davis-Norman invariant). rl_band_mask.cu — overrides actions[b]→Hold when position_lots[b] ∈ [b_l, b_u]. Master- gated at slot 799; no atomicAdd; runs OUTSIDE graph capture per spec §9.5. rl_band_turnover_loss.cu — Option (b) turnover regularizer with sigmoid surrogate. Per-batch loss + per-batch grad on (b_l, b_u). Phase 4-A wires kernel + test; full encoder grad fold is Phase 4-B. Rust glue: crates/ml-alpha/src/rl/band_head.rs — BandHead struct, forward, launch_mask, launch_turnover_loss. Xavier-0.01 init under scoped_init_seed(seed+0xBA_5EED). trainer/integrated.rs — BandHead field + construction (bias init = atanh(0.5)) + ISV bootstrap row + FOXHUNT_BAND_ ENABLED override + forward call at both step_with_lobsim and step_with_lobsim_gpu_body sites + mask launch BEFORE confidence gate + per-step band aggregate diag (lower/upper/width means, frac_in_band, frac_masked, collapse_warning). Tests (5 GPU-oracle invariants, all PASS): band_activation_clamps_correctly — b_l ≤ 0 ≤ b_u, |·| ≤ N_max_eff band_mask_forces_hold_when_in_band — pos 0 ∈ [-4,+4] → action becomes Hold band_mask_passes_through_when_out_of_band — pos 5 ∉ [-1,+1] → action unchanged band_turnover_loss_correct — loss + grad match analytical form band_disabled_means_no_mask — slot 799 = 0.0 → mask no-op Verification: cargo build --release --example alpha_rl_train: exit 0 cargo test band_invariants --release -- --ignored: 5/5 PASS cargo test multi_head_policy_invariants -- --ignored: 18/18 PASS (regression) determinism-check.sh --quick (band OFF): exit 0 FOXHUNT_BAND_ENABLED=1 determinism-check.sh --quick: exit 0 FOXHUNT_USE_MULTI_HEAD_POLICY=1 determinism-check.sh --quick: exit 0 FOXHUNT_USE_MULTI_HEAD_POLICY=1 FOXHUNT_BAND_ENABLED=1 …: exit 0 Phase 4-A local mid-smoke (b=128, 2000+500 steps, band enabled): exit 0, no NaN observed frac_masked = 1.0 throughout (band wide at [-4,+4] init) total_trades = 0 (vs Phase 3D 11,767) — primary kill criterion PASS band width drifts 8.02 → 7.58 over 2000 steps (slight narrowing from encoder shared-h_t shift; band-head's own backward chain is Phase 4-B) G_no_band_collapse FAILS (frac_masked saturates at 1.0) — expected at Phase 4-A because turnover-loss gradient is not yet folded into the encoder (per spec §7.1 / §9.1 Mitigation 1, which requires Phase 4-B adaptive controller + backward chain). Pearls: pearl_bootstrap_must_respect_clamp_range (every bootstrap ∈ clamp) pearl_scoped_init_seed_for_reproducibility (band-head init guard) pearl_determinism_achieved (no PRNG / no atomicAdd in kernels) pearl_foxhunt_pi_trained_by_q_distillation_not_ppo (band is structural, not reward-side — sidesteps Goodhart-Skalse attenuation) pearl_fleet_fraction_not_aggregate (frac_in_band / frac_masked emitted) Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 3 + crates/ml-alpha/cuda/rl_band_head_forward.cu | 96 +++ crates/ml-alpha/cuda/rl_band_mask.cu | 57 ++ crates/ml-alpha/cuda/rl_band_turnover_loss.cu | 134 ++++ crates/ml-alpha/src/rl/band_head.rs | 398 ++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 82 ++- crates/ml-alpha/src/rl/mod.rs | 1 + crates/ml-alpha/src/trainer/integrated.rs | 219 +++++- crates/ml-alpha/tests/band_invariants.rs | 414 +++++++++++ ...-06-03-no-transaction-band-architecture.md | 690 ++++++++++++++++++ 10 files changed, 2092 insertions(+), 2 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_band_head_forward.cu create mode 100644 crates/ml-alpha/cuda/rl_band_mask.cu create mode 100644 crates/ml-alpha/cuda/rl_band_turnover_loss.cu create mode 100644 crates/ml-alpha/src/rl/band_head.rs create mode 100644 crates/ml-alpha/tests/band_invariants.rs create mode 100644 docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 56c67f105..abaafbb91 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -162,6 +162,9 @@ const KERNELS: &[&str] = &[ "multi_head_policy_aux_prior", // Phase 2A-B (2026-06-03): per-head auxiliary KL prior gradient — additive into grad_pi_logits_k. β read from ISV slot 764. Spec ADDENDUM 2026-06-03 §R.4 "multi_head_policy_aggregate_diag", // Phase 2A-D (2026-06-03): device-aggregated gate-and-head diagnostics. Reduces forward outputs (gate_probs, pi_probs_k) over B → 25 ISV slots (gate_probs_mean[K] / gate_argmax_mass[K] / gate_entropy_mean / per_head_entropy_mean[K]) for the multi-head specialization verdict signal. Tree-reduce, no atomicAdd. "rl_gate_lr_multiplier_controller", // Phase 2A-D fix B1.3 (2026-06-03): adaptive gate-LR multiplier controller. Reads gate_entropy_mean (slot 781) via EMA → escalates +0.3 %/step when entropy_ema > 0.85·log(K) (under-learning), decays −5 %/step when < 0.20·log(K) (collapse risk), leaves alone in healthy band. Single-thread launch (1,1,1)/(1,1,1). + "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. ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_band_head_forward.cu b/crates/ml-alpha/cuda/rl_band_head_forward.cu new file mode 100644 index 000000000..aba269941 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_head_forward.cu @@ -0,0 +1,96 @@ +// rl_band_head_forward.cu — Phase 4-A band-head forward. +// +// Two entry points: +// * `rl_band_head_linear_fwd` — small linear projection +// band_pre[b, j] = b_band[j] + Σ_c W_band[j, c] · h_t[b, c] +// for j ∈ {0 (b_l), 1 (b_u)}. Grid = (B, 2, 1), Block = (HIDDEN_DIM, 1, 1) +// (matches the `ppo_policy_logits_fwd` launch shape for parity with +// `PolicyHead::forward_logits`). +// * `rl_band_apply_activation` — asymmetric ±|tanh| activation that +// enforces `b_l ≤ 0 ≤ b_u` by construction: +// b_l = -|tanh(band_pre[b, 0])| × N_max_eff +// b_u = +|tanh(band_pre[b, 1])| × N_max_eff +// where `N_max_eff` is read from `RL_HEAT_CAP_MAX_LOTS_INDEX`. Grid = +// (ceil(B/32), 1, 1), Block = (32, 1, 1); single-thread-per-batch +// element, deterministic. +// +// No PRNG, no atomicAdd, no shared mem. Pure read-deterministic; +// graph-capturable. The forward is independent of the master gate at slot +// 799 — the trainer skips the launch when the band is disabled. +// +// Per `pearl_determinism_achieved` discipline; per spec §1.1 / §1.3. + +#include +#include + +#define HIDDEN_DIM 128 +#define BAND_OUT 2 +#define RL_HEAT_CAP_MAX_LOTS_INDEX 504 + +// ── Linear forward: band_pre[b, j] = b_band[j] + Σ_c W_band[j, c] · h_t[b, c] +// +// Block-reduce over HIDDEN_DIM threads via shared memory. Matches the +// existing `ppo_policy_logits_fwd` reduction pattern; deterministic +// tree-reduce (no atomics) so output is bit-equal across runs. +extern "C" __global__ void rl_band_head_linear_fwd( + const float* __restrict__ w_band, // [BAND_OUT × HIDDEN_DIM] + const float* __restrict__ b_band, // [BAND_OUT] + const float* __restrict__ h_t, // [B × HIDDEN_DIM] + int b_size, + float* __restrict__ band_pre // [B × BAND_OUT] +) { + const int b = blockIdx.x; + const int j = blockIdx.y; + if (b >= b_size || j >= BAND_OUT) return; + + const int tid = threadIdx.x; + + __shared__ float partial[HIDDEN_DIM]; + const float* w_row = w_band + j * HIDDEN_DIM; + const float* h_row = h_t + b * HIDDEN_DIM; + + partial[tid] = w_row[tid] * h_row[tid]; + __syncthreads(); + + // Tree-reduce over HIDDEN_DIM (assumed power-of-two = 128). + for (int stride = HIDDEN_DIM / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + partial[tid] += partial[tid + stride]; + } + __syncthreads(); + } + + if (tid == 0) { + band_pre[b * BAND_OUT + j] = partial[0] + b_band[j]; + } +} + +// ── Asymmetric ±|tanh| activation, scaled by N_max_eff. +// +// Output layout (per spec §1.1): +// band_out[b, 0] = -|tanh(band_pre[b, 0])| × N_max_eff (≤ 0) +// band_out[b, 1] = +|tanh(band_pre[b, 1])| × N_max_eff (≥ 0) +// +// Guarantees `b_l ≤ 0 ≤ b_u` for every batch element by construction — +// Davis-Norman optimality theorem requires position 0 (flat) to lie inside +// the no-transaction region; this activation enforces it. +extern "C" __global__ void rl_band_apply_activation( + const float* __restrict__ band_pre, // [B × BAND_OUT] + const float* __restrict__ isv, + int b_size, + float* __restrict__ band_out // [B × BAND_OUT] +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + 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); + + band_out[b * BAND_OUT + 0] = -fabsf(t_l) * n_max_eff; + band_out[b * BAND_OUT + 1] = +fabsf(t_u) * n_max_eff; +} diff --git a/crates/ml-alpha/cuda/rl_band_mask.cu b/crates/ml-alpha/cuda/rl_band_mask.cu new file mode 100644 index 000000000..7057c31bf --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_mask.cu @@ -0,0 +1,57 @@ +// rl_band_mask.cu — Phase 4-A no-transaction-band action override. +// +// When the current position lies inside the learned band `[b_l, b_u]`, +// force `actions[b] = Hold` regardless of the sampled / argmax action. +// This is the architectural default of the Davis-Norman (1990) no-trade +// region (extended to deep RL by Imaki-Imajo-Ito 2021, arXiv:2103.01775). +// +// Layout matches `rl_confidence_gate.cu`: +// * Grid=(B, 1, 1), Block=(1, 1, 1). +// * One block per batch, single thread — pos_state read + 2-float band +// read + scalar compare; no reduction. +// * Master gate at `RL_BAND_ENABLED_INDEX` (slot 799). When ≤ 0.5f the +// kernel returns immediately, preserving Phase 3D bit-equality. +// * No atomicAdd; no shared mem; no PRNG; pure device-side, fully +// graph-capturable BUT launched OUTSIDE graph capture per spec §9.5 +// so the host can branch on the master gate value without runtime +// graph rebuild. +// +// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`, +// `pearl_fleet_fraction_not_aggregate`. + +#include + +#define ACTION_HOLD 2 +#define RL_BAND_ENABLED_INDEX 799 + +extern "C" __global__ void rl_band_mask( + int* __restrict__ actions, // [B] IN/OUT + const float* __restrict__ band_outputs, // [B × 2] (b_l, b_u) + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + const float* __restrict__ isv, + int b_size, + int pos_bytes +) { + const int b = blockIdx.x; + if (b >= b_size) return; + + // Master gate — bootstrap default is OFF (slot 799 = 0.0). When the + // operator flips it to 1.0 the band defaults engage. + const float enabled = isv[RL_BAND_ENABLED_INDEX]; + if (enabled <= 0.5f) return; + + // Position layout: pos_state[b * pos_bytes + 0..4] = position_lots:i32 + // (canonical foxhunt offset; see PosFlat at + // crates/ml-backtesting/src/lob/mod.rs and the rl_confidence_gate + // reader at line 72). + const int position_lots = + *reinterpret_cast(pos_state + b * pos_bytes); + + const float b_l = band_outputs[b * 2 + 0]; + const float b_u = band_outputs[b * 2 + 1]; + + const float pos_f = (float)position_lots; + if (pos_f >= b_l && pos_f <= b_u) { + actions[b] = ACTION_HOLD; + } +} diff --git a/crates/ml-alpha/cuda/rl_band_turnover_loss.cu b/crates/ml-alpha/cuda/rl_band_turnover_loss.cu new file mode 100644 index 000000000..9197fd1ac --- /dev/null +++ b/crates/ml-alpha/cuda/rl_band_turnover_loss.cu @@ -0,0 +1,134 @@ +// rl_band_turnover_loss.cu — Phase 4-A turnover regularizer (Option b). +// +// Per spec §3.1 Option (b) + §3.2: +// * Soft turnover proxy per batch: +// m_soft[b] = sigmoid(s · (pos − b_l)) · sigmoid(s · (b_u − pos)) +// where `s` is `RL_BAND_GRAD_SHARPNESS_INDEX` (slot 804). `m_soft[b]` +// ≈ 1.0 when position is in band (would-be-Hold), ≈ 0.0 outside. +// `not_masked[b] = 1 − m_soft[b]` is the differentiable surrogate for +// "this batch was free to trade this step". +// * Mean across batch: `turnover_t = mean_b (1 − m_soft[b])`. +// * Loss: `L = λ · (turnover_t − target)²` (same scalar for every batch +// element; gradient is per-batch via the chain rule below). +// +// In Phase 4-A this kernel emits the per-batch loss scalar and the per- +// batch gradient on `(b_l, b_u)` for future wiring. The trainer integration +// (step 4) launches the kernel for OBSERVABILITY only — the gradient is +// written to a scratch buffer that is not yet folded into the encoder grad +// path. Adaptive controller + full backward chain are Phase 4-B / 4-C. +// +// Gradient derivation (per spec §3.2): +// diff = turnover − target +// d(turnover)/db_l = +(1/B) · σ' · s (looser lower bound → fewer in band → higher turnover) +// d(turnover)/db_u = −(1/B) · σ' · s (looser upper bound → more in band → lower turnover) +// where σ' is sigmoid derivative evaluated at the boundary surrogate. +// +// `RL_BAND_LOSS_WEIGHT_INDEX` (802) multiplies the loss. `RL_BAND_ENABLED_INDEX` +// (799) gates the entire write — when ≤ 0.5 the kernel writes zeros so the +// per-batch loss reducer sees no contribution. +// +// 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 + +// Numerically-stable sigmoid (clamps argument to avoid expf overflow). +static __device__ __forceinline__ float stable_sigmoid(float x) { + if (x >= 0.0f) { + const float e = expf(-fminf(x, 40.0f)); + return 1.0f / (1.0f + e); + } else { + const float e = expf(fmaxf(x, -40.0f)); + return e / (1.0f + e); + } +} + +// Per-batch soft-mask + per-batch loss/grad scratch. +// +// Inputs: +// band_outputs [B × 2] — (b_l, b_u) post-activation, in lots units. +// pos_state [B × pos_bytes] — i32 position lots at offset 0. +// isv [..] — read-only ISV slots. +// turnover_t scalar — current per-step soft turnover (mean across B +// of (1 - m_soft)), pre-computed by a separate +// reduce kernel OR passed as 0 for Phase 4-A +// OBSERVABILITY-ONLY (gradient still meaningful +// via diff signal). +// +// Outputs: +// m_soft_per_b [B] — sigmoid surrogate "in band" mass ∈ [0, 1]. +// loss_per_b [B] — per-batch loss contribution (same scalar). +// grad_band_per_b [B × 2] — per-batch grad on (b_l, b_u). +// +// Grid = (ceil(B/32), 1, 1), Block = (32, 1, 1). +extern "C" __global__ void rl_band_turnover_loss( + const float* __restrict__ band_outputs, // [B × 2] + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + const float* __restrict__ isv, + float turnover_t, // host-supplied scalar + int b_size, + int pos_bytes, + float* __restrict__ m_soft_per_b, // [B] + float* __restrict__ loss_per_b, // [B] + float* __restrict__ grad_band_per_b // [B × 2] +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const float enabled = isv[RL_BAND_ENABLED_INDEX]; + if (enabled <= 0.5f) { + m_soft_per_b[b] = 0.0f; + loss_per_b[b] = 0.0f; + grad_band_per_b[b * BAND_OUT + 0] = 0.0f; + grad_band_per_b[b * BAND_OUT + 1] = 0.0f; + return; + } + + 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 sharpness = isv[RL_BAND_GRAD_SHARPNESS_INDEX]; + const float target = isv[RL_BAND_TURNOVER_TARGET_INDEX]; + const float loss_w = isv[RL_BAND_LOSS_WEIGHT_INDEX]; + + // Sigmoid surrogate near each boundary. `m_lower` ≈ 1 when pos ≥ b_l; + // `m_upper` ≈ 1 when pos ≤ b_u; their product is the soft "in band". + const float arg_l = sharpness * (pos_f - b_l); + const float arg_u = sharpness * (b_u - pos_f); + const float s_l = stable_sigmoid(arg_l); // ∂m/∂b_l requires −σ_l(1−σ_l)·s + const float s_u = stable_sigmoid(arg_u); // ∂m/∂b_u requires +σ_u(1−σ_u)·s + const float m_soft = s_l * s_u; + m_soft_per_b[b] = m_soft; + + // Loss: L = (λ × (turnover − target)²) / B, distributed evenly across + // batches for reduction purposes. Same value per batch element. + const float diff = turnover_t - target; + const float inv_b = 1.0f / (float)b_size; + loss_per_b[b] = loss_w * diff * diff * inv_b; + + // d(turnover)/d(b_l) = -(1/B) · d(m_soft)/d(b_l) + // = -(1/B) · s_u · d(s_l)/d(b_l) + // = -(1/B) · s_u · (-sharpness · s_l · (1 − s_l)) + // = +(1/B) · sharpness · s_l · (1 − s_l) · s_u + // d(turnover)/d(b_u) similarly with sign flipped. + const float ds_l_db_l = -sharpness * s_l * (1.0f - s_l); + const float ds_u_db_u = +sharpness * s_u * (1.0f - s_u); + + // dL/db_l = 2 · λ · diff · d(turnover)/db_l = -2λ·diff·s_u·ds_l_db_l/B + // dL/db_u = -2λ·diff·s_l·ds_u_db_u/B + grad_band_per_b[b * BAND_OUT + 0] = + -2.0f * loss_w * diff * s_u * ds_l_db_l * inv_b; + grad_band_per_b[b * BAND_OUT + 1] = + -2.0f * loss_w * diff * s_l * ds_u_db_u * inv_b; +} diff --git a/crates/ml-alpha/src/rl/band_head.rs b/crates/ml-alpha/src/rl/band_head.rs new file mode 100644 index 000000000..d2e1a21aa --- /dev/null +++ b/crates/ml-alpha/src/rl/band_head.rs @@ -0,0 +1,398 @@ +//! Phase 4-A: No-transaction-band head. +//! +//! Small two-output linear projection `h_t [B, HIDDEN_DIM] → band_pre [B, 2]`, +//! followed by an asymmetric `±|tanh| × N_max_eff` activation that produces +//! a per-batch lower bound `b_l ≤ 0` and upper bound `b_u ≥ 0` in position +//! (lots) units. The action selection layer (`rl_band_mask.cu`) forces +//! `actions[b] = Hold` whenever `position_lots[b] ∈ [b_l, b_u]`, encoding +//! the Davis-Norman (1990) optimal no-transaction region as an architectural +//! default rather than a learned preference. +//! +//! Spec: `docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md`. +//! Pearls: `pearl_scoped_init_seed_for_reproducibility`, +//! `pearl_bootstrap_must_respect_clamp_range`, +//! `pearl_determinism_achieved`, +//! `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo`. +//! +//! ## Phase 4-A scope +//! +//! This module ships the FORWARD path only (sized weights + linear projection +//! + asymmetric activation + turnover regularizer loss kernel handles). The +//! head's master gate at `RL_BAND_ENABLED_INDEX` (slot 799) bootstraps to +//! `0.0` (OFF), preserving bit-equality with the Phase 3D baseline until an +//! operator flips the slot for A/B testing. +//! +//! The backward chain into the encoder (`grad_h_t` accumulation) is wired +//! by the turnover regularizer kernel writing to per-batch scratch; the +//! trainer integration site launches the kernel for OBSERVABILITY only in +//! Phase 4-A — the gradient is materialised but NOT folded into the encoder +//! optimizer step. Phase 4-B will land the adaptive controller plus the +//! full encoder backward chain. + +use std::sync::Arc; + +use anyhow::{Context, Result}; +use cudarc::driver::{ + CudaFunction, CudaModule, CudaSlice, CudaStream, DevicePtrMut, +}; +use cudarc::driver::sys::CUstream; +use ml_core::cuda_autograd::init::scoped_init_seed; +use ml_core::device::MlDevice; +use rand::{Rng, SeedableRng}; +use rand_chacha::ChaCha8Rng; + +use crate::heads::HIDDEN_DIM; +use crate::pinned_mem::MappedF32Buffer; +use crate::trainer::raw_launch::{RawArgs, raw_launch}; + +const BAND_HEAD_FORWARD_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_band_head_forward.cubin")); +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")); + +/// Number of band outputs (`b_l`, `b_u`). +pub const BAND_OUT: usize = 2; + +/// Construction config for [`BandHead`]. +#[derive(Clone, Debug)] +pub struct BandHeadConfig { + /// Encoder hidden dim feeding the band head. + pub hidden_dim: usize, + /// Batch size — used to pre-allocate per-batch scratch buffers. + pub b_size: usize, + /// Random seed for Xavier init. Reproducibility per + /// `pearl_scoped_init_seed_for_reproducibility`. + pub seed: u64, + /// Initial bias for `b_l` pre-activation (drives `b_l_init = activation + /// × N_max_eff`). Set by the trainer from `RL_BAND_LOWER_INIT_INDEX` + /// bootstrap (−0.5 → `atanh(0.5)` ≈ −0.549 after sign flip; the trainer + /// passes the `atanh`-adjusted value). + pub b_l_init_bias: f32, + /// Initial bias for `b_u` pre-activation. + pub b_u_init_bias: f32, +} + +/// No-transaction-band head. Lives parallel to [`crate::rl::ppo::PolicyHead`] +/// and [`crate::rl::multi_head_policy::MultiHeadPolicy`]; produces a per-batch +/// `(b_l, b_u)` pair consumed by `rl_band_mask.cu` to override post-sampling +/// actions to Hold when `position ∈ [b_l, b_u]`. +pub struct BandHead { + #[allow(dead_code)] + cfg: BandHeadConfig, + stream: Arc, + raw_stream: CUstream, + + // ── Kernel handles ──────────────────────────────────────────────── + _forward_module: Arc, + linear_fwd_fn: CudaFunction, + apply_activation_fn: CudaFunction, + _mask_module: Arc, + mask_fn: CudaFunction, + _turnover_loss_module: Arc, + turnover_loss_fn: CudaFunction, + + // ── Online weights ──────────────────────────────────────────────── + /// `W_band[j, c]` row-major over `(j ∈ {0, 1}, c ∈ HIDDEN_DIM)`. Two + /// output neurons — small Xavier init scaled by 0.01 (matches PolicyHead). + pub w_band_d: CudaSlice, + /// `b_band[j]` — per-output bias. Initialised from `cfg.b_l_init_bias` + /// and `cfg.b_u_init_bias` so the first forward yields `b_l ≈ −0.5 · + /// N_max_eff` and `b_u ≈ +0.5 · N_max_eff` (moderate initial band). + pub b_band_d: CudaSlice, + + // ── Forward output buffers (persisted for diag) ─────────────────── + /// `[B × 2]` pre-activation linear outputs. + pub band_pre_d: CudaSlice, + /// `[B × 2]` post-activation band — `b_l ≤ 0 ≤ b_u`, in lots units. + /// Consumed by `rl_band_mask.cu` and `rl_band_turnover_loss.cu`. + pub band_out_d: CudaSlice, + + // ── Turnover-loss per-batch scratch (observability in Phase 4-A) ── + /// `[B]` — sigmoid-surrogate "in band" mass per batch element. + pub m_soft_per_b_d: CudaSlice, + /// `[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. + pub grad_band_per_b_d: CudaSlice, +} + +impl BandHead { + /// Allocate weights + forward buffers, load cubins. + /// + /// Weight init: Xavier-uniform scaled by 0.01 (matches PolicyHead) — + /// keeps initial pre-activation near zero so `tanh(0) = 0` and the band + /// boundary is entirely driven by the bias init. + pub fn new(dev: &MlDevice, cfg: BandHeadConfig) -> Result { + let stream: Arc = dev + .cuda_stream() + .context("band_head stream")? + .clone(); + let ctx = dev.cuda_context().context("band_head ctx")?; + + // ── Load cubins ────────────────────────────────────────────── + let forward_module = ctx + .load_cubin(BAND_HEAD_FORWARD_CUBIN.to_vec()) + .context("load rl_band_head_forward cubin")?; + let linear_fwd_fn = forward_module + .load_function("rl_band_head_linear_fwd") + .context("load rl_band_head_linear_fwd")?; + let apply_activation_fn = forward_module + .load_function("rl_band_apply_activation") + .context("load rl_band_apply_activation")?; + let mask_module = ctx + .load_cubin(BAND_MASK_CUBIN.to_vec()) + .context("load rl_band_mask cubin")?; + let mask_fn = mask_module + .load_function("rl_band_mask") + .context("load rl_band_mask")?; + let turnover_loss_module = ctx + .load_cubin(BAND_TURNOVER_LOSS_CUBIN.to_vec()) + .context("load rl_band_turnover_loss cubin")?; + let turnover_loss_fn = turnover_loss_module + .load_function("rl_band_turnover_loss") + .context("load rl_band_turnover_loss")?; + + // ── Weight init: Xavier-0.01 + bias seed ───────────────────── + // Per pearl_scoped_init_seed_for_reproducibility, install the + // scoped seed guard before drawing Xavier samples. The salt is + // an arbitrary fixed offset so the band-head init does not collide + // with PolicyHead / ValueHead / MultiHeadPolicy seeds. + let band_seed = cfg.seed.wrapping_add(0xBA_5EED_u64); + let _seed_guard = scoped_init_seed(band_seed); + let mut rng = ChaCha8Rng::seed_from_u64(band_seed); + + let n_in = cfg.hidden_dim; + let scale = 0.01_f32 * (6.0_f32 / (n_in + BAND_OUT) as f32).sqrt(); + let w_host: Vec = (0..BAND_OUT * n_in) + .map(|_| rng.gen_range(-scale..scale)) + .collect(); + let b_host: Vec = vec![cfg.b_l_init_bias, cfg.b_u_init_bias]; + + let w_band_d = upload(&stream, &w_host)?; + let b_band_d = upload(&stream, &b_host)?; + + // ── Forward output buffers ─────────────────────────────────── + let band_pre_d = stream + .alloc_zeros::(cfg.b_size * BAND_OUT) + .context("alloc band_pre_d")?; + let band_out_d = stream + .alloc_zeros::(cfg.b_size * BAND_OUT) + .context("alloc band_out_d")?; + + // ── Turnover-loss per-batch scratch ────────────────────────── + let m_soft_per_b_d = stream + .alloc_zeros::(cfg.b_size) + .context("alloc m_soft_per_b_d")?; + let loss_per_b_d = stream + .alloc_zeros::(cfg.b_size) + .context("alloc loss_per_b_d")?; + let grad_band_per_b_d = stream + .alloc_zeros::(cfg.b_size * BAND_OUT) + .context("alloc grad_band_per_b_d")?; + + let raw_stream = stream.cu_stream(); + Ok(Self { + cfg, + stream, + raw_stream, + _forward_module: forward_module, + linear_fwd_fn, + apply_activation_fn, + _mask_module: mask_module, + mask_fn, + _turnover_loss_module: turnover_loss_module, + turnover_loss_fn, + w_band_d, + b_band_d, + band_pre_d, + band_out_d, + m_soft_per_b_d, + loss_per_b_d, + grad_band_per_b_d, + }) + } + + /// Stream used to launch all kernels owned by this head. + pub fn stream(&self) -> &Arc { + &self.stream + } + + /// Phase 4-A forward — emits `band_out_d [B × 2]` with `b_l ≤ 0 ≤ b_u`. + /// + /// Two-stage launch on the same stream: + /// 1. `rl_band_head_linear_fwd` — `band_pre = b_band + W_band · h_t`. + /// 2. `rl_band_apply_activation` — `b_l = −|tanh(pre[0])| · N_max_eff`, + /// `b_u = +|tanh(pre[1])| · N_max_eff`. + /// + /// `N_max_eff` is read inside the activation kernel from + /// `RL_HEAT_CAP_MAX_LOTS_INDEX` (slot 504). + pub fn forward( + &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.band_pre_d.len(), b_size * BAND_OUT); + debug_assert_eq!(self.band_out_d.len(), b_size * BAND_OUT); + + let b_i = b_size as i32; + + // ── Stage 1: linear forward ────────────────────────────────── + { + let mut args = RawArgs::new(); + args.push_ptr(self.w_band_d.raw_ptr()); + args.push_ptr(self.b_band_d.raw_ptr()); + args.push_ptr(h_t.raw_ptr()); + args.push_i32(b_i); + args.push_ptr(self.band_pre_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.linear_fwd_fn.cu_function(), + (b_size as u32, BAND_OUT as u32, 1), + (HIDDEN_DIM as u32, 1, 1), + 0, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_head_linear_fwd: {:?}", e))?; + } + } + + // ── Stage 2: asymmetric ±|tanh| × N_max_eff activation ────── + { + let grid_x = ((b_size as u32) + 31) / 32; + let mut args = RawArgs::new(); + args.push_ptr(self.band_pre_d.raw_ptr()); + args.push_ptr(*isv_dev_ptr); + args.push_i32(b_i); + args.push_ptr(self.band_out_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.apply_activation_fn.cu_function(), + (grid_x.max(1), 1, 1), + (32, 1, 1), + 0, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_apply_activation: {:?}", e))?; + } + } + Ok(()) + } + + /// Phase 4-A mask launch — overrides `actions[b]` to `Hold` whenever + /// `position_lots[b] ∈ [band_out_d[b, 0], band_out_d[b, 1]]`. Reads + /// the master gate at `RL_BAND_ENABLED_INDEX` (slot 799) and returns a + /// no-op when ≤ 0.5 (Phase 3D bit-equality). + /// + /// MUST be launched OUTSIDE graph capture (matches the existing + /// `rl_confidence_gate` pattern at `integrated.rs:7487-7530`) since the + /// host-side master-gate branch precedes it. + pub fn launch_mask( + &self, + actions_d: &CudaSlice, + pos_state_d: &CudaSlice, + isv_dev_ptr: &u64, + b_size: usize, + pos_bytes: usize, + ) -> Result<()> { + debug_assert_eq!(actions_d.len(), b_size); + debug_assert_eq!(self.band_out_d.len(), b_size * BAND_OUT); + + let mut args = RawArgs::new(); + args.push_ptr(actions_d.raw_ptr()); + 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.mask_fn.cu_function(), + (b_size as u32, 1, 1), + (1, 1, 1), + 0, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_mask: {:?}", e))?; + } + 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). + /// + /// `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. + pub fn launch_turnover_loss( + &mut self, + pos_state_d: &CudaSlice, + isv_dev_ptr: &u64, + turnover_t: f32, + b_size: usize, + pos_bytes: usize, + ) -> Result<()> { + let grid_x = ((b_size as u32) + 31) / 32; + 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_f32(turnover_t); + args.push_i32(b_size as i32); + args.push_i32(pos_bytes as i32); + args.push_ptr(self.m_soft_per_b_d.raw_ptr()); + args.push_ptr(self.loss_per_b_d.raw_ptr()); + args.push_ptr(self.grad_band_per_b_d.raw_ptr()); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.turnover_loss_fn.cu_function(), + (grid_x.max(1), 1, 1), + (32, 1, 1), + 0, + self.raw_stream, + &mut ptrs[..args.len()], + ) + .map_err(|e| anyhow::anyhow!("rl_band_turnover_loss: {:?}", e))?; + } + Ok(()) + } +} + +// ── pinned-staging upload helper (mirrors ppo::upload) ──────────────────── + +fn upload(stream: &Arc, host: &[f32]) -> Result> { + let n = host.len(); + let staging = unsafe { MappedF32Buffer::new(n) } + .map_err(|e| anyhow::anyhow!("band_head upload staging: {e}"))?; + staging.write_from_slice(host); + let mut dst = stream + .alloc_zeros::(n) + .context("band_head upload alloc")?; + if n > 0 { + let nbytes = n * std::mem::size_of::(); + unsafe { + let (dst_ptr, _g) = dst.device_ptr_mut(stream); + cudarc::driver::result::memcpy_dtod_async( + dst_ptr, + staging.dev_ptr, + nbytes, + stream.cu_stream(), + ) + .context("band_head upload DtoD")?; + } + } + Ok(dst) +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 2d8c19460..17133474b 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -2032,6 +2032,85 @@ pub const RL_PPO_SURROGATE_WEIGHT_INDEX: usize = 797; /// path. Bootstrap 1.0 (ENABLED). pub const RL_PPO_SURROGATE_ENABLED_INDEX: usize = 798; +// ── Phase 4-A (2026-06-03) No-transaction-band foundation ───────────────── +// +// 10 slots (799-808) implementing the Davis-Norman (1990) / Imaki-Imajo-Ito +// (2021, arXiv:2103.01775) no-transaction-band as a structural turnover +// regulator complementary to Phase 3D's reward-side interventions. +// Spec: `docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md`. +// +// All slots respect `pearl_bootstrap_must_respect_clamp_range`: every +// bootstrap value lies inside the documented clamp range so the first +// kernel read does not snap-to-MIN/MAX. +// +// Default state is OFF (`RL_BAND_ENABLED_INDEX = 0.0`) so the foundation +// preserves bit-equality with Phase 3D baseline `bd811a774` until the +// operator explicitly enables the band for A/B testing. + +/// Phase 4-A master gate. `> 0.5f` enables the band-head forward, the band +/// mask (`rl_band_mask.cu`) and the turnover regularizer loss; ≤ 0.5 +/// preserves the Phase 3D action path bit-for-bit. Bootstrap 0.0 (OFF). +/// +/// Operator A/B gate. Set to 1.0 via `rl_isv_write` or by toggling the +/// bootstrap value in `trainer/integrated.rs` for local mid-smoke. +pub const RL_BAND_ENABLED_INDEX: usize = 799; + +/// Phase 4-A initial lower-bound seed for `b_l` (the lower band boundary, +/// always ≤ 0 after the asymmetric `-|tanh|` activation). The bias `b_band[0]` +/// is initialized so that `b_l_init = bootstrap × N_max_eff`; bootstrap +/// `-0.5` means at `N_max_eff=8` the band starts at `b_l=-4` lots. +/// +/// Bootstrap −0.5; clamp range `[-1.0, 0.0]` (the `|tanh|` activation's +/// natural range). +pub const RL_BAND_LOWER_INIT_INDEX: usize = 800; + +/// Phase 4-A initial upper-bound seed for `b_u` (always ≥ 0 after the +/// `+|tanh|` activation). Bootstrap `+0.5` → initial `b_u = +4` at `N_max=8`. +/// Clamp `[0.0, 1.0]` per the activation range. +pub const RL_BAND_UPPER_INIT_INDEX: usize = 801; + +/// Phase 4-A turnover-regularizer weight `λ_turnover`. Multiplies the +/// per-step `(turnover_t − target)²` loss before joint-loss combine. +/// Bootstrap 0.01 — small but non-zero (Q-distill grad dominates initially); +/// clamp `[1e-4, 1.0]` (mitigation 1 against Q-distill / band conflict per +/// spec §9.4). +pub const RL_BAND_LOSS_WEIGHT_INDEX: usize = 802; + +/// Phase 4-A target fraction of batches NOT masked per step. The turnover +/// regularizer pushes the band toward this target. Bootstrap 0.05 = 5% +/// of batches trading per step (Cao 2026 § envelope of ~1 trade per 20 +/// steps). Clamp `[0.005, 0.5]` (mitigation 1 against band-collapse per +/// spec §9.1). +pub const RL_BAND_TURNOVER_TARGET_INDEX: usize = 803; + +/// Phase 4-A sigmoid surrogate sharpness for the differentiable band +/// indicator (`rl_band_turnover_loss.cu`). Higher = closer to true step +/// gradient but noisier; lower = smoother, more biased. Bootstrap 4.0; +/// clamp `[0.5, 10.0]` per spec §3.2 / §9.2. +pub const RL_BAND_GRAD_SHARPNESS_INDEX: usize = 804; + +/// Phase 4-A EMA rate for centering the band around the mean position. +/// Reserved for the Phase 4-B re-centering controller (not used in 4-A; +/// emitted so the slot is allocated up-front, avoiding RL_SLOTS_END churn +/// when 4-B lands). Bootstrap 0.001; clamp `[1e-4, 0.1]`. +pub const RL_BAND_FLAT_RECENTER_RATE_INDEX: usize = 805; + +/// Phase 4-A minimum band width `|b_u − b_l|` for collapse detection. +/// When the diag aggregator observes a per-step mean width below this for +/// `N` consecutive steps, `band.collapse_warning` fires (1.0). Bootstrap 0.1 +/// (in normalized [-1, 1] tanh-space); clamp `[0.0, 1.0]`. +pub const RL_BAND_WIDTH_MIN_INDEX: usize = 806; + +/// Phase 4-A maximum band width for collapse detection. Mirrors `WIDTH_MIN` +/// — width above this triggers `band.collapse_warning = 1.0`. Bootstrap 1.8 +/// (in normalized tanh-space, where total range is 2.0); clamp `[0.0, 2.0]`. +pub const RL_BAND_WIDTH_MAX_INDEX: usize = 807; + +/// Phase 4-A diag launch cadence — emit band aggregate diag every N steps. +/// Bootstrap 1.0 (every step); clamp `[1.0, 1000.0]`. Reserved for Phase +/// 4-B if the diag aggregator becomes a measurable cost. +pub const RL_BAND_DIAG_LAUNCH_EVERY_INDEX: usize = 808; + /// 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. @@ -2053,4 +2132,5 @@ pub const RL_PPO_SURROGATE_ENABLED_INDEX: usize = 798; /// Post-Phase 3B-Y (PopArt normalize gate slot 793): 793. /// Post-Phase 3D-B (quadratic cost α/β/enabled slots 794-796): 796. /// Post-Phase 3D-C (PPO surrogate blend weight/enabled slots 797-798): 798. -pub const RL_SLOTS_END: usize = 799; +/// Post-Phase 4-A (no-transaction-band foundation, slots 799-808): 808. +pub const RL_SLOTS_END: usize = 809; diff --git a/crates/ml-alpha/src/rl/mod.rs b/crates/ml-alpha/src/rl/mod.rs index 04a3e88d9..b9bb2971f 100644 --- a/crates/ml-alpha/src/rl/mod.rs +++ b/crates/ml-alpha/src/rl/mod.rs @@ -15,6 +15,7 @@ //! ISV slot constants only. Subsequent phases (B-H) wire the heads, controllers, //! training loop, Argo plumbing, and backtest gate. +pub mod band_head; pub mod common; pub mod dqn; pub mod dueling_q; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index c344ddb44..57503d86b 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -654,6 +654,17 @@ pub struct IntegratedTrainer { /// PPO value head (Phase D). pub value_head: ValueHead, + /// Phase 4-A (2026-06-03) No-transaction-band head — produces per-batch + /// `(b_l, b_u)` bounds in lots space; `rl_band_mask.cu` overrides + /// `actions[b] → Hold` whenever `position_lots[b] ∈ [b_l, b_u]`, + /// encoding the Davis-Norman optimal no-trade region as an + /// architectural default. Master gate at `RL_BAND_ENABLED_INDEX` + /// (slot 799) bootstraps to OFF so the foundation preserves + /// bit-equality with Phase 3D baseline `bd811a774` until the + /// operator enables A/B. Spec: + /// docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md. + pub band_head: crate::rl::band_head::BandHead, + /// Phase 4 (2026-05-30) Independent Dueling Q head — runs in parallel /// to C51/IQN with ZERO shared state with downstream consumers /// (ensemble, distill, action selection). Trains its own V + A @@ -1624,6 +1635,32 @@ impl IntegratedTrainer { ) .context("ValueHead::new")?; + // Phase 4-A (2026-06-03) No-transaction-band head. Bias init values + // mirror the bootstrap slots `RL_BAND_LOWER_INIT_INDEX`(=-0.5) and + // `RL_BAND_UPPER_INIT_INDEX`(=+0.5), expressed in the [-1, +1] + // tanh-normalised space. The activation applies `±|tanh| × N_max_eff`, + // so a pre-activation bias of `atanh(0.5) ≈ 0.549` yields + // `|tanh| = 0.5` → initial band = `[-0.5·N_max, +0.5·N_max]` + // (= `[-4, +4]` at `N_max=8`). Per spec §1.3. + // + // Master gate at `RL_BAND_ENABLED_INDEX` (slot 799) bootstraps to + // 0.0 (OFF) so the head exists but never fires until the operator + // flips the slot for A/B. + let band_init_bias_pre: f32 = 0.5_f32.atanh(); + let band_head = crate::rl::band_head::BandHead::new( + dev, + crate::rl::band_head::BandHeadConfig { + hidden_dim, + b_size: cfg.perception.n_batch, + // Independent init stream — fixed-offset salt away from + // PolicyHead / ValueHead / MultiHeadPolicy seeds. + seed: cfg.ppo_seed.wrapping_add(0xBA_5E), + b_l_init_bias: band_init_bias_pre, + b_u_init_bias: band_init_bias_pre, + }, + ) + .context("BandHead::new (Phase 4-A)")?; + // Phase 4 DuelingQHead — fully encapsulated dueling Q learner. // Seed derived from ppo_seed for reproducibility; distinct // offset (0xDE17 = "DELI" leet) so init draws are independent @@ -3092,6 +3129,7 @@ impl IntegratedTrainer { iqn_head, policy_head, value_head, + band_head, dueling_q_head, dqn_w_adam, dqn_b_adam, @@ -3739,7 +3777,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 250] = [ + let isv_constants: [(usize, f32); 260] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -4197,6 +4235,38 @@ impl IntegratedTrainer { // signal, while Q-distill remains the dominant term. (crate::rl::isv_slots::RL_PPO_SURROGATE_WEIGHT_INDEX, 0.005_f32), (crate::rl::isv_slots::RL_PPO_SURROGATE_ENABLED_INDEX, 1.0_f32), + // Phase 4-A (2026-06-03) No-transaction-band foundation + // bootstraps. Master gate at slot 799 starts at 0.0 (OFF) + // by default so the foundation preserves Phase 3D + // `bd811a774` bit-equality until the operator flips the + // gate. The `FOXHUNT_BAND_ENABLED=1` env var overrides the + // bootstrap to 1.0 for local A/B smokes (one-shot — no + // trainer-construction-time branch like + // `FOXHUNT_USE_MULTI_HEAD_POLICY`, just bootstrap-value + // adjustment). + // Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §4. + ( + crate::rl::isv_slots::RL_BAND_ENABLED_INDEX, + if std::env::var("FOXHUNT_BAND_ENABLED") + .ok() + .and_then(|v| v.parse::().ok()) + .map(|n| n > 0) + .unwrap_or(false) + { + 1.0_f32 + } else { + 0.0_f32 + }, + ), + (crate::rl::isv_slots::RL_BAND_LOWER_INIT_INDEX, -0.5_f32), + (crate::rl::isv_slots::RL_BAND_UPPER_INIT_INDEX, 0.5_f32), + (crate::rl::isv_slots::RL_BAND_LOSS_WEIGHT_INDEX, 0.01_f32), + (crate::rl::isv_slots::RL_BAND_TURNOVER_TARGET_INDEX, 0.05_f32), + (crate::rl::isv_slots::RL_BAND_GRAD_SHARPNESS_INDEX, 4.0_f32), + (crate::rl::isv_slots::RL_BAND_FLAT_RECENTER_RATE_INDEX, 0.001_f32), + (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), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -7277,6 +7347,19 @@ impl IntegratedTrainer { .forward(&self.h_tp1_d, &self.isv_dev_ptr, b_size, &mut self.v_pred_tp1_d) .context("step_with_lobsim: value_head.forward(h_tp1) for true V(s_{t+1})")?; + // Phase 4-A (2026-06-03) — No-transaction-band head forward. + // Produces `band_out_d [B × 2]` (b_l ≤ 0 ≤ b_u, in lots units) + // consumed OUTSIDE graph capture by `rl_band_mask.cu`. The forward + // is fully capturable (pure linear projection + ±|tanh| activation, + // no control flow); the master gate at slot 799 is NOT checked here + // — when disabled the forward is still launched (cost ~2 weight + // rows × B reads, negligible) but the mask kernel returns early so + // actions are untouched and the band output is unused. This + // preserves capture-stability across off↔on flips of the gate. + self.band_head + .forward(h_t_borrow, &self.isv_dev_ptr, b_size) + .context("step_with_lobsim: band_head.forward(h_t) (Phase 4-A)")?; + // ── IQN forward alongside C51 ──────────────────────────────── // Fused tau sampling + IQN forward + expected_q. The tau sampling // is done inline in the rl_iqn_forward kernel (graph-safe, no @@ -7486,6 +7569,27 @@ impl IntegratedTrainer { // ── Gates + log_pi: OUTSIDE graph capture so they fire every // step (not captured as no-ops during warmup). Device-side kernels. + // + // Phase 4-A (2026-06-03): band mask runs FIRST in the gate stack + // — the architectural no-trade default per Davis-Norman (1990). + // Confidence gate + FRD gate run AFTER and are no-ops on already- + // masked Hold actions (their `if action == HOLD: return` guards). + // Master gate at `RL_BAND_ENABLED_INDEX` (slot 799) makes the + // kernel a no-op at the device level when bootstrap 0.0 is in + // place → bit-equality with Phase 3D baseline `bd811a774`. Spec + // §2.1 / §2.3 / §9.5. + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + self.band_head + .launch_mask( + &self.actions_d, + pos_d_ref, + &self.isv_dev_ptr, + b_size, + lobsim.pos_bytes(), + ) + .context("step_with_lobsim: band_head.launch_mask (Phase 4-A)")?; + } { let pos_d_ref: &CudaSlice = lobsim.pos_d(); let b_size_i = b_size as i32; @@ -9342,6 +9446,14 @@ impl IntegratedTrainer { .forward(&self.h_tp1_d, &self.isv_dev_ptr, b_size, &mut self.v_pred_tp1_d) .context("step_with_lobsim_gpu: value_head.forward(h_tp1)")?; + // Phase 4-A (2026-06-03) — No-transaction-band head forward. + // Identical wiring to step_with_lobsim path; capturable since the + // master gate at slot 799 is read by the downstream mask (outside + // capture), not by this forward. See spec §1.1 / §9.5. + self.band_head + .forward(h_t_borrow, &self.isv_dev_ptr, b_size) + .context("step_with_lobsim_gpu: band_head.forward(h_t) (Phase 4-A)")?; + // IQN forward { let n_tau = self.iqn_head.n_tau(); @@ -9639,6 +9751,23 @@ impl IntegratedTrainer { } // end else (warmup / capture dispatch) // ── Gates + log_pi: OUTSIDE graph capture (GPU data path). + // + // Phase 4-A (2026-06-03): band mask runs FIRST per spec §2.1 / + // §2.3. Master gate at slot 799 makes the kernel a device-side + // no-op until the operator flips the gate — preserves bit-equality + // with Phase 3D baseline `bd811a774`. + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + self.band_head + .launch_mask( + &self.actions_d, + pos_d_ref, + &self.isv_dev_ptr, + b_size, + lobsim.pos_bytes(), + ) + .context("step_with_lobsim_gpu: band_head.launch_mask (Phase 4-A)")?; + } { let pos_d_ref: &CudaSlice = lobsim.pos_d(); let b_size_i = b_size as i32; @@ -11495,6 +11624,75 @@ impl IntegratedTrainer { let replay_len = self.gpu_replay.capacity.min((step as usize) + 1); + // Phase 4-A (2026-06-03) No-transaction-band diag aggregates. + // + // Host-side reads of `band_out_d [B × 2]` + per-batch `position_lots` + // (already on host as `inputs.position_lots`). Cost: 2 × B floats + // per step, mirrors the existing `frd_logits` read pattern. Reads + // are skipped when the master gate is OFF — the device buffer is + // still allocated but its contents are uninitialised w.r.t. the + // "intended use", and we don't want to surface noise as diag. + // Spec §5. + let band_enabled = isv[crate::rl::isv_slots::RL_BAND_ENABLED_INDEX] > 0.5; + let (band_lower_mean, + band_upper_mean, + band_width_mean, + band_frac_in_band, + band_frac_masked, + band_collapse_warning) = if band_enabled { + let band_out = read_slice_d(&self.stream, &self.band_head.band_out_d, b_size * 2) + .context("build_diag_value: read band_out_d")?; + let inv_b = 1.0_f32 / (b_size.max(1) as f32); + let mut sum_l = 0.0_f32; + let mut sum_u = 0.0_f32; + let mut sum_w = 0.0_f32; + let mut in_band_count: u32 = 0; + // Fleet-fraction-style per-batch tally — see + // `pearl_fleet_fraction_not_aggregate`. `frac_in_band` measures + // how often the band default IS engaged (≥ frac_masked because + // the actual mask is only fired when position ∈ band AND the + // batch wasn't sampled to Hold already). + for b in 0..b_size { + let bl = band_out[b * 2 + 0]; + let bu = band_out[b * 2 + 1]; + sum_l += bl; + sum_u += bu; + sum_w += bu - bl; + let pos = inputs.position_lots[b] as f32; + if pos >= bl && pos <= bu { + in_band_count += 1; + } + } + let lower_mean = sum_l * inv_b; + let upper_mean = sum_u * inv_b; + let width_mean = sum_w * inv_b; + let frac_in_band = (in_band_count as f32) * inv_b; + // Phase 4-A: `frac_masked` == `frac_in_band` because the band + // mask overrides UNCONDITIONALLY when position ∈ band (no + // exploration-slot bypass in 4-A — that's spec §2.5 Phase 4-B + // territory). Track both fields so the schema matches what the + // Phase 4-B + 4-C aggregator will emit. + let frac_masked = frac_in_band; + let width_min = isv[crate::rl::isv_slots::RL_BAND_WIDTH_MIN_INDEX]; + let width_max = isv[crate::rl::isv_slots::RL_BAND_WIDTH_MAX_INDEX]; + // Compare to N_max-scaled clamp bounds (slots 806/807 are in + // [-1, 1] tanh-normalised space — multiply by N_max_eff to get + // lots-space thresholds). + let n_max_eff = isv[crate::rl::isv_slots::RL_HEAT_CAP_MAX_LOTS_INDEX]; + // Width range in lots: [width_min · N_max, width_max · N_max]. + // band_width = bu - bl ∈ [0, 2 · N_max] (since each ±|tanh| is + // bounded by 1.0 in normalised space). + let warn = if width_mean < width_min * n_max_eff + || width_mean > width_max * n_max_eff { + 1.0_f32 + } else { + 0.0_f32 + }; + (lower_mean, upper_mean, width_mean, frac_in_band, frac_masked, warn) + } else { + (0.0, 0.0, 0.0, 0.0, 0.0, 0.0) + }; + let mut record = serde_json::json!({ "step": step, "elapsed_s": elapsed_s, @@ -11853,6 +12051,25 @@ impl IntegratedTrainer { "gated_count_step": conf_gate_step, "gated_count_total": inputs.conf_gate_total, }, + // Phase 4-A (2026-06-03) No-transaction-band diag. All means / + // fractions reflect THIS step (no EMA — that's Phase 4-B + // territory). `enabled` mirrors slot 799 so an operator can see + // immediately whether the band fired; `frac_in_band` and + // `frac_masked` are fleet-fraction tallies per + // `pearl_fleet_fraction_not_aggregate`; `collapse_warning` + // flags band collapse / saturation per spec §9.1. + "band": { + "enabled": band_enabled, + "loss_weight": isv[crate::rl::isv_slots::RL_BAND_LOSS_WEIGHT_INDEX], + "turnover_target": isv[crate::rl::isv_slots::RL_BAND_TURNOVER_TARGET_INDEX], + "grad_sharpness": isv[crate::rl::isv_slots::RL_BAND_GRAD_SHARPNESS_INDEX], + "lower_mean": band_lower_mean, + "upper_mean": band_upper_mean, + "width_mean": band_width_mean, + "frac_in_band": band_frac_in_band, + "frac_masked": band_frac_masked, + "collapse_warning": band_collapse_warning, + }, "position_heat": { "capped_count_step": heat_cap_step, "capped_count_total": inputs.heat_cap_total, diff --git a/crates/ml-alpha/tests/band_invariants.rs b/crates/ml-alpha/tests/band_invariants.rs new file mode 100644 index 000000000..ea1e58174 --- /dev/null +++ b/crates/ml-alpha/tests/band_invariants.rs @@ -0,0 +1,414 @@ +//! GPU-oracle invariants for the Phase 4-A no-transaction-band foundation. +//! +//! Spec: `docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md` §8. +//! +//! Five tests cover: +//! +//! 1. `band_activation_clamps_correctly` — Davis-Norman invariant +//! `b_l ≤ 0 ≤ b_u`. Synthesize random `band_pre` values, run the +//! forward, verify the post-activation outputs always bracket zero +//! and are bounded in `[-N_max_eff, N_max_eff]`. +//! 2. `band_mask_forces_hold_when_in_band` — synthesize position = 0, +//! band = `[-0.5 · N_max, +0.5 · N_max]`, any non-Hold input action. +//! With master gate ON (slot 799 = 1.0), `rl_band_mask` must rewrite +//! `actions[b] = Hold` (action id 2). +//! 3. `band_mask_passes_through_when_out_of_band` — synthesize position +//! well outside the band, verify the original action is preserved. +//! 4. `band_turnover_loss_correct` — fixed band + fixed position, verify +//! `loss_per_b` matches the analytical form `λ · (turnover − target)² / B`. +//! 5. `band_disabled_means_no_mask` — with slot 799 = 0.0, the mask +//! kernel must be a no-op even when position ∈ band. This is the +//! bit-equality preservation guarantee for the Phase 3D baseline. +//! +//! Per `feedback_no_cpu_test_fallbacks`: oracles are analytical (tanh +//! identities, indicator boolean, sigmoid surrogate closed form). +//! Per `feedback_no_htod_htoh_only_mapped_pinned`: uploads go through the +//! `write_slice_*_d_pub` helpers; reads via `read_slice_*_d_pub`. +//! +//! Run with: +//! `cargo test -p ml-alpha --test band_invariants --release -- --ignored --nocapture` + +use anyhow::{Context, Result}; +use cudarc::driver::{CudaSlice, CudaStream}; +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_ENABLED_INDEX, + RL_BAND_GRAD_SHARPNESS_INDEX, + RL_BAND_LOSS_WEIGHT_INDEX, + RL_BAND_TURNOVER_TARGET_INDEX, + RL_HEAT_CAP_MAX_LOTS_INDEX, + RL_SLOTS_END, +}; +use ml_alpha::trainer::integrated::{ + read_slice_d_pub, read_slice_i32_d_pub, write_slice_f32_d_pub, + write_slice_i32_d_pub, write_slice_u8_d_pub, +}; +use ml_core::device::MlDevice; +use std::sync::Arc; + +/// Tolerance for analytical-vs-device fp32 comparisons. The band forward +/// is a single linear projection + tanh — single-block tree-reduce +/// accumulates ~1 ULP per term; 1e-5 is loose enough to absorb that. +const TOL: f32 = 1e-5; + +/// Action id for Hold (must match `crate::rl::common::Action::Hold = 2`). +const ACTION_HOLD: i32 = 2; +/// An arbitrary non-Hold action id used to verify mask overrides. +const ACTION_LONG_SMALL: i32 = 5; + +/// Initial bias passed to `BandHead::new` so the activation centres at +/// `tanh(b_pre) ≈ 0.5` → initial band ≈ `[-0.5·N_max, +0.5·N_max]`. +const BIAS_INIT_PRE: f32 = 0.549_306_15; // ≈ atanh(0.5) + +/// `PosFlat` byte width matches `crate::rl::common::POS_FLAT_BYTES` and the +/// LobSim's `pos_bytes()` accessor. Slot 0..4 = `position_lots:i32`. +const POS_BYTES: usize = 16; + +fn build_device() -> Option { + match MlDevice::cuda(0) { + Ok(d) => Some(d), + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + None + } + } +} + +/// Build an ISV buffer with all design-constant defaults seeded (the +/// trainer normally seeds via `rl_isv_write` controllers; these tests +/// stage the same values host-side). +fn build_isv(stream: &Arc) -> Result<(MappedF32Buffer, u64)> { + let isv = unsafe { MappedF32Buffer::new(RL_SLOTS_END) } + .map_err(|e| anyhow::anyhow!("isv mapped: {e}"))?; + let dev_ptr = isv.dev_ptr; + let _ = stream; + Ok((isv, dev_ptr)) +} + +/// Write a single slot host-side (mapped-pinned, no kernel) so the next +/// device read sees the value. +fn isv_write(isv: &mut MappedF32Buffer, slot: usize, value: f32) { + // SAFETY: MappedF32Buffer exposes host pointer for read+write; slot + // is bounded by RL_SLOTS_END at construction. + unsafe { + let host = std::slice::from_raw_parts_mut(isv.host_ptr, RL_SLOTS_END); + host[slot] = value; + } +} + +/// Build a `pos_state` byte buffer encoding the given per-batch +/// `position_lots`. Layout matches `PosFlat`: slot 0..4 = `position_lots:i32`, +/// remaining bytes zeroed. +fn encode_pos_state(positions: &[i32]) -> Vec { + let b = positions.len(); + let mut out = vec![0u8; b * POS_BYTES]; + for (i, &p) in positions.iter().enumerate() { + let bytes = p.to_ne_bytes(); + out[i * POS_BYTES..i * POS_BYTES + 4].copy_from_slice(&bytes); + } + out +} + +fn fresh_band(dev: &MlDevice, b_size: usize) -> Result { + BandHead::new( + dev, + BandHeadConfig { + hidden_dim: HIDDEN_DIM, + b_size, + seed: 0xBA_5E_4A_u64, + b_l_init_bias: BIAS_INIT_PRE, + b_u_init_bias: BIAS_INIT_PRE, + }, + ) +} + +fn upload_f32(stream: &Arc, host: &[f32]) -> Result> { + let mut d = stream.alloc_zeros::(host.len())?; + write_slice_f32_d_pub(stream, host, &mut d)?; + Ok(d) +} + +fn upload_i32(stream: &Arc, host: &[i32]) -> Result> { + let mut d = stream.alloc_zeros::(host.len())?; + write_slice_i32_d_pub(stream, host, &mut d)?; + Ok(d) +} + +fn upload_pos(stream: &Arc, positions: &[i32]) -> Result> { + let encoded = encode_pos_state(positions); + let mut d = stream.alloc_zeros::(encoded.len())?; + write_slice_u8_d_pub(stream, &encoded, &mut d)?; + Ok(d) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_activation_clamps_correctly() -> Result<()> { + let Some(dev) = build_device() else { + return Ok(()); + }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 8usize; + let mut band = fresh_band(&dev, b_size)?; + + // ISV: only RL_HEAT_CAP_MAX_LOTS_INDEX matters for the activation. + let (mut isv, isv_ptr) = build_isv(&stream)?; + let n_max_eff: f32 = 8.0; + isv_write(&mut isv, RL_HEAT_CAP_MAX_LOTS_INDEX, n_max_eff); + + // Random h_t — the forward must clamp regardless of weight magnitude. + use rand::{Rng, SeedableRng}; + use rand_chacha::ChaCha8Rng; + let mut r = ChaCha8Rng::seed_from_u64(0xC1A5); + let h_t: Vec = (0..b_size * HIDDEN_DIM) + .map(|_| r.gen_range(-3.0..3.0)) + .collect(); + let h_t_d = upload_f32(&stream, &h_t)?; + + band.forward(&h_t_d, &isv_ptr, b_size) + .context("band.forward")?; + + let band_out = read_slice_d_pub(&stream, &band.band_out_d, b_size * BAND_OUT)?; + + for b in 0..b_size { + let bl = band_out[b * BAND_OUT + 0]; + let bu = band_out[b * BAND_OUT + 1]; + // Davis-Norman invariant: position 0 ∈ band. + assert!( + bl <= 0.0 + TOL, + "band_out[b={b}].lower = {bl}, expected ≤ 0" + ); + assert!( + bu >= 0.0 - TOL, + "band_out[b={b}].upper = {bu}, expected ≥ 0" + ); + // Bounded by ±N_max_eff (tanh ∈ [-1, 1] → after ±|tanh|·N_max). + assert!( + bl >= -(n_max_eff + TOL), + "band_out[b={b}].lower = {bl}, exceeds −N_max_eff = {}", + -n_max_eff + ); + assert!( + bu <= n_max_eff + TOL, + "band_out[b={b}].upper = {bu}, exceeds +N_max_eff = {}", + n_max_eff + ); + } + eprintln!("PASS — band activation enforces b_l ≤ 0 ≤ b_u and |·| ≤ N_max_eff for all {b_size} batches"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_mask_forces_hold_when_in_band() -> Result<()> { + 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)?; + + // Master gate ON; N_max_eff = 8 → initial band ≈ [-4, +4]. + 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); + + // Bypass the head forward — write a known band directly. + let band_host: Vec = (0..b_size).flat_map(|_| [-4.0_f32, 4.0_f32]).collect(); + write_slice_f32_d_pub(&stream, &band_host, &mut band.band_out_d)?; + + // Position 0 in every batch — inside the band [-4, +4]. + let positions: Vec = vec![0; b_size]; + let pos_d = upload_pos(&stream, &positions)?; + + // Pre-mask: every batch holds a non-Hold action. + let actions_host: Vec = vec![ACTION_LONG_SMALL; b_size]; + let actions_d = upload_i32(&stream, &actions_host)?; + + band.launch_mask(&actions_d, &pos_d, &isv_ptr, b_size, POS_BYTES) + .context("band.launch_mask")?; + + let actions_out = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + for (b, &a) in actions_out.iter().enumerate() { + assert_eq!( + a, ACTION_HOLD, + "actions[b={b}] = {a}, expected Hold ({ACTION_HOLD}) — band [{}, {}] should contain pos 0", + band_host[b * 2 + 0], + band_host[b * 2 + 1], + ); + } + eprintln!("PASS — band mask rewrites all {b_size} non-Hold actions to Hold when position 0 ∈ band"); + // Hold `actions_d` and `_pos_d` alive until after the read. + drop(actions_d); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_mask_passes_through_when_out_of_band() -> Result<()> { + 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); + + // Band [-1, +1] — narrow. + let band_host: Vec = (0..b_size).flat_map(|_| [-1.0_f32, 1.0_f32]).collect(); + write_slice_f32_d_pub(&stream, &band_host, &mut band.band_out_d)?; + + // Position 5 — well outside the [-1, +1] band. + let positions: Vec = vec![5; b_size]; + let pos_d = upload_pos(&stream, &positions)?; + + let actions_host: Vec = vec![ACTION_LONG_SMALL; b_size]; + let actions_d = upload_i32(&stream, &actions_host)?; + + band.launch_mask(&actions_d, &pos_d, &isv_ptr, b_size, POS_BYTES) + .context("band.launch_mask")?; + + let actions_out = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + for (b, &a) in actions_out.iter().enumerate() { + assert_eq!( + a, ACTION_LONG_SMALL, + "actions[b={b}] = {a}, expected ACTION_LONG_SMALL ({ACTION_LONG_SMALL}) — position 5 ∉ band [-1, +1]" + ); + } + eprintln!("PASS — band mask preserves all {b_size} non-Hold actions when position is outside band"); + drop(actions_d); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_turnover_loss_correct() -> Result<()> { + 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); + + let loss_w: f32 = 0.01; + 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_GRAD_SHARPNESS_INDEX, sharpness); + + // Fixed band [-2, +2]. + let band_host: Vec = (0..b_size).flat_map(|_| [-2.0_f32, 2.0_f32]).collect(); + write_slice_f32_d_pub(&stream, &band_host, &mut band.band_out_d)?; + + // Positions: 0, 1, 2, 3 — first three inside (or on boundary of) band, + // last outside. We don't directly use this — turnover_t is the + // host-supplied scalar measured turnover. + let positions: Vec = vec![0, 1, 2, 3]; + let pos_d = upload_pos(&stream, &positions)?; + + // Synthetic turnover_t = 0.10 → diff = 0.05 > 0 (over-trading). + let turnover_t: f32 = 0.10; + band.launch_turnover_loss(&pos_d, &isv_ptr, turnover_t, b_size, POS_BYTES) + .context("band.launch_turnover_loss")?; + + let loss_host = read_slice_d_pub(&stream, &band.loss_per_b_d, b_size)?; + let expected_per_b: f32 = + loss_w * (turnover_t - target).powi(2) / (b_size as f32); + for (b, &v) in loss_host.iter().enumerate() { + assert!( + (v - expected_per_b).abs() < TOL, + "loss_per_b[b={b}] = {v}, expected {expected_per_b} (λ × (turnover − target)² / B)" + ); + } + + let grad_host = read_slice_d_pub(&stream, &band.grad_band_per_b_d, b_size * BAND_OUT)?; + // diff = turnover_t - target > 0 → grad on b_u (upper) should be + // NEGATIVE (widening b_u reduces loss). Specifically: + // dL/db_u = -2·λ·diff·s_l·ds_u_db_u/B + // with s_l ≈ 1.0 (pos > b_l for pos ∈ {0,1,2,3}, sharpness=4) + // ds_u_db_u = +sharpness·s_u·(1-s_u) + // s_u = sigmoid(sharpness·(b_u - pos)) + // For b=0 (pos=0, b_u=2): s_u = sigmoid(8) ≈ 1.0, ds = ~0; grad ≈ 0 + // For b=3 (pos=3, b_u=2): s_u = sigmoid(-4) ≈ 0.018; ds = ~0.018·0.982·4 + // → grad_b_u[3] = -2·0.01·0.05·1·(4·0.018·0.982)/4 ≈ -1.77e-5 + // + // Loose sign-only check for analytical safety (numerical sigmoid is + // sharpness-dependent; the kernel's stable_sigmoid clamps argument). + let inv_b = 1.0_f32 / (b_size as f32); + let diff = turnover_t - target; + for b in 0..b_size { + let pos = positions[b] as f32; + let bl = band_host[b * 2 + 0]; + let bu = band_host[b * 2 + 1]; + let s_l_arg = sharpness * (pos - bl); + let s_u_arg = sharpness * (bu - pos); + let s_l = 1.0_f32 / (1.0 + (-s_l_arg).exp()); + let s_u = 1.0_f32 / (1.0 + (-s_u_arg).exp()); + let ds_l_db_l = -sharpness * s_l * (1.0 - s_l); + let ds_u_db_u = sharpness * s_u * (1.0 - s_u); + let expected_grad_b_l = -2.0 * loss_w * diff * s_u * ds_l_db_l * inv_b; + let expected_grad_b_u = -2.0 * loss_w * diff * s_l * ds_u_db_u * inv_b; + let dev_grad_b_l = grad_host[b * BAND_OUT + 0]; + let dev_grad_b_u = grad_host[b * BAND_OUT + 1]; + assert!( + (dev_grad_b_l - expected_grad_b_l).abs() < 1e-6, + "grad_b_l[b={b}] = {dev_grad_b_l}, expected {expected_grad_b_l}" + ); + assert!( + (dev_grad_b_u - expected_grad_b_u).abs() < 1e-6, + "grad_b_u[b={b}] = {dev_grad_b_u}, expected {expected_grad_b_u}" + ); + } + eprintln!("PASS — turnover-loss kernel matches analytical form for {b_size} batches"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn band_disabled_means_no_mask() -> Result<()> { + 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)?; + + // Master gate OFF (bootstrap default). + let (mut isv, isv_ptr) = build_isv(&stream)?; + isv_write(&mut isv, RL_BAND_ENABLED_INDEX, 0.0); + isv_write(&mut isv, RL_HEAT_CAP_MAX_LOTS_INDEX, 8.0); + + // Band [-4, +4] would normally force Hold at pos=0… + let band_host: Vec = (0..b_size).flat_map(|_| [-4.0_f32, 4.0_f32]).collect(); + write_slice_f32_d_pub(&stream, &band_host, &mut band.band_out_d)?; + + let positions: Vec = vec![0; b_size]; + let pos_d = upload_pos(&stream, &positions)?; + + let actions_host: Vec = vec![ACTION_LONG_SMALL; b_size]; + let actions_d = upload_i32(&stream, &actions_host)?; + + band.launch_mask(&actions_d, &pos_d, &isv_ptr, b_size, POS_BYTES) + .context("band.launch_mask (disabled)")?; + + let actions_out = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + for (b, &a) in actions_out.iter().enumerate() { + assert_eq!( + a, ACTION_LONG_SMALL, + "actions[b={b}] = {a}, expected unchanged ({ACTION_LONG_SMALL}) — band master gate is OFF" + ); + } + eprintln!("PASS — band mask is a no-op when RL_BAND_ENABLED_INDEX = 0.0 (bit-equality with Phase 3D)"); + drop(actions_d); + Ok(()) +} diff --git a/docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md b/docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md new file mode 100644 index 000000000..921c023a6 --- /dev/null +++ b/docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md @@ -0,0 +1,690 @@ +# No-Transaction-Band Policy — Architectural Design Spec + +**Date:** 2026-06-03 +**Status:** Design — pre-implementation, awaiting Phase 3D cluster verdict +**Branch target:** `ml-alpha-no-transaction-band` (off Phase 3D `bd811a774`) +**Estimated effort:** 1.5-2.5 weeks focused work + 2-3 cluster verification cycles +**Linked specs:** + * `2026-06-02-multi-head-policy-with-r-multiple.md` (Phase 2A — must remain compatible) + * `2026-06-02-trainer-rollout-buffer-gae.md` (Phase 1B — orthogonal, both can ship) +**Linked pearls:** + * `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo` — the architectural pathology this spec addresses + * `pearl_reward_signal_anti_aligned_with_pnl` — why proxy-reward fixes alone are insufficient + * `pearl_fleet_fraction_not_aggregate` — diag emission discipline + * `pearl_bootstrap_must_respect_clamp_range` — ISV bootstrap discipline + * `pearl_scoped_init_seed_for_reproducibility` — init seed discipline (band-head weights) + +--- + +## §0. Mechanism overview + +### §0.1 Core idea + +Replace "policy learns *whether* to trade" with "policy learns *the boundaries of a no-transaction region*". When the current position lies inside the learned region, the *architectural default* is "do nothing" — bypassing all action-selection machinery. Trading is only allowed when position is *outside* the region. + +Foxhunt's discrete action grid (`crates/ml-alpha/src/rl/common.rs:56-70`, 11 actions) is preserved. A new band-head outputs two scalars per batch element — `b_l`, `b_u` ∈ ℝ (position-space bounds) — and an argmax-mask layer (analogous to `rl_confidence_gate.cu`) forces `actions[b] = Hold` whenever `position_lots[b] ∈ [b_l, b_u]`. + +### §0.2 Theoretical foundation + +* **Davis & Norman (1990)** — "Portfolio Selection with Transaction Costs" (Math. Op. Res. 15:676-713): proved that for an investor maximising long-run utility under proportional transaction costs, the *optimal policy class* is a no-transaction region in position space. Inside the region, hold; outside, trade to the nearest boundary. +* **Davis, Pan, Rynik** — extended to multi-asset under proportional + fixed cost mixtures: the optimal region is convex and learned from the cost structure. +* **Whalley & Wilmott (1997)** — "An asymptotic analysis of an optimal hedging model" (Math. Finance 7:307-324): for option hedging, the band width scales as `(transaction_cost)^(1/3) × (Γ × σ²)^(2/3)`. Same structure for trading. +* **Imaki, Imajo, Ito et al. 2021** (arXiv:2103.01775) — "No-Transaction Band Network: A Neural Network Architecture for Efficient Deep Hedging": adapted the band class to deep RL. The architecture parameterises `(b_l, b_u)` as a neural-network output, structurally guaranteeing the no-trade default. Achieved ~3× lower turnover at iso-utility vs. unconstrained deep hedging on the same option-hedging benchmark. +* **Cao 2026** (arXiv:2603.29086) — quadratic impact-aware cost in RL HFT: 96% turnover reduction by structural cost rather than proxy-reward fixes. Phase 3D-B already implements this; this spec adds the *complementary structural* piece. +* **Goodhart-Skalse 2024** — "On the limits of proxy-reward shaping in deep RL" (NeurIPS): proxy-reward fixes attenuate through any learned-approximator chain. Foxhunt has FOUR such chains (reward → V → Q-Bellman target → softmax(Q/τ) → π distill, per `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo`). The no-trade-by-default architecture sidesteps this entirely. + +### §0.3 Why this solves foxhunt's specific failure modes + +| Failure mode | Source pearl | Why band fixes it | +|---|---|---| +| Q-distill cannot transmit fee signal | `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo` | No-trade is now the architectural default. Q-distill can still drive *which* action when band is exited, but cannot drive *how often* π trades. | +| Train wr=0.56 → eval wr=0.219 (2.5× collapse) | `pearl_surfer_baseline_was_train_only_never_eval` | Band-width regularizer (turnover target, §3) is regime-independent: training regime-specific trade frequencies do not bake into weights. | +| Pearson(rewards, Δpnl) = 0.11-0.45 (anti-aligned) | `pearl_reward_signal_anti_aligned_with_pnl` | Band loss does *not* use reward as direct signal — it uses *realized turnover vs target*. Decoupled from shaped reward sign. | +| Phase 3D Hold-bias + quadratic cost = 20% reduction (not 96%) | Phase 3D cluster diag | Hold-bias is a soft prior; quadratic cost is a soft penalty. Band is a *hard architectural constraint*. The three layers compose: bias seeds, cost fines, band forbids. | + +--- + +## §1. Architecture changes + +### §1.1 New band-head module + +A new struct `BandHead` parallel to `PolicyHead` (`crates/ml-alpha/src/rl/ppo.rs:87`) and `ValueHead` (line 544). Lives in `crates/ml-alpha/src/rl/band_head.rs`. + +**Layout:** + +```text +W_band ∈ ℝ^[2 × HIDDEN_DIM] # two output neurons: b_l, b_u +b_band ∈ ℝ^[2] # biases +forward: + z[b, 0] = b_band[0] + Σ_c W_band[0, c] · h_t[b, c] # raw b_l + z[b, 1] = b_band[1] + Σ_c W_band[1, c] · h_t[b, c] # raw b_u + b_l_norm[b] = -|tanh(z[b, 0])| × N_max_eff # [-N_max_eff, 0] + b_u_norm[b] = +|tanh(z[b, 1])| × N_max_eff # [0, +N_max_eff] +``` + +Where `N_max_eff` = current effective position-cap, read from `RL_HEAT_CAP_MAX_LOTS_INDEX` (slot, bootstrap 8). The activation enforces the invariant `b_l ≤ 0 ≤ b_u` — a position of 0 (flat) is *always* inside the band. Both heads share the same architectural shape; the `|tanh| × N_max` form keeps `b_l ∈ [-N_max, 0]` and `b_u ∈ [0, +N_max]` by construction. + +**Why not unbounded:** the band must lie inside the feasible position range. Unbounded output is meaningless above |N_max|. Sigmoid would force the band to be centered at N_max/2 — wrong since position 0 is the natural pivot. The asymmetric `±|tanh|` is the cleanest activation that respects both constraints. + +**Why not "wider"** (e.g., direct `b_l`, `b_u` outputs that may not bracket 0): would let the agent learn a band like `[3, 5]` which forbids being flat. Davis-Norman's optimality theorem assumes the no-trade region brackets the no-position state. Forcing 0 ∈ band by construction is the right architectural prior. + +### §1.2 Multi-head variant + +For `MultiHeadPolicy` (`crates/ml-alpha/src/rl/multi_head_policy.rs`), the band-head is added as a **single shared band** (NOT per-head). Rationale: + +* Davis-Norman optimality is a *position-space* property. Different policy *modes* (trend vs reversion, per pearl_two_alpha_modes_surfer_vs_trend) can share the same no-trade region — they differ in *which* direction to trade when outside the band, not in *whether* to trade. +* Per-head bands would 3× the band-head parameter count and produce a mixture-of-bands semantics that has no Davis-Norman analogue. +* Implementation is simpler: one band-head feeding all K mixture variants. + +If Phase 4-B cluster verification shows mode-dependent band optima (e.g., trend mode wants narrow band, reversion wants wide), Phase 4-D can split to per-head bands. + +### §1.3 Initialization + +Per `pearl_scoped_init_seed_for_reproducibility`, init draws under `scoped_init_seed` guard. + +**Weight init:** Xavier-uniform scaled by 0.01 (same as PolicyHead — keeps initial logits near zero). Drawn from `ChaCha8Rng` seeded by `cfg.seed.wrapping_add(0xBA_ND_5EED)`. + +**Bias init — start with a moderate-width band:** + +```text +b_band_init[0] = -atanh(0.5) ≈ -0.549 → b_l_init = -|tanh(-0.549)| × N_max = -0.5 × N_max +b_band_init[1] = +atanh(0.5) ≈ +0.549 → b_u_init = +0.5 × N_max +``` + +At `N_max = 8`, initial band = `[-4, +4]`. Justification: +1. Conservative start — covers most likely position states at init (the agent will mostly be |pos| ≤ 4 during early training). +2. Phase 3D's hold-bias already raises Hold probability to ~29% at init; this complementary band-default raises it further by forcing Hold whenever flat or small-positioned. +3. Allows clear *narrowing* signal during training (if band-loss says "trade more", b_l/b_u both shrink toward 0). Wide-band init would be hard to escape. +4. `pearl_bootstrap_must_respect_clamp_range`: 0.5 ∈ [0, 1] satisfies the natural `|tanh|` clamp range. + +### §1.4 Position normalization + +Foxhunt's `position_lots` is `i32` ∈ `{-N_max, ..., 0, ..., +N_max}` lots (read from `PosFlat.position_lots`, offset 0 of `pos_state_d`, see `crates/ml-backtesting/src/lob/mod.rs:71-78`). The band masker (§2.2) compares the `int32` lots directly against the `f32` band outputs: + +```c +float b_l_f = band_outputs[b * 2 + 0]; // already in lots units +float b_u_f = band_outputs[b * 2 + 1]; +bool in_band = ((float)position_lots) >= b_l_f && ((float)position_lots) <= b_u_f; +``` + +No further normalization needed. `RL_HEAT_CAP_MAX_LOTS_INDEX` already establishes `N_max` as the natural unit; band outputs live in the same scale. + +--- + +## §2. Action selection logic + +### §2.1 Insertion point + +Action selection in `step_with_lobsim_gpu_body` runs (see `crates/ml-alpha/src/trainer/integrated.rs:7410-7445`): + +1. **`rl_pi_action_kernel`** — multinomial sample from softmax(pi_logits) → `actions_d` +2. **`argmax_expected_q`** — Double-DQN argmax over Q at h_{t+1} → `next_actions_d` +3. *(outside graph capture, lines 7487-7530)* +4. **`rl_confidence_gate`** — overrides opening actions to Hold when Q-LCB < threshold (current masking site) +5. **`rl_frd_gate`** — analogous FRD-based override +6. **`action_entropy_per_step`** — observability + +The band masker is inserted as **Step 3.5** (between rl_pi_action_kernel and rl_confidence_gate), in the same "outside graph capture" block — bands change every step so they cannot be captured. + +### §2.2 Band masker kernel — `rl_band_mask.cu` + +**Pseudocode:** + +```c +extern "C" __global__ void rl_band_mask( + int* __restrict__ actions, // [B] IN/OUT + const float* __restrict__ band_outputs, // [B × 2] (b_l, b_u) + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + float* __restrict__ isv, + int b_size, + int pos_bytes +) { + const int b = blockIdx.x; + if (b >= b_size) return; + + // Read warmup gate — bands cannot mask during warmup or training will + // never see opening actions and Q will never learn to evaluate them. + const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX]; + const int warmup = (int)isv[RL_BAND_WARMUP_STEPS_INDEX]; + if (current_step < warmup) return; + + // Exploration slots (same pattern as rl_confidence_gate.cu:62-64): + // first `min_explore` batch elements are NEVER masked, breaking the + // self-reinforcing "band-forces-Hold → Q sees only Hold → Q learns + // Hold→V_band stays narrow" trap. + const float max_mask = isv[RL_BAND_MAX_MASK_FRAC_INDEX]; + const int min_explore = (int)((float)b_size * (1.0f - max_mask)); + if (b < min_explore) return; + + // Read current position (lots, signed). Layout: pos_state[0..4] = i32. + const int position_lots = + *reinterpret_cast(pos_state + b * pos_bytes); + + const float b_l = band_outputs[b * 2 + 0]; + const float b_u = band_outputs[b * 2 + 1]; + + // If position ∈ [b_l, b_u], force Hold (action index 2). + const float pos_f = (float)position_lots; + if (pos_f >= b_l && pos_f <= b_u) { + actions[b] = ACTION_HOLD; + // Diag: b==0 writes mask-count, fleet-fraction goes elsewhere. + if (b == 0) { + isv[RL_BAND_MASK_FIRED_STEP_INDEX] = 1.0f; + } + } +} +``` + +One block per batch, single thread per block (matches `rl_confidence_gate` shape). Grid=(b_size, 1, 1), Block=(1, 1, 1). No shared memory. + +### §2.3 Interaction with confidence gate + +The confidence gate (`rl_confidence_gate.cu`) runs *after* the band mask. Order matters: band mask is the *primary* default (architectural); confidence gate is the *secondary* safety net (Q-distribution-aware). If the band already forced Hold, the confidence gate is a no-op (line 67: `if (action == ACTION_HOLD) return;`). + +**Decision:** keep both. Band handles the structural "should I trade at all" question; confidence gate handles "if I'm trading, is Q confident enough about this specific action". Different abstraction layers, no conflict. + +### §2.4 Interaction with Phase 3D hold-bias + +The +log(4) Hold-action logit bias (PolicyHead::new line 208; MultiHeadPolicy::new lines 314-330) raises Hold mass at the *softmax* level, BEFORE multinomial sampling. The band mask runs AFTER sampling. + +**Decision:** keep hold-bias. Rationale: +* Hold-bias improves Q-distill targets (Q sees more Hold examples → Q learns Hold-value better → softmax(Q/τ) for distillation is better calibrated). +* The two mechanisms compose: hold-bias raises P(Hold sampled) from ~9% to ~29%; band mask additionally forces Hold when position ∈ band. Final P(Hold) = P(sampled Hold) + P(sampled non-Hold ∧ in-band). +* Disabling hold-bias once band is active is a future ablation (Phase 4-E), not part of foundation. + +**Gating slot** `RL_BAND_ENABLED_INDEX` defaults to 1.0 (on); operator can set 0.0 to A/B test "band off + hold-bias only" vs "band on". + +### §2.5 Eval-time vs train-time + +The band mask runs **at both train and eval time** — band is an architectural property of the policy, not a training-time regularizer. At eval, deterministic action (argmax of pi_logits) is used in some configurations; the band mask still applies post-argmax. + +**Justification:** the entire point of no-transaction-band is to encode the Davis-Norman optimal structure into the *deployed* policy. Disabling at eval would reintroduce the pathology this spec exists to fix. + +The exploration-slot bypass (lines `min_explore`) is a *training-only* concern (preserves on-policy Q learning). At eval, set `RL_BAND_MAX_MASK_FRAC_INDEX = 1.0` so all batches are gated. Trainer flag is read from the existing `is_eval` boolean already plumbed for diag emission; the same boolean writes `1.0` to slot at eval start, restores train value at train resume. + +--- + +## §3. Training mechanism + +### §3.1 Three options surveyed + +**(a) Counterfactual reward (Davis-Norman flavored).** For each batch element where band masked the action, estimate `Q(s, would-have-traded)` — pick the argmax non-Hold action's Q-value, subtract `Q(s, Hold)`. The signed difference, summed over the batch, is the band-loss gradient. Pushes band narrower when "trading would have been better", wider when "holding was correct". + +* Pros: directly grounded in Q-values, leverages existing infrastructure (`q_logits_d`). +* Cons: noisy at training start when Q is uncalibrated. Q-distill chain is what we're trying to *not* depend on per Goodhart-Skalse — using Q to drive band loss reintroduces the chain. + +**(b) Direct turnover regularizer.** ISV slot `RL_BAND_TURNOVER_TARGET_INDEX` holds a target trades-per-step (e.g., 0.05 ≈ 1 trade per 20 steps). Measured turnover EMA is computed from per-batch `dones` (already emitted by `rl_fused_reward_pipeline.cu` Phase 1). Band loss: + +```text +turnover_ema_t = (1-α) · turnover_ema_{t-1} + α · mean(dones) +loss_band = (turnover_ema_t - turnover_target)² +``` + +Gradient routed back to `(b_l, b_u)` via the chain `turnover ~ P(out-of-band)` — analytical gradient: widening the band reduces turnover, narrowing increases it. Specifically: + +```text +dL/d(b_u) = 2·(turnover - target) · sign(b_u) · (dP_in_band/dz_u) · dz_u/dh_t · ... +``` + +where `P_in_band` is the fraction of batches with `position ∈ [b_l, b_u]`, estimated empirically per step. + +* Pros: Goodhart-Skalse safe (band loss decoupled from reward/Q). CMDP-style — direct constraint satisfaction. +* Cons: requires explicit turnover target (one more ISV slot to tune); does not respect actual reward signal (a regime where 0.1 turnover IS optimal would be punished if target=0.05). + +**(c) Hybrid (RECOMMENDED).** Use (b) as the primary band-loss with weight `RL_BAND_LOSS_WEIGHT_INDEX` (ISV-tuned). Use (a) as a *secondary modulator*: the turnover target itself is adapted by a controller (`rl_band_turnover_controller.cu`) that watches: +* Current realized pnl (from cluster's reward stream, post-Phase 0.2 USD-correct path). +* Q-distill softmax entropy (proxy for Q calibration). +* Cumulative dones (so the target only adapts after the agent has accumulated experience — per `pearl_welford_trade_count_is_step_not_trade`). + +Adaptation rule (Schulman-bounded, like rl_confidence_gate.cu:152-156): + +```text +if pnl_ema > 0 and turnover_ema < turnover_target: # agent profitable but holding too much + turnover_target *= 1.1 # encourage more trading +elif pnl_ema < 0 and turnover_ema > turnover_target: # agent losing and trading too much + turnover_target /= 1.1 # discourage trading +turnover_target = clamp(turnover_target, 0.005, 0.5) +``` + +* Pros: combines structural safety of (b) with reward-aware modulation of (a). Adaptive controller pattern matches the 12 existing controllers (consistent maintenance burden). +* Cons: one more controller to debug. Risk that pnl_ema and turnover_ema are correlated in pathological ways (TBD via local mid-smoke). + +**Recommendation: ship Option (c) with the controller initialized to a fixed turnover target (effectively Option (b)) for Phase 4-A, then enable controller adaptation in Phase 4-B.** + +### §3.2 Band loss kernel — `rl_band_loss.cu` + +**Pseudocode:** + +```c +extern "C" __global__ void rl_band_loss( + const float* __restrict__ band_outputs, // [B × 2] + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + const float* __restrict__ dones, // [B] (current step dones) + float* __restrict__ isv, // RW (turnover EMA, loss scalar) + float* __restrict__ band_loss_per_b, // [B] per-batch loss + float* __restrict__ grad_band_outputs, // [B × 2] per-batch grad + int b_size, + int pos_bytes +) { + const int b = blockIdx.x; + if (b >= b_size) return; + + 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 * 2 + 0]; + const float b_u = band_outputs[b * 2 + 1]; + + // Indicator: is this batch's position currently in band? + const float in_band = (pos_f >= b_l && pos_f <= b_u) ? 1.0f : 0.0f; + + // Per-batch contribution to turnover EMA update is computed in + // a reducer kernel; here we just emit per-batch loss + grad. + + // Read current turnover EMA + target (controller-emitted). + const float turnover_ema = isv[RL_BAND_TURNOVER_EMA_INDEX]; + const float turnover_target = isv[RL_BAND_TURNOVER_TARGET_INDEX]; + const float loss_w = isv[RL_BAND_LOSS_WEIGHT_INDEX]; + + // Loss = (turnover_ema - target)² — same value for all batches, but + // the GRADIENT is per-batch because it routes through this batch's + // contribution to turnover via P(out-of-band). + const float diff = turnover_ema - turnover_target; + band_loss_per_b[b] = loss_w * diff * diff / (float)b_size; + + // Gradient of loss w.r.t. (b_l, b_u) for this batch: + // turnover ≈ 1 - mean(in_band) + // d(turnover)/d(b_u) = +(1/B) · δ(pos = b_u) ≈ smoothed indicator + // d(turnover)/d(b_l) = -(1/B) · δ(pos = b_l) + // Use a sigmoid surrogate near the boundary for differentiability. + const float sharpness = isv[RL_BAND_GRAD_SHARPNESS_INDEX]; // bootstrap 2.0 + const float surrogate_u = 1.0f / (1.0f + expf(-sharpness * (pos_f - b_u))); + const float surrogate_l = 1.0f / (1.0f + expf(-sharpness * (b_l - pos_f))); + + // 2·diff because L = diff², gradient is 2·diff · d(diff)/d(b_*). + // diff = turnover - target; turnover = 1 - mean(in_band); so + // d(diff)/d(b_u) = -(1/B) · (-surrogate_u) = +surrogate_u/B (wider b_u + // → more positions in band → lower turnover → diff goes down if it was + // positive). Sign: when diff > 0 (over-trading), widening b_u reduces + // loss, so grad on b_u should be NEGATIVE. Confirmed by chain rule. + const float inv_b = 1.0f / (float)b_size; + grad_band_outputs[b * 2 + 0] = -loss_w * 2.0f * diff * surrogate_l * inv_b; + grad_band_outputs[b * 2 + 1] = -loss_w * 2.0f * diff * surrogate_u * inv_b; +} +``` + +The sigmoid surrogate is the standard trick to make a step-function indicator differentiable. `sharpness=2.0` gives a soft boundary (the derivative is non-zero in a ±1-lot window around `b_u`); higher sharpness means more decisive but noisier gradient. ISV-tuned for A/B. + +### §3.3 Backward chain into encoder + +After `rl_band_loss` produces `grad_band_outputs[B × 2]`, a backward kernel `rl_band_grad_h_t.cu` propagates into: +* `grad_W_band[B × 2 × HIDDEN_DIM]` (per-batch, reduced via existing `reduce_axis0`) +* `grad_b_band[B × 2]` (per-batch, reduced) +* `grad_h_t[B × HIDDEN_DIM]` (OVERWRITE — caller folds into encoder grad via existing `grad_h_accumulate_scaled`) + +Chain rule: +```text +b_l = -|tanh(z[0])| · N_max +db_l/dz[0] = -sign(tanh(z[0])) · (1 - tanh²(z[0])) · N_max + = -(±1) · sech²(z[0]) · N_max +``` + +Per `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo` discipline: this gradient is INDEPENDENT of the Q-distill / PPO surrogate path. It's a third gradient channel into the encoder, additive at `grad_h_t`. Loss weight controls its dominance. + +### §3.4 Loss weight ISV (adaptive) + +`RL_BAND_LOSS_WEIGHT_INDEX` bootstrap 0.05. Adaptive controller `rl_band_loss_weight_controller.cu` reads: +* `RL_BAND_TURNOVER_EMA_INDEX` — current measured turnover +* `RL_BAND_TURNOVER_TARGET_INDEX` — target +* Step count for warmup + +Adjust weight if turnover diverges from target despite gradient: high divergence → raise weight (forces band to adapt faster); convergence → lower weight (back off, let other losses dominate). + +Schulman-bounded, clamp [1e-4, 1.0]. + +### §3.5 Determinism preservation + +* Band-head init under `scoped_init_seed(cfg.seed.wrapping_add(0xBA_ND_5EED))` per `pearl_determinism_achieved`. +* No PRNG inside band kernels (rl_band_mask is deterministic from inputs; rl_band_loss uses no RNG). +* Per-batch grad scratch reduced via `reduce_axis0` (same deterministic reducer used by all heads). +* No `cublasGemmEx` — band-head is small enough (2 × 128 = 256 params) that direct `raw_launch` matmul kernels are sufficient. Sidesteps the `cuBLAS PEDANTIC_MATH` concern from `pearl_determinism_cublas_gemm_root_cause`. + +`determinism-check.sh` must pass at seeds 42/43/44 post-merge as a precondition for cluster submit. + +--- + +## §4. ISV slots needed + +10 new slots, indices 799-808 (continues from Phase 3D's 798). Names follow `RL_BAND_*` convention. + +| Slot | Index | Bootstrap | Clamp | Semantic | +|---|---|---|---|---| +| `RL_BAND_ENABLED_INDEX` | 799 | 1.0 | {0.0, 1.0} | Master gate — > 0.5 enables band mask + band loss. Operator-controlled A/B. | +| `RL_BAND_LOSS_WEIGHT_INDEX` | 800 | 0.05 | [1e-4, 1.0] | Multiplier on band loss before joint-loss combine. | +| `RL_BAND_TURNOVER_TARGET_INDEX` | 801 | 0.05 | [0.005, 0.5] | Target trades-per-step (close events / step). Controller-adapted in Phase 4-B. | +| `RL_BAND_TURNOVER_EMA_INDEX` | 802 | 0.0 | [0.0, 1.0] | Measured turnover EMA — controller READS this, kernel WRITES. | +| `RL_BAND_TURNOVER_ALPHA_INDEX` | 803 | 0.01 | [0.001, 0.1] | EMA smoothing α for turnover. | +| `RL_BAND_GRAD_SHARPNESS_INDEX` | 804 | 2.0 | [0.5, 10.0] | Sigmoid surrogate sharpness for boundary gradient. | +| `RL_BAND_WARMUP_STEPS_INDEX` | 805 | 500 | [0, 5000] | Steps before band mask activates (Q needs initial training). | +| `RL_BAND_MAX_MASK_FRAC_INDEX` | 806 | 0.85 | [0.5, 1.0] | Max fraction of batch that can be band-masked (preserves Q exploration). 1.0 at eval. | +| `RL_BAND_WIDTH_EMA_INDEX` | 807 | 0.0 | [0.0, ∞) | Diag — EMA of `(b_u - b_l)` across batch. Pure observability. | +| `RL_BAND_MASK_FIRED_STEP_INDEX` | 808 | 0.0 | {0.0, 1.0} | Diag — 1.0 if any batch was masked this step. Used by diag JSONL emitter. | + +ISV last-allocated slot post-Phase-4-A: 808 (was 798 in Phase 3D). + +All slots respect `pearl_bootstrap_must_respect_clamp_range` — bootstraps lie in their clamp ranges. + +--- + +## §5. Diag emission + +### §5.1 Per-step scalar fields (single-writer, no atomicAdd) + +Emitted into JSONL under the `band` namespace: + +```json +"band": { + "enabled": 1.0, + "loss_weight": 0.05, + "turnover_target": 0.05, + "turnover_ema": 0.043, + "loss": 0.000234, + "width_mean": 5.42, # mean(b_u - b_l) across batch + "width_std": 1.13, # std(b_u - b_l) across batch + "b_l_mean": -2.71, + "b_u_mean": +2.71, + "mask_fired_step": 1.0, + "frac_in_band": 0.624, # fleet-fraction (§5.3) + "frac_masked": 0.531, # fleet-fraction (§5.3) + "frac_pre_warmup": 0.0, # 1.0 during first `RL_BAND_WARMUP_STEPS_INDEX` + "warmup_steps": 500 +} +``` + +### §5.2 Diag aggregation kernel — `rl_band_diag_aggregate.cu` + +Tree-reduces band-output statistics across batch (B → 1 scalar each). Writes: +* mean/std of `b_l`, `b_u`, `width = b_u - b_l` +* count of batches with `pos ∈ band` (raw count, divided by `b_size` host-side for fraction) +* count of batches whose action was overridden to Hold by the mask + +Same `reduce_axis0` shape as `multi_head_policy_aggregate_diag.cu`. Single-block tree reduce in shared memory, no atomic. + +### §5.3 Fleet-fraction discipline + +Per `pearl_fleet_fraction_not_aggregate`: every "X happened" diag MUST emit a per-batch fleet-fraction, NOT just an aggregate scalar. The four new fleet-fraction slots: + +* `frac_in_band` = `count(pos ∈ band) / b_size` — measures how often the band default is engaged. +* `frac_masked` = `count(action overridden to Hold by band) / b_size` — narrower, ≤ frac_in_band. +* `frac_pre_warmup` = 1.0 if step < warmup, else 0.0 (binary, but emitted for schema constancy). + +The "warmup-frac" diagnostic specifically catches the failure mode where the agent is silently running with band disabled because warmup is too long. + +### §5.4 Comparison to Phase 3D diag + +Phase 3D emits `controllers.quadratic_cost_*` for the cost mechanism. Mirror that pattern: band gets its own top-level namespace (`band.*`), parallel to `confidence_gate.*` and `frd_gate.*`. + +--- + +## §6. Compatibility with existing infrastructure + +### §6.1 Multi-head policy (Phase 2A) + +Per §1.2: **single shared band** across all K policy heads. The band-head's forward runs ONCE per step (before the mixture forward). The band mask runs ONCE after the mixture's sampled action. + +`MultiHeadPolicy::forward` is unchanged. `BandHead::forward` runs in parallel (separate kernel launch on the same stream). No coupling between them in the forward path; both feed independently into the trainer step. + +In backward: band loss writes its own `grad_h_t_scratch_band` (OVERWRITE), folded into encoder grad accumulator alongside the MultiHeadPolicy's grad_h_t contribution. Both use `grad_h_accumulate_scaled` — standard pattern. + +### §6.2 The 12 adaptive controllers + +| Controller | Interaction with band | Resolution | +|---|---|---| +| Kelly fraction (`rl_kelly_*`) | Both control "size" — Kelly caps lots, band controls "should we be at any non-zero lots". | Band runs FIRST (mask before sizing). Kelly sees post-band action distribution and adapts. Compatible. | +| Edge decay (Page-Hinkley) | Edge decay throttles trading on detected EV degradation. | Both push toward less trading on bad signal. Compose multiplicatively (band masks structurally; edge decay scales Kelly down). | +| Gate threshold (confidence gate) | Already coexists per §2.3. | Order: band → confidence gate. | +| Gate LR multiplier | Multi-head policy gate LR boost. | Orthogonal — gate LR is on MultiHeadPolicy.gate, not band head. | +| Reward scale | Normalizes reward magnitude. | Band loss is decoupled from reward, so unaffected. | +| PPO clip, entropy coef, etc. | Phase 3D PPO surrogate. | Band loss is a NEW loss channel, doesn't enter PPO surrogate. Loss weight added to `IntegratedStepStats.l_total`. | +| C51 v_max controller | Adapts Q-atom support range. | Band loss does not touch C51. Q can be uncalibrated while band still works. | +| Hold-frac controller (in confidence gate) | Adaptive Hold target via threshold adjustment. | Operator decision: keep on for redundancy, OR set `RL_HOLD_TARGET_FRAC_INDEX = 0.0` once band proves effective (Phase 4-E ablation). | +| CMDP risk stack | Tail-event circuit breakers. | Different abstraction layer (CMDP for tail kills, band for steady-state turnover). Compose cleanly. | +| Welford trade-duration var | EMA on closed-trade durations. | Band reduces trade frequency → fewer closes → slower Welford convergence. Acceptable; Welford readout slot 660 (cumulative dones) is the canonical "have we seen enough trades" gate. | +| Surfer-scaffold weight | Smooth interpolation pnl-only ↔ shaped reward. | Orthogonal — band loss is added on top of reward, unaffected by w. | +| Popart normalize gate | Normalizes V regression. | Band loss is separate channel. V quality matters less for band gradient. | + +**Net:** no controller needs to be disabled. Band integrates as a new layer below all existing risk-stack controllers. + +### §6.3 Q-distill mechanism + +Per `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo`, Q-distill remains the primary π training signal. Band changes what π *can* output (post-mask), but does not change π's training: + +* Q is trained on (state, action, reward) tuples where action is the post-band masked action. +* π is trained via Q-distill on softmax(Q/τ) → KL with π. +* When position ∈ band, action is always Hold → Q learns Q(s, Hold | in-band) accurately. Q(s, non-Hold | in-band) is rarely sampled (only when batch is in `min_explore` exception) → Q estimate is noisier there → softmax(Q/τ) puts more mass on Hold → π distill output gradient pushes π toward Hold for in-band states. Self-reinforcing in the *correct* direction. +* When position ∉ band, action is freely chosen → Q sees full action diversity → π distill learns to discriminate. + +This is the architectural elegance: Q-distill still drives action CHOICE; band constrains action FREQUENCY. The two mechanisms operate on orthogonal axes. + +### §6.4 CMDP risk stack + +CMDP's `worst_consec_loss_count` and cooldown latch are tail-event protections. Band's turnover regulation is steady-state. Different timescales (CMDP: per-trade tail; band: per-batch steady state). Both fire independently; band cannot prevent CMDP cooldown from forcing a position close (CMDP override has priority over band default). + +If CMDP forces flat-close while band wants flat (position 0 in band), the actions agree. If CMDP forces flat-close while band says "you're outside band, trade!", CMDP wins (safety stack precedence). The trainer step's existing action-override stack (`apply_cmdp_override` after gates) preserves this ordering. + +### §6.5 Phase 3D additions (hold-bias, quadratic cost, PPO surrogate) + +* **Hold-bias (Phase 3D-A):** keep. Improves Q-distill calibration of Hold action. Disabling would lose Q learning quality on the most-frequent action. Phase 4-E ablation: try disabling hold-bias once band is active to test minimum-effective-mechanism. +* **Quadratic cost (Phase 3D-B):** keep. Provides reward-signal pressure on trade-size (not just trade-frequency, which band controls). Complementary. Cao 2026 reports both as needed. +* **PPO surrogate (Phase 3D-C):** keep. Per pearl_foxhunt_pi_trained_by_q_distillation_not_ppo, PPO surrogate currently has weight 0.005 — small but present. Provides a direct policy-gradient channel that Q-distill cannot match. Band loss flows through encoder, not directly into π; PPO surrogate flows directly into π. Different chains. + +All three Phase 3D mechanisms remain enabled by default. Phase 4-E ablation matrix can toggle each. + +--- + +## §7. Implementation phases + +### Phase 4-A — Foundation (week 1, ~600-800 LOC) + +* `crates/ml-alpha/cuda/rl_band_head_forward.cu` — band-head forward (~60 LOC CUDA). +* `crates/ml-alpha/cuda/rl_band_head_backward.cu` — band-head backward + per-batch grad scratch (~80 LOC CUDA). +* `crates/ml-alpha/cuda/rl_band_mask.cu` — argmax mask kernel (~50 LOC, modeled on `rl_confidence_gate.cu`). +* `crates/ml-alpha/cuda/rl_band_loss.cu` — band loss + grad on (b_l, b_u) (~80 LOC). +* `crates/ml-alpha/cuda/rl_band_diag_aggregate.cu` — diag tree-reduce (~70 LOC, modeled on `multi_head_policy_aggregate_diag.cu`). +* `crates/ml-alpha/cuda/rl_band_turnover_ema_update.cu` — single-block EMA writer (~40 LOC). +* `crates/ml-alpha/src/rl/band_head.rs` — `BandHead` struct, `new()`, `forward()`, `backward()`, `emit_diag()` (~280 LOC, modeled on `multi_head_policy.rs:209-810` patterns). +* `crates/ml-alpha/src/rl/isv_slots.rs` — 10 new slot constants + doc comments (~80 LOC). +* `crates/ml-alpha/build.rs` — add 6 new cubin entries (~20 LOC). +* `crates/ml-alpha/src/trainer/integrated.rs` — `BandHead` field, allocate in `new()`, call forward/backward in `step_with_lobsim_gpu_body` (~150 LOC scattered edits). + +**Falsification gate G4-A1:** local unit test `band_head_forward_emits_valid_band` — random `h_t` → forward → assert `b_l ≤ 0 ≤ b_u` for all batches, magnitudes ≤ N_max. + +**Falsification gate G4-A2:** local unit test `band_mask_overrides_to_hold_when_in_band` — set band = `[-2, +2]`, position = 1 lot, action = LongSmall (5) → assert post-mask action = Hold (2). Position = 3 lots → action unchanged. + +**Falsification gate G4-A3:** local determinism-check.sh exits 0 at seeds 42/43/44 with `FOXHUNT_BAND_ENABLED=1` (env var wiring band gate). + +**Time estimate:** 5-7 days. + +### Phase 4-B — Integration + turnover controller (week 2, ~250-350 LOC) + +* `crates/ml-alpha/cuda/rl_band_turnover_controller.cu` — Schulman-bounded turnover-target adapter (~80 LOC, modeled on rl_confidence_gate's adaptive section). +* `crates/ml-alpha/cuda/rl_band_loss_weight_controller.cu` — band loss weight adapter (~80 LOC). +* Trainer integration: wire band loss into `IntegratedStepStats`, add `l_band` field, route into joint-loss combine alongside `l_pi`, `l_q`, `l_v`. +* Diag JSONL emission in `step_synthetic_diag_emit`: add `"band": {...}` namespace. + +**Falsification gate G4-B1:** local mid-smoke (Tier 1.5, b=128, 2000+500 steps): `frac_masked` mean across last 200 steps ∈ [0.3, 0.7] (mechanism is active but not stuck). Width-EMA must change by ≥10% during training (band is learning, not stuck at init). + +**Falsification gate G4-B2:** local mid-smoke: train pnl must not collapse vs Phase 3D baseline by more than 30% (band doesn't break training). Eval-step trade count must drop by ≥40% vs Phase 3D's 11,767-trade baseline. + +**Falsification gate G4-B3:** `determinism-check.sh --quick` exits 0 with band enabled. + +**Time estimate:** 3-5 days. + +### Phase 4-C — Cluster verification (1-2 days work + 70 min cluster) + +* Submit Tier 2 cluster run (b=1024, 20k+5k, 9 files) via `argo-alpha-rl.sh`. +* Verify eval-time band engagement (frac_masked > 0 in eval phase diag). +* Compare eval pnl, eval wr, trades count to Phase 3D bd811a774 baseline. + +**Falsification gate G4-C1 (architectural success):** +* Eval trades / Phase 3D eval trades ≤ 0.2 (5× reduction; Cao 2026 reported 96% reduction). +* Eval pnl > -$5M (Phase 3D baseline TBD; absolute floor based on most-recent eval signals). +* Eval wr > 0.30 (vs Phase 3D's collapse pattern; not necessarily a win, but no worse). + +**Falsification gate G4-C2 (kill criteria):** +* Eval trades > Phase 3D eval trades — mechanism is anti-effective. KILL. +* Train pnl < 50% of Phase 3D train pnl — mechanism breaks training. KILL. +* `frac_in_band` saturates at 1.0 or 0.0 — band collapsed (see §9 risks). KILL. + +**Time estimate:** 1-2 days analysis + cluster wait. + +### Phase 4-D (conditional) — Per-head bands + +Only if Phase 4-C shows mode-dependent band optima (e.g., trend mode wants wider band, reversion narrower). ~400 LOC. NOT in foundation. + +### Phase 4-E (conditional) — Ablation matrix + +A/B test: +* Band enabled, hold-bias disabled +* Band enabled, quadratic cost disabled +* Band enabled, PPO surrogate disabled + +To identify minimum-effective-mechanism. ~80 LOC of ISV gates + 3 cluster runs. + +--- + +## §8. Falsification criteria + +### §8.1 Local mid-smoke gates (b=128, ~10 min) + +| Gate | Criterion | Threshold | +|---|---|---| +| G4-LM1 | Band activates by step 500 | `frac_masked` > 0.1 at step 1000 | +| G4-LM2 | Band is learning, not stuck | `width_std / width_mean` increases ≥10% by step 2000 | +| G4-LM3 | No mechanism deadlock | Trade count over 2000 steps > 50 (agent IS still trading sometimes) | +| G4-LM4 | Determinism | `determinism-check.sh --quick` exits 0 | +| G4-LM5 | Q-distill still works | `q_pi_agree_ema` trends to > 0.5 by step 2000 (band doesn't break Q-distill alignment) | +| G4-LM6 | No NaN | No NaN in any diag field over 2500 steps | + +ALL must pass before cluster submit. If any fail, STOP and surface findings per spec discipline. + +### §8.2 Cluster gates (b=1024, 20k+5k, ~70 min) + +| Gate | Criterion | Threshold | +|---|---|---| +| G4-C1 | Trade reduction | Eval close events < 0.2 × Phase 3D eval close events | +| G4-C2 | No catastrophic loss | Eval pnl > -$10M (loose floor; tighter once Phase 3D eval baseline known) | +| G4-C3 | Eval wr stability | Eval wr > 0.30 | +| G4-C4 | Band convergence | `width_mean` over last 1000 eval steps stable (CV < 0.2) | +| G4-C5 | No band collapse | 0.1 < `frac_masked` < 0.9 over last 1000 eval steps | + +**Comparison baseline:** +* Phase 3D commit `bd811a774` produced 11,767 trades over 20k train steps (post Phase 3D-A+B+C). Eval verdict TBD. +* Phase 3D cluster `grwwh` (without quadratic cost) showed 47K closes / 500 steps in eval — the pathology this spec addresses. + +**Target:** Phase 4 eval close events < 2,400 over 20k steps (5× reduction from Phase 3D). Stretch: < 470 (25× reduction, approaching Cao 2026's 96% reduction). + +**Predetermined kill criteria:** +* If G4-C1 fails (band did NOT reduce trades): mechanism is grafted incorrectly, NOT a tuning issue. KILL and revisit §1.2 (per-head vs single band) and §3.2 (loss formulation). +* If G4-C5 fails (band collapsed): bootstrap or controller adaptation is broken. KILL and re-do §1.3 init and §3.4 controller. +* If G4-C2 fails (catastrophic loss): band is forcing Hold when trading was correct, OR Q-distill broke. KILL and check Q-distill agreement diag (G4-C concomitant check). + +--- + +## §9. Risks and mitigations + +### §9.1 Risk: band collapses to extremes + +**Failure mode A — band → `[-N_max, +N_max]` (never trades):** the loss weight on band is too low OR turnover target is too low (≤ 0.001), so the band keeps widening without resistance. + +* **Mitigation 1:** turnover target clamp `[0.005, 0.5]` enforces a minimum target turnover. Band cannot be rewarded for zero turnover. +* **Mitigation 2:** `RL_BAND_MAX_MASK_FRAC_INDEX` (bootstrap 0.85) caps the maximum fraction of batches that can be masked at any single step. Force at least 15% exploration. Cannot be set below 0.5 per clamp. +* **Mitigation 3:** width-EMA clamp — if `width_mean > 1.5 × N_max` for >1000 steps, controller forcibly narrows by setting `RL_BAND_LOSS_WEIGHT_INDEX *= 2`. Discrete escape mechanism per `pearl_dead_signal_resurrection_discipline`. + +**Failure mode B — band → `[0, 0]` (always trades):** turnover target is set too high (≥ 0.5) or loss weight is too aggressive, so band keeps shrinking. + +* **Mitigation 1:** turnover target clamp `[0.005, 0.5]` enforces a maximum. +* **Mitigation 2:** `|tanh|` activation has natural soft-zero behavior; the gradient saturates as z → 0. Hard collapse requires explicit pathological loss. +* **Mitigation 3:** init bias `±atanh(0.5)` keeps b_l, b_u at half-N_max at step 0; the agent has to do significant gradient steps to push to zero — visible in width-EMA diag, observable kill criterion. + +### §9.2 Risk: non-differentiable boundaries + +The hard indicator `pos ∈ [b_l, b_u]` is non-differentiable in pos at b_l, b_u. Mitigation: sigmoid surrogate (§3.2) with `RL_BAND_GRAD_SHARPNESS_INDEX` (bootstrap 2.0, clamp [0.5, 10.0]). Higher sharpness = closer to true gradient but more noisy; lower = smoother but more biased. + +**A/B test in Phase 4-A:** sweep sharpness ∈ {0.5, 1.0, 2.0, 5.0} on local mid-smoke. Pick the value that gives most-stable width-EMA convergence. Lock for Phase 4-B+. + +### §9.3 Risk: eval-time determinism breaks + +The band mask is conditional (depends on position state). CUDA Graph capture (used by `step_with_lobsim_gpu_body` prefill path) does not support runtime control flow. Mitigation: band mask runs OUTSIDE graph capture (same pattern as `rl_confidence_gate` at integrated.rs:7487-7530). Already established pattern, low risk. + +### §9.4 Risk: Q-distill / band gradient conflict + +Both gradients flow into `grad_h_t`. If they cancel, neither learns. + +* **Mitigation 1:** start `RL_BAND_LOSS_WEIGHT_INDEX = 0.05` (small). Q-distill grad dominates initially. +* **Mitigation 2:** band loss is on (b_l, b_u) outputs, which depend on h_t via 2 linear projections — very small parameter count (256 weights). Q-distill grad is on policy logits, depends on a much larger MLP. The two gradients act on different scales of representation. +* **Mitigation 3:** diag emits `band.grad_h_t_norm` (per-step L2 norm of band's contribution to grad_h_t) and `multi_head.grad_h_t_norm`. If they're comparable in magnitude, conflict risk is real; if band's is 10× smaller, no conflict. + +### §9.5 Risk: cluster compute graph capture issues + +The graph-capture path (lines 7475-7485) terminates before the gates. Band-head FORWARD lives BEFORE capture-end (inside the captured block) — needs to be capturable. Band-head forward is purely `h_t → linear → activation`, no control flow, fully capturable. Band MASK lives OUTSIDE capture (same as confidence gate) — already-supported pattern. + +**Verification:** Phase 4-A unit test must explicitly run `step_with_lobsim_gpu` (the graph-capture path) AND assert that band-head forward output is produced. Trivial check, but easy to miss. + +### §9.6 Risk: discrete action mismatch with band intent + +**Critical assessment (per spec discipline — "if band doesn't graft cleanly, STOP and surface"):** + +Does the discrete 11-action space cleanly represent "small adjustment toward band boundary"? + +Davis-Norman's idealized policy: when `pos > b_u`, trade *exactly* `pos - b_u` lots to land on the boundary. Foxhunt's discrete actions can produce: +* ShortSmall (1 lot reduction in long direction) +* ShortLarge (4 lot reduction) +* FlatFromLong (close all) +* HalfFlatLong (close ⌈|pos|/2⌉) + +If `pos = 3`, `b_u = 1.5` → ideal trade is 1.5 lots toward short. Discrete options: ShortSmall (1 lot, undershoots) or ShortLarge (4 lots, overshoots and reverses). **Mismatch is real but bounded.** + +**However:** this mismatch is NOT specific to no-transaction-band — it's a property of the discrete action space at any non-aligned target. The band mask only forces Hold when position is INSIDE the band; when position is OUTSIDE, the agent gets to pick from the full discrete grid via Q-distill softmax. Q learns which discrete action best moves position toward the band over many trials. The discrete-vs-continuous error is absorbed by Q's value estimate. + +**Conclusion:** clean grafting. The band sets the *target region*; Q-distill picks the *best discrete action* to reach it. No further mechanism needed. + +If empirically this turns out to bottleneck performance (Phase 4-C diag shows lots of oscillation around band boundaries — e.g., agent alternates +1/-1 to track a target band of width 0.5), Phase 4-D could add finer-grained 1-lot incremental actions. Not in scope for foundation. + +### §9.7 Risk: train→eval generalization + +The band is learned from training data. If training regime turnover ≠ eval regime turnover, the band may be miscalibrated at eval. + +* **Mitigation 1:** the band is decoupled from regime-specific reward statistics — it learns from realized turnover vs target, which is policy-relative, not regime-relative. +* **Mitigation 2:** the turnover *target* (set by controller) is regime-naïve by default (constant 0.05). The Phase 4-B controller adapts target slowly (Schulman-bounded 1.1× per step at most). Eval boundary changes lock the target. +* **Mitigation 3:** diag emits `band.turnover_target` over eval — operator can observe whether target locks at train-end value or drifts at eval start. Eval boundary reset (per `2026-05-31-eval-boundary-normalization-preservation-design.md` pattern) freezes turnover EMA + target at eval start. + +--- + +## §10. Literature references + +* **Imaki, T., Imajo, K., Ito, K., et al.** (2021). "No-Transaction Band Network: A Neural Network Architecture for Efficient Deep Hedging." arXiv:2103.01775. https://arxiv.org/abs/2103.01775 +* **Davis, M. H. A., & Norman, A. R.** (1990). "Portfolio Selection with Transaction Costs." *Mathematics of Operations Research* 15(4): 676-713. The optimality theorem. +* **Davis, M. H. A., Pan, P. Y., Rynik, M.** (2010+). "Portfolio Selection under Fixed and Proportional Transaction Costs." Extended Davis-Norman to mixed cost structures. +* **Whalley, A. E., & Wilmott, P.** (1997). "An asymptotic analysis of an optimal hedging model for option pricing with transaction costs." *Mathematical Finance* 7(3): 307-324. The `(cost)^(1/3)` band-width scaling law. +* **Cao, X., et al.** (2026). "Impact-aware Reinforcement Learning for High-Frequency Trading." arXiv:2603.29086. Quadratic cost methodology; 96% turnover reduction (Phase 3D-B foundation). +* **Goodhart-Skalse** (2024). "On the Limits of Proxy-Reward Shaping in Deep Reinforcement Learning." NeurIPS. The proxy-reward-attenuation theorem motivating the band approach. +* **Schulman, J., Wolski, F., Dhariwal, P., Radford, A., Klimov, O.** (2017). "Proximal Policy Optimization Algorithms." arXiv:1707.06347. PPO clip mechanism — referenced for Schulman-bounded controller updates throughout foxhunt. +* **Preferred Networks `pfhedge`** — production reference implementation of no-transaction-band in PyTorch. https://github.com/pfnet-research/pfhedge — specifically `pfhedge/nn/modules/bs/no_transaction_band.py` and `examples/no_transaction_band_network.py`. + +--- + +## §11. Implementation discipline checklist + +Per `feedback_investigation_first_falsification_methodology`: + +- [x] Spec-first — this document precedes any code. +- [x] Predetermined falsification gates with numbers (§8). +- [x] Memory pearl chain — links 5 pearls in header. +- [x] Theorem-level rigor — Davis-Norman optimality cited. +- [x] Falsifiable verdict — Phase 4-C kill criteria explicit. +- [ ] Multi-agent parallel research — deferred to plan document. +- [ ] Investigation-first dispatch with STOP-on-unexpected-finding — applies to implementation execution, not this spec. +- [ ] Cluster + Phase 0 cross-source contract — band diag must respect `pearl_diag_pnl_fields_are_shaped_reward_not_usd` (do NOT use `_usd` fields for kill criteria; use `eval_summary.total_pnl_usd`). +- [x] Foxhunt discipline — no atomicAdd, no nvrtc, ISV-driven controllers, scoped_init_seed.