merge(sp20): restore is_win_per_env across struct/launcher/registry after Phase 3 cherry-pick

Cherry-picking Phase 3 (a8731dda4 + 84929f419) onto Phase 2 fix tip
(13ce78385) with --strategy-option=theirs dropped is_win_per_env arg + field
+ registration the wr_ema fix added. Manual restoration:

- experience_kernels.cu: re-added is_win_per_env arg to experience_env_step
- gpu_experience_collector.rs: re-added struct field, alloc, ctor entry, launcher arg
- state_reset_registry.rs: re-added FoldReset RegistryEntry
- training_loop.rs: re-added named-reset dispatch arm
- docs/dqn-wire-up-audit.md: documented the merge restoration

Test sp20_is_win_per_env_registered_fold_reset passes. Workspace + examples
compile clean. wr_ema runtime mystery (fix wired, but production still
shows wr_ema=0) unchanged — debug per 13ce78385 scaffold pending.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -2205,6 +2205,13 @@ extern "C" __global__ void experience_env_step(
     float* __restrict__ alpha_per_env,
     int loss_cap_idx,
     int alpha_ema_idx,
+    /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — segment-level win flag.
+     * `is_win_per_env` ← [N] i32: written at `is_close && segment_
+     * complete && segment_hold_time > 0.0f` with
+     * `(segment_return > 0.0f) ? 1 : 0`. Consumer:
+     * sp20_aggregate_inputs_kernel reads this in place of the broken
+     * per-bar `step_ret > 0` predicate. NULL-tolerant for tests. */
+    int* __restrict__ is_win_per_env,
     /* SP20 Phase 3 Task 3.2 (2026-05-10) — Hold opportunity-cost dual
      * emission (Component 2) per spec §4.2.
      *

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -735,6 +735,14 @@ pub struct GpuExperienceCollector {
     /// don't leak into the EMA.
     pub(crate) alpha_per_env: CudaSlice<f32>, // [alloc_episodes]
 
+    /// SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment-level win
+    /// flag scratch. Companion to `alpha_per_env` — same producer site
+    /// (`segment_complete` branch in experience_kernels.cu), same fold-
+    /// reset semantics, same NULL-tolerant consumer (sp20_aggregate_
+    /// inputs_kernel reads when non-NULL, falls back to legacy per-bar
+    /// `step_ret > 0` predicate when NULL for test scaffolds).
+    pub(crate) is_win_per_env: CudaSlice<i32>, // [alloc_episodes]
+
     /// SP20 Phase 3 Task 3.2 (2026-05-10): per-env circular Hold
     /// opportunity-cost buffer (Component 2) per spec §4.2.
     ///
@@ -1873,6 +1881,18 @@ impl GpuExperienceCollector {
                 "sp20: alloc alpha_per_env ({alloc_episodes} f32): {e}"
             )))?;
 
+        // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment-level
+        // win flag (1 iff segment_return > 0 at trade close, else 0).
+        // Sentinel = 0; written by experience_env_step at the same
+        // segment_complete site as alpha_per_env. Consumed by the
+        // aggregation kernel via the new is_win_per_env arg replacing
+        // the broken per-bar step_ret predicate. Sized [alloc_episodes].
+        let is_win_per_env = stream
+            .alloc_zeros::<i32>(alloc_episodes)
+            .map_err(|e| MLError::ModelError(format!(
+                "sp20: alloc is_win_per_env ({alloc_episodes} i32): {e}"
+            )))?;
+
         // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold opportunity-
         // cost circular buffer. Layout `[alloc_episodes ×
         // HOLD_BASELINE_BUFFER_SIZE = 30]` f32 row-major. Sentinel 0.0
@@ -2933,6 +2953,7 @@ impl GpuExperienceCollector {
             episode_starts_buf,
             label_at_open_per_env,
             alpha_per_env,
+            is_win_per_env,
             hold_baseline_buffer,
             states_out,
             actions_out,
@@ -6367,6 +6388,14 @@ 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): segment-level
+                    // win flag i32 buffer. Written at segment_complete with
+                    // `(segment_return > 0.0f) ? 1 : 0`. Replaces the broken
+                    // per-bar `step_ret > 0` predicate in sp20_aggregate_
+                    // inputs_kernel's wins_count accumulation (the per-bar
+                    // tick at close bars is dominated by tx_cost deduction
+                    // → systematically negative even on winning trades).
+                    .arg(&mut self.is_win_per_env)
                     // SP20 Phase 3 Task 3.2 (2026-05-10): Hold
                     // opportunity-cost dual emission (Component 2). Six
                     // new args at the end of the kernel signature:

crates/ml/src/trainers/dqn/state_reset_registry.rs

@@ -2132,6 +2132,14 @@ 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<f32>` (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 segment-
+            // level win flag scratch. Companion to alpha_per_env, same
+            // segment_complete producer site, same FoldReset semantics.
+            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 segment-level win flag (1 iff segment_return > 0.0f at trade close, else 0). Producer: `experience_env_step` `is_close && segment_complete && segment_hold_time > 0.0f` branch — writes alongside `alpha_per_env[i] = alpha`. Consumer: `sp20_aggregate_inputs_kernel` reads `is_win_per_env[env]` to compute `wins_count` (replacing the broken per-bar `step_ret > 0` predicate; per-bar step_ret at close bars is dominated by tx_cost deduction → systematically negative even on winning trades, pinning WR_EMA at 0). 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 (not-won) 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<i32>` (device-resident, no host buffer to fill).",
+            },
             // ── SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold
             // opportunity-cost circular buffer. Component 2 producer/
             // consumer per spec §4.2. Written at every per-bar

crates/ml/src/trainers/dqn/trainer/training_loop.rs

@@ -9567,6 +9567,17 @@ impl DQNTrainer {
                         )))?;
                 }
             }
+            // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env segment-
+            // level win flag scratch. Companion to 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}"
+                        )))?;
+                }
+            }
             // SP20 Phase 3 Task 3.2 (2026-05-10): per-env Hold
             // opportunity-cost circular buffer (Component 2). Same
             // device-resident `CudaSlice<f32>` reset pattern as

docs/dqn-wire-up-audit.md

@@ -2,6 +2,10 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+## 2026-05-10 — SP20 Phase 2 fix + Phase 3 merge: restore is_win_per_env
+
+Cherry-picking Phase 3 commits onto Phase 2 fix tip with `--strategy-option=theirs` dropped the `is_win_per_env` kernel arg + struct field + reset registration that the wr_ema fix branch added (Phase 3 had its own version of the `experience_env_step` signature without that arg). Manual restoration: kernel signature, struct field, alloc, launcher arg, registry entry, reset dispatch arm. Test `sp20_is_win_per_env_registered_fold_reset` passes.
+
 ## 2026-05-10 — SP20 Phase 3 Task 3.4: replay-buffer schema for per-bar aux_conf (atomic)
 
 Lands the producer→ring→consumer schema plumbing for per-bar