diff --git a/crates/ml-alpha/cuda/compute_advantage_return.cu b/crates/ml-alpha/cuda/compute_advantage_return.cu index 24d717cc8..f2d643c4a 100644 --- a/crates/ml-alpha/cuda/compute_advantage_return.cu +++ b/crates/ml-alpha/cuda/compute_advantage_return.cu @@ -22,17 +22,20 @@ // Element-wise, trivially parallel. One thread per batch entry; no // reduction, no atomics. Per `feedback_no_atomicadd` not needed. -#define RL_GAMMA_INDEX 400 +#define RL_GAMMA_INDEX 400 +#define RL_HOLD_PRIOR_INDEX 574 +#define HOLD_ACTION 2 extern "C" __global__ void compute_advantage_return( - const float* __restrict__ isv, // ISV bus (≥ RL_GAMMA_INDEX + 1) + const float* __restrict__ isv, // ISV bus const float* __restrict__ rewards, // [b_size] const float* __restrict__ dones, // [b_size] 0.0 / 1.0 const float* __restrict__ v_t, // [b_size] V(s_t) const float* __restrict__ v_tp1, // [b_size] V(s_{t+1}) float* __restrict__ returns, // [b_size] OUT float* __restrict__ advantages, // [b_size] OUT - int b_size + int b_size, + const int* __restrict__ actions // [b_size] taken action indices ) { const int b = blockIdx.x * blockDim.x + threadIdx.x; if (b >= b_size) return; @@ -45,14 +48,22 @@ extern "C" __global__ void compute_advantage_return( const float ret = r + gamma * (1.0f - done) * vtp1; returns[b] = ret; - // Mask advantage to zero for non-done steps. With sparse rewards - // (5-6% of batch has trade closings), non-done advantages are - // dominated by V-estimation noise (γV(s')-V(s) ≈ 0 when V≈0). - // Only done-steps carry real reward signal. Masking prevents the - // 94% noise from diluting the gradient. The /B normalization in - // ppo_clipped_surrogate_bwd ensures the effective gradient from - // the ~6% real signals is properly scaled. - advantages[b] = done * (ret - vt); + + if (done > 0.5f) { + // Done-step: real reward signal from trade closing. + advantages[b] = ret - vt; + } else { + // Non-done step: surfer Hold prior. + // Small positive advantage for Hold, small negative for trading. + // This gives π continuous gradient that reinforces waiting as + // the default. Without this, done-gating makes Hold invisible + // to π (zero gradient → abandoned). + // + // ISV-driven magnitude so it adapts with reward_scale. + const float hold_prior = isv[RL_HOLD_PRIOR_INDEX]; + const int a = actions[b]; + advantages[b] = (a == HOLD_ACTION) ? hold_prior : -hold_prior; + } } // Advantage normalization: zero mean, unit variance across the batch. diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 470d86704..cd172e7cf 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1063,5 +1063,11 @@ pub const RL_OUTCOME_AUX_LAMBDA_INDEX: usize = 572; /// stable advantages = better q_pi_agree. Bootstrap: 0.005. pub const RL_TARGET_TAU_MAX_INDEX: usize = 573; +/// Per-step Hold prior magnitude. On non-done steps, Hold gets +/// +hold_prior advantage, non-Hold gets -hold_prior. Surfer +/// philosophy: waiting is the low-energy default. ISV-driven so it +/// can adapt with reward_scale. Bootstrap: 0.01. +pub const RL_HOLD_PRIOR_INDEX: usize = 574; + /// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 574; +pub const RL_SLOTS_END: usize = 575; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index ae6d82ae8..fa33a3e84 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -2639,7 +2639,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 103] = [ + let isv_constants: [(usize, f32); 104] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -2756,6 +2756,7 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_KURT_NOISE_FLOOR_INDEX, 1.0), (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_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), @@ -3645,6 +3646,7 @@ impl IntegratedTrainer { args.push_ptr(returns_d.raw_ptr()); args.push_ptr(advantages_d.raw_ptr()); args.push_i32(b_size_i); + args.push_ptr(self.actions_d.raw_ptr()); let mut ptrs = args.build_arg_ptrs(); unsafe { raw_launch( @@ -6161,6 +6163,7 @@ impl IntegratedTrainer { args.push_ptr(self.returns_d.raw_ptr()); args.push_ptr(self.advantages_d.raw_ptr()); args.push_i32(b_size_i); + args.push_ptr(self.actions_d.raw_ptr()); let mut ptrs = args.build_arg_ptrs(); unsafe { raw_launch(