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