diff --git a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu b/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu index c72a8d710..8cd6689af 100644 --- a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu +++ b/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu @@ -64,6 +64,7 @@ extern "C" __global__ void rl_fused_reward_pipeline( int* __restrict__ unit_entry_step, // [B*MAX_UNITS] OUT int* __restrict__ unit_lots, // [B*MAX_UNITS] OUT float* __restrict__ unit_initial_r, // [B*MAX_UNITS] OUT + float* __restrict__ unit_peak_unrealized_r, // [B*MAX_UNITS] OUT float* __restrict__ unit_trail_distance, // [B*MAX_UNITS] OUT unsigned char* __restrict__ unit_active, // [B*MAX_UNITS] OUT int* __restrict__ pyramid_units_count, // [B] OUT @@ -142,6 +143,7 @@ extern "C" __global__ void rl_fused_reward_pipeline( unit_entry_step[base + 0] = current_step; unit_lots[base + 0] = current_lots; unit_initial_r[base + 0] = bootstrap_d; + unit_peak_unrealized_r[base + 0] = 0.0f; // sentinel: first obs replaces unit_trail_distance[base + 0] = bootstrap_d; unit_active[base + 0] = 1; pyramid_units_count[b] = 1; @@ -161,6 +163,7 @@ extern "C" __global__ void rl_fused_reward_pipeline( const int delta_lots = abs_curr - abs_prev_u; unit_lots[base + slot] = (current_lots > 0) ? delta_lots : -delta_lots; unit_initial_r[base + slot] = bootstrap_d; + unit_peak_unrealized_r[base + slot] = 0.0f; // sentinel: first obs replaces unit_trail_distance[base + slot] = bootstrap_d; unit_active[base + slot] = 1; pyramid_units_count[b] = count + 1; diff --git a/crates/ml-alpha/cuda/rl_trade_context_update.cu b/crates/ml-alpha/cuda/rl_trade_context_update.cu index 6e242d828..c29a5f4d8 100644 --- a/crates/ml-alpha/cuda/rl_trade_context_update.cu +++ b/crates/ml-alpha/cuda/rl_trade_context_update.cu @@ -5,20 +5,23 @@ // unit (anchor) and current market state: // // [0] time_in_trade_norm = (current_step - oldest_entry_step) / 1000.0 -// [1] drawdown_R = min(0, unrealized_R) — LOSS-ONLY signal +// [1] drawdown_R = unrealized_R - peak(unrealized_R) — symmetric exit signal // [2] pos_magnitude_norm = |position_lots| / 8.0 (MAX_UNITS × max_order) // [3] entry_distance_sigma = (mid - oldest_entry_price) / mean_abs_pnl_ema // -// Feature [1] is asymmetric (loss-only) to prevent the value-leak -// overfitting attractor. If Q sees `unrealized_R > 0` as a state -// feature, it learns "open winning position = high V" → closing -// realizes < V(holding) → never close (observed alpha-rl-wfan-f0 -// step 2000: dones=0 sustained, pnl frozen). With drawdown-only: -// - Losing positions: visible drawdown → agent learns to cut losses -// - Winning positions: invisible to state → close decision driven -// by general policy (not state-conditioned on profit) -// Realized PnL on close still gives Q the gradient to learn "take -// profitable trades" via standard Bellman backup. +// Feature [1] is symmetric drawdown-from-peak (always ≤ 0). Tracks +// the per-unit `peak_unrealized_r` across the trade lifecycle; the +// observed feature is `unrealized_r - peak`, which fires on: +// - Losing positions: monotone drawdown below the entry baseline +// - Winning-then-retracing positions: drawdown from the high-water +// mark even while still net-positive +// This avoids the one-sided clip artifact (12× W/L bias) where a +// LOSS-ONLY signal pushes the agent to exit losers fast and ride +// winners indefinitely. Realized PnL on close still drives Bellman +// learning of "take profitable trades". +// Sentinel-zero bootstrap (per `pearl_first_observation_bootstrap`): +// `peak == 0` on a freshly opened/added/reversed slot; the first +// observation directly replaces the peak with `unrealized_r`. // // When no unit is active (flat): all outputs = 0. // @@ -38,6 +41,7 @@ extern "C" __global__ void rl_trade_context_update( 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] + float* __restrict__ unit_peak_unrealized_r, // [B × MAX_UNITS] IN/OUT const int* __restrict__ unit_lots, // [B × MAX_UNITS] const unsigned char* __restrict__ pos_state, // [B × pos_bytes] const float* __restrict__ bid_px, // [BOOK_LEVELS] @@ -91,10 +95,18 @@ extern "C" __global__ void rl_trade_context_update( ? (mid - entry_px) / mean_abs_pnl : 0.0f; + // Drawdown-from-peak: symmetric exit signal (always ≤ 0). + // Per `pearl_first_observation_bootstrap`: sentinel peak == 0 → + // first observation replaces directly. + const float peak_prev = unit_peak_unrealized_r[base + oldest]; + const float peak_new = (peak_prev == 0.0f) + ? unrealized_r + : fmaxf(peak_prev, unrealized_r); + unit_peak_unrealized_r[base + oldest] = peak_new; + const float drawdown_r = unrealized_r - peak_new; // always ≤ 0 + trade_context[out_base + 0] = time_norm; - // Asymmetric clip: only show LOSSES to the agent. Hides positive - // unrealized to break the value-leak overfit (see header comment). - trade_context[out_base + 1] = fminf(0.0f, unrealized_r); + trade_context[out_base + 1] = drawdown_r; trade_context[out_base + 2] = pos_mag_norm; trade_context[out_base + 3] = entry_dist_sigma; } diff --git a/crates/ml-alpha/cuda/rl_unit_state_update.cu b/crates/ml-alpha/cuda/rl_unit_state_update.cu index 483665816..fc5337208 100644 --- a/crates/ml-alpha/cuda/rl_unit_state_update.cu +++ b/crates/ml-alpha/cuda/rl_unit_state_update.cu @@ -47,6 +47,7 @@ extern "C" __global__ void rl_unit_state_update( int* __restrict__ unit_entry_step, // [B * MAX_UNITS] int* __restrict__ unit_lots, // [B * MAX_UNITS] float* __restrict__ unit_initial_r, // [B * MAX_UNITS] + float* __restrict__ unit_peak_unrealized_r, // [B * MAX_UNITS] float* __restrict__ unit_trail_distance, // [B * MAX_UNITS] unsigned char* __restrict__ unit_active, // [B * MAX_UNITS] int* __restrict__ pyramid_units_count, // [B] @@ -98,6 +99,7 @@ extern "C" __global__ void rl_unit_state_update( unit_entry_step[base + 0] = current_step; unit_lots[base + 0] = curr_pos; unit_initial_r[base + 0] = bootstrap_d; + unit_peak_unrealized_r[base + 0] = 0.0f; // sentinel: first obs replaces unit_trail_distance[base + 0] = bootstrap_d; unit_active[base + 0] = 1; pyramid_units_count[b] = 1; @@ -117,6 +119,7 @@ extern "C" __global__ void rl_unit_state_update( const int delta_lots = abs_curr - abs_prev; unit_lots[base + slot] = (curr_pos > 0) ? delta_lots : -delta_lots; unit_initial_r[base + slot] = bootstrap_d; + unit_peak_unrealized_r[base + slot] = 0.0f; // sentinel: first obs replaces unit_trail_distance[base + slot] = bootstrap_d; unit_active[base + slot] = 1; pyramid_units_count[b] = count + 1; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 73d64934f..c82e2ac88 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -683,6 +683,7 @@ pub struct IntegratedTrainer { pub unit_entry_step_d: CudaSlice, // [B * 4] pub unit_lots_d: CudaSlice, // [B * 4] pub unit_initial_r_d: CudaSlice, // [B * 4] + pub unit_peak_unrealized_r_d: CudaSlice, // [B * 4] — peak(unrealized_R) per active slot pub unit_trail_distance_d: CudaSlice, // [B * 4] pub unit_active_d: CudaSlice, // [B * 4] pub pyramid_units_count_d: CudaSlice, // [B] @@ -1768,6 +1769,9 @@ impl IntegratedTrainer { let unit_initial_r_d = stream .alloc_zeros::(b_size * 4) .context("alloc unit_initial_r_d")?; + let unit_peak_unrealized_r_d = stream + .alloc_zeros::(b_size * 4) + .context("alloc unit_peak_unrealized_r_d")?; let unit_trail_distance_d = stream .alloc_zeros::(b_size * 4) .context("alloc unit_trail_distance_d")?; @@ -2405,6 +2409,7 @@ impl IntegratedTrainer { unit_entry_step_d, unit_lots_d, unit_initial_r_d, + unit_peak_unrealized_r_d, unit_trail_distance_d, unit_active_d, pyramid_units_count_d, @@ -3521,6 +3526,7 @@ impl IntegratedTrainer { unit_entry_step_d: &mut CudaSlice, unit_lots_d: &mut CudaSlice, unit_initial_r_d: &mut CudaSlice, + unit_peak_unrealized_r_d: &mut CudaSlice, unit_trail_distance_d: &mut CudaSlice, unit_active_d: &mut CudaSlice, pyramid_units_count_d: &mut CudaSlice, @@ -3533,6 +3539,7 @@ impl IntegratedTrainer { debug_assert_eq!(unit_entry_step_d.len(), b_size * 4); debug_assert_eq!(unit_lots_d.len(), b_size * 4); debug_assert_eq!(unit_initial_r_d.len(), b_size * 4); + debug_assert_eq!(unit_peak_unrealized_r_d.len(), b_size * 4); debug_assert_eq!(unit_trail_distance_d.len(), b_size * 4); debug_assert_eq!(unit_active_d.len(), b_size * 4); debug_assert_eq!(pyramid_units_count_d.len(), b_size); @@ -3547,6 +3554,7 @@ impl IntegratedTrainer { args.push_ptr(unit_entry_step_d.raw_ptr()); args.push_ptr(unit_lots_d.raw_ptr()); args.push_ptr(unit_initial_r_d.raw_ptr()); + args.push_ptr(unit_peak_unrealized_r_d.raw_ptr()); args.push_ptr(unit_trail_distance_d.raw_ptr()); args.push_ptr(unit_active_d.raw_ptr()); args.push_ptr(pyramid_units_count_d.raw_ptr()); @@ -5934,6 +5942,7 @@ impl IntegratedTrainer { args.push_ptr(self.unit_entry_step_d.raw_ptr()); args.push_ptr(self.unit_lots_d.raw_ptr()); args.push_ptr(self.unit_initial_r_d.raw_ptr()); + args.push_ptr(self.unit_peak_unrealized_r_d.raw_ptr()); args.push_ptr(self.unit_trail_distance_d.raw_ptr()); args.push_ptr(self.unit_active_d.raw_ptr()); args.push_ptr(self.pyramid_units_count_d.raw_ptr()); @@ -6030,6 +6039,7 @@ impl IntegratedTrainer { args.push_ptr(self.unit_entry_price_d.raw_ptr()); args.push_ptr(self.unit_entry_step_d.raw_ptr()); args.push_ptr(self.unit_initial_r_d.raw_ptr()); + args.push_ptr(self.unit_peak_unrealized_r_d.raw_ptr()); args.push_ptr(self.unit_lots_d.raw_ptr()); args.push_ptr(pos_d_ref.raw_ptr()); args.push_ptr(bid_px_d.raw_ptr()); diff --git a/crates/ml-alpha/tests/trade_management_kernels.rs b/crates/ml-alpha/tests/trade_management_kernels.rs index 3561579a7..7c08928a2 100644 --- a/crates/ml-alpha/tests/trade_management_kernels.rs +++ b/crates/ml-alpha/tests/trade_management_kernels.rs @@ -318,6 +318,7 @@ fn unit_state_transitions() -> Result<()> { let mut unit_entry_step_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_lots_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_initial_r_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let mut unit_peak_unrealized_r_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_trail_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_active_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut pyramid_count_d = stream.alloc_zeros::(b_size)?; @@ -332,6 +333,7 @@ fn unit_state_transitions() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d, @@ -377,6 +379,7 @@ fn unit_state_transitions() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d, @@ -655,6 +658,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { let mut unit_entry_step_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_lots_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_initial_r_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; + let mut unit_peak_unrealized_r_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_trail_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut unit_active_d = stream.alloc_zeros::(b_size * MAX_UNITS)?; let mut pyramid_count_d = stream.alloc_zeros::(b_size)?; @@ -669,6 +673,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d, @@ -689,6 +694,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d, @@ -720,6 +726,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d, @@ -746,6 +753,7 @@ fn pyramid_add_populates_next_unit_slot() -> Result<()> { &mut unit_entry_step_d, &mut unit_lots_d, &mut unit_initial_r_d, + &mut unit_peak_unrealized_r_d, &mut unit_trail_d, &mut unit_active_d, &mut pyramid_count_d,