gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat(rl): V_pred structurally clamped to C51 atom span (G.1)

Browse Source 
v_head_fwd now reads V_MIN/V_MAX from ISV slots 485/484 (same slots
the C51 atom support adapter writes) and clamps the linear output to
that range at the kernel boundary. Bounds advantage magnitude
(|returns − V_pred|) by 2 × V_MAX regardless of stale-V state.

Defensive fix per pearl_clamp_v_target_at_atom_span +
pearl_c51_atom_span_must_track_clamp_range — protects against the
canonical reward_scale↔V-head response-time pathology where V's stale
predictions amplify into PPO surrogate + V regression spikes when the
controller adapts reward_scale aggressively. In the F.5 200-step local
smoke this clamp didn't bite (V_pred stayed within bounds at the
short run length), but the structural protection matters for longer
production runs where V can drift before the controllers catch up.

Hard-saturated clamp (no straight-through estimator) — the gradient
at the boundary is zero in the "push further out" direction, normal
toward the interior. V can always learn back into bounds when its raw
output drifts out (target is inside bounds → grad pulls V back in),
but cannot push the prediction outside support.

API surface change: `ValueHead::forward(h_t, b_size, v_pred)` →
`ValueHead::forward(h_t, isv, b_size, v_pred)`. The 3 call sites in
IntegratedTrainer (step_synthetic + step_with_lobsim h_t/h_tp1) now
pass `&self.isv_d`.

Verification (RTX 3050 Ti):
  * cargo check -p ml-alpha → clean
  * integrated_trainer_smoke 1/1 → ok
  * frd_head 10/10 + trade_management_kernels 5/5 → no regression
  * audit-rust-consts → 0 flags

Independent finding from the smoke diag: the OBSERVED chronic spike
pattern (|l_pi|>30, l_v>100) traces to `rl_reward_clamp_controller`
widening WIN/LOSS bounds to 41.3 (vs seeds 1.0/3.0) when MARGIN hits
its MAX_MARGIN=5 ceiling. That's a separate failure mode addressed
in the next commit (structural cap on scaled reward magnitude).
This commit is contained in:
jgrusewski
2026-05-24 19:47:25 +02:00
parent 7df7c81d37
commit bc6e5bcde4
3 changed files with 44 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

31
crates/ml-alpha/cuda/v_head_fwd_bwd.cu
View File

@@ -5,6 +5,23 @@
// head is a single linear layer `V(s) = b_v + Σ_c w_v[c] × h_t[b, c]`
// supervised by MSE against the bootstrap return target.
//
// G.1 (2026-05-24) — V_pred is now structurally clamped to the ISV-driven
// C51 atom span [V_MIN, V_MAX] (slots 485/484) at the fwd kernel's output.
// Rationale: F.5 smoke diagnosed chronic |l_pi|>30 + l_v>100 spikes
// arising from the reward_scale↔V-head response-time mismatch — when
// reward_scale adapts down to 0.003 to track a fresh trade-magnitude
// regime, V's stale predictions of O(10) produced advantage = returns -
// V_pred ≈ -10, blowing up PPO surrogate + V regression alike. The
// structural clamp bounds V_pred to the same support the C51 distribution
// uses, so advantage cannot exceed 2 × V_MAX regardless of stale-V state.
//
// The clamp is hard-saturated (no straight-through estimator): the
// gradient at the boundary is zero in the "push further out" direction,
// non-zero toward the interior. This means V can always learn back into
// bounds when its raw output drifts out (target is inside bounds → grad
// pulls V back in), but cannot push the prediction outside support. Per
// `pearl_clamp_v_target_at_atom_span` + `pearl_c51_atom_span_must_track_clamp_range`.
//
// Forward (`v_head_fwd`):
// v_pred[b] = b_v + Σ_c w_v[c] × h_t[b, c]
//
@@ -46,6 +63,12 @@
#define HIDDEN_DIM 128
// G.1 ISV slot indices for V_pred clamp bounds — same slots used by the
// C51 atom-support kernel (rl_atom_support_update) so the V head and the
// distributional Q head share a single canonical support range.
#define RL_C51_V_MAX_INDEX 484
#define RL_C51_V_MIN_INDEX 485
// ─────────────────────────────────────────────────────────────────────
// v_head_fwd: per-batch scalar V projection.
@@ -67,6 +90,7 @@ extern "C" __global__ void v_head_fwd(
const float* __restrict__ w_v, // [HIDDEN_DIM]
const float* __restrict__ b_v, // [1]
const float* __restrict__ h_t, // [B * HIDDEN_DIM]
const float* __restrict__ isv, // ISV bus (G.1 — reads V_MIN/V_MAX)
int B,
float* __restrict__ v_pred // [B]
) {
@@ -86,6 +110,13 @@ extern "C" __global__ void v_head_fwd(
for (int i = 0; i < HIDDEN_DIM; ++i) {
acc += s_w[i] * h_t[batch * HIDDEN_DIM + i];
}
// G.1 structural clamp to ISV-driven C51 atom span. Bounds
// advantage magnitude (= |returns - V_pred|) by 2 × V_MAX
// regardless of stale-V state, breaking the F.5-diagnosed
// reward_scale↔V-head response-time pathology.
const float v_min = isv[RL_C51_V_MIN_INDEX];
const float v_max = isv[RL_C51_V_MAX_INDEX];
acc = fmaxf(v_min, fminf(acc, v_max));
v_pred[batch] = acc;
}
}

11
crates/ml-alpha/src/rl/ppo.rs
View File

@@ -435,11 +435,19 @@ impl ValueHead {
self.cfg.hidden_dim
}
/// Phase E.2 forward — V(s) = b + Σ_c w[c] × h_t[b, c]. Writes
/// Phase E.2 forward — V(s) = b + Σ_c w[c] × h_t[b, c], clamped to
/// the ISV-driven C51 atom span `[V_MIN, V_MAX]` (G.1). Writes
/// `v_pred [B]`.
///
/// G.1 — `isv` is the trainer's device ISV bus; the kernel reads
/// slots `RL_C51_V_MIN_INDEX = 485` and `RL_C51_V_MAX_INDEX = 484`
/// to clamp the prediction. The clamp prevents stale V_pred (during
/// reward_scale controller transitions) from amplifying advantage
/// magnitude into the PPO surrogate + V regression loss.
pub fn forward(
&self,
h_t: &CudaSlice<f32>,
isv: &CudaSlice<f32>,
b_size: usize,
v_pred: &mut CudaSlice<f32>,
) -> Result<()> {
@@ -457,6 +465,7 @@ impl ValueHead {
.arg(&self.w_d)
.arg(&self.b_d)
.arg(h_t)
.arg(isv)
.arg(&b_i)
.arg(v_pred);
unsafe {

6
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -2438,7 +2438,7 @@ impl IntegratedTrainer {
.forward_logits(h_t_borrow, b_size, &mut pi_logits_d)
.context("policy_head.forward_logits")?;
self.value_head
.forward(h_t_borrow, b_size, &mut v_pred_d)
.forward(h_t_borrow, &self.isv_d, b_size, &mut v_pred_d)
.context("value_head.forward")?;
// R7d: Online Q at sampled_h_tp1 for the Double-DQN argmax that
@@ -3326,10 +3326,10 @@ impl IntegratedTrainer {
.forward(&self.h_tp1_d, b_size, &mut q_logits_tp1_d)
.context("step_with_lobsim: dqn_head.forward(h_tp1) for Double-DQN argmax")?;
self.value_head
.forward(h_t_borrow, b_size, &mut v_pred_d)
.forward(h_t_borrow, &self.isv_d, b_size, &mut v_pred_d)
.context("step_with_lobsim: value_head.forward(h_t)")?;
self.value_head
.forward(&self.h_tp1_d, b_size, &mut v_pred_tp1_d)
.forward(&self.h_tp1_d, &self.isv_d, b_size, &mut v_pred_tp1_d)
.context("step_with_lobsim: value_head.forward(h_tp1) for true V(s_{t+1})")?;
// ── Step 2b: Forward π logits for log_pi_old. ─────────────────