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

fix(sp20): replay buffer records magnitude INTENT (Bellman consistency)

Browse Source 
Parallel to the 2026-05-08 Class C direction fix but on the **magnitude axis**.
Class C correctly chose REALIZED for direction because env-gated trades (capital
floor / trail / broker cap → Flat) mean no trade happens — recording as Flat is
honest. For magnitude, 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 in trade_physics.cuh:1099-1117 from |position|/max_position;
when the policy intends Full and Kelly clamps to 0.35× max (bucket → Quarter),
recording Quarter loses the policy's intent. Q(s, mag=Full) never receives the
reward gradient from Full-intent trades that actually executed.

One-line swap inside the Class C encoding block at experience_kernels.cu:2689-
2697: `actual_mag_core * b2_size * b3_size` → `mag_idx * b2_size * b3_size`.
`mag_idx` is the local intent magnitude already decoded at line ~2460 from the
original action_idx BEFORE any Kelly/CVaR/Q-gap clamping. Direction axis keeps
Class C semantics (`actual_dir_core` → gated trades record as Flat). Order/
urgency remain intent passthrough.

`actual_mag_core` remains consumed downstream by Task 2.X per-magnitude trade-
lifecycle instrumentation at the seg_mag_bin site below — direction-branch Q
learns from realized; magnitude-branch Q learns from intent. The two axes have
different env-enforcement semantics so they take different actions at the
replay-write site.

Verification: new CPU oracle test crates/ml/tests/sp20_magnitude_intent_record_
test.rs pins the encoding contract with 4 invariants (Full→Quarter scenario,
all-pairs sweep, order/urgency passthrough, direction Class C unchanged). Per
pearl_tests_must_prove_not_lock_observations the test asserts invariants
("recorded mag == intent mag, regardless of realized") not observed values, so
it cannot become a bug-lock if the kernel's compute path changes — only if the
encoding contract itself is broken. All 4 tests pass; cargo check clean.

Audit-doc entry added to docs/dqn-wire-up-audit.md as the SP20 magnitude intent
fix, parallel to the Class C direction fix above with the asymmetry between
the two axes spelled out.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 20:37:05 +02:00
parent 4e21c38b85
commit e96f7fcdd1
3 changed files with 440 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

298
crates/ml/tests/sp20_magnitude_intent_record_test.rs Normal file
View File

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