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

test(sp20): failing test pinning WR_EMA bug — wins predicate must use segment P&L not per-bar step_ret

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-10 01:20:13 +02:00
parent e1aef53373
commit 5a29c37cd8
6 changed files with 360 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

15
crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs
View File

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

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

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

184
crates/ml/tests/sp20_aggregate_inputs_test.rs
View File

@@ -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).",
);
}
}

8
crates/ml/tests/sp20_phase1_4_wireup_test.rs
View File

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