diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 5f04110e7..b288ef90e 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -88,6 +88,7 @@ const KERNELS: &[&str] = &[ "rl_encoder_context_broadcast", // P2: broadcast per-batch context (16 dims) into encoder input [B,K,56] at cols 40-55 "rl_gate_threshold_controller", // adaptive gate thresholds from trade frequency (dones EMA) "rl_reward_shaping", // surfer-philosophy: entry cost + short-hold penalty + hold bonus + "rl_min_hold_check", // hard minimum hold time — override close actions to Hold before N steps ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_min_hold_check.cu b/crates/ml-alpha/cuda/rl_min_hold_check.cu new file mode 100644 index 000000000..bf9214d78 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_min_hold_check.cu @@ -0,0 +1,45 @@ +// rl_min_hold_check.cu — ISV-driven minimum hold time. +// +// Overrides closing actions (a3/a4/a9/a10) to Hold (a2) when the +// current position has been held for fewer than min_hold steps. +// Hard constraint — the agent cannot exit before the minimum. +// +// Runs AFTER action selection, BEFORE trail_stop_check (which can +// still force-exit on genuine stop-loss breaches regardless of +// hold time — safety overrides patience). +// +// Per `feedback_no_atomicadd`: per-batch element-wise. +// Per `feedback_cpu_is_read_only`: pure device-side. + +#include + +#define ACTION_HOLD 2 +#define ACTION_FLAT_FROM_LONG 3 +#define ACTION_FLAT_FROM_SHORT 4 +#define ACTION_HALF_FLAT_LONG 9 +#define ACTION_HALF_FLAT_SHORT 10 +#define RL_MIN_HOLD_STEPS_INDEX 536 + +extern "C" __global__ void rl_min_hold_check( + int* __restrict__ actions, // [B] IN/OUT + const int* __restrict__ steps_since_done, // [B] + const float* __restrict__ isv, + int b_size +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const int min_hold = (int)isv[RL_MIN_HOLD_STEPS_INDEX]; + if (min_hold <= 0) return; + + const int hold = steps_since_done[b]; + if (hold >= min_hold) return; + + const int action = actions[b]; + if (action == ACTION_FLAT_FROM_LONG || + action == ACTION_FLAT_FROM_SHORT || + action == ACTION_HALF_FLAT_LONG || + action == ACTION_HALF_FLAT_SHORT) { + actions[b] = ACTION_HOLD; + } +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index e5a862a04..c353ed8bc 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -926,6 +926,11 @@ pub const RL_SHORT_HOLD_PENALTY_INDEX: usize = 534; /// Scales with sqrt(hold_time). Default: 0.50. pub const RL_HOLD_BONUS_INDEX: usize = 535; +/// Hard minimum hold time (steps). Closing actions are overridden to +/// Hold when steps_since_done < this value. Trail stops still fire +/// regardless (safety overrides patience). Default: 20. +pub const RL_MIN_HOLD_STEPS_INDEX: usize = 536; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 536; +pub const RL_SLOTS_END: usize = 537; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 341a53276..5bea9071c 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -204,6 +204,8 @@ const RL_GATE_THRESHOLD_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_gate_threshold_controller.cubin")); const RL_REWARD_SHAPING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_shaping.cubin")); +const RL_MIN_HOLD_CHECK_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_min_hold_check.cubin")); const RL_TRADE_CONTEXT_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_trade_context_update.cubin")); const RL_MULTIRES_FEATURES_UPDATE_CUBIN: &[u8] = @@ -467,6 +469,8 @@ pub struct IntegratedTrainer { rl_gate_threshold_controller_fn: CudaFunction, _rl_reward_shaping_module: Arc, rl_reward_shaping_fn: CudaFunction, + _rl_min_hold_check_module: Arc, + rl_min_hold_check_fn: CudaFunction, _rl_trade_context_update_module: Arc, rl_trade_context_update_fn: CudaFunction, _rl_multires_features_update_module: Arc, @@ -990,6 +994,12 @@ impl IntegratedTrainer { let rl_reward_shaping_fn = rl_reward_shaping_module .load_function("rl_reward_shaping") .context("load rl_reward_shaping")?; + let rl_min_hold_check_module = ctx + .load_cubin(RL_MIN_HOLD_CHECK_CUBIN.to_vec()) + .context("load rl_min_hold_check cubin")?; + let rl_min_hold_check_fn = rl_min_hold_check_module + .load_function("rl_min_hold_check") + .context("load rl_min_hold_check")?; let rl_trade_context_update_module = ctx .load_cubin(RL_TRADE_CONTEXT_UPDATE_CUBIN.to_vec()) .context("load rl_trade_context_update cubin")?; @@ -1419,6 +1429,8 @@ impl IntegratedTrainer { rl_gate_threshold_controller_fn, _rl_reward_shaping_module: rl_reward_shaping_module, rl_reward_shaping_fn, + _rl_min_hold_check_module: rl_min_hold_check_module, + rl_min_hold_check_fn, _rl_trade_context_update_module: rl_trade_context_update_module, rl_trade_context_update_fn, _rl_multires_features_update_module: rl_multires_features_update_module, @@ -1564,7 +1576,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 72] = [ + let isv_constants: [(usize, f32); 73] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1643,6 +1655,7 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_SHORT_HOLD_MIN_STEPS_INDEX, 20.0), (crate::rl::isv_slots::RL_SHORT_HOLD_PENALTY_INDEX, 0.5), (crate::rl::isv_slots::RL_HOLD_BONUS_INDEX, 0.50), + (crate::rl::isv_slots::RL_MIN_HOLD_STEPS_INDEX, 20.0), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_1_INDEX, 1.0), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_2_INDEX, 10.0), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_3_INDEX, 600.0), @@ -3875,7 +3888,29 @@ impl IntegratedTrainer { let pos_bytes_i = lobsim.pos_bytes() as i32; let b_size_i = b_size as i32; - // ── SP20 P1+P5 trail mutate (a7/a8) — BEFORE actions_to_market_targets. + // Min hold check — override closing actions to Hold when + // hold time < ISV-driven minimum. Trail stops still fire + // (they run AFTER this, safety overrides patience). + { + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&self.rl_min_hold_check_fn); + launch + .arg(&mut self.actions_d) + .arg(&self.steps_since_done_d) + .arg(&self.isv_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_min_hold_check launch")?; + } + } + + // Trail mutate (a7/a8) — BEFORE actions_to_market_targets. // Mutates unit_trail_distance for active units if the chosen action // is TrailTighten/Loosen. Non-trail actions pass through. {