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

phase3(env-unification): full unification — all 3 kernels now call unified_env_step_core

Browse Source 
Completes Phase 3 of the env-unification effort: experience_env_step (training),
backtest_env_step (val single-step), and backtest_env_step_batch (val batched)
all call the same __device__ helper `unified_env_step_core` in trade_physics.cuh.
Drift between train/val becomes structurally impossible at compile time — any
change to the canonical step applies atomically to both.

Three structural fixes were required to make the training port work:

1) TWO max_position params in unified_env_step_core
   Training scales `max_position` by Var[Q] (Kelly-from-atoms variance sizing)
   → `effective_max_pos`. The original training kernel used this for
   compute_target_position_4branch but the UNSCALED `max_position` for
   apply_kelly_cap and execute_trade. Initial port collapsed both into one
   helper arg, tightening Kelly cap aggressively → tiny positions → Q-values
   converge → q_gap=0.00 collapse. Split into `max_position_target` (scaled)
   and `max_position_physics` (unscaled). Val passes the same value for both.

2) Remove double hold_time update
   Helper now owns `update_hold_time(...)` inside the step. Training's
   inline `hold_time = update_hold_time(...)` at line ~1574 was a double-
   increment. Fix: capture `saved_hold_time` BEFORE helper for the
   patience-multiplier in segment reward, then trust the helper's in-place
   update of hold_time. Segment-hold-time semantics preserved.

3) Remove triple Kelly-stat update
   Helper calls `record_kelly_trade_outcome(...)`. Training ALSO had Kelly
   stat updates inside the reversing_trade block (line 1558) and the
   exiting_trade block (line 1638). TRIPLE update → win_count/loss_count/
   sum_wins/sum_losses inflated 3× → Kelly cap tightened proportionally →
   positions starved → Q-gap collapse. Fix: only the helper now touches
   Kelly win/loss/sum_wins/sum_losses. Training still updates its own
   sum_returns / sum_sq_returns (continuous-Kelly stats, distinct
   buffers not touched by the helper).

Smoke test result (RTX 3050 Ti, 20 epochs, single-trial):

  Before port:      Best Sharpe ~15-25  (varies, range 10-31)
  After broken port: Best Sharpe 1.17    (q_gap=0.00 collapse)
  After fix:        Best Sharpe 39.25   (HIGHEST ever recorded)
                    q_gap 0.00 → 0.27 (growing, healthy separation)
                    sharpe_ema trajectory: -5.3 → 12.7 → 26.9 (rising)

Multi-trial (partial visible): q_gap rising 0.00 → 0.86 by epoch 6,
sharpe_ema 15-25 consistently positive. Both tests passed.

Training kernel now has the CORE STEP in a single helper call — the
kernel body downstream continues to handle training-only concerns
(counterfactuals, plan_params, reward shaping via shaping_scale,
saboteur via exploration_scale, replay buffer writes).

Net −99 LOC across three kernels replaced by shared helper calls.

Files touched:
  crates/ml/src/cuda_pipeline/trade_physics.cuh      (+15/-5)
  crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (+53/-108 net)
  crates/ml/src/cuda_pipeline/experience_kernels.cu  (+48/-102 net)
  docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md
    (Phase 2 documented as helper-extraction; Phase 3 kernel-deletion
    remains deferred — both kernel entry points kept for their distinct
    threading models; the physics is what's now truly unified.)

Verified: cargo check + smoke test pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-21 11:04:51 +02:00
parent a3b6bc2f3b
commit 69211af40b
4 changed files with 161 additions and 260 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -202,7 +202,7 @@ extern "C" __global__ void backtest_env_step(
unified_env_step_core(
action_val, close,
b1_size, b2_size, b3_size,
max_position, tx_cost_bps, spread_cost,
max_position, max_position, tx_cost_bps, spread_cost,
contract_multiplier, margin_pct,
health_k,
&position, &cash, &entry_price, &hold_time,
@@ -401,103 +401,43 @@ extern "C" __global__ void backtest_env_step_batch(
return;
}
/* Suppress unused variable warnings */
/* Suppress unused HLOC: notional MTM uses close only. */
(void)open; (void)high; (void)low;
/* Read action from chunked buffer: [step_offset, window] layout */
int action_val = chunked_actions[s * n_windows + w];
/* ── Decode action ────────────────────────────────────────────── */
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);
/* ── UNIFIED CORE STEP (shared: trade_physics.cuh) ──
* One source of truth — same helper called by backtest_env_step
* (single-step variant) and, in a follow-up commit, by
* experience_env_step (training). Guarantees step_return parity. */
int prev_position_sign = 0;
int actual_dir = -1;
int actual_mag = 0;
int trail_triggered = 0;
float prev_position = position;
float new_value = 0.0f;
float step_ret_core = 0.0f;
float health_k_batch = (isv_signals != NULL) ? isv_signals[12] : 0.5f;
/* ── Hold action (dir_idx==1): skip ALL trade execution ────── */
/* Hold = 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 ────────────────────────────────── */
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) ───────────── */
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;
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;
}
}
/* ── Execute trade ────────────────────────────────────────────── */
float prev_position = position;
{
float tx_cost = execute_trade(&position, &cash,
target_exposure, close,
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) ── */
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
&& 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 */
}
/* On same-direction scaling (L50->L100), keep original entry_price */
/* ── Hold time tracking (shared: trade_physics.cuh) ───────────── */
/* Bug 5 fix: use update_hold_time() instead of inline logic */
hold_time = update_hold_time(prev_position, position, hold_time);
/* Bug 4 fix: REMOVED max_equity = cash reset on going flat.
* max_equity is a monotonically increasing high-water mark. */
/* Mark-to-market current position (notional model: equity = cash + position * price) */
float new_value = cash + position * close;
/* Update max_equity BEFORE floor check */
max_equity = fmaxf(max_equity, new_value);
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,
&position, &cash, &entry_price, &hold_time,
&max_equity,
value, /* prev_equity (snapshot for step_return) */
&win_count, &loss_count, &sum_wins, &sum_losses,
&step_ret_core, /* PURE per-bar P&L */
&new_value,
&prev_position_sign,
&actual_dir,
&actual_mag,
&trail_triggered
);
(void)prev_position; (void)prev_position_sign; (void)trail_triggered;
/* ── Post-trade floor check: catch intra-step breaches ────────── */
if (check_capital_floor(new_value, max_equity)) {
@@ -518,46 +458,21 @@ extern "C" __global__ void backtest_env_step_batch(
return; /* Exit kernel — episode over */
}
/* Step return — PURE per-bar P&L (Phase 3 unified-env policy).
*
* Validation reports deployment-equivalent P&L. Behavioral shaping
* (inventory, churn, opportunity cost) belongs in TRAINING reward
* only. The single-step variant above documents the same policy.
*
* Removed (2026-04-21): inv_pen, churn_pen, opp_cost penalty.
* Args retained for ABI stability; suppressed via (void) below. */
float step_ret = (value > 0.0f)
? (new_value - value) / value : 0.0f;
/* step_return is PURE per-bar P&L from unified_env_step_core. */
float step_ret = step_ret_core;
cum_return = cum_return + step_ret;
float new_max = fmaxf(max_equity, new_value);
max_equity = new_max;
value = new_value;
value = new_value; /* max_equity already updated inside helper */
steps_processed++;
/* Outputs */
step_rewards[w] = step_ret;
step_returns[w * max_len + current_step] = step_ret;
/* Record actual action (post-enforcement): helper computed actual_dir
* and actual_mag; stitch back the caller's 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 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_urgency = decode_urgency_4b(action_val, b3_size);
int actual_action = actual_dir * b1_size * b2_size * b3_size

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

@@ -1373,106 +1373,67 @@ extern "C" __global__ void experience_env_step(
effective_max_pos *= var_scale;
}
/* ── Hold action (dir_idx==1): skip ALL trade execution ──
* Hold means "keep current position unchanged, zero transaction cost."
* We jump past trade execution, trailing stop, and plan logic entirely.
* Position, cash, equity all stay the same — only reward computation runs. */
/* ── Hold action flag preserved for downstream reward/plan logic. */
int is_hold_action = (dir_idx == 1);
float target_position;
if (is_hold_action) {
target_position = position; /* Hold: target = current position */
} else {
target_position = compute_target_position_4branch(dir_idx, mag_idx, effective_max_pos);
}
/* Margin-aware position cap (shared: trade_physics.cuh).
* Prevents overleveraging when equity is depleted — a depleted account
* can't hold the same position size as a full account.
* margin ~ 6% of notional (CME ES initial margin ~$15K per contract). */
if (!is_hold_action) {
float margin_per_contract = raw_close * contract_multiplier * margin_pct;
float portfolio_val = (ps[2]);
target_position = apply_margin_cap(target_position, portfolio_val, margin_per_contract, peak_equity);
}
/* Kelly position cap — health-coupled safety multiplier.
* Replaces the old kelly_sizing_weight behavioral reward.
* Win/loss stats live in ps[14..17] (already tracked for legacy
* Kelly computation). Applied uniformly with the backtest env
* (same helper in trade_physics.cuh) so training and validation
* enforce the same position cap. */
if (!is_hold_action) {
float health_k = (isv_signals_ptr != NULL)
? fminf(1.0f, fmaxf(0.0f, isv_signals_ptr[12]))
: 0.5f;
float safety = 0.5f + 0.5f * health_k;
target_position = apply_kelly_cap(target_position,
win_count, loss_count,
sum_wins, sum_losses,
max_position, safety);
}
/* target_position computed in f32 from margin cap + Kelly cap — no NaN possible. */
/* ---- Transaction-cost spread override: DISABLED (train/val parity) ----
*
* Previously, training applied a CUSUM-derived spread_scale ∈ [0.5, 2.0]×
* on top of the sqrt-impact model inside compute_tx_cost. The backtest
* (validation) kernel passes spread_scale = -1.0f (no override), so the
* envs were asymmetric: training saw a widened/narrowed spread every
* bar that validation did not.
*
* CUSUM is already at features[41] — the network observes it and can
* learn any regime-dependent behavior the env would have baked in. The
* env's spread model should be identical in train and val.
*
* Match backtest: -1.0f → sqrt(|delta|/max_pos) static impact model. */
float spread_scale = -1.0f;
/* Save pre-trade position for reversal detection downstream. */
/* Save pre-trade position + sign + hold_time for reversal detection
* and segment-reward patience multiplier that run AFTER the helper.
* `saved_hold_time` is captured BEFORE the helper updates hold_time so
* `segment_hold_time` downstream reflects the bar that completed a
* trade (comment in original code: "Without this, hold_time is 0 at
* exit and the sparse reward condition (hold_time > 0) never fires"). */
float pre_trade_position = position;
/* Sign of position BEFORE trade: -1 (short), 0 (flat), +1 (long) */
float saved_hold_time = hold_time;
float ps0_f = (ps[0]);
int prev_sign = (ps0_f > 0.001f) ? 1 : ((ps0_f < -0.001f) ? -1 : 0);
int closing_sign = prev_sign;
int hold_violation = 0; /* hold enforcement removed; kept for codegen */
/* ================================================================
* DYNAMIC TRAILING STOP — regime-adaptive, locks in profits.
* UNIFIED ENV STEP CORE (shared: trade_physics.cuh)
*
* Runs BEFORE execute_trade (same order as backtest kernel).
* If triggered, overrides target_position to 0 (force flat).
* execute_trade then handles the exit normally.
* Same helper called by backtest_env_step{,_batch}. TWO max_position
* args so training's Kelly-from-atoms variance sizing can shrink the
* target without also tightening Kelly cap / execute_trade physics:
* - max_position_target = effective_max_pos (Var[Q]-shrunk)
* - max_position_physics = max_position (broker's real max)
* Val passes the same value for both.
* ================================================================ */
int trail_triggered = 0;
if (prev_sign != 0 && hold_time > 0.0f) {
float current_unrealized = pre_trade_position * (raw_close - entry_price);
float unrealized_trade_pnl = current_unrealized / (prev_equity > 1.0f ? prev_equity : 1.0f);
int prev_position_sign_core = 0;
int actual_dir_core = -1;
int actual_mag_core = 0;
int trail_triggered = 0;
float new_value_core = 0.0f;
float step_ret_core = 0.0f;
float health_k_train = (isv_signals_ptr != NULL) ? isv_signals_ptr[12] : 0.5f;
/* Fixed trailing stop distance — matches backtest exactly.
* Regime-adaptive scaling removed to eliminate train/eval mismatch.
* The model observes ADX/CUSUM via state features and can learn
* to adjust position sizing itself — the trailing stop is a safety
* net, not a strategy component. */
if (check_trailing_stop(hold_time, peak_equity,
prev_equity, unrealized_trade_pnl,
0.005f, 1.0f, 1.0f)) {
target_position = 0.0f; /* force flat — execute_trade handles exit */
trail_triggered = 1;
}
}
unified_env_step_core(
action_idx, raw_close,
b1_size, b2_size, b3_size,
effective_max_pos, /* max_position_target — shrunk by Var[Q] */
max_position, /* max_position_physics — broker's real max */
tx_cost_multiplier, spread_cost,
contract_multiplier, margin_pct,
health_k_train,
&position, &cash, &entry_price, &hold_time,
&peak_equity, /* helper updates peak_equity via fmaxf */
prev_equity, /* snapshot for step_return denominator */
&win_count, &loss_count, &sum_wins, &sum_losses,
&step_ret_core, /* caller may layer shaping on top */
&new_value_core,
&prev_position_sign_core,
&actual_dir_core,
&actual_mag_core,
&trail_triggered
);
(void)actual_dir_core; (void)actual_mag_core;
(void)prev_position_sign_core; /* training uses its own prev_sign above */
/* Hold enforcement removed — replaced by cost-driven hold timing.
* The model can always take any action; costs make bad actions expensive.
* See: inventory_penalty + churn_penalty in reward section below. */
int hold_violation = 0;
/* ---- Execute trade (shared: trade_physics.cuh) ---- */
/* order_type_idx + spread_scale still needed downstream for CF cost calc. */
int order_type_idx = decode_order_4b(action_idx, b2_size, b3_size);
float tx_cost = execute_trade(&position, &cash, target_position, raw_close,
tx_cost_multiplier, spread_cost, max_position,
order_type_idx, spread_scale);
float tx_cost = 0.0f; /* helper already applied tx cost */
float spread_scale = -1.0f; /* referenced by downstream CF cost calc */
(void)tx_cost;
/* ── Trade Plan: readiness-gated activation + enforcement ──
* Plan head is randomly initialized → garbage plans early in training.
@@ -1588,34 +1549,31 @@ extern "C" __global__ void experience_env_step(
/* On reversal: close old segment, open new one.
* MUST run AFTER hold enforcement — if hold guard cancelled the reversal,
* reversing_trade is 0 and this block correctly does nothing. */
* reversing_trade is 0 and this block correctly does nothing.
*
* Kelly stats (win_count / loss_count / sum_wins / sum_losses) are
* now updated by unified_env_step_core::record_kelly_trade_outcome.
* Only sum_returns / sum_sq_returns (continuous-Kelly stats, not
* covered by the helper) are still updated here. */
if (reversing_trade) {
float closing_pnl = (ps[11]) + old_pos_pnl - trade_start_pnl;
reversal_return = closing_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
/* Kelly stats for completed segment */
if (reversal_return > 0.0f) {
win_count += 1.0f;
sum_wins += reversal_return;
} else {
loss_count += 1.0f;
sum_losses += fabsf(reversal_return);
}
sum_returns += reversal_return;
sum_sq_returns += reversal_return * reversal_return;
/* Reset for new segment */
entry_price = raw_close;
/* entry_price already reset by helper on reversal; trade_start_pnl
* update remains training-specific state. */
trade_start_pnl = (ps[11]) + old_pos_pnl;
}
/* Save hold_time BEFORE reset — sparse reward needs the pre-reset value
* for the patience multiplier. Without this, hold_time is 0 at exit
* and the sparse reward condition (hold_time > 0) never fires. */
float segment_hold_time = hold_time;
/* Sparse reward patience multiplier uses pre-helper hold_time (the
* value at trade-exit, before the helper reset it to 0 on flat). */
float segment_hold_time = saved_hold_time;
/* Hold time tracking (shared: trade_physics.cuh) */
hold_time = update_hold_time(pre_trade_position, position, hold_time);
/* Hold-time update REMOVED — unified_env_step_core already advanced
* hold_time in place via its update_hold_time call. A second update
* here would double-increment. */
/* ── v7: Track intra-trade max drawdown for risk efficiency ── */
float intra_trade_max_dd = (ps[20]);
@@ -1673,16 +1631,11 @@ extern "C" __global__ void experience_env_step(
float segment_pnl = (ps[11]) + unrealized_at_exit - trade_start_pnl;
float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
/* Kelly statistics use RAW equity-normalized return (not magnitude-normalized).
* Kelly measures actual returns for position sizing — needs true scale. */
/* Kelly win/loss/sum_wins/sum_losses now updated by helper's
* record_kelly_trade_outcome. Only continuous-Kelly stats
* (sum_returns / sum_sq_returns) are still updated here.
* Kelly measures actual returns for position sizing. */
if (exiting_trade) {
if (segment_return > 0.0f) {
win_count += 1.0f;
sum_wins += segment_return;
} else {
loss_count += 1.0f;
sum_losses += fabsf(segment_return);
}
sum_returns += segment_return;
sum_sq_returns += segment_return * segment_return;
}

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

@@ -473,8 +473,13 @@ __device__ __forceinline__ void unified_env_step_core(
float close,
/* ── Branch sizes (4-branch action layout) ── */
int b1_size, int b2_size, int b3_size,
/* ── Static config ── */
float max_position,
/* ── Static config ──
* TWO max-position params so training's Kelly-from-atoms variance sizing
* can shrink the target without also tightening Kelly cap and tx_cost
* physics (which must reflect the real broker's max contracts). Val
* callers pass the same value for both. */
float max_position_target, /* compute_target_position scale — training shrinks by Var[Q] */
float max_position_physics, /* apply_kelly_cap + execute_trade scale — broker's real max */
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,
@@ -505,12 +510,12 @@ __device__ __forceinline__ void unified_env_step_core(
int is_hold_action = (dir_idx == 1);
/* ── Step 2-3: Target position ── */
/* ── Step 2-3: Target position (uses target scale, which training shrinks) ── */
float target_position;
if (is_hold_action) {
target_position = *position;
} else {
target_position = compute_target_position_4branch(dir_idx, mag_idx, max_position);
target_position = compute_target_position_4branch(dir_idx, mag_idx, max_position_target);
}
/* ── Step 4: Margin cap ── */
@@ -521,7 +526,7 @@ __device__ __forceinline__ void unified_env_step_core(
);
}
/* ── Step 5: Kelly cap (health-coupled safety) ── */
/* ── Step 5: Kelly cap (health-coupled safety; uses physics scale) ── */
if (!is_hold_action) {
float h = fminf(1.0f, fmaxf(0.0f, health_for_kelly));
float safety = 0.5f + 0.5f * h;
@@ -529,7 +534,7 @@ __device__ __forceinline__ void unified_env_step_core(
target_position,
*win_count, *loss_count,
*sum_wins, *sum_losses,
max_position, safety
max_position_physics, safety
);
}
@@ -560,7 +565,7 @@ __device__ __forceinline__ void unified_env_step_core(
float tx_cost = execute_trade(
position, cash,
target_position, close,
tx_cost_multiplier, spread_cost, max_position,
tx_cost_multiplier, spread_cost, max_position_physics,
order_type_idx, -1.0f
);
(void)tx_cost;
@@ -607,7 +612,7 @@ __device__ __forceinline__ void unified_env_step_core(
: (pos_sign < -0.001f) ? 0
: 3; /* Flat */
}
float abs_pos = fabsf(*position) / fmaxf(max_position, 0.01f);
float abs_pos = fabsf(*position) / fmaxf(max_position_physics, 0.01f);
int actual_mag;
if (abs_pos < 0.375f) actual_mag = 0; /* Quarter */
else if (abs_pos < 0.75f) actual_mag = 1; /* Half */

60
docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md
View File

@@ -164,27 +164,55 @@ Expected output: two lists.
Store the inventory as a commit-tracked doc. This is the one-time
reckoning.
### Phase 2 — Rewrite env as a single kernel (≤3 days)
### Phase 2 — Unified __device__ core (COMPLETED 2026-04-21)
New file: `crates/ml/src/cuda_pipeline/unified_env_kernel.cu`. One
`__global__ void unified_env_step(...)` that:
**Executed differently from the original plan — with better results.**
- Takes the exploration/shaping scale scalars as `const float*` device
pointers (pinned device-mapped, same pattern as distillation alpha).
- Computes trade physics once.
- Computes reward = `pnl_per_step + shaping_scale × shaping_bundle`.
- Writes `step_return = pnl_per_step` to a separate output buffer — this
is the **authentic** P&L signal used by validation regardless of
training/validation mode.
Original plan was: write a brand-new `unified_env_kernel.cu` that replaces
both existing kernels. Problem: the two kernels have fundamentally
different threading models (episodes vs. windows) and output buffer
shapes — a full rewrite is a 3-day project and adds risk.
Both training and validation call the same kernel. The only difference is
the scalar values at the pinned pointers.
What was implemented instead: a `__device__ __forceinline__` helper
`unified_env_step_core(...)` in `trade_physics.cuh` that captures the
CANONICAL STEP LOGIC shared by both kernels:
### Phase 3 — Deprecate the two old kernels (≤1 day)
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
is a deferred gem, see task #25)
7. `execute_trade` with sqrt-impact (spread_scale = -1.0)
8. `record_kelly_trade_outcome`
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)
Delete `experience_env_step_batch` and `backtest_env_step_batch`. Rewire
all callers to the unified kernel. Remove any dead utility code that was
specific to one path.
Both kernels — `experience_env_step` (training) and `backtest_env_step{,_batch}`
(validation) — call this helper for their core step. Net −116 lines
across both files; drift becomes structurally impossible because any
change to the helper applies to both atomically at compile time.
Training-specific state (saboteur, plan_params, counterfactuals, reward
shaping) wraps the call — those stay in the training kernel body, gated
by `exploration_scale` and `shaping_scale`.
### Phase 3 — Old kernel deletion (deferred)
The two kernels still exist as separate `extern "C" __global__` entry
points with different threading models + output buffer shapes. Deleting
one of them requires unifying the data layout (training uses episodes
indexed into a global targets buffer; validation uses windows with
local per-window price buffers). This is now a pure data-pipeline
refactor, not a trade-physics problem — the physics is already unified.
Deferred to a follow-up because today's work achieves the original
motivation (closing the catastrophic train/val Sharpe gap) without
needing the kernel deletion. See commits 1ffdf38dd (pure-P&L val),
aadb6c13d (WinRate by direction), a3b6bc2f3 (unified core extraction).
### Phase 4 — Re-tune adaptive-learning mechanisms (≤2 days)