From 65c328d3f1c1dab35598c247c05e27ab1a9f2461 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 4 Jun 2026 09:03:01 +0200 Subject: [PATCH] =?UTF-8?q?fix(ml-alpha):=20Phase=204-A3=20=E2=80=94=20ban?= =?UTF-8?q?d=20position-zero=20exception=20(opens-from-flat=20allowed)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- crates/ml-alpha/cuda/rl_band_mask.cu | 19 ++++++ crates/ml-alpha/tests/band_invariants.rs | 87 +++++++++++++++++++++--- 2 files changed, 97 insertions(+), 9 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_band_mask.cu b/crates/ml-alpha/cuda/rl_band_mask.cu index 44c5eeed0..fe1971528 100644 --- a/crates/ml-alpha/cuda/rl_band_mask.cu +++ b/crates/ml-alpha/cuda/rl_band_mask.cu @@ -62,6 +62,25 @@ extern "C" __global__ void rl_band_mask( const int position_lots = *reinterpret_cast(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]; diff --git a/crates/ml-alpha/tests/band_invariants.rs b/crates/ml-alpha/tests/band_invariants.rs index 00ec48190..54dc3564a 100644 --- a/crates/ml-alpha/tests/band_invariants.rs +++ b/crates/ml-alpha/tests/band_invariants.rs @@ -228,8 +228,11 @@ fn band_mask_forces_hold_when_in_band() -> Result<()> { let band_host: Vec = (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 = 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 = 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 = (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 = (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 = vec![0; b_size]; + let positions: Vec = vec![1; b_size]; let pos_d = upload_pos(&stream, &positions)?; let actions_host: Vec = 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 = (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 = vec![0; b_size]; + let pos_d = upload_pos(&stream, &positions)?; + + let actions_host: Vec = 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<()> {