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

feat(ml-alpha): Phase 4-A2 — band exploration bypass (slot 813)

Browse Source 
Adds RL_BAND_MAX_MASK_FRAC_INDEX (bootstrap 0.85) and modifies
rl_band_mask.cu to never mask the first (1 - max_mask) · B batches
per step. Resolves the dead-signal trap from Phase 4-B at 552d91bf4:
band saturated wide because all positions stayed at 0 (deep in
band ⇒ sigmoid surrogate gradient ~0 ⇒ no learning signal).

## Mechanism

* Slot 813 RL_BAND_MAX_MASK_FRAC_INDEX (bootstrap 0.85). RL_SLOTS_END=814.
* rl_band_mask.cu: deterministic first-N-batches bypass (preserves
  determinism per pearl_determinism_achieved).
* New test band_mask_respects_exploration_bypass; existing tests
  updated.

## Smoke result (FOXHUNT_BAND_ENABLED=1, seed 42, b=128, 2000+500)

The bypass mechanism works (15% of batches never see the mask).
The band IS learning (width 8.02 → 0.085 by step 1999, controller
drove turnover_target to MAX 0.20 as designed).

BUT: full policy collapse to bimodal `[0,0,109,0,0,0,0,0,19,0,0]`
from step 100 onward — 109 masked → Hold, 19 bypassed → all pick
TrailLoosen. action_entropy 0.45, total_trades=8.

## Architectural finding (Phase 4-A3 prerequisite)

The ±|tanh| activation guarantees b_l ≤ 0 ≤ b_u — position 0 is
ALWAYS in band. But foxhunt agents start flat (position 0), so
every batch starts masked, never opens trades, positions never
move, band has no exposure to non-zero positions, no gradient
signal to "open up". Davis-Norman semantics (band constrains
EXISTING positions) misapplied to "stay flat forever".

Phase 4-A3 will add a position-zero exception: band only
constrains already-positioned batches; opens from flat are
always allowed.

## Linked
* Phase 4-A: e41a73208 (band foundation)
* Phase 4-B: 552d91bf4 (controller + backward chain)
* Spec §2.2 + §9.1 Mitigation 2

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-06-04 08:17:28 +02:00
parent 552d91bf45
commit 1c23ff368a
4 changed files with 156 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
crates/ml-alpha/cuda/rl_band_mask.cu
View File

@@ -23,6 +23,7 @@
#define ACTION_HOLD 2
#define RL_BAND_ENABLED_INDEX 799
#define RL_BAND_MAX_MASK_FRAC_INDEX 813
extern "C" __global__ void rl_band_mask(
int* __restrict__ actions, // [B] IN/OUT
@@ -40,6 +41,20 @@ extern "C" __global__ void rl_band_mask(
const float enabled = isv[RL_BAND_ENABLED_INDEX];
if (enabled <= 0.5f) return;
// Phase 4-A2 exploration bypass — guarantee at least
// `(1 − max_mask) · B` batches are NEVER masked, so positions vary
// and the sigmoid surrogate gradient on (b_l, b_u) can flow into the
// band head. Without this the band collapses to the dead-signal trap
// observed at 552d91bf4 (Phase 4-B smoke): all positions stuck at 0,
// deep-in-band sigmoid argument, ~zero gradient, band stays wide.
//
// Bypass is DETERMINISTIC (lowest-index batches bypass) — random
// sampling would break `pearl_determinism_achieved` bit-equality.
// Spec §2.2 (kernel) + §9.1 Mitigation 2 (justification).
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;
// 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

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

@@ -2166,6 +2166,26 @@ pub const RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX: usize = 811;
/// `[0.0, 1.0]`.
pub const RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX: usize = 812;
/// Phase 4-A2 (2026-06-04) — Maximum fraction of batches that the band
/// mask may force to Hold at any single step. The first
/// `(1 − max_mask_frac) · B` batches always bypass the mask, guaranteeing
/// positions vary so the sigmoid surrogate gradient on (b_l, b_u) can
/// flow. Resolves the dead-signal trap surfaced in the Phase 4-B smoke at
/// `552d91bf4` where unconditional masking ⇒ all positions stuck at 0 ⇒
/// deep-in-band sigmoid argument ⇒ ~zero gradient ⇒ band stays wide
/// despite the controller correctly driving the turnover target to MAX.
///
/// Bypass is **deterministic** (the lowest-index batches always bypass)
/// to preserve `pearl_determinism_achieved` bit-equality. The frac
/// aggregator (slot 812) remains geometric — measuring band coverage,
/// not actual mask firings — because the controller adapts on the
/// architectural property (band coverage), not the diluted post-bypass
/// number. Bypass is an orthogonal exploration tax on the mask kernel.
///
/// Bootstrap 0.85 (15% guaranteed exploration per spec §9.1 Mitigation
/// 2). Clamp `[0.0, 1.0]` (1.0 = full masking, no bypass — eval mode).
pub const RL_BAND_MAX_MASK_FRAC_INDEX: usize = 813;
/// 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.
@@ -2189,4 +2209,5 @@ pub const RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX: usize = 812;
/// 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.
/// Post-Phase 4-B (adaptive turnover controller + backward chain, slots 809-812): 812.
pub const RL_SLOTS_END: usize = 813;
/// Post-Phase 4-A2 (band mask exploration bypass, slot 813): 813.
pub const RL_SLOTS_END: usize = 814;

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

@@ -3794,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); 264] = [
let isv_constants: [(usize, f32); 265] = [
// Static seeds for the adaptive reward-clamp controller —
// these are the initial values that
// `rl_reward_clamp_controller` will replace once it
@@ -4295,6 +4295,12 @@ impl IntegratedTrainer {
(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),
// Phase 4-A2 (2026-06-04) band mask exploration bypass.
// 0.85 = 15% of batches always escape masking → positions
// vary → sigmoid surrogate gradient flows → band can learn.
// Closes the dead-signal trap from Phase 4-B smoke at
// 552d91bf4. Spec §2.2 + §9.1 Mitigation 2.
(crate::rl::isv_slots::RL_BAND_MAX_MASK_FRAC_INDEX, 0.85_f32),
];
for (slot, value) in isv_constants.iter() {
let slot_i32 = *slot as i32;
@@ -11814,12 +11820,27 @@ impl IntegratedTrainer {
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;
// Phase 4-A2 (2026-06-04): `frac_masked` discounts the
// exploration bypass — the lowest-index `min_explore`
// batches NEVER get masked even when their position is in
// band. Verdict gate `frac_masked at step 1999 ∈ [0.3, 0.9]`
// depends on this being the post-bypass true mask rate, not
// the geometric `frac_in_band` (which now overstates).
let max_mask = isv[crate::rl::isv_slots::RL_BAND_MAX_MASK_FRAC_INDEX];
let min_explore = ((b_size as f32) * (1.0_f32 - max_mask)) as usize;
// Re-tally in-band only for batches >= min_explore (the masked
// population). Mirrors the kernel's deterministic-lowest-index
// bypass — bypass batches contribute 0 to actual mask count.
let mut masked_count: u32 = 0;
for b in min_explore..b_size {
let bl = band_out[b * 2 + 0];
let bu = band_out[b * 2 + 1];
let pos = inputs.position_lots[b] as f32;
if pos >= bl && pos <= bu {
masked_count += 1;
}
}
let frac_masked = (masked_count as f32) * inv_b;
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

92
crates/ml-alpha/tests/band_invariants.rs
View File

@@ -39,6 +39,7 @@ use ml_alpha::rl::isv_slots::{
RL_BAND_FRAC_NOT_MASKED_OBSERVED_INDEX,
RL_BAND_GRAD_SHARPNESS_INDEX,
RL_BAND_LOSS_WEIGHT_INDEX,
RL_BAND_MAX_MASK_FRAC_INDEX,
RL_BAND_TURNOVER_EMA_INDEX,
RL_BAND_TURNOVER_TARGET_ADAPTIVE_INDEX,
RL_HEAT_CAP_MAX_LOTS_INDEX,
@@ -214,9 +215,14 @@ fn band_mask_forces_hold_when_in_band() -> Result<()> {
let mut band = fresh_band(&dev, b_size)?;
// Master gate ON; N_max_eff = 8 → initial band ≈ [-4, +4].
// `max_mask = 1.0` disables the Phase 4-A2 exploration bypass so the
// test exercises the full-masking case (every batch eligible to be
// masked). The bypass behaviour is covered by
// `band_mask_respects_exploration_bypass` below.
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_MAX_MASK_FRAC_INDEX, 1.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();
@@ -242,7 +248,7 @@ fn band_mask_forces_hold_when_in_band() -> Result<()> {
band_host[b * 2 + 1],
);
}
eprintln!("PASS — band mask rewrites all {b_size} non-Hold actions to Hold when position 0 ∈ band");
eprintln!("PASS — band mask rewrites all {b_size} non-Hold actions to Hold when position 0 ∈ band (max_mask=1.0)");
// Hold `actions_d` and `_pos_d` alive until after the read.
drop(actions_d);
Ok(())
@@ -261,6 +267,8 @@ fn band_mask_passes_through_when_out_of_band() -> Result<()> {
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);
// No exploration bypass — exercise the mask code-path directly.
isv_write(&mut isv, RL_BAND_MAX_MASK_FRAC_INDEX, 1.0);
// Band [-1, +1] — narrow.
let band_host: Vec<f32> = (0..b_size).flat_map(|_| [-1.0_f32, 1.0_f32]).collect();
@@ -288,6 +296,88 @@ fn band_mask_passes_through_when_out_of_band() -> Result<()> {
Ok(())
}
#[test]
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
fn band_mask_respects_exploration_bypass() -> Result<()> {
// Phase 4-A2 (2026-06-04): the band mask must NEVER touch the first
// `min_explore = (1 − max_mask) · B` batches even when their position
// is inside the band. Without this exploration bypass the band
// collapses to the dead-signal trap observed at 552d91bf4 — masking
// forces position = 0, deep-in-band sigmoid argument, ~zero gradient,
// band stays wide forever.
//
// Setup: 100 batches, max_mask = 0.75 → min_explore = 25. All
// positions = 0 placed inside band [-0.5, +0.5]. Without bypass every
// batch would be rewritten to Hold; with bypass, the first 25
// preserve their original (non-Hold) action and the remaining 75 are
// rewritten to Hold.
//
// We deliberately use 0.75 (not the bootstrap value 0.85) because
// 0.75 is exactly representable in f32 (= 3/4), so `b_size × (1 −
// max_mask) = 100 × 0.25 = 25.0` rounds identically host-side and
// device-side. With 0.85 the f32 round-off makes `100 × 0.15` evaluate
// to `14.99999…` → floor = 14, an unobvious off-by-one. The
// bootstrap-default 0.85 path is exercised by the integration smoke
// (Step 6 of the spec).
let Some(dev) = build_device() else {
return Ok(());
};
let stream = dev.cuda_stream()?.clone();
let b_size = 100usize;
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_MAX_MASK_FRAC_INDEX, 0.75);
// Narrow band [-0.5, +0.5] so every position (= 0) sits inside.
let band_host: Vec<f32> = (0..b_size).flat_map(|_| [-0.5_f32, 0.5_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 (exploration bypass)")?;
let actions_out = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
// min_explore = (int)(100.0 × (1.0 − 0.75)) = (int)(25.0) = 25.
let min_explore: usize = ((b_size as f32) * (1.0 - 0.75)) as usize;
assert_eq!(min_explore, 25, "min_explore arithmetic guard (0.75 is exact in f32)");
let mut bypassed = 0usize;
let mut masked = 0usize;
for (b, &a) in actions_out.iter().enumerate() {
if b < min_explore {
assert_eq!(
a, ACTION_LONG_SMALL,
"actions[b={b}] = {a}, expected ACTION_LONG_SMALL ({ACTION_LONG_SMALL}) \
— batch b={b} < min_explore={min_explore} must bypass mask"
);
bypassed += 1;
} else {
assert_eq!(
a, ACTION_HOLD,
"actions[b={b}] = {a}, expected Hold ({ACTION_HOLD}) \
— batch b={b} ≥ min_explore={min_explore} must be masked (pos 0 ∈ [-0.5, +0.5])"
);
masked += 1;
}
}
assert_eq!(bypassed, 25, "expected 25 bypassed batches, got {bypassed}");
assert_eq!(masked, 75, "expected 75 masked batches, got {masked}");
eprintln!(
"PASS — exploration bypass: {bypassed}/{b_size} batches preserved, {masked}/{b_size} masked (max_mask=0.75)"
);
drop(actions_d);
Ok(())
}
#[test]
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
fn band_turnover_loss_correct() -> Result<()> {