From 6e0cb3d62ccd0eb92c7aaa1f7f4ab23971d23247 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 20 Apr 2026 17:30:25 +0200 Subject: [PATCH] =?UTF-8?q?fix:=20CRITICAL=20=E2=80=94=20dense=20micro-rew?= =?UTF-8?q?ard=20+=20OFI=20embed=20read=20stale=20state=20positions?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Task 3 refactored experience_state_gather (the WRITER) to use assemble_state() with canonical layout (OFI at [42..62)). But env_step and ofi_embed_build_input (the READERS) still hardcoded OFI at [66..84) — the OLD pre-refactor layout. This meant 7 locations were reading MTF/portfolio features as if they were OFI: 1. Line 1745: dense micro-reward ofi_cur = state+66 → actually MTF[4] 2. Line 1778: book_aggression = state[82] → actually plan_isv region 3. Line 1983: ps[30..37] OFI delta storage for NEXT bar — storing MTF data 4. Lines 5833/5836/5838/5840: ofi_embed_build_input — feeds 18→10 MLP into Mamba2 temporal SSM and attention. Entire temporal pipeline was training on MTF features dressed as OFI. Symptoms explained: - WinRate=20.9% on validation (anti-correlated): dense micro-reward computes quality=sign_pos × garbage_MTF_deltas, systematically rewarding wrong direction - mean_reward=+0.004 but Sharpe_raw=-0.0004: shaped reward exploits garbage signal, real portfolio loses money - grad_norm=23560 at epoch 2: gradients chasing noise - Q-value explosion to ±10 in one epoch: learning contradictions Fix: replaced all hardcoded 66/74/82/83 with SL_OFI_START from state_layout.cuh. Both reader kernels now use the same canonical layout as assemble_state(). Verified locally: smoke test passes, OFI_DIAG shows correct non-zero values (raw_mean=-0.36, delta_mean=-0.21, log_dur=-0.23). Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 28 +++++++++++-------- 1 file changed, 16 insertions(+), 12 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 310b8f350..df6a5c572 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1738,11 +1738,11 @@ extern "C" __global__ void experience_env_step( * C51 atoms now model distribution of TRADE OUTCOMES, not per-bar noise. * ──────────────────────────────────────────────────────────────────── */ if (!segment_complete && fabsf(position) > 0.001f && micro_reward_scale > 0.0f - && full_batch_states != NULL && full_state_dim >= 74) { + && full_batch_states != NULL && full_state_dim >= SL_OFI_START + SL_OFI_DIM) { /* ── Dense micro-reward: OFI momentum + MBP-10 price confirmation ── * Replaces flat holding cost with informed per-bar signal using * order flow imbalance deltas and mid-price mark-to-market. */ - const float* ofi_cur = full_batch_states + (long long)i * full_state_dim + 66; + const float* ofi_cur = full_batch_states + (long long)i * full_state_dim + SL_OFI_START; float sign_pos = (position > 0.0f) ? 1.0f : -1.0f; /* OFI momentum: delta from previous bar (stored in ps[30..37]) */ @@ -1774,8 +1774,11 @@ extern "C" __global__ void experience_env_step( float cost_tolerance = fmaxf(capital_ratio * fmaxf(sharpe_ema, 0.0f), 0.1f); float spread_penalty = spread_at_entry / (fmaxf(atr_abs, 1e-6f) * cost_tolerance); - /* Book aggression from state vector (precomputed from 10-level MBP-10) */ - float book_aggression = (full_state_dim >= 83) ? full_batch_states[(long long)i * full_state_dim + 82] : 0.0f; + /* Book aggression from state vector (precomputed from 10-level MBP-10). + * Located at OFI[16] = SL_OFI_START + 16 in canonical layout. */ + float book_aggression = (full_state_dim > SL_OFI_START + 16) + ? full_batch_states[(long long)i * full_state_dim + SL_OFI_START + 16] + : 0.0f; /* Informed flow intensity */ float vol_informed = ofi_cur[6] * ofi_cur[7]; /* VPIN x Kyle's_lambda */ @@ -1979,8 +1982,8 @@ extern "C" __global__ void experience_env_step( /* ---- 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. */ - if (full_batch_states != NULL && full_state_dim >= 74) { - const float* ofi_cur = full_batch_states + (long long)i * full_state_dim + 66; + if (full_batch_states != NULL && full_state_dim >= SL_OFI_START + SL_OFI_DIM) { + 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]; @@ -5815,8 +5818,9 @@ extern "C" __global__ void plan_noise_inject( /* ── OFI embedding MLP input construction ───────────────────────────────── * * ofi_embed_build_input — Extract 18-dim OFI input from states_buf. - * raw OFI at state[66..74), delta OFI at state[74..82), - * book_aggression at state[82], log_bar_duration at state[83]. + * Canonical layout via state_layout.cuh: + * raw OFI at state[SL_OFI_START..SL_OFI_START+8), delta at [+8..+16), + * book_aggression at [+16], log_bar_duration at [+17]. * Output: row-major [B, 18] for cuBLAS sgemm_f32 (row-major trick). * Grid: ceil(B/256), Block: 256. */ @@ -5830,14 +5834,14 @@ extern "C" __global__ void ofi_embed_build_input( const float* s = states + (long long)b * state_dim; /* Raw OFI [0..8) */ for (int k = 0; k < 8; k++) - output[b * 18 + k] = (state_dim > 66 + k) ? s[66 + k] : 0.0f; + output[b * 18 + k] = (state_dim > SL_OFI_START + k) ? s[SL_OFI_START + k] : 0.0f; /* Delta OFI [8..16) */ for (int k = 0; k < 8; k++) - output[b * 18 + 8 + k] = (state_dim > 74 + k) ? s[74 + k] : 0.0f; + output[b * 18 + 8 + k] = (state_dim > SL_OFI_START + 8 + k) ? s[SL_OFI_START + 8 + k] : 0.0f; /* Book aggression [16] */ - output[b * 18 + 16] = (state_dim > 82) ? s[82] : 0.0f; + output[b * 18 + 16] = (state_dim > SL_OFI_START + 16) ? s[SL_OFI_START + 16] : 0.0f; /* Log bar duration [17] */ - output[b * 18 + 17] = (state_dim > 83) ? s[83] : 0.0f; + output[b * 18 + 17] = (state_dim > SL_OFI_START + 17) ? s[SL_OFI_START + 17] : 0.0f; } /* ================================================================== */