feat(sp15-p1.3.b): wire dd_state per-step launch + drop equity-recompute bug + env-0 canonical observable
Path A of the blocked 1.3.b investigation: fixes two architectural issues atomically and wires the launcher. (1) Bug fix: dd_state_kernel.cu was recomputing new_equity = PS_PREV_EQUITY + pnl_step and writing it back, but experience_env_step already maintains PS_PREV_EQUITY (experience_kernels.cu:3473-3475) — wiring as-is would silently double-accumulate equity every step. Kernel now READS PS_PREV_EQUITY / PS_PEAK_EQUITY only; does not modify them. pnl_step parameter dropped from both kernel and launcher signatures. (2) Per-env shape decision: kernel is single-thread/single-block; production has N envs but DD ISV slots [401..407) are scalars. Picks 'env 0 as canonical observable' — kernel reads pos_state[0 * PS_STRIDE + ...]. Per-env redesign (per-env tiles + reduction kernel) deferred to Phase 1.3.b-followup if L40S smoke shows single-env DD aggregation is insufficient. (3) Wire-up: launch added at gpu_experience_collector.rs step 5b in launch_timestep_loop, immediately after env_step writes PS_PREV_EQUITY, outside the exp-fwd graph capture region (which ends at line ~3829, well before env_step). Atomic per feedback_no_partial_refactor: kernel signature change + oracle test update + launcher call site update all in this commit. Eliminates the Phase 1.3 orphan launcher per feedback_wire_everything_up. Downstream Phase 3.3 / 3.5.2 / 3.5.4 / 3.5.5 readers will receive live DD values when their consumer wiring lands. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -2,10 +2,34 @@
|
||||
//
|
||||
// SP15 Phase 1.3 — per-step drawdown state tracking.
|
||||
//
|
||||
// Reads existing PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) from
|
||||
// the position state buffer (canonical equity tracking — no new ISV slot
|
||||
// needed; invariant verified during plan v2 critical review). Writes 6
|
||||
// ISV slots per spec §6.3:
|
||||
// READ-ONLY on the position-state buffer's equity fields. Reads existing
|
||||
// PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) — both maintained
|
||||
// by `experience_env_step` (`experience_kernels.cu:3473-3475`, which
|
||||
// writes ps[PS_PEAK_EQUITY]=peak_equity and
|
||||
// ps[PS_PREV_EQUITY]=new_portfolio_value at the end of each env step).
|
||||
// This kernel does NOT recompute equity and does NOT write back to
|
||||
// PS_PEAK_EQUITY / PS_PREV_EQUITY — the env-step kernel is the canonical
|
||||
// equity writer. SP15 Phase 1.3.b atomic fix (Path A): the original
|
||||
// kernel additionally recomputed `new_equity = PS_PREV_EQUITY + pnl_step`
|
||||
// and wrote it back, which would double-accumulate equity once wired
|
||||
// after env_step (which already writes the same slot). Per
|
||||
// `feedback_no_quickfixes.md` the recompute is removed outright (no
|
||||
// guards, no "already updated" sentinels). The `pnl_step` parameter is
|
||||
// dropped from the kernel signature accordingly.
|
||||
//
|
||||
// Per-env shape: kernel is single-thread / single-block ([1,1,1] grid)
|
||||
// and the 6 ISV slots [401..407) are scalars. Production has N envs.
|
||||
// SP15 Phase 1.3.b chooses **env 0 as the canonical DD observable**:
|
||||
// the kernel reads `pos_state[0 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY]` and
|
||||
// `pos_state[0 * PORTFOLIO_STRIDE + PS_PREV_EQUITY]`. The DD ISV slots
|
||||
// therefore reflect env-0's equity history, which matches the
|
||||
// representative-env semantics the DD-aware reward terms expect (one DD
|
||||
// signal per training instance). A per-env redesign (per-env DD tile +
|
||||
// reduction kernel writing aggregate DD into ISV) is deferred to Phase
|
||||
// 1.3.b-followup if L40S smoke shows the env-0 single-observable
|
||||
// aggregation is insufficient.
|
||||
//
|
||||
// Writes 6 ISV slots per spec §6.3:
|
||||
//
|
||||
// ISV[DD_CURRENT_INDEX=401] : current_dd = max(0, (peak − equity) / peak)
|
||||
// ISV[DD_MAX_INDEX=402] : running max of current_dd
|
||||
@@ -29,51 +53,58 @@
|
||||
// `#define`d in state_layout.cuh — use the macros so the kernel tracks
|
||||
// any future portfolio-state reshuffle automatically.
|
||||
//
|
||||
// Allowed-write rule: the kernel writes ISV slots that the registry has
|
||||
// FoldReset entries for (sentinels documented in state_reset_registry.rs)
|
||||
// — pos_state[PS_*] writes go through the same slots that the env-step
|
||||
// kernel already owns, so this kernel does NOT introduce new pos_state
|
||||
// ownership; it only updates the existing high-water mark and per-step
|
||||
// equity slots which the env-step kernel updates from gross PnL normally.
|
||||
// Phase 1.3 lands the kernel + launcher + registry only (per
|
||||
// `feedback_no_partial_refactor.md`); production wiring lands in a
|
||||
// follow-up atomic commit alongside the HEALTH_DIAG composer.
|
||||
// Allowed-write rule: the kernel writes ONLY to ISV slots that the
|
||||
// registry has FoldReset entries for (sentinels documented in
|
||||
// state_reset_registry.rs). It performs zero writes on the position
|
||||
// state buffer (read-only contract), so it does not contend with
|
||||
// env_step's pos_state ownership and does not introduce new pos_state
|
||||
// ownership. The DD ISV slots [401..407) are owned solely by this
|
||||
// kernel.
|
||||
|
||||
#include "state_layout.cuh"
|
||||
|
||||
// PS_STRIDE is defined in state_layout.cuh (== 43 floats per env). Env 0
|
||||
// offset is therefore `0 * PS_STRIDE = 0`. The Rust constant
|
||||
// `PORTFOLIO_STRIDE` mirrors this same value in
|
||||
// `gpu_experience_collector.rs` (cross-checked at constructor init);
|
||||
// the layout-fingerprint check guards drift.
|
||||
|
||||
extern "C" __global__ void dd_state_kernel(
|
||||
float pnl_step,
|
||||
float dd_budget,
|
||||
float* __restrict__ isv,
|
||||
float* __restrict__ pos_state
|
||||
const float* __restrict__ pos_state
|
||||
) {
|
||||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||||
|
||||
// Update equity (PS_PREV_EQUITY = slot 9): new_equity = prev_equity + pnl_step.
|
||||
const float prev_equity = pos_state[PS_PREV_EQUITY];
|
||||
const float new_equity = prev_equity + pnl_step;
|
||||
pos_state[PS_PREV_EQUITY] = new_equity;
|
||||
|
||||
// Update peak (PS_PEAK_EQUITY = slot 7) and recovery counters.
|
||||
float peak = pos_state[PS_PEAK_EQUITY];
|
||||
if (new_equity > peak) {
|
||||
peak = new_equity;
|
||||
pos_state[PS_PEAK_EQUITY] = peak;
|
||||
// New high-water mark: zero recovery + persistence counters.
|
||||
isv[403] = 0.0f; // DD_RECOVERY_BARS_INDEX
|
||||
isv[404] = 0.0f; // DD_PERSISTENCE_INDEX
|
||||
} else {
|
||||
// Below peak: increment both counters by 1.
|
||||
isv[403] = isv[403] + 1.0f; // DD_RECOVERY_BARS_INDEX
|
||||
isv[404] = isv[404] + 1.0f; // DD_PERSISTENCE_INDEX
|
||||
}
|
||||
// Env 0 as canonical DD observable (see kernel header for rationale).
|
||||
// Read prev_equity (PS_PREV_EQUITY = slot 9) and peak (PS_PEAK_EQUITY = slot 7)
|
||||
// from the env-0 row of the portfolio state buffer. Both are written
|
||||
// by experience_env_step earlier this step — by launching this kernel
|
||||
// immediately after env_step on the same stream we get the live values.
|
||||
const float prev_equity = pos_state[0 * PS_STRIDE + PS_PREV_EQUITY];
|
||||
const float peak = pos_state[0 * PS_STRIDE + PS_PEAK_EQUITY];
|
||||
|
||||
// current_dd = max(0, (peak − equity) / peak), guarded against peak ≤ 0.
|
||||
const float current_dd = (peak > 0.0f)
|
||||
? fmaxf(0.0f, (peak - new_equity) / peak)
|
||||
? fmaxf(0.0f, (peak - prev_equity) / peak)
|
||||
: 0.0f;
|
||||
isv[401] = current_dd; // DD_CURRENT_INDEX
|
||||
|
||||
// Recovery + persistence counters: zero on a fresh new-high-water-mark
|
||||
// step, increment otherwise. We detect a new HWM by checking whether
|
||||
// current_dd is exactly zero (peak == prev_equity in the no-DD case).
|
||||
// env_step already updated peak to max(prev_peak, prev_equity) before
|
||||
// we read; if equity ≥ peak then current_dd == 0 ↔ new HWM (or flat
|
||||
// pre-trade equity), which is the canonical "no drawdown right now"
|
||||
// condition the recovery/persistence counters must zero on.
|
||||
if (current_dd <= 0.0f) {
|
||||
isv[403] = 0.0f; // DD_RECOVERY_BARS_INDEX
|
||||
isv[404] = 0.0f; // DD_PERSISTENCE_INDEX
|
||||
} else {
|
||||
isv[403] = isv[403] + 1.0f; // DD_RECOVERY_BARS_INDEX
|
||||
isv[404] = isv[404] + 1.0f; // DD_PERSISTENCE_INDEX
|
||||
}
|
||||
|
||||
// Running max of current_dd.
|
||||
const float prior_max = isv[402]; // DD_MAX_INDEX
|
||||
const float new_max = (current_dd > prior_max) ? current_dd : prior_max;
|
||||
|
||||
@@ -831,10 +831,18 @@ pub fn launch_sp15_cost_net_sharpe(
|
||||
}
|
||||
|
||||
/// SP15 Phase 1.3 (2026-05-06): per-step drawdown state tracking kernel.
|
||||
/// Reads existing PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9)
|
||||
/// from the position state buffer (canonical equity tracking — no new
|
||||
/// ISV equity slot; invariant verified during plan v2 critical review).
|
||||
/// Writes 6 ISV slots per spec §6.3:
|
||||
/// READ-ONLY on the position state buffer's equity fields. Reads existing
|
||||
/// PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) from env 0's row
|
||||
/// (`pos_state[0 * PORTFOLIO_STRIDE + …]`); does NOT write back to those
|
||||
/// slots — `experience_env_step` is the canonical equity writer (see
|
||||
/// `experience_kernels.cu:3473-3475`). SP15 Phase 1.3.b atomic fix
|
||||
/// (Path A): the original kernel additionally recomputed
|
||||
/// `new_equity = PS_PREV_EQUITY + pnl_step` and wrote it back, which
|
||||
/// would silently double-accumulate equity once wired after env_step
|
||||
/// (which already writes the same slot). Per `feedback_no_quickfixes.md`
|
||||
/// the recompute is removed outright; the `pnl_step` parameter is
|
||||
/// dropped from the kernel signature accordingly. Writes 6 ISV slots
|
||||
/// per spec §6.3:
|
||||
/// ISV[DD_CURRENT_INDEX=401], ISV[DD_MAX_INDEX=402],
|
||||
/// ISV[DD_RECOVERY_BARS_INDEX=403], ISV[DD_PERSISTENCE_INDEX=404],
|
||||
/// ISV[CALMAR_INDEX=405] (floored max_dd; host composes
|
||||
@@ -850,25 +858,31 @@ pub static SP15_DD_STATE_CUBIN: &[u8] =
|
||||
/// drive the kernel directly without the trainer struct; production
|
||||
/// callers invoke the same launcher.
|
||||
///
|
||||
/// Per-step semantics (spec §6.3):
|
||||
/// 1. new_equity = pos_state[PS_PREV_EQUITY] + pnl_step (writes back).
|
||||
/// 2. If new_equity > pos_state[PS_PEAK_EQUITY], update peak and zero
|
||||
/// DD_RECOVERY_BARS / DD_PERSISTENCE; else increment both by 1.
|
||||
/// 3. current_dd = max(0, (peak − new_equity) / peak); writes to
|
||||
/// Per-step semantics (spec §6.3, post-Phase-1.3.b):
|
||||
/// 1. Read `prev_equity = pos_state[0 * PORTFOLIO_STRIDE + PS_PREV_EQUITY]`
|
||||
/// and `peak = pos_state[0 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY]`
|
||||
/// (env 0 as canonical DD observable; per-env redesign deferred to
|
||||
/// Phase 1.3.b-followup).
|
||||
/// 2. current_dd = max(0, (peak − prev_equity) / peak); writes to
|
||||
/// DD_CURRENT_INDEX. Updates DD_MAX_INDEX = max(prior, current_dd).
|
||||
/// 3. If current_dd ≤ 0 (new high-water mark or pre-trade flat),
|
||||
/// zero DD_RECOVERY_BARS / DD_PERSISTENCE; else increment both by 1.
|
||||
/// 4. dd_pct = clip(current_dd / max(dd_budget, 1e-4), 0, 1) →
|
||||
/// DD_PCT_INDEX. CALMAR_INDEX = max(dd_max, 1e-4) (denominator
|
||||
/// floor; host composer divides mean_pnl by this value).
|
||||
///
|
||||
/// `pnl_step` is the per-step gross PnL delta (signed). `dd_budget` is
|
||||
/// the configured drawdown limit (e.g. 0.20 = 20%). `isv` MUST be the
|
||||
/// ISV bus (≥443 f32 slots — slots 401-406 written, 405 floored).
|
||||
/// `pos_state` MUST point to a portfolio-state buffer with at least
|
||||
/// `PS_STRIDE` f32 slots (PS_PEAK_EQUITY=7 and PS_PREV_EQUITY=9 are
|
||||
/// read + written).
|
||||
/// The kernel does NOT write to the position state buffer — env_step is
|
||||
/// the sole writer of PS_PREV_EQUITY / PS_PEAK_EQUITY. Wire callers
|
||||
/// MUST launch this kernel AFTER the env_step launch on the same stream
|
||||
/// (see `gpu_experience_collector::launch_timestep_loop` step 5b).
|
||||
///
|
||||
/// `dd_budget` is the configured drawdown limit (e.g. 0.20 = 20%).
|
||||
/// `isv` MUST be the ISV bus (≥443 f32 slots — slots 401-406 written,
|
||||
/// 405 floored). `pos_state` MUST point to a portfolio-state buffer
|
||||
/// with at least `PORTFOLIO_STRIDE` f32 slots (PS_PEAK_EQUITY=7 and
|
||||
/// PS_PREV_EQUITY=9 are READ from the env-0 row).
|
||||
pub fn launch_sp15_dd_state(
|
||||
stream: &Arc<CudaStream>,
|
||||
pnl_step: f32,
|
||||
dd_budget: f32,
|
||||
isv: cudarc::driver::sys::CUdeviceptr,
|
||||
pos_state: cudarc::driver::sys::CUdeviceptr,
|
||||
@@ -887,7 +901,6 @@ pub fn launch_sp15_dd_state(
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernel)
|
||||
.arg(&pnl_step)
|
||||
.arg(&dd_budget)
|
||||
.arg(&isv)
|
||||
.arg(&pos_state)
|
||||
|
||||
@@ -4378,6 +4378,39 @@ impl GpuExperienceCollector {
|
||||
)))?;
|
||||
}
|
||||
|
||||
// ── 5b. SP15 Phase 1.3.b: per-step drawdown state tracking ──
|
||||
// Closes the orphan launcher gap from Phase 1.3 per
|
||||
// `feedback_wire_everything_up.md`. Reads env-0's freshly-
|
||||
// written PS_PREV_EQUITY / PS_PEAK_EQUITY (env_step above
|
||||
// wrote both at experience_kernels.cu:3473-3475) and emits
|
||||
// 6 ISV slots [401..407): DD_CURRENT, DD_MAX, DD_RECOVERY_BARS,
|
||||
// DD_PERSISTENCE, CALMAR (denominator floor), DD_PCT.
|
||||
//
|
||||
// Ordering: launched on the same stream as env_step → CUDA
|
||||
// serializes the two on-stream, so the dd_state read sees
|
||||
// env_step's writes (no event sync required). Outside the
|
||||
// exp-fwd graph capture region (which ends at line ~3829,
|
||||
// well before this point), so per
|
||||
// `pearl_no_host_branches_in_captured_graph.md` no graph
|
||||
// capture interaction.
|
||||
//
|
||||
// Per-env shape: kernel is single-thread / single-block and
|
||||
// reads env-0 specifically as the canonical DD observable.
|
||||
// Per-env redesign deferred to Phase 1.3.b-followup.
|
||||
//
|
||||
// dd_budget = config.dd_threshold (0.02 default) — same
|
||||
// signal env_step uses for its w_dd penalty term, kept in
|
||||
// sync so the DD_PCT slot is comparable across consumers.
|
||||
{
|
||||
let pos_state_ptr = self.portfolio_states.raw_ptr();
|
||||
crate::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_state(
|
||||
&self.stream,
|
||||
config.dd_threshold,
|
||||
self.isv_signals_dev_ptr,
|
||||
pos_state_ptr,
|
||||
)?;
|
||||
}
|
||||
|
||||
// ── 6. Advance GPU-resident step counter ────────────────────
|
||||
// Single-thread kernel: step_counter_gpu[0] += 1.
|
||||
// Runs on the same stream — serialized after env_step, before
|
||||
|
||||
@@ -221,14 +221,22 @@ mod gpu {
|
||||
|
||||
/// Test 1.3 — synthetic equity curve drives DD slots correctly.
|
||||
///
|
||||
/// PnL steps: [+100, +10, −20, +15, −10, +5]
|
||||
/// Cumulative equity (starting from 0 prev_equity):
|
||||
/// step 1: 0 + 100 = 100 → peak 100, dd=0, rec=0
|
||||
/// step 2: 100 + 10 = 110 → peak 110, dd=0, rec=0
|
||||
/// step 3: 110 − 20 = 90 → dd=(110-90)/110≈0.182, rec=1
|
||||
/// step 4: 90 + 15 = 105 → dd=(110-105)/110≈0.045, rec=2
|
||||
/// step 5: 105 − 10 = 95 → dd=(110-95)/110≈0.136, rec=3
|
||||
/// step 6: 95 + 5 = 100 → dd=(110-100)/110≈0.091, rec=4
|
||||
/// SP15 Phase 1.3.b contract (post Path A): the kernel is READ-ONLY
|
||||
/// on `pos_state` and consumes the env-0 row's `PS_PREV_EQUITY` /
|
||||
/// `PS_PEAK_EQUITY` slots. In production these are maintained by
|
||||
/// `experience_env_step`; the oracle test simulates that role by
|
||||
/// writing prev_equity and the running peak directly into the env-0
|
||||
/// portfolio-state row before each `launch_sp15_dd_state` call. This
|
||||
/// matches the live-system invariant that env_step writes equity
|
||||
/// first, then dd_state reads.
|
||||
///
|
||||
/// Equity curve (cumulative): [+100, +10, -20, +15, -10, +5]
|
||||
/// step 1: equity=100, peak=100, dd=0, rec=0
|
||||
/// step 2: equity=110, peak=110, dd=0, rec=0
|
||||
/// step 3: equity=90, peak=110, dd=(110-90)/110 ≈ 0.182, rec=1
|
||||
/// step 4: equity=105, peak=110, dd=(110-105)/110 ≈ 0.045, rec=2
|
||||
/// step 5: equity=95, peak=110, dd=(110-95)/110 ≈ 0.136, rec=3
|
||||
/// step 6: equity=100, peak=110, dd=(110-100)/110 ≈ 0.091, rec=4
|
||||
///
|
||||
/// Expected: max_dd ≈ 0.182 (hit at step 3), current_dd ≈ 0.091
|
||||
/// (never recovers), recovery_bars = 4 at end. dd_pct = 0.091 / 0.20
|
||||
@@ -236,9 +244,21 @@ mod gpu {
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn dd_state_kernel_tracks_drawdown_correctly() {
|
||||
// PORTFOLIO_STRIDE / PS_PEAK_EQUITY / PS_PREV_EQUITY constants
|
||||
// mirror state_layout.cuh — the kernel reads
|
||||
// `pos_state[0 * PORTFOLIO_STRIDE + PS_*]`. Update here if the
|
||||
// CUDA-side layout shifts (the layout fingerprint check enforces
|
||||
// synchrony in production).
|
||||
const PORTFOLIO_STRIDE: usize = 43;
|
||||
const PS_PEAK_EQUITY: usize = 7;
|
||||
const PS_PREV_EQUITY: usize = 9;
|
||||
|
||||
let stream = make_test_stream();
|
||||
|
||||
let pnl_steps: [f32; 6] = [100.0, 10.0, -20.0, 15.0, -10.0, 5.0];
|
||||
// Equity at the END of each step (cumulative running sum starting
|
||||
// from 0 prev_equity — same series the original test used, just
|
||||
// pre-summed to match the new "env_step writes equity first" contract).
|
||||
let equity_per_step: [f32; 6] = [100.0, 110.0, 90.0, 105.0, 95.0, 100.0];
|
||||
let dd_budget: f32 = 0.20;
|
||||
|
||||
// ISV bus sized to SP15_SLOT_END (slots 401-406 are written).
|
||||
@@ -246,18 +266,25 @@ mod gpu {
|
||||
.expect("alloc MappedF32Buffer for isv bus");
|
||||
isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]);
|
||||
|
||||
// Position state buffer — conservative size 64 floats covers
|
||||
// PS_STRIDE=43 plus headroom (the kernel only touches PS_PEAK_EQUITY
|
||||
// at slot 7 and PS_PREV_EQUITY at slot 9). Cold-start at zero so
|
||||
// the first step's new_equity = 0 + 100 = 100.
|
||||
let pos_state_buf = unsafe { MappedF32Buffer::new(64) }
|
||||
// Position state buffer sized for one env (PORTFOLIO_STRIDE) plus
|
||||
// headroom. The kernel only reads env-0's row.
|
||||
let pos_len = PORTFOLIO_STRIDE.max(64);
|
||||
let pos_state_buf = unsafe { MappedF32Buffer::new(pos_len) }
|
||||
.expect("alloc MappedF32Buffer for pos_state");
|
||||
pos_state_buf.write_from_slice(&vec![0.0f32; 64]);
|
||||
|
||||
for &pnl_step in &pnl_steps {
|
||||
// Step the env-0 equity curve manually (mirrors what env_step
|
||||
// does in production); after each write the dd_state kernel
|
||||
// reads the live values.
|
||||
let mut peak: f32 = 0.0;
|
||||
for &equity in &equity_per_step {
|
||||
peak = peak.max(equity);
|
||||
let mut pos_init = vec![0.0f32; pos_len];
|
||||
pos_init[PS_PREV_EQUITY] = equity;
|
||||
pos_init[PS_PEAK_EQUITY] = peak;
|
||||
pos_state_buf.write_from_slice(&pos_init);
|
||||
|
||||
launch_sp15_dd_state(
|
||||
&stream,
|
||||
pnl_step,
|
||||
dd_budget,
|
||||
isv_buf.dev_ptr,
|
||||
pos_state_buf.dev_ptr,
|
||||
|
||||
File diff suppressed because one or more lines are too long
Reference in New Issue
Block a user