fix(ml-alpha): Phase 4-A3 — band position-zero exception (opens-from-flat allowed)
Davis-Norman (1990) is a theorem about MANAGING existing hedge
positions; it tells the agent NOT to micro-adjust within the band. It
says nothing about whether to open from flat.
The `±|tanh|` activation in rl_band_head_forward.cu guarantees
`b_l ≤ 0 ≤ b_u` (the Davis-Norman invariant). Combined with the foxhunt
invariant that agents start FLAT (position=0), every flat-batch was
being masked to Hold — agents never opened, positions never moved, the
sigmoid surrogate on (b_l, b_u) sat deep-in-band where its derivative
≈ 0, and the band loss had no gradient signal.
Verified empirically at 1c23ff368 (Phase 4-A2 smoke): action_hist
collapsed to bimodal [0,0,109,0,0,0,0,0,19,0,0] from step 100 onward,
total_trades = 8 over 2000 steps.
Fix (rl_band_mask.cu): immediately after the master-gate and
exploration-bypass checks, return when `position_lots == 0`. Flat
positions get free choice — opens from flat are NOT band-constrained.
Reads the same canonical offset as the existing band-position check.
Tests (band_invariants.rs):
* New: `band_mask_allows_opens_from_flat` — master gate ON, exploration
bypass DISABLED, all positions=0, band [-0.5, +0.5]; pre-A3 this
would mask every batch to Hold. Asserts the position-zero exception
preserves every batch's original (non-Hold) action.
* Updated: `band_mask_forces_hold_when_in_band` now uses position=1 in
band [-4, +4] (position=0 would unconditionally bypass post-A3).
* Updated: `band_mask_respects_exploration_bypass` now uses position=1
in band [-2, +2] so the exploration-bypass path can be exercised
without the position-zero exception pre-empting it.
11/11 band_invariants PASS. Determinism preserved across band-OFF,
band-ON, and multi-head+band-ON (`./scripts/determinism-check.sh
--quick` exit 0 in all three modes — the new check is integer-equality
and the master gate at slot 799 still defaults OFF, so bit-equality
with the Phase 3D baseline is intact when the band is disabled.
Single-seed smoke (FOXHUNT_BAND_ENABLED=1, seed 42, b=128):
* total_trades 8 → 1911 (trades restored)
* positions: 88/128 non-zero at step 1999 (was 0/128)
* band.width_mean trajectory 4.32 (step 100) → 7.37 (step 1999) —
band NOT collapsing now (was → 0.085 in 4-A2)
* band.frac_masked = 0.797 ∈ (0, 0.85) — band doing meaningful work
* eval pnl = -$793k vs -$8M kill threshold
* NaN count = 0
Falsification gates (Phase 4-A3):
* G_mechanism: PASS (exit 0, no NaN)
* G_positions_become_nonzero: PASS (88/128 non-zero)
* G_trade_recovery: PASS (1911 ∈ [500, 8000]) ← PRIMARY
* G_no_regression: PASS (-$794k > -$8M)
* G_band_active: PASS (frac_masked 0.797 ∈ (0, 0.85))
* G_action_diversity: FAIL (action_entropy 0.65 < 1.0) — see below
The action-diversity gate fails because the policy still concentrates
on Hold (action 2) + TrailLoosen (action 8). However this is NOT the
4-A2 "band can't train" collapse: trades flow, band trains, positions
move. The bimodal pattern toward TrailLoosen is a separate Q-distill
or controller side-effect that pre-existed the band — surfaced by the
spec's STOP-on-surprise rule #2 for follow-up rather than fixed here.
Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §10.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
@@ -62,6 +62,25 @@ extern "C" __global__ void rl_band_mask(
|
||||
const int position_lots =
|
||||
*reinterpret_cast<const int*>(pos_state + b * pos_bytes);
|
||||
|
||||
// Phase 4-A3 (2026-06-04): Davis-Norman position-zero exception.
|
||||
// The Davis-Norman (1990) no-transaction band is a theorem about
|
||||
// MANAGING an EXISTING hedge position — it tells the agent NOT to
|
||||
// micro-adjust within the band. It says nothing about whether to
|
||||
// open a position from flat.
|
||||
//
|
||||
// The `±|tanh|` activation in `rl_band_head_forward.cu` guarantees
|
||||
// `b_l ≤ 0 ≤ b_u` (the Davis-Norman invariant). Combined with the
|
||||
// foxhunt invariant that agents start FLAT (position = 0), every
|
||||
// single flat-batch would be masked to Hold — agents never open,
|
||||
// positions never move, the sigmoid surrogate on (b_l, b_u) sits
|
||||
// deep-in-band where its derivative ≈ 0, and the band loss has no
|
||||
// gradient signal. Verified empirically at 1c23ff368 (Phase 4-A2):
|
||||
// action_hist collapsed to bimodal [0,0,109,0,0,0,0,0,19,0,0] from
|
||||
// step 100 onward, total_trades = 8 over 2000 steps.
|
||||
//
|
||||
// Flat positions get free choice — opens are NOT band-constrained.
|
||||
if (position_lots == 0) return;
|
||||
|
||||
const float b_l = band_outputs[b * 2 + 0];
|
||||
const float b_u = band_outputs[b * 2 + 1];
|
||||
|
||||
|
||||
@@ -228,8 +228,11 @@ fn band_mask_forces_hold_when_in_band() -> Result<()> {
|
||||
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];
|
||||
// Phase 4-A3: position 0 ALWAYS bypasses the mask (Davis-Norman
|
||||
// position-zero exception — opens from flat are unconstrained). To
|
||||
// exercise the in-band masking path the test must use a NON-ZERO
|
||||
// position. Position 1 sits comfortably inside the [-4, +4] band.
|
||||
let positions: Vec<i32> = vec![1; b_size];
|
||||
let pos_d = upload_pos(&stream, &positions)?;
|
||||
|
||||
// Pre-mask: every batch holds a non-Hold action.
|
||||
@@ -243,12 +246,12 @@ fn band_mask_forces_hold_when_in_band() -> Result<()> {
|
||||
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",
|
||||
"actions[b={b}] = {a}, expected Hold ({ACTION_HOLD}) — band [{}, {}] should contain pos 1",
|
||||
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 (max_mask=1.0)");
|
||||
eprintln!("PASS — band mask rewrites all {b_size} non-Hold actions to Hold when position 1 ∈ band (max_mask=1.0)");
|
||||
// Hold `actions_d` and `_pos_d` alive until after the read.
|
||||
drop(actions_d);
|
||||
Ok(())
|
||||
@@ -307,7 +310,8 @@ fn band_mask_respects_exploration_bypass() -> Result<()> {
|
||||
// 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
|
||||
// positions = 1 (non-zero, so the Phase 4-A3 position-zero exception
|
||||
// doesn't fire) placed inside band [-2, +2]. 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.
|
||||
@@ -331,11 +335,13 @@ fn band_mask_respects_exploration_bypass() -> Result<()> {
|
||||
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();
|
||||
// Band [-2, +2] so position = 1 sits inside. Position 1 is non-zero
|
||||
// so the Phase 4-A3 position-zero exception does not pre-empt the
|
||||
// exploration-bypass path under test.
|
||||
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)?;
|
||||
|
||||
let positions: Vec<i32> = vec![0; b_size];
|
||||
let positions: Vec<i32> = vec![1; b_size];
|
||||
let pos_d = upload_pos(&stream, &positions)?;
|
||||
|
||||
let actions_host: Vec<i32> = vec![ACTION_LONG_SMALL; b_size];
|
||||
@@ -364,7 +370,7 @@ fn band_mask_respects_exploration_bypass() -> Result<()> {
|
||||
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])"
|
||||
— batch b={b} ≥ min_explore={min_explore} must be masked (pos 1 ∈ [-2, +2])"
|
||||
);
|
||||
masked += 1;
|
||||
}
|
||||
@@ -378,6 +384,69 @@ fn band_mask_respects_exploration_bypass() -> Result<()> {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn band_mask_allows_opens_from_flat() -> Result<()> {
|
||||
// Phase 4-A3 (2026-06-04): Davis-Norman position-zero exception.
|
||||
//
|
||||
// The Davis-Norman (1990) no-transaction band is a theorem about
|
||||
// managing an EXISTING hedge position. The `±|tanh|` activation
|
||||
// guarantees `b_l ≤ 0 ≤ b_u`, so position 0 is ALWAYS inside the
|
||||
// band by construction. Without this exception every flat-batch
|
||||
// (= the foxhunt starting state) gets masked to Hold, agents never
|
||||
// open, positions never move, the sigmoid surrogate sits deep-
|
||||
// in-band where its derivative is ~0, and the band loss gradient
|
||||
// dies. Verified empirically at 1c23ff368: action_hist collapsed
|
||||
// to [0,0,109,0,0,0,0,0,19,0,0], total_trades = 8 over 2000 steps.
|
||||
//
|
||||
// Setup: master gate ON, exploration bypass DISABLED (max_mask=1.0)
|
||||
// so every batch is eligible to be masked. All positions = 0, band
|
||||
// [-0.5, +0.5] (which would force-Hold every batch under the
|
||||
// pre-A3 kernel). Expectation: with the position-zero exception,
|
||||
// every batch's original (non-Hold) action is preserved.
|
||||
let Some(dev) = build_device() else {
|
||||
return Ok(());
|
||||
};
|
||||
let stream = dev.cuda_stream()?.clone();
|
||||
let b_size = 16usize;
|
||||
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);
|
||||
// Disable exploration bypass — isolate the position-zero exception.
|
||||
isv_write(&mut isv, RL_BAND_MAX_MASK_FRAC_INDEX, 1.0);
|
||||
|
||||
// Narrow band [-0.5, +0.5] — pre-A3 this would force-Hold every
|
||||
// batch since position 0 ∈ [-0.5, +0.5].
|
||||
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)?;
|
||||
|
||||
// Every batch is flat (position = 0) — the foxhunt starting state.
|
||||
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 (opens-from-flat)")?;
|
||||
|
||||
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}) \
|
||||
— Phase 4-A3 position-zero exception: opens from flat must NEVER be masked"
|
||||
);
|
||||
}
|
||||
eprintln!(
|
||||
"PASS — Phase 4-A3 position-zero exception: all {b_size}/{b_size} flat batches preserved (max_mask=1.0, band=[-0.5,+0.5])"
|
||||
);
|
||||
drop(actions_d);
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn band_turnover_loss_correct() -> Result<()> {
|
||||
|
||||
Reference in New Issue
Block a user