From e82049c7791b9b9bb1aeaea283f74df44f264d30 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 5 Jun 2026 01:28:32 +0200 Subject: [PATCH] =?UTF-8?q?feat(ml-alpha):=20Phase=207b=20F5=20Option=20A?= =?UTF-8?q?=20=E2=80=94=20gate-aware=20Hold=20suppression=20for=20confiden?= =?UTF-8?q?ce=20+=20FRD=20gates?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit When F5 state-conditional action availability mask (slot 823) is engaged AND the position is flat, F5 has deliberately forced the agent into an opening action by masking Hold (and other state-invalid actions) in `pi_logits` pre-sample. Two downstream gates — `rl_confidence_gate` (slot 512..) and `rl_frd_gate` (slot 516..) — silently overrode the sampled action back to Hold at the same flat-state, neutralizing F5's surfer→trend choice-set forcing and preventing F5-G1 (`hold_frac_flat == 0` from flat) from being achieved at the agent layer. Both kernels patched with the identical interop pattern: const bool f5_engaged = (isv[RL_F5_STATE_MASK_ENABLED_INDEX] > 0.5f); if (f5_engaged) skip the Hold override; The existing `if (position_lots != 0) return;` early-return in both kernels already restricts the override site to flat batches, so the check on the F5 master gate alone is sufficient — no new position read needed. When F5 is disabled (slot 823 == 0.0, the production default) the original code path runs verbatim, preserving Phase 7a bit-equality. Scope (per Phase 7b follow-up): * `rl_confidence_gate.cu` — Hold-override at `conf < threshold` gated on `!f5_engaged`; fired-count increments only when the override actually fires. * `rl_frd_gate.cu` — short-circuits before entry-quality computation when `f5_engaged`; the long/short Hold-override branches are skipped entirely. NOT touched (per Option A scope constraint): * `rl_session_risk_check.cu` — hard DD-limit safety. When F5-driven losses trip the session PnL EMA breaker, it correctly overrides opens-from-flat back to Hold; this is the documented "pro trader stops trading when down their limit" semantics and must persist regardless of F5 state. * `rl_min_hold_check.cu` — only fires for in-position batches; not on the F5-relevant flat-state path. Verification (RTX 3050, `test_data/futures-baseline-mid/`): * `cargo build --release --example alpha_rl_train -p ml-alpha` — clean. * `./scripts/determinism-check.sh --quick` (defaults, F5 OFF) — DETERMINISTIC across 200 train rows: same-seed bit-equal at slot 823 = 0.0. * `FOXHUNT_BAND_ENABLED=1 ./scripts/determinism-check.sh --quick` (F5 OFF) — DETERMINISTIC across 200 train rows. * `phase_5_invariants + band_invariants + multi_head_policy_invariants + eval_diag_emission --ignored` — 33/33 PASS (3 + 11 + 18 + 1). * F5 ON 50-step smoke (slot 823 temporarily set to 1.0 then reverted before commit): hold_frac_flat = 0.0 at steps 0–3 (F5 mask + Option A patches successfully forcing opens from flat — modal_action_id ∈ {0, 5, 4}, none Hold). At step 4 onward hold_frac_flat → 1.0 due to `rl_session_risk_check` tripping after cumulative losses (-$292k by step 49 at b=128 b_size) drop the session PnL EMA below the bootstrap limit (-50.0 at slot 540). This is the documented hard-safety override (out of scope per task brief) — the first 3–4 steps confirm the Option A mechanism works at both gates before the safety fires. * No NaN; training stable. Bootstrap of slot 823 reverted to 0.0 (OFF) before committing. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/rl_confidence_gate.cu | 18 +++++++++++++++++- crates/ml-alpha/cuda/rl_frd_gate.cu | 18 ++++++++++++++++++ 2 files changed, 35 insertions(+), 1 deletion(-) diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu index 6a2562e03..b8862c596 100644 --- a/crates/ml-alpha/cuda/rl_confidence_gate.cu +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -34,6 +34,15 @@ #define RL_HOLD_TARGET_FRAC_INDEX 575 #define RL_HOLD_FRAC_EMA_INDEX 576 #define RL_CONF_GATE_MAX_HOLD_FRAC_INDEX 584 +// Phase 7b F5 (2026-06-05) Option A — interop with state-conditional +// action availability mask. When F5 master gate (slot 823) is engaged +// AND position is flat, F5 has DELIBERATELY forced the agent into an +// opening action by masking Hold (and other invalid actions) in +// pi_logits. Overriding that back to Hold here neutralizes F5's +// surfer→trend choice-set forcing and prevents F5-G1 from achieving +// `hold_frac_flat == 0`. The suppression is gated on F5-engaged AND +// flat — in-position safety overrides are preserved verbatim. +#define RL_F5_STATE_MASK_ENABLED_INDEX 823 #define THRESHOLD_MIN 0.05f #define THRESHOLD_MAX 0.95f #define HOLD_EMA_ALPHA 0.1f @@ -118,7 +127,14 @@ extern "C" __global__ void rl_confidence_gate( const float lcb = (mu - v_min) - lambda * sigma; const float conf = fmaxf(0.0f, fminf(lcb / span, 1.0f)); - if (conf < threshold) { + // Phase 7b F5 (2026-06-05) Option A — suppress Hold override when + // F5 is engaged and we are flat. F5 has masked Hold (and other + // invalid actions) in pi_logits, so the sampled action is a + // deliberate F5-allowed opening. The early-return at line above + // (`position_lots != 0 → return`) already restricts us to flat + // here, so checking only the F5 master gate is sufficient. + const bool f5_engaged = (isv[RL_F5_STATE_MASK_ENABLED_INDEX] > 0.5f); + if (conf < threshold && !f5_engaged) { actions[b] = ACTION_HOLD; isv[RL_CONF_GATE_FIRED_COUNT_INDEX] += 1.0f; } diff --git a/crates/ml-alpha/cuda/rl_frd_gate.cu b/crates/ml-alpha/cuda/rl_frd_gate.cu index bb18c8b3e..a7d378cdc 100644 --- a/crates/ml-alpha/cuda/rl_frd_gate.cu +++ b/crates/ml-alpha/cuda/rl_frd_gate.cu @@ -35,6 +35,15 @@ #define RL_FRD_GATE_FIRED_COUNT_INDEX 518 #define RL_GATE_WARMUP_STEPS_INDEX 524 #define RL_STEP_COUNTER_ISV_INDEX 548 +// Phase 7b F5 (2026-06-05) Option A — interop with state-conditional +// action availability mask. When F5 master gate (slot 823) is engaged +// AND position is flat, F5 has DELIBERATELY forced the agent into an +// opening action by masking Hold (and other invalid actions) in +// pi_logits. Overriding that back to Hold here neutralizes F5's +// surfer→trend choice-set forcing and prevents F5-G1 from achieving +// `hold_frac_flat == 0`. The suppression is gated on F5-engaged AND +// flat — in-position safety overrides are preserved verbatim. +#define RL_F5_STATE_MASK_ENABLED_INDEX 823 extern "C" __global__ void rl_frd_gate( int* __restrict__ actions, // [B] IN/OUT @@ -60,6 +69,15 @@ extern "C" __global__ void rl_frd_gate( const int position_lots = *(const int*)(pos_state + b * pos_bytes); if (position_lots != 0) return; + // Phase 7b F5 (2026-06-05) Option A — short-circuit before quality + // computation when F5 is engaged and we are flat. F5 has masked + // Hold in pi_logits so the sampled opening is deliberate; overriding + // here would neutralize F5-G1 (`hold_frac_flat == 0` from flat). + // The early `position_lots != 0 → return` above already restricts + // us to flat, so checking the F5 master gate alone is sufficient. + const bool f5_engaged = (isv[RL_F5_STATE_MASK_ENABLED_INDEX] > 0.5f); + if (f5_engaged) return; + // Horizon 2 logits for this batch. const float* h2 = frd_logits + b * (FRD_N_HORIZONS * FRD_N_ATOMS) + FRD_H2_OFFSET;