From 1c23ff368a0f8d7168863c406d21c5f8feffaadc Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 4 Jun 2026 08:17:28 +0200 Subject: [PATCH] =?UTF-8?q?feat(ml-alpha):=20Phase=204-A2=20=E2=80=94=20ba?= =?UTF-8?q?nd=20exploration=20bypass=20(slot=20813)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- crates/ml-alpha/cuda/rl_band_mask.cu | 15 ++++ crates/ml-alpha/src/rl/isv_slots.rs | 23 +++++- crates/ml-alpha/src/trainer/integrated.rs | 35 +++++++-- crates/ml-alpha/tests/band_invariants.rs | 92 ++++++++++++++++++++++- 4 files changed, 156 insertions(+), 9 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_band_mask.cu b/crates/ml-alpha/cuda/rl_band_mask.cu index 7057c31bf..44c5eeed0 100644 --- a/crates/ml-alpha/cuda/rl_band_mask.cu +++ b/crates/ml-alpha/cuda/rl_band_mask.cu @@ -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 diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 6d032bdde..082a40262 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -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; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 190568834..6123ddcd2 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -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 diff --git a/crates/ml-alpha/tests/band_invariants.rs b/crates/ml-alpha/tests/band_invariants.rs index 7ce0cdf5e..00ec48190 100644 --- a/crates/ml-alpha/tests/band_invariants.rs +++ b/crates/ml-alpha/tests/band_invariants.rs @@ -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 = (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 = (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 = (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 = 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 (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<()> {