diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 014d60284..33b304715 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -349,6 +349,16 @@ fn main() { // segment_complete branch; the test wrapper has no production // callers. "sp12_reward_math_test_kernel.cu", + // SP20 Phase 2 Task 2.1 (2026-05-09): standalone test wrapper for + // the 4-quadrant directional-correctness reward function defined + // in `sp20_reward.cuh`. Wraps `sp20_compute_event_reward` so the + // Rust GPU oracle tests in `tests/sp20_event_reward_test.rs` can + // drive each quadrant of the truth table on synthetic inputs with + // analytically-known expected outputs per + // `feedback_no_cpu_test_fallbacks`. Producer call site lands in + // Task 2.2 inside `experience_env_step::segment_complete` branch; + // the test wrapper has no production callers. + "sp20_event_reward_test_kernel.cu", // SP4 Layer A Task A5 (2026-04-30): first end-to-end SP4 producer. // Single-block kernel reads `denoise_target_q_buf`, computes // p99(|target_q|) via `sp4_histogram_p99<256>`, writes the step diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 0c8fd95d5..32f58d679 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -86,6 +86,18 @@ * Used by both experience_state_gather and backtest_state_gather. */ #include "state_layout.cuh" +/* SP20 Phase 2 Task 2.1 (2026-05-09): 4-quadrant directional-correctness + * event reward at trade close. Header-only `__device__ __forceinline__` + * function so the test wrapper kernel + * (`sp20_event_reward_test_kernel.cu`) and the production caller (Task + * 2.2's trade-close site replacement) share the exact math. Per + * `feedback_no_partial_refactor.md` the function ships in this header + * BEFORE the production call site lands — Task 2.1 oracle tests + * exercise the math in isolation; Task 2.2 wires it into + * `experience_env_step`'s segment_complete branch atomically with the + * SP12 v3 reward block deletion + alpha plumbing. */ +#include "sp20_reward.cuh" + /* ------------------------------------------------------------------ */ /* BF16 support — common_device_functions.cuh is prepended by build.rs */ /* ------------------------------------------------------------------ */ diff --git a/crates/ml/src/cuda_pipeline/sp20_event_reward_test_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_event_reward_test_kernel.cu new file mode 100644 index 000000000..d5c60161e --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_event_reward_test_kernel.cu @@ -0,0 +1,47 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP20 Phase 2 Task 2.1 (2026-05-09) — sp20_event_reward GPU oracle test wrapper. + * + * Standalone test wrapper that exercises the `sp20_compute_event_reward` + * device function defined in `sp20_reward.cuh`. One thread per kernel + * launch is sufficient — the function is pure register-only arithmetic + * with no shared state. The wrapper mirrors Task 2.2's production call + * site bit-for-bit so the tests verify the exact math the production + * reward chain executes (no CPU reference impl per + * `feedback_no_cpu_test_fallbacks`). + * + * Outputs go to a mapped-pinned buffer with `__threadfence_system()` + * before thread exit so the host can read via `read_volatile` after a + * stream sync. Same pattern as `sp12_reward_math_test_kernel.cu` / + * `thompson_test_kernel.cu` / `sp4_histogram_p99_test_kernel.cu`. + * + * Test-only: no production callers. Cubin built by `crates/ml/build.rs` + * (kernel name `sp20_event_reward_test_kernel`). The production call + * site lands in Task 2.2 inside `experience_env_step`'s + * `segment_complete` branch. + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +#include "sp20_reward.cuh" /* sp20_compute_event_reward */ + +extern "C" __global__ void sp20_event_reward_test_kernel( + /* Single (input, expected-output) triple per launch. */ + float close_pnl, + int label_at_open_sign, + float loss_cap, + float* __restrict__ out_r_event) +{ + /* Single-thread launch — the device function has no parallelism to + * exploit (each test invocation drives one input → output pair). */ + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + out_r_event[0] = sp20_compute_event_reward( + close_pnl, label_at_open_sign, loss_cap + ); + + /* Make the write PCIe-visible to the host pointer (mapped-pinned + * buffer). The host reads via `read_volatile` after a stream sync; + * the fence guarantees the write has actually drained from the GPU + * write buffers before the sync returns. */ + __threadfence_system(); +} diff --git a/crates/ml/src/cuda_pipeline/sp20_reward.cuh b/crates/ml/src/cuda_pipeline/sp20_reward.cuh new file mode 100644 index 000000000..2e7de5ff2 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_reward.cuh @@ -0,0 +1,130 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP20 Phase 2 Task 2.1 (2026-05-09) — sp20_compute_event_reward. + * + * Component 1 / 4-quadrant directional-correctness reward at trade close. + * Per spec `docs/superpowers/specs/2026-05-09-sp19-20-wr-first-design.md` §4.1. + * + * Forces "right reason for right outcome" via the SP19 multi-horizon label + * sign agreement check. Loss cap is adaptive via + * `ISV[LOSS_CAP_INDEX = 510]` (ramps from −1 to −2 with WR_EMA per spec + * §4.1; produced by `sp20_controllers_compute_kernel`). + * + * 4-quadrant truth table: + * + * pnl_sign dir_match R_event + * ─────────────────────────────────── + * +1 true +win_cap = +1.0 (right reason, right outcome) + * +1 false +0.5 (wrong reason, right outcome) + * -1 true -0.5 (right reason, wrong outcome) + * -1 false loss_cap (wrong reason, wrong outcome — + * ISV-driven asymmetric cap; + * [-2.0, -1.0] range per + * pearl_audit_unboundedness_for_ + * implicit_asymmetry) + * 0 * 0.0 (no information — close_pnl == 0) + * + * `label_at_open_sign` follows the SP20 sign convention: `+1` = aux + * predicted up at trade open, `-1` = aux predicted not-up, `0` = sentinel + * (no info / aux skip-window). The dir_match check requires + * `pnl_sign == label_at_open_sign` for AGREEMENT — when the sentinel `0` + * lands here paired with non-zero `pnl_sign`, dir_match is necessarily + * false (the trade closed but we have no aux signal to confirm the + * reasoning was correct). This is the wrong-reason quadrant: a +0.5 + * shoulder reward on a winning trade with no aux signal, or a `loss_cap` + * penalty on a losing trade with no aux signal. Safer default than + * over-rewarding a no-signal close. + * + * Pearl/feedback notes: + * - `pearl_one_unbounded_signal_per_reward.md` — the only multiplicand + * here is the constant cap (win_cap = +1.0, the +0.5 shoulder, the + * -0.5 right-reason-wrong-outcome, or `loss_cap ∈ [-2, -1]`). Every + * branch returns a structurally bounded value, so the entire output + * sits in `[loss_cap, +1]` ⊆ `[-2, +1]` regardless of `close_pnl`'s + * magnitude. + * - `pearl_event_driven_reward_density_alignment.md` — fires only at + * trade close (`is_close == 1`); per-bar density is handled by + * Component 2 (Phase 3.2 Hold opportunity-cost dual emission). + * - `feedback_isv_for_adaptive_bounds.md` — `loss_cap` is read from + * `ISV[LOSS_CAP_INDEX]` by the caller and passed in; this header is + * decoupled from slot indices. + * - `feedback_no_partial_refactor.md` — the function is shipped in this + * header BEFORE the production caller (Task 2.2) so Task 2.1's + * GPU oracle tests can exercise the math in isolation; Task 2.2 + * lands the call site at the SP12 v3 trade-close block replacement + * atomically with the per-env alpha plumbing. + * + * Single-block / single-thread / register-only callers — no shared + * memory, no atomics, no host branches. Safe to inline at the + * `experience_env_step` site. + * ══════════════════════════════════════════════════════════════════════════ */ +#ifndef SP20_REWARD_CUH +#define SP20_REWARD_CUH + +#include + +/* Win cap is structurally `+1.0` per spec §4.1 — the upper bound of + * the 4-quadrant table. Loss cap is adaptive (caller passes the ISV + * value); win cap stays a compile-time constant because the upside is + * bounded by directional-correctness alone (no asymmetric loss-aversion + * concept on the win side). + * + * Hardcoded inline rather than macro'd because this is the single + * canonical anchor and a one-shot read in the only-caller pattern. Per + * `feedback_no_quickfixes` carve-out for Invariant-1 anchors. */ +#define SP20_WIN_CAP 1.0f + +/* `+0.5` is the wrong-reason-right-outcome shoulder. `-0.5` is the + * right-reason-wrong-outcome shoulder. Both are spec §4.1 anchors + * (table rows 2 and 3). The 0.5 fraction is calibrated to the + * `win_cap = 1.0` baseline as "half credit" — the spec discusses the + * shoulder reward as a partial signal that survives ablations of the + * label-agreement check. Same Invariant-1 carve-out. */ +#define SP20_SHOULDER_RIGHT_OUTCOME 0.5f +#define SP20_SHOULDER_WRONG_OUTCOME -0.5f + +/* SP20 Component 1 — 4-quadrant directional reward at trade close. + * + * Args: + * close_pnl — segment realized P&L (any unit; only the + * sign matters per spec §4.1). + * label_at_open_sign — `{-1, 0, +1}` aux next-bar label sign at + * trade open (Task 2.0 buffer). + * loss_cap — `ISV[LOSS_CAP_INDEX]`; ramps from −1 to −2 + * with WR_EMA per spec §4.1 + * (`sp20_controllers_compute_kernel` output). + * + * Returns: bounded scalar `R_event ∈ [loss_cap, SP20_WIN_CAP]`. */ +__device__ __forceinline__ float sp20_compute_event_reward( + float close_pnl, + int label_at_open_sign, + float loss_cap) +{ + /* close_pnl == 0 ⇒ no information. + * Use exact equality against 0.0f rather than a tolerance epsilon — + * `close_pnl` is computed downstream from a discrete realized-P&L + * accumulator (`PS_REALIZED_PNL`) plus continuous unrealized at exit, + * and a true-zero close is the natural "didn't move" case (e.g., + * tiny lots fully offset by tx cost). A near-zero (1e-6) close is + * still informative — it pushes through to either the `right_outcome` + * or `wrong_outcome` quadrant by sign. */ + int pnl_sign; + if (close_pnl > 0.0f) pnl_sign = +1; + else if (close_pnl < 0.0f) pnl_sign = -1; + else pnl_sign = 0; + if (pnl_sign == 0) return 0.0f; + + /* Direction match: requires `label_at_open_sign` to MATCH `pnl_sign`. + * When the SP20 sentinel 0 lands here (no info / aux skip-window + * at trade open), dir_match is necessarily false because pnl_sign is + * already known non-zero; the trade routes to the wrong-reason + * quadrant. This is the safer default per the audit-doc rationale. */ + const bool dir_match = (pnl_sign == label_at_open_sign); + + if (pnl_sign > 0 && dir_match) return SP20_WIN_CAP; /* +1.0 */ + if (pnl_sign > 0 && !dir_match) return SP20_SHOULDER_RIGHT_OUTCOME; /* +0.5 */ + if (pnl_sign < 0 && dir_match) return SP20_SHOULDER_WRONG_OUTCOME; /* -0.5 */ + /* pnl_sign < 0 && !dir_match ⇒ wrong reason, wrong outcome. */ + return loss_cap; +} + +#endif /* SP20_REWARD_CUH */ diff --git a/crates/ml/tests/sp20_event_reward_test.rs b/crates/ml/tests/sp20_event_reward_test.rs new file mode 100644 index 000000000..aa7f95446 --- /dev/null +++ b/crates/ml/tests/sp20_event_reward_test.rs @@ -0,0 +1,254 @@ +#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. + +//! SP20 Phase 2 Task 2.1 — `sp20_compute_event_reward` GPU oracle tests. +//! +//! Validates the 4-quadrant directional-correctness event reward defined +//! in `crates/ml/src/cuda_pipeline/sp20_reward.cuh` (spec §4.1). Each +//! test drives a synthetic `(close_pnl, label_at_open_sign, loss_cap)` +//! triple through the production device function via the standalone +//! wrapper `sp20_event_reward_test_kernel.cu` and verifies the +//! GPU-computed result against the analytically-known expected value. +//! Per `feedback_no_cpu_test_fallbacks`: GPU oracle only, no CPU +//! reference impl. +//! +//! All tests are `#[ignore = "requires GPU"]`-gated to match every other +//! GPU oracle test in this crate (sp4 / sp5 / sp11 / sp12). Run on a +//! GPU host: +//! +//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \ +//! cargo test -p ml --test sp20_event_reward_test --features cuda \ +//! -- --ignored --nocapture +//! +//! ## 4-quadrant truth table (spec §4.1) +//! +//! | `pnl_sign` | `dir_match` | `R_event` | Test | +//! |------------|--------------------|---------------|--------------------------------------------------------------| +//! | `+1` | `true` | `+1.0` | `right_reason_right_outcome` | +//! | `+1` | `false` | `+0.5` | `wrong_reason_right_outcome` | +//! | `-1` | `true` | `-0.5` | `right_reason_wrong_outcome` | +//! | `-1` | `false` | `loss_cap` | `wrong_reason_wrong_outcome_loss_cap_at_minus_1` / `..._minus_2` | +//! | `0` | `*` | `0.0` | `zero_pnl_returns_zero` | + +#![cfg(feature = "cuda")] + +use std::sync::Arc; + +use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg}; +use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; + +/// Test-only cubin built by `crates/ml/build.rs`. +const SP20_EVENT_REWARD_CUBIN: &[u8] = include_bytes!(concat!( + env!("OUT_DIR"), + "/sp20_event_reward_test_kernel.cubin" +)); + +/// Resolve a CUDA stream against device 0. Mirrors `make_test_stream` in +/// every other oracle test in this crate. +fn make_test_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() +} + +/// Load the test kernel function from the embedded cubin. +fn load_test_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP20_EVENT_REWARD_CUBIN.to_vec()) + .expect("load sp20_event_reward_test_kernel cubin"); + module + .load_function("sp20_event_reward_test_kernel") + .expect("load sp20_event_reward_test_kernel function") +} + +/// Drive one `(close_pnl, label_at_open_sign, loss_cap) → R_event` +/// triple through the device function. Single-thread / single-block +/// launch — the function is pure register arithmetic, no parallelism +/// to exploit. +/// +/// Returns the GPU-computed `R_event` scalar. +fn launch_sp20_event_reward( + stream: &Arc, + f: &CudaFunction, + close_pnl: f32, + label_at_open_sign: i32, + loss_cap: f32, +) -> f32 { + let out = unsafe { MappedF32Buffer::new(1) }.expect("alloc out_r_event (1 f32)"); + out.write_from_slice(&[f32::NAN]); // poison-fill so test fails loud if kernel didn't write + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let out_dev = out.dev_ptr; + + unsafe { + stream + .launch_builder(f) + .arg(&close_pnl) + .arg(&label_at_open_sign) + .arg(&loss_cap) + .arg(&out_dev) + .launch(cfg) + .expect("launch sp20_event_reward_test_kernel"); + } + stream + .synchronize() + .expect("sync after sp20_event_reward_test_kernel"); + + out.read_all()[0] +} + +// ── 4-quadrant truth table ───────────────────────────────────────────── + +/// Quadrant 1: `pnl_sign=+1`, `dir_match=true` ⇒ `R_event = +SP20_WIN_CAP = +1.0`. +/// "Right reason for right outcome" — the canonical maximum reward. +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_quadrant_right_reason_right_outcome() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ 0.5, /* label_at_open_sign =*/ 1, /* loss_cap =*/ -1.0, + ); + assert!( + (r - 1.0).abs() < 1e-6, + "expected R_event = +1.0 (right reason, right outcome); got {r}", + ); +} + +/// Quadrant 2: `pnl_sign=+1`, `dir_match=false` ⇒ `R_event = +0.5`. +/// "Wrong reason, right outcome" — partial reward for a winning trade +/// the aux head did NOT predict (lucky win). +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_quadrant_wrong_reason_right_outcome() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ 0.3, /* label_at_open_sign =*/ -1, /* loss_cap =*/ -1.0, + ); + assert!( + (r - 0.5).abs() < 1e-6, + "expected R_event = +0.5 (wrong reason, right outcome); got {r}", + ); +} + +/// Quadrant 3: `pnl_sign=-1`, `dir_match=true` ⇒ `R_event = -0.5`. +/// "Right reason, wrong outcome" — partial penalty for a losing trade +/// the aux head DID predict correctly (correct call, market disagreed). +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_quadrant_right_reason_wrong_outcome() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ -0.2, /* label_at_open_sign =*/ -1, /* loss_cap =*/ -1.0, + ); + assert!( + (r - (-0.5)).abs() < 1e-6, + "expected R_event = -0.5 (right reason, wrong outcome); got {r}", + ); +} + +/// Quadrant 4 cold-start: `pnl_sign=-1`, `dir_match=false`, +/// `loss_cap=-1.0` (early training, WR_EMA below 50%) ⇒ `R_event = -1.0`. +/// "Wrong reason, wrong outcome" — full loss cap. +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_1() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ -0.4, /* label_at_open_sign =*/ 1, /* loss_cap =*/ -1.0, + ); + assert!( + (r - (-1.0)).abs() < 1e-6, + "expected R_event = -1.0 (wrong reason, wrong outcome at cold-start cap); got {r}", + ); +} + +/// Quadrant 4 mature: `pnl_sign=-1`, `dir_match=false`, `loss_cap=-2.0` +/// (WR_EMA at or above 55%) ⇒ `R_event = -2.0`. Full mature loss cap. +/// Validates that `loss_cap` is correctly threaded through, not +/// hard-coded inside the device function. +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_2() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ -0.4, /* label_at_open_sign =*/ 1, /* loss_cap =*/ -2.0, + ); + assert!( + (r - (-2.0)).abs() < 1e-6, + "expected R_event = -2.0 (wrong reason, wrong outcome at mature cap); got {r}", + ); +} + +/// Edge case: `close_pnl == 0.0` ⇒ `R_event = 0.0` regardless of +/// `label_at_open_sign` / `loss_cap` (no information case per spec +/// §4.1). Validates the early-return path. +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_zero_pnl_returns_zero() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ 0.0, /* label_at_open_sign =*/ 1, /* loss_cap =*/ -1.0, + ); + assert!( + r.abs() < 1e-6, + "expected R_event = 0.0 (close_pnl == 0 / no information); got {r}", + ); +} + +// ── Sentinel-handling sanity (Task 2.0 contract) ─────────────────────── + +/// `label_at_open_sign == 0` (sentinel: no aux signal at trade open) + +/// `pnl_sign == +1` ⇒ `dir_match = false` ⇒ `R_event = +0.5` (wrong- +/// reason right-outcome shoulder). Validates the Task 2.0 sentinel +/// contract: when the aux skip-window or no-position-yet sentinel +/// reaches the reward function, the wrong-reason path lands the +0.5 +/// shoulder on a winning trade. +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_sentinel_label_winning_trade_lands_shoulder() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ 0.7, /* label_at_open_sign =*/ 0, /* loss_cap =*/ -1.0, + ); + assert!( + (r - 0.5).abs() < 1e-6, + "expected R_event = +0.5 (sentinel label + winning trade ⇒ wrong-reason shoulder); got {r}", + ); +} + +/// `label_at_open_sign == 0` (sentinel) + `pnl_sign == -1` ⇒ +/// `dir_match = false` ⇒ `R_event = loss_cap` (wrong-reason +/// wrong-outcome quadrant). Validates the safer-default semantic for +/// no-aux-signal losing trades (full loss cap, NOT the right-reason +/// -0.5 shoulder). +#[test] +#[ignore = "requires GPU"] +fn sp20_event_reward_sentinel_label_losing_trade_lands_loss_cap() { + let stream = make_test_stream(); + let f = load_test_kernel(&stream); + let r = launch_sp20_event_reward( + &stream, &f, + /* close_pnl =*/ -0.4, /* label_at_open_sign =*/ 0, /* loss_cap =*/ -1.5, + ); + assert!( + (r - (-1.5)).abs() < 1e-6, + "expected R_event = loss_cap = -1.5 (sentinel label + losing trade ⇒ wrong-reason wrong-outcome); got {r}", + ); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 77dd85868..9819c87d1 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,128 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP20 Phase 2 Task 2.1: sp20_compute_event_reward device function + GPU oracle tests (additive) + +Lands the SP20 Component 1 reward math as a header-only `__device__ +__forceinline__` function in a new `sp20_reward.cuh` header, plus the +GPU oracle test scaffold for the 4-quadrant truth table per spec §4.1. +The function has no production callers in this commit — Task 2.2's +trade-close site replacement wires it in atomically with the SP12 v3 +reward block deletion. + +### Why a dedicated header (`sp20_reward.cuh`) instead of inlining in `experience_kernels.cu` + +The plan called for adding `sp20_compute_event_reward` "in +`experience_kernels.cu`". The user-decision review of Phase 2's six +gaps revised this to "place it near the other SP20 helpers in that +file" — i.e., the spirit was "don't put it in `trade_physics.cuh` (which +is for trade physics)", not "literally inside the .cu file". + +The test wrapper (`sp20_event_reward_test_kernel.cu`) needs to call +the SAME math the production caller does — per +`feedback_no_cpu_test_fallbacks` GPU oracle tests share the production +function bit-for-bit. Three options: + + - **Inline in `experience_kernels.cu` only** — test kernel cannot + reach the function (it's in a `.cu` file, not a header). Breaks + `feedback_no_quickfixes` (would force code duplication in the test + kernel). + - **Inline in `trade_physics.cuh`** — bloats trade_physics with SP20 + reward logic (rejected by user). + - **New `sp20_reward.cuh` header** — clean SP20 namespace, included + by both `experience_kernels.cu` (Task 2.2) and the test wrapper + (this commit). **Chosen.** + +Mirrors the `compute_asymmetric_capped_pnl` / `compute_min_hold_penalty` +pattern in `trade_physics.cuh` for shared SP12 reward helpers. + +### 4-quadrant truth table (spec §4.1) + +| `pnl_sign` | `dir_match` | `R_event` | Constant | +|------------|--------------|-------------------|--------------------------------| +| `+1` | `true` | `+1.0` | `SP20_WIN_CAP` | +| `+1` | `false` | `+0.5` | `SP20_SHOULDER_RIGHT_OUTCOME` | +| `-1` | `true` | `-0.5` | `SP20_SHOULDER_WRONG_OUTCOME` | +| `-1` | `false` | `loss_cap` | ISV[LOSS_CAP_INDEX] (caller) | +| `0` | `*` | `0.0` | early-return | + +`label_at_open_sign == 0` (sentinel: aux skip-window or +no-position-open at trade-open per Task 2.0) maps to `dir_match = +false` because the comparison `pnl_sign == 0` is never true once +`pnl_sign` is filtered to `±1`. This routes sentinel inputs to the +wrong-reason quadrant — `+0.5` shoulder on winning trades, `loss_cap` +on losing trades. Safer default than over-rewarding a no-signal close +(documented in `sp20_reward.cuh` header comment). + +### Tests + +`crates/ml/tests/sp20_event_reward_test.rs` — 8 GPU oracle tests +(all `#[ignore = "requires GPU"]`): + + 1. `quadrant_right_reason_right_outcome` — pnl=+, label=+ → +1.0 + 2. `quadrant_wrong_reason_right_outcome` — pnl=+, label=− → +0.5 + 3. `quadrant_right_reason_wrong_outcome` — pnl=−, label=− → −0.5 + 4. `quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_1` — cold-start cap + 5. `quadrant_wrong_reason_wrong_outcome_loss_cap_at_minus_2` — mature cap + 6. `zero_pnl_returns_zero` — early-return path + 7. `sentinel_label_winning_trade_lands_shoulder` — Task 2.0 contract + 8. `sentinel_label_losing_trade_lands_loss_cap` — Task 2.0 contract + +Tests 7-8 lock the Task 2.0 sentinel-label contract: when the +FoldReset-zeroed `label_at_open_per_env` slot reaches the reward +function (no `entering_trade` fired this fold OR the aux landed in +skip-window), the wrong-reason quadrant fires. This pins the Task 2.0 +"safer default" semantic against future regressions that might add a +"sentinel = pass-through" early-return. + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp20_reward.cuh` | NEW | Header with `sp20_compute_event_reward` | +| `crates/ml/src/cuda_pipeline/sp20_event_reward_test_kernel.cu` | NEW | Standalone GPU oracle test wrapper | +| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +1 include | Pre-includes for Task 2.2 caller | +| `crates/ml/build.rs` | +1 cubin entry | Compile entry for nvcc | +| `crates/ml/tests/sp20_event_reward_test.rs` | NEW | 8 GPU oracle tests | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification + +``` +SQLX_OFFLINE=true cargo check -p ml --tests --features cuda # clean +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \ + --test sp20_event_reward_test --features cuda -- --ignored --nocapture +``` + +### Phase 2 Task 2.1 → Task 2.2 forward references + +Task 2.2 atomically: + - replaces the SP12 v3 inlined reward block at + `experience_kernels.cu:3216-3380` (segment_complete branch only) + with the SP20 4-quadrant reward via `sp20_compute_event_reward`; + - adds the `is_close` consumer of `label_at_open_per_env[i]` (Task + 2.0's per-env scratch); + - adds per-env `alpha_per_env[N]` plumbing through the SP20 + aggregation kernel signature, so `alpha_ema` (ISV[ALPHA_EMA_INDEX]) + becomes non-zero post-trade-close (replacing the Phase 1.4 `0.0` + forward-reference placeholder); + - updates the 2 placeholder pin tests (`sp20_aggregate_inputs_test:: + alpha_and_per_bar_hold_reward_are_phase_2_and_3_2_placeholders` + and `sp20_phase1_4_wireup_test:: + alpha_and_hold_reward_emas_stay_at_zero_phase_2_3_2_placeholders`) + to assert the new contract: `alpha_ema` becomes non-zero after at + least one trade close; `hold_reward_ema` stays at sentinel until + Phase 3.2. + +Per-bar SP18 D-leg call sites at `experience_kernels.cu:3722` and +`:3833` are RETAINED pending Phase 3.2 Component 2 replacement (per +`feedback_no_stubs.md`: working signal stays until replacement lands). + +`compute_sp18_hold_opportunity_cost`, `compute_asymmetric_capped_pnl`, +and `compute_min_hold_penalty` device functions in `trade_physics.cuh` +are RETAINED — they remain called by per-bar SP18 sites and by the +`sp12_reward_math_test_kernel.cu` test wrapper. + ## 2026-05-09 — SP20 Phase 2 Task 2.0: per-env trade-open label-sign infrastructure (additive) Lands the producer-side infrastructure for Phase 2's 4-quadrant reward