From 69211af40b757c1bcde885507c3a73c74b5e509d Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 21 Apr 2026 11:04:51 +0200 Subject: [PATCH] =?UTF-8?q?phase3(env-unification):=20full=20unification?= =?UTF-8?q?=20=E2=80=94=20all=203=20kernels=20now=20call=20unified=5Fenv?= =?UTF-8?q?=5Fstep=5Fcore?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- .../src/cuda_pipeline/backtest_env_kernel.cu | 161 ++++------------ .../src/cuda_pipeline/experience_kernels.cu | 179 +++++++----------- crates/ml/src/cuda_pipeline/trade_physics.cuh | 21 +- ...2026-04-21-unified-train-val-env-design.md | 60 ++++-- 4 files changed, 161 insertions(+), 260 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu index 389588ee1..310ba3443 100644 --- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu @@ -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 diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 6984eccac..187eec99d 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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; } diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh index cf5d67a93..6cc551f2c 100644 --- a/crates/ml/src/cuda_pipeline/trade_physics.cuh +++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh @@ -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 */ diff --git a/docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md b/docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md index e6d19123f..68b1f9a8e 100644 --- a/docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md +++ b/docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md @@ -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)