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:
jgrusewski
2026-06-04 09:03:01 +02:00
parent 1c23ff368a
commit 65c328d3f1
2 changed files with 97 additions and 9 deletions

View File

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

View File

@@ -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<()> {