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