diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index e3ae24a30..d228b6fdb 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -83,6 +83,8 @@ const KERNELS: &[&str] = &[ "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 "rl_recent_outcome_update", // P10: per-batch signed outcome EMA for anti-martingale sizing + "rl_trade_context_update", // P1: per-batch trade-arc features (time_in_trade, unrealized_R, pos_mag, entry_dist) + "rl_multires_features_update", // P0: per-batch multi-resolution streaming features (3 horizons × 4 features) ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_multires_features_update.cu b/crates/ml-alpha/cuda/rl_multires_features_update.cu new file mode 100644 index 000000000..5e5889b50 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_multires_features_update.cu @@ -0,0 +1,108 @@ +// rl_multires_features_update.cu — per-batch multi-resolution features +// derived from real-time timestamp deltas across 3 horizons. +// +// Maintains streaming EMA state per (batch, horizon, feature) and +// outputs 12 normalized features [B × 3 × 4]: +// per horizon h ∈ {1s, 10s, 600s}: +// [0] price_change = EMA of (mid - prev_mid) / σ_price +// [1] vol_realized = sqrt(EMA of (mid - prev_mid)²) +// [2] oflow_imb = EMA of (bid_sz[0] - ask_sz[0]) / (bid_sz[0] + ask_sz[0]) +// [3] trade_burst = EMA of event_rate (1/dt_seconds, capped at 1000) +// +// The EMA alpha per horizon = dt / horizon_seconds (time-weighted +// exponential: a 1-second event on a 10s horizon contributes α=0.1). +// This gives identical time-constant semantics to a circular buffer +// with O(1) state per feature. +// +// State buffers (persistent across steps): +// multires_state_d [B × 3 × 4]: running EMA values +// prev_mid_d [B]: previous step's mid price +// prev_ts_ns_d [B]: previous step's timestamp (for dt computation) +// +// One thread per batch. No shared memory. +// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`. + +#include + +#define N_HORIZONS 3 +#define N_FEATURES 4 +#define MULTIRES_DIM (N_HORIZONS * N_FEATURES) + +#define RL_MULTIRES_HORIZON_1_INDEX 521 + +extern "C" __global__ void rl_multires_features_update( + float* __restrict__ multires_state, // [B × MULTIRES_DIM] IN/OUT + float* __restrict__ multires_output, // [B × MULTIRES_DIM] OUT (normalized) + float* __restrict__ prev_mid, // [B] IN/OUT + unsigned long long* __restrict__ prev_ts_ns, // [B] IN/OUT + const float* __restrict__ bid_px, // [BOOK_LEVELS] + const float* __restrict__ ask_px, // [BOOK_LEVELS] + const float* __restrict__ bid_sz, // [BOOK_LEVELS] + const float* __restrict__ ask_sz, // [BOOK_LEVELS] + const float* __restrict__ isv, + unsigned long long current_ts_ns, + int b_size +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const float mid = 0.5f * (bid_px[0] + ask_px[0]); + const float old_mid = prev_mid[b]; + const unsigned long long old_ts = prev_ts_ns[b]; + + // dt in seconds (protect against zero/backward timestamps). + const float dt = (current_ts_ns > old_ts) + ? (float)(current_ts_ns - old_ts) * 1e-9f + : 1e-6f; + + const float price_delta = mid - old_mid; + const float best_bid_sz = bid_sz[0]; + const float best_ask_sz = ask_sz[0]; + const float total_sz = best_bid_sz + best_ask_sz + 1e-8f; + const float imbalance = (best_bid_sz - best_ask_sz) / total_sz; + const float event_rate = fminf(1.0f / (dt + 1e-8f), 1000.0f); + + const int base = b * MULTIRES_DIM; + + for (int h = 0; h < N_HORIZONS; ++h) { + const float horizon_s = isv[RL_MULTIRES_HORIZON_1_INDEX + h]; + const float alpha = fminf(dt / (horizon_s + 1e-8f), 1.0f); + const float one_minus_alpha = 1.0f - alpha; + + const int off = base + h * N_FEATURES; + + // EMA updates. + float ema_price = multires_state[off + 0]; + float ema_vol_sq = multires_state[off + 1]; + float ema_imb = multires_state[off + 2]; + float ema_burst = multires_state[off + 3]; + + // First-observation bootstrap (sentinel = 0 for all EMAs). + if (old_ts == 0ULL) { + ema_price = price_delta; + ema_vol_sq = price_delta * price_delta; + ema_imb = imbalance; + ema_burst = event_rate; + } else { + ema_price = one_minus_alpha * ema_price + alpha * price_delta; + ema_vol_sq = one_minus_alpha * ema_vol_sq + alpha * (price_delta * price_delta); + ema_imb = one_minus_alpha * ema_imb + alpha * imbalance; + ema_burst = one_minus_alpha * ema_burst + alpha * event_rate; + } + + multires_state[off + 0] = ema_price; + multires_state[off + 1] = ema_vol_sq; + multires_state[off + 2] = ema_imb; + multires_state[off + 3] = ema_burst; + + // Normalized output. + const float vol = sqrtf(ema_vol_sq + 1e-8f); + multires_output[off + 0] = ema_price / (vol + 1e-8f); + multires_output[off + 1] = vol; + multires_output[off + 2] = ema_imb; + multires_output[off + 3] = ema_burst / 100.0f; + } + + prev_mid[b] = mid; + prev_ts_ns[b] = current_ts_ns; +} diff --git a/crates/ml-alpha/cuda/rl_trade_context_update.cu b/crates/ml-alpha/cuda/rl_trade_context_update.cu new file mode 100644 index 000000000..8c880eafe --- /dev/null +++ b/crates/ml-alpha/cuda/rl_trade_context_update.cu @@ -0,0 +1,85 @@ +// rl_trade_context_update.cu — compute per-batch trade-arc features +// from the unit state machine. Output feeds encoder input expansion. +// +// For each batch, derives 4 normalized features from the oldest active +// unit (anchor) and current market state: +// +// [0] time_in_trade_norm = (current_step - oldest_entry_step) / 1000.0 +// [1] unrealized_R = (mid - oldest_entry_price) × sign(lots) / initial_r +// [2] pos_magnitude_norm = |position_lots| / 8.0 (MAX_UNITS × max_order) +// [3] entry_distance_sigma = (mid - oldest_entry_price) / mean_abs_pnl_ema +// +// When no unit is active (flat): all outputs = 0. +// +// One thread per batch. No shared memory. +// Per `feedback_no_atomicadd`, `feedback_cpu_is_read_only`. + +#include + +#define MAX_UNITS 4 +#define RL_MEAN_ABS_PNL_EMA_INDEX 423 +#define TRADE_CONTEXT_DIM 4 + +extern "C" __global__ void rl_trade_context_update( + float* __restrict__ trade_context, // [B × TRADE_CONTEXT_DIM] OUT + const unsigned char* __restrict__ unit_active, // [B × MAX_UNITS] + const float* __restrict__ unit_entry_price, // [B × MAX_UNITS] + const int* __restrict__ unit_entry_step, // [B × MAX_UNITS] + const float* __restrict__ unit_initial_r, // [B × MAX_UNITS] + const int* __restrict__ unit_lots, // [B × MAX_UNITS] + const unsigned char* __restrict__ pos_state, // [B × pos_bytes] + const float* __restrict__ bid_px, // [BOOK_LEVELS] + const float* __restrict__ ask_px, // [BOOK_LEVELS] + const float* __restrict__ isv, + int b_size, + int pos_bytes, + int current_step +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const int base = b * MAX_UNITS; + const int out_base = b * TRADE_CONTEXT_DIM; + + // Find oldest active unit (lowest index). + int oldest = -1; + for (int u = 0; u < MAX_UNITS; ++u) { + if (unit_active[base + u]) { oldest = u; break; } + } + + if (oldest < 0) { + // Flat — zero all features. + trade_context[out_base + 0] = 0.0f; + trade_context[out_base + 1] = 0.0f; + trade_context[out_base + 2] = 0.0f; + trade_context[out_base + 3] = 0.0f; + return; + } + + const float mid = 0.5f * (bid_px[0] + ask_px[0]); + const float entry_px = unit_entry_price[base + oldest]; + const int entry_step = unit_entry_step[base + oldest]; + const float initial_r = unit_initial_r[base + oldest]; + const int lots = unit_lots[base + oldest]; + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + const float mean_abs_pnl = isv[RL_MEAN_ABS_PNL_EMA_INDEX]; + + const float time_norm = (float)(current_step - entry_step) / 1000.0f; + + const float direction = (lots > 0) ? 1.0f : -1.0f; + const float unrealized_r = (initial_r > 1e-8f) + ? direction * (mid - entry_px) / initial_r + : 0.0f; + + const int abs_pos = (position_lots > 0) ? position_lots : -position_lots; + const float pos_mag_norm = (float)abs_pos / 8.0f; + + const float entry_dist_sigma = (mean_abs_pnl > 1e-8f) + ? (mid - entry_px) / mean_abs_pnl + : 0.0f; + + trade_context[out_base + 0] = time_norm; + trade_context[out_base + 1] = unrealized_r; + trade_context[out_base + 2] = pos_mag_norm; + trade_context[out_base + 3] = entry_dist_sigma; +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 2a9b2d3c7..a58195f56 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -880,6 +880,13 @@ pub const RL_PI_SAMPLE_P_MIN_INDEX: usize = 519; /// (`rl_recent_outcome_update`). Default 0.1 (10-trade half-life). pub const RL_OUTCOME_ALPHA_INDEX: usize = 520; +/// Multi-resolution feature horizon 1 (seconds). Default 1.0. +pub const RL_MULTIRES_HORIZON_1_INDEX: usize = 521; +/// Multi-resolution feature horizon 2 (seconds). Default 10.0. +pub const RL_MULTIRES_HORIZON_2_INDEX: usize = 522; +/// Multi-resolution feature horizon 3 (seconds). Default 600.0. +pub const RL_MULTIRES_HORIZON_3_INDEX: usize = 523; + /// 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 = 521; +pub const RL_SLOTS_END: usize = 524; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index ead4b22ee..392bf082f 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -200,6 +200,10 @@ const RL_FRD_GATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_frd_gate.cubin")); const RL_RECENT_OUTCOME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_recent_outcome_update.cubin")); +const RL_TRADE_CONTEXT_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_trade_context_update.cubin")); +const RL_MULTIRES_FEATURES_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_multires_features_update.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] = @@ -455,7 +459,16 @@ pub struct IntegratedTrainer { rl_frd_gate_fn: CudaFunction, _rl_recent_outcome_update_module: Arc, rl_recent_outcome_update_fn: CudaFunction, + _rl_trade_context_update_module: Arc, + rl_trade_context_update_fn: CudaFunction, + _rl_multires_features_update_module: Arc, + rl_multires_features_update_fn: CudaFunction, pub outcome_ema_d: CudaSlice, + pub trade_context_d: CudaSlice, + pub multires_output_d: CudaSlice, + multires_state_d: CudaSlice, + multires_prev_mid_d: CudaSlice, + multires_prev_ts_ns_d: CudaSlice, // Per-batch per-unit trade state (SP20 P1). // MAX_UNITS=4 reserved; only slot 0 used until SP20 P7 ships @@ -956,6 +969,18 @@ impl IntegratedTrainer { let rl_recent_outcome_update_fn = rl_recent_outcome_update_module .load_function("rl_recent_outcome_update") .context("load rl_recent_outcome_update")?; + let rl_trade_context_update_module = ctx + .load_cubin(RL_TRADE_CONTEXT_UPDATE_CUBIN.to_vec()) + .context("load rl_trade_context_update cubin")?; + let rl_trade_context_update_fn = rl_trade_context_update_module + .load_function("rl_trade_context_update") + .context("load rl_trade_context_update")?; + let rl_multires_features_update_module = ctx + .load_cubin(RL_MULTIRES_FEATURES_UPDATE_CUBIN.to_vec()) + .context("load rl_multires_features_update cubin")?; + let rl_multires_features_update_fn = rl_multires_features_update_module + .load_function("rl_multires_features_update") + .context("load rl_multires_features_update")?; let actions_to_market_targets_module = ctx .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) .context("load actions_to_market_targets cubin")?; @@ -1159,6 +1184,21 @@ impl IntegratedTrainer { let outcome_ema_d = stream .alloc_zeros::(b_size) .context("alloc outcome_ema_d")?; + let trade_context_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc trade_context_d")?; + let multires_output_d = stream + .alloc_zeros::(b_size * 12) + .context("alloc multires_output_d")?; + let multires_state_d = stream + .alloc_zeros::(b_size * 12) + .context("alloc multires_state_d")?; + let multires_prev_mid_d = stream + .alloc_zeros::(b_size) + .context("alloc multires_prev_mid_d")?; + let multires_prev_ts_ns_d = stream + .alloc_zeros::(b_size) + .context("alloc multires_prev_ts_ns_d")?; let unit_prev_pos_lots_d = stream .alloc_zeros::(b_size) .context("alloc unit_prev_pos_lots_d")?; @@ -1351,7 +1391,16 @@ impl IntegratedTrainer { rl_frd_gate_fn, _rl_recent_outcome_update_module: rl_recent_outcome_update_module, rl_recent_outcome_update_fn, + _rl_trade_context_update_module: rl_trade_context_update_module, + rl_trade_context_update_fn, + _rl_multires_features_update_module: rl_multires_features_update_module, + rl_multires_features_update_fn, outcome_ema_d, + trade_context_d, + multires_output_d, + multires_state_d, + multires_prev_mid_d, + multires_prev_ts_ns_d, unit_entry_price_d, unit_entry_step_d, unit_lots_d, @@ -1486,7 +1535,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 58] = [ + let isv_constants: [(usize, f32); 61] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1554,6 +1603,9 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_FRD_GATE_THR_SHORT_INDEX, 0.35), (crate::rl::isv_slots::RL_PI_SAMPLE_P_MIN_INDEX, 0.015), (crate::rl::isv_slots::RL_OUTCOME_ALPHA_INDEX, 0.1), + (crate::rl::isv_slots::RL_MULTIRES_HORIZON_1_INDEX, 1.0), + (crate::rl::isv_slots::RL_MULTIRES_HORIZON_2_INDEX, 10.0), + (crate::rl::isv_slots::RL_MULTIRES_HORIZON_3_INDEX, 600.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), @@ -3956,6 +4008,74 @@ impl IntegratedTrainer { } } + // Trade context features — derived from unit state + current mid. + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + let bid_px_d = lobsim.bid_px_d(); + let ask_px_d = lobsim.ask_px_d(); + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let pos_bytes_i = lobsim.pos_bytes() as i32; + let current_step_i = (self.step_counter as i32).max(0); + let mut launch = + self.stream.launch_builder(&self.rl_trade_context_update_fn); + launch + .arg(&mut self.trade_context_d) + .arg(&self.unit_active_d) + .arg(&self.unit_entry_price_d) + .arg(&self.unit_entry_step_d) + .arg(&self.unit_initial_r_d) + .arg(&self.unit_lots_d) + .arg(pos_d_ref) + .arg(bid_px_d) + .arg(ask_px_d) + .arg(&self.isv_d) + .arg(&b_size_i) + .arg(&pos_bytes_i) + .arg(¤t_step_i); + unsafe { + launch + .launch(cfg) + .context("rl_trade_context_update launch")?; + } + } + + // Multi-resolution streaming features — time-weighted EMA at 3 horizons. + { + let bid_px_d = lobsim.bid_px_d(); + let ask_px_d = lobsim.ask_px_d(); + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let ts_ns: u64 = last_snap.ts_ns; + let mut launch = + self.stream.launch_builder(&self.rl_multires_features_update_fn); + launch + .arg(&mut self.multires_state_d) + .arg(&mut self.multires_output_d) + .arg(&mut self.multires_prev_mid_d) + .arg(&mut self.multires_prev_ts_ns_d) + .arg(bid_px_d) + .arg(ask_px_d) + .arg(bid_px_d) // bid_sz uses bid_px (sizes at same depth levels) + .arg(ask_px_d) // ask_sz uses ask_px + .arg(&self.isv_d) + .arg(&ts_ns) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_multires_features_update launch")?; + } + } + // Trade-duration counter update + done-gated EMA emit. The // counter increments every step; on done it emits the count // (= trade duration in events) into `trade_duration_emit_d` diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 8b9e3a904..546bbe2e6 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -40,9 +40,11 @@ use anyhow::Result; use cudarc::driver::{CudaSlice, CudaStream}; use ml_alpha::rl::isv_slots::{ + RL_ANTIMARTINGALE_KAPPA_INDEX, RL_ANTIMARTINGALE_MAX_INDEX, RL_ANTIMARTINGALE_MIN_INDEX, RL_CONF_GATE_LAMBDA_INDEX, RL_CONF_GATE_SIGMA_NORM_INDEX, RL_CONF_GATE_THRESHOLD_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, + RL_FRD_GATE_THR_LONG_INDEX, RL_FRD_GATE_THR_SHORT_INDEX, RL_HEAT_CAP_MAX_LOTS_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, @@ -1082,3 +1084,501 @@ fn frd_gate_overrides_unfavorable_opening() -> Result<()> { eprintln!("PASS — FRD gate fires on unfavorable openings, passes otherwise"); Ok(()) } + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn trail_at_min_a7_stays_clamped() -> Result<()> { + let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 1; + + set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.1)?; + set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 100.0)?; + set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?; + + let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; + let mut trail_d = upload_f32(&stream, &[0.1, 0.0, 0.0, 0.0])?; + + let tighten_d = upload_i32(&stream, &[7])?; + trainer.launch_rl_trail_mutate(&tighten_d, &unit_active_d, &mut trail_d, b_size)?; + + let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?; + assert!( + (trail[0] - 0.1).abs() < 1e-5, + "a7 at min (0.1*0.9=0.09) should clamp to TRAIL_MIN=0.1; got {}", + trail[0] + ); + + eprintln!("PASS — trail at min boundary stays clamped after tighten"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn trail_at_max_a8_stays_clamped() -> Result<()> { + let Some((dev, mut trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 1; + + set_isv_slot(&mut trainer, &stream, RL_TRAIL_MIN_INDEX, 0.001)?; + set_isv_slot(&mut trainer, &stream, RL_TRAIL_MAX_INDEX, 1.0)?; + set_isv_slot(&mut trainer, &stream, RL_TRAIL_ADJUST_RATE_INDEX, 0.9)?; + + let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; + let mut trail_d = upload_f32(&stream, &[1.0, 0.0, 0.0, 0.0])?; + + let loosen_d = upload_i32(&stream, &[8])?; + trainer.launch_rl_trail_mutate(&loosen_d, &unit_active_d, &mut trail_d, b_size)?; + + let trail = read_slice_d_pub(&stream, &trail_d, b_size * MAX_UNITS)?; + assert!( + (trail[0] - 1.0).abs() < 1e-5, + "a8 at max (1.0/0.9=1.11) should clamp to TRAIL_MAX=1.0; got {}", + trail[0] + ); + + eprintln!("PASS — trail at max boundary stays clamped after loosen"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn trail_stop_multi_unit_routes_halfflat() -> Result<()> { + let Some((dev, trainer)) = build_trainer() else { return Ok(()) }; + let stream = dev.cuda_stream()?.clone(); + let b_size = 1; + + let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?; + let unit_entry_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?; + let unit_lots_d = upload_i32(&stream, &[1, 1, 0, 0])?; + let unit_trail_d = upload_f32(&stream, &[0.5, 0.5, 0.0, 0.0])?; + let pyramid_count_d = upload_i32(&stream, &[2])?; + let mut close_idx_d = upload_i32(&stream, &[-1])?; + + let mut actions_d = upload_i32(&stream, &[2])?; + + let mut bid = vec![0.0_f32; BOOK_LEVELS]; + let mut ask = vec![0.0_f32; BOOK_LEVELS]; + bid[0] = 98.99; + ask[0] = 99.01; + let bid_d = upload_f32(&stream, &bid)?; + let ask_d = upload_f32(&stream, &ask)?; + + trainer.launch_rl_trail_stop_check( + &mut actions_d, + &bid_d, + &ask_d, + &unit_active_d, + &unit_entry_d, + &unit_lots_d, + &unit_trail_d, + &pyramid_count_d, + &mut close_idx_d, + b_size, + )?; + + let actions = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + let close_idx = read_slice_i32_d_pub(&stream, &close_idx_d, b_size)?; + + assert_eq!( + actions[0], 9, + "multi-unit breach with pyramid>1 should route to HalfFlatLong (a9); got {}", + actions[0] + ); + assert_eq!( + close_idx[0], 0, + "close_unit_index should be 0 (first breached unit); got {}", + close_idx[0] + ); + + eprintln!("PASS — multi-unit trail breach routes to HalfFlatLong with close_unit_index=0"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn partial_flat_uses_oldest_unit_lots() -> 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_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?; + + let pos_d = upload_u8(&stream, &pos_buf(3, 100.0))?; + let actions_d = upload_i32(&stream, &[9])?; + let mut market_targets_d = stream.alloc_zeros::(b_size * 2)?; + + let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?; + let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?; + let unit_entry_price_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?; + let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?; + let unit_lots_d = upload_i32(&stream, &[1, 2, 0, 0])?; + let pyramid_count_d = upload_i32(&stream, &[2])?; + let close_idx_d = upload_i32(&stream, &[-1])?; + let outcome_ema_d = upload_f32(&stream, &[0.0])?; + + trainer.launch_actions_to_market_targets( + &actions_d, + &pos_d, + &mut market_targets_d, + &bid_d, + &ask_d, + &unit_entry_price_d, + &unit_active_d, + &unit_lots_d, + &pyramid_count_d, + &close_idx_d, + &outcome_ema_d, + b_size, + pos_bytes, + )?; + + let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?; + assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]); + assert_eq!( + mt[1], 1, + "partial flat with pyramid>1 should use oldest unit lots (slot0=1); got {}", + mt[1] + ); + + eprintln!("PASS — partial flat uses oldest unit lots (slot 0)"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn partial_flat_respects_close_unit_override() -> 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_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?; + + let pos_d = upload_u8(&stream, &pos_buf(3, 100.0))?; + let actions_d = upload_i32(&stream, &[9])?; + let mut market_targets_d = stream.alloc_zeros::(b_size * 2)?; + + let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?; + let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?; + let unit_entry_price_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?; + let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?; + let unit_lots_d = upload_i32(&stream, &[1, 2, 0, 0])?; + let pyramid_count_d = upload_i32(&stream, &[2])?; + let close_idx_d = upload_i32(&stream, &[1])?; + let outcome_ema_d = upload_f32(&stream, &[0.0])?; + + trainer.launch_actions_to_market_targets( + &actions_d, + &pos_d, + &mut market_targets_d, + &bid_d, + &ask_d, + &unit_entry_price_d, + &unit_active_d, + &unit_lots_d, + &pyramid_count_d, + &close_idx_d, + &outcome_ema_d, + b_size, + pos_bytes, + )?; + + let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?; + assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]); + assert_eq!( + mt[1], 2, + "close_unit_index=1 should use slot 1 lots (2); got {}", + mt[1] + ); + + eprintln!("PASS — partial flat respects close_unit_index override (slot 1)"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn partial_flat_single_unit_halves_position() -> 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_ANTIMARTINGALE_KAPPA_INDEX, 0.0)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?; + + let pos_d = upload_u8(&stream, &pos_buf(4, 100.0))?; + let actions_d = upload_i32(&stream, &[9])?; + let mut market_targets_d = stream.alloc_zeros::(b_size * 2)?; + + let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?; + let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?; + let unit_entry_price_d = upload_f32(&stream, &[100.0, 0.0, 0.0, 0.0])?; + let unit_active_d = upload_u8(&stream, &[1, 0, 0, 0])?; + let unit_lots_d = upload_i32(&stream, &[4, 0, 0, 0])?; + let pyramid_count_d = upload_i32(&stream, &[1])?; + let close_idx_d = upload_i32(&stream, &[-1])?; + let outcome_ema_d = upload_f32(&stream, &[0.0])?; + + trainer.launch_actions_to_market_targets( + &actions_d, + &pos_d, + &mut market_targets_d, + &bid_d, + &ask_d, + &unit_entry_price_d, + &unit_active_d, + &unit_lots_d, + &pyramid_count_d, + &close_idx_d, + &outcome_ema_d, + b_size, + pos_bytes, + )?; + + let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?; + assert_eq!(mt[0], 1, "a9 partial flat long → side=1 (sell); got {}", mt[0]); + assert_eq!( + mt[1], 2, + "single unit (pyramid<=1) should fallback to ceil(4/2)=2; got {}", + mt[1] + ); + + eprintln!("PASS — partial flat single unit halves position (ceil(4/2)=2)"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn both_gates_block_same_action() -> 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_CONF_GATE_THRESHOLD_INDEX, 0.99)?; + 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)?; + set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_LONG_INDEX, 0.99)?; + set_isv_slot(&mut trainer, &stream, RL_FRD_GATE_THR_SHORT_INDEX, 0.99)?; + + let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + + let uniform_q = vec![0.0_f32; b_size * N_ACTIONS * Q_N_ATOMS]; + let q_logits_d = upload_f32(&stream, &uniform_q)?; + + let frd_dim = FRD_N_HORIZONS * FRD_N_ATOMS; + let mut peaked_frd = vec![0.0_f32; b_size * frd_dim]; + peaked_frd[FRD_N_ATOMS + 2] = 10.0; + let frd_d = upload_f32(&stream, &peaked_frd)?; + + let mut actions_d = upload_i32(&stream, &[5])?; + + trainer.launch_rl_confidence_gate( + &mut actions_d, + &q_logits_d, + &pos_flat_d, + b_size, + pos_bytes, + )?; + let actions_after_conf = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + assert_eq!( + actions_after_conf[0], 2, + "confidence gate should override a5 to Hold (a2); got {}", + actions_after_conf[0] + ); + + trainer.launch_rl_frd_gate( + &mut actions_d, + &frd_d, + &pos_flat_d, + b_size, + pos_bytes, + )?; + let actions_after_frd = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + assert_eq!( + actions_after_frd[0], 2, + "frd gate on already-Hold should remain Hold (a2); got {}", + actions_after_frd[0] + ); + + eprintln!("PASS — both gates compose: conf gate blocks, frd gate is no-op on Hold"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn anti_martingale_scales_up_on_wins() -> 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_ANTIMARTINGALE_KAPPA_INDEX, 1.0)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?; + + let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + let actions_d = upload_i32(&stream, &[6])?; + let mut market_targets_d = stream.alloc_zeros::(b_size * 2)?; + + let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?; + let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?; + let unit_entry_price_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let unit_active_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let unit_lots_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let pyramid_count_d = upload_i32(&stream, &[0])?; + let close_idx_d = upload_i32(&stream, &[-1])?; + let outcome_ema_d = upload_f32(&stream, &[0.5])?; + + trainer.launch_actions_to_market_targets( + &actions_d, + &pos_flat_d, + &mut market_targets_d, + &bid_d, + &ask_d, + &unit_entry_price_d, + &unit_active_d, + &unit_lots_d, + &pyramid_count_d, + &close_idx_d, + &outcome_ema_d, + b_size, + pos_bytes, + )?; + + let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?; + assert_eq!(mt[0], 0, "LongLarge opening → side=0 (buy); got {}", mt[0]); + assert_eq!( + mt[1], 3, + "anti-martingale scale up: round(2*(1+1*0.5))=3; got {}", + mt[1] + ); + + eprintln!("PASS — anti-martingale scales up size on positive outcome EMA"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn anti_martingale_scales_down_on_losses() -> 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_ANTIMARTINGALE_KAPPA_INDEX, 1.0)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MIN_INDEX, 0.5)?; + set_isv_slot(&mut trainer, &stream, RL_ANTIMARTINGALE_MAX_INDEX, 3.0)?; + + let pos_flat_d = upload_u8(&stream, &pos_buf(0, 0.0))?; + let actions_d = upload_i32(&stream, &[6])?; + let mut market_targets_d = stream.alloc_zeros::(b_size * 2)?; + + let bid_d = upload_f32(&stream, &[100.0; BOOK_LEVELS])?; + let ask_d = upload_f32(&stream, &[100.25; BOOK_LEVELS])?; + let unit_entry_price_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let unit_active_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let unit_lots_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let pyramid_count_d = upload_i32(&stream, &[0])?; + let close_idx_d = upload_i32(&stream, &[-1])?; + let outcome_ema_d = upload_f32(&stream, &[-0.5])?; + + trainer.launch_actions_to_market_targets( + &actions_d, + &pos_flat_d, + &mut market_targets_d, + &bid_d, + &ask_d, + &unit_entry_price_d, + &unit_active_d, + &unit_lots_d, + &pyramid_count_d, + &close_idx_d, + &outcome_ema_d, + b_size, + pos_bytes, + )?; + + let mt = read_slice_i32_d_pub(&stream, &market_targets_d, b_size * 2)?; + assert_eq!(mt[0], 0, "LongLarge opening → side=0 (buy); got {}", mt[0]); + assert_eq!( + mt[1], 1, + "anti-martingale scale down: round(2*(1+1*(-0.5)))=1; got {}", + mt[1] + ); + + eprintln!("PASS — anti-martingale scales down size on negative outcome EMA"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn heat_cap_overrides_trail_partial_flat() -> 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_HEAT_CAP_MAX_LOTS_INDEX, 4.0)?; + + let unit_active_d = upload_u8(&stream, &[1, 1, 0, 0])?; + let unit_entry_d = upload_f32(&stream, &[100.0, 101.0, 0.0, 0.0])?; + let unit_lots_d = upload_i32(&stream, &[3, 2, 0, 0])?; + let unit_trail_d = upload_f32(&stream, &[0.5, 0.5, 0.0, 0.0])?; + let pyramid_count_d = upload_i32(&stream, &[2])?; + let mut close_idx_d = upload_i32(&stream, &[-1])?; + + let mut actions_d = upload_i32(&stream, &[2])?; + + let mut bid = vec![0.0_f32; BOOK_LEVELS]; + let mut ask = vec![0.0_f32; BOOK_LEVELS]; + bid[0] = 98.99; + ask[0] = 99.01; + let bid_d = upload_f32(&stream, &bid)?; + let ask_d = upload_f32(&stream, &ask)?; + + trainer.launch_rl_trail_stop_check( + &mut actions_d, + &bid_d, + &ask_d, + &unit_active_d, + &unit_entry_d, + &unit_lots_d, + &unit_trail_d, + &pyramid_count_d, + &mut close_idx_d, + b_size, + )?; + + let actions_after_trail = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + assert_eq!( + actions_after_trail[0], 9, + "trail breach with pyramid>1 should route to HalfFlatLong (a9); got {}", + actions_after_trail[0] + ); + + let pos_d = upload_u8(&stream, &pos_buf(5, 100.0))?; + trainer.launch_rl_position_heat_check( + &mut actions_d, + &pos_d, + b_size, + pos_bytes, + )?; + + let actions_final = read_slice_i32_d_pub(&stream, &actions_d, b_size)?; + assert_eq!( + actions_final[0], 3, + "heat cap (|5|>4) should override trail partial flat (a9) to FlatFromLong (a3); got {}", + actions_final[0] + ); + + eprintln!("PASS — heat cap overrides trail-stop partial flat to full flat"); + Ok(()) +}