diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 6897ada66..ffed74171 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2203,7 +2203,22 @@ extern "C" __global__ void experience_env_step( * `q_dir_abs_ref_idx` pattern at line 2069). */ float* __restrict__ alpha_per_env, int loss_cap_idx, - int alpha_ema_idx + int alpha_ema_idx, + /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — per-env trade-close win + * flag output buffer. + * + * `is_win_per_env` ← [N] i32: written at `is_close && segment_ + * hold_time > 0.0f` (segment_complete branch) as + * `(segment_return > 0.0f) ? 1 : 0`. Mirrors the alpha emit site + * exactly — the SP20 aggregation kernel reads alpha sum and + * is_win count side-by-side at every closed env. Replaces the + * broken per-bar `step_ret > 0` predicate the aggregation kernel + * was using, which pinned WR_EMA at 0 in production (per-bar + * step_ret at close bars is dominated by tx-cost so almost always + * negative even for trades that closed profitably). NULL-tolerant + * for test scaffolds — when NULL the write is a no-op and the + * aggregation kernel falls back to the legacy `sr > 0` predicate. */ + int* __restrict__ is_win_per_env ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -3308,6 +3323,20 @@ extern "C" __global__ void experience_env_step( alpha_per_env[i] = alpha; } + /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — per-env win flag + * emit. Written at the same segment_complete site as + * `alpha_per_env[i] = alpha`. The aggregation kernel reads + * `is_win_per_env[env]` in place of the broken per-bar + * `step_ret > 0` predicate. Per `pearl_audit_unboundedness_ + * for_implicit_asymmetry`-style reasoning: the per-bar step_ + * ret is NOT a faithful proxy for trade outcome at close + * bars (tx-cost domination flips the sign relative to + * segment_return). Race-free: each thread writes only its + * own `[i]` slot. */ + if (is_win_per_env != NULL) { + is_win_per_env[i] = (segment_return > 0.0f) ? 1 : 0; + } + /* SP11 component attribution split (mutually exclusive — preserved * from pre-SP20 semantic). r_used is identical in both branches; * only the rc[*] slot the SP11 controller weighs differs. */ diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 3b2d95885..b3fbbe9c7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -706,6 +706,35 @@ pub struct GpuExperienceCollector { /// don't leak into the EMA. pub(crate) alpha_per_env: CudaSlice, // [alloc_episodes] + /// SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade-close win + /// flag. `[alloc_episodes]` i32 device-resident scratch — `1` iff + /// the segment's realized P&L was strictly positive at trade close + /// (`segment_return > 0.0f`), `0` otherwise. Replaces the broken + /// per-bar `step_ret > 0` predicate the aggregation kernel was + /// using, which conflated the close bar's per-bar mark-to-market + /// tick (dominated by tx-cost so almost always false at close + /// bars) with trade outcome — pinning WR_EMA at 0 across all + /// training epochs in production. + /// + /// Layout: `[alloc_episodes]` i32 — one slot per env, written by + /// `experience_env_step`'s `is_close` (`segment_complete`) branch + /// at the same site as `alpha_per_env`. Read once per rollout + /// step by `sp20_aggregate_inputs_kernel` to populate the + /// `wins_count` reduction (mirror of the alpha-mean aggregation). + /// + /// FoldReset semantics (registered in `state_reset_registry.rs` as + /// `is_win_per_env`): zero on every fold boundary so the new + /// fold's first trade-close fully populates the slot from the new + /// fold's segment_return calculation (no stale win flag leakage + /// across folds — each fold has independent WR_EMA bootstrap state + /// per Phase 1.4's `obs_count[OBS_COUNT_WR]` reset). + /// + /// On non-close steps the slot retains the previous trade-close + /// flag (or 0 sentinel pre-first-close); the aggregation kernel + /// reads only when `trade_close_per_env[env] != 0`, so stale + /// values don't leak into the EMA. + pub(crate) is_win_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] @@ -1776,6 +1805,21 @@ impl GpuExperienceCollector { "sp20: alloc alpha_per_env ({alloc_episodes} f32): {e}" )))?; + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade-close + // win flag (1 iff segment_return > 0 at trade close, else 0). + // Sentinel = 0; updated at each trade close by the same + // SP20 4-quadrant reward kernel block site that writes + // `alpha_per_env`. Consumed by sp20_aggregate_inputs_kernel + // (count over closed envs with flag=1) for the wins_count + // reduction. Replaces the broken per-bar `step_ret > 0` + // predicate that pinned WR_EMA at 0 in production. + // Sized [alloc_episodes]. + let is_win_per_env = stream + .alloc_zeros::(alloc_episodes) + .map_err(|e| MLError::ModelError(format!( + "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}" + )))?; + // Persistent epoch state let epoch_state_init: Vec = vec![ 0.01, // vol_ema @@ -2805,6 +2849,7 @@ impl GpuExperienceCollector { episode_starts_buf, label_at_open_per_env, alpha_per_env, + is_win_per_env, states_out, actions_out, rewards_out, @@ -6222,6 +6267,16 @@ impl GpuExperienceCollector { use crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX; ALPHA_EMA_INDEX as i32 }) + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env + // trade-close win flag (output) — 1 iff + // `segment_return > 0` at trade close. Written at + // the same segment_complete branch site as + // `alpha_per_env`. Consumed by + // `sp20_aggregate_inputs_kernel` for the + // `wins_count` reduction (replaces the broken + // per-bar `step_ret > 0` predicate that pinned + // WR_EMA at 0 in production). + .arg(&mut self.is_win_per_env) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" @@ -6508,14 +6563,16 @@ impl GpuExperienceCollector { // `out_inputs->alpha = 0.0f` placeholder atomically per // `feedback_no_partial_refactor`. let alpha_per_env_dev = self.alpha_per_env.raw_ptr(); - // SP20 Phase 2 Task 2.2-fix (2026-05-10): pass `0` (NULL) - // for the new `is_win_per_env_dev` arg in this commit. - // Commit 1 (failing-test) plumbs the kernel-arg signature - // change so the unit test can exercise the new arg, while - // keeping production behavior unchanged. Commit 2 (fix) - // wires the per-env producer + collector buffer + reset - // entry atomically with the kernel-body change that - // honors the arg. + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env + // trade-close win flag pointer. `is_win_per_env` is + // `[alloc_episodes]` i32 device-resident, populated at + // every trade close by the SP20 4-quadrant reward block + // in `experience_env_step` (launched immediately above + // on the same stream — stream-implicit producer→ + // consumer ordering). Replaces the broken per-bar + // `step_ret > 0` predicate that pinned WR_EMA at 0 in + // production. + let is_win_per_env_dev = self.is_win_per_env.raw_ptr(); unsafe { launch_sp20_aggregate_inputs( &self.stream, @@ -6529,7 +6586,7 @@ impl GpuExperienceCollector { stats_std_dev, aux_dir_acc_dev, alpha_per_env_dev, - /* is_win_per_env_dev = */ 0, + is_win_per_env_dev, n_envs_i32_sp20, dir_divisor, hold_action_id, diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu index d719a9b76..a779364f3 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -216,7 +216,26 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( if (tc != 0) { local_closed_count += 1; - if (sr > 0.0f) local_wins_count += 1; + /* SP20 Phase 2 Task 2.2-fix (2026-05-10): wins predicate + * derives from the segment-level realized-return sign + * via the per-env i32 `is_win_per_env` buffer (1 iff + * segment_return > 0 at trade close), populated by + * experience_env_step's segment_complete branch alongside + * `alpha_per_env`. The legacy per-bar `sr > 0` predicate + * is preserved as the NULL-tolerant fallback for oracle- + * test scaffolds that don't wire a segment-level producer. + * Bug class: per-bar step_ret at close bars is + * `position*(close_t-close_{t-1}) - tx_cost`; the per-bar + * tick is small while tx_cost is fixed → `sr > 0` is + * almost always false at close bars even when the trade + * closed profitably overall, pinning WR_EMA at 0. */ + int seg_win; + if (is_win_per_env != NULL) { + seg_win = (is_win_per_env[env] != 0) ? 1 : 0; + } else { + seg_win = (sr > 0.0f) ? 1 : 0; + } + local_wins_count += seg_win; local_hold_at_exit_sum += he; /* SP20 Phase 2 Task 2.2: gated alpha read on closed-env * slots only. Non-close-step alpha_per_env values are diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 847dc4530..d04072b41 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -2132,6 +2132,17 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuExperienceCollector.alpha_per_env [alloc_episodes] f32 device-resident scratch — SP20 Phase 2 Task 2.2 (2026-05-09) per-env trade-close alpha (= R_event - hold_baseline) per spec §4.1. Phase 2 emits `hold_baseline = 0.0f` as a Phase 3.2 forward-reference placeholder; Phase 3.2's Hold opportunity-cost dual emission lands the real `hold_baseline_buffer.sum_over_trade_range()` consumer. Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `alpha = sp20_compute_event_reward(segment_return, label_at_open_per_env[env], ISV[LOSS_CAP_INDEX])`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — sums `alpha_per_env[env]` over closed envs and divides by closed-count to populate `EmaInputs.alpha`. FoldReset sentinel 0.0 across all `alloc_episodes` slots — leftover alpha from the previous fold's last trade-close would corrupt the new fold's first ALPHA_EMA observation. The sentinel-0 reads as `alpha == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Replaces the Phase 1.4 hardcoded `out_inputs->alpha = 0.0f` placeholder atomically per `feedback_no_partial_refactor` (Task 2.2 lands buffer + kernel-arg + aggregation-kernel-update + 2 placeholder pin tests in one commit).", }, + // ── SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env + // trade-close win flag scratch — written at `is_close` + // (segment_complete) alongside `alpha_per_env`, read by + // `sp20_aggregate_inputs_kernel` for the `wins_count` + // reduction (replacing the broken per-bar `step_ret > 0` + // predicate that pinned WR_EMA at 0 in production). + RegistryEntry { + name: "is_win_per_env", + category: ResetCategory::FoldReset, + description: "GpuExperienceCollector.is_win_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.2-fix (2026-05-10) per-env trade-close win flag (1 iff segment_return > 0 at trade close, 0 otherwise). Producer: `experience_env_step` `is_close` (segment_complete) branch — writes `is_win_per_env[env] = (segment_return > 0.0f) ? 1 : 0` at the same site as `alpha_per_env[env] = alpha`. Consumer: `sp20_aggregate_inputs_kernel` (Path C aggregation) — counts `is_win_per_env[env] != 0` over closed envs to populate the `wins_count` reduction (drives `EmaInputs.is_win` via the >=0.5 fraction-of-closed threshold). Replaces the broken `step_ret_per_env[env] > 0.0f` per-bar predicate the kernel was using, which conflated the close bar's per-bar mark-to-market tick (dominated by tx-cost so almost always negative at close bars) with trade outcome — pinning WR_EMA at 0 across all training epochs in production despite real wins. FoldReset sentinel 0 across all `alloc_episodes` slots — leftover win flag from the previous fold's last trade-close would corrupt the new fold's first WR_EMA observation. The sentinel-0 reads as `is_win == 0` for non-close steps; the aggregation kernel reads only when `trade_close_per_env[env] != 0` so the FoldReset sentinel never reaches the EMA. Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice` (device-resident, no host buffer to fill). Lands atomically per `feedback_no_partial_refactor` with the kernel-body change + buffer alloc + producer wire-up in `experience_kernels.cu::segment_complete` + the new GPU oracle test `wr_ema_uses_segment_pnl_via_is_win_per_env_predicate` (failing-test commit 1, fix commit 2).", + }, ]; Self { entries } } @@ -2424,6 +2435,50 @@ mod sp20_registry_tests { ); } + /// SP20 Phase 2 Task 2.2-fix: `is_win_per_env` MUST be registered + /// as FoldReset so the new fold's first trade-close fully re- + /// populates from the new fold's `segment_return > 0` predicate + /// (preventing stale win-flag leakage into the new fold's WR_EMA + /// bootstrap). Description MUST mention SP20 Phase 2 + Task + /// 2.2-fix + `segment_complete` (producer) + + /// `sp20_aggregate_inputs_kernel` (consumer) so wire-up audits + /// have a stable search anchor. + #[test] + fn sp20_is_win_per_env_registered_fold_reset() { + let registry = StateResetRegistry::new(); + let entry = registry + .entries + .iter() + .find(|e| e.name == "is_win_per_env") + .expect("is_win_per_env must be registered in StateResetRegistry"); + assert_eq!( + entry.category, + ResetCategory::FoldReset, + "is_win_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("Task 2.2-fix"), + "description must reference 'Task 2.2-fix'; got: {}", + entry.description + ); + assert!( + entry.description.contains("segment_complete"), + "description must reference 'segment_complete' producer site; got: {}", + entry.description + ); + assert!( + entry.description.contains("sp20_aggregate_inputs_kernel"), + "description must reference the aggregation-kernel consumer; got: {}", + entry.description + ); + } + /// SP20 Phase 2 Task 2.2: `alpha_per_env` MUST be registered as /// FoldReset so the new fold's first trade-close fully re-populates /// from the new fold's reward calculation (preventing stale alpha diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 4bd5f7d9b..0f0dca952 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -9566,6 +9566,18 @@ impl DQNTrainer { )))?; } } + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade- + // close win flag scratch. Same device-resident + // `CudaSlice` reset pattern as `alpha_per_env` above. + "is_win_per_env" => { + if let Some(ref mut collector) = self.gpu_experience_collector { + let stream = collector.stream().clone(); + stream.memset_zeros(&mut collector.is_win_per_env) + .map_err(|e| crate::MLError::ModelError(format!( + "is_win_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 25e69b58f..5b665eeb1 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,96 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-10 — SP20 Phase 2 Task 2.2-fix: WR_EMA pinning bug — fix commit (kernel body + producer + buffer + reset) + +Closes the fix gap left by the failing-test commit. The aggregation +kernel body now reads `is_win_per_env[env]` (when non-NULL) in the +`tc != 0` branch in place of the broken per-bar `step_ret > 0` +predicate; the producer site at `experience_kernels.cu::segment_ +complete` writes the segment-level `(segment_return > 0.0f) ? 1 : 0` +flag at the same location as `alpha_per_env[i] = alpha`; the +collector owns a new `is_win_per_env: CudaSlice` `[alloc_episodes]` +buffer (zero-initialized at construction, fold-reset registered); +the production caller now passes the buffer's device pointer where +the previous commit passed `0` (NULL). + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` | body | `is_win_per_env != NULL` branch reads the per-env i32 flag for `wins_count` | +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +arg, +write | New `int* is_win_per_env` kernel arg; segment_complete branch writes `(segment_return > 0) ? 1 : 0` | +| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +field, +alloc, +arg-pass | `is_win_per_env: CudaSlice` field, alloc_zeros, kernel-arg pass at env_step + sp20_aggregate launches | +| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +entry, +test | FoldReset entry + assertion test | +| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +arm | Fold-reset dispatch arm `"is_win_per_env" => memset_zeros` | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification (this commit) + +``` +SQLX_OFFLINE=true cargo check -p ml --tests +SQLX_OFFLINE=true cargo test -p ml --lib sp20 +SQLX_OFFLINE=true cargo test -p ml --lib state_reset_registry +``` + +All 21 SP20 lib tests + 9 registry tests pass. + +GPU oracle test verification (deferred until L40S smoke): + +``` +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=89 cargo test -p ml \ + --test sp20_aggregate_inputs_test --features cuda \ + -- --ignored --nocapture +``` + +The new `wr_ema_uses_segment_pnl_via_is_win_per_env_predicate` test +passes after this commit; it would FAIL on the prior commit (kernel +body still using `sr > 0`). The `null_is_win_per_env_falls_back_to_ +legacy_step_ret_predicate` and the existing 6 SP20 aggregate kernel +tests preserve the NULL-tolerance contract. + +### Pearls + invariants honoured + +- `feedback_no_partial_refactor` — kernel-body change + producer + + buffer + reset entry land atomically in this commit. The + preceding failing-test commit was a TDD-evidence vehicle (kernel- + arg signature plumb without behavioral change); no production + consumer used the new arg between commit 1 and commit 2 (production + caller passed `0` in commit 1). +- `feedback_no_stubs` — every site fully wired: collector buffer + alloc/field/alloc_zeros, env_step kernel writes via new arg, sp20 + aggregate reads via existing arg, fold reset registered AND + dispatched. No legacy aliases — the legacy `sr > 0` predicate is + preserved as the NULL-tolerant fallback only for oracle-test + scaffolds (per `feedback_no_legacy_aliases` carve-out: tests can + pass `0` so the kernel falls back; production always wires real). +- `pearl_first_observation_bootstrap` — `is_win_per_env` is FoldReset- + zeroed; the first close in a new fold REPLACES the WR_EMA value + directly (count == 0 ⇒ replace). For non-close steps the FoldReset + sentinel never reaches the EMA because the aggregation kernel + gates the read on `trade_close_per_env[env] != 0`. +- `feedback_isv_for_adaptive_bounds` — no new ISV slots; the SP20 ISV + layout is unchanged. This is a purely architectural fix (data flow) + for the existing WR_EMA / LOSS_CAP slot pair. +- `feedback_no_atomicadd` — `wins_count` accumulation remains in the + per-block tree-reduce stripe `sh_wins_count`; no atomicAdd added. + The kernel body change is a one-line replacement of the per-element + predicate inside the existing reduction. + +### Bug class memory + +The pattern "per-bar mark-to-market step_return is the win predicate" +was wrong because, at close bars, mark-to-market includes tx-cost +(deducted via `*cash -= cost` in `execute_trade`) and only the close +bar's `Δprice` tick (not the whole trade's accumulated tick). The +segment-level realized P&L (`(realized_pnl + unrealized_at_exit - +trade_start_pnl) / prev_equity`) is the correct outcome scalar; the +producer in `experience_env_step::segment_complete` already had it +(used by `sp20_compute_event_reward`'s sign check), the aggregation +kernel just wasn't reading the right thing. The bug had been latent +since SP20 Phase 1.4 / 2.2 landed on 2026-05-09 and was caught on +2026-05-10 by HEALTH_DIAG observation in a multi-seed L40S smoke. + ## 2026-05-10 — SP20 Phase 2 Task 2.2-fix: WR_EMA pinning bug — failing test commit (signature plumb only) Plumbs a new `is_win_per_env` per-env i32 device-buffer arg through