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

feat(sp20): Phase 2 Task 2.0 — per-env trade-open label-sign infra

Browse Source 
Lands the producer-side infrastructure for SP20 Phase 2's 4-quadrant
reward kernel (Task 2.2):

* New `[alloc_episodes]` `i32` device buffer `label_at_open_per_env`
  on `GpuExperienceCollector`, allocated alongside `episode_starts_buf`.
* Two new `experience_env_step` kernel args (`aux_label_per_env` input,
  `label_at_open_per_env` output), threaded into the launch site.
* Write site at the `entering_trade` branch maps the upstream aux
  next-bar label (`{0,1,-1}` from `aux_sign_label_per_step_kernel`) to
  the SP20 spec sign convention (`{-1,+1,0}`) per spec §4.1.
* StateResetRegistry entry `label_at_open_per_env` (FoldReset) +
  `reset_named_state` dispatch arm via `stream.memset_zeros`.
* Registry invariant test `sp20_label_at_open_per_env_registered_fold_reset`
  pins the FoldReset category + key description anchors.
* Audit-doc entry documents the design rationale (why per-env buffer
  vs. derived-at-trade-close read), sign mapping, FoldReset semantics,
  kernel-arg threading, and Task 2.1/2.2 forward references.

Producer-only this commit. The consumer (`is_close` branch reading
`label_at_open_per_env[i]` and passing to `sp20_compute_event_reward`)
lands atomically with the SP12 v3 reward block replacement in Task 2.2,
per `feedback_no_partial_refactor`.

NULL-tolerant kernel arg pair lets test scaffolds without aux-head
wiring continue to work; the FoldReset-zeroed buffer reads as sentinel
0 at the consumer (which Task 2.2 maps to "no info" / wrong-reason
quadrant — safer default than over-rewarding a no-signal trade).

Verification:
  SQLX_OFFLINE=true cargo check -p ml          # clean
  SQLX_OFFLINE=true cargo test -p ml --lib sp20 # 18/18 (+1 new)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 23:54:01 +02:00
parent abd7e533bc
commit 29a2615f6a
5 changed files with 268 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -2156,7 +2156,27 @@ extern "C" __global__ void experience_env_step(
* between on-policy and CF: if on-policy at (i,t) was an
* early-return, the CF slot at (i,t) is also skipped. NULL-tolerant
* for any test scaffold that hasn't wired the buffer yet. */
int* __restrict__ slot_completed_normally_out
int* __restrict__ slot_completed_normally_out,
/* SP20 Phase 2 Task 2.0 (2026-05-09) — aux next-bar label tile
* (input) + per-env trade-open label-sign scratch (output). Both
* `[N]`-sized i32 device buffers. NULL-tolerant for test scaffolds
* without aux-head wiring (write site is gated on both being
* non-NULL).
*
* `aux_label_per_env` ← `exp_aux_nb_label_buf`: per-env aux next-
* bar label written by `aux_sign_label_per_step_kernel` BEFORE this
* kernel runs (same stream, stream-implicit ordering). Values:
* `1` = "up", `0` = "not-up", `-1` = skip-window (lookahead past
* total_bars).
*
* `label_at_open_per_env`: this kernel's `entering_trade` branch
* maps the aux label to the SP20 sign convention (`{-1, 0, +1}`)
* and writes it here. Slot persists across rollout steps until the
* next `entering_trade` overwrites it; consumed at `is_close` by
* the SP20 reward kernel (Task 2.2 wires the read site).
* FoldReset-zeroed by `state_reset_registry.rs::label_at_open_per_env`. */
const int* __restrict__ aux_label_per_env,
int* __restrict__ label_at_open_per_env
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -2887,6 +2907,39 @@ extern "C" __global__ void experience_env_step(
if (entering_trade) {
entry_price = raw_close;
trade_start_pnl = (ps[PS_REALIZED_PNL]);
/* SP20 Phase 2 Task 2.0 (2026-05-09): capture aux next-bar label
* sign at trade open into per-env scratch. Consumed at trade
* close (`is_close == 1`) by `sp20_compute_event_reward`'s
* `dir_match` quadrant check (spec §4.1).
*
* Mapping from `aux_label_per_env[i]` (kernel emits `0`/`1`/`-1`
* per `aux_sign_label_per_step_kernel.cu:77` — `0`=not-up,
* `1`=up, `-1`=skip-window) to the SP20 sign convention
* (`{-1, 0, +1}`):
* `1` → `+1` (aux predicted up at trade open)
* `0` → `-1` (aux predicted not-up — down or flat)
* `-1` → `0` (aux skip-window: lookahead past total_bars,
* no signal — sentinel 0 is the "no info"
* case in `sp20_compute_event_reward`).
*
* NULL-tolerant — the kernel arg pair is optional so test
* scaffolds without aux-head wiring don't break. When NULL the
* write is a no-op; the FoldReset-zeroed buffer reads as
* sentinel 0 at the trade-close consumer (Task 2.2 wires the
* read site).
*
* Race-free by construction: each thread writes only its own
* `[i]` slot — no inter-thread contention. Pure GPU buffer
* (kernel-to-kernel only); no host visibility needed. */
if (aux_label_per_env != NULL && label_at_open_per_env != NULL) {
const int aux_l = aux_label_per_env[i];
int label_at_open_sign;
if (aux_l == 1) label_at_open_sign = +1;
else if (aux_l == 0) label_at_open_sign = -1;
else label_at_open_sign = 0; /* aux_l == -1 (skip) */
label_at_open_per_env[i] = label_at_open_sign;
}
}
/* On reversal: close old segment, open new one.

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

@@ -641,6 +641,43 @@ pub struct GpuExperienceCollector {
portfolio_states: CudaSlice<f32>, // [alloc_episodes * PORTFOLIO_STRIDE] (43 floats per episode)
episode_starts_buf: CudaSlice<i32>,// [alloc_episodes]
/// SP20 Phase 2 Task 2.0 (2026-05-09): per-env label sign captured at trade
/// open, consumed at trade close by `sp20_compute_event_reward`'s
/// `dir_match` check (spec §4.1).
///
/// Layout: `[alloc_episodes]` i32 — one slot per env, persists across
/// the rollout and replay-window inside `experience_env_step`. Values:
/// `+1` — aux next-bar predicted "up" at trade open
/// `-1` — aux next-bar predicted "not-up" (down/flat) at trade open
/// `0` — sentinel: trade-open was during aux skip-window (look-ahead
/// past `total_bars`) or no position currently open. The 4-quadrant
/// reward maps `pnl_sign != label_at_open_sign` to "wrong reason"
/// and `pnl_sign == 0` (close_pnl == 0) to no-info; sentinel 0
/// input flows through naturally because no `pnl_sign != 0` case
/// can match `label_at_open_sign == 0` (mismatch ⇒ "wrong reason").
/// This lands the wrong-reason quadrant for sentinel inputs, which
/// is the safer default than over-rewarding a trade with no
/// reasoning signal.
///
/// Producer: `experience_env_step` writes to this slot in the
/// `entering_trade` branch by mapping
/// `exp_aux_nb_label_buf[env]` (which is the per-step aux next-bar
/// label written by `aux_sign_label_per_step_kernel` — values
/// `{-1, 0, +1}`) to `{0, -1, +1}` respectively (the kernel's `0` =
/// "not-up" maps to `-1` for the spec sign convention; the kernel's
/// `-1` = "skip" maps to sentinel `0`).
///
/// Consumer: `experience_env_step` reads the slot in the `is_close`
/// branch and passes it to `sp20_compute_event_reward(close_pnl,
/// label_at_open_sign, isv[LOSS_CAP_INDEX])`.
///
/// FoldReset semantics (registered in `state_reset_registry.rs` as
/// `label_at_open_per_env`): zero on every fold boundary so the new
/// fold's first `entering_trade` branch fully populates the slot from
/// the new fold's aux head outputs (no stale label leakage across
/// folds — each fold has independent aux next-bar predictions).
pub(crate) label_at_open_per_env: CudaSlice<i32>, // [alloc_episodes]
// Output buffers [alloc_episodes * alloc_timesteps, ...]
states_out: CudaSlice<f32>, // #30 [alloc_episodes * alloc_timesteps * STATE_DIM] f32
actions_out: CudaSlice<i32>, // [alloc_episodes * alloc_timesteps]
@@ -1689,6 +1726,17 @@ impl GpuExperienceCollector {
.alloc_zeros::<i32>(alloc_episodes)
.map_err(|e| MLError::ModelError(format!("alloc episode_starts_buf: {e}")))?;
// SP20 Phase 2 Task 2.0 (2026-05-09): per-env label sign captured
// at trade open, consumed at trade close. Sentinel = 0 ("no
// position currently open or aux was in skip-window at entry");
// first entering_trade observation overwrites with mapped aux
// label sign in {-1, 0, +1}. Sized [alloc_episodes].
let label_at_open_per_env = stream
.alloc_zeros::<i32>(alloc_episodes)
.map_err(|e| MLError::ModelError(format!(
"sp20: alloc label_at_open_per_env ({alloc_episodes} i32): {e}"
)))?;
// Persistent epoch state
let epoch_state_init: Vec<f32> = vec![
0.01, // vol_ema
@@ -2716,6 +2764,7 @@ impl GpuExperienceCollector {
curiosity_weights,
portfolio_states,
episode_starts_buf,
label_at_open_per_env,
states_out,
actions_out,
rewards_out,
@@ -6100,6 +6149,16 @@ impl GpuExperienceCollector {
// `rewards_out` while skipping early-return slots.
.arg(&mut self.r_discipline_per_sample)
.arg(&mut self.slot_completed_normally_per_sample)
// SP20 Phase 2 Task 2.0 (2026-05-09): aux next-bar
// label tile (input — already populated by the
// `aux_sign_label_per_step_kernel` launch in step
// 3a-β above on the same stream; stream-implicit
// producer→consumer ordering) + per-env trade-open
// label-sign scratch (output — written at the
// `entering_trade` branch, consumed at `is_close`
// by Task 2.2's `sp20_compute_event_reward` call).
.arg(&self.exp_aux_nb_label_buf)
.arg(&mut self.label_at_open_per_env)
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_env_step t={t}: {e}"

46
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -2114,6 +2114,14 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "SP20 Phase 1.4 mapped-pinned [8] i32 per-EMA observation counter for `sp20_emas_compute_kernel`. counter==0 ⇒ next firing observation REPLACES sentinel directly (Pearl-A bootstrap); counter>0 ⇒ Wiener-blend with α=0.4 floor. FoldReset zeros all 8 counters so each fold's first per-EMA firing observation re-bootstraps from sentinel — necessary because some EMAs legitimately observe 0.0 and the count-based bootstrap is the correct sentinel discriminator per `pearl_first_observation_bootstrap`.",
},
// ── SP20 Phase 2 Task 2.0 (2026-05-09): per-env trade-open
// label sign scratch — written at `entering_trade`, read at
// `is_close` by the SP20 4-quadrant reward kernel.
RegistryEntry {
name: "label_at_open_per_env",
category: ResetCategory::FoldReset,
description: "GpuExperienceCollector.label_at_open_per_env [alloc_episodes] i32 device-resident scratch — SP20 Phase 2 Task 2.0 (2026-05-09) per-env aux next-bar label sign captured at trade open. Layout: tile[env] in {-1, 0, +1} where +1 = aux predicted up at trade open, -1 = aux predicted not-up, 0 = sentinel (no position currently open or trade-open landed in aux skip-window). Producer: `experience_env_step` `entering_trade` branch — reads `exp_aux_nb_label_buf[env]` (kernel emits `0`/`1`/`-1` per `aux_sign_label_per_step_kernel.cu:77`) and maps `0`→`-1`, `1`→`+1`, `-1`→`0` to match the SP20 spec sign convention (spec §4.1: `label_at_open = sign(SP19_blended_label[trade_open_bar])`). Consumer: `experience_env_step` `is_close` branch — passes `label_at_open_per_env[env]` to `sp20_compute_event_reward(close_pnl, label_at_open_sign, isv[LOSS_CAP_INDEX])`. FoldReset sentinel 0 across all `alloc_episodes` slots — leftover label from the previous fold's open trade would corrupt the new fold's first trade-close `dir_match` check (the new fold's aux head has independent calibration and would not have produced that label). Reset path: `cudarc::driver::CudaStream::memset_zeros` on the `CudaSlice<i32>` (device-resident, no host buffer to fill). Per `pearl_first_observation_bootstrap` the sentinel-0 input also flows through `sp20_compute_event_reward` to the wrong-reason quadrant on the corner case where a trade closes WITHOUT a prior `entering_trade` having fired this fold (trade carry-over from before the reset is impossible because portfolio_states is also FoldReset).",
},
];
Self { entries }
}
@@ -2367,4 +2375,42 @@ mod sp20_registry_tests {
);
}
}
/// SP20 Phase 2 Task 2.0: `label_at_open_per_env` MUST be registered as
/// FoldReset so the new fold's first `entering_trade` branch fully
/// re-populates from the new fold's aux head (preventing stale label
/// leakage across folds — each fold has independent aux head
/// calibration). Description MUST mention SP20 Phase 2 + Task 2.0 +
/// `entering_trade` + `is_close` so wire-up audits have a stable
/// search anchor.
#[test]
fn sp20_label_at_open_per_env_registered_fold_reset() {
let registry = StateResetRegistry::new();
let entry = registry
.entries
.iter()
.find(|e| e.name == "label_at_open_per_env")
.expect("label_at_open_per_env must be registered in StateResetRegistry");
assert_eq!(
entry.category,
ResetCategory::FoldReset,
"label_at_open_per_env must be FoldReset (got {:?})",
entry.category
);
assert!(
entry.description.contains("SP20 Phase 2"),
"description must reference 'SP20 Phase 2'; got: {}",
entry.description
);
assert!(
entry.description.contains("entering_trade"),
"description must reference 'entering_trade' producer site; got: {}",
entry.description
);
assert!(
entry.description.contains("is_close"),
"description must reference 'is_close' consumer site; got: {}",
entry.description
);
}
}

16
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -9538,6 +9538,22 @@ impl DQNTrainer {
collector.sp20_emas_obs_count_buf.host_slice_mut().fill(0);
}
}
// SP20 Phase 2 Task 2.0 (2026-05-09): per-env trade-open label
// sign scratch. Device-resident `CudaSlice<i32>` (NOT mapped-
// pinned, because the kernel ONLY reads/writes from device
// threads — no host visibility needed). Reset path:
// `stream.memset_zeros(&mut buf)` — async GPU memset, same
// pattern as `gpu_attention.rs::reset_eval_v_range_state`'s
// `attn_m/attn_v` zeroing.
"label_at_open_per_env" => {
if let Some(ref mut collector) = self.gpu_experience_collector {
let stream = collector.stream().clone();
stream.memset_zeros(&mut collector.label_at_open_per_env)
.map_err(|e| crate::MLError::ModelError(format!(
"label_at_open_per_env memset_zeros: {e}"
)))?;
}
}
// SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
// OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
// a stateful EMA) — rewrite the constructor's value at fold

93
docs/dqn-wire-up-audit.md
View File

@@ -2,6 +2,99 @@
**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.0: per-env trade-open label-sign infrastructure (additive)
Lands the producer-side infrastructure for Phase 2's 4-quadrant reward
kernel: a per-env `[alloc_episodes]` `i32` device buffer
(`label_at_open_per_env`) populated at every `entering_trade` branch
in `experience_env_step` from the upstream `aux_sign_label_per_step_kernel`
output. The Task 2.2 follow-up commit will land the consumer (the
`is_close` branch reading `label_at_open_per_env[i]` and passing it to
`sp20_compute_event_reward(close_pnl, label_at_open_sign,
isv[LOSS_CAP_INDEX])`).
### Why a separate per-env buffer (NOT a derived-at-trade-close read)
The 4-quadrant reward semantically requires the label sampled at
**trade-open**, not at close-bar. Reading `aux_label_per_env[i]` at
`is_close` would sample a different bar than the model's entry decision
was based on — that turns the dir_match check into "did the model close
when the immediate-next-bar's prediction agrees with the close P&L"
instead of "did the model close when the entry-bar's prediction agreed
with the close P&L". The trade-open sample is the load-bearing variant
per spec §4.1.
### Layout / sign mapping
| Slot value | Meaning |
|------------|------------------------------------------------------------------|
| `+1` | Aux next-bar predicted "up" at trade-open |
| `-1` | Aux next-bar predicted "not-up" (down/flat) at trade-open |
| `0` | Sentinel: no position open OR aux skip-window at trade-open |
`aux_sign_label_per_step_kernel.cu:77` emits `0`/`1` for not-up/up and
`-1` for skip-window. The Task 2.0 producer at `experience_env_step`'s
`entering_trade` branch maps `aux_l ∈ {1, 0, -1}` → `{+1, -1, 0}` to
match the SP20 spec sign convention (spec §4.1: `label_at_open =
sign(SP19_blended_label[trade_open_bar])`).
### FoldReset semantics
Registry entry `label_at_open_per_env` (FoldReset). Reset path:
`stream.memset_zeros(&mut buf)` (async GPU memset — same pattern as
`gpu_attention.rs::reset_eval_v_range_state`'s `attn_m/attn_v`
zeroing). Without this reset, leftover label from the previous fold's
open trade would corrupt the new fold's first trade-close `dir_match`
check (each fold has independent aux head calibration).
### Kernel-arg threading
`experience_env_step` signature gains 2 new args at the end (no
re-ordering of existing args):
```c
const int* __restrict__ aux_label_per_env, // [N] input
int* __restrict__ label_at_open_per_env // [N] output
```
Both are NULL-tolerant — test scaffolds without aux-head wiring
continue to work; the NULL-guarded write is a no-op and the
FoldReset-zeroed buffer reads as sentinel 0 at the consumer (Task 2.2).
### Files modified
| File | Status | Purpose |
|------|--------|---------|
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | +2 args, +write at L2887 | Producer: aux_l → label_at_open_sign mapping at `entering_trade` |
| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | +field, +alloc, +launch args | Buffer + alloc + threading into env_step launch |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | +entry + test | FoldReset registration + invariant test |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | +match arm | `reset_named_state` dispatch arm |
| `docs/dqn-wire-up-audit.md` | This entry | Audit log |
### Verification
```
SQLX_OFFLINE=true cargo check -p ml
SQLX_OFFLINE=true cargo test -p ml --lib sp20
SQLX_OFFLINE=true cargo test -p ml --lib sp20_label_at_open_per_env_registered_fold_reset
```
CPU tests (Task 2.0): registry invariant test + 18 Task 0/Phase 1 tests
all green. GPU oracle tests for the producer + consumer land in Task 2.1
(reward function math) and Task 2.2 (full trade-close path swap).
### Phase 2 → Task 2.1 / 2.2 forward references
- Task 2.1: `sp20_compute_event_reward` device function definition + 6
unit tests via `sp12_reward_math_test_kernel.cu`-style pattern.
- Task 2.2: replaces SP12 v3 inlined reward block at the
`segment_complete && segment_hold_time > 0.0f` site with the SP20
4-quadrant reward, threading `label_at_open_per_env[i]` as the
`label_at_open_sign` arg. Same atomic commit lands the
per-env `alpha_per_env[N]` plumbing through to the SP20 aggregation
kernel's `alpha` field (which currently emits 0.0 as a Phase 2
forward-reference placeholder).
## 2026-05-09 — volume_bar_size in cache key + OFI front-month filter fix
`crates/ml/src/trainers/dqn/data_loading.rs` updated in two places: