diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 92f9e16e7..5dd03fa0e 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -70,6 +70,7 @@ const KERNELS: &[&str] = &[ "rl_pi_action_kernel", // audit Option B: π drives action selection via multinomial sampling from softmax(pi_logits); Q becomes pure critic "rl_reward_clamp_controller", // audit 2026-05-24: adaptive [-LOSS, +WIN] clamp from positive-tail EMA; replaces static [-3, +1] that crushed winning-trade signal in rmgm5 "rl_atom_support_update", // audit 2026-05-24 followup: refreshes atom_supports_d from ISV V_MIN/V_MAX so C51 atom span adapts with reward clamp (Q learning was capped at V_MAX=1.0) + "rl_kl_reference_grad", "rl_q_pi_distill_grad", // audit 2026-05-24 vj5f6 followup: KL(softmax(E_Q/τ) || π_new) gradient ADDED to pi_grad_logits — couples Q's improved C51 calibration to action selection (was decoupled per Option B) "rl_q_distill_lambda_controller",// audit 2026-05-24 rljzl followup: adaptive λ_distill via Schulman bounded step on KL_EMA vs target "rl_unit_state_update", // SP20 P1+P5 audit fix: per-unit trade state machine — detects open/close/reverse transitions, sets up unit slot 0 entry+trail diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu index a7ab9598f..1bf5214f4 100644 --- a/crates/ml-alpha/cuda/rl_confidence_gate.cu +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -31,6 +31,12 @@ #define RL_CONF_GATE_FIRED_COUNT_INDEX 515 #define RL_GATE_WARMUP_STEPS_INDEX 524 #define RL_STEP_COUNTER_ISV_INDEX 548 +#define RL_HOLD_TARGET_FRAC_INDEX 575 +#define RL_HOLD_FRAC_EMA_INDEX 576 +#define THRESHOLD_MIN 0.05f +#define THRESHOLD_MAX 0.95f +#define HOLD_EMA_ALPHA 0.1f +#define THRESHOLD_ADJUST_RATE 1.1f extern "C" __global__ void rl_confidence_gate( int* __restrict__ actions, // [B] IN/OUT @@ -98,7 +104,39 @@ extern "C" __global__ void rl_confidence_gate( if (conf < threshold) { actions[b] = ACTION_HOLD; - // Diag: count fires. Single-thread kernel so direct write is safe. isv[RL_CONF_GATE_FIRED_COUNT_INDEX] += 1.0f; } + + // ── Adaptive threshold controller (block 0 only). ─────────────── + // Count Hold actions in the batch (post-gate), compute hold_frac, + // EMA it, and adjust threshold to maintain the target Hold fraction. + // Runs once per step via the single-block launch. + if (b == 0) { + // Count post-gate Hold actions across the batch. + int hold_count = 0; + for (int i = 0; i < b_size; i++) { + if (actions[i] == ACTION_HOLD) hold_count++; + } + const float hold_frac = (float)hold_count / (float)b_size; + + // EMA of Hold fraction. + const float prev_ema = isv[RL_HOLD_FRAC_EMA_INDEX]; + const float new_ema = (prev_ema == 0.0f) + ? hold_frac + : (1.0f - HOLD_EMA_ALPHA) * prev_ema + HOLD_EMA_ALPHA * hold_frac; + isv[RL_HOLD_FRAC_EMA_INDEX] = new_ema; + + // Schulman-style bounded adjustment: if Hold fraction is below + // target, raise threshold (gate more aggressively). If above + // target, lower threshold (allow more trading). + const float hold_target = isv[RL_HOLD_TARGET_FRAC_INDEX]; + float new_threshold = threshold; + if (new_ema < hold_target * 0.9f) { + new_threshold = threshold * THRESHOLD_ADJUST_RATE; + } else if (new_ema > hold_target * 1.1f) { + new_threshold = threshold / (1.0f + (THRESHOLD_ADJUST_RATE - 1.0f) * 0.01f); + } + new_threshold = fmaxf(THRESHOLD_MIN, fminf(new_threshold, THRESHOLD_MAX)); + isv[RL_CONF_GATE_THRESHOLD_INDEX] = new_threshold; + } } diff --git a/crates/ml-alpha/cuda/rl_kl_reference_grad.cu b/crates/ml-alpha/cuda/rl_kl_reference_grad.cu new file mode 100644 index 000000000..a69b2c602 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_kl_reference_grad.cu @@ -0,0 +1,77 @@ +// rl_kl_reference_grad.cu — KL divergence gradient from a hold-heavy +// reference policy. Proven approach for preventing policy collapse with +// sparse rewards (RLHF/InstructGPT pattern). +// +// Adds β × (π_θ(a|s) - π_ref(a)) to pi_grad_logits for every batch +// element, every step. Unlike entropy bonus (pushes toward uniform) or +// done-gated advantages (only fires on 6% of batch), this provides +// continuous gradient signal that specifically preserves Hold as the +// default action. +// +// π_ref encodes the surfer philosophy: +// Hold = 50%, each trading action = 5% (uniform over 10 non-Hold) +// +// The gradient ∂KL(π_θ||π_ref)/∂logit_a = π_θ(a) - π_ref(a): +// - If π(Hold) < 50%: gradient pushes Hold probability UP +// - If π(Hold) > 50%: gradient gently pushes Hold DOWN (allows trading) +// - If π(BuyL1) > 5%: gradient pushes it DOWN toward the prior +// +// β is ISV-driven so it can adapt with training progress. +// Grid=(B,1,1), Block=(N_ACTIONS,1,1). One thread per action. + +#define N_ACTIONS 11 +#define ACTION_HOLD 2 +#define RL_KL_REF_BETA_INDEX 578 + +extern "C" __global__ void rl_kl_reference_grad( + const float* __restrict__ pi_logits, // [B × N_ACTIONS] + float* __restrict__ isv, // ISV bus (reads β) + float* __restrict__ pi_grad_logits, // [B × N_ACTIONS] — ADD to + int B +) { + const int b = blockIdx.x; + const int a = threadIdx.x; + if (b >= B || a >= N_ACTIONS) return; + + const float beta = isv[RL_KL_REF_BETA_INDEX]; + if (beta <= 0.0f) return; + + // Reference policy: Hold=50%, rest=5% each. + const float pi_ref = (a == ACTION_HOLD) ? 0.5f : 0.05f; + + // Compute π_θ(a|s) via softmax. + const int base = b * N_ACTIONS; + + __shared__ float s_logits[11]; + __shared__ float s_max; + __shared__ float s_sumexp; + + s_logits[a] = pi_logits[base + a]; + __syncthreads(); + + if (a == 0) { + float m = s_logits[0]; + for (int i = 1; i < N_ACTIONS; ++i) m = fmaxf(m, s_logits[i]); + s_max = m; + } + __syncthreads(); + + float exp_val = expf(s_logits[a] - s_max); + __shared__ float s_exp[11]; + s_exp[a] = exp_val; + __syncthreads(); + + if (a == 0) { + float sum = 0.0f; + for (int i = 0; i < N_ACTIONS; ++i) sum += s_exp[i]; + s_sumexp = sum; + } + __syncthreads(); + + const float pi_theta = s_exp[a] / s_sumexp; + + // ∂KL/∂logit_a = π_θ(a) - π_ref(a) + // Additive to existing PPO + distillation gradients. + // /B for batch-size invariance. + pi_grad_logits[base + a] += beta * (pi_theta - pi_ref) / (float)B; +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index cd172e7cf..0f6b69e32 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1069,5 +1069,18 @@ pub const RL_TARGET_TAU_MAX_INDEX: usize = 573; /// can adapt with reward_scale. Bootstrap: 0.01. pub const RL_HOLD_PRIOR_INDEX: usize = 574; +/// Target fraction of actions that should be Hold (surfer default). +/// The gate threshold adapts to maintain this fraction. +/// Bootstrap: 0.5 (50% Hold is the surfer minimum). +pub const RL_HOLD_TARGET_FRAC_INDEX: usize = 575; + +/// EMA of actual Hold fraction (updated by gate kernel each step). +pub const RL_HOLD_FRAC_EMA_INDEX: usize = 576; + +/// KL reference policy penalty weight. Controls how strongly π is +/// pulled toward the hold-heavy reference [Hold=50%, rest=5%]. +/// Higher β = stronger Hold preservation. Bootstrap: 0.5. +pub const RL_KL_REF_BETA_INDEX: usize = 578; + /// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 575; +pub const RL_SLOTS_END: usize = 579; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index b541062f4..65c8c5a63 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -599,6 +599,8 @@ pub struct IntegratedTrainer { // Q→π distillation gradient (audit 2026-05-24 vj5f6 followup). _rl_q_pi_distill_grad_module: Arc, rl_q_pi_distill_grad_fn: CudaFunction, + _rl_kl_reference_module: Arc, + rl_kl_reference_grad_fn: CudaFunction, // C51+IQN ensemble action-value (audit 2026-05-25). _rl_ensemble_action_value_module: Arc, rl_ensemble_action_value_fn: CudaFunction, @@ -1361,6 +1363,12 @@ impl IntegratedTrainer { let rl_q_pi_distill_grad_fn = rl_q_pi_distill_grad_module .load_function("rl_q_pi_distill_grad") .context("load rl_q_pi_distill_grad")?; + let rl_kl_reference_module = ctx + .load_cubin(include_bytes!(concat!(env!("OUT_DIR"), "/rl_kl_reference_grad.cubin")).to_vec()) + .context("load rl_kl_reference_grad cubin")?; + let rl_kl_reference_grad_fn = rl_kl_reference_module + .load_function("rl_kl_reference_grad") + .context("load rl_kl_reference_grad")?; let rl_ensemble_action_value_module = ctx .load_cubin(RL_ENSEMBLE_ACTION_VALUE_CUBIN.to_vec()) .context("load rl_ensemble_action_value cubin")?; @@ -2310,6 +2318,8 @@ impl IntegratedTrainer { rl_atom_support_update_fn, _rl_q_pi_distill_grad_module: rl_q_pi_distill_grad_module, rl_q_pi_distill_grad_fn, + _rl_kl_reference_module: rl_kl_reference_module, + rl_kl_reference_grad_fn, _rl_ensemble_action_value_module: rl_ensemble_action_value_module, rl_ensemble_action_value_fn, _rl_q_distill_lambda_controller_module: rl_q_distill_lambda_controller_module, @@ -2639,7 +2649,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 104] = [ + let isv_constants: [(usize, f32); 106] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -2757,6 +2767,8 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_TAU_BOOTSTRAP_INDEX, 0.005), (crate::rl::isv_slots::RL_TARGET_TAU_MAX_INDEX, 0.005), (crate::rl::isv_slots::RL_HOLD_PRIOR_INDEX, 0.1), + (crate::rl::isv_slots::RL_HOLD_TARGET_FRAC_INDEX, 0.5), + (crate::rl::isv_slots::RL_KL_REF_BETA_INDEX, 0.5), (crate::rl::isv_slots::RL_EPS_BOOTSTRAP_INDEX, 0.2), (crate::rl::isv_slots::RL_ROLLOUT_BOOTSTRAP_INDEX, 2048.0), (crate::rl::isv_slots::RL_REWARD_SCALE_BOOTSTRAP_INDEX, 1.0), @@ -4064,10 +4076,27 @@ impl IntegratedTrainer { } } - // NOTE: q_distill_lambda controller is now fused into - // rl_fused_controllers (fired in the per-step block). Reads - // KL_EMA from the previous step's backward pass — one-step - // delay is negligible for a 0.998/1.2 step-size controller. + // KL penalty to hold-heavy reference policy. Fires on ALL batch + // elements (not done-gated). Provides continuous gradient that + // keeps π near [Hold=50%, rest=5%]. Proven RLHF pattern for + // preventing policy collapse with sparse rewards. + { + let b_size_i = b_size as i32; + let mut args = RawArgs::new(); + args.push_ptr(self.pi_logits_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); + args.push_ptr(self.ss_pi_grad_logits_d.raw_ptr()); + args.push_i32(b_size_i); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_kl_reference_grad_fn.cu_function(), + (b_size as u32, 1, 1), (N_ACTIONS as u32, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_kl_reference_grad: {:?}", e))?; + } + } self.policy_head .backward_to_w_b_h(