From 82481db06d9d1bbf58fa364bddb450325341ad43 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 25 May 2026 10:16:05 +0200 Subject: [PATCH] =?UTF-8?q?feat(rl):=20gate=20warmup=20=E2=80=94=20confide?= =?UTF-8?q?nce=20+=20FRD=20gates=20inactive=20for=20first=2010k=20steps?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Both gates now read RL_GATE_WARMUP_STEPS_INDEX (slot 524, default 10000) and return early when current_step < warmup. During warmup, the agent opens positions freely, collects reward signal, and calibrates Q. After warmup, gates activate and filter low-quality entries. Without warmup, gates blocked 100% of opening actions from step 0 (uniform Q → zero confidence → permanent Hold attractor → zero trades → zero reward → Q never learns). Confirmed by 50k-step L40S run with zero dones across 800k action decisions. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/rl_confidence_gate.cu | 7 +++++- crates/ml-alpha/cuda/rl_frd_gate.cu | 7 +++++- crates/ml-alpha/src/rl/isv_slots.rs | 8 ++++++- crates/ml-alpha/src/trainer/integrated.rs | 23 ++++++++++++------- .../tests/trade_management_kernels.rs | 10 ++++++++ 5 files changed, 44 insertions(+), 11 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu index f799c12ba..bf40c024e 100644 --- a/crates/ml-alpha/cuda/rl_confidence_gate.cu +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -29,6 +29,7 @@ #define RL_CONF_GATE_LAMBDA_INDEX 513 #define RL_CONF_GATE_SIGMA_NORM_INDEX 514 #define RL_CONF_GATE_FIRED_COUNT_INDEX 515 +#define RL_GATE_WARMUP_STEPS_INDEX 524 extern "C" __global__ void rl_confidence_gate( int* __restrict__ actions, // [B] IN/OUT @@ -37,11 +38,15 @@ extern "C" __global__ void rl_confidence_gate( const unsigned char* __restrict__ pos_state, // [B × pos_bytes] float* __restrict__ isv, int b_size, - int pos_bytes + int pos_bytes, + int current_step ) { const int b = blockIdx.x; if (b >= b_size) return; + const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX]; + if (current_step < warmup) return; + const int action = actions[b]; const bool is_opening = (action == 0 || action == 1 || action == 5 || action == 6); if (!is_opening) return; diff --git a/crates/ml-alpha/cuda/rl_frd_gate.cu b/crates/ml-alpha/cuda/rl_frd_gate.cu index 2000dbb37..39680999e 100644 --- a/crates/ml-alpha/cuda/rl_frd_gate.cu +++ b/crates/ml-alpha/cuda/rl_frd_gate.cu @@ -33,6 +33,7 @@ #define RL_FRD_GATE_THR_LONG_INDEX 516 #define RL_FRD_GATE_THR_SHORT_INDEX 517 #define RL_FRD_GATE_FIRED_COUNT_INDEX 518 +#define RL_GATE_WARMUP_STEPS_INDEX 524 extern "C" __global__ void rl_frd_gate( int* __restrict__ actions, // [B] IN/OUT @@ -40,11 +41,15 @@ extern "C" __global__ void rl_frd_gate( const unsigned char* __restrict__ pos_state, // [B × pos_bytes] float* __restrict__ isv, int b_size, - int pos_bytes + int pos_bytes, + int current_step ) { const int b = blockIdx.x; if (b >= b_size) return; + const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX]; + if (current_step < warmup) return; + const int action = actions[b]; const bool is_long_open = (action == 5 || action == 6); diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index a58195f56..0e57d8286 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -887,6 +887,12 @@ pub const RL_MULTIRES_HORIZON_2_INDEX: usize = 522; /// Multi-resolution feature horizon 3 (seconds). Default 600.0. pub const RL_MULTIRES_HORIZON_3_INDEX: usize = 523; +/// Gate warmup: number of steps before confidence + FRD gates +/// activate. During warmup, gates pass all actions through so the +/// agent can open positions, collect reward signal, and calibrate Q. +/// Default 10000. +pub const RL_GATE_WARMUP_STEPS_INDEX: usize = 524; + /// 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 = 524; +pub const RL_SLOTS_END: usize = 525; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index bffdc67f8..a76005fbd 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1535,7 +1535,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 61] = [ + let isv_constants: [(usize, f32); 62] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1603,6 +1603,7 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_FRD_GATE_THR_SHORT_INDEX, 0.35), (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_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), @@ -2248,6 +2249,7 @@ impl IntegratedTrainer { pos_state_d: &CudaSlice, b_size: usize, pos_bytes: usize, + current_step: i32, ) -> Result<()> { debug_assert_eq!(actions_d.len(), b_size); debug_assert_eq!(q_logits_d.len(), b_size * N_ACTIONS * Q_N_ATOMS); @@ -2267,7 +2269,8 @@ impl IntegratedTrainer { .arg(pos_state_d) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i); + .arg(&pos_bytes_i) + .arg(¤t_step); unsafe { launch .launch(cfg) @@ -2286,6 +2289,7 @@ impl IntegratedTrainer { pos_state_d: &CudaSlice, b_size: usize, pos_bytes: usize, + current_step: i32, ) -> Result<()> { debug_assert_eq!(actions_d.len(), b_size); let frd_out_dim = crate::rl::frd::FRD_OUT_DIM; @@ -2305,7 +2309,8 @@ impl IntegratedTrainer { .arg(pos_state_d) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i); + .arg(&pos_bytes_i) + .arg(¤t_step); unsafe { launch .launch(cfg) @@ -3758,6 +3763,7 @@ impl IntegratedTrainer { }; let b_size_i = b_size as i32; let pos_bytes_i = lobsim.pos_bytes() as i32; + let current_step_i = (self.step_counter as i32).max(0); let mut launch = self.stream.launch_builder(&self.rl_confidence_gate_fn); launch .arg(&mut self.actions_d) @@ -3766,7 +3772,8 @@ impl IntegratedTrainer { .arg(pos_d_ref) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i); + .arg(&pos_bytes_i) + .arg(¤t_step_i); unsafe { launch .launch(cfg) @@ -3774,9 +3781,7 @@ impl IntegratedTrainer { } } - // ── FRD gate — override opening actions to Hold when FRD - // head's horizon-2 favorable mass is below threshold. Fires - // after confidence gate, same position/action preconditions. + // FRD gate — same warmup as confidence gate. { let pos_d_ref: &CudaSlice = lobsim.pos_d(); let cfg = LaunchConfig { @@ -3786,6 +3791,7 @@ impl IntegratedTrainer { }; let b_size_i = b_size as i32; let pos_bytes_i = lobsim.pos_bytes() as i32; + let current_step_i = (self.step_counter as i32).max(0); let mut launch = self.stream.launch_builder(&self.rl_frd_gate_fn); launch .arg(&mut self.actions_d) @@ -3793,7 +3799,8 @@ impl IntegratedTrainer { .arg(pos_d_ref) .arg(&self.isv_d) .arg(&b_size_i) - .arg(&pos_bytes_i); + .arg(&pos_bytes_i) + .arg(¤t_step_i); unsafe { launch .launch(cfg) diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 546bbe2e6..8f2a63536 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -918,6 +918,7 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -942,6 +943,7 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?; assert_eq!( @@ -959,6 +961,7 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_long_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?; assert_eq!( @@ -975,6 +978,7 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?; assert_eq!( @@ -1013,6 +1017,7 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -1036,6 +1041,7 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?; assert_eq!( @@ -1056,6 +1062,7 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?; assert_eq!( @@ -1073,6 +1080,7 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { &pos_long_d, b_size, pos_bytes, + 99999, )?; let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?; assert_eq!( @@ -1386,6 +1394,7 @@ fn both_gates_block_same_action() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions_after_conf = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!( @@ -1400,6 +1409,7 @@ fn both_gates_block_same_action() -> Result<()> { &pos_flat_d, b_size, pos_bytes, + 99999, )?; let actions_after_frd = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; assert_eq!(