feat(rl): FRD gate — override entries when forward-return is unfavorable
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
98
crates/ml-alpha/cuda/rl_frd_gate.cu
Normal file
98
crates/ml-alpha/cuda/rl_frd_gate.cu
Normal file
@@ -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 <stdint.h>
|
||||
#include <math.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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<CudaModule>,
|
||||
rl_confidence_gate_fn: CudaFunction,
|
||||
_rl_frd_gate_module: Arc<CudaModule>,
|
||||
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<i32>,
|
||||
frd_logits_d: &CudaSlice<f32>,
|
||||
pos_state_d: &CudaSlice<u8>,
|
||||
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<u8> = 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
|
||||
|
||||
@@ -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(())
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user