feat(rl): adaptive gate threshold controller from trade frequency
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
80
crates/ml-alpha/cuda/rl_gate_threshold_controller.cu
Normal file
80
crates/ml-alpha/cuda/rl_gate_threshold_controller.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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));
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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<CudaModule>,
|
||||
rl_recent_outcome_update_fn: CudaFunction,
|
||||
_rl_gate_threshold_controller_module: Arc<CudaModule>,
|
||||
rl_gate_threshold_controller_fn: CudaFunction,
|
||||
_rl_trade_context_update_module: Arc<CudaModule>,
|
||||
rl_trade_context_update_fn: CudaFunction,
|
||||
_rl_multires_features_update_module: Arc<CudaModule>,
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user