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

feat(ml-alpha): Phase 4-A — No-transaction-band foundation (ISV slot 799 gate)

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-06-04 00:50:27 +02:00
parent bd811a7748
commit e41a732081
10 changed files with 2092 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
crates/ml-alpha/build.rs
View File

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

96
crates/ml-alpha/cuda/rl_band_head_forward.cu Normal file
View File

@@ -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 <stdint.h>
#include <math.h>
#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;
}

57
crates/ml-alpha/cuda/rl_band_mask.cu Normal file
View File

@@ -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 <stdint.h>
#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<const int*>(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;
}
}

134
crates/ml-alpha/cuda/rl_band_turnover_loss.cu Normal file
View File

@@ -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 <stdint.h>
#include <math.h>
#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<const int*>(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;
}

398
crates/ml-alpha/src/rl/band_head.rs Normal file
View File

@@ -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<CudaStream>,
raw_stream: CUstream,
// ── Kernel handles ────────────────────────────────────────────────
_forward_module: Arc<CudaModule>,
linear_fwd_fn: CudaFunction,
apply_activation_fn: CudaFunction,
_mask_module: Arc<CudaModule>,
mask_fn: CudaFunction,
_turnover_loss_module: Arc<CudaModule>,
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<f32>,
/// `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<f32>,
// ── Forward output buffers (persisted for diag) ───────────────────
/// `[B × 2]` pre-activation linear outputs.
pub band_pre_d: CudaSlice<f32>,
/// `[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<f32>,
// ── 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<f32>,
/// `[B]` — per-batch loss contribution (same scalar value, ×1/B).
pub loss_per_b_d: CudaSlice<f32>,
/// `[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<f32>,
}
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<Self> {
let stream: Arc<CudaStream> = 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<f32> = (0..BAND_OUT * n_in)
.map(|_| rng.gen_range(-scale..scale))
.collect();
let b_host: Vec<f32> = 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::<f32>(cfg.b_size * BAND_OUT)
.context("alloc band_pre_d")?;
let band_out_d = stream
.alloc_zeros::<f32>(cfg.b_size * BAND_OUT)
.context("alloc band_out_d")?;
// ── Turnover-loss per-batch scratch ──────────────────────────
let m_soft_per_b_d = stream
.alloc_zeros::<f32>(cfg.b_size)
.context("alloc m_soft_per_b_d")?;
let loss_per_b_d = stream
.alloc_zeros::<f32>(cfg.b_size)
.context("alloc loss_per_b_d")?;
let grad_band_per_b_d = stream
.alloc_zeros::<f32>(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<CudaStream> {
&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<f32>,
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<i32>,
pos_state_d: &CudaSlice<u8>,
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<u8>,
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<CudaStream>, host: &[f32]) -> Result<CudaSlice<f32>> {
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::<f32>(n)
.context("band_head upload alloc")?;
if n > 0 {
let nbytes = n * std::mem::size_of::<f32>();
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)
}

82
crates/ml-alpha/src/rl/isv_slots.rs
View File

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

1
crates/ml-alpha/src/rl/mod.rs
View File

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

219
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -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::<u8>().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<u8> = 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<u8> = 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<u8> = 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<u8> = 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,

414
crates/ml-alpha/tests/band_invariants.rs Normal file
View File

@@ -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<MlDevice> {
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<CudaStream>) -> 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<u8> {
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> {
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<CudaStream>, host: &[f32]) -> Result<CudaSlice<f32>> {
let mut d = stream.alloc_zeros::<f32>(host.len())?;
write_slice_f32_d_pub(stream, host, &mut d)?;
Ok(d)
}
fn upload_i32(stream: &Arc<CudaStream>, host: &[i32]) -> Result<CudaSlice<i32>> {
let mut d = stream.alloc_zeros::<i32>(host.len())?;
write_slice_i32_d_pub(stream, host, &mut d)?;
Ok(d)
}
fn upload_pos(stream: &Arc<CudaStream>, positions: &[i32]) -> Result<CudaSlice<u8>> {
let encoded = encode_pos_state(positions);
let mut d = stream.alloc_zeros::<u8>(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<f32> = (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<f32> = (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<i32> = vec![0; b_size];
let pos_d = upload_pos(&stream, &positions)?;
// Pre-mask: every batch holds a non-Hold action.
let actions_host: Vec<i32> = 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<f32> = (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<i32> = vec![5; b_size];
let pos_d = upload_pos(&stream, &positions)?;
let actions_host: Vec<i32> = 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<f32> = (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<i32> = 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<f32> = (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<i32> = vec![0; b_size];
let pos_d = upload_pos(&stream, &positions)?;
let actions_host: Vec<i32> = 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(())
}

690
docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md Normal file
View File

@@ -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<const int*>(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<const int*>(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.