diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index e98bff8ff..9177857e1 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -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 diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu new file mode 100644 index 000000000..f799c12ba --- /dev/null +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -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 +#include + +#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; + } +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index cf65d1680..83003d82c 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -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; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 26e7edcbb..af76ef110 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -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, rl_position_heat_check_fn: CudaFunction, + _rl_confidence_gate_module: Arc, + 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, + q_logits_d: &CudaSlice, + pos_state_d: &CudaSlice, + 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 = 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 diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index abdae6048..9f1c35854 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -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(()) +}