feat(ml-alpha): Phase A — surfer-scaffold force-pin (slot 824 + kernel gate + env flag + diag)

Adds ISV slot 824 (RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX) and a matching
kernel early-return in rl_surfer_scaffold_controller.cu so that when
FOXHUNT_PIN_SURFER_SCAFFOLD=1 is set, the controller leaves slot 753 at
its 0.0 pure-pnl bootstrap instead of overwriting it every step.
Without the pin, the 0.0 bootstrap was cosmetic — the unconditional
write re-enabled all four Phase-5 non-potential shaping terms each step,
giving Pearson(reward,pnl)=0.28 vs the 0.70 gate.  Pin=OFF leaves
behaviour fully unchanged (new branch only fires when slot 824 > 0.5).

- isv_slots.rs: slot 824 constant + RL_SLOTS_END 824→825 + test
- rl_surfer_scaffold_controller.cu: #define + early-return guard
- integrated.rs: isv_constants 275→276 + env-gated bootstrap entry
- eval_diag_emission.rs: rewards.surfer_scaffold_force_pin diag field +
  EXPECTED_LEAVES 746→747

Build: SQLX_OFFLINE=true cargo build -p ml-alpha --profile=dev-release OK
Tests: 69 passed / 0 failed (force_pin_slot_allocated_below_end + all existing)

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-06-05 09:18:39 +02:00
parent 7376b1c670
commit d8447475c9
4 changed files with 58 additions and 3 deletions

View File

@@ -55,6 +55,7 @@
#define RL_SURFER_BREAK_EVEN_WR_INDEX 754
#define RL_SURFER_K_SHARPNESS_INDEX 755
#define RL_SURFER_WARMUP_TRADES_INDEX 756
#define RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX 824
__device__ static inline float sigmoidf(float x) {
return 1.0f / (1.0f + expf(-x));
@@ -64,6 +65,13 @@ extern "C" __global__ void rl_surfer_scaffold_controller(float* isv) {
// Single-thread kernel — launched (1,1,1)/(1,1,1).
if (threadIdx.x != 0 || blockIdx.x != 0) return;
// Phase A force-pin (2026-06-05): when set, leave slot 753 at its
// bootstrap value (pure-pnl) instead of overwriting it. Fixes the
// cosmetic-bootstrap bug (pearl_reward_misalign_blocks_f4_slot753_override).
if (isv[RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX] > 0.5f) {
return; // do NOT write isv[RL_SURFER_SCAFFOLD_WEIGHT_INDEX]
}
const float wr_ema = isv[RL_WIN_RATE_EMA_INDEX];
const float n_trades = isv[RL_CUMULATIVE_DONES_INDEX];
const float frac_decay = isv[RL_EDGE_PH_FRAC_ALERTED_INDEX];

View File

@@ -2345,6 +2345,22 @@ pub const RL_F2_DIAG_TARGET_ARGMAX_INDEX: usize = 822;
/// to 1.0 to engage Phase 7b. Clamp `{0.0, 1.0}` (binary).
pub const RL_F5_STATE_MASK_ENABLED_INDEX: usize = 823;
/// Phase A reward-pnl alignment (2026-06-05) — force-pin gate for the
/// surfer-scaffold weight controller. When `> 0.5`,
/// `rl_surfer_scaffold_controller.cu` returns WITHOUT writing slot 753,
/// so slot 753 holds its bootstrap value (pure-pnl = 0.0) for the whole
/// run instead of being overwritten to w≈0.97 every step.
///
/// Root cause it fixes: the slot-753 0.0 bootstrap was cosmetic — the
/// controller's unconditional write re-enabled all four Phase-5
/// non-potential shaping terms, giving Pearson(reward,pnl)=0.28 (gate
/// 0.70). See `pearl_reward_misalign_blocks_f4_slot753_override`.
///
/// Bootstrap 0.0 (OFF — controller writes as before). Set to 1.0 via the
/// `FOXHUNT_PIN_SURFER_SCAFFOLD=1` env flag to pin pure-pnl. Clamp
/// `{0.0, 1.0}` (binary).
pub const RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX: usize = 824;
/// 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.
@@ -2373,4 +2389,19 @@ pub const RL_F5_STATE_MASK_ENABLED_INDEX: usize = 823;
/// Post-Phase 7a F2 (Q-centered distill target, slots 817-818): 818.
/// Post-Phase 7a F2-diag (kernel-side counters, slots 819-822): 822.
/// Post-Phase 7b F5 (state-conditional action mask gate, slot 823): 823.
pub const RL_SLOTS_END: usize = 824;
/// Post-Phase A reward-pnl alignment (surfer-scaffold force-pin, slot 824): 824.
pub const RL_SLOTS_END: usize = 825;
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn force_pin_slot_allocated_below_end() {
// Phase A (2026-06-05) — reward-pnl alignment force-pin.
assert_eq!(RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX, 824);
assert_eq!(RL_SLOTS_END, 825);
assert!(RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX < RL_SLOTS_END);
assert_ne!(RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX, RL_F5_STATE_MASK_ENABLED_INDEX);
}
}

View File

@@ -3862,7 +3862,7 @@ impl IntegratedTrainer {
// (slot, value) pair — pure device write, no HtoD per
// `feedback_no_htod_htoh_only_mapped_pinned`.
{
let isv_constants: [(usize, f32); 275] = [
let isv_constants: [(usize, f32); 276] = [
// Static seeds for the adaptive reward-clamp controller —
// these are the initial values that
// `rl_reward_clamp_controller` will replace once it
@@ -4414,6 +4414,19 @@ impl IntegratedTrainer {
// (confidence-gate interaction) covers this; orthogonal
// to F5 kernel correctness.
(crate::rl::isv_slots::RL_F5_STATE_MASK_ENABLED_INDEX, 1.0),
// Phase A reward-pnl alignment (2026-06-05) — surfer-scaffold
// force-pin. Env-gated per the FOXHUNT_USE_ROLLOUT precedent
// (commit 779c03b9d): unset/0 → 0.0 (controller writes slot
// 753 as before, behaviour unchanged); =1 → 1.0 (controller
// skips its write, slot 753 holds its 0.0 pure-pnl bootstrap).
(crate::rl::isv_slots::RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX, {
let pin: bool = std::env::var("FOXHUNT_PIN_SURFER_SCAFFOLD")
.ok()
.and_then(|v| v.parse::<u8>().ok())
.map(|n| n > 0)
.unwrap_or(false);
if pin { 1.0_f32 } else { 0.0_f32 }
}),
];
for (slot, value) in isv_constants.iter() {
let slot_i32 = *slot as i32;
@@ -12574,6 +12587,8 @@ impl IntegratedTrainer {
isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX],
"surfer_scaffold_weight":
isv[crate::rl::isv_slots::RL_SURFER_SCAFFOLD_WEIGHT_INDEX],
"surfer_scaffold_force_pin":
isv[crate::rl::isv_slots::RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX],
// Phase 3D-B (2026-06-03): quadratic-cost slot snapshot.
// No per-step batch-aggregated cost is emitted here — the
// reward kernel applies the penalty per-batch element

View File

@@ -284,7 +284,8 @@ fn eval_diag_jsonl_emitted_with_train_schema_parity() -> Result<()> {
// leaves from upstream commits not captured in the prior EXPECTED_LEAVES
// ledger between B-11-β bump 671→672 and Phase 5 0463e44e0; those 8 are
// not Phase 7a additions but had not been resynced before this commit).
const EXPECTED_LEAVES: usize = 746;
// +1 Phase A (2026-06-05): rewards.surfer_scaffold_force_pin (slot 824).
const EXPECTED_LEAVES: usize = 747;
anyhow::ensure!(
train_paths.len() == EXPECTED_LEAVES,
"train leaf count {} != expected {EXPECTED_LEAVES} (schema drifted; \