feat(rl): confidence gate — override low-certainty openings to Hold
New rl_confidence_gate.cu kernel computes C51 distributional Lower
Confidence Bound (μ - λσ) / σ_norm for the chosen action. When
position is flat and the selected action is an opening (a0/a1/a5/a6),
overrides to Hold if conf < threshold.
Fires after π action selection, before trail/heat/market pipeline.
Only gates on flat positions — existing positions pass through
unconditionally regardless of Q uncertainty.
ISV slots: 512 threshold (0.10), 513 λ (1.0), 514 σ_norm (1.0),
515 fired_count (diag). RL_SLOTS_END → 516.
GPU oracle test: 4 cases (low-conf gate, high-conf pass, non-flat
bypass, non-opening bypass).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -80,6 +80,7 @@ const KERNELS: &[&str] = &[
|
||||
"rl_frd_layer2_bwd", // SP20 P3 F.3b: FRD head layer-2 backward — dW2 (per-batch scratch), db2 (per-batch scratch), dhidden (per-batch overwrite); 1 block per batch, 64 threads
|
||||
"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)
|
||||
];
|
||||
|
||||
// Cache bust v31 — five new reduce / derive kernels populate the input
|
||||
|
||||
96
crates/ml-alpha/cuda/rl_confidence_gate.cu
Normal file
96
crates/ml-alpha/cuda/rl_confidence_gate.cu
Normal file
@@ -0,0 +1,96 @@
|
||||
// rl_confidence_gate.cu — override opening actions to Hold when the
|
||||
// C51 distributional Q is insufficiently confident about the chosen
|
||||
// action.
|
||||
//
|
||||
// For each batch where position==0 and the chosen action is an opening
|
||||
// action (a0, a1, a5, a6):
|
||||
// 1. Extract q_logits[b, a*, 0..Q_N_ATOMS]
|
||||
// 2. Numerically-stable softmax → probs
|
||||
// 3. μ = Σ probs[i] × atom_supports[i]
|
||||
// 4. σ² = Σ probs[i] × (atom_supports[i] - μ)²
|
||||
// 5. conf = clamp((μ - λ×√σ²) / σ_norm, 0, 1)
|
||||
// 6. If conf < threshold → override actions[b] = 2 (Hold)
|
||||
//
|
||||
// One block per batch, single thread (Q_N_ATOMS=21 is small enough
|
||||
// for sequential softmax + reduce). No shared memory needed.
|
||||
//
|
||||
// Per `feedback_no_atomicadd`: fired-count written by thread 0 via
|
||||
// simple sequential count (b_size=1 in practice; multi-batch uses
|
||||
// block-serial approach consistent with other diag counters).
|
||||
// Per `feedback_cpu_is_read_only`: pure device-side.
|
||||
|
||||
#include <stdint.h>
|
||||
#include <math.h>
|
||||
|
||||
#define Q_N_ATOMS 21
|
||||
#define N_ACTIONS 11
|
||||
#define ACTION_HOLD 2
|
||||
#define RL_CONF_GATE_THRESHOLD_INDEX 512
|
||||
#define RL_CONF_GATE_LAMBDA_INDEX 513
|
||||
#define RL_CONF_GATE_SIGMA_NORM_INDEX 514
|
||||
#define RL_CONF_GATE_FIRED_COUNT_INDEX 515
|
||||
|
||||
extern "C" __global__ void rl_confidence_gate(
|
||||
int* __restrict__ actions, // [B] IN/OUT
|
||||
const float* __restrict__ q_logits, // [B × N_ACTIONS × Q_N_ATOMS]
|
||||
const float* __restrict__ atom_supports, // [Q_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_opening = (action == 0 || action == 1 || action == 5 || action == 6);
|
||||
if (!is_opening) return;
|
||||
|
||||
const int position_lots = *(const int*)(pos_state + b * pos_bytes);
|
||||
if (position_lots != 0) return;
|
||||
|
||||
const float threshold = isv[RL_CONF_GATE_THRESHOLD_INDEX];
|
||||
const float lambda = isv[RL_CONF_GATE_LAMBDA_INDEX];
|
||||
const float sigma_norm = isv[RL_CONF_GATE_SIGMA_NORM_INDEX];
|
||||
|
||||
// q_logits layout: [B, N_ACTIONS, Q_N_ATOMS]
|
||||
const float* logits = q_logits + b * N_ACTIONS * Q_N_ATOMS + action * Q_N_ATOMS;
|
||||
|
||||
// Numerically-stable softmax.
|
||||
float max_l = logits[0];
|
||||
for (int i = 1; i < Q_N_ATOMS; ++i) {
|
||||
if (logits[i] > max_l) max_l = logits[i];
|
||||
}
|
||||
float sum_exp = 0.0f;
|
||||
float probs[Q_N_ATOMS];
|
||||
for (int i = 0; i < Q_N_ATOMS; ++i) {
|
||||
probs[i] = expf(logits[i] - max_l);
|
||||
sum_exp += probs[i];
|
||||
}
|
||||
const float inv_sum = 1.0f / sum_exp;
|
||||
for (int i = 0; i < Q_N_ATOMS; ++i) {
|
||||
probs[i] *= inv_sum;
|
||||
}
|
||||
|
||||
// μ and σ² from the categorical distribution.
|
||||
float mu = 0.0f;
|
||||
for (int i = 0; i < Q_N_ATOMS; ++i) {
|
||||
mu += probs[i] * atom_supports[i];
|
||||
}
|
||||
float var = 0.0f;
|
||||
for (int i = 0; i < Q_N_ATOMS; ++i) {
|
||||
const float diff = atom_supports[i] - mu;
|
||||
var += probs[i] * diff * diff;
|
||||
}
|
||||
const float sigma = sqrtf(var + 1e-8f);
|
||||
|
||||
// Lower confidence bound.
|
||||
const float lcb = mu - lambda * sigma;
|
||||
const float conf = fmaxf(0.0f, fminf(lcb / (sigma_norm + 1e-8f), 1.0f));
|
||||
|
||||
if (conf < threshold) {
|
||||
actions[b] = ACTION_HOLD;
|
||||
// Diag: count fires. Single-thread kernel so direct write is safe.
|
||||
isv[RL_CONF_GATE_FIRED_COUNT_INDEX] += 1.0f;
|
||||
}
|
||||
}
|
||||
@@ -837,6 +837,26 @@ pub const RL_ANTIMARTINGALE_MIN_INDEX: usize = 510;
|
||||
/// (never exceed 2× base order size).
|
||||
pub const RL_ANTIMARTINGALE_MAX_INDEX: usize = 511;
|
||||
|
||||
/// Confidence gate threshold. Opening actions are overridden to Hold
|
||||
/// when `conf < THR`. Default: 0.10 (10% — only gates when the Q
|
||||
/// distribution is extremely uncertain about the chosen action).
|
||||
pub const RL_CONF_GATE_THRESHOLD_INDEX: usize = 512;
|
||||
|
||||
/// Confidence gate λ — number of standard deviations subtracted from
|
||||
/// the mean in the lower-confidence-bound formula:
|
||||
/// `conf = clamp((μ - λ×σ) / σ_norm, 0, 1)`. Default: 1.0.
|
||||
pub const RL_CONF_GATE_LAMBDA_INDEX: usize = 513;
|
||||
|
||||
/// Confidence gate σ_norm — normalization divisor for the LCB signal.
|
||||
/// Converts the raw (μ - λσ) value to [0, 1] scale. Default: 1.0
|
||||
/// (= the full atom span, so conf=1 when LCB equals V_MAX).
|
||||
pub const RL_CONF_GATE_SIGMA_NORM_INDEX: usize = 514;
|
||||
|
||||
/// Diagnostic: number of batch entries where the confidence gate
|
||||
/// fired (overrode to Hold) this step. Written by
|
||||
/// `rl_confidence_gate`; read by diag JSONL.
|
||||
pub const RL_CONF_GATE_FIRED_COUNT_INDEX: usize = 515;
|
||||
|
||||
/// 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 = 512;
|
||||
pub const RL_SLOTS_END: usize = 516;
|
||||
|
||||
@@ -194,6 +194,8 @@ const RL_TRAIL_MUTATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_mutate.cubin"));
|
||||
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_TRAIL_STOP_CHECK_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_trail_stop_check.cubin"));
|
||||
const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] =
|
||||
@@ -443,6 +445,8 @@ pub struct IntegratedTrainer {
|
||||
rl_trail_stop_check_fn: CudaFunction,
|
||||
_rl_position_heat_check_module: Arc<CudaModule>,
|
||||
rl_position_heat_check_fn: CudaFunction,
|
||||
_rl_confidence_gate_module: Arc<CudaModule>,
|
||||
rl_confidence_gate_fn: CudaFunction,
|
||||
|
||||
// Per-batch per-unit trade state (SP20 P1).
|
||||
// MAX_UNITS=4 reserved; only slot 0 used until SP20 P7 ships
|
||||
@@ -925,6 +929,12 @@ impl IntegratedTrainer {
|
||||
let rl_position_heat_check_fn = rl_position_heat_check_module
|
||||
.load_function("rl_position_heat_check")
|
||||
.context("load rl_position_heat_check")?;
|
||||
let rl_confidence_gate_module = ctx
|
||||
.load_cubin(RL_CONFIDENCE_GATE_CUBIN.to_vec())
|
||||
.context("load rl_confidence_gate cubin")?;
|
||||
let rl_confidence_gate_fn = rl_confidence_gate_module
|
||||
.load_function("rl_confidence_gate")
|
||||
.context("load rl_confidence_gate")?;
|
||||
let actions_to_market_targets_module = ctx
|
||||
.load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec())
|
||||
.context("load actions_to_market_targets cubin")?;
|
||||
@@ -1311,6 +1321,8 @@ impl IntegratedTrainer {
|
||||
rl_trail_stop_check_fn,
|
||||
_rl_position_heat_check_module: rl_position_heat_check_module,
|
||||
rl_position_heat_check_fn,
|
||||
_rl_confidence_gate_module: rl_confidence_gate_module,
|
||||
rl_confidence_gate_fn,
|
||||
unit_entry_price_d,
|
||||
unit_entry_step_d,
|
||||
unit_lots_d,
|
||||
@@ -1445,7 +1457,7 @@ impl IntegratedTrainer {
|
||||
// (slot, value) pair — pure device write, no HtoD per
|
||||
// `feedback_no_htod_htoh_only_mapped_pinned`.
|
||||
{
|
||||
let isv_constants: [(usize, f32); 51] = [
|
||||
let isv_constants: [(usize, f32); 54] = [
|
||||
// Static seeds for the adaptive reward-clamp controller —
|
||||
// these are the initial values that
|
||||
// `rl_reward_clamp_controller` will replace once it
|
||||
@@ -1506,6 +1518,9 @@ 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_LAMBDA_INDEX, 1.0),
|
||||
(crate::rl::isv_slots::RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0),
|
||||
(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),
|
||||
@@ -2134,6 +2149,45 @@ impl IntegratedTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Launch `rl_confidence_gate` — for opening actions on flat
|
||||
/// positions, compute C51 distributional LCB and override to Hold
|
||||
/// if below threshold. Public for GPU-oracle tests.
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
pub fn launch_rl_confidence_gate(
|
||||
&self,
|
||||
actions_d: &mut CudaSlice<i32>,
|
||||
q_logits_d: &CudaSlice<f32>,
|
||||
pos_state_d: &CudaSlice<u8>,
|
||||
b_size: usize,
|
||||
pos_bytes: usize,
|
||||
) -> Result<()> {
|
||||
debug_assert_eq!(actions_d.len(), b_size);
|
||||
debug_assert_eq!(q_logits_d.len(), b_size * N_ACTIONS * Q_N_ATOMS);
|
||||
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_confidence_gate_fn);
|
||||
launch
|
||||
.arg(actions_d)
|
||||
.arg(q_logits_d)
|
||||
.arg(&self.atom_supports_d)
|
||||
.arg(pos_state_d)
|
||||
.arg(&self.isv_d)
|
||||
.arg(&b_size_i)
|
||||
.arg(&pos_bytes_i);
|
||||
unsafe {
|
||||
launch
|
||||
.launch(cfg)
|
||||
.context("rl_confidence_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,
|
||||
@@ -3557,6 +3611,36 @@ impl IntegratedTrainer {
|
||||
}
|
||||
}
|
||||
|
||||
// ── Confidence gate — override opening actions to Hold when C51
|
||||
// distributional Q has high uncertainty for the chosen action.
|
||||
// Fires AFTER π samples an action but BEFORE trail/heat/market
|
||||
// pipeline. Reads pos_state from lobsim (pre-snapshot = current
|
||||
// position) to only gate on flat entries.
|
||||
{
|
||||
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_confidence_gate_fn);
|
||||
launch
|
||||
.arg(&mut self.actions_d)
|
||||
.arg(&q_logits_d)
|
||||
.arg(&self.atom_supports_d)
|
||||
.arg(pos_d_ref)
|
||||
.arg(&self.isv_d)
|
||||
.arg(&b_size_i)
|
||||
.arg(&pos_bytes_i);
|
||||
unsafe {
|
||||
launch
|
||||
.launch(cfg)
|
||||
.context("rl_confidence_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
|
||||
|
||||
@@ -40,6 +40,7 @@
|
||||
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,
|
||||
};
|
||||
@@ -52,6 +53,8 @@ use ml_core::device::MlDevice;
|
||||
use std::sync::Arc;
|
||||
|
||||
const MAX_UNITS: usize = 4;
|
||||
const N_ACTIONS: usize = 11;
|
||||
const Q_N_ATOMS: usize = 21;
|
||||
|
||||
/// PosFlat byte layout (mirrors `ml_backtesting::lob::PosFlat`):
|
||||
/// offset 0: position_lots: i32
|
||||
@@ -868,3 +871,101 @@ fn pyramid_threshold_gates_add_in_actions_kernel() -> Result<()> {
|
||||
eprintln!("PASS — pyramid threshold gates add/replace/no-op in actions kernel");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
#[test]
|
||||
#[ignore = "requires CUDA (MlDevice::cuda(0))"]
|
||||
fn confidence_gate_overrides_low_confidence_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 gate params: threshold=0.5, lambda=1.0, sigma_norm=1.0
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_THRESHOLD_INDEX, 0.5)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_LAMBDA_INDEX, 1.0)?;
|
||||
set_isv_slot(&mut trainer, &stream, RL_CONF_GATE_SIGMA_NORM_INDEX, 1.0)?;
|
||||
|
||||
// Flat position (opening action should be gated).
|
||||
let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?;
|
||||
|
||||
// Case 1: Uniform Q logits → high variance, low confidence → gate fires.
|
||||
// With uniform logits, softmax → 1/21 each. μ ≈ midpoint of atom span.
|
||||
// σ will be large relative to μ → LCB = μ - 1×σ likely < 0 → conf = 0.
|
||||
let uniform_logits = vec![0.0_f32; b_size * N_ACTIONS * Q_N_ATOMS];
|
||||
let q_logits_d = upload_f32(&stream, &uniform_logits)?;
|
||||
let mut actions_d = upload_i32(&stream, &[5])?; // LongSmall — opening
|
||||
|
||||
trainer.launch_rl_confidence_gate(
|
||||
&mut actions_d,
|
||||
&q_logits_d,
|
||||
&pos_flat_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?;
|
||||
assert_eq!(
|
||||
actions[0], 2,
|
||||
"uniform Q (high uncertainty) should gate opening to Hold (a2); got {}",
|
||||
actions[0]
|
||||
);
|
||||
|
||||
// Case 2: Peaked Q logits for action 5 → high confidence → no gate.
|
||||
// Set logits for action 5 to a peaked distribution at the high end
|
||||
// of the atom span → μ high, σ low → LCB/σ_norm > threshold.
|
||||
let mut peaked_logits = vec![0.0_f32; b_size * N_ACTIONS * Q_N_ATOMS];
|
||||
// Action 5 logits: peaked at atom 18 (near V_MAX).
|
||||
let action5_offset = 5 * Q_N_ATOMS;
|
||||
peaked_logits[action5_offset + 18] = 10.0; // dominant atom
|
||||
let q_peaked_d = upload_f32(&stream, &peaked_logits)?;
|
||||
let mut actions_d2 = upload_i32(&stream, &[5])?;
|
||||
|
||||
trainer.launch_rl_confidence_gate(
|
||||
&mut actions_d2,
|
||||
&q_peaked_d,
|
||||
&pos_flat_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions = read_slice_i32_d_pub(&stream, &actions_d2, b_size)?;
|
||||
assert_eq!(
|
||||
actions[0], 5,
|
||||
"peaked Q (high confidence) should NOT gate; got {}",
|
||||
actions[0]
|
||||
);
|
||||
|
||||
// Case 3: Non-flat position → gate should NOT fire regardless of Q.
|
||||
let pos_long_d = upload_u8(&stream, &pos_buf(2, 100.0))?;
|
||||
let mut actions_d3 = upload_i32(&stream, &[5])?;
|
||||
trainer.launch_rl_confidence_gate(
|
||||
&mut actions_d3,
|
||||
&q_logits_d, // uniform (would gate if flat)
|
||||
&pos_long_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions = read_slice_i32_d_pub(&stream, &actions_d3, b_size)?;
|
||||
assert_eq!(
|
||||
actions[0], 5,
|
||||
"non-flat position → gate must not fire; got {}",
|
||||
actions[0]
|
||||
);
|
||||
|
||||
// Case 4: Hold action (a2) → not an opening → gate must not fire.
|
||||
let mut actions_d4 = upload_i32(&stream, &[2])?;
|
||||
trainer.launch_rl_confidence_gate(
|
||||
&mut actions_d4,
|
||||
&q_logits_d,
|
||||
&pos_flat_d,
|
||||
b_size,
|
||||
pos_bytes,
|
||||
)?;
|
||||
let actions = read_slice_i32_d_pub(&stream, &actions_d4, b_size)?;
|
||||
assert_eq!(
|
||||
actions[0], 2,
|
||||
"Hold action → not an opening → unchanged; got {}",
|
||||
actions[0]
|
||||
);
|
||||
|
||||
eprintln!("PASS — confidence gate fires on low-confidence openings, passes otherwise");
|
||||
Ok(())
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user