From 88abf8185c0f8ebf2dc2bb07d3bbd0e41a63e530 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 25 May 2026 10:26:10 +0200 Subject: [PATCH] feat(rl): adaptive gate threshold controller from trade frequency MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit New rl_gate_threshold_controller.cu watches dones EMA and adjusts confidence + FRD gate thresholds via Schulman bounded step: - dones_ema < target×0.5 → relax thresholds (×0.95) - dones_ema > target×2.0 → tighten thresholds (÷0.95) ISV-driven: target (0.02), conf bounds (0.001-0.50), FRD bounds (0.05-0.50), adjust rate (0.95). Runs per step after warmup. Also lowers default thresholds: conf 0.10→0.01, FRD 0.35→0.15. Post-warmup, the controller adapts these based on actual trade flow instead of relying on static defaults. ISV slots 525-531. RL_SLOTS_END → 532. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 1 + .../cuda/rl_gate_threshold_controller.cu | 80 +++++++++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 18 ++++- crates/ml-alpha/src/trainer/integrated.rs | 49 +++++++++++- 4 files changed, 143 insertions(+), 5 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_gate_threshold_controller.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 3d6676bd5..e511f7d48 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -86,6 +86,7 @@ const KERNELS: &[&str] = &[ "rl_trade_context_update", // P1: per-batch trade-arc features (time_in_trade, unrealized_R, pos_mag, entry_dist) "rl_multires_features_update", // P0: per-batch multi-resolution streaming features (3 horizons × 4 features) "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) ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu new file mode 100644 index 000000000..8de10f656 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu @@ -0,0 +1,80 @@ +// rl_gate_threshold_controller.cu — adaptive gate threshold controller. +// +// Watches trade frequency (dones EMA) and adjusts confidence + FRD +// gate thresholds to maintain a target trade rate. When trades dry up +// (dones_ema < target), thresholds DECREASE (more permissive). When +// trades are abundant (dones_ema > target), thresholds INCREASE +// (more selective). Schulman bounded discrete step pattern. +// +// Runs once per step, single thread. Reads/writes ISV. +// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`. + +#include + +#define RL_CONF_GATE_THRESHOLD_INDEX 512 +#define RL_FRD_GATE_THR_LONG_INDEX 516 +#define RL_FRD_GATE_THR_SHORT_INDEX 517 +#define RL_GATE_WARMUP_STEPS_INDEX 524 +#define RL_GATE_DONES_EMA_INDEX 525 +#define RL_GATE_DONES_TARGET_INDEX 526 +#define RL_GATE_CONF_MIN_INDEX 527 +#define RL_GATE_CONF_MAX_INDEX 528 +#define RL_GATE_FRD_MIN_INDEX 529 +#define RL_GATE_FRD_MAX_INDEX 530 +#define RL_GATE_ADJUST_RATE_INDEX 531 + +extern "C" __global__ void rl_gate_threshold_controller( + float* __restrict__ isv, + const float* __restrict__ dones, // [B] + int b_size, + int current_step +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX]; + if (current_step < warmup) return; + + // Update dones EMA (count of done events this step / b_size). + float done_count = 0.0f; + for (int b = 0; b < b_size; ++b) { + if (dones[b] > 0.5f) done_count += 1.0f; + } + const float done_rate = done_count / (float)b_size; + const float alpha = 0.01f; + const float prev_ema = isv[RL_GATE_DONES_EMA_INDEX]; + const float dones_ema = (prev_ema == 0.0f && done_rate > 0.0f) + ? done_rate + : (1.0f - alpha) * prev_ema + alpha * done_rate; + isv[RL_GATE_DONES_EMA_INDEX] = dones_ema; + + const float target = isv[RL_GATE_DONES_TARGET_INDEX]; + const float adjust = isv[RL_GATE_ADJUST_RATE_INDEX]; + + // Schulman bounded step: if dones_ema < target, DECREASE thresholds + // (more permissive); if > target, INCREASE (more selective). + float conf_thr = isv[RL_CONF_GATE_THRESHOLD_INDEX]; + float frd_long = isv[RL_FRD_GATE_THR_LONG_INDEX]; + float frd_short = isv[RL_FRD_GATE_THR_SHORT_INDEX]; + + if (dones_ema < target * 0.5f) { + // Trades well below target — relax gates. + conf_thr *= adjust; + frd_long *= adjust; + frd_short *= adjust; + } else if (dones_ema > target * 2.0f) { + // Trades well above target — tighten gates. + conf_thr /= adjust; + frd_long /= adjust; + frd_short /= adjust; + } + + // Bilateral clamp. + const float conf_min = isv[RL_GATE_CONF_MIN_INDEX]; + const float conf_max = isv[RL_GATE_CONF_MAX_INDEX]; + const float frd_min = isv[RL_GATE_FRD_MIN_INDEX]; + const float frd_max = isv[RL_GATE_FRD_MAX_INDEX]; + + isv[RL_CONF_GATE_THRESHOLD_INDEX] = fmaxf(conf_min, fminf(conf_thr, conf_max)); + isv[RL_FRD_GATE_THR_LONG_INDEX] = fmaxf(frd_min, fminf(frd_long, frd_max)); + isv[RL_FRD_GATE_THR_SHORT_INDEX] = fmaxf(frd_min, fminf(frd_short, frd_max)); +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 0e57d8286..2f1ed77ef 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -893,6 +893,22 @@ pub const RL_MULTIRES_HORIZON_3_INDEX: usize = 523; /// Default 10000. pub const RL_GATE_WARMUP_STEPS_INDEX: usize = 524; +/// Gate threshold controller: dones EMA (streaming trade frequency). +pub const RL_GATE_DONES_EMA_INDEX: usize = 525; +/// Gate threshold controller: target dones rate per step. Default 0.02 +/// (1 trade per 50 steps at b=16 — ~2% of steps have a done event). +pub const RL_GATE_DONES_TARGET_INDEX: usize = 526; +/// Gate threshold controller: confidence gate minimum. +pub const RL_GATE_CONF_MIN_INDEX: usize = 527; +/// Gate threshold controller: confidence gate maximum. +pub const RL_GATE_CONF_MAX_INDEX: usize = 528; +/// Gate threshold controller: FRD gate minimum. +pub const RL_GATE_FRD_MIN_INDEX: usize = 529; +/// Gate threshold controller: FRD gate maximum. +pub const RL_GATE_FRD_MAX_INDEX: usize = 530; +/// Gate threshold controller: Schulman step rate. Default 0.95. +pub const RL_GATE_ADJUST_RATE_INDEX: usize = 531; + /// 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 = 525; +pub const RL_SLOTS_END: usize = 532; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index a76005fbd..11accb937 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -200,6 +200,8 @@ const RL_FRD_GATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_frd_gate.cubin")); const RL_RECENT_OUTCOME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_recent_outcome_update.cubin")); +const RL_GATE_THRESHOLD_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_gate_threshold_controller.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] = @@ -459,6 +461,8 @@ pub struct IntegratedTrainer { rl_frd_gate_fn: CudaFunction, _rl_recent_outcome_update_module: Arc, rl_recent_outcome_update_fn: CudaFunction, + _rl_gate_threshold_controller_module: Arc, + rl_gate_threshold_controller_fn: CudaFunction, _rl_trade_context_update_module: Arc, rl_trade_context_update_fn: CudaFunction, _rl_multires_features_update_module: Arc, @@ -969,6 +973,12 @@ impl IntegratedTrainer { let rl_recent_outcome_update_fn = rl_recent_outcome_update_module .load_function("rl_recent_outcome_update") .context("load rl_recent_outcome_update")?; + let rl_gate_threshold_controller_module = ctx + .load_cubin(RL_GATE_THRESHOLD_CONTROLLER_CUBIN.to_vec()) + .context("load rl_gate_threshold_controller cubin")?; + let rl_gate_threshold_controller_fn = rl_gate_threshold_controller_module + .load_function("rl_gate_threshold_controller") + .context("load rl_gate_threshold_controller")?; let rl_trade_context_update_module = ctx .load_cubin(RL_TRADE_CONTEXT_UPDATE_CUBIN.to_vec()) .context("load rl_trade_context_update cubin")?; @@ -1391,6 +1401,8 @@ impl IntegratedTrainer { rl_frd_gate_fn, _rl_recent_outcome_update_module: rl_recent_outcome_update_module, rl_recent_outcome_update_fn, + _rl_gate_threshold_controller_module: rl_gate_threshold_controller_module, + rl_gate_threshold_controller_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, @@ -1535,7 +1547,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 62] = [ + let isv_constants: [(usize, f32); 68] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1596,14 +1608,20 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_ANTIMARTINGALE_KAPPA_INDEX, 1.0), (crate::rl::isv_slots::RL_ANTIMARTINGALE_MIN_INDEX, 0.5), (crate::rl::isv_slots::RL_ANTIMARTINGALE_MAX_INDEX, 2.0), - (crate::rl::isv_slots::RL_CONF_GATE_THRESHOLD_INDEX, 0.10), + (crate::rl::isv_slots::RL_CONF_GATE_THRESHOLD_INDEX, 0.01), (crate::rl::isv_slots::RL_CONF_GATE_LAMBDA_INDEX, 1.0), (crate::rl::isv_slots::RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0), - (crate::rl::isv_slots::RL_FRD_GATE_THR_LONG_INDEX, 0.35), - (crate::rl::isv_slots::RL_FRD_GATE_THR_SHORT_INDEX, 0.35), + (crate::rl::isv_slots::RL_FRD_GATE_THR_LONG_INDEX, 0.15), + (crate::rl::isv_slots::RL_FRD_GATE_THR_SHORT_INDEX, 0.15), (crate::rl::isv_slots::RL_PI_SAMPLE_P_MIN_INDEX, 0.015), (crate::rl::isv_slots::RL_OUTCOME_ALPHA_INDEX, 0.1), (crate::rl::isv_slots::RL_GATE_WARMUP_STEPS_INDEX, 10000.0), + (crate::rl::isv_slots::RL_GATE_DONES_TARGET_INDEX, 0.02), + (crate::rl::isv_slots::RL_GATE_CONF_MIN_INDEX, 0.001), + (crate::rl::isv_slots::RL_GATE_CONF_MAX_INDEX, 0.50), + (crate::rl::isv_slots::RL_GATE_FRD_MIN_INDEX, 0.05), + (crate::rl::isv_slots::RL_GATE_FRD_MAX_INDEX, 0.50), + (crate::rl::isv_slots::RL_GATE_ADJUST_RATE_INDEX, 0.95), (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), @@ -4263,6 +4281,29 @@ impl IntegratedTrainer { self.launch_rl_controllers_per_step() .context("launch_rl_controllers_per_step")?; + // Gate threshold controller — adapts conf + FRD thresholds from + // trade frequency. Runs after other controllers, before reward scale. + { + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let current_step_i = (self.step_counter as i32).max(0); + let mut launch = self.stream.launch_builder(&self.rl_gate_threshold_controller_fn); + launch + .arg(&self.isv_d) + .arg(&self.dones_d) + .arg(&b_size_i) + .arg(¤t_step_i); + unsafe { + launch + .launch(cfg) + .context("rl_gate_threshold_controller launch")?; + } + } + // Apply the reward scale on device (in place on rewards_d) + // adaptive clamp + per-step pre-clamp diag + per-step // positive-tail max for the adaptive clamp controller.