From 7a2adeb37483a40ad9ceff6e54bc4367eed4a4c6 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 24 Apr 2026 11:19:25 +0200 Subject: [PATCH] =?UTF-8?q?refactor(dqn-v2):=20Invariant=208=20=E2=80=94?= =?UTF-8?q?=20named=20constants=20for=20portfolio=20state=20ps[0..PS=5FSTR?= =?UTF-8?q?IDE=3D38)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- .../src/cuda_pipeline/experience_kernels.cu | 266 +++++++++--------- crates/ml/src/cuda_pipeline/state_layout.cuh | 44 +++ .../src/cuda_pipeline/trade_stats_kernel.cu | 3 +- docs/dqn-named-dims.md | 60 ++-- 4 files changed, 214 insertions(+), 159 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 3c89aaf4a..8088c5d21 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -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; } } diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index b4e1423b6..b240d618d 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -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"); diff --git a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu index 85a0f4026..1f3c0fa69 100644 --- a/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu +++ b/crates/ml/src/cuda_pipeline/trade_stats_kernel.cu @@ -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]; } } diff --git a/docs/dqn-named-dims.md b/docs/dqn-named-dims.md index 828c1ea7d..244d33d76 100644 --- a/docs/dqn-named-dims.md +++ b/docs/dqn-named-dims.md @@ -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