diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 6562bee45..c858f76dd 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -11,10 +11,25 @@ // [6] cum_return - cumulative log return // [7] step_count - number of completed steps // -// Trade physics (action decoding, tx costs, capital floor) -// are shared with the training kernel via trade_physics.cuh. +// The 8-slot layout matches `SL_PORTFOLIO_BASE_DIM` — this buffer is read +// directly by `backtest_state_gather` (experience_kernels.cu) which indexes +// with stride 8. Changing this stride would corrupt window boundaries in +// the gather kernel. +// +// Kelly win/loss statistics used by the shared apply_kelly_cap helper +// live in a SEPARATE kelly_stats buffer of stride 4 per window: +// [0] win_count, [1] loss_count, [2] sum_wins, [3] sum_losses +// This keeps the gather semantics clean (the network doesn't see Kelly +// stats) while letting the env kernel apply a Kelly-constrained +// position cap and record completed-trade outcomes via the shared +// record_kelly_trade_outcome helper in trade_physics.cuh. +// +// Trade physics (action decoding, tx costs, capital floor, Kelly cap, +// Kelly stat updates) are ALL shared with the training kernel via +// trade_physics.cuh — no duplicated step logic between the two envs. #define PORTFOLIO_STATE_SIZE 8 +#define KELLY_STATS_SIZE 4 #include "trade_physics.cuh" // Capital floor breach: full episode restart for this window. @@ -77,7 +92,10 @@ extern "C" __global__ void backtest_env_step( float churn_penalty_scale, // scale factor for churn penalty float contract_multiplier, // e.g. 50.0 for ES, 20.0 for NQ float margin_pct, // e.g. 0.06 (6% initial margin) - float opp_cost_scale // opportunity cost scale (0.0 = disabled during backtest) + float opp_cost_scale, // opportunity cost scale (0.0 = disabled during backtest) + /* ── Kelly position cap (shared with training via trade_physics.cuh) ── */ + float* kelly_stats, // [n_windows * KELLY_STATS_SIZE=4]: win_count, loss_count, sum_wins, sum_losses + const float* __restrict__ isv_signals // [13] pinned; [12]=health. NULL → 0.5 fallback. ) { __shared__ float shmem_pf[256 * PORTFOLIO_STATE_SIZE]; @@ -118,6 +136,13 @@ extern "C" __global__ void backtest_env_step( float hold_time = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 5]; float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6]; + // Kelly stats from separate buffer (stride 4 per window) + int ks = w * KELLY_STATS_SIZE; + float win_count = kelly_stats[ks + 0]; + float loss_count = kelly_stats[ks + 1]; + float sum_wins = kelly_stats[ks + 2]; + float sum_losses = kelly_stats[ks + 3]; + // ── Capital floor circuit breaker (shared: trade_physics.cuh) ──────── if (check_capital_floor(value, max_equity)) { // Force flat: close any open position at current price (notional model) @@ -165,6 +190,20 @@ extern "C" __global__ void backtest_env_step( margin_per_contract, max_equity); } + // ── Kelly position cap (shared: trade_physics.cuh) ─────────────────── + // Health-coupled safety multiplier: safety = 0.5 + 0.5 × health. + // Same helper as training kernel — one source of truth for the cap. + if (!is_hold_action) { + float health_k = (isv_signals != NULL) + ? fminf(1.0f, fmaxf(0.0f, isv_signals[12])) + : 0.5f; + float safety = 0.5f + 0.5f * health_k; + target_exposure = apply_kelly_cap(target_exposure, + win_count, loss_count, + sum_wins, sum_losses, + max_position, safety); + } + // Suppress unused variable warnings (void)open; (void)high; (void)low; @@ -199,6 +238,14 @@ extern "C" __global__ void backtest_env_step( (void)tx_cost; } + // ── Kelly stats update on trade exit/reversal (shared: trade_physics.cuh) ── + // Records the completed-trade return into (win_count, loss_count, + // sum_wins, sum_losses). Uses pre-trade equity (`value`) so returns + // are measured on the same scale as training's segment_return. + record_kelly_trade_outcome(prev_position, position, + entry_price, close, value, + &win_count, &loss_count, &sum_wins, &sum_losses); + // Update entry_price only on NEW trade entry or reversal (for trailing stop) if (fabsf(prev_position) < 0.001f && fabsf(position) > 0.001f) { @@ -288,6 +335,12 @@ extern "C" __global__ void backtest_env_step( portfolio_state[ps + 6] = new_cum_return; portfolio_state[ps + 7] = portfolio_state[ps + 7] + 1.0f; + // Kelly stats — separate buffer, updated by record_kelly_trade_outcome + kelly_stats[ks + 0] = win_count; + kelly_stats[ks + 1] = loss_count; + kelly_stats[ks + 2] = sum_wins; + kelly_stats[ks + 3] = sum_losses; + // Outputs step_rewards[w] = step_ret; step_returns[w * max_len + current_step] = step_ret; @@ -358,13 +411,17 @@ extern "C" __global__ void backtest_env_step_batch( float churn_penalty_scale, float contract_multiplier, float margin_pct, - float opp_cost_scale /* opportunity cost scale (0.0 = disabled during backtest) */ + float opp_cost_scale, /* opportunity cost scale (0.0 = disabled during backtest) */ + /* ── Kelly position cap (shared with training via trade_physics.cuh) ── */ + float* kelly_stats, /* [n_windows * KELLY_STATS_SIZE=4] win_count, loss_count, sum_wins, sum_losses */ + const float* __restrict__ isv_signals /* [13] pinned; [12]=health. NULL → 0.5 fallback. */ ) { int w = blockIdx.x * blockDim.x + threadIdx.x; if (w >= n_windows) return; int wlen = window_lens[w]; int ps = w * PORTFOLIO_STATE_SIZE; + int ks = w * KELLY_STATS_SIZE; /* Load portfolio into registers — persistent across the step loop */ float value = portfolio_state[ps + 0]; @@ -375,6 +432,13 @@ extern "C" __global__ void backtest_env_step_batch( float hold_time = portfolio_state[ps + 5]; float cum_return = portfolio_state[ps + 6]; + /* Kelly stats from dedicated buffer — kept in registers, updated on trade + * exits below via record_kelly_trade_outcome (shared: trade_physics.cuh). */ + float win_count = kelly_stats[ks + 0]; + float loss_count = kelly_stats[ks + 1]; + float sum_wins = kelly_stats[ks + 2]; + float sum_losses = kelly_stats[ks + 3]; + int steps_processed = 0; for (int s = 0; s < chunk_len; s++) { @@ -440,6 +504,18 @@ extern "C" __global__ void backtest_env_step_batch( margin_per_contract, max_equity); } + /* ── Kelly position cap (shared: trade_physics.cuh) ───────────── */ + if (!is_hold_action) { + float health_k = (isv_signals != NULL) + ? fminf(1.0f, fmaxf(0.0f, isv_signals[12])) + : 0.5f; + float safety = 0.5f + 0.5f * health_k; + target_exposure = apply_kelly_cap(target_exposure, + win_count, loss_count, + sum_wins, sum_losses, + max_position, safety); + } + /* ── Trailing stop ────────────────────────────────────────────── */ if (!is_hold_action) { float trade_ret = 0.0f; @@ -466,6 +542,11 @@ extern "C" __global__ void backtest_env_step_batch( (void)tx_cost; } + /* ── Kelly stats update on trade exit/reversal (shared: trade_physics.cuh) ── */ + record_kelly_trade_outcome(prev_position, position, + entry_price, close, value, + &win_count, &loss_count, &sum_wins, &sum_losses); + /* ── Entry price tracking (matches original exactly) ──────────── */ /* Bug 6 fix: add entry_price = 0.0f when going flat */ if (fabsf(prev_position) < 0.001f @@ -588,4 +669,10 @@ extern "C" __global__ void backtest_env_step_batch( portfolio_state[ps + 6] = cum_return; portfolio_state[ps + 7] = portfolio_state[ps + 7] + (float)steps_processed; } + /* Kelly stats — always written back (accumulate across episodes, even + * if a capital-floor breach reset portfolio_state). */ + kelly_stats[ks + 0] = win_count; + kelly_stats[ks + 1] = loss_count; + kelly_stats[ks + 2] = sum_wins; + kelly_stats[ks + 3] = sum_losses; } diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index f73d416b0..9f6afdf65 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -83,6 +83,15 @@ static BACKTEST_DQN_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/exp /// (301 vs 2406 for 154K bars) outweighs the marginal staleness. const DQN_BACKTEST_CHUNK_SIZE: usize = 512; +/// Per-window Kelly-stats buffer width (f32 slots). Stored in a buffer +/// separate from `portfolio_buf` so the `backtest_state_gather` kernel's +/// stride-8 indexing into portfolio state stays correct while the env +/// kernel still has the stats it needs for the shared `apply_kelly_cap` +/// and `record_kelly_trade_outcome` helpers in `trade_physics.cuh`. +/// Slots: win_count, loss_count, sum_wins, sum_losses. Must equal +/// `KELLY_STATS_SIZE` in `backtest_env_kernel.cu`. +const BACKTEST_KELLY_STATS_SIZE: usize = 4; + // ── DQN backtest forward config ─────────────────────────────────────────────── /// Configuration for cuBLAS-based DQN forward pass in backtest evaluation. @@ -269,6 +278,9 @@ pub struct GpuBacktestEvaluator { // Mutable state (written by kernels each step) portfolio_buf: CudaSlice, // [n_windows * 8] + /// Kelly win/loss stats — separate from portfolio_buf so the gather + /// kernel's stride-8 indexing stays correct. [n_windows * 4]. + kelly_stats_buf: CudaSlice, /// Plan/ISV features per window [n_windows * 6 f32]: plan_progress, pnl_vs_target, pnl_vs_stop, conviction, conviction_drift, regime_shift plan_isv_buf: CudaSlice, step_rewards_buf: CudaSlice, // [n_windows] @@ -479,6 +491,15 @@ impl GpuBacktestEvaluator { let portfolio_buf = stream.clone_htod(&portfolio_init) .map_err(|e| MLError::ModelError(format!("portfolio upload: {e}")))?; + // Kelly stats buffer — zero-initialized (no trades yet). The Kelly + // priors + warmup floor in trade_physics.cuh handle the cold start + // gracefully: during the first ~10 completed trades the cap is + // dominated by the 50% warmup floor, then transitions to pure + // data-driven Kelly. + let kelly_stats_buf = stream + .alloc_zeros::(n_windows * BACKTEST_KELLY_STATS_SIZE) + .map_err(|e| MLError::ModelError(format!("kelly_stats alloc: {e}")))?; + // TODO(task-4-followup): integrate real plan/ISV computation with backtest // env step kernel. For now, plan_isv_buf is zero-filled — validation sees // zeros in state positions [86..92) while training sees real plan/ISV. @@ -538,6 +559,7 @@ impl GpuBacktestEvaluator { features_buf, window_lens_buf, portfolio_buf, + kelly_stats_buf, plan_isv_buf, step_rewards_buf, step_returns_buf, @@ -856,6 +878,11 @@ impl GpuBacktestEvaluator { self.stream.memset_zeros(&mut self.actions_history_buf) .map_err(|e| MLError::ModelError(format!("reset actions_history: {e}")))?; + // Reset Kelly stats — each evaluation starts cold. Kelly priors + + // warmup floor in trade_physics.cuh handle the cold-start period. + self.stream.memset_zeros(&mut self.kelly_stats_buf) + .map_err(|e| MLError::ModelError(format!("reset kelly_stats: {e}")))?; + Ok(()) } @@ -1044,6 +1071,9 @@ impl GpuBacktestEvaluator { .arg(&self.config.contract_multiplier) .arg(&self.config.margin_pct) .arg(&self.config.opportunity_cost_scale) + // Kelly cap shared with training: stats buffer + ISV[12]=health + .arg(&self.kelly_stats_buf) + .arg(&self.isv_signals_ptr) .launch(LaunchConfig { grid_dim: (env_blocks.max(1), 1, 1), block_dim: (256, 1, 1), @@ -1366,6 +1396,9 @@ impl GpuBacktestEvaluator { .arg(&self.config.contract_multiplier) .arg(&self.config.margin_pct) .arg(&self.config.opportunity_cost_scale) + // Kelly cap shared with training: stats buffer + ISV[12]=health + .arg(&self.kelly_stats_buf) + .arg(&self.isv_signals_ptr) .launch(env_cfg) .map_err(|e| { MLError::ModelError(format!("backtest_env_step launch step {step}: {e}")) diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index 5b64ae0e0..d7f928cfe 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -312,6 +312,53 @@ __device__ __forceinline__ float apply_kelly_cap( return fmaxf(-cap, fminf(cap, target_position)); } +/* ── Record a completed-trade outcome into Kelly statistics ─────────── + * + * Shared by both training (experience_kernels.cu) and validation + * (backtest_env_kernel.cu) env steps so the Kelly stats feeding + * apply_kelly_cap are computed from identical trade-exit semantics in + * both environments. Eliminates the stats-update duplication that would + * otherwise drift — exactly the mismatch we're trying to fix with the + * shared trade-physics module. + * + * A "trade completion" is: positioned → flat, OR direction reversal. + * Both represent a realized P&L event on the prior position. Partial + * scaling (L50 → L100) is NOT counted — no prior position was closed. + * + * The realized trade return is computed as + * (prev_position × (close − entry_price)) / pre_trade_equity + * normalized by PRE-trade equity so winners and losers are measured on + * the same scale. Break-even trades (return exactly 0) are ignored. + * + * Stats are passed by pointer so the caller's register-resident floats + * are updated in place. Caller owns writing back to global/shared memory. */ +__device__ __forceinline__ void record_kelly_trade_outcome( + float prev_position, /* position before the trade executed this step */ + float curr_position, /* position after execute_trade this step */ + float entry_price, /* entry price recorded for the closed trade */ + float close, /* mark-to-market close price this step */ + float pre_trade_equity,/* portfolio value BEFORE this step's P&L */ + float* win_count, /* in/out: winning trades counter */ + float* loss_count, /* in/out: losing trades counter */ + float* sum_wins, /* in/out: cumulative winning returns */ + float* sum_losses /* in/out: cumulative |losing returns| */ +) { + int prev_pos_active = (fabsf(prev_position) > 0.001f); + int is_exit = prev_pos_active && (fabsf(curr_position) < 0.001f); + int is_reversal = (prev_position * curr_position < 0.0f); + if ((is_exit || is_reversal) && entry_price > 0.0f && pre_trade_equity > 1.0f) { + float realized = prev_position * (close - entry_price); + float trade_return = realized / pre_trade_equity; + if (trade_return > 0.0f) { + *win_count += 1.0f; + *sum_wins += trade_return; + } else if (trade_return < 0.0f) { + *loss_count += 1.0f; + *sum_losses += fabsf(trade_return); + } + } +} + /* ── Drawdown penalty: smooth ramp from threshold to floor ───────────── */ /* Returns a penalty in [-5*w_dd, 0]. Applied every step so the model */ /* learns to reduce position size DURING drawdown. */