From 65ad9debbb48b0e9ffa850d3472ae8fe874a55a1 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 20 Apr 2026 14:49:03 +0200 Subject: [PATCH] =?UTF-8?q?refactor:=20experience=5Fstate=5Fgather=20uses?= =?UTF-8?q?=20assemble=5Fstate()=20=E2=80=94=20fixes=20OFI/MTF=20collision?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The kernel now writes to local arrays (market[], portfolio[], plan_isv[], mtf[], ofi[]) then calls assemble_state() from state_layout.cuh to produce the canonical layout. This eliminates the hardcoded ofi_start=66 that collided with MTF features and ensures training/validation use identical state vectors. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 217 +++++++++--------- crates/ml/src/cuda_pipeline/state_layout.cuh | 4 +- 2 files changed, 109 insertions(+), 112 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index f5e38bf86..9bd60f2c8 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -84,6 +84,10 @@ __device__ __forceinline__ float asymmetric_soft_clamp(float x) { * capital floor). Single source of truth — also used by backtest_env_kernel. */ #include "trade_physics.cuh" +/* Canonical state layout — single source of truth for state assembly. + * Used by both experience_state_gather and backtest_state_gather. */ +#include "state_layout.cuh" + /* ------------------------------------------------------------------ */ /* BF16 support — common_device_functions.cuh is prepended by build.rs */ /* ------------------------------------------------------------------ */ @@ -423,23 +427,25 @@ extern "C" __global__ void experience_state_gather( float* out = batch_states + (long long)i * state_dim; - /* Zero the entire state vector first — any feature slots not explicitly - * written below stay zero instead of containing stale GPU memory. - * Without this, features beyond market_dim (portfolio, MTF, OFI) are - * uninitialized when state_dim > market_dim but those slots aren't - * reached by the feature-writing code below. */ - for (int k = 0; k < state_dim; k++) - out[k] = 0.0f; - - /* Out-of-data or negative index: state is already zeroed above. */ + /* Out-of-data or negative index: zero the output and return early. */ if (bar_idx < 0 || bar_idx >= total_bars) { + for (int k = 0; k < state_dim; k++) + out[k] = 0.0f; return; } - /* -- Market features: [0 .. market_dim) -- bf16 input → f32 output */ + /* ── Local arrays for canonical state assembly ── + * Each feature group is written to a local array, then assemble_state() + * places them at the correct offsets — single source of truth. */ + + /* -- Market features → local market[] array -- */ + float market[SL_MARKET_DIM]; const float* mf_row = market_features + (long long)bar_idx * market_dim; - for (int k = 0; k < market_dim; k++) - out[k] = (float)mf_row[k]; + for (int k = 0; k < market_dim && k < SL_MARKET_DIM; k++) + market[k] = (float)mf_row[k]; + /* Zero remaining if market_dim < SL_MARKET_DIM */ + for (int k = market_dim; k < SL_MARKET_DIM; k++) + market[k] = 0.0f; /* ── #10 Mirror universe: negate ALL directional features ── * Features that encode direction (returns, momentum, slopes) must be @@ -457,40 +463,39 @@ extern "C" __global__ void experience_state_gather( * [5] RSI: mirror = 1.0 - RSI */ if (mirror_active && market_dim >= 42) { /* OHLCV returns */ - out[0] = -out[0]; - out[1] = -out[1]; - out[2] = -out[2]; - out[3] = -out[3]; + market[0] = -market[0]; + market[1] = -market[1]; + market[2] = -market[2]; + market[3] = -market[3]; /* RSI: flip around 0.5 */ - out[5] = 1.0f - out[5]; + market[5] = 1.0f - market[5]; /* MACD + Bollinger position */ - out[6] = -out[6]; - out[7] = -out[7]; + market[6] = -market[6]; + market[7] = -market[7]; /* Price patterns: returns + SMA ratios + regression slope */ - out[10] = -out[10]; - out[11] = -out[11]; - out[12] = -out[12]; - out[13] = -out[13]; - out[14] = -out[14]; - out[15] = -out[15]; + market[10] = -market[10]; + market[11] = -market[11]; + market[12] = -market[12]; + market[13] = -market[13]; + market[14] = -market[14]; + market[15] = -market[15]; /* CUSUM direction */ - out[41] = -out[41]; + market[41] = -market[41]; } /* ── #13 Vol normalization: scale return features [0..3] by 1/vol ── */ if (vol_normalizer > 0.0f && market_dim >= 4) { float inv_vol = 1.0f / vol_normalizer; for (int k = 0; k < 4; k++) { - float v = (float)out[k] * inv_vol; - out[k] = v; + market[k] = market[k] * inv_vol; } } /* ── #23 Causal feature masking: zero out masked features ── */ if (feature_mask != NULL) { - for (int k = 0; k < market_dim; k++) { + for (int k = 0; k < market_dim && k < SL_MARKET_DIM; k++) { if (feature_mask[k] < 0.5f) { - out[k] = 0.0f; + market[k] = 0.0f; } } } @@ -500,34 +505,29 @@ extern "C" __global__ void experience_state_gather( * Deterministic regardless of execution history. */ if (feature_noise_scale > 0.0f) { int timestep_val = raw_t; - for (int k = 0; k < market_dim; k++) { + for (int k = 0; k < market_dim && k < SL_MARKET_DIM; k++) { /* Two Philox hashes → Box-Muller Gaussian */ float u1 = philox_uniform(i, timestep_val, 1000 + k * 2); float u2 = philox_uniform(i, timestep_val, 1000 + k * 2 + 1); u1 = fmaxf(u1, 1e-6f); float noise = sqrtf(-2.0f * logf(u1)) * cosf(6.2831853f * u2); - float v = (float)out[k] + noise * feature_noise_scale; - out[k] = v; + market[k] = market[k] + noise * feature_noise_scale; } } - /* -- Portfolio features: [market_dim .. market_dim+8) -- + /* -- Portfolio features → local portfolio[] and plan_isv[] arrays -- * * 8 informative features that let the Q-network SEE its trading context. * Without these, the model is blind to risk, P&L, and position state. * - * slot +0: position / max_position — normalized position [-1, +1] - * slot +1: unrealized_pnl / equity — how is THIS trade doing? - * slot +2: drawdown — (peak - equity) / peak [0, 1] - * slot +3: hold_time / 100.0 — normalized holding duration - * slot +4: realized_pnl / equity — overall session P&L - * slot +5: distance_to_floor — (equity - floor) / equity [0, 1] - * slot +6: trade_return — trade P&L / equity - * slot +7: cash_ratio — cash / equity - * - * max_position is read from kernel arg (line ~370). We use a constant - * approximation here (max_pos not available in gather kernel). - * The network learns the actual scale via the position feature. + * portfolio[0]: position — raw position (signed contracts) + * portfolio[1]: unrealized_pnl / equity — how is THIS trade doing? + * portfolio[2]: drawdown — (peak - equity) / peak [0, 1] + * portfolio[3]: hold_time / 100.0 — normalized holding duration + * portfolio[4]: realized_pnl / equity — overall session P&L + * portfolio[5]: distance_to_floor — (equity - floor) / equity [0, 1] + * portfolio[6]: trade_return — trade P&L / equity + * portfolio[7]: cash_ratio — cash / equity */ const float* ps = portfolio_states + (long long)i * PORTFOLIO_STRIDE; float position = ps[0]; @@ -577,49 +577,50 @@ extern "C" __global__ void experience_state_gather( ? (f_equity - f_floor_val) / f_equity : 0.0f; - int portfolio_base = market_dim; - if (portfolio_base + 13 < state_dim) { - out[portfolio_base + 0] = f_position; /* raw position */ - out[portfolio_base + 1] = f_unrealized_pnl / f_equity; /* trade P&L signal */ - out[portfolio_base + 2] = f_drawdown; /* risk: how deep are we? */ - out[portfolio_base + 3] = f_hold_time / 100.0f; /* how long in trade? */ - out[portfolio_base + 4] = f_realized_pnl / f_equity; /* session P&L */ - out[portfolio_base + 5] = f_floor_dist; /* distance to game over */ - out[portfolio_base + 6] = f_trade_return; /* this trade's return */ - out[portfolio_base + 7] = f_cash / f_equity; /* available capital */ - /* Plan-aware features: model sees its own plan progress. - * Enables temporal reasoning: "I'm 80% through my plan, close to target." */ - float plan_tgt_bars = ps[23]; - float plan_profit = ps[24]; - float plan_stop = ps[25]; - float plan_conv = ps[27]; - out[portfolio_base + 8] = (plan_tgt_bars > 0.5f) - ? fminf(f_hold_time / plan_tgt_bars, 2.0f) /* plan progress [0, 2] */ - : 0.0f; - out[portfolio_base + 9] = (plan_profit > 1e-6f) - ? fminf(f_unrealized_pnl / (plan_profit * f_equity + 1e-6f), 2.0f) /* pnl vs target */ - : 0.0f; - out[portfolio_base + 10] = (plan_stop > 1e-6f) - ? fminf(-f_unrealized_pnl / (plan_stop * f_equity + 1e-6f), 2.0f) /* pnl vs stop */ - : 0.0f; - out[portfolio_base + 11] = plan_conv; /* conviction at entry [0, 1] */ + float portfolio[SL_PORTFOLIO_BASE_DIM]; + portfolio[0] = f_position; /* raw position */ + portfolio[1] = f_unrealized_pnl / f_equity; /* trade P&L signal */ + portfolio[2] = f_drawdown; /* risk: how deep are we? */ + portfolio[3] = f_hold_time / 100.0f; /* how long in trade? */ + portfolio[4] = f_realized_pnl / f_equity; /* session P&L */ + portfolio[5] = f_floor_dist; /* distance to game over */ + portfolio[6] = f_trade_return; /* this trade's return */ + portfolio[7] = f_cash / f_equity; /* available capital */ - /* P11: 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). */ - out[portfolio_base + 12] = (plan_params_ptr != NULL && ps[23] > 0.5f) - ? (plan_params_ptr[i * 6 + 4] - ps[27]) /* conviction drift */ - : 0.0f; + /* 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]; - /* P12: 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. */ - out[portfolio_base + 13] = (isv_signals_ptr != NULL && ps[23] > 0.5f) - ? fabsf(isv_signals_ptr[11] - ps[29]) /* regime shift */ - : 0.0f; - } + plan_isv[0] = (plan_tgt_bars > 0.5f) + ? fminf(f_hold_time / plan_tgt_bars, 2.0f) /* plan progress [0, 2] */ + : 0.0f; + plan_isv[1] = (plan_profit > 1e-6f) + ? fminf(f_unrealized_pnl / (plan_profit * f_equity + 1e-6f), 2.0f) /* pnl vs target */ + : 0.0f; + plan_isv[2] = (plan_stop > 1e-6f) + ? fminf(-f_unrealized_pnl / (plan_stop * f_equity + 1e-6f), 2.0f) /* pnl vs stop */ + : 0.0f; + plan_isv[3] = plan_conv; /* conviction at entry [0, 1] */ - /* -- Multi-timeframe features: [market_dim+8 .. market_dim+8+16) -- + /* 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 */ + : 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 */ + : 0.0f; + + /* -- Multi-timeframe features → local mtf[] array -- * * 4 lookback windows x 4 features = 16 GPU-native features. * Computed directly from the market data already on GPU — no CPU. @@ -632,14 +633,15 @@ extern "C" __global__ void experience_state_gather( * +2: volume_trend = volume_now / avg_volume_N (>1 = increasing) * +3: momentum = close_now vs midpoint of N-bar range (0=bottom, 1=top) */ + float mtf[SL_MTF_DIM]; + for (int k = 0; k < SL_MTF_DIM; k++) mtf[k] = 0.0f; + const int lookbacks[4] = {5, 15, 60, 240}; - int mtf_base = market_dim + 14; /* after 14 portfolio features (8 base + 6 plan) */ for (int lb = 0; lb < 4; lb++) { int Nlb = lookbacks[lb]; int past_idx = bar_idx - Nlb; - int slot = mtf_base + lb * 4; - if (past_idx >= 0 && slot + 3 < state_dim) { + if (past_idx >= 0) { const float* now_row = market_features + (long long)bar_idx * market_dim; const float* past_row = market_features + (long long)past_idx * market_dim; /* Float arithmetic — bf16 close prices (~5000) subtracted produce @@ -650,7 +652,7 @@ extern "C" __global__ void experience_state_gather( /* Return over N bars */ float f_ret = (f_close_past > 0.0f) ? (f_close_now - f_close_past) / f_close_past : 0.0f; float f_scaled_ret = f_ret * 100.0f; - out[slot + 0] = fminf(10.0f, fmaxf(-10.0f, f_scaled_ret)); + mtf[lb * 4 + 0] = fminf(10.0f, fmaxf(-10.0f, f_scaled_ret)); /* Volatility: scan high/low over window */ float f_max_val = f_close_now; @@ -667,40 +669,35 @@ extern "C" __global__ void experience_state_gather( } } float f_range = (f_close_now > 0.0f) ? (f_max_val - f_min_val) / f_close_now : 0.0f; - out[slot + 1] = fminf(10.0f, fmaxf(0.0f, f_range * 100.0f)); + mtf[lb * 4 + 1] = fminf(10.0f, fmaxf(0.0f, f_range * 100.0f)); /* Volume trend */ float f_avg_vol = (vol_count > 0) ? f_vol_sum / (float)vol_count : 1.0f; float f_cur_vol = (market_dim > 4) ? (float)now_row[4] : 1.0f; float f_vol_ratio = (f_avg_vol > 0.0f) ? f_cur_vol / f_avg_vol : 1.0f; - out[slot + 2] = fminf(5.0f, fmaxf(0.0f, f_vol_ratio)); + mtf[lb * 4 + 2] = fminf(5.0f, fmaxf(0.0f, f_vol_ratio)); /* Momentum: position within range [0=bottom, 1=top] */ float f_range_size = f_max_val - f_min_val; - out[slot + 3] = (f_range_size > 0.0f) - ? (f_close_now - f_min_val) / f_range_size - : 0.5f; - } else { - /* Not enough history — zero pad */ - if (slot + 3 < state_dim) { - out[slot + 0] = 0.0f; - out[slot + 1] = 0.0f; - out[slot + 2] = 0.0f; - out[slot + 3] = 0.0f; - } + mtf[lb * 4 + 3] = (f_range_size > 0.0f) + ? (f_close_now - f_min_val) / f_range_size + : 0.5f; } + /* else: mtf[lb*4..lb*4+3] already zeroed above */ } - /* -- OFI features: raw(8) + deltas(8) + book_aggression(1) + log_duration(1) -- */ - /* Written at state[66..84). ofi_features layout: [total_bars, ofi_dim] where - * ofi_dim=20: [raw_ofi(8), deltas(8), book_agg(1), log_dur(1), accel(1), toxicity(1)] */ + /* -- OFI features → local ofi[] array -- */ + float ofi[SL_OFI_DIM]; + for (int k = 0; k < SL_OFI_DIM; k++) ofi[k] = 0.0f; if (ofi_features != NULL && ofi_dim > 0 && bar_idx >= 0 && bar_idx < total_bars) { const float* ofi_row = ofi_features + (long long)bar_idx * ofi_dim; - int ofi_start = 66; /* OFI starts at state index 66 */ - int ofi_copy = (ofi_dim < 18) ? ofi_dim : 18; /* copy up to 18 features */ - for (int k = 0; k < ofi_copy && (ofi_start + k) < state_dim; k++) - out[ofi_start + k] = ofi_row[k]; + int ofi_copy = (ofi_dim < SL_OFI_DIM) ? ofi_dim : SL_OFI_DIM; + for (int k = 0; k < ofi_copy; k++) + ofi[k] = ofi_row[k]; } + + /* ── Final assembly: canonical layout via state_layout.cuh ── */ + assemble_state(out, market, ofi, mtf, portfolio, plan_isv); } /* ================================================================== */ diff --git a/crates/ml/src/cuda_pipeline/state_layout.cuh b/crates/ml/src/cuda_pipeline/state_layout.cuh index 2d98868bb..334b46a91 100644 --- a/crates/ml/src/cuda_pipeline/state_layout.cuh +++ b/crates/ml/src/cuda_pipeline/state_layout.cuh @@ -25,9 +25,9 @@ #define SL_STATE_DIM_PADDED 128 // ── Compile-time checks ── -_Static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, +static_assert(SL_PADDING_START + SL_PADDING_DIM == SL_STATE_DIM, "State layout dimensions must sum to SL_STATE_DIM"); -_Static_assert(SL_STATE_DIM % 8 == 0, +static_assert(SL_STATE_DIM % 8 == 0, "SL_STATE_DIM must be 8-aligned for tensor core cuBLAS"); // ── Shared state assembly function ──