gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(sp20): WR_EMA pinning bug — wire segment-level is_win_per_env producer + kernel-body read

Browse Source 
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`.

Atomic changes:
- sp20_aggregate_inputs_kernel.cu: when is_win_per_env != NULL, read
  the per-env i32 flag for wins_count; legacy `sr > 0` predicate
  preserved as NULL-tolerant fallback for oracle-test scaffolds.
- experience_kernels.cu: new `int* is_win_per_env` kernel arg;
  segment_complete branch writes `(segment_return > 0.0f) ? 1 : 0`
  at the same site as `alpha_per_env[i] = alpha`. Race-free per-i
  write (mirrors the alpha producer pattern). NULL-tolerant.
- gpu_experience_collector.rs: `is_win_per_env: CudaSlice<i32>`
  `[alloc_episodes]` field, alloc_zeros at construction, kernel-arg
  pass at experience_env_step launch site, raw_ptr() pass at
  sp20_aggregate_inputs launch site.
- state_reset_registry.rs: FoldReset entry + assertion test
  (sp20_is_win_per_env_registered_fold_reset). The
  every_fold_and_soft_reset_entry_has_dispatch_arm test passing
  confirms the dispatch arm is present.
- training_loop.rs: `"is_win_per_env" => memset_zeros` arm.
- docs/dqn-wire-up-audit.md: 2026-05-10 fix-commit entry.

Bug pattern for the memory pearl: per-bar mark-to-market step_return
is NOT the trade's segment-level win/loss outcome at close bars
because tx_cost (deducted via `*cash -= cost` in
trade_physics.cuh::execute_trade) plus the per-bar Δprice tick
swamp out segment-accumulated unrealized gains. The segment-level
realized P&L (`(realized_pnl + unrealized_at_exit -
trade_start_pnl) / prev_equity`) is the correct outcome scalar;
the producer in segment_complete already had it (used by
sp20_compute_event_reward's sign check for alpha), the aggregation
kernel just wasn't reading the right thing. Bug latent since SP20
Phase 1.4 / 2.2 landed on 2026-05-09; caught 2026-05-10 by
HEALTH_DIAG observation in a multi-seed L40S smoke (alpha_ema
populated correctly while wr_ema pinned at 0.0000 across 17 epochs).

Verification:
- SQLX_OFFLINE=true cargo check -p ml --tests passes
- 21 SP20 lib tests pass; 9 state_reset_registry tests pass
- The new GPU oracle test
  wr_ema_uses_segment_pnl_via_is_win_per_env_predicate passes
  after this commit (would FAIL on prior commit's kernel-body)
- The NULL-tolerance pin
  null_is_win_per_env_falls_back_to_legacy_step_ret_predicate
  preserves the pre-fix oracle-test contract

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-10 01:27:36 +02:00
parent 5a29c37cd8
commit 64bbbe4181
6 changed files with 273 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

31
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -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. */

75
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -706,6 +706,35 @@ 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 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<i32>, // [alloc_episodes]
// Output buffers [alloc_episodes * alloc_timesteps, ...]
states_out: CudaSlice<f32>, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32
actions_out: CudaSlice<i32>, // [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::<i32>(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<f32> = 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,

21
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu
View File

@@ -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

55
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -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<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
// 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<i32>` (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

12
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -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<i32>` 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

90
docs/dqn-wire-up-audit.md
View File

@@ -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<i32>` `[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<i32>` 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