fix: CRITICAL — dense micro-reward + OFI embed read stale state positions

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) <noreply@anthropic.com>

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