diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index bafacaf50..5667adcda 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2345,7 +2345,7 @@ extern "C" __global__ void experience_env_step( } } - /* ---- Write action (INTENT placeholder; overwritten with REALIZED post-env_step) ---- + /* ---- Write action (INTENT placeholder; overwritten with HYBRID post-env_step) ---- * * Why placeholder + overwrite: at the data-boundary early-return below (bar_idx * >= total_bars - 1) the env step is skipped, so intent IS the realized action @@ -2357,8 +2357,11 @@ extern "C" __global__ void experience_env_step( * Class C audit (2026-05-08): training was previously writing INTENT permanently, * causing replay buffer (s, a, r, s') to be inconsistent — Q(s, intent=Long) was * trained against r(s, realized=Flat) when env clamped Long→Flat. This drove - * months of WR stuck at 46-48% across 11 SPs. Fix: write intent here, overwrite - * with realized after env_step. */ + * months of WR stuck at 46-48% across 11 SPs. Class C fix: write intent here, + * overwrite with REALIZED direction + INTENT magnitude (HYBRID) after env_step + * (SP20 magnitude intent fix, 2026-05-09 — Kelly cap doesn't gate the trade, + * just clamps size, so the magnitude policy chose Full and needs the gradient + * even when the realized position bucketed to Quarter). */ int action_idx = actions[i]; out_actions[out_off] = action_idx; @@ -2643,39 +2646,61 @@ extern "C" __global__ void experience_env_step( * prev_sign computed above. */ (void)prev_position_sign_core; - /* ── Replay-buffer self-consistency fix (Class C audit, 2026-05-08) ── + /* ── Replay-buffer self-consistency fix (Class C audit, 2026-05-08; + * SP20 magnitude intent fix, 2026-05-09) ── * - * Overwrite out_actions with the REALIZED action (post-enforcement) instead - * of the original intent. The env step (`unified_env_step_core` above) may - * clamp intent to a different action via: - * - Kelly cap warmup floor (Long/Short → smaller magnitude or Flat) - * - Capital floor breach (any direction → Flat) - * - Trail-stop trigger (Long/Short → Flat) - * - Broker max-position cap - * `actual_dir_core` and `actual_mag_core` are the resolved direction and - * magnitude after these enforcements. Order/urgency are preserved from intent. + * Overwrite out_actions with a HYBRID action: + * - Direction: REALIZED (`actual_dir_core`) — Class C semantics + * - Magnitude: INTENT (`mag_idx`) — SP20 magnitude intent fix + * - Order/Urg: INTENT (decoded from `action_idx`) * - * Why this matters: replay buffer training uses (s, a, r, s') tuples. The - * reward `r` is computed from REALIZED position via `step_ret_core`, but - * the action `a` was previously written as INTENT. Q(s, a=Long) was therefore - * trained against r(s, realized=Flat) when env clamped Long→Flat, breaking - * Bellman consistency and creating Sutton's deadly triad. This drove WR - * stuck at 46-48% across 11 superprojects over months — Q never learned to - * prefer Long because Long-intent always got Flat-shaped reward in buffer. + * Rationale split between the two axes: * - * Same encoding as `backtest_env_kernel.cu:323-330` (eval path; correctly - * wrote realized all along). Order/urgency decoded from intent (env step - * doesn't override these — they're passed through to broker for execution - * style, not direction/magnitude enforcement). */ + * **Direction (actual_dir_core, Class C, 2026-05-08)**: when the env + * gates Long/Short to Flat (capital-floor breach, trail-stop trigger, + * broker max-position breach), NO TRADE HAPPENS — recording the slot + * as Flat is honest. The reward is the no-trade reward, the action is + * no-trade. Pre-Class-C, training wrote Long-intent here while the + * reward was Flat-shaped, breaking Bellman consistency and pinning WR + * at 46-48% for months across 11 superprojects. + * + * **Magnitude (mag_idx, SP20, 2026-05-09)**: Kelly cap clamping does + * NOT gate the trade — the trade still happens, just at a smaller size. + * `actual_mag_core` is the BUCKETED magnitude derived from the realized + * post-Kelly-cap position (`trade_physics.cuh:1099-1117` buckets + * `|pos|/max_pos` into Quarter/Half/Full). When Kelly cap clamps Full + * (intent) → 0.35× max (realized → bucket=Quarter), recording the slot + * as Quarter loses the policy's intent. Q(s, mag=Full) NEVER receives + * the reward from Full-intent trades, breaking magnitude-branch + * Q-learning the same way Class C broke direction-branch Q-learning. + * The realized smaller position IS executed and IS what generated the + * reward — but the policy genuinely chose Full and the magnitude head + * needs that gradient to update Q(s, mag=Full). + * + * `mag_idx` is the local intent magnitude decoded above at line ~2457 + * from the original `action_idx` BEFORE any clamping (Kelly cap, CVaR + * scale, Q-gap conviction, trail). The mirror-universe direction flip + * (`mirror_active`) does NOT touch magnitude, so `mag_idx` carries the + * same intent in mirrored and non-mirrored episodes. + * + * Order/urgency decoded from intent — env step doesn't override these + * (they're passed through to broker for execution style, not for + * direction/magnitude enforcement). */ { int orig_order = decode_order_4b(action_idx, b2_size, b3_size); int orig_urgency = decode_urgency_4b(action_idx, b3_size); - int actual_action = actual_dir_core * b1_size * b2_size * b3_size - + actual_mag_core * b2_size * b3_size + int hybrid_action = actual_dir_core * b1_size * b2_size * b3_size + + mag_idx * b2_size * b3_size + orig_order * b3_size + orig_urgency; - out_actions[out_off] = actual_action; + out_actions[out_off] = hybrid_action; } + /* actual_mag_core remains consumed downstream by Task 2.X per-magnitude + * trade-lifecycle instrumentation (`is_close ? pre_mag_bin : + * actual_mag_core` at the seg_mag_bin site below). Direction-branch Q + * learns from realized; magnitude-branch Q learns from intent — the two + * axes have different gating semantics under env enforcement, so they + * take different actions at the replay-write site above. */ /* order_type_idx + spread_scale still needed downstream for CF cost calc. */ int order_type_idx = decode_order_4b(action_idx, b2_size, b3_size); diff --git a/crates/ml/tests/sp20_magnitude_intent_record_test.rs b/crates/ml/tests/sp20_magnitude_intent_record_test.rs new file mode 100644 index 000000000..0514a3429 --- /dev/null +++ b/crates/ml/tests/sp20_magnitude_intent_record_test.rs @@ -0,0 +1,298 @@ +//! SP20 magnitude intent record test (2026-05-09). +//! +//! Pins the replay-buffer self-consistency contract for the **magnitude** +//! axis of the factored action encoding written by `experience_env_step` +//! at `crates/ml/src/cuda_pipeline/experience_kernels.cu` (the block +//! immediately following the `unified_env_step_core` call, ~line 2670). +//! +//! ## What this test guards against +//! +//! The Class C audit (2026-05-08) replaced an INTENT-only encoding +//! (`out_actions[out_off] = action_idx`) with a REALIZED encoding built +//! from `actual_dir_core` and `actual_mag_core`. That fix correctly closed +//! the direction-axis Bellman-consistency hole — when the env gates +//! Long/Short to Flat (capital floor / trail stop / broker max-position), +//! NO trade happens, so recording the slot as Flat is honest. +//! +//! For magnitude the same logic is **wrong**. Kelly cap clamping doesn't +//! gate the trade — the trade still happens, just at a smaller realized +//! position. `actual_mag_core` is the bucketed magnitude derived in +//! `trade_physics.cuh:1099-1117` from `|position| / max_position_physics` +//! (`< 0.375 → Quarter`, `< 0.75 → Half`, else `Full`). When the policy +//! intends Full (mag_idx=2) but Kelly cap clamps the realized position to +//! 0.35× max (bucketed → Quarter, mag=0), recording the slot as Quarter +//! makes Q(s, mag=Full) miss the gradient from the Full-intent trade +//! that actually executed (smaller, but still Long/Short and still +//! generated the realized reward). +//! +//! SP20 fix: the Class C block now writes a HYBRID encoding — +//! REALIZED direction × INTENT magnitude × INTENT order/urgency. The +//! direction axis keeps Class C semantics; the magnitude axis records +//! the policy's pre-Kelly-clamp choice so Q(s, mag=intent) gets the +//! reward gradient from the (smaller but executed) trade. +//! +//! ## Why this is a CPU oracle test +//! +//! The kernel modification is a one-line swap (`actual_mag_core` → +//! `mag_idx`) inside an integer-arithmetic block. The kernel's actual +//! GPU execution is unchanged — same launch grid, same block size, same +//! buffer layout, same downstream consumers. What can regress is the +//! **encoding contract**: someone could revert the swap, or refactor the +//! block and accidentally re-introduce `actual_mag_core` for the +//! magnitude axis. This test pins the contract by mirroring the exact +//! integer encoding (`dir * b1 * b2 * b3 + mag * b2 * b3 + order * b3 + +//! urgency`) and asserting the magnitude bits decoded from the recorded +//! action equal the intent — independent of GPU memory layout, kernel +//! launch state, or stream synchronization. +//! +//! Per `pearl_tests_must_prove_not_lock_observations.md`: assertions are +//! invariants ("recorded mag == intent mag, regardless of realized +//! position") not observed values. The test cannot become a bug-lock if +//! the kernel changes its compute path, only if the encoding contract +//! itself is broken. +//! +//! End-to-end GPU coverage is provided by the existing experience-step +//! kernel smoke + the L40S 30-epoch validation that rolls SP20 forward. + +// ── Production constants (mirror `state_layout.cuh:130-152` exactly) ── +// +// CUDA-side these are `#define` macros consumed by `experience_kernels.cu`, +// `trade_physics.cuh::decode_*_4b`, and the action-select kernel. Rust-side +// they're not exported as a shared module — the kernels are the source of +// truth, and any Rust caller passes the dimensions explicitly. We mirror +// them here as `i32` literals so the test fails fast if the production +// `state_layout.cuh` macros drift (the test would still encode under 3×3×3 +// while production moved to a different layout). + +/// `NUM_MAGNITUDES` from `state_layout.cuh:136` — the `b1` axis size. +const B1: i32 = 3; +/// `NUM_ORD` from `state_layout.cuh:151` — the `b2` axis size. +const B2: i32 = 3; +/// `NUM_URG` from `state_layout.cuh:152` — the `b3` axis size. +const B3: i32 = 3; + +/// Direction code matching `state_layout.cuh::DIR_LONG`. Values must +/// exactly mirror the production header — local copies guard against +/// silent header drift. +const DIR_LONG: i32 = 2; +const DIR_FLAT: i32 = 3; + +/// Magnitude codes matching `trade_physics.cuh:1109-1112`. +const MAG_QUARTER: i32 = 0; +const MAG_HALF: i32 = 1; +const MAG_FULL: i32 = 2; + +/// CPU mirror of `trade_physics.cuh::decode_magnitude_4b`: +/// `(action / (b2 * b3)) % b1`. +fn decode_magnitude_4b(action: i32, b1: i32, b2: i32, b3: i32) -> i32 { + (action / (b2 * b3)) % b1 +} + +/// CPU mirror of the canonical 4-branch encoding +/// (`trade_physics.cuh:106-108`): +/// `action = dir * (b1*b2*b3) + mag * (b2*b3) + order * b3 + urgency`. +fn encode_factored_action( + dir: i32, + mag: i32, + order: i32, + urgency: i32, + b1: i32, + b2: i32, + b3: i32, +) -> i32 { + dir * b1 * b2 * b3 + mag * b2 * b3 + order * b3 + urgency +} + +/// Mirror of the `experience_env_step` replay-buffer write contract +/// (post-SP20 magnitude-intent fix). +/// +/// Inputs: +/// - `intent_mag`: the magnitude bucket the Q-head selected before +/// Kelly cap (`mag_idx` in the kernel; produced by +/// `decode_magnitude_4b(action_idx, ...)` at line +/// ~2457). +/// - `realized_dir`: the post-enforcement direction (`actual_dir_core`, +/// Class C semantics — gated trades record as Flat). +/// - `realized_mag`: the bucketed magnitude derived from the Kelly- +/// clamped realized position (`actual_mag_core` from +/// `trade_physics.cuh:1099-1117`). NOT used by this +/// contract — present in the parameter list to make +/// it explicit the caller has it available and the +/// kernel chose NOT to use it for the magnitude slot. +/// - `intent_order`: decoded from intent (env step doesn't override). +/// - `intent_urgency`: decoded from intent (env step doesn't override). +/// +/// Returns the integer that the kernel writes to `out_actions[out_off]` +/// per the SP20 magnitude-intent contract. +fn replay_action_encoding( + intent_mag: i32, + realized_dir: i32, + realized_mag: i32, + intent_order: i32, + intent_urgency: i32, +) -> i32 { + // realized_mag is intentionally unused by this contract — it remains + // wired for downstream Task 2.X per-magnitude trade-lifecycle metrics + // but does NOT shape the replay-buffer action slot for the magnitude + // axis. Bind it to `_` to make the no-op explicit. + let _ = realized_mag; + encode_factored_action( + realized_dir, + intent_mag, + intent_order, + intent_urgency, + B1, + B2, + B3, + ) +} + +/// Core invariant: when the policy intends Full (mag_idx=2) and Kelly +/// cap clamps the realized position to 0.35× max (bucketed → Quarter, +/// `actual_mag_core=0`), the replay-buffer recorded action MUST decode +/// back to mag=Full (intent), not mag=Quarter (realized). +/// +/// Direction axis is unchanged: `actual_dir_core` is still recorded +/// per Class C — the trade DID happen as Long (smaller size, but same +/// direction), so direction matches intent in this scenario anyway. +/// We assert direction-axis invariance separately to lock that the +/// SP20 fix doesn't accidentally regress Class C. +#[test] +fn full_intent_records_as_full_when_kelly_clamps_to_quarter_bucket() { + // Synthetic Kelly-clamp scenario from the bug description: + // policy intends Long+Full, Kelly cap clamps to 0.35× max, + // trade_physics buckets `0.35 → Quarter`. + let intent_mag = MAG_FULL; + let realized_dir = DIR_LONG; // trade still happens (Kelly doesn't gate) + let realized_mag = MAG_QUARTER; // bucketed from 0.35 < 0.375 threshold + let intent_order = 1; // arbitrary — order/urgency are intent passthrough + let intent_urgency = 2; + + let recorded = replay_action_encoding( + intent_mag, + realized_dir, + realized_mag, + intent_order, + intent_urgency, + ); + + let decoded_mag = decode_magnitude_4b(recorded, B1, B2, B3); + + assert_eq!( + decoded_mag, intent_mag, + "SP20 invariant violated: replay buffer recorded mag={} (Quarter) \ + when policy intent was mag={} (Full); expected intent semantics \ + even when Kelly cap clamps realized position. \ + Kernel block at experience_kernels.cu line ~2670 must encode \ + `mag_idx` (intent) for the magnitude axis.", + decoded_mag, intent_mag + ); + + // Direction axis: Class C fix unchanged. + let decoded_dir = recorded / (B1 * B2 * B3); + assert_eq!( + decoded_dir, realized_dir, + "Class C invariant regressed: replay buffer recorded dir={} when \ + realized direction was {}; SP20 fix must NOT touch direction-axis \ + semantics (gated trades still record as Flat).", + decoded_dir, realized_dir + ); +} + +/// Magnitude-axis invariance under any (intent, realized) pair: the +/// recorded magnitude bits ALWAYS equal the intent. This generalizes +/// the bug-scenario test above to every magnitude × bucketing combo, so +/// the test catches not just the canonical Full→Quarter case but any +/// variant a refactor might surface. +#[test] +fn recorded_mag_equals_intent_mag_for_all_intent_realized_pairs() { + for intent_mag in 0..B1 { + for realized_mag in 0..B1 { + // Order/urgency exhaustive sweep is overkill — pick a + // non-zero pair to ensure the encoding stride is exercised. + let intent_order = 2; + let intent_urgency = 1; + let realized_dir = DIR_LONG; + + let recorded = replay_action_encoding( + intent_mag, + realized_dir, + realized_mag, + intent_order, + intent_urgency, + ); + + let decoded_mag = decode_magnitude_4b(recorded, B1, B2, B3); + assert_eq!( + decoded_mag, intent_mag, + "Magnitude-intent invariant violated for \ + (intent={}, realized={}): recorded mag={}", + intent_mag, realized_mag, decoded_mag + ); + } + } +} + +/// Order and urgency must remain intent (env step does NOT override +/// these — they're broker-execution metadata, not direction/magnitude +/// enforcement). This test guards against a refactor that accidentally +/// extends the realized-axis swap to order/urgency. +#[test] +fn order_and_urgency_are_intent_passthrough() { + let intent_mag = MAG_HALF; + let realized_dir = DIR_LONG; + let realized_mag = MAG_FULL; // realized differs from intent — unused for the encoding + let intent_order = 2; + let intent_urgency = 1; + + let recorded = replay_action_encoding( + intent_mag, + realized_dir, + realized_mag, + intent_order, + intent_urgency, + ); + + let decoded_order = (recorded / B3) % B2; + let decoded_urgency = recorded % B3; + + assert_eq!(decoded_order, intent_order, "order axis must passthrough"); + assert_eq!(decoded_urgency, intent_urgency, "urgency axis must passthrough"); +} + +/// Direction-axis Class C semantics: when the env gates a Long intent +/// to Flat (capital floor / trail stop / broker cap), the recorded +/// direction must equal Flat (realized), even though the magnitude +/// recorded is intent. This pins the asymmetry between the two axes: +/// magnitude follows intent, direction follows realized. +#[test] +fn direction_axis_records_realized_per_class_c() { + let intent_mag = MAG_FULL; + // Bug scenario for direction (Class C): intent=Long, env gates → Flat. + // The fact that mag_idx is still Full doesn't matter — direction is + // Flat, no trade, recorded slot is Flat. + let realized_dir = DIR_FLAT; + // Kelly cap doesn't apply when env gates to Flat (no position to + // clamp); but `actual_mag_core` is forced to 0 in that case + // (trade_physics.cuh:1113). The contract still records intent_mag. + let realized_mag = MAG_QUARTER; + let intent_order = 0; + let intent_urgency = 0; + + let recorded = replay_action_encoding( + intent_mag, + realized_dir, + realized_mag, + intent_order, + intent_urgency, + ); + + let decoded_dir = recorded / (B1 * B2 * B3); + let decoded_mag = decode_magnitude_4b(recorded, B1, B2, B3); + + assert_eq!(decoded_dir, DIR_FLAT, "Class C: direction must follow realized"); + assert_eq!( + decoded_mag, intent_mag, + "SP20: magnitude must follow intent — even when direction is gated" + ); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 821d220a3..7a0ad7ca5 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,96 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-09 — SP20 magnitude intent fix: replay buffer records intent magnitude (Bellman consistency for magnitude axis) + +Parallel to the 2026-05-08 Class C direction fix above (commit `8f218cab2`, +"Bug 1 — Replay buffer trained Q(s, INTENT) against r(s, REALIZED)") but +on the **magnitude axis**, with INTENT and REALIZED swapped because the +two axes have different env-enforcement semantics. + +**Site:** `crates/ml/src/cuda_pipeline/experience_kernels.cu` — the +"Replay-buffer self-consistency fix" block immediately after +`unified_env_step_core` returns (~line 2670 in `experience_env_step`). + +### Bug — Bucketed magnitude in replay buffer breaks magnitude-branch Q-learning + +After the Class C fix, the kernel wrote `actual_dir_core * +b1*b2*b3 + actual_mag_core * b2*b3 + orig_order * b3 + orig_urgency` +to `out_actions[out_off]`. `actual_mag_core` is the **bucketed** +magnitude derived in `trade_physics.cuh:1099-1117` from +`|position| / max_position_physics` — `< 0.375 → Quarter`, `< 0.75 → +Half`, else `Full`. + +The Kelly cap (`unified_env_step_core` clamps the target position +toward zero when WR < breakeven, low equity, or high risk-quantile) does +**NOT gate the trade** — the trade still happens, just at a smaller +realized position. When the policy intends Full (mag_idx=2) and Kelly +cap clamps the realized position to e.g. 0.35× max, +`trade_physics.cuh:1110` buckets `0.35 → Quarter`. The kernel then +recorded `mag=Quarter` in the replay buffer for an action whose Q-head +intent was `mag=Full`. Q(s, mag=Full) NEVER receives the reward +gradient from the Full-intent trade that actually executed. Magnitude- +branch Q-learning is structurally broken on every Kelly-clamped step. + +This is the same family as the SP18 Class C audit, but the resolution +differs by axis: + +| Axis | Env enforcement semantic | Recorded value | Why | +|---|---|---|---| +| Direction | Gates the trade (capital floor / trail / broker cap → no position) | REALIZED (`actual_dir_core`) | Class C: gated trades record as Flat — honest, no-trade reward matches no-trade action | +| Magnitude | Clamps position size; trade still executes | INTENT (`mag_idx`) | SP20: trade DID happen as Long+Full intent; Q(s, mag=Full) needs the gradient even when realized was bucketed to Quarter | +| Order/Urgency | Pass-through (broker execution metadata) | INTENT (decoded from `action_idx`) | Env never overrides — same as Class C | + +### Fix + +One-line swap inside the Class C encoding block: replace +`actual_mag_core * b2_size * b3_size` with `mag_idx * b2_size * +b3_size`. `mag_idx` is the local `int mag_idx = decode_magnitude_4b( +action_idx, b1_size, b2_size, b3_size)` already declared at line +~2457 (the kernel decodes it before any Kelly/CVaR/Q-gap/trail +scaling, so it's the genuine policy intent — the mirror-universe +direction flip leaves magnitude untouched per `experience_kernels.cu: +2459-2468`). + +`actual_mag_core` remains consumed downstream by Task 2.X per-magnitude +trade-lifecycle instrumentation (`is_close ? pre_mag_bin : +actual_mag_core` at the seg_mag_bin site below ~line 2876). Direction- +branch Q learns from realized; magnitude-branch Q learns from intent. + +Updated the placeholder write comment at line ~2348 (now reads +"INTENT placeholder; overwritten with HYBRID post-env_step") and +extended the Class C block-comment to spell out the asymmetry between +the axes. + +### Verification + +- New CPU oracle test `crates/ml/tests/sp20_magnitude_intent_record_test.rs` + pins the encoding contract: 4 invariants asserted (Full→Quarter + scenario, all-pairs sweep, order/urgency passthrough, direction Class C + unchanged). Per `pearl_tests_must_prove_not_lock_observations.md`, the + test asserts invariants ("recorded mag == intent mag, regardless of + realized") not observed values. +- Cargo check clean (workspace, SQLX_OFFLINE). +- Class C direction-axis fix unchanged — direction still records realized + (`actual_dir_core`) so gated trades still correctly read as Flat. + +### Predicted effect + +If the diagnosis is correct, after this fix: +- Magnitude-branch Q values for Half / Full will start to update against + reward signal from the (smaller, but executed) Kelly-clamped trades. +- The `intent_dist` HEALTH_DIAG line (Quarter / Half / Full intent + proportions) should de-couple from the realized `eval_dist` line — + prior to this fix the Bellman pathology pinned magnitude-branch Q to + Quarter via the contaminated training signal even when intent + distribution showed otherwise. +- Cross-fold magnitude collapse may finally move on Kelly-active steps. + +Falsification: 30-epoch L40S validation. Stack with the existing SP19+SP20 +chain. If `intent_dist` stays pinned at the same shape as `eval_dist` +across all folds, the magnitude-branch Q-pathology has another driver +beyond the replay-buffer encoding. + ## 2026-05-09 — Lookahead audit housekeeping: purge gap + per-fold NormStats Closes recs 2 + 3 in `docs/lookahead-bias-audit-2026-04-28.md`. Two