From 5a29c37cd8d9a0f962034377ef1ee5bd991545b7 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 01:20:13 +0200 Subject: [PATCH] =?UTF-8?q?test(sp20):=20failing=20test=20pinning=20WR=5FE?= =?UTF-8?q?MA=20bug=20=E2=80=94=20wins=20predicate=20must=20use=20segment?= =?UTF-8?q?=20P&L=20not=20per-bar=20step=5Fret?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Plumbs new is_win_per_env i32 device-buffer arg through the sp20_aggregate_inputs kernel signature and Rust launcher, NULL-tolerant to preserve the legacy `step_ret_per_env > 0` predicate for existing oracle-test scaffolds. Production caller passes 0 (NULL) in this commit; commit 2 wires the per-env producer + collector buffer + reset entry atomically with the kernel-body change that honors the arg. Adds two new GPU oracle tests in sp20_aggregate_inputs_test: - wr_ema_uses_segment_pnl_via_is_win_per_env_predicate (FAILS without fix): step_ret all negative (tx-cost dominated close bars — the realistic production case), is_win_per_env = [1,1,1,0]. With the legacy predicate wins_count = 0 ⇒ is_win = 0 (the production bug). With the fix wins_count = 3 ⇒ is_win = 1. - null_is_win_per_env_falls_back_to_legacy_step_ret_predicate: pins the NULL-tolerance contract so existing oracle tests + Phase 1.4 wireup test scaffolds keep working unchanged. Bug root cause: WR_EMA pinned at 0 across all training epochs in production (HEALTH_DIAG observed `wr_ema=0.0000` while `alpha_ema` evolved with real values). The aggregation kernel's win predicate was `step_ret_per_env[env] > 0.0f` at the close bar, which is the per-bar mark-to-market `(new_value - prev_equity) / prev_equity`, NOT the trade's segment-level outcome. At close bars `step_ret = position * (close_t - close_{t-1}) - tx_cost`; the per-bar tick is small while tx_cost is fixed → `step_ret < 0` is dominant across closing bars even for trades that closed profitably overall. Result: wins_count = 0 always, is_win = 0 always, WR_EMA pinned at 0, LOSS_CAP pinned at -1.0 (cold-start cap, never ramping to -2.0 per spec §4.1). Fix arrives in the next commit: experience_kernels.cu segment_complete branch writes `is_win_per_env[i] = (segment_return > 0.0f) ? 1 : 0` (the segment-level realized-return signed-flag), passed through to the aggregation kernel which uses the new per-env i32 buffer in place of the per-bar `sr > 0` predicate when wired. Per `feedback_no_partial_refactor` the kernel signature change + all callers (production + 2 test files) migrate atomically in this commit; the kernel-body behavioral change + producer wire-up + reset registry + audit-doc entry land atomically in the immediately following commit. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../cuda_pipeline/gpu_experience_collector.rs | 9 + .../cuda_pipeline/sp20_aggregate_inputs.rs | 15 ++ .../sp20_aggregate_inputs_kernel.cu | 28 ++- crates/ml/tests/sp20_aggregate_inputs_test.rs | 184 ++++++++++++++++++ crates/ml/tests/sp20_phase1_4_wireup_test.rs | 8 + docs/dqn-wire-up-audit.md | 117 +++++++++++ 6 files changed, 360 insertions(+), 1 deletion(-) diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 1263469f1..3b2d95885 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -6508,6 +6508,14 @@ 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. unsafe { launch_sp20_aggregate_inputs( &self.stream, @@ -6521,6 +6529,7 @@ impl GpuExperienceCollector { stats_std_dev, aux_dir_acc_dev, alpha_per_env_dev, + /* is_win_per_env_dev = */ 0, n_envs_i32_sp20, dir_divisor, hold_action_id, diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs index 0804124ae..c4f585094 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs @@ -134,6 +134,16 @@ pub unsafe fn launch_sp20_aggregate_inputs( // kernel falls back to `alpha=0.0f` in that case, matching Phase // 1.4 placeholder behavior). alpha_per_env_dev: u64, + // SP20 Phase 2 Task 2.2-fix (2026-05-10): per-env trade-close win + // flag (i32, 1 = segment_return > 0, 0 = otherwise). When `0` + // (NULL), the kernel falls back to the legacy `step_ret_per_env > + // 0` predicate. Production callers pass the + // `GpuExperienceCollector.is_win_per_env` device pointer; tests + // that don't wire a segment-level producer pass `0` to preserve + // the pre-fix oracle-test contract. Fixes WR_EMA pinned at 0 — + // the per-bar `step_ret` at close bars is dominated by tx-cost so + // the legacy predicate was essentially always false. + is_win_per_env_dev: u64, n_envs: i32, dir_divisor: i32, hold_action_id: i32, @@ -152,6 +162,10 @@ pub unsafe fn launch_sp20_aggregate_inputs( * gracefully falls back to alpha=0.0f. This preserves the Phase 1.4 * test scaffold contract for `sp20_phase1_4_wireup_test` paths * that don't wire the SP20 reward producer. */ + /* is_win_per_env_dev is allowed to be 0 (NULL) — the kernel falls + * back to `step_ret > 0` per-bar predicate (legacy aggregate-kernel + * oracle-test contract). Production caller passes the per-env + * device buffer populated at segment_complete from segment_return. */ debug_assert!(out_inputs_dev != 0); debug_assert!(n_envs > 0, "n_envs must be > 0, got {n_envs}"); debug_assert!(env_stride > 0, "env_stride must be > 0, got {env_stride}"); @@ -176,6 +190,7 @@ pub unsafe fn launch_sp20_aggregate_inputs( .arg(&aux_logits_std_dev) .arg(&aux_dir_acc_dev) .arg(&alpha_per_env_dev) + .arg(&is_win_per_env_dev) .arg(&n_envs) .arg(&dir_divisor) .arg(&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 80b514caa..d719a9b76 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -12,7 +12,16 @@ * * is_close = 1 if any env's trade_close_per_env[env] != 0, else 0. * is_win = 1 if (is_close == 1) AND fraction of closed envs with - * step_ret_per_env[env] > 0 is >= 0.5, else 0. + * `is_win_per_env[env] != 0` is >= 0.5, else 0. + * SP20 Phase 2 Task 2.2-fix (2026-05-10): the predicate + * now reads the segment-level per-env `is_win_per_env` + * flag (1 iff `segment_return > 0` at trade close) rather + * than the per-bar `step_ret_per_env > 0` proxy. The + * per-bar step_return at close bars is dominated by + * tx-cost so `sr > 0` was almost always false, pinning + * `WR_EMA` at 0 in production despite real wins. NULL- + * tolerant fallback to the legacy predicate preserved for + * oracle tests without a segment-level producer wired. * trade_duration = round(mean(hold_at_exit_per_env[env]) for env * where trade_close_per_env[env] != 0) * if is_close == 1, else 0. @@ -141,6 +150,23 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( * kernel emits `out_inputs->alpha = 0.0f` (matches Phase 1.4 * behavior, used by tests that haven't wired the alpha producer). */ const float* __restrict__ alpha_per_env, /* [n_envs] or NULL */ + /* SP20 Phase 2 Task 2.2-fix (2026-05-10) — per-env trade-close + * win flag. `[n_envs]` i32 device buffer; the env's slot is read + * only when `trade_close_per_env[env * env_stride] != 0`. Each + * slot is `1` if the segment's realized P&L was strictly positive + * at trade close, else `0`. Replaces the broken `sr > 0` per-bar + * step-return predicate (which conflates the close bar's per-bar + * MTM tick with trade outcome — at close bars the per-bar tick is + * dominated by tx-cost so the `sr > 0` predicate is essentially + * always false, pinning WR_EMA at 0). + * + * NULL-tolerant — when NULL the kernel falls back to the legacy + * `step_ret_per_env[env] > 0` predicate (matches the pre-2.2-fix + * behavior used by aggregate-kernel oracle tests that don't wire + * the segment-level producer). Production callers pass non-NULL + * with the per-env buffer populated by `experience_env_step`'s + * segment_complete branch as `(segment_return > 0.0f) ? 1 : 0`. */ + const int* __restrict__ is_win_per_env, /* [n_envs] or NULL */ int n_envs, /* Action decoding parameters (production packed factored action * encoding). `dir_divisor = NUM_MAGNITUDES * NUM_ORD * NUM_URG`, diff --git a/crates/ml/tests/sp20_aggregate_inputs_test.rs b/crates/ml/tests/sp20_aggregate_inputs_test.rs index 70b111bcd..d27254573 100644 --- a/crates/ml/tests/sp20_aggregate_inputs_test.rs +++ b/crates/ml/tests/sp20_aggregate_inputs_test.rs @@ -145,6 +145,7 @@ mod gpu { std_buf.dev_ptr, dir_buf.dev_ptr, /* alpha_per_env_dev = */ 0, // NULL — test scaffold without SP20 reward producer + /* is_win_per_env_dev = */ 0, // NULL — falls back to `sr > 0` legacy predicate /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -213,6 +214,81 @@ mod gpu { std_buf.dev_ptr, dir_buf.dev_ptr, alpha_buf.dev_ptr, // SP20 Phase 2 Task 2.2 alpha producer + /* is_win_per_env_dev = */ 0, // NULL — alpha-only variant; legacy `sr > 0` predicate + /* n_envs = */ n as i32, + DIR_DIVISOR, + HOLD_ACTION_ID, + out_buf.dev_ptr, + ).expect("launch aggregate kernel"); + } + stream.synchronize().expect("sync after aggregate"); + read_ema_inputs(&out_buf) + } + + /// SP20 Phase 2 Task 2.2-fix (2026-05-10) variant of `run_kernel` + /// that exercises the new `is_win_per_env` per-env trade-close win + /// flag (i32, 1 = segment_return > 0, 0 = otherwise). Used by + /// `wr_ema_uses_segment_pnl_via_is_win_per_env_predicate` to + /// validate the WR_EMA-fix contract: when `is_win_per_env` is wired, + /// `wins_count` derives from segment-level outcome rather than the + /// per-bar `step_ret > 0` predicate at close (which is dominated by + /// tx-cost so almost always false at close bars, pinning WR_EMA at + /// 0 in production). + fn run_kernel_with_is_win( + trade_close: &[i32], + step_ret: &[f32], + hold_at_exit: &[f32], + actions: &[i32], + is_win_per_env: &[i32], + aux_logits_p50: f32, + aux_logits_std: f32, + aux_dir_acc: f32, + ) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + let stream = make_test_stream(); + let kernel = load_aggregation_kernel(&stream); + let n = trade_close.len(); + assert!(n > 0, "test n_envs must be > 0"); + assert_eq!(step_ret.len(), n); + assert_eq!(hold_at_exit.len(), n); + assert_eq!(actions.len(), n); + assert_eq!(is_win_per_env.len(), n); + + let tc_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc tc"); + tc_buf.write_from_slice(trade_close); + let sr_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc sr"); + sr_buf.write_from_slice(step_ret); + let he_buf = unsafe { MappedF32Buffer::new(n) }.expect("alloc he"); + he_buf.write_from_slice(hold_at_exit); + let act_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc act"); + act_buf.write_from_slice(actions); + let win_buf = unsafe { MappedI32Buffer::new(n) }.expect("alloc is_win"); + win_buf.write_from_slice(is_win_per_env); + + let p50_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc p50"); + p50_buf.write_from_slice(&[aux_logits_p50]); + let std_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc std"); + std_buf.write_from_slice(&[aux_logits_std]); + let dir_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc dir"); + dir_buf.write_from_slice(&[aux_dir_acc]); + + let out_buf = unsafe { MappedF32Buffer::new(SP20_EMA_INPUTS_F32_LEN) } + .expect("alloc out"); + out_buf.write_from_slice(&[f32::NAN; SP20_EMA_INPUTS_F32_LEN]); + + unsafe { + launch_sp20_aggregate_inputs( + &stream, + &kernel, + tc_buf.dev_ptr, + sr_buf.dev_ptr, + he_buf.dev_ptr, + act_buf.dev_ptr, + /* env_stride = */ 1, + p50_buf.dev_ptr, + std_buf.dev_ptr, + dir_buf.dev_ptr, + /* alpha_per_env_dev = */ 0, + win_buf.dev_ptr, // SP20 Phase 2 Task 2.2-fix is_win producer /* n_envs = */ n as i32, DIR_DIVISOR, HOLD_ACTION_ID, @@ -414,4 +490,112 @@ mod gpu { values must NOT leak into the EMA on non-close steps", ); } + + /// SP20 Phase 2 Task 2.2-fix (2026-05-10) — WR_EMA pinning bug. + /// + /// **Root cause** (production HEALTH_DIAG observed `wr_ema = 0.0000` + /// at every epoch while `alpha_ema` evolved with real values): + /// the aggregation kernel's win predicate was + /// `step_ret_per_env[env] > 0.0f` at the close bar, which is the + /// per-bar mark-to-market change `(new_value - prev_equity) / + /// prev_equity` of the closing bar, NOT the trade segment outcome. + /// + /// At close bars `step_ret_core ≈ position * Δprice * mult − + /// tx_cost`. The per-bar tick `Δprice` is small (one bar) while + /// `tx_cost` is fixed at trade-close → `step_ret_core < 0` is the + /// dominant case across closing bars even for trades that closed + /// profitably overall. Result: `wins_count = 0` always, `is_win = + /// 0` always, `WR_EMA` pinned at 0 across all training epochs, + /// LOSS_CAP pinned at -1.0 (the WR=0.5 cold-start cap, never + /// ramping up to -2.0 per spec §4.1). + /// + /// **Fix**: derive `is_win` from the segment-level realized P&L + /// signed by trade close, plumbed through a new `is_win_per_env` + /// per-env i32 buffer that the kernel reads in place of the + /// per-bar `step_ret > 0` predicate. This test pins the contract: + /// when `is_win_per_env` is wired (non-NULL) the kernel's + /// `wins_count` MUST equal the count of closed envs with + /// `is_win_per_env[env] != 0`, regardless of `step_ret`'s sign. + /// + /// **Why this test catches the bug**: the scenario is the realistic + /// production case where `step_ret < 0` for all 4 closing envs (the + /// per-bar tick at close is dominated by tx-cost) BUT 3 of 4 + /// trades were profitable overall (`is_win_per_env = [1, 1, 1, + /// 0]`). With the legacy `sr > 0` predicate `wins_count = 0` ⇒ + /// `is_win = 0`. With the fix `wins_count = 3` ⇒ `2 * 3 >= 4` ⇒ + /// `is_win = 1`. + /// + /// **TDD note**: in commit 1 the kernel signature accepts + /// `is_win_per_env` but the kernel BODY still uses `sr > 0`, so + /// this test FAILS at the assertion (kernel emits `is_win = 0` + /// despite `is_win_per_env = [1,1,1,0]`). Commit 2's kernel body + /// change closes the gap and this test passes. Per the + /// `feedback_no_partial_refactor` exception for TDD-style commits + /// — the gap is closed in the IMMEDIATE next commit. + #[test] + #[ignore = "requires GPU"] + fn wr_ema_uses_segment_pnl_via_is_win_per_env_predicate() { + // 4 envs, all closing this step. + let trade_close = vec![1_i32; 4]; + + // step_ret all negative (tx-cost dominated per-bar return at + // close bars — the realistic production case). With the legacy + // `sr > 0` predicate, `wins_count = 0` so `is_win = 0` (bug). + let step_ret = vec![-0.001_f32; 4]; + + let hold_at_exit = vec![5.0_f32; 4]; + let actions = vec![encode_dir(0); 4]; // not Hold + + // 3 of 4 segment-level wins (segment_return > 0). With the + // fix, the kernel reads is_win_per_env in the close-gated path + // and accumulates `wins_count = 3`. + let is_win_per_env = vec![1_i32, 1, 1, 0]; + + let (ic, iw, _, _, _, _, _, _, _) = run_kernel_with_is_win( + &trade_close, &step_ret, &hold_at_exit, &actions, + &is_win_per_env, 0.0, 0.0, 0.0, + ); + + assert_eq!(ic, 1, "is_close = 1 (4 envs closed)"); + assert_eq!( + iw, 1, + "is_win MUST be 1 — segment-level wins (is_win_per_env = \ + [1,1,1,0]) put 3 of 4 closed envs in the win bucket; \ + 2*3 >= 4 ⇒ majority threshold met. The legacy `step_ret \ + > 0` predicate would have given is_win = 0 (all step_rets \ + negative due to tx-cost domination at close bars), pinning \ + WR_EMA at 0 in production.", + ); + } + + /// SP20 Phase 2 Task 2.2-fix — the kernel's NULL-tolerance for + /// `is_win_per_env` MUST preserve the legacy `sr > 0` predicate + /// for the aggregate-kernel oracle tests + Phase 1.4 wireup test + /// scaffolds that don't wire a segment-level producer. This test + /// pins that contract: with `is_win_per_env` NULL and step_ret all + /// positive (fake "all wins"), the legacy predicate emits + /// `wins_count = n_envs` ⇒ `is_win = 1`. + #[test] + #[ignore = "requires GPU"] + fn null_is_win_per_env_falls_back_to_legacy_step_ret_predicate() { + // All envs close with positive step_ret (the contrived oracle- + // test signal that the legacy predicate honors). + let trade_close = vec![1_i32; 4]; + let step_ret = vec![0.5_f32; 4]; // legacy predicate fires + let hold_at_exit = vec![3.0_f32; 4]; + let actions = vec![encode_dir(0); 4]; + + let (ic, iw, _, _, _, _, _, _, _) = + run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, + 0.0, 0.0, 0.0); + + assert_eq!(ic, 1); + assert_eq!( + iw, 1, + "NULL is_win_per_env falls back to `sr > 0` legacy \ + predicate — all 4 step_rets positive ⇒ wins_count = 4 ⇒ \ + is_win = 1 (the pre-fix oracle-test contract preserved \ + via NULL-tolerance).", + ); + } } diff --git a/crates/ml/tests/sp20_phase1_4_wireup_test.rs b/crates/ml/tests/sp20_phase1_4_wireup_test.rs index 699ddf2f4..7f381c0e2 100644 --- a/crates/ml/tests/sp20_phase1_4_wireup_test.rs +++ b/crates/ml/tests/sp20_phase1_4_wireup_test.rs @@ -170,6 +170,14 @@ mod gpu { // collector path. Tests in this file // assert the placeholder contract // (alpha_ema stays at sentinel 0.0). + /* is_win_per_env_dev = */ 0, // SP20 Phase 2 Task 2.2-fix + // (2026-05-10) — wireup test stays + // NULL; the kernel falls back to + // the `sr > 0` legacy predicate so + // existing wireup-test assertions + // (e.g. WR_EMA = 1.0 with one + // winning trade at step_ret > 0) + // remain valid. n_envs as i32, DIR_DIVISOR, HOLD_ACTION_ID, diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 67c92e12e..25e69b58f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,123 @@ **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 — failing test commit (signature plumb only) + +Plumbs a new `is_win_per_env` per-env i32 device-buffer arg through +the `sp20_aggregate_inputs_kernel` signature + `launch_sp20_aggregate_ +inputs` Rust launcher, NULL-tolerant so existing oracle-test scaffolds +that don't wire a segment-level producer keep using the legacy `step_ +ret_per_env > 0` predicate. Production caller (gpu_experience_collector) +passes `0` (NULL) in this commit; the kernel BODY still uses the +legacy predicate. The atomic kernel-body change + per-env producer + +collector buffer + reset entry land in the immediately following +commit per `feedback_no_partial_refactor`. + +### Bug context + +Production HEALTH_DIAG observed `wr_ema=0.0000` at every epoch across +17+ epochs of an L40S smoke (`train-multi-seed-4cqts`), while +`alpha_ema` evolved with real values (0.055 → -0.032 → +0.022). Both +EMAs share the same close gate (`is_close == 1`), so the gate IS +firing — alpha aggregation is correct. The pinning is in the +`is_win` aggregation predicate. + +### Root cause (verified by code analysis) + +`sp20_aggregate_inputs_kernel.cu:184,193` accumulates `wins_count` +via `step_ret_per_env[env * env_stride] > 0.0f` at trade-close bars. +But `step_ret_per_sample[out_off]` is `step_ret_core` from +`unified_env_step_core` (per-bar `(new_value - prev_equity) / +prev_equity`), NOT segment-level realized P&L. + +At a close bar with mark-to-market accounting: + `new_value - prev_equity = position * (close_t - close_{t-1}) - tx_cost` + +The per-bar tick `Δprice` is small (one bar's price movement) while +`tx_cost` is fixed at trade close. Result: `step_ret < 0` is the +dominant case at close bars even when the trade itself was profitable +(`segment_pnl > 0`). The predicate is broken. + +Confirmation: `alpha_ema` reads `R_event` derived from `segment_ +return = (realized_pnl + unrealized_at_exit - trade_start_pnl) / +prev_equity` at line 3187 — the proper trade P&L. So the producer +HAS the correct trade-outcome signal; the aggregation kernel just +isn't reading the right thing for the wins counter. + +### Fix architecture (this commit + next) + +This commit (commit 1, failing test): +- Adds `const int* is_win_per_env` arg to kernel signature (NULL- + tolerant — falls back to legacy `sr > 0` predicate when 0). +- Adds `is_win_per_env_dev: u64` arg to launcher. +- Updates 2 test files (`sp20_aggregate_inputs_test.rs`, `sp20_phase + 1_4_wireup_test.rs`) to pass `0` for the new arg, preserving their + oracle-test contract. +- Updates production caller (gpu_experience_collector) to pass `0` + in this commit (the producer + buffer arrive in commit 2). +- Adds 2 new GPU oracle tests in `sp20_aggregate_inputs_test`: + - `wr_ema_uses_segment_pnl_via_is_win_per_env_predicate` — + constructs the realistic production scenario (step_ret all + negative due to tx-cost domination, but 3 of 4 trades profitable + overall via `is_win_per_env = [1,1,1,0]`); asserts `is_win = 1`. + FAILS with the legacy `sr > 0` predicate, PASSES once commit 2 + lands the kernel-body change. + - `null_is_win_per_env_falls_back_to_legacy_step_ret_predicate` — + pins the NULL-tolerance contract for oracle tests. + +Commit 2 (next, atomic kernel-body fix): +- Allocates `is_win_per_env: CudaSlice` `[alloc_episodes]` in + `GpuExperienceCollector` (mirror of `alpha_per_env`). +- Wires `experience_env_step::segment_complete` to write + `is_win_per_env[i] = (segment_return > 0.0f) ? 1 : 0` at the same + site as `alpha_per_env[i] = alpha`. +- Updates kernel body to read `is_win_per_env[env]` in the + `tc != 0` branch when `is_win_per_env != NULL`, falling back to + `sr > 0.0f` when NULL. +- Registers `is_win_per_env` as `FoldReset` in + `state_reset_registry.rs` with the matching dispatch arm in + `training_loop.rs::reset_state` (mirror of `alpha_per_env`'s + reset path). + +### Pearls + invariants honoured + +- `feedback_no_partial_refactor` — kernel signature + all callers + migrate atomically in commit 1; kernel-body behavioral change + + producer + buffer + reset registry land atomically in commit 2. + The 2-commit split is for TDD evidence (commit 2's body change is + the minimal diff that makes commit 1's failing test pass) and is + the exception carve-out per the user-requested test-first + protocol. +- `feedback_no_stubs` — the legacy `sr > 0` predicate REMAINS as the + NULL-tolerant fallback; existing oracle tests that don't wire a + segment-level producer keep working. Production caller is wired + for-real in commit 2. +- `pearl_audit_unboundedness_for_implicit_asymmetry` — the bug class + is "implicit assumption that per-bar return tracks trade outcome"; + the fix preserves segment-level boundedness via direct sign read + rather than indirectly via the per-bar proxy. + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` | +arg | Kernel signature: `is_win_per_env` (unused in body — commit 2 adds the read site) | +| `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` | +arg | Launcher signature mirrors kernel | +| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +0-arg | Production caller passes `0` (NULL) — commit 2 wires the buffer | +| `crates/ml/tests/sp20_aggregate_inputs_test.rs` | +2 tests, +1 helper | Failing test for the bug + NULL-tolerance pin | +| `crates/ml/tests/sp20_phase1_4_wireup_test.rs` | +0-arg | Wireup test passes `0` (NULL) — preserves WR_EMA = 1.0 expectation via legacy predicate | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml --tests +SQLX_OFFLINE=true cargo test -p ml --lib sp20_aggregate_inputs +``` + +Lib tests pass; new GPU oracle test fails on current code (kernel +body still uses legacy predicate). Commit 2 closes the gap. + ## 2026-05-09 — SP20 Phase 2 Task 2.2: SP12 v3 reward block replacement + per-env alpha plumbing (atomic) Lands the SP20 4-quadrant reward at the `experience_env_step::