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

gem(trail): regime-adaptive trailing stop — symmetric train/val wiring

Browse Source 
Restores regime adaptation to the 0.5% trailing stop using ADX (trend
strength, feat idx 40) and CUSUM (directional persistence, feat idx 41).
Previously removed to eliminate train/val asymmetry (val kernel had no
features buffer); this commit wires features to both kernels so the
stop fires on identical thresholds in training and backtest evaluation.

Implementation (V7 methodology, all three steps satisfied):
  * Step 1 signal: trending regimes need wider stop to ride trend;
    volatile regimes need wider stop to avoid noise exits. Original
    formulation from reward v5 (commit 263997ad3).
  * Step 2 coverage: no existing mechanism adapts exit *distance* to
    regime — Kelly cap scales entry *size*, policy chooses exit *timing*
    but not exit *distance*. So this is a genuine gap.
  * Step 3 measurement: shared compute_regime_trail_scales helper
    guarantees bit-identical scale computation across training and val;
    multi-trial smoke (5 trials × 20 epochs) passes 5/5 with
    median_q_gap=2.00, Best Sharpe 27.05 in last trial. No regression
    vs fixed-width (pre: median_q_gap=2.24; post: 2.00, within variance).

Shared helper (trade_physics.cuh::compute_regime_trail_scales):
  trend_scale = min(2.5, 1 + max(ADX   - 0.25, 0) * 2.0)
  vol_scale   = min(2.5, 1 + max(|CUS| - 0.50, 0) * 2.0)
Pass NULL features to fall back to 1.0/1.0 (fixed-width). Thresholds
match existing regime_conditional convention (adx>0.25=Trending).

Kernel-side:
  * unified_env_step_core: add vol_scale / trend_scale params; forward
    to check_trailing_stop (which already accepted these but callers
    passed 1.0/1.0). Updated step-6 docstring.
  * experience_env_step (training): computes scales from existing
    features+market_dim inputs before calling the helper.
  * backtest_env_step + backtest_env_step_batch (val): new `features` +
    `market_dim` params plumbed through the single-step and batched
    launches.

Rust-side (gpu_backtest_evaluator.rs):
  * launch_env_step: pass self.features_buf + feature_dim as i32
    (features were already uploaded for the state_gather kernel).
  * Batched launcher: same.

Net effect: train/val env kernels now see identical regime-adaptive
trailing stops — no drift, no measurement gap, symmetric physics.

Verified:
  cargo check/build clean (release); multi-trial smoke 5/5 pass.

Closes task #25.
This commit is contained in:
jgrusewski
2026-04-21 14:46:38 +02:00
parent c823865008
commit cb2015ab2f
4 changed files with 93 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

11
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

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

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

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

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

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