From 050407dbd5ad671da0da82385eb0387a64db1b37 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 23 Mar 2026 10:13:27 +0100 Subject: [PATCH] =?UTF-8?q?fix:=203=20critical=20findings=20=E2=80=94=20st?= =?UTF-8?q?ate=20blindness,=20cold-start,=20reward=20scaling?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Finding 1: STATE BLINDNESS — Q-network couldn't see risk Portfolio features expanded from 3 (useless: pv/pv≈1, cash/pv≈1) to 8: position, unrealized_pnl/equity, drawdown, hold_time/100, realized_pnl/equity, distance_to_floor, trade_return, cash_ratio. PORTFOLIO_DIM 3→8 across all NVRTC injection sites. State dim: 48→56 (without OFI), 56→64 (with OFI). Finding 2: Q-GAP COLD-START — model couldn't trade to learn Q-gap threshold ramps linearly from 0.0 to target over first 10 epochs. At epoch 0, all Q-values ≈ 0 → Q-gap < threshold → no trades → no signal. Added current_epoch field to DQNTrainer. Finding 3: TRADE REWARD SCALING — idle penalty dominated trade signal trade_return scaling: ×100 → ×1000 (1 ES tick = 0.36 reward, was 0.036). idle penalty max: 0.05 → 0.01 (was competing with 1-tick trade reward). Now: trade reward (0.36) >> idle penalty (0.01). Correct incentive. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 75 +++++++++++++++---- .../src/cuda_pipeline/gpu_action_selector.rs | 4 +- .../cuda_pipeline/gpu_backtest_evaluator.rs | 6 +- .../cuda_pipeline/gpu_curiosity_trainer.rs | 2 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 2 +- .../cuda_pipeline/gpu_experience_collector.rs | 2 +- .../ml/src/cuda_pipeline/gpu_ppo_collector.rs | 2 +- crates/ml/src/hyperopt/adapters/dqn.rs | 2 +- .../src/trainers/dqn/trainer/constructor.rs | 1 + crates/ml/src/trainers/dqn/trainer/mod.rs | 2 + .../src/trainers/dqn/trainer/training_loop.rs | 11 ++- 11 files changed, 82 insertions(+), 27 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 3e211e91b..062ba4586 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -167,31 +167,71 @@ extern "C" __global__ void experience_state_gather( for (int k = 0; k < market_dim; k++) out[k] = mf_row[k]; - /* -- Portfolio features: [market_dim .. market_dim+3) -- + /* -- Portfolio features: [market_dim .. market_dim+8) -- * - * slot +0: normalised portfolio value = portfolio_value / denom - * slot +1: raw position (contract units; network learns the scale) - * slot +2: normalised cash = cash / denom + * 8 informative features that let the Q-network SEE its trading context. + * Without these, the model is blind to risk, P&L, and position state. * - * denom = max(portfolio_value, 1.0) avoids division by zero on an - * empty account. This mirrors the current_norm / pos_norm pattern - * from the monolithic kernel (spread is omitted in Phase 3). */ + * slot +0: position / max_position — normalized position [-1, +1] + * slot +1: unrealized_pnl / equity — how is THIS trade doing? + * slot +2: drawdown — (peak - equity) / peak [0, 1] + * slot +3: hold_time / 100.0 — normalized holding duration + * slot +4: realized_pnl / equity — overall session P&L + * slot +5: distance_to_floor — (equity - floor) / equity [0, 1] + * slot +6: trade_return — trade P&L / equity + * slot +7: cash_ratio — cash / equity + * + * max_position is read from kernel arg (line ~370). We use a constant + * approximation here (max_pos not available in gather kernel). + * The network learns the actual scale via the position feature. + */ const float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE; float position = ps[0]; float cash = ps[1]; float portfolio_value = ps[2]; + float peak_equity = ps[7]; + float prev_equity = ps[9]; + float hold_time = ps[10]; + float realized_pnl = ps[11]; + float entry_price = ps[12]; + float trade_start_pnl = ps[13]; - float denom = (portfolio_value > 1.0f) ? portfolio_value : 1.0f; + float equity = (portfolio_value > 1.0f) ? portfolio_value : 1.0f; + float drawdown = (peak_equity > 1.0f) + ? (peak_equity - portfolio_value) / peak_equity + : 0.0f; + drawdown = (drawdown > 0.0f) ? drawdown : 0.0f; + + /* Unrealized P&L: current position mark-to-market minus entry cost */ + float unrealized_pnl = (entry_price > 0.0f && position != 0.0f) + ? position * (market_features[(long long)bar_idx * market_dim] - entry_price) + : 0.0f; + + /* Trade return since entry */ + float trade_return = (trade_start_pnl != 0.0f || realized_pnl != 0.0f) + ? (realized_pnl - trade_start_pnl + unrealized_pnl) / equity + : 0.0f; + + /* Capital floor distance (75% of peak = capital floor) */ + float floor = peak_equity * 0.75f; + float floor_dist = (equity > floor && equity > 1.0f) + ? (equity - floor) / equity + : 0.0f; int portfolio_base = market_dim; - if (portfolio_base + 2 < state_dim) { - out[portfolio_base + 0] = portfolio_value / denom; - out[portfolio_base + 1] = position; - out[portfolio_base + 2] = cash / denom; + if (portfolio_base + 7 < state_dim) { + out[portfolio_base + 0] = position; /* raw position */ + out[portfolio_base + 1] = unrealized_pnl / equity; /* trade P&L signal */ + out[portfolio_base + 2] = drawdown; /* risk: how deep are we? */ + out[portfolio_base + 3] = hold_time / 100.0f; /* how long in trade? */ + out[portfolio_base + 4] = realized_pnl / equity; /* session P&L */ + out[portfolio_base + 5] = floor_dist; /* distance to game over */ + out[portfolio_base + 6] = trade_return; /* this trade's return */ + out[portfolio_base + 7] = cash / equity; /* available capital */ } /* -- Zero-pad remaining dimensions (tensor-core alignment) -- */ - int filled = market_dim + 3; + int filled = market_dim + 8; for (int k = filled; k < state_dim; k++) out[k] = 0.0f; } @@ -589,9 +629,10 @@ extern "C" __global__ void experience_env_step( if (exiting_trade && hold_time > 0.0f) { float trade_pnl = ps[11] + raw_pnl - trade_start_pnl; float trade_return = trade_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f); - /* Scale: a 0.1% trade return becomes 0.1 reward. + /* Scale: 1 ES tick ($12.50 on $35K) = 0.036% return × 1000 = 0.36 reward. + * Must be stronger than idle penalty (max 0.05) to incentivize quality trades. * Clamp to [-1, +1] so exceptional trades don't dominate. */ - trade_completion_reward = fmaxf(-1.0f, fminf(1.0f, trade_return * 100.0f)); + trade_completion_reward = fmaxf(-1.0f, fminf(1.0f, trade_return * 1000.0f)); } /* Idle counter (flat bars) */ @@ -600,7 +641,9 @@ extern "C" __global__ void experience_env_step( } else { flat_counter = 0.0f; } - float idle_penalty_t = fminf(0.05f, flat_counter * 0.001f); + /* Idle penalty: tiny nudge (max 0.01), must be << trade completion reward. + * Old max 0.05 competed with 1-tick trade reward 0.036 — wrong incentive. */ + float idle_penalty_t = fminf(0.01f, flat_counter * 0.0002f); /* ==== Component 5: Position-time decay (losing positions only) ==== */ float hold_time_penalty = 0.0f; diff --git a/crates/ml/src/cuda_pipeline/gpu_action_selector.rs b/crates/ml/src/cuda_pipeline/gpu_action_selector.rs index 41f0989c5..f8374edee 100644 --- a/crates/ml/src/cuda_pipeline/gpu_action_selector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_action_selector.rs @@ -18,9 +18,9 @@ static EPSILON_GREEDY_PTX: OnceLock> = OnceLock::new(); fn compile_kernel_ptx(context: &CudaContext) -> Result { let defines = "\ - #define STATE_DIM 48\n\ + #define STATE_DIM 56\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n"; + #define PORTFOLIO_DIM 8\n"; let common_src = include_str!("common_device_functions.cuh"); let kernel_src = include_str!("epsilon_greedy_kernel.cu"); let full_source = format!("{defines}{common_src}\n{kernel_src}"); diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 731aed809..94c560063 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -93,7 +93,7 @@ fn compile_ppo_forward_ptx(context: &CudaContext) -> Result { let defines = "\ #define STATE_DIM 48\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n"; + #define PORTFOLIO_DIM 8\n"; let common_src = include_str!("common_device_functions.cuh"); let kernel_src = include_str!("backtest_forward_ppo_kernel.cu"); let full_source = format!("{defines}{common_src}\n{kernel_src}"); @@ -112,7 +112,7 @@ fn compile_backtest_dqn_ptx(context: &CudaContext) -> Result { let defines = "\ #define STATE_DIM 48\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n"; + #define PORTFOLIO_DIM 8\n"; let kernel_src = include_str!("experience_kernels.cu"); let full_source = format!("{defines}{kernel_src}"); crate::cuda_pipeline::compile_ptx_for_device(&full_source, context) @@ -566,7 +566,7 @@ impl GpuBacktestEvaluator { // Portfolio dimension is always 3: (normalised value, position, spread_cost). // state_dim = feature_dim + 3, then 8-aligned for H100 tensor core HMMA dispatch. // The CUDA gather_states kernel zero-pads positions [feat_dim+3 .. state_dim). - const PORTFOLIO_DIM: usize = 3; + const PORTFOLIO_DIM: usize = 8; let state_dim = (feature_dim + PORTFOLIO_DIM + 7) & !7; let states_buf = stream .alloc_zeros::(n_windows * state_dim) diff --git a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs index a474a244a..6d4772d5b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs @@ -56,7 +56,7 @@ fn compile_curiosity_training_ptx(context: &CudaContext) -> Result let defines = "\ #define STATE_DIM 48\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n"; + #define PORTFOLIO_DIM 8\n"; let common_src = include_str!("common_device_functions.cuh"); let kernel_src = include_str!("curiosity_training_kernel.cu"); let full_source = format!("{defines}{common_src}\n{kernel_src}"); diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 9cf8e4f77..71db9297e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -2271,7 +2271,7 @@ fn compile_training_kernels( let dim_overrides = format!( "#define STATE_DIM {state_dim}\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n", + #define PORTFOLIO_DIM 8\n", state_dim = config.state_dim, ); diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 8b947ce05..27ec2eaa9 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -1369,7 +1369,7 @@ fn compile_experience_kernels( let dim_overrides = format!( "#define STATE_DIM {state_dim}\n\ #define MARKET_DIM {market_dim}\n\ - #define PORTFOLIO_DIM 3\n" + #define PORTFOLIO_DIM 8\n" ); let full_source = format!("{dim_overrides}\n{kernel_src}"); diff --git a/crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs b/crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs index 0cb66f5a3..9c58b3385 100644 --- a/crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs @@ -288,7 +288,7 @@ impl GpuPpoExperienceCollector { let defines = "\ #define STATE_DIM 48\n\ #define MARKET_DIM 42\n\ - #define PORTFOLIO_DIM 3\n"; + #define PORTFOLIO_DIM 8\n"; let common_src = include_str!("common_device_functions.cuh"); let kernel_src = include_str!("ppo_experience_kernel.cu"); let full_source = format!("{defines}{common_src}\n{kernel_src}"); diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index aed01af5a..48cea917a 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -2627,7 +2627,7 @@ impl HyperparameterOptimizable for DQNTrainer { // penalize VRAM estimate for replay buffer capacity. let vram_buffer = if budget.gpu_memory_mb <= 8192 { 0 } else { clamped_buffer_size }; // Aligned state_dim: 56 with OFI (53→56), 48 without (45→48) - let aligned_state_dim: usize = if self.mbp10_data_dir.is_some() { 56 } else { 48 }; + let aligned_state_dim: usize = if self.mbp10_data_dir.is_some() { 64 } else { 56 }; let vram_check = budget.trial_fits_vram( params.hidden_dim_base, params.batch_size, diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index b40b51f68..42bbf294d 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -619,6 +619,7 @@ impl DQNTrainer { // Wave 16 Portfolio Features enable_action_masking, max_position, + current_epoch: 0, multi_asset_portfolio, stress_tester, diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 35a99e98c..1f27c5e21 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -105,6 +105,8 @@ pub struct DQNTrainer { pub enable_action_masking: bool, /// Maximum position size for action masking (default: 2.0) pub max_position: f64, + /// Current epoch (for Q-gap warm-up ramp) + pub(crate) current_epoch: usize, /// Entropy regularizer for preventing policy collapse (None if disabled) // entropy_regularizer removed — SAC-style entropy is computed directly on Q-value tensors in DQN::compute_loss_internal /// Multi-asset portfolio tracker (None if single-asset mode) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 025f06e9f..c60b27d71 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -77,6 +77,7 @@ impl DQNTrainer { // Training loop for epoch in 0..self.hyperparams.epochs { + self.current_epoch = epoch; self.reset_epoch_state(epoch); log_epoch_start(epoch + 1, self.hyperparams.epochs, self.hyperparams.learning_rate); @@ -784,7 +785,15 @@ impl DQNTrainer { fill_spread_cost_frac: 0.50, fill_spread_capture_frac: 0.50, fill_simulation_enabled: self.median_vol > 0.0, - q_gap_threshold: self.hyperparams.q_gap_threshold as f32, + // Q-gap cold-start fix: ramp threshold from 0 to target over 10 epochs. + // At epoch 0, Q-values are random — forcing Q-gap > 0.1 prevents ALL + // trades (except 10% epsilon), starving the model of learning signal. + // Linear ramp: epoch 0→0.0, epoch 5→0.05, epoch 10→full threshold. + q_gap_threshold: { + let target = self.hyperparams.q_gap_threshold as f32; + let epoch = self.current_epoch.min(10) as f32; + target * (epoch / 10.0) + }, dsr_eta: self.hyperparams.dsr_eta as f32, n_steps: self.hyperparams.n_steps as i32, ..Default::default()