gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(rl): drawdown-from-peak — symmetric exit signal (replaces clip artifact)

Browse Source 
Yesterday's fix `trade_context[1] = fminf(0.0f, unrealized_R)` created a
structural 12× W/L bias: the agent saw losses fast (closed them), couldn't
see wins (rode blindly until max_hold/session). The reproducibility across
folds (f0=f1=f2 at step 500) was the SAME exploit pattern in each window,
not real edge.

This replaces the asymmetric clip with drawdown-from-peak:

  trade_context[1] = unrealized_R - peak(unrealized_R)  (always ≤ 0)

The peak tracks per-active-unit and resets on slot activation. Symmetric:
fires on losing positions (monotone drawdown from entry) AND on
winning-then-retracing positions (drawdown from high-water mark even
while still net-positive).

Sentinel-zero bootstrap per pearl_first_observation_bootstrap: peak==0
on freshly opened/added/reversed slot; first observation directly
replaces peak with unrealized_R.

Implementation:
- New CudaSlice unit_peak_unrealized_r_d [B×MAX_UNITS], allocated zero
- rl_unit_state_update + rl_fused_reward_pipeline write 0.0f sentinel
  on each OPEN/REVERSE/PYRAMID-ADD activation
- rl_trade_context_update reads/updates peak inline, outputs drawdown

Local smoke (RTX 3050 Ti, b=128, 500 steps):
- avg_hold rose from 17.6 → 26-32 steps (wave-scale-ier)
- W/L magnitude ratio is no longer pinned at 12× — varies 0.67-27 across
  the run, sometimes L>W (genuinely symmetric signal)
- No crash, training trajectory healthy

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-28 11:59:17 +02:00
parent 78b781a523
commit 4df0867622
5 changed files with 50 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu
View File

@@ -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;

40
crates/ml-alpha/cuda/rl_trade_context_update.cu
View File

@@ -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;
}

3
crates/ml-alpha/cuda/rl_unit_state_update.cu
View File

@@ -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;

10
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -683,6 +683,7 @@ pub struct IntegratedTrainer {
pub unit_entry_step_d: CudaSlice<i32>, // [B * 4]
pub unit_lots_d: CudaSlice<i32>, // [B * 4]
pub unit_initial_r_d: CudaSlice<f32>, // [B * 4]
pub unit_peak_unrealized_r_d: CudaSlice<f32>, // [B * 4] — peak(unrealized_R) per active slot
pub unit_trail_distance_d: CudaSlice<f32>, // [B * 4]
pub unit_active_d: CudaSlice<u8>, // [B * 4]
pub pyramid_units_count_d: CudaSlice<i32>, // [B]
@@ -1768,6 +1769,9 @@ impl IntegratedTrainer {
let unit_initial_r_d = stream
.alloc_zeros::<f32>(b_size * 4)
.context("alloc unit_initial_r_d")?;
let unit_peak_unrealized_r_d = stream
.alloc_zeros::<f32>(b_size * 4)
.context("alloc unit_peak_unrealized_r_d")?;
let unit_trail_distance_d = stream
.alloc_zeros::<f32>(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<i32>,
unit_lots_d: &mut CudaSlice<i32>,
unit_initial_r_d: &mut CudaSlice<f32>,
unit_peak_unrealized_r_d: &mut CudaSlice<f32>,
unit_trail_distance_d: &mut CudaSlice<f32>,
unit_active_d: &mut CudaSlice<u8>,
pyramid_units_count_d: &mut CudaSlice<i32>,
@@ -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());

8
crates/ml-alpha/tests/trade_management_kernels.rs
View File

@@ -318,6 +318,7 @@ fn unit_state_transitions() -> Result<()> {
let mut unit_entry_step_d = stream.alloc_zeros::<i32>(b_size * MAX_UNITS)?;
let mut unit_lots_d = stream.alloc_zeros::<i32>(b_size * MAX_UNITS)?;
let mut unit_initial_r_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_peak_unrealized_r_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_trail_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_active_d = stream.alloc_zeros::<u8>(b_size * MAX_UNITS)?;
let mut pyramid_count_d = stream.alloc_zeros::<i32>(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::<i32>(b_size * MAX_UNITS)?;
let mut unit_lots_d = stream.alloc_zeros::<i32>(b_size * MAX_UNITS)?;
let mut unit_initial_r_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_peak_unrealized_r_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_trail_d = stream.alloc_zeros::<f32>(b_size * MAX_UNITS)?;
let mut unit_active_d = stream.alloc_zeros::<u8>(b_size * MAX_UNITS)?;
let mut pyramid_count_d = stream.alloc_zeros::<i32>(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,