diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 310ba3443..445578902 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -63,6 +63,11 @@ extern "C" __global__ void backtest_env_step( // Market data (read-only, uploaded once) const float* __restrict__ prices, // [n_windows * max_len * 4] (OHLC) const int* __restrict__ window_lens, // [n_windows] + // Pre-uploaded feature tensor [n_windows * max_len * market_dim] — read-only. + // Needed for regime-adaptive trailing stop (ADX at idx 40, CUSUM at idx 41). + // Pass NULL (and market_dim <= 41) to fall back to fixed-width stop. + const float* __restrict__ features, + int market_dim, // Actions from model for current step const int* __restrict__ actions, // [n_windows] factored action index @@ -199,12 +204,23 @@ extern "C" __global__ void backtest_env_step( float step_ret_core = 0.0f; float health_k = (isv_signals != NULL) ? isv_signals[12] : 0.5f; + /* Regime-adaptive trailing-stop scales — same helper as training so the + * stop fires on identical ADX/CUSUM thresholds. features buffer is + * [n_windows, max_len, market_dim] so the flat bar index is w*max_len+step. */ + float trail_vol_scale = 1.0f; + float trail_trend_scale = 1.0f; + compute_regime_trail_scales( + features, (long long)w * max_len + current_step, market_dim, + &trail_vol_scale, &trail_trend_scale + ); + unified_env_step_core( action_val, close, b1_size, b2_size, b3_size, max_position, max_position, tx_cost_bps, spread_cost, contract_multiplier, margin_pct, health_k, + trail_vol_scale, trail_trend_scale, &position, &cash, &entry_price, &hold_time, &max_equity, value, /* prev_equity (snapshot) */ @@ -303,6 +319,8 @@ extern "C" __global__ void backtest_env_step( extern "C" __global__ void backtest_env_step_batch( const float* __restrict__ prices, const int* __restrict__ window_lens, + const float* __restrict__ features, /* [n_windows, max_len, market_dim] or NULL */ + int market_dim, /* <=41 disables regime-adaptive trail */ const int* __restrict__ chunked_actions, /* [chunk_len * n_windows] */ float* portfolio_state, float* step_rewards, @@ -420,12 +438,22 @@ extern "C" __global__ void backtest_env_step_batch( float step_ret_core = 0.0f; float health_k_batch = (isv_signals != NULL) ? isv_signals[12] : 0.5f; + /* Regime-adaptive trailing stop — shared helper keeps formulas + * identical to the single-step kernel and the training kernel. */ + float trail_vol_scale_b = 1.0f; + float trail_trend_scale_b = 1.0f; + compute_regime_trail_scales( + features, (long long)w * max_len + current_step, market_dim, + &trail_vol_scale_b, &trail_trend_scale_b + ); + unified_env_step_core( action_val, close, b1_size, b2_size, b3_size, max_position, max_position, tx_cost_bps, spread_cost, contract_multiplier, margin_pct, health_k_batch, + trail_vol_scale_b, trail_trend_scale_b, &position, &cash, &entry_price, &hold_time, &max_equity, value, /* prev_equity (snapshot for step_return) */ diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 12dc24399..0eb54ece1 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1407,6 +1407,16 @@ extern "C" __global__ void experience_env_step( float step_ret_core = 0.0f; float health_k_train = (isv_signals_ptr != NULL) ? isv_signals_ptr[12] : 0.5f; + /* Regime-adaptive trailing-stop scales — shared helper guarantees + * bit-identical computation with val kernel (features_buf layout differs + * but per-bar ADX/CUSUM resolve to the same floats). */ + float trail_vol_scale = 1.0f; + float trail_trend_scale = 1.0f; + compute_regime_trail_scales( + features, (long long)bar_idx, market_dim, + &trail_vol_scale, &trail_trend_scale + ); + unified_env_step_core( action_idx, raw_close, b1_size, b2_size, b3_size, @@ -1415,6 +1425,7 @@ extern "C" __global__ void experience_env_step( tx_cost_multiplier, spread_cost, contract_multiplier, margin_pct, health_k_train, + trail_vol_scale, trail_trend_scale, &position, &cash, &entry_price, &hold_time, &peak_equity, /* helper updates peak_equity via fmaxf */ prev_equity, /* snapshot for step_return denominator */ diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index be69138e1..b5c5c92b7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -1043,11 +1043,14 @@ impl GpuBacktestEvaluator { let start_step_i32 = chunk_start as i32; let chunk_len_i32 = chunk_len as i32; let env_blocks = ((n as u32) + 255) / 256; + let feature_dim_i32 = self.feature_dim as i32; unsafe { self.stream .launch_builder(&self.env_batch_kernel) .arg(&self.prices_buf) .arg(&self.window_lens_buf) + .arg(&self.features_buf) // regime-adaptive trail: ADX/CUSUM source + .arg(&feature_dim_i32) .arg(ch_actions) .arg(&self.portfolio_buf) .arg(&self.step_rewards_buf) @@ -1371,12 +1374,15 @@ impl GpuBacktestEvaluator { let b0_i32 = self.b0_size; let b1_i32 = self.b1_size; let b2_i32 = self.b2_size; + let feature_dim_i32 = self.feature_dim as i32; unsafe { self.stream .launch_builder(&self.env_kernel) .arg(&self.prices_buf) .arg(&self.window_lens_buf) + .arg(&self.features_buf) // regime-adaptive trail: ADX/CUSUM source + .arg(&feature_dim_i32) .arg(&self.actions_buf) .arg(&self.portfolio_buf) .arg(&self.step_rewards_buf) diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index 6cc551f2c..20d93dc8c 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -436,6 +436,42 @@ __device__ __forceinline__ int check_trailing_stop( return 0; } +/* ── Regime-adaptive trailing-stop scales from ADX/CUSUM features ────── + * Both training and val call this with their own features[] + bar_idx so + * the computation is bit-identical. Feature layout: index 40 = ADX + * (normalized, threshold 0.25 for trending), index 41 = CUSUM direction + * (in [-1,+1], |value| > 0.5 indicates regime change). + * + * Caller passes NULL features to fall back to 1.0/1.0 (fixed-width stop), + * which also covers kernels built before features were plumbed. */ +__device__ __forceinline__ void compute_regime_trail_scales( + const float* __restrict__ features, + long long bar_flat_idx, /* already resolved to features[bar_flat_idx * market_dim + ...] */ + int market_dim, + float* vol_scale_out, + float* trend_scale_out +) { + if (features == NULL || market_dim <= 41) { + *vol_scale_out = 1.0f; + *trend_scale_out = 1.0f; + return; + } + float adx = features[bar_flat_idx * (long long)market_dim + 40]; + float cusum = features[bar_flat_idx * (long long)market_dim + 41]; + /* Defensive clamp — feature pipeline guarantees [0,1]/[-1,1] but guard + * against out-of-range from data gaps to keep the stop bounded. */ + adx = fminf(fmaxf(adx, 0.0f), 1.0f); + float cusum_abs = fminf(fmaxf(fabsf(cusum), 0.0f), 1.0f); + /* Formulas mirror the pre-unification reward v5 (commit 263997ad3), + * translated to the normalized feature convention: + * trend_scale widens when ADX > 0.25 (trending), cap 2.5 + * vol_scale widens when |CUSUM| > 0.5 (regime-change), cap 2.5 */ + float trend_scale = 1.0f + fmaxf(adx - 0.25f, 0.0f) * 2.0f; + float vol_scale = 1.0f + fmaxf(cusum_abs - 0.5f, 0.0f) * 2.0f; + *trend_scale_out = fminf(trend_scale, 2.5f); + *vol_scale_out = fminf(vol_scale, 2.5f); +} + /* ══════════════════════════════════════════════════════════════════════ * UNIFIED ENV STEP CORE (Phase 3 — single source of truth) * @@ -447,8 +483,9 @@ __device__ __forceinline__ int check_trailing_stop( * 3. Target position via compute_target_position_4branch * 4. Margin cap (apply_margin_cap) * 5. Kelly cap with health-coupled safety (apply_kelly_cap) - * 6. Trailing stop check (fixed scales 0.005/1.0/1.0 — regime-adaptive - * removed earlier to eliminate train/eval mismatch) + * 6. Trailing stop check — base 0.5%, regime-adaptive via vol_scale and + * trend_scale (caller resolves from ADX/CUSUM features; identical + * formulas in both training and val to preserve train/eval parity) * 7. execute_trade with sqrt-impact tx_cost (spread_scale = -1.0) * 8. record_kelly_trade_outcome * 9. entry_price update on entry/reversal/flat @@ -485,6 +522,14 @@ __device__ __forceinline__ void unified_env_step_core( float contract_multiplier, float margin_pct, float health_for_kelly, /* clamped to [0,1] before passing — usually isv_signals[12] */ + /* ── Regime-adaptive trailing stop scales (caller resolves from features) ── + * Trail distance = 0.005f * vol_scale * trend_scale. Pass 1.0f/1.0f for + * fixed-width behaviour. Caller formulas (kept symmetric train↔val): + * vol_scale = min(2.5, 1 + max(|CUSUM| - 0.5, 0) * 2.0) + * trend_scale = min(2.5, 1 + max(ADX - 0.25, 0) * 2.0) + * with ADX, CUSUM read from features[bar * market_dim + 40/41]. */ + float vol_scale, + float trend_scale, /* ── Per-thread mutable state (in/out) ── */ float* position, float* cash, @@ -548,7 +593,7 @@ __device__ __forceinline__ void unified_env_step_core( } if (check_trailing_stop( *hold_time, *max_equity, prev_equity, - trade_ret, 0.005f, 1.0f, 1.0f + trade_ret, 0.005f, vol_scale, trend_scale )) { target_position = 0.0f; trail_triggered = 1;