From e9ecacbdfa5ad150a76e2ce600b5505e63c17074 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 24 May 2026 21:43:15 +0200 Subject: [PATCH] =?UTF-8?q?feat(rl):=20FRD=20gate=20=E2=80=94=20override?= =?UTF-8?q?=20entries=20when=20forward-return=20is=20unfavorable?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit New rl_frd_gate.cu kernel reads the FRD head's horizon-2 (medium, ~300 ticks) categorical distribution. For long openings, sums probability mass in the positive tail (atoms > +0.5σ); for short openings, sums the negative tail (atoms < -0.5σ). Overrides to Hold when favorable mass < threshold. Fires after confidence gate, before trail/heat/market pipeline. Same preconditions: only gates flat positions with opening actions. ISV slots: 516 THR_LONG (0.35), 517 THR_SHORT (0.35), 518 fired_count (diag). RL_SLOTS_END → 519. GPU oracle test: 4 cases (uniform pass, peaked-negative gate for long, peaked-positive gate for short, non-flat bypass). Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 1 + crates/ml-alpha/cuda/rl_frd_gate.cu | 98 +++++++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 16 ++- crates/ml-alpha/src/trainer/integrated.rs | 81 +++++++++++++- .../tests/trade_management_kernels.rs | 104 +++++++++++++++++- 5 files changed, 296 insertions(+), 4 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_frd_gate.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 9177857e1..955d672b8 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -81,6 +81,7 @@ const KERNELS: &[&str] = &[ "rl_frd_layer1_bwd", // SP20 P3 F.3c: FRD head layer-1 backward — dW1 (per-batch scratch), db1 (per-batch scratch), dh_t (per-batch overwrite); applies ReLU mask via cached post-ReLU hidden; 1 block per batch, 128 threads "rl_position_heat_check", // SP20 P6: position heat cap — force-flat when |position_lots| exceeds ISV-driven max; last defense before actions_to_market_targets "rl_confidence_gate", // P8: override opening actions to Hold when C51 distributional Q is insufficiently confident (LCB < threshold) + "rl_frd_gate", // P9: override opening actions to Hold when FRD head predicts insufficient favorable probability mass ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_frd_gate.cu b/crates/ml-alpha/cuda/rl_frd_gate.cu new file mode 100644 index 000000000..2000dbb37 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_frd_gate.cu @@ -0,0 +1,98 @@ +// rl_frd_gate.cu — override opening actions to Hold when the FRD +// (Forward Return Distribution) head predicts insufficient favorable +// probability mass in the entry direction. +// +// For each batch where position==0 and the chosen action is an opening: +// - Long (a5/a6): compute entry_quality = Σ softmax(h2_logits)[i] +// for atoms i where bucket_center > +0.5σ. If < threshold → Hold. +// - Short (a0/a1): compute entry_quality = Σ softmax(h2_logits)[i] +// for atoms i where bucket_center < -0.5σ. If < threshold → Hold. +// +// FRD logits layout: [B, FRD_N_HORIZONS × FRD_N_ATOMS] where horizons +// are stacked contiguously. Horizon 2 (medium, ~300 ticks) is at +// offset h=1 (0-based) × FRD_N_ATOMS. +// +// Bucket centers span [-RANGE_SIGMA, +RANGE_SIGMA] uniformly over +// FRD_N_ATOMS atoms. With RANGE_SIGMA=3.0 and N=21: +// atom_i center = -3.0 + i × (6.0 / 20) = -3.0 + i × 0.3 +// > +0.5 when i >= 12 (center = -3.0 + 12*0.3 = +0.6) +// < -0.5 when i <= 8 (center = -3.0 + 8*0.3 = -0.6) +// +// One block per batch, single thread. No shared memory. +// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`. + +#include +#include + +#define FRD_N_HORIZONS 3 +#define FRD_N_ATOMS 21 +#define FRD_H2_OFFSET (1 * FRD_N_ATOMS) +#define FRD_FAVORABLE_LONG_START 12 +#define FRD_FAVORABLE_SHORT_END 9 +#define ACTION_HOLD 2 +#define RL_FRD_GATE_THR_LONG_INDEX 516 +#define RL_FRD_GATE_THR_SHORT_INDEX 517 +#define RL_FRD_GATE_FIRED_COUNT_INDEX 518 + +extern "C" __global__ void rl_frd_gate( + int* __restrict__ actions, // [B] IN/OUT + const float* __restrict__ frd_logits, // [B × FRD_N_HORIZONS × FRD_N_ATOMS] + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + float* __restrict__ isv, + int b_size, + int pos_bytes +) { + const int b = blockIdx.x; + if (b >= b_size) return; + + const int action = actions[b]; + + const bool is_long_open = (action == 5 || action == 6); + const bool is_short_open = (action == 0 || action == 1); + if (!is_long_open && !is_short_open) return; + + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + if (position_lots != 0) return; + + // Horizon 2 logits for this batch. + const float* h2 = frd_logits + b * (FRD_N_HORIZONS * FRD_N_ATOMS) + FRD_H2_OFFSET; + + // Numerically-stable softmax over h2 atoms. + float max_l = h2[0]; + for (int i = 1; i < FRD_N_ATOMS; ++i) { + if (h2[i] > max_l) max_l = h2[i]; + } + float sum_exp = 0.0f; + float probs[FRD_N_ATOMS]; + for (int i = 0; i < FRD_N_ATOMS; ++i) { + probs[i] = expf(h2[i] - max_l); + sum_exp += probs[i]; + } + const float inv_sum = 1.0f / sum_exp; + for (int i = 0; i < FRD_N_ATOMS; ++i) { + probs[i] *= inv_sum; + } + + float entry_quality = 0.0f; + if (is_long_open) { + // Favorable for long: atoms where bucket > +0.5σ. + for (int i = FRD_FAVORABLE_LONG_START; i < FRD_N_ATOMS; ++i) { + entry_quality += probs[i]; + } + const float thr = isv[RL_FRD_GATE_THR_LONG_INDEX]; + if (entry_quality < thr) { + actions[b] = ACTION_HOLD; + isv[RL_FRD_GATE_FIRED_COUNT_INDEX] += 1.0f; + } + } else { + // Favorable for short: atoms where bucket < -0.5σ. + for (int i = 0; i < FRD_FAVORABLE_SHORT_END; ++i) { + entry_quality += probs[i]; + } + const float thr = isv[RL_FRD_GATE_THR_SHORT_INDEX]; + if (entry_quality < thr) { + actions[b] = ACTION_HOLD; + isv[RL_FRD_GATE_FIRED_COUNT_INDEX] += 1.0f; + } + } +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 83003d82c..dfd6eecfc 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -857,6 +857,20 @@ pub const RL_CONF_GATE_SIGMA_NORM_INDEX: usize = 514; /// `rl_confidence_gate`; read by diag JSONL. pub const RL_CONF_GATE_FIRED_COUNT_INDEX: usize = 515; +/// FRD gate threshold for long entries. Opening long actions are +/// overridden to Hold when the FRD horizon-2 favorable mass (atoms +/// with center > +0.5σ) is below this value. Default: 0.35. +pub const RL_FRD_GATE_THR_LONG_INDEX: usize = 516; + +/// FRD gate threshold for short entries. Opening short actions are +/// overridden to Hold when the FRD horizon-2 favorable mass (atoms +/// with center < -0.5σ) is below this value. Default: 0.35. +pub const RL_FRD_GATE_THR_SHORT_INDEX: usize = 517; + +/// Diagnostic: number of batch entries where the FRD gate fired +/// (overrode to Hold) this step. +pub const RL_FRD_GATE_FIRED_COUNT_INDEX: usize = 518; + /// 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 = 516; +pub const RL_SLOTS_END: usize = 519; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index af76ef110..9d0666b52 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -196,6 +196,8 @@ const RL_POSITION_HEAT_CHECK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_position_heat_check.cubin")); const RL_CONFIDENCE_GATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_confidence_gate.cubin")); +const RL_FRD_GATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_frd_gate.cubin")); const RL_TRAIL_STOP_CHECK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_stop_check.cubin")); const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] = @@ -447,6 +449,8 @@ pub struct IntegratedTrainer { rl_position_heat_check_fn: CudaFunction, _rl_confidence_gate_module: Arc, rl_confidence_gate_fn: CudaFunction, + _rl_frd_gate_module: Arc, + rl_frd_gate_fn: CudaFunction, // Per-batch per-unit trade state (SP20 P1). // MAX_UNITS=4 reserved; only slot 0 used until SP20 P7 ships @@ -935,6 +939,12 @@ impl IntegratedTrainer { let rl_confidence_gate_fn = rl_confidence_gate_module .load_function("rl_confidence_gate") .context("load rl_confidence_gate")?; + let rl_frd_gate_module = ctx + .load_cubin(RL_FRD_GATE_CUBIN.to_vec()) + .context("load rl_frd_gate cubin")?; + let rl_frd_gate_fn = rl_frd_gate_module + .load_function("rl_frd_gate") + .context("load rl_frd_gate")?; let actions_to_market_targets_module = ctx .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) .context("load actions_to_market_targets cubin")?; @@ -1323,6 +1333,8 @@ impl IntegratedTrainer { rl_position_heat_check_fn, _rl_confidence_gate_module: rl_confidence_gate_module, rl_confidence_gate_fn, + _rl_frd_gate_module: rl_frd_gate_module, + rl_frd_gate_fn, unit_entry_price_d, unit_entry_step_d, unit_lots_d, @@ -1457,7 +1469,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 54] = [ + let isv_constants: [(usize, f32); 56] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1521,6 +1533,8 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_CONF_GATE_THRESHOLD_INDEX, 0.10), (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_KL_TARGET_INDEX, 0.01), (crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99), (crate::rl::isv_slots::RL_PLATEAU_PATIENCE_INDEX, 1000.0), @@ -2188,6 +2202,44 @@ impl IntegratedTrainer { Ok(()) } + /// Launch `rl_frd_gate` — for opening actions on flat positions, + /// check FRD head's horizon-2 favorable probability mass. Override + /// to Hold if below threshold. Public for GPU-oracle tests. + pub fn launch_rl_frd_gate( + &self, + actions_d: &mut CudaSlice, + frd_logits_d: &CudaSlice, + pos_state_d: &CudaSlice, + b_size: usize, + pos_bytes: usize, + ) -> Result<()> { + debug_assert_eq!(actions_d.len(), b_size); + let frd_out_dim = crate::rl::frd::FRD_OUT_DIM; + debug_assert_eq!(frd_logits_d.len(), b_size * frd_out_dim); + debug_assert_eq!(pos_state_d.len(), b_size * pos_bytes); + let cfg = LaunchConfig { + grid_dim: (b_size as u32, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let pos_bytes_i = pos_bytes as i32; + let mut launch = self.stream.launch_builder(&self.rl_frd_gate_fn); + launch + .arg(actions_d) + .arg(frd_logits_d) + .arg(pos_state_d) + .arg(&self.isv_d) + .arg(&b_size_i) + .arg(&pos_bytes_i); + unsafe { + launch + .launch(cfg) + .context("rl_frd_gate launch")?; + } + Ok(()) + } + /// Launch `rl_unit_state_update` — detects per-batch position /// transitions vs `prev_pos_lots_d` and updates per-unit state /// (entry_price, entry_step, lots, initial_r, trail_distance, @@ -3641,6 +3693,33 @@ 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. + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + let cfg = LaunchConfig { + grid_dim: (b_size as u32, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let pos_bytes_i = lobsim.pos_bytes() as i32; + let mut launch = self.stream.launch_builder(&self.rl_frd_gate_fn); + launch + .arg(&mut self.actions_d) + .arg(&self.frd_logits_d) + .arg(pos_d_ref) + .arg(&self.isv_d) + .arg(&b_size_i) + .arg(&pos_bytes_i); + unsafe { + launch + .launch(cfg) + .context("rl_frd_gate launch")?; + } + } + // ── Step 5 (Phase R6): GPU-pure env step. ───────────────────── // No per-batch host loop — actions land in lobsim's // market_targets_d via the `actions_to_market_targets` kernel diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 9f1c35854..0b5deb881 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -41,8 +41,8 @@ use anyhow::Result; use cudarc::driver::{CudaSlice, CudaStream}; use ml_alpha::rl::isv_slots::{ RL_CONF_GATE_LAMBDA_INDEX, RL_CONF_GATE_SIGMA_NORM_INDEX, RL_CONF_GATE_THRESHOLD_INDEX, - RL_PYRAMID_THRESHOLD_INDEX, RL_TRAIL_ADJUST_RATE_INDEX, RL_TRAIL_K_INIT_INDEX, - RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX, + RL_FRD_GATE_THR_LONG_INDEX, RL_FRD_GATE_THR_SHORT_INDEX, RL_PYRAMID_THRESHOLD_INDEX, + RL_TRAIL_ADJUST_RATE_INDEX, RL_TRAIL_K_INIT_INDEX, RL_TRAIL_MAX_INDEX, RL_TRAIL_MIN_INDEX, }; use ml_alpha::trainer::integrated::{ read_slice_d_pub, read_slice_i32_d_pub, read_slice_u8_d_pub, write_slice_f32_d_pub, @@ -55,6 +55,8 @@ use std::sync::Arc; const MAX_UNITS: usize = 4; const N_ACTIONS: usize = 11; const Q_N_ATOMS: usize = 21; +const FRD_N_HORIZONS: usize = 3; +const FRD_N_ATOMS: usize = 21; /// PosFlat byte layout (mirrors `ml_backtesting::lob::PosFlat`): /// offset 0: position_lots: i32 @@ -969,3 +971,101 @@ fn confidence_gate_overrides_low_confidence_opening() -> Result<()> { eprintln!("PASS — confidence gate fires on low-confidence openings, passes otherwise"); Ok(()) } + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn frd_gate_overrides_unfavorable_opening() -> Result<()> { + let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + let pos_bytes = POS_BYTES; + let b_size = 1; + + set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_LONG_INDEX, 0.35)?; + set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_SHORT_INDEX, 0.35)?; + + let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + let frd_dim = FRD_N_HORIZONS * FRD_N_ATOMS; + + // Case 1: Uniform FRD logits → each atom gets 1/21 ≈ 0.048. + // Favorable mass for long = 9 atoms × 0.048 = 0.43 > 0.35 → NO gate. + let uniform_frd = vec![0.0_f32; b_size * frd_dim]; + let frd_d = upload_f32(&stream, &uniform_frd)?; + let mut actions_d = upload_i32(&stream, &[5])?; + + trainer.launch_rl_frd_gate( + &mut actions_d, + &frd_d, + &pos_flat_d, + b_size, + pos_bytes, + )?; + let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + assert_eq!( + actions[0], 5, + "uniform FRD (9/21=0.43 > 0.35) should NOT gate long; got {}", + actions[0] + ); + + // Case 2: FRD peaked at negative atoms for h2 → unfavorable for long. + // Peak at atom 2 (center = -3.0 + 2*0.3 = -2.4σ) → nearly all mass + // in negative tail → favorable long mass ≈ 0. + let mut peaked_neg = vec![0.0_f32; b_size * frd_dim]; + // h2 is at offset FRD_N_ATOMS (horizon index 1). + peaked_neg[FRD_N_ATOMS + 2] = 10.0; + let frd_neg_d = upload_f32(&stream, &peaked_neg)?; + let mut actions_d2 = upload_i32(&stream, &[6])?; // LongLarge + + trainer.launch_rl_frd_gate( + &mut actions_d2, + &frd_neg_d, + &pos_flat_d, + b_size, + pos_bytes, + )?; + let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?; + assert_eq!( + actions[0], 2, + "FRD peaked negative (unfavorable for long) should gate to Hold; got {}", + actions[0] + ); + + // Case 3: Short opening, FRD peaked at positive atoms → unfavorable. + let mut peaked_pos = vec![0.0_f32; b_size * frd_dim]; + peaked_pos[FRD_N_ATOMS + 18] = 10.0; // h2 peaked at atom 18 (+2.4σ) + let frd_pos_d = upload_f32(&stream, &peaked_pos)?; + let mut actions_d3 = upload_i32(&stream, &[0])?; // ShortLarge + + trainer.launch_rl_frd_gate( + &mut actions_d3, + &frd_pos_d, + &pos_flat_d, + b_size, + pos_bytes, + )?; + let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?; + assert_eq!( + actions[0], 2, + "FRD peaked positive (unfavorable for short) should gate to Hold; got {}", + actions[0] + ); + + // Case 4: Non-flat position → gate must not fire. + let pos_long_d = upload_u8(&stream, &pos_buf(2, 100.0))?; + let mut actions_d4 = upload_i32(&stream, &[5])?; + trainer.launch_rl_frd_gate( + &mut actions_d4, + &frd_neg_d, // would gate if flat + &pos_long_d, + b_size, + pos_bytes, + )?; + let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?; + assert_eq!( + actions[0], 5, + "non-flat position → FRD gate must not fire; got {}", + actions[0] + ); + + eprintln!("PASS — FRD gate fires on unfavorable openings, passes otherwise"); + Ok(()) +}