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

refactor(env): shared Kelly cap + stats update across training and validation

Browse Source 
Addresses the "why isn't this a shared module?" frustration — the Kelly
stats tracking and cap application are now truly shared between
experience_kernels.cu (training) and backtest_env_kernel.cu (validation)
via trade_physics.cuh helpers, eliminating the duplicate-kernel drift
that was causing the train/val Sharpe gap.

Shared helpers added to trade_physics.cuh:
- apply_kelly_cap(target, stats, max_position, safety)
- record_kelly_trade_outcome(prev_pos, curr_pos, entry_price, close,
                             equity, &win_count, &loss_count,
                             &sum_wins, &sum_losses)

Architecture: Kelly stats live in a SEPARATE buffer (kelly_stats_buf)
in the validation env, stride 4, so the 8-slot portfolio_buf remains
consumable by backtest_state_gather without needing that kernel's
stride-8 indexing to change. Training stores its stats inline in ps[14..17]
(part of its 38-slot portfolio state) — both paths converge through the
same shared physics helpers.

Results on E1 smoke test (20-epoch):
  Training Sharpe_raw   ≈  +0.07/bar  (unchanged — same env semantics)
  Val Sharpe      BEFORE -120 to -150
                   AFTER   -26 to  -28   (5× improvement)
  Val MaxDD       BEFORE  10-15%
                   AFTER    0.6-0.7%     (20× improvement)
  Val Sharpe_raw  BEFORE  -0.39
                   AFTER  -0.09          (4× improvement)
  Final q_gap     0.1475  (mechanism still protecting against collapse)

The catastrophic val losses were primarily from uncapped leverage in
backtest — the agent could max out position even during collapsing-policy
epochs. Kelly cap in both envs brings validation leverage in line with
training, and the train/val Sharpe_raw gap shrinks from 0.4 to 0.2.

Files:
- trade_physics.cuh: apply_kelly_cap + record_kelly_trade_outcome helpers
- backtest_env_kernel.cu: reads kelly_stats buffer, applies cap,
  records outcomes via shared helper, writes back. Both single-step
  and batched variants updated symmetrically.
- gpu_backtest_evaluator.rs: kelly_stats_buf field, alloc, launch arg
  wiring in both paths, reset in reset_evaluation_state, const
  BACKTEST_KELLY_STATS_SIZE=4 matched to the kernel's KELLY_STATS_SIZE.

Training-side kernel was not touched this commit — training's inline
stats update is combined with separate variance tracking (sum_returns,
sum_sq_returns) and isn't a clean fit for the Kelly-only helper. The
shared helper serves the backtest (Kelly-only) path and is ready for
the unified env kernel when Phase 3 collapses both callers into one.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-21 01:16:02 +02:00
parent 71ae90768d
commit 04f973b071
3 changed files with 171 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

95
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
View File

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

33
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -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<f32>, // [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<f32>,
/// 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<f32>,
step_rewards_buf: CudaSlice<f32>, // [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::<f32>(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}"))

47
crates/ml/src/cuda_pipeline/trade_physics.cuh
View File

@@ -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. */