diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 5667adcda..0c8fd95d5 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2156,7 +2156,27 @@ extern "C" __global__ void experience_env_step( * between on-policy and CF: if on-policy at (i,t) was an * early-return, the CF slot at (i,t) is also skipped. NULL-tolerant * for any test scaffold that hasn't wired the buffer yet. */ - int* __restrict__ slot_completed_normally_out + int* __restrict__ slot_completed_normally_out, + /* SP20 Phase 2 Task 2.0 (2026-05-09) — aux next-bar label tile + * (input) + per-env trade-open label-sign scratch (output). Both + * `[N]`-sized i32 device buffers. NULL-tolerant for test scaffolds + * without aux-head wiring (write site is gated on both being + * non-NULL). + * + * `aux_label_per_env` ← `exp_aux_nb_label_buf`: per-env aux next- + * bar label written by `aux_sign_label_per_step_kernel` BEFORE this + * kernel runs (same stream, stream-implicit ordering). Values: + * `1` = "up", `0` = "not-up", `-1` = skip-window (lookahead past + * total_bars). + * + * `label_at_open_per_env`: this kernel's `entering_trade` branch + * maps the aux label to the SP20 sign convention (`{-1, 0, +1}`) + * and writes it here. Slot persists across rollout steps until the + * next `entering_trade` overwrites it; consumed at `is_close` by + * the SP20 reward kernel (Task 2.2 wires the read site). + * FoldReset-zeroed by `state_reset_registry.rs::label_at_open_per_env`. */ + const int* __restrict__ aux_label_per_env, + int* __restrict__ label_at_open_per_env ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -2887,6 +2907,39 @@ extern "C" __global__ void experience_env_step( if (entering_trade) { entry_price = raw_close; trade_start_pnl = (ps[PS_REALIZED_PNL]); + + /* SP20 Phase 2 Task 2.0 (2026-05-09): capture aux next-bar label + * sign at trade open into per-env scratch. Consumed at trade + * close (`is_close == 1`) by `sp20_compute_event_reward`'s + * `dir_match` quadrant check (spec §4.1). + * + * Mapping from `aux_label_per_env[i]` (kernel emits `0`/`1`/`-1` + * per `aux_sign_label_per_step_kernel.cu:77` — `0`=not-up, + * `1`=up, `-1`=skip-window) to the SP20 sign convention + * (`{-1, 0, +1}`): + * `1` → `+1` (aux predicted up at trade open) + * `0` → `-1` (aux predicted not-up — down or flat) + * `-1` → `0` (aux skip-window: lookahead past total_bars, + * no signal — sentinel 0 is the "no info" + * case in `sp20_compute_event_reward`). + * + * NULL-tolerant — the kernel arg pair is optional so test + * scaffolds without aux-head wiring don't break. When NULL the + * write is a no-op; the FoldReset-zeroed buffer reads as + * sentinel 0 at the trade-close consumer (Task 2.2 wires the + * read site). + * + * Race-free by construction: each thread writes only its own + * `[i]` slot — no inter-thread contention. Pure GPU buffer + * (kernel-to-kernel only); no host visibility needed. */ + if (aux_label_per_env != NULL && label_at_open_per_env != NULL) { + const int aux_l = aux_label_per_env[i]; + int label_at_open_sign; + if (aux_l == 1) label_at_open_sign = +1; + else if (aux_l == 0) label_at_open_sign = -1; + else label_at_open_sign = 0; /* aux_l == -1 (skip) */ + label_at_open_per_env[i] = label_at_open_sign; + } } /* On reversal: close old segment, open new one. diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index ca854f8c4..f5a91f3a3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -641,6 +641,43 @@ pub struct GpuExperienceCollector { portfolio_states: CudaSlice, // [alloc_episodes * PORTFOLIO_STRIDE] (43 floats per episode) episode_starts_buf: CudaSlice,// [alloc_episodes] + /// SP20 Phase 2 Task 2.0 (2026-05-09): per-env label sign captured at trade + /// open, consumed at trade close by `sp20_compute_event_reward`'s + /// `dir_match` check (spec §4.1). + /// + /// Layout: `[alloc_episodes]` i32 — one slot per env, persists across + /// the rollout and replay-window inside `experience_env_step`. Values: + /// `+1` — aux next-bar predicted "up" at trade open + /// `-1` — aux next-bar predicted "not-up" (down/flat) at trade open + /// `0` — sentinel: trade-open was during aux skip-window (look-ahead + /// past `total_bars`) or no position currently open. The 4-quadrant + /// reward maps `pnl_sign != label_at_open_sign` to "wrong reason" + /// and `pnl_sign == 0` (close_pnl == 0) to no-info; sentinel 0 + /// input flows through naturally because no `pnl_sign != 0` case + /// can match `label_at_open_sign == 0` (mismatch ⇒ "wrong reason"). + /// This lands the wrong-reason quadrant for sentinel inputs, which + /// is the safer default than over-rewarding a trade with no + /// reasoning signal. + /// + /// Producer: `experience_env_step` writes to this slot in the + /// `entering_trade` branch by mapping + /// `exp_aux_nb_label_buf[env]` (which is the per-step aux next-bar + /// label written by `aux_sign_label_per_step_kernel` — values + /// `{-1, 0, +1}`) to `{0, -1, +1}` respectively (the kernel's `0` = + /// "not-up" maps to `-1` for the spec sign convention; the kernel's + /// `-1` = "skip" maps to sentinel `0`). + /// + /// Consumer: `experience_env_step` reads the slot in the `is_close` + /// branch and passes it to `sp20_compute_event_reward(close_pnl, + /// label_at_open_sign, isv[LOSS_CAP_INDEX])`. + /// + /// FoldReset semantics (registered in `state_reset_registry.rs` as + /// `label_at_open_per_env`): zero on every fold boundary so the new + /// fold's first `entering_trade` branch fully populates the slot from + /// the new fold's aux head outputs (no stale label leakage across + /// folds — each fold has independent aux next-bar predictions). + pub(crate) label_at_open_per_env: CudaSlice, // [alloc_episodes] + // Output buffers [alloc_episodes * alloc_timesteps, ...] states_out: CudaSlice, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32 actions_out: CudaSlice, // [alloc_episodes * alloc_timesteps] @@ -1689,6 +1726,17 @@ impl GpuExperienceCollector { .alloc_zeros::(alloc_episodes) .map_err(|e| MLError::ModelError(format!("alloc episode_starts_buf: {e}")))?; + // SP20 Phase 2 Task 2.0 (2026-05-09): per-env label sign captured + // at trade open, consumed at trade close. Sentinel = 0 ("no + // position currently open or aux was in skip-window at entry"); + // first entering_trade observation overwrites with mapped aux + // label sign in {-1, 0, +1}. Sized [alloc_episodes]. + let label_at_open_per_env = stream + .alloc_zeros::(alloc_episodes) + .map_err(|e| MLError::ModelError(format!( + "sp20: alloc label_at_open_per_env ({alloc_episodes} i32): {e}" + )))?; + // Persistent epoch state let epoch_state_init: Vec = vec![ 0.01, // vol_ema @@ -2716,6 +2764,7 @@ impl GpuExperienceCollector { curiosity_weights, portfolio_states, episode_starts_buf, + label_at_open_per_env, states_out, actions_out, rewards_out, @@ -6100,6 +6149,16 @@ impl GpuExperienceCollector { // `rewards_out` while skipping early-return slots. .arg(&mut self.r_discipline_per_sample) .arg(&mut self.slot_completed_normally_per_sample) + // SP20 Phase 2 Task 2.0 (2026-05-09): aux next-bar + // label tile (input — already populated by the + // `aux_sign_label_per_step_kernel` launch in step + // 3a-β above on the same stream; stream-implicit + // producer→consumer ordering) + per-env trade-open + // label-sign scratch (output — written at the + // `entering_trade` branch, consumed at `is_close` + // by Task 2.2's `sp20_compute_event_reward` call). + .arg(&self.exp_aux_nb_label_buf) + .arg(&mut self.label_at_open_per_env) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 44187bfdb..ebcc18868 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2114,6 +2114,14 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "SP20 Phase 1.4 mapped-pinned [8] i32 per-EMA observation counter for `sp20_emas_compute_kernel`. counter==0 ⇒ next firing observation REPLACES sentinel directly (Pearl-A bootstrap); counter>0 ⇒ Wiener-blend with α=0.4 floor. FoldReset zeros all 8 counters so each fold's first per-EMA firing observation re-bootstraps from sentinel — necessary because some EMAs legitimately observe 0.0 and the count-based bootstrap is the correct sentinel discriminator per `pearl_first_observation_bootstrap`.", }, + // ── SP20 Phase 2 Task 2.0 (2026-05-09): per-env trade-open + // label sign scratch — written at `entering_trade`, read at + // `is_close` by the SP20 4-quadrant reward kernel. + RegistryEntry { + name: "label_at_open_per_env", + category: ResetCategory::FoldReset, + description: "GpuExperienceCollector.label_at_open_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.0 (2026-05-09) per-env aux next-bar label sign captured at trade open. Layout: tile[env] in {-1, 0, +1} where +1 = aux predicted up at trade open, -1 = aux predicted not-up, 0 = sentinel (no position currently open or trade-open landed in aux skip-window). Producer: `experience_env_step` `entering_trade` branch — reads `exp_aux_nb_label_buf[env]` (kernel emits `0`/`1`/`-1` per `aux_sign_label_per_step_kernel.cu:77`) and maps `0`→`-1`, `1`→`+1`, `-1`→`0` to match the SP20 spec sign convention (spec §4.1: `label_at_open = sign(SP19_blended_label[trade_open_bar])`). Consumer: `experience_env_step` `is_close` branch — passes `label_at_open_per_env[env]` to `sp20_compute_event_reward(close_pnl, label_at_open_sign, isv[LOSS_CAP_INDEX])`. FoldReset sentinel 0 across all `alloc_episodes` slots — leftover label from the previous fold's open trade would corrupt the new fold's first trade-close `dir_match` check (the new fold's aux head has independent calibration and would not have produced that label). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Per `pearl_first_observation_bootstrap` the sentinel-0 input also flows through `sp20_compute_event_reward` to the wrong-reason quadrant on the corner case where a trade closes WITHOUT a prior `entering_trade` having fired this fold (trade carry-over from before the reset is impossible because portfolio_states is also FoldReset).", + }, ]; Self { entries } } @@ -2367,4 +2375,42 @@ mod sp20_registry_tests { ); } } + + /// SP20 Phase 2 Task 2.0: `label_at_open_per_env` MUST be registered as + /// FoldReset so the new fold's first `entering_trade` branch fully + /// re-populates from the new fold's aux head (preventing stale label + /// leakage across folds — each fold has independent aux head + /// calibration). Description MUST mention SP20 Phase 2 + Task 2.0 + + /// `entering_trade` + `is_close` so wire-up audits have a stable + /// search anchor. + #[test] + fn sp20_label_at_open_per_env_registered_fold_reset() { + let registry = StateResetRegistry::new(); + let entry = registry + .entries + .iter() + .find(|e| e.name == "label_at_open_per_env") + .expect("label_at_open_per_env must be registered in StateResetRegistry"); + assert_eq!( + entry.category, + ResetCategory::FoldReset, + "label_at_open_per_env must be FoldReset (got {:?})", + entry.category + ); + assert!( + entry.description.contains("SP20 Phase 2"), + "description must reference 'SP20 Phase 2'; got: {}", + entry.description + ); + assert!( + entry.description.contains("entering_trade"), + "description must reference 'entering_trade' producer site; got: {}", + entry.description + ); + assert!( + entry.description.contains("is_close"), + "description must reference 'is_close' consumer site; got: {}", + entry.description + ); + } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 309747a73..298b0d7bb 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -9538,6 +9538,22 @@ impl DQNTrainer { collector.sp20_emas_obs_count_buf.host_slice_mut().fill(0); } } + // SP20 Phase 2 Task 2.0 (2026-05-09): per-env trade-open label + // sign scratch. Device-resident `CudaSlice` (NOT mapped- + // pinned, because the kernel ONLY reads/writes from device + // threads — no host visibility needed). Reset path: + // `stream.memset_zeros(&mut buf)` — async GPU memset, same + // pattern as `gpu_attention.rs::reset_eval_v_range_state`'s + // `attn_m/attn_v` zeroing. + "label_at_open_per_env" => { + if let Some(ref mut collector) = self.gpu_experience_collector { + let stream = collector.stream().clone(); + stream.memset_zeros(&mut collector.label_at_open_per_env) + .map_err(|e| crate::MLError::ModelError(format!( + "label_at_open_per_env memset_zeros: {e}" + )))?; + } + } // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots. // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT // a stateful EMA) — rewrite the constructor's value at fold diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 721a87464..77dd85868 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,99 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP20 Phase 2 Task 2.0: per-env trade-open label-sign infrastructure (additive) + +Lands the producer-side infrastructure for Phase 2's 4-quadrant reward +kernel: a per-env `[alloc_episodes]` `i32` device buffer +(`label_at_open_per_env`) populated at every `entering_trade` branch +in `experience_env_step` from the upstream `aux_sign_label_per_step_kernel` +output. The Task 2.2 follow-up commit will land the consumer (the +`is_close` branch reading `label_at_open_per_env[i]` and passing it to +`sp20_compute_event_reward(close_pnl, label_at_open_sign, +isv[LOSS_CAP_INDEX])`). + +### Why a separate per-env buffer (NOT a derived-at-trade-close read) + +The 4-quadrant reward semantically requires the label sampled at +**trade-open**, not at close-bar. Reading `aux_label_per_env[i]` at +`is_close` would sample a different bar than the model's entry decision +was based on — that turns the dir_match check into "did the model close +when the immediate-next-bar's prediction agrees with the close P&L" +instead of "did the model close when the entry-bar's prediction agreed +with the close P&L". The trade-open sample is the load-bearing variant +per spec §4.1. + +### Layout / sign mapping + +| Slot value | Meaning | +|------------|------------------------------------------------------------------| +| `+1` | Aux next-bar predicted "up" at trade-open | +| `-1` | Aux next-bar predicted "not-up" (down/flat) at trade-open | +| `0` | Sentinel: no position open OR aux skip-window at trade-open | + +`aux_sign_label_per_step_kernel.cu:77` emits `0`/`1` for not-up/up and +`-1` for skip-window. The Task 2.0 producer at `experience_env_step`'s +`entering_trade` branch maps `aux_l ∈ {1, 0, -1}` → `{+1, -1, 0}` to +match the SP20 spec sign convention (spec §4.1: `label_at_open = +sign(SP19_blended_label[trade_open_bar])`). + +### FoldReset semantics + +Registry entry `label_at_open_per_env` (FoldReset). Reset path: +`stream.memset_zeros(&mut buf)` (async GPU memset — same pattern as +`gpu_attention.rs::reset_eval_v_range_state`'s `attn_m/attn_v` +zeroing). Without this reset, leftover label from the previous fold's +open trade would corrupt the new fold's first trade-close `dir_match` +check (each fold has independent aux head calibration). + +### Kernel-arg threading + +`experience_env_step` signature gains 2 new args at the end (no +re-ordering of existing args): + +```c +const int* __restrict__ aux_label_per_env, // [N] input +int* __restrict__ label_at_open_per_env // [N] output +``` + +Both are NULL-tolerant — test scaffolds without aux-head wiring +continue to work; the NULL-guarded write is a no-op and the +FoldReset-zeroed buffer reads as sentinel 0 at the consumer (Task 2.2). + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +2 args, +write at L2887 | Producer: aux_l → label_at_open_sign mapping at `entering_trade` | +| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +field, +alloc, +launch args | Buffer + alloc + threading into env_step launch | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +entry + test | FoldReset registration + invariant test | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +match arm | `reset_named_state` dispatch arm | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml +SQLX_OFFLINE=true cargo test -p ml --lib sp20 +SQLX_OFFLINE=true cargo test -p ml --lib sp20_label_at_open_per_env_registered_fold_reset +``` + +CPU tests (Task 2.0): registry invariant test + 18 Task 0/Phase 1 tests +all green. GPU oracle tests for the producer + consumer land in Task 2.1 +(reward function math) and Task 2.2 (full trade-close path swap). + +### Phase 2 → Task 2.1 / 2.2 forward references + +- Task 2.1: `sp20_compute_event_reward` device function definition + 6 + unit tests via `sp12_reward_math_test_kernel.cu`-style pattern. +- Task 2.2: replaces SP12 v3 inlined reward block at the + `segment_complete && segment_hold_time > 0.0f` site with the SP20 + 4-quadrant reward, threading `label_at_open_per_env[i]` as the + `label_at_open_sign` arg. Same atomic commit lands the + per-env `alpha_per_env[N]` plumbing through to the SP20 aggregation + kernel's `alpha` field (which currently emits 0.0 as a Phase 2 + forward-reference placeholder). + ## 2026-05-09 — volume_bar_size in cache key + OFI front-month filter fix `crates/ml/src/trainers/dqn/data_loading.rs` updated in two places: