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

phase3(env-unification): extract unified_env_step_core to trade_physics.cuh + port val

Browse Source 
Core of the Phase 3 unification: instead of writing a brand-new
unified_env_kernel.cu (3-day rewrite per the design doc), extract the
canonical step logic into a single __device__ __forceinline__ helper
that BOTH kernels call. Drift between train/val becomes structurally
impossible — any change to the helper applies to both atomically.

unified_env_step_core (trade_physics.cuh) encapsulates:
   1. Action decode (4-branch: dir, mag, order, urgency)
   2. Hold action passthrough
   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 0.005/1.0/1.0 — regime-adaptive remains
      a deferred gem, see trade_physics.cuh comment block)
   7. execute_trade with sqrt-impact (spread_scale = -1.0)
   8. record_kelly_trade_outcome (Kelly stats update)
   9. entry_price update (new entry / reversal / flat)
  10. hold_time tick via update_hold_time
  11. step_return = (new_value − prev_equity) / prev_equity (PURE P&L)
  12. Post-enforcement actual_dir + actual_mag (for actions_history)

Pass-by-pointer for all mutable state (position, cash, entry_price,
hold_time, max_equity, Kelly stats). Outputs: step_return, new_value,
prev_position_sign, actual_dir, actual_mag, trail_triggered.

Ported backtest_env_step (single-step variant) to call the helper —
replaced ~100 lines of inline step logic with a single call. Kept the
capital-floor pre-check / post-check / actions_history stitching as
per-kernel logic (output buffer formats differ between train and val,
so these stay per-kernel).

Files touched:
  crates/ml/src/cuda_pipeline/trade_physics.cuh      (+172)
  crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (-99 / +27 net)

Follow-up in a separate commit:
  - Port backtest_env_step_batch (same refactor, batched variant)
  - Port experience_env_step to call unified_env_step_core
    (subset — training's body has many more layers: counterfactuals,
    plan_params, reward shaping bundle — all of which STAY in the
    caller; only the core step gets unified)

Verified: cargo check passes. Smoke test in flight to confirm
behavioral equivalence (should be bit-identical — same arithmetic in
same order, just refactored into a helper).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-21 10:33:21 +02:00
parent bd022154f4
commit a3b6bc2f3b
2 changed files with 220 additions and 139 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -183,108 +183,40 @@ extern "C" __global__ void backtest_env_step(
return; return;
} }
// ── Decode action (4-branch: trade_physics.cuh) ─────────────────────── /* Suppress unused HLOC: kernel only uses close (notional MTM model). */
int action_val = actions[w];
int dir_idx = decode_direction_4b(action_val, b1_size, b2_size, b3_size);
int mag_idx = decode_magnitude_4b(action_val, b1_size, b2_size, b3_size);
int order_type_idx = decode_order_4b(action_val, b2_size, b3_size);
// ── Hold action (dir_idx==1): skip ALL trade execution ──────────────
// Hold means "keep current position unchanged, zero transaction cost."
// Matches training kernel (experience_env_step) Hold logic.
int is_hold_action = (dir_idx == 1);
float target_exposure;
if (is_hold_action) {
target_exposure = position; // Hold: target = current position
} else {
target_exposure = compute_target_position_4branch(dir_idx, mag_idx, max_position);
}
// ── Margin-aware position cap (shared: trade_physics.cuh) ────────────
// margin = price * multiplier * margin_pct (e.g. 5000 * 50 * 0.06 = $15K for ES)
if (!is_hold_action) {
float margin_per_contract = close * contract_multiplier * margin_pct;
target_exposure = apply_margin_cap(target_exposure, value,
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; (void)open; (void)high; (void)low;
// ── Trailing stop (shared: trade_physics.cuh) ──────────────────────── /* ── UNIFIED CORE STEP (shared: trade_physics.cuh::unified_env_step_core) ──
// Exit when profit retreats from peak. Uses 0.5% base distance. * One source of truth for the canonical env step. The training kernel
// vol_scale=1.0, trend_scale=1.0: backtest has no ADX/CUSUM features, * (experience_env_step) calls the same helper for the same computation,
// so regime-adaptive scaling is disabled (neutral = no widening). * guaranteeing step_return parity at the source level. */
if (!is_hold_action) { int action_val = actions[w];
// Portfolio-weighted trade return: unrealized P&L / equity int prev_position_sign = 0;
// Matches training kernel's unrealized_trade_pnl computation int actual_dir = -1, actual_mag = 0;
float trade_ret = 0.0f; int trail_triggered = 0;
if (fabsf(position) > 0.001f
&& entry_price > 0.0f
&& value > 1.0f) {
float unrealized = position * (close - entry_price);
trade_ret = unrealized / value;
}
if (check_trailing_stop(hold_time,
max_equity, value,
trade_ret, 0.005f, 1.0f, 1.0f)) {
target_exposure = 0.0f; // Force flat — trailing stop triggered
}
}
// ── Execute trade (shared: trade_physics.cuh) ────────────────────────
float prev_position = position; float prev_position = position;
{ float new_value = 0.0f;
float tx_cost = execute_trade(&position, &cash, float step_ret_core = 0.0f;
target_exposure, close, float health_k = (isv_signals != NULL) ? isv_signals[12] : 0.5f;
tx_cost_bps, spread_cost, max_position,
order_type_idx, -1.0f);
(void)tx_cost;
}
// ── Kelly stats update on trade exit/reversal (shared: trade_physics.cuh) ── unified_env_step_core(
// Records the completed-trade return into (win_count, loss_count, action_val, close,
// sum_wins, sum_losses). Uses pre-trade equity (`value`) so returns b1_size, b2_size, b3_size,
// are measured on the same scale as training's segment_return. max_position, tx_cost_bps, spread_cost,
record_kelly_trade_outcome(prev_position, position, contract_multiplier, margin_pct,
entry_price, close, value, health_k,
&win_count, &loss_count, &sum_wins, &sum_losses); &position, &cash, &entry_price, &hold_time,
&max_equity,
// Update entry_price only on NEW trade entry or reversal (for trailing stop) value, /* prev_equity (snapshot) */
if (fabsf(prev_position) < 0.001f &win_count, &loss_count, &sum_wins, &sum_losses,
&& fabsf(position) > 0.001f) { &step_ret_core, /* pure P&L (caller may add shaping) */
entry_price = close; // new entry from flat &new_value,
} else if (prev_position * position < 0.0f) { &prev_position_sign,
entry_price = close; // reversal &actual_dir,
} else if (fabsf(position) < 0.001f) { &actual_mag,
entry_price = 0.0f; // went flat &trail_triggered
} );
// On same-direction scaling (L50->L100), keep original entry_price (void)prev_position; (void)prev_position_sign; (void)trail_triggered;
// ── Hold time tracking (shared: trade_physics.cuh) ───────────────────
hold_time = update_hold_time(prev_position, position, hold_time);
// Mark-to-market current position (notional model: equity = cash + position * price)
float new_value = cash + position * close;
// Update max_equity BEFORE floor check — prevents stale peak from
// missing breaches or triggering false ones after profitable trades.
max_equity = fmaxf(max_equity, new_value);
// ── Post-trade floor check: catch intra-step breaches ──────────── // ── Post-trade floor check: catch intra-step breaches ────────────
// The pre-trade check (top of kernel) catches breaches from the // The pre-trade check (top of kernel) catches breaches from the
@@ -313,32 +245,17 @@ extern "C" __global__ void backtest_env_step(
return; return;
} }
// Step return — PURE per-bar P&L (Phase 3 unified-env policy). /* step_return is PURE per-bar P&L computed by unified_env_step_core
// * (Phase 3 unified-env policy). Shaping args kept for ABI stability
// Validation reports what *deployment* would experience: the actual * but ignored. */
// change in portfolio equity. Behavioral shaping (inventory penalty, float step_ret = step_ret_core;
// churn penalty, opportunity cost) belongs in TRAINING reward only —
// those terms steer the policy but should not be subtracted from the
// measurement we use to judge whether the policy is profitable in
// production. Equivalent to running experience_env_step with
// shaping_scale = 0.
//
// Removed (2026-04-21):
// inventory_penalty = holding_cost_rate * |position| // shaping
// churn_penalty = scale * (threshold - hold_time)/.. // shaping
// opp_cost penalty (was opp_cost_scale * is_flat * 0) // dead, scale=0
// These are still suppressed via the unused-args contract — the kernel
// signature retains the parameters for ABI stability but ignores them.
float step_ret = (value > 0.0f)
? (new_value - value) / value
: 0.0f;
(void)holding_cost_rate; (void)holding_cost_rate;
(void)churn_threshold; (void)churn_threshold;
(void)churn_penalty_scale; (void)churn_penalty_scale;
(void)opp_cost_scale; (void)opp_cost_scale;
float new_cum_return = cum_return + step_ret; float new_cum_return = cum_return + step_ret;
float new_max = fmaxf(max_equity, new_value); float new_max = max_equity; /* already updated inside unified_env_step_core */
// Write portfolio state // Write portfolio state
portfolio_state[ps + 0] = new_value; portfolio_state[ps + 0] = new_value;
@@ -359,28 +276,10 @@ extern "C" __global__ void backtest_env_step(
// Outputs // Outputs
step_rewards[w] = step_ret; step_rewards[w] = step_ret;
step_returns[w * max_len + current_step] = step_ret; step_returns[w * max_len + current_step] = step_ret;
// Record the ACTUAL action taken (post-enforcement, post-trailing-stop). /* Record ACTUAL action taken (post-enforcement, post-trailing-stop).
// Re-encode position -> exposure_idx -> factored action to ensure metrics * actual_dir + actual_mag already computed by unified_env_step_core;
// kernel counts real position changes, not model-requested actions. * we just stitch back the caller's original order/urgency. */
{ {
/* Re-encode actual position as 4-branch factored action.
* Recover direction and magnitude from position, preserve order/urgency.
* Direction(4): Short(0), Hold(1), Long(2), Flat(3).
* Hold is only recorded when the model chose Hold and position is unchanged.
* Near-zero position → Flat(3), not Hold(1). */
float pos_sign = position;
int actual_dir;
if (is_hold_action && fabsf(position - prev_position) < 0.001f) {
actual_dir = 1; /* Hold: model chose Hold and position unchanged */
} else {
actual_dir = (pos_sign > 0.001f) ? 2 : (pos_sign < -0.001f) ? 0 : 3; /* Flat=3 */
}
float abs_pos = fabsf(position) / fmaxf(max_position, 0.01f);
int actual_mag;
if (abs_pos < 0.375f) actual_mag = 0; /* ~0.25 -> Quarter */
else if (abs_pos < 0.75f) actual_mag = 1; /* ~0.50 -> Half */
else actual_mag = 2; /* ~1.00 -> Full */
if (actual_dir == 1 || actual_dir == 3) actual_mag = 0; /* Hold/Flat: magnitude irrelevant */
int orig_order = decode_order_4b(action_val, b2_size, b3_size); int orig_order = decode_order_4b(action_val, b2_size, b3_size);
int orig_urgency = decode_urgency_4b(action_val, b3_size); int orig_urgency = decode_urgency_4b(action_val, b3_size);
int actual_action = actual_dir * b1_size * b2_size * b3_size int actual_action = actual_dir * b1_size * b2_size * b3_size

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

@@ -436,4 +436,186 @@ __device__ __forceinline__ int check_trailing_stop(
return 0; return 0;
} }
/* ══════════════════════════════════════════════════════════════════════
* UNIFIED ENV STEP CORE (Phase 3 — single source of truth)
*
* Captures the canonical environment step that BOTH experience_env_step
* (training) and backtest_env_step (validation) execute identically:
*
* 1. Action decode (4-branch: dir, mag, order)
* 2. Hold action passthrough
* 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)
* 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
* 10. hold_time tick via update_hold_time
* 11. step_return = (new_value − old_value) / old_value (PURE per-bar P&L)
*
* Pass-by-pointer for in/out state. Capital-floor check is done by caller
* BEFORE invoking this helper (different output buffer formats between
* training/val make in-helper handling awkward).
*
* Saboteur and shaping are NOT applied here — those are training-only
* concerns the caller layers on top by gating their own additional code
* with exploration_scale / shaping_scale before/after this call.
*
* Both kernels MUST call this for their step_return computation. Any
* change here propagates to both atomically — drift becomes structurally
* impossible.
* ══════════════════════════════════════════════════════════════════════ */
__device__ __forceinline__ void unified_env_step_core(
/* ── Action + price (caller-provided) ── */
int action_val,
float close,
/* ── Branch sizes (4-branch action layout) ── */
int b1_size, int b2_size, int b3_size,
/* ── Static config ── */
float max_position,
float tx_cost_multiplier, /* training: includes saboteur/curriculum scaling; val: tx_cost_bps as-is */
float spread_cost, /* training: scaled by saboteur/curriculum; val: as-is */
float contract_multiplier,
float margin_pct,
float health_for_kelly, /* clamped to [0,1] before passing — usually isv_signals[12] */
/* ── Per-thread mutable state (in/out) ── */
float* position,
float* cash,
float* entry_price,
float* hold_time,
float* max_equity,
float prev_equity, /* read-only snapshot for step_return denominator + trailing stop */
/* ── Kelly stats (in/out) ── */
float* win_count, float* loss_count,
float* sum_wins, float* sum_losses,
/* ── Outputs ── */
float* step_return_out, /* pure per-bar P&L = (new_value - prev_equity) / prev_equity */
float* new_value_out, /* post-trade equity = cash + position * close */
int* prev_position_sign_out, /* sign of position BEFORE the trade — for caller's reversal logic */
int* actual_dir_out, /* post-enforcement direction (Hold / Long / Short / Flat) — for caller's actions_history */
int* actual_mag_out, /* post-enforcement magnitude bucket (0/1/2 for Quarter/Half/Full; 0 if Hold/Flat) */
int* trail_triggered_out /* 1 if trailing stop forced position to 0 */
) {
/* ── Step 1: Decode action (4-branch) ── */
int dir_idx = decode_direction_4b(action_val, b1_size, b2_size, b3_size);
int mag_idx = decode_magnitude_4b(action_val, b1_size, b2_size, b3_size);
int order_type_idx = decode_order_4b(action_val, b2_size, b3_size);
int is_hold_action = (dir_idx == 1);
/* ── Step 2-3: Target position ── */
float target_position;
if (is_hold_action) {
target_position = *position;
} else {
target_position = compute_target_position_4branch(dir_idx, mag_idx, max_position);
}
/* ── Step 4: Margin cap ── */
if (!is_hold_action) {
float margin_per_contract = close * contract_multiplier * margin_pct;
target_position = apply_margin_cap(
target_position, prev_equity, margin_per_contract, *max_equity
);
}
/* ── Step 5: Kelly cap (health-coupled safety) ── */
if (!is_hold_action) {
float h = fminf(1.0f, fmaxf(0.0f, health_for_kelly));
float safety = 0.5f + 0.5f * h;
target_position = apply_kelly_cap(
target_position,
*win_count, *loss_count,
*sum_wins, *sum_losses,
max_position, safety
);
}
/* ── Step 6: Trailing stop ── */
int trail_triggered = 0;
if (!is_hold_action && fabsf(*position) > 0.001f) {
float trade_ret = 0.0f;
if (*entry_price > 0.0f && prev_equity > 1.0f) {
float unrealized = (*position) * (close - *entry_price);
trade_ret = unrealized / prev_equity;
}
if (check_trailing_stop(
*hold_time, *max_equity, prev_equity,
trade_ret, 0.005f, 1.0f, 1.0f
)) {
target_position = 0.0f;
trail_triggered = 1;
}
}
*trail_triggered_out = trail_triggered;
/* ── Step 7: Execute trade — spread_scale = -1.0 (sqrt impact, no override) ── */
float prev_position = *position;
int prev_sign = (prev_position > 0.001f) ? 1
: (prev_position < -0.001f) ? -1 : 0;
*prev_position_sign_out = prev_sign;
float tx_cost = execute_trade(
position, cash,
target_position, close,
tx_cost_multiplier, spread_cost, max_position,
order_type_idx, -1.0f
);
(void)tx_cost;
/* ── Step 8: Kelly stats update on trade exit/reversal ── */
record_kelly_trade_outcome(
prev_position, *position,
*entry_price, close, prev_equity,
win_count, loss_count, sum_wins, sum_losses
);
/* ── Step 9: entry_price update ── */
if (fabsf(prev_position) < 0.001f && fabsf(*position) > 0.001f) {
*entry_price = close; /* new entry from flat */
} else if (prev_position * (*position) < 0.0f) {
*entry_price = close; /* reversal */
} else if (fabsf(*position) < 0.001f) {
*entry_price = 0.0f; /* went flat */
}
/* same-sign scaling: keep original entry_price */
/* ── Step 10: hold_time tick ── */
*hold_time = update_hold_time(prev_position, *position, *hold_time);
/* ── Step 11: Mark-to-market + step_return ── */
float new_value = (*cash) + (*position) * close;
*new_value_out = new_value;
*max_equity = fmaxf(*max_equity, new_value);
*step_return_out = (prev_equity > 0.0f)
? (new_value - prev_equity) / prev_equity
: 0.0f;
/* ── Action re-encoding for actions_history (post-enforcement direction)
* Caller stitches order_type/urgency back if they need the full action,
* since the post-enforcement direction is the only thing this helper
* can determine. */
int actual_dir;
if (is_hold_action && fabsf(*position - prev_position) < 0.001f) {
actual_dir = 1; /* Hold (model chose Hold and position unchanged) */
} else {
float pos_sign = *position;
actual_dir = (pos_sign > 0.001f) ? 2
: (pos_sign < -0.001f) ? 0
: 3; /* Flat */
}
float abs_pos = fabsf(*position) / fmaxf(max_position, 0.01f);
int actual_mag;
if (abs_pos < 0.375f) actual_mag = 0; /* Quarter */
else if (abs_pos < 0.75f) actual_mag = 1; /* Half */
else actual_mag = 2; /* Full */
if (actual_dir == 1 || actual_dir == 3) actual_mag = 0; /* Hold/Flat: mag irrelevant */
*actual_dir_out = actual_dir;
*actual_mag_out = actual_mag;
}
#endif /* TRADE_PHYSICS_CUH */ #endif /* TRADE_PHYSICS_CUH */