gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(rl): gate warmup — confidence + FRD gates inactive for first 10k steps

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-25 10:16:05 +02:00
parent 50b78e2430
commit 82481db06d
5 changed files with 44 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
crates/ml-alpha/cuda/rl_confidence_gate.cu
View File

@@ -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;

7
crates/ml-alpha/cuda/rl_frd_gate.cu
View File

@@ -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);

8
crates/ml-alpha/src/rl/isv_slots.rs
View File

@@ -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;

23
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -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<u8>,
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(&current_step);
unsafe {
launch
.launch(cfg)
@@ -2286,6 +2289,7 @@ impl IntegratedTrainer {
pos_state_d: &CudaSlice<u8>,
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(&current_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(&current_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<u8> = 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(&current_step_i);
unsafe {
launch
.launch(cfg)

10
crates/ml-alpha/tests/trade_management_kernels.rs
View File

@@ -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!(