feat(ml-alpha): Phase 7b F5 Option A — gate-aware Hold suppression for confidence + FRD gates

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 <noreply@anthropic.com>

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;
     }

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;