diff --git a/.claude/scheduled_tasks.lock b/.claude/scheduled_tasks.lock new file mode 100644 index 000000000..0dfa75159 --- /dev/null +++ b/.claude/scheduled_tasks.lock @@ -0,0 +1 @@ +{"sessionId":"edd8cbc4-f0a4-4090-9e9a-424c797dee3d","pid":1676949,"procStart":"6609584","acquiredAt":1779464404408} \ No newline at end of file diff --git a/crates/ml-alpha/cuda/actions_to_market_targets.cu b/crates/ml-alpha/cuda/actions_to_market_targets.cu index 5fc2608d6..1d56f49df 100644 --- a/crates/ml-alpha/cuda/actions_to_market_targets.cu +++ b/crates/ml-alpha/cuda/actions_to_market_targets.cu @@ -18,8 +18,10 @@ // 4 FlatFromShort side=0 (buy), size=|position_lots| (only if pos < 0) // 5 LongSmall side=0 (buy), size=1 // 6 LongLarge side=0 (buy), size=2 -// 7 TrailTighten side=2 (no-op),size=0 — handled as ISV mutation, not a fill -// 8 TrailLoosen side=2 (no-op),size=0 — handled as ISV mutation, not a fill +// 7 TrailTighten side=2 (no-op),size=0 — handled by rl_trail_mutate kernel +// 8 TrailLoosen side=2 (no-op),size=0 — handled by rl_trail_mutate kernel +// 9 HalfFlatLong side=1 (sell), size=max(1, position_lots/2) if pos>0 (SP20 P4) +// 10 HalfFlatShort side=0 (buy), size=max(1, |position_lots|/2) if pos<0 (SP20 P4) // // Actions 3/4 require reading the current `position_lots` from the // Pos struct (`offset 0`, i32). Other actions are pure constant @@ -68,11 +70,28 @@ extern "C" __global__ void actions_to_market_targets( side = 0; size = 1; } else if (action == 6) { side = 0; size = 2; + } else if (action == 9) { + // SP20 P4 HalfFlatLong: close ⌈|pos|/2⌉ lots of long position. + // No-op if not long. + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + if (position_lots > 0) { + side = 1; + const int half = (position_lots + 1) / 2; // round-up so min 1 + size = (half > 0) ? half : 1; + } + } else if (action == 10) { + // SP20 P4 HalfFlatShort: close ⌈|pos|/2⌉ lots of short position. + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + if (position_lots < 0) { + side = 0; + const int abs_pos = -position_lots; + const int half = (abs_pos + 1) / 2; + size = (half > 0) ? half : 1; + } } - // actions 7, 8 (TrailTighten/Loosen): no fill. These actions - // operate via ISV trail-distance mutation, not market-order - // submission — the side=2 no-op is correct here. The ISV - // mutation path lives outside this kernel. + // actions 7, 8 (TrailTighten/Loosen): no fill — handled by + // rl_trail_mutate kernel (SP20 P5). The side=2 no-op default is + // correct here. market_targets[b * 2 + 0] = side; market_targets[b * 2 + 1] = size; diff --git a/crates/ml-alpha/cuda/argmax_expected_q.cu b/crates/ml-alpha/cuda/argmax_expected_q.cu index 2011a7893..e934be30c 100644 --- a/crates/ml-alpha/cuda/argmax_expected_q.cu +++ b/crates/ml-alpha/cuda/argmax_expected_q.cu @@ -22,7 +22,7 @@ // Per `feedback_no_atomicadd`: thread 0 writes to `next_actions[b]` // after __syncthreads. -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define Q_N_ATOMS 21 // One block per batch (grid_dim.x = b_size). N_ACTIONS threads per diff --git a/crates/ml-alpha/cuda/bellman_target_projection.cu b/crates/ml-alpha/cuda/bellman_target_projection.cu index 497ada90f..a1ab9c9f7 100644 --- a/crates/ml-alpha/cuda/bellman_target_projection.cu +++ b/crates/ml-alpha/cuda/bellman_target_projection.cu @@ -41,7 +41,7 @@ // block per batch, small block size). #define Q_N_ATOMS 21 -#define N_ACTIONS 9 +#define N_ACTIONS 11 // V_MIN/V_MAX/DELTA_Z are now ISV-driven per audit 2026-05-24 followup // so adaptive reward clamps also lift Q's distributional support. // Slots RL_C51_V_MIN_INDEX/RL_C51_V_MAX_INDEX are ratchet (monotone- diff --git a/crates/ml-alpha/cuda/dqn_distributional_q.cu b/crates/ml-alpha/cuda/dqn_distributional_q.cu index a3673e1ec..1cf5da068 100644 --- a/crates/ml-alpha/cuda/dqn_distributional_q.cu +++ b/crates/ml-alpha/cuda/dqn_distributional_q.cu @@ -50,7 +50,7 @@ // header without diving into the loss kernel body. #define HIDDEN_DIM 128 -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define Q_N_ATOMS 21 // V_MIN / V_MAX / DELTA_Z are kept here as documentation only — the diff --git a/crates/ml-alpha/cuda/log_pi_at_action.cu b/crates/ml-alpha/cuda/log_pi_at_action.cu index 566a17070..4baf79e76 100644 --- a/crates/ml-alpha/cuda/log_pi_at_action.cu +++ b/crates/ml-alpha/cuda/log_pi_at_action.cu @@ -15,7 +15,7 @@ // batches, no atomics. N_ACTIONS = 9 fits in a per-thread sequential // loop comfortably. -#define N_ACTIONS 9 +#define N_ACTIONS 11 // Inputs: // pi_logits [b_size, N_ACTIONS] row-major diff --git a/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu b/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu index e0322d386..a12ce44ef 100644 --- a/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu +++ b/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu @@ -59,7 +59,7 @@ // Backward: same layout. Each thread writes its own logit's gradient. #define HIDDEN_DIM 128 -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define RL_PPO_CLIP_INDEX 402 #define RL_ENTROPY_COEF_INDEX 403 // ISV slot carrying the adaptive importance-ratio clamp ceiling diff --git a/crates/ml-alpha/cuda/rl_action_kernel.cu b/crates/ml-alpha/cuda/rl_action_kernel.cu index 19bfed0ee..b4da974a5 100644 --- a/crates/ml-alpha/cuda/rl_action_kernel.cu +++ b/crates/ml-alpha/cuda/rl_action_kernel.cu @@ -33,7 +33,7 @@ #include -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define Q_N_ATOMS 21 __device__ static uint32_t xorshift32(uint32_t* state) { diff --git a/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu b/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu index 053a87dd0..b161a804c 100644 --- a/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu +++ b/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu @@ -45,7 +45,7 @@ // hit zero (a fully-trained policy may need no exploration pressure). #define RL_ENTROPY_COEF_INDEX 403 -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define COEF_MIN 0.0f #define COEF_MAX 0.05f // ISV-driven entropy-target fraction per `feedback_isv_for_adaptive_bounds`. diff --git a/crates/ml-alpha/cuda/rl_pi_action_kernel.cu b/crates/ml-alpha/cuda/rl_pi_action_kernel.cu index d90bbbf54..7ee649fbc 100644 --- a/crates/ml-alpha/cuda/rl_pi_action_kernel.cu +++ b/crates/ml-alpha/cuda/rl_pi_action_kernel.cu @@ -46,7 +46,7 @@ #include -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define P_MIN 0.02f __device__ static uint32_t xorshift32(uint32_t* state) { diff --git a/crates/ml-alpha/cuda/rl_q_pi_agree_b.cu b/crates/ml-alpha/cuda/rl_q_pi_agree_b.cu index 45aa9924f..1fcd48331 100644 --- a/crates/ml-alpha/cuda/rl_q_pi_agree_b.cu +++ b/crates/ml-alpha/cuda/rl_q_pi_agree_b.cu @@ -31,7 +31,7 @@ #define RL_Q_ARG_VS_PI_AGREE_INDEX 407 #define Q_N_ATOMS 21 -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define AGREE_EMA_ALPHA 0.05f extern "C" __global__ void rl_q_pi_agree_b( diff --git a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu index 55ad7b004..599645d43 100644 --- a/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu +++ b/crates/ml-alpha/cuda/rl_q_pi_distill_grad.cu @@ -34,7 +34,7 @@ #include -#define N_ACTIONS 9 +#define N_ACTIONS 11 #define Q_N_ATOMS 21 #define RL_Q_DISTILL_LAMBDA_INDEX 486 #define RL_Q_DISTILL_TEMPERATURE_INDEX 487 diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index e1207a6b9..4451adcd3 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -38,6 +38,14 @@ use anyhow::{Context, Result}; use clap::Parser; use data::providers::databento::dbn_parser::InstrumentFilter; use ml_alpha::cfc::snap_features::Mbp10RawInput; +// Use the canonical Action enum size, NOT a literal. Caught 2026-05-24 +// during SP20 P4 dogfood: a `0..9` range survived the N_ACTIONS=9→11 +// bump and silently dropped HalfFlat samples (a9, a10 showed 0% in +// diag despite P_MIN=0.02 floor guaranteeing 2% each). Per the +// emerging meta-pattern: any constant that mirrors a kernel-side +// structural dimension MUST reference the Rust const, not duplicate +// the literal value. +use ml_alpha::rl::common::N_ACTIONS; use ml_alpha::data::loader::{ discover_mbp10_files_sorted, MultiHorizonLoader, MultiHorizonLoaderConfig, DEFAULT_OUTCOME_LABEL_COST_ES, @@ -439,7 +447,10 @@ fn main() -> Result<()> { // α = 1/1000 → half-life ≈ 690 steps. Tracks the recent ~1k steps' // action distribution; entropy of normalised windowed_act_hist is // a real exploration signal. - let mut windowed_act_hist: [f32; 9] = [0.0; 9]; + // SP20 P4: action histogram width MUST track the Action enum + // size. Always reference the const (NOT a literal) — see + // import-site comment for the dogfood-caught literal-drift bug. + let mut windowed_act_hist: [f32; N_ACTIONS] = [0.0; N_ACTIONS]; const WINDOWED_ACT_ALPHA: f32 = 1.0 / 1000.0; // Per-step staging buffers for the B×K snapshot tensor that @@ -534,9 +545,9 @@ fn main() -> Result<()> { ) .context("diag: read prev_position_lots_d")?; - let mut act_hist = [0u32; 9]; + let mut act_hist = [0u32; N_ACTIONS]; for &a in &actions_host { - if (0..9).contains(&a) { + if (a as usize) < N_ACTIONS { act_hist[a as usize] += 1; } } @@ -549,7 +560,7 @@ fn main() -> Result<()> { // actual policy exploration variety over the recent run. let total_actions: u32 = act_hist.iter().sum(); if total_actions > 0 { - for i in 0..9 { + for i in 0..N_ACTIONS { let p_step = (act_hist[i] as f32) / (total_actions as f32); windowed_act_hist[i] = (1.0 - WINDOWED_ACT_ALPHA) * windowed_act_hist[i] diff --git a/crates/ml-alpha/src/rl/common.rs b/crates/ml-alpha/src/rl/common.rs index 535afc3a9..fd466de10 100644 --- a/crates/ml-alpha/src/rl/common.rs +++ b/crates/ml-alpha/src/rl/common.rs @@ -7,7 +7,7 @@ use cudarc::driver::CudaSlice; /// systems stay comparable. See `action.rs` (Phase C/D) for the /// enum-style mapping {Short-large, Short-small, Hold, Flat-from-long, /// Flat-from-short, Long-small, Long-large, Trail-tighten, Trail-loosen}. -pub const N_ACTIONS: usize = 9; +pub const N_ACTIONS: usize = 11; /// C51 distributional Q-head atom count. Matches existing /// `pearl_per_branch_c51_atom_span` infrastructure (21 atoms over the @@ -22,21 +22,24 @@ pub const Q_V_MIN: f32 = -1.0; /// C51 atom support `v_max`. See [`Q_V_MIN`]. pub const Q_V_MAX: f32 = 1.0; -/// Discrete 9-action grid identifiers. Matches the layout used by the -/// existing `crates/ml` DQN (`sp5_isv_slots.rs`) so the two systems' -/// policies stay directly comparable — same index → same trade decision. +/// Discrete 11-action grid identifiers. Was 9 pre-SP20 P4; extended +/// with HalfFlatLong/HalfFlatShort to enable per-unit partial-close +/// routing (consumed by SP20 P5 trail-stop check + future P7 +/// pyramid-aware actions_to_market_targets logic). #[derive(Clone, Copy, Debug, PartialEq, Eq)] #[repr(u32)] pub enum Action { - ShortLarge = 0, - ShortSmall = 1, - Hold = 2, - FlatFromLong = 3, - FlatFromShort= 4, - LongSmall = 5, - LongLarge = 6, - TrailTighten = 7, - TrailLoosen = 8, + ShortLarge = 0, + ShortSmall = 1, + Hold = 2, + FlatFromLong = 3, + FlatFromShort = 4, + LongSmall = 5, + LongLarge = 6, + TrailTighten = 7, + TrailLoosen = 8, + HalfFlatLong = 9, // SP20 P4: close ⌈|pos|/2⌉ of long position (min 1 lot) + HalfFlatShort = 10, // SP20 P4: close ⌈|pos|/2⌉ of short position (min 1 lot) } impl Action { @@ -45,16 +48,18 @@ impl Action { /// kernel/loader boundary, this is the type-safety net. pub fn try_from_u32(v: u32) -> Result { match v { - 0 => Ok(Action::ShortLarge), - 1 => Ok(Action::ShortSmall), - 2 => Ok(Action::Hold), - 3 => Ok(Action::FlatFromLong), - 4 => Ok(Action::FlatFromShort), - 5 => Ok(Action::LongSmall), - 6 => Ok(Action::LongLarge), - 7 => Ok(Action::TrailTighten), - 8 => Ok(Action::TrailLoosen), - _ => Err(format!("invalid action {v}; valid range 0..{}", N_ACTIONS)), + 0 => Ok(Action::ShortLarge), + 1 => Ok(Action::ShortSmall), + 2 => Ok(Action::Hold), + 3 => Ok(Action::FlatFromLong), + 4 => Ok(Action::FlatFromShort), + 5 => Ok(Action::LongSmall), + 6 => Ok(Action::LongLarge), + 7 => Ok(Action::TrailTighten), + 8 => Ok(Action::TrailLoosen), + 9 => Ok(Action::HalfFlatLong), + 10 => Ok(Action::HalfFlatShort), + _ => Err(format!("invalid action {v}; valid range 0..{}", N_ACTIONS)), } } } diff --git a/scripts/audit-manifest/actions.txt b/scripts/audit-manifest/actions.txt index 72585f2bb..59d410920 100644 --- a/scripts/audit-manifest/actions.txt +++ b/scripts/audit-manifest/actions.txt @@ -5,3 +5,5 @@ # (rl_trail_mutate + rl_trail_stop_check kernels). TrailTighten TrailLoosen +HalfFlatLong +HalfFlatShort