refactor(dqn-v2): Invariant 8 — named constants for portfolio state ps[0..PS_STRIDE=38)

Replace raw ps[N] access across all CUDA kernels with PS_* named
constants defined in state_layout.cuh. 32 constants total covering
the full 38-slot portfolio state: position/cash/value, DSR stats,
Kelly stats, plan fields, and OFI scratch range.

Notable deviations from pre-written plan mapping: the plan's
PS_VALUE/PS_POSITION/PS_CASH values (0,1,2) had the wrong semantics.
Code wins (feedback_trust_code_not_docs): ps[0]=position, ps[1]=cash,
ps[2]=portfolio_value. All 148 raw ps[digit] accesses in
experience_kernels.cu and trade_stats_kernel.cu migrated. In-comment
range references (e.g. "ps[30..37]") left as documentation.

trade_stats_kernel.cu: migrated base+14 offset to PS_KELLY_WIN_COUNT.
docs/dqn-named-dims.md: PS_* table corrected to match actual code layout.

No behavioural change; pure refactor.

Plan 1 Task 4A. Spec §3 Invariant 8.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-24 11:19:25 +02:00
parent 46e99c5ccc
commit 7a2adeb374
4 changed files with 214 additions and 159 deletions

View File

@@ -528,15 +528,15 @@ extern "C" __global__ void experience_state_gather(
* portfolio[7]: cash_ratio — cash / equity
*/
const float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
float position = ps[0];
float cash = ps[1];
float portfolio_value = ps[2];
float peak_equity = ps[7];
float prev_equity = ps[9];
float hold_time = ps[10];
float realized_pnl = ps[11];
float entry_price = ps[12];
float trade_start_pnl = ps[13];
float position = ps[PS_POSITION];
float cash = ps[PS_CASH];
float portfolio_value = ps[PS_PORTFOLIO_VALUE];
float peak_equity = ps[PS_PEAK_EQUITY];
float prev_equity = ps[PS_PREV_EQUITY];
float hold_time = ps[PS_HOLD_TIME];
float realized_pnl = ps[PS_REALIZED_PNL];
float entry_price = ps[PS_ENTRY_PRICE];
float trade_start_pnl = ps[PS_TRADE_START_PNL];
/* ── Portfolio features computed in FLOAT to prevent bf16 overflow ──
* BF16 max ~65504. ES position × price_diff can exceed this (2.0 × 5000 = 10000,
@@ -588,10 +588,10 @@ extern "C" __global__ void experience_state_gather(
/* Plan-aware + ISV features: model sees its own plan progress.
* Enables temporal reasoning: "I'm 80% through my plan, close to target." */
float plan_isv[SL_PORTFOLIO_PLAN_DIM];
float plan_tgt_bars = ps[23];
float plan_profit = ps[24];
float plan_stop = ps[25];
float plan_conv = ps[27];
float plan_tgt_bars = ps[PS_PLAN_TARGET_BARS];
float plan_profit = ps[PS_PLAN_PROFIT_TARGET];
float plan_stop = ps[PS_PLAN_STOP_LOSS];
float plan_conv = ps[PS_PLAN_CONVICTION];
plan_isv[0] = (plan_tgt_bars > 0.5f)
? fminf(f_hold_time / plan_tgt_bars, 2.0f) /* plan progress [0, 2] */
@@ -606,16 +606,16 @@ extern "C" __global__ void experience_state_gather(
/* Conviction drift — current conviction minus entry conviction.
* Negative drift = model's thesis is weakening → exit signal.
* Only meaningful when a plan is active (ps[23] > 0.5). */
plan_isv[4] = (plan_params_ptr != NULL && ps[23] > 0.5f)
? (plan_params_ptr[i * 6 + 4] - ps[27]) /* conviction drift */
* Only meaningful when a plan is active (ps[PS_PLAN_TARGET_BARS] > 0.5). */
plan_isv[4] = (plan_params_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f)
? (plan_params_ptr[i * 6 + 4] - ps[PS_PLAN_CONVICTION]) /* conviction drift */
: 0.0f;
/* Regime shift since entry — |regime_stability_now - regime_stability_at_entry|.
* Large shift = regime changed mid-trade, plan may be invalid.
* Entry regime stored in ps[29] during plan activation. */
plan_isv[5] = (isv_signals_ptr != NULL && ps[23] > 0.5f)
? fabsf(isv_signals_ptr[11] - ps[29]) /* regime shift */
* Entry regime stored in ps[PS_PLAN_ENTRY_REGIME] during plan activation. */
plan_isv[5] = (isv_signals_ptr != NULL && ps[PS_PLAN_TARGET_BARS] > 0.5f)
? fabsf(isv_signals_ptr[11] - ps[PS_PLAN_ENTRY_REGIME]) /* regime shift */
: 0.0f;
/* -- Multi-timeframe features → local mtf[] array --
@@ -1551,23 +1551,23 @@ extern "C" __global__ void experience_env_step(
* require rewriting the shared header used by backtest_env_kernel too.
* We load bf16 → float here, compute, then store float → bf16 at end. */
float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE;
float position = (ps[0]);
float cash = (ps[1]);
/* ps[2] = portfolio_value (updated at end) */
float position = (ps[PS_POSITION]);
float cash = (ps[PS_CASH]);
/* ps[PS_PORTFOLIO_VALUE] = portfolio_value (updated at end) */
/* ps[3:6]: reserved (formerly trade clustering tracking, retained for slot stability) */
float peak_equity = (ps[7]);
float flat_counter = (ps[8]);
float prev_equity = (ps[9]);
float hold_time = (ps[10]);
/* ps[11] = realized_pnl (cumulative, updated at end) */
float entry_price = (ps[12]); /* price when trade was entered */
float trade_start_pnl = (ps[13]); /* realized_pnl snapshot at trade entry */
float win_count = (ps[14]); /* Kelly: number of profitable trade exits */
float loss_count = (ps[15]); /* Kelly: number of losing trade exits */
float sum_wins = (ps[16]); /* Kelly: cumulative profit from winners */
float sum_losses = (ps[17]); /* Kelly: cumulative |loss| from losers */
float sum_returns = (ps[18]); /* Kelly: cumulative net returns (for mu) */
float sum_sq_returns = (ps[19]); /* Kelly: cumulative squared returns (for sigma^2) */
float peak_equity = (ps[PS_PEAK_EQUITY]);
float flat_counter = (ps[PS_FLAT_COUNTER]);
float prev_equity = (ps[PS_PREV_EQUITY]);
float hold_time = (ps[PS_HOLD_TIME]);
/* ps[PS_REALIZED_PNL] = realized_pnl (cumulative, updated at end) */
float entry_price = (ps[PS_ENTRY_PRICE]); /* price when trade was entered */
float trade_start_pnl = (ps[PS_TRADE_START_PNL]); /* realized_pnl snapshot at trade entry */
float win_count = (ps[PS_KELLY_WIN_COUNT]); /* Kelly: number of profitable trade exits */
float loss_count = (ps[PS_KELLY_LOSS_COUNT]); /* Kelly: number of losing trade exits */
float sum_wins = (ps[PS_KELLY_SUM_WINS]); /* Kelly: cumulative profit from winners */
float sum_losses = (ps[PS_KELLY_SUM_LOSSES]); /* Kelly: cumulative |loss| from losers */
float sum_returns = (ps[PS_KELLY_SUM_RETURNS]); /* Kelly: cumulative net returns (for mu) */
float sum_sq_returns = (ps[PS_KELLY_SUM_SQ_RETURNS]); /* Kelly: cumulative squared returns (for sigma^2) */
/* Pre-trade capital floor: skip trade execution on blown accounts.
* When the floor triggers, write done=1 AND reset the portfolio to
@@ -1575,28 +1575,28 @@ extern "C" __global__ void experience_env_step(
* the reset, the blown portfolio persists and every subsequent step
* hits the floor again (producing fake 92% MaxDD from stuck episodes). */
{
float portfolio_val = (ps[2]);
float portfolio_val = (ps[PS_PORTFOLIO_VALUE]);
if (check_capital_floor(portfolio_val, peak_equity)) {
out_rewards[out_off] = -10.0f;
out_dones[out_off] = 1.0f;
/* Reset portfolio for next episode (fresh capital) */
float init_cap = peak_equity; /* use peak as initial for next episode */
ps[0] = 0.0f; /* position = flat */
ps[1] = (init_cap); /* cash = initial_capital */
ps[2] = (init_cap); /* portfolio_value */
ps[7] = (init_cap); /* peak_equity */
ps[8] = 0.0f; /* flat_counter */
ps[9] = (init_cap); /* prev_equity */
ps[10] = 0.0f; /* hold_time */
ps[11] = 0.0f; /* realized_pnl */
ps[12] = 0.0f; /* entry_price */
ps[13] = 0.0f; /* trade_start_pnl */
ps[14] = 0.0f; /* win_count (reset Kelly) */
ps[15] = 0.0f; /* loss_count */
ps[16] = 0.0f; /* sum_wins */
ps[17] = 0.0f; /* sum_losses */
ps[18] = 0.0f; /* sum_returns */
ps[19] = 0.0f; /* sum_sq_returns */
ps[PS_POSITION] = 0.0f; /* position = flat */
ps[PS_CASH] = (init_cap); /* cash = initial_capital */
ps[PS_PORTFOLIO_VALUE] = (init_cap); /* portfolio_value */
ps[PS_PEAK_EQUITY] = (init_cap); /* peak_equity */
ps[PS_FLAT_COUNTER] = 0.0f; /* flat_counter */
ps[PS_PREV_EQUITY] = (init_cap); /* prev_equity */
ps[PS_HOLD_TIME] = 0.0f; /* hold_time */
ps[PS_REALIZED_PNL] = 0.0f; /* realized_pnl */
ps[PS_ENTRY_PRICE] = 0.0f; /* entry_price */
ps[PS_TRADE_START_PNL] = 0.0f; /* trade_start_pnl */
ps[PS_KELLY_WIN_COUNT] = 0.0f; /* win_count (reset Kelly) */
ps[PS_KELLY_LOSS_COUNT] = 0.0f; /* loss_count */
ps[PS_KELLY_SUM_WINS] = 0.0f; /* sum_wins */
ps[PS_KELLY_SUM_LOSSES] = 0.0f; /* sum_losses */
ps[PS_KELLY_SUM_RETURNS] = 0.0f; /* sum_returns */
ps[PS_KELLY_SUM_SQ_RETURNS] = 0.0f; /* sum_sq_returns */
current_timesteps[i] = 0; /* reset episode timer */
return;
}
@@ -1641,12 +1641,12 @@ extern "C" __global__ void experience_env_step(
}
/* Kelly criterion: scale by optimal fraction after sufficient samples.
* win_count/loss_count already declared earlier from ps[14]/ps[15]. */
* win_count/loss_count already declared earlier from ps[PS_KELLY_WIN_COUNT]/ps[PS_KELLY_LOSS_COUNT]. */
{
float total_trades = win_count + loss_count;
if (total_trades >= 20.0f) {
float sum_wins_val = (ps[16]);
float sum_losses_val = (ps[17]);
float sum_wins_val = (ps[PS_KELLY_SUM_WINS]);
float sum_losses_val = (ps[PS_KELLY_SUM_LOSSES]);
float avg_win = sum_wins_val / fmaxf(win_count, 1.0f);
float avg_loss = sum_losses_val / fmaxf(loss_count, 1.0f);
float win_rate = win_count / total_trades;
@@ -1693,7 +1693,7 @@ extern "C" __global__ void experience_env_step(
* exit and the sparse reward condition (hold_time > 0) never fires"). */
float pre_trade_position = position;
float saved_hold_time = hold_time;
float ps0_f = (ps[0]);
float ps0_f = (ps[PS_POSITION]);
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 */
@@ -1788,36 +1788,36 @@ extern "C" __global__ void experience_env_step(
if (was_flat_plan && now_positioned_plan && plan_enabled) {
const float* pp = plan_params_ptr + i * 6;
ps[23] = pp[0]; /* target_bars */
ps[24] = pp[1]; /* profit_target */
ps[25] = pp[2]; /* stop_loss */
ps[26] = pp[3]; /* scale_aggression */
ps[27] = pp[4]; /* conviction */
ps[28] = pp[5]; /* asymmetry */
ps[PS_PLAN_TARGET_BARS] = pp[0]; /* target_bars */
ps[PS_PLAN_PROFIT_TARGET] = pp[1]; /* profit_target */
ps[PS_PLAN_STOP_LOSS] = pp[2]; /* stop_loss */
ps[PS_PLAN_SCALE_AGGRESSION] = pp[3]; /* scale_aggression */
ps[PS_PLAN_CONVICTION] = pp[4]; /* conviction */
ps[PS_PLAN_ASYMMETRY] = pp[5]; /* asymmetry */
/* Store entry regime_stability for drift detection (P12). */
ps[29] = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f;
ps[21] = 0.0f; /* reset max P&L for new trade */
ps[PS_PLAN_ENTRY_REGIME] = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f;
ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f; /* reset max P&L for new trade */
/* Apply conviction to position size — scale by readiness for smooth ramp */
position *= fmaxf(ps[27] * readiness, 0.1f);
position *= fmaxf(ps[PS_PLAN_CONVICTION] * readiness, 0.1f);
}
int has_plan = (ps[23] > 0.5f);
int has_plan = (ps[PS_PLAN_TARGET_BARS] > 0.5f);
/* Recursive plan revision (N17): living plan document.
* When model is highly confident (readiness ≥ 0.7) and conviction changes
* significantly, update the stored plan. The model learns WHEN to revise. */
if (has_plan && fabsf(position) > 0.001f && readiness >= 0.7f && plan_params_ptr != NULL) {
float conv_now = plan_params_ptr[i * 6 + 4];
float conv_entry = ps[27];
float conv_entry = ps[PS_PLAN_CONVICTION];
if (conv_now > conv_entry * 1.3f) {
/* Conviction increased 30%+ → thesis strengthening → extend plan */
float new_target = plan_params_ptr[i * 6 + 0];
ps[23] = fmaxf(ps[23], new_target); /* only extend, never shorten from conviction increase */
ps[24] = fmaxf(ps[24], plan_params_ptr[i * 6 + 1]); /* widen profit target */
ps[PS_PLAN_TARGET_BARS] = fmaxf(ps[PS_PLAN_TARGET_BARS], new_target); /* only extend, never shorten from conviction increase */
ps[PS_PLAN_PROFIT_TARGET] = fmaxf(ps[PS_PLAN_PROFIT_TARGET], plan_params_ptr[i * 6 + 1]); /* widen profit target */
} else if (conv_now < conv_entry * 0.4f) {
/* Conviction dropped 60%+ → thesis collapsing → shorten plan to exit soon */
ps[23] = fminf(ps[23], hold_time + 3.0f); /* exit within 3 bars */
ps[PS_PLAN_TARGET_BARS] = fminf(ps[PS_PLAN_TARGET_BARS], hold_time + 3.0f); /* exit within 3 bars */
}
/* Between 0.4x and 1.3x → plan stays as-is (normal fluctuation) */
}
@@ -1833,7 +1833,7 @@ extern "C" __global__ void experience_env_step(
/* ISV-driven thresholds: adapt to training dynamics + regime */
float isv_instability = (isv_signals_ptr != NULL) ? isv_signals_ptr[1] : 0.0f; /* grad_norm */
float regime_stab = (isv_signals_ptr != NULL) ? isv_signals_ptr[11] : 1.0f; /* regime_stability */
float plan_stop = ps[25]; /* learned stop from plan head */
float plan_stop = ps[PS_PLAN_STOP_LOSS]; /* learned stop from plan head */
/* Catastrophic stop: base 2% tightened by instability, widened by plan's learned stop */
float base_stop = fmaxf(plan_stop, 0.005f); /* at least 0.5%, or plan's learned value */
@@ -1913,7 +1913,7 @@ extern "C" __global__ void experience_env_step(
if (entering_trade) {
entry_price = raw_close;
trade_start_pnl = (ps[11]);
trade_start_pnl = (ps[PS_REALIZED_PNL]);
}
/* On reversal: close old segment, open new one.
@@ -1925,7 +1925,7 @@ extern "C" __global__ void experience_env_step(
* 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;
float closing_pnl = (ps[PS_REALIZED_PNL]) + old_pos_pnl - trade_start_pnl;
reversal_return = closing_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
sum_returns += reversal_return;
@@ -1933,7 +1933,7 @@ extern "C" __global__ void experience_env_step(
/* entry_price already reset by helper on reversal; trade_start_pnl
* update remains training-specific state. */
trade_start_pnl = (ps[11]) + old_pos_pnl;
trade_start_pnl = (ps[PS_REALIZED_PNL]) + old_pos_pnl;
}
/* Sparse reward patience multiplier uses pre-helper hold_time (the
@@ -1945,14 +1945,14 @@ extern "C" __global__ void experience_env_step(
* here would double-increment. */
/* ── v7: Track intra-trade max drawdown for risk efficiency ── */
float intra_trade_max_dd = (ps[20]);
float intra_trade_max_dd = (ps[PS_INTRA_TRADE_MAX_DD]);
if (entering_trade) {
intra_trade_max_dd = 0.0f;
ps[21] = 0.0f; /* reset max P&L for new trade */
ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f; /* reset max P&L for new trade */
}
if (reversing_trade) {
intra_trade_max_dd = 0.0f;
ps[21] = 0.0f; /* reset max P&L for reversed trade */
ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f; /* reset max P&L for reversed trade */
}
if (fabsf(position) > 0.001f && entry_price > 0.0f) {
float unrealized = (position > 0.0f)
@@ -1964,7 +1964,7 @@ extern "C" __global__ void experience_env_step(
/* Track best unrealized P&L during trade (for hindsight plan labels) */
float pnl_pct = (raw_close - entry_price) / fmaxf(fabsf(entry_price), 1.0f)
* ((position > 0.0f) ? 1.0f : -1.0f);
ps[21] = fmaxf(ps[21], pnl_pct); /* max P&L during trade */
ps[PS_INTRA_TRADE_MAX_PNL] = fmaxf(ps[PS_INTRA_TRADE_MAX_PNL], pnl_pct); /* max P&L during trade */
}
/* Idle counter (flat bars) — used for state features, no penalty */
@@ -2000,7 +2000,7 @@ extern "C" __global__ void experience_env_step(
* unrealized at exit = position * (raw_close - entry_price) * contract_multiplier.
* raw_pnl was removed (used raw_next = future price). Use raw_close directly. */
float unrealized_at_exit = pre_trade_position * (raw_close - entry_price) * contract_multiplier;
float segment_pnl = (ps[11]) + unrealized_at_exit - trade_start_pnl;
float segment_pnl = (ps[PS_REALIZED_PNL]) + unrealized_at_exit - trade_start_pnl;
float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
/* Kelly win/loss/sum_wins/sum_losses now updated by helper's
@@ -2132,10 +2132,10 @@ extern "C" __global__ void experience_env_step(
float sign_pos = (position > 0.0f) ? 1.0f : -1.0f;
/* OFI momentum: delta from previous bar (stored in ps[30..37]) */
float delta_depth = ofi_cur[1] - ps[31]; /* depth_imbalance_L1 */
float delta_vpin = ofi_cur[6] - ps[36]; /* VPIN */
float delta_queue = ofi_cur[2] - ps[32]; /* queue_pressure */
float delta_tarr = ofi_cur[5] - ps[35]; /* trade_arrival_rate */
float delta_depth = ofi_cur[1] - ps[PS_OFI_PREV_BASE + 1]; /* depth_imbalance_L1 */
float delta_vpin = ofi_cur[6] - ps[PS_OFI_PREV_BASE + 6]; /* VPIN */
float delta_queue = ofi_cur[2] - ps[PS_OFI_PREV_BASE + 2]; /* queue_pressure */
float delta_tarr = ofi_cur[5] - ps[PS_OFI_PREV_BASE + 5]; /* trade_arrival_rate */
float flow_momentum = sign_pos * 0.25f * (delta_depth + delta_vpin + delta_queue + delta_tarr);
/* ATR for normalization (same decode as vol normalization above) */
@@ -2315,32 +2315,32 @@ extern "C" __global__ void experience_env_step(
int done = trade_complete | capital_breach | data_end;
/* ---- Update full portfolio state (PORTFOLIO_STRIDE=38) ---- */
ps[0] = (position);
ps[1] = (cash);
ps[2] = (new_portfolio_value);
ps[PS_POSITION] = (position);
ps[PS_CASH] = (cash);
ps[PS_PORTFOLIO_VALUE] = (new_portfolio_value);
/* ps[3:6]: reserved (formerly clustering stats, retained for slot stability) */
ps[7] = (peak_equity);
ps[8] = (flat_counter);
ps[9] = (new_portfolio_value); /* prev_equity = current equity for next step */
ps[10] = (hold_time);
ps[PS_PEAK_EQUITY] = (peak_equity);
ps[PS_FLAT_COUNTER] = (flat_counter);
ps[PS_PREV_EQUITY] = (new_portfolio_value); /* prev_equity = current equity for next step */
ps[PS_HOLD_TIME] = (hold_time);
/* realized_pnl accumulates. On reversal bars, use old_pos_pnl
* (saved before position was overwritten at ps[0] = position).
* (saved before position was overwritten at ps[PS_POSITION] = position).
* For non-reversal bars, raw_pnl (= current position's PnL) is correct. */
if (reversing_trade) {
ps[11] = ((ps[11]) + old_pos_pnl); /* old position's PnL */
ps[PS_REALIZED_PNL] = ((ps[PS_REALIZED_PNL]) + old_pos_pnl); /* old position's PnL */
} else {
ps[11] = ((ps[11]) + raw_pnl); /* normal: current position's PnL */
ps[PS_REALIZED_PNL] = ((ps[PS_REALIZED_PNL]) + raw_pnl); /* normal: current position's PnL */
}
ps[12] = (entry_price); /* preserved across bars of a trade */
ps[13] = (trade_start_pnl); /* realized_pnl snapshot at trade entry */
ps[14] = (win_count); /* Kelly: accumulated across episode */
ps[15] = (loss_count);
ps[16] = (sum_wins);
ps[17] = (sum_losses);
ps[18] = (sum_returns); /* Kelly continuous: sum returns */
ps[19] = (sum_sq_returns); /* Kelly continuous: sum returns^2 */
ps[20] = (intra_trade_max_dd); /* v7: intra-trade max drawdown */
/* ps[22]: reserved (formerly clustering stats, retained for slot stability) */
ps[PS_ENTRY_PRICE] = (entry_price); /* preserved across bars of a trade */
ps[PS_TRADE_START_PNL] = (trade_start_pnl); /* realized_pnl snapshot at trade entry */
ps[PS_KELLY_WIN_COUNT] = (win_count); /* Kelly: accumulated across episode */
ps[PS_KELLY_LOSS_COUNT] = (loss_count);
ps[PS_KELLY_SUM_WINS] = (sum_wins);
ps[PS_KELLY_SUM_LOSSES] = (sum_losses);
ps[PS_KELLY_SUM_RETURNS] = (sum_returns); /* Kelly continuous: sum returns */
ps[PS_KELLY_SUM_SQ_RETURNS] = (sum_sq_returns); /* Kelly continuous: sum returns^2 */
ps[PS_INTRA_TRADE_MAX_DD] = (intra_trade_max_dd); /* v7: intra-trade max drawdown */
/* ps[PS_PREV_MID]: reserved (formerly clustering stats, retained for slot stability) */
/* ---- Store current OFI in ps[30..37] for next bar's delta computation ---- */
/* Also store prev mid-price and prev close for micro-reward / DSR. */
@@ -2348,7 +2348,7 @@ extern "C" __global__ void experience_env_step(
const float* ofi_cur = full_batch_states + (long long)i * full_state_dim + SL_OFI_START;
for (int k = 0; k < 8; k++) {
/* Update previous OFI for next bar's delta computation */
ps[30 + k] = ofi_cur[k];
ps[PS_OFI_PREV_BASE + k] = ofi_cur[k];
}
}
ps[PREV_MID_SLOT] = tgt[5]; /* MBP-10 mid-price at this decision point */
@@ -2583,41 +2583,41 @@ extern "C" __global__ void experience_env_step(
if (data_end || capital_breach) {
/* Hard reset: new data window or capital wipeout — full restart */
float init_cap = (peak_equity); /* use peak as starting capital for next ep */
ps[0] = 0.0f; /* position = flat */
ps[1] = init_cap; /* cash */
ps[2] = init_cap; /* portfolio_value */
ps[7] = init_cap; /* peak_equity */
ps[8] = 0.0f; /* flat_counter */
ps[9] = init_cap; /* prev_equity */
ps[10] = 0.0f; /* hold_time */
ps[11] = 0.0f; /* realized_pnl */
ps[12] = 0.0f; /* entry_price */
ps[13] = 0.0f; /* trade_start_pnl */
ps[14] = 0.0f; /* win/loss/Kelly counters */
ps[15] = 0.0f;
ps[16] = 0.0f;
ps[17] = 0.0f;
ps[18] = 0.0f;
ps[19] = 0.0f;
ps[20] = 0.0f; /* intra_trade_max_dd */
ps[21] = 0.0f; /* intra_trade_max_pnl */
ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f;
ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f;
ps[29] = 0.0f; /* plan slots zeroed */
ps[PS_POSITION] = 0.0f; /* position = flat */
ps[PS_CASH] = init_cap; /* cash */
ps[PS_PORTFOLIO_VALUE] = init_cap; /* portfolio_value */
ps[PS_PEAK_EQUITY] = init_cap; /* peak_equity */
ps[PS_FLAT_COUNTER] = 0.0f; /* flat_counter */
ps[PS_PREV_EQUITY] = init_cap; /* prev_equity */
ps[PS_HOLD_TIME] = 0.0f; /* hold_time */
ps[PS_REALIZED_PNL] = 0.0f; /* realized_pnl */
ps[PS_ENTRY_PRICE] = 0.0f; /* entry_price */
ps[PS_TRADE_START_PNL] = 0.0f; /* trade_start_pnl */
ps[PS_KELLY_WIN_COUNT] = 0.0f; /* win/loss/Kelly counters */
ps[PS_KELLY_LOSS_COUNT] = 0.0f;
ps[PS_KELLY_SUM_WINS] = 0.0f;
ps[PS_KELLY_SUM_LOSSES] = 0.0f;
ps[PS_KELLY_SUM_RETURNS] = 0.0f;
ps[PS_KELLY_SUM_SQ_RETURNS] = 0.0f;
ps[PS_INTRA_TRADE_MAX_DD] = 0.0f; /* intra_trade_max_dd */
ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f; /* intra_trade_max_pnl */
ps[PS_PLAN_TARGET_BARS] = 0.0f; ps[PS_PLAN_PROFIT_TARGET] = 0.0f; ps[PS_PLAN_STOP_LOSS] = 0.0f;
ps[PS_PLAN_SCALE_AGGRESSION] = 0.0f; ps[PS_PLAN_CONVICTION] = 0.0f; ps[PS_PLAN_ASYMMETRY] = 0.0f;
ps[PS_PLAN_ENTRY_REGIME] = 0.0f; /* plan slots zeroed */
current_timesteps[i] = 0;
} else {
/* Soft reset (trade_complete): keep equity, clear trade state only.
* The model starts next segment with accumulated gains/losses.
* Position is already 0 (model chose Flat). Cash, equity,
* win/loss counters, DSR state — all preserved. */
ps[10] = 0.0f; /* hold_time — new trade starts fresh */
ps[12] = 0.0f; /* entry_price — no active trade */
ps[13] = (ps[11]); /* trade_start_pnl = current realized_pnl */
ps[20] = 0.0f; /* intra_trade_max_dd — reset for next trade */
ps[21] = 0.0f; /* intra_trade_max_pnl — reset for next trade */
ps[23] = 0.0f; ps[24] = 0.0f; ps[25] = 0.0f;
ps[26] = 0.0f; ps[27] = 0.0f; ps[28] = 0.0f;
ps[29] = 0.0f; /* plan slots zeroed on trade complete */
ps[PS_HOLD_TIME] = 0.0f; /* hold_time — new trade starts fresh */
ps[PS_ENTRY_PRICE] = 0.0f; /* entry_price — no active trade */
ps[PS_TRADE_START_PNL] = (ps[PS_REALIZED_PNL]); /* trade_start_pnl = current realized_pnl */
ps[PS_INTRA_TRADE_MAX_DD] = 0.0f; /* intra_trade_max_dd — reset for next trade */
ps[PS_INTRA_TRADE_MAX_PNL] = 0.0f; /* intra_trade_max_pnl — reset for next trade */
ps[PS_PLAN_TARGET_BARS] = 0.0f; ps[PS_PLAN_PROFIT_TARGET] = 0.0f; ps[PS_PLAN_STOP_LOSS] = 0.0f;
ps[PS_PLAN_SCALE_AGGRESSION] = 0.0f; ps[PS_PLAN_CONVICTION] = 0.0f; ps[PS_PLAN_ASYMMETRY] = 0.0f;
ps[PS_PLAN_ENTRY_REGIME] = 0.0f; /* plan slots zeroed on trade complete */
current_timesteps[i] = 0;
}
}

View File

@@ -24,6 +24,50 @@
// ── cuBLAS alignment ──
#define SL_STATE_DIM_PADDED 128
// ────────────────────────────────────────────────────────────────────────────
// Portfolio state slot indices (ps[0..PORTFOLIO_STRIDE=38)).
// Invariant 8: every slot has a named constant. Raw ps[N] access outside the
// definition site in experience_kernels.cu is a lint violation.
//
// Slots 3-6 and 22 also carry named aliases defined locally in
// experience_kernels.cu (DSR_A_SLOT, DSR_B_SLOT, DSR_TRADE_COUNT_SLOT,
// PREV_CLOSE_SLOT, PREV_MID_SLOT). The PS_* names below are canonical for
// cross-file reference; the local aliases remain for in-file use.
// ────────────────────────────────────────────────────────────────────────────
#define PS_POSITION 0 // current contract position (signed)
#define PS_CASH 1 // cash balance
#define PS_PORTFOLIO_VALUE 2 // mark-to-market total (cash + pos*price)
#define PS_DSR_A 3 // EMA of realised trade returns (DSR)
#define PS_DSR_B 4 // EMA of squared trade returns (DSR)
#define PS_DSR_TRADE_COUNT 5 // completed trades counter (DSR warmup)
#define PS_PREV_CLOSE 6 // prev bar raw_close for DSR per-bar returns
#define PS_PEAK_EQUITY 7 // high-water mark (init to initial_capital)
#define PS_FLAT_COUNTER 8 // consecutive flat steps
#define PS_PREV_EQUITY 9 // equity at previous step
#define PS_HOLD_TIME 10 // consecutive steps with position
#define PS_REALIZED_PNL 11 // cumulative realised PnL
#define PS_ENTRY_PRICE 12 // price when trade was entered
#define PS_TRADE_START_PNL 13 // realized_pnl snapshot at trade entry
#define PS_KELLY_WIN_COUNT 14 // Kelly: number of profitable trade exits
#define PS_KELLY_LOSS_COUNT 15 // Kelly: number of losing trade exits
#define PS_KELLY_SUM_WINS 16 // Kelly: cumulative profit from winners
#define PS_KELLY_SUM_LOSSES 17 // Kelly: cumulative |loss| from losers
#define PS_KELLY_SUM_RETURNS 18 // Kelly: cumulative net returns (for mu)
#define PS_KELLY_SUM_SQ_RETURNS 19 // Kelly: cumulative squared returns (sigma^2)
#define PS_INTRA_TRADE_MAX_DD 20 // worst unrealised drawdown during trade (v8)
#define PS_INTRA_TRADE_MAX_PNL 21 // best unrealised P&L during trade (hindsight)
#define PS_PREV_MID 22 // prev bar MBP-10 mid-price
#define PS_PLAN_TARGET_BARS 23 // plan: max hold bars (>0.5 = plan active)
#define PS_PLAN_PROFIT_TARGET 24 // plan: raw profit threshold
#define PS_PLAN_STOP_LOSS 25 // plan: raw stop threshold
#define PS_PLAN_SCALE_AGGRESSION 26 // plan: position ramp speed
#define PS_PLAN_CONVICTION 27 // plan: conviction at entry [0,1]
#define PS_PLAN_ASYMMETRY 28 // plan: profit/stop asymmetry
#define PS_PLAN_ENTRY_REGIME 29 // plan: regime_stability at entry
// Slots 30-37: prev-bar OFI scratch for delta computation (see experience_kernels.cu)
#define PS_OFI_PREV_BASE 30 // first OFI prev-bar slot (stride 8: [30..37])
#define PS_STRIDE 38 // total portfolio state stride (== PORTFOLIO_STRIDE)
// ── Compile-time checks ──
static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM,
"State layout dimensions must sum to SL_STATE_DIM");

View File

@@ -10,6 +10,7 @@
*/
#define PORTFOLIO_STRIDE 38
#include "state_layout.cuh"
extern "C" __global__ void trade_stats_reduce(
const float* __restrict__ portfolio_states, // [N * PORTFOLIO_STRIDE]
@@ -34,7 +35,7 @@ extern "C" __global__ void trade_stats_reduce(
for (int i = tid; i < N; i += stride) {
int base = i * PORTFOLIO_STRIDE;
for (int k = 0; k < 6; k++) {
local_sums[k] = local_sums[k] + portfolio_states[base + 14 + k];
local_sums[k] = local_sums[k] + portfolio_states[base + PS_KELLY_WIN_COUNT + k];
}
}

View File

@@ -21,42 +21,48 @@ Defined in `crates/ml/src/cuda_pipeline/state_layout.cuh` and mirrored in Rust v
| `SL_PORTFOLIO_START` | 90 | |
| `SL_PLAN_ISV_START` | 98 | |
## Portfolio state slots (ps[0..30))
## Portfolio state slots (ps[0..PS_STRIDE=38))
Defined in `crates/ml/src/cuda_pipeline/experience_kernels.cu` header comments. Migrated to named constants during Task 4 of this plan.
Defined in `crates/ml/src/cuda_pipeline/state_layout.cuh` (PS_* constants).
Consumer code uses the named constants; raw `ps[N]` access outside definition
sites is a lint violation (Invariant 8). Slots 3-6 and 22 also carry legacy
named aliases in `experience_kernels.cu` (DSR_A_SLOT etc.) retained for
in-file use; PS_* names are the canonical cross-file form.
| Index | Constant | Meaning |
|---|---|---|
| 0 | `PS_VALUE` | Equity |
| 1 | `PS_POSITION` | Current position |
| 2 | `PS_CASH` | Available cash |
| 3 | `PS_ENTRY_PRICE` | Entry price |
| 4 | `PS_MAX_EQUITY` | Peak equity (for drawdown) |
| 5 | `PS_HOLD_TIME` | Bars held |
| 6 | `PS_CUM_RETURN` | Cumulative return |
| 7 | `PS_STEP_COUNT` | Bars in window |
| 8 | `PS_SPREAD_COST` | Current spread cost |
| 0 | `PS_POSITION` | Current contract position (signed) |
| 1 | `PS_CASH` | Cash balance |
| 2 | `PS_PORTFOLIO_VALUE` | Mark-to-market total (cash + pos×price) |
| 3 | `PS_DSR_A` | EMA of realised trade returns (DSR) |
| 4 | `PS_DSR_B` | EMA of squared trade returns (DSR) |
| 5 | `PS_DSR_TRADE_COUNT` | Completed trades counter (DSR warmup) |
| 6 | `PS_PREV_CLOSE` | Prev bar raw_close for DSR per-bar returns |
| 7 | `PS_PEAK_EQUITY` | High-water mark (init to initial_capital) |
| 8 | `PS_FLAT_COUNTER` | Consecutive flat steps |
| 9 | `PS_PREV_EQUITY` | Equity at previous step |
| 10 | `PS_HOLD_TIME_TOTAL` | Total hold time |
| 10 | `PS_HOLD_TIME` | Consecutive steps with position |
| 11 | `PS_REALIZED_PNL` | Cumulative realised PnL |
| 12 | `PS_WIN_COUNT_UNUSED` | (slot retained, not currently used) |
| 13 | `PS_LOSS_COUNT_UNUSED` | (slot retained, not currently used) |
| 14 | `PS_KELLY_WIN_COUNT` | Kelly stats: winning trades |
| 15 | `PS_KELLY_LOSS_COUNT` | Kelly stats: losing trades |
| 16 | `PS_KELLY_SUM_WINS` | Kelly stats: sum of winning returns |
| 17 | `PS_KELLY_SUM_LOSSES` | Kelly stats: sum of \|losing returns\| |
| 18 | `PS_RESERVED_18` | reserved |
| 19 | `PS_RESERVED_19` | reserved |
| 20 | `PS_INTRA_TRADE_MAX_DD` | worst intra-trade drawdown (v8 reward) |
| 21 | `PS_INTRA_TRADE_MAX_PNL` | best intra-trade unrealised PnL (hindsight) |
| 22 | `PS_RESERVED_22` | (was interval_sum, slot retained) |
| 12 | `PS_ENTRY_PRICE` | Price when trade was entered |
| 13 | `PS_TRADE_START_PNL` | Realised PnL snapshot at trade entry |
| 14 | `PS_KELLY_WIN_COUNT` | Kelly: number of profitable trade exits |
| 15 | `PS_KELLY_LOSS_COUNT` | Kelly: number of losing trade exits |
| 16 | `PS_KELLY_SUM_WINS` | Kelly: cumulative profit from winners |
| 17 | `PS_KELLY_SUM_LOSSES` | Kelly: cumulative \|loss\| from losers |
| 18 | `PS_KELLY_SUM_RETURNS` | Kelly: cumulative net returns (for mu) |
| 19 | `PS_KELLY_SUM_SQ_RETURNS` | Kelly: cumulative squared returns (sigma^2) |
| 20 | `PS_INTRA_TRADE_MAX_DD` | Worst intra-trade drawdown (v8 reward) |
| 21 | `PS_INTRA_TRADE_MAX_PNL` | Best intra-trade unrealised PnL (hindsight) |
| 22 | `PS_PREV_MID` | Prev bar MBP-10 mid-price |
| 23 | `PS_PLAN_TARGET_BARS` | plan: max hold bars (>0.5 = plan active) |
| 24 | `PS_PLAN_PROFIT_TARGET` | plan: raw profit threshold |
| 25 | `PS_PLAN_STOP_LOSS` | plan: raw stop threshold |
| 26 | `PS_PLAN_SCALE_AGGRESSION` | plan: position ramp speed |
| 27 | `PS_PLAN_CONVICTION` | plan: conviction at entry [0,1] |
| 27 | `PS_PLAN_CONVICTION` | plan: conviction at entry [0,1] |
| 28 | `PS_PLAN_ASYMMETRY` | plan: profit/stop asymmetry |
| 29 | `PS_PLAN_ENTRY_REGIME` | plan: regime_stability at entry (for P12) |
| 29 | `PS_PLAN_ENTRY_REGIME` | plan: regime_stability at entry |
| 30-37 | `PS_OFI_PREV_BASE + k` | Prev-bar OFI scratch for delta computation |
| 38 | `PS_STRIDE` | Total portfolio state stride (=PORTFOLIO_STRIDE) |
## Plan_isv dimensions (plan_isv[0..6))
@@ -109,3 +115,7 @@ From the trade_plan MLP output.
| 0 | `MAG_QUARTER` | 0.25× max_position |
| 1 | `MAG_HALF` | 0.50× max_position |
| 2 | `MAG_FULL` | 1.00× max_position |
## Commit history
- Task 4A (ps[0..PS_STRIDE) constants): commit <SHA>