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

refactor: experience_state_gather uses assemble_state() — fixes OFI/MTF collision

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-20 14:49:03 +02:00
parent 62650aa53b
commit 65ad9debbb
2 changed files with 109 additions and 112 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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);
}
/* ================================================================== */

4
crates/ml/src/cuda_pipeline/state_layout.cuh
View File

@@ -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 ──