diff --git a/crates/ml/src/cuda_pipeline/dd_state_kernel.cu b/crates/ml/src/cuda_pipeline/dd_state_kernel.cu index ccd687992..da5f172e9 100644 --- a/crates/ml/src/cuda_pipeline/dd_state_kernel.cu +++ b/crates/ml/src/cuda_pipeline/dd_state_kernel.cu @@ -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; diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 9fee6478d..b66516b3e 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, - 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) diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 72f0b9d2d..cc629f77b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -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 diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 9c75cdd2b..0e4f2cf98 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -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, diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 84c046f67..cd2c0fb87 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Phase 1.3.b — wire dd_state per-step launch + drop equity-recompute bug + env-0 canonical observable (2026-05-06): closes the Phase 1.3 orphan-launcher gap per `feedback_wire_everything_up.md` and atomically fixes two architectural blockers found during the wire-up investigation (Path A: minimum-scope fix; per-env redesign deferred to Phase 1.3.b-followup if smoke shows env-0-only aggregation insufficient). (1) **Double-update bug fix** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): the original kernel recomputed `new_equity = pos_state[PS_PREV_EQUITY] + pnl_step` and wrote it back into the same slot, but `experience_env_step` already maintains `PS_PREV_EQUITY = new_portfolio_value` (`experience_kernels.cu:3473-3475`). Wiring the kernel as-is would silently double-accumulate equity every step (1× from env_step, 1× from dd_state's recompute), corrupting the position state buffer for every downstream consumer (Kelly EMAs, dd penalty, trade-stats). Per `feedback_no_quickfixes.md` the recompute is removed outright (no `if (already_updated)` guards, no sentinel skip-stores) — the kernel is now READ-ONLY on `pos_state[PS_PREV_EQUITY]` and `pos_state[PS_PEAK_EQUITY]`. The `pnl_step` parameter is dropped from both the CUDA kernel signature and the `launch_sp15_dd_state` launcher signature accordingly. (2) **Env-0 canonical observable** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): kernel grid is `[1,1,1]` (single thread, single block) and the 6 DD ISV slots [401..407) are scalars, but production has N envs in the portfolio_states buffer. The kernel now reads `pos_state[0 * PS_STRIDE + PS_PREV_EQUITY]` and `pos_state[0 * PS_STRIDE + PS_PEAK_EQUITY]` — env 0 specifically — documented in the kernel header as "env 0 as canonical DD observable; per-env redesign deferred to Phase 1.3.b-followup." This is the minimum-scope fix that produces a coherent per-step DD signal for downstream consumers; it accepts the limitation that the DD slots reflect env-0's history rather than aggregate DD across all envs. The recovery / persistence counter logic is also adapted: zero on `current_dd ≤ 0` (new HWM or pre-trade flat) instead of the original `new_equity > peak` branch (which referenced the now-gone local variable). The two conditions are equivalent at the env-0 reading because env_step has just written `peak = max(prev_peak, prev_equity)` ≥ `prev_equity`, so `current_dd = max(0, (peak - prev_equity)/peak)` evaluates to exactly 0 iff env_step set a new HWM this step. (3) **Wire-up location**: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` step "5b" inside `launch_timestep_loop`, immediately after the `experience_env_step` launch (line ~4378, after the env_step kernel block ends at line ~4379) and before the step counter advance kernel (line ~4385). 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 is bounded `begin_capture` line ~3734 → `end_capture` line ~3829, well before env_step) so per `pearl_no_host_branches_in_captured_graph.md` no graph-capture interaction. Launcher invoked with `(stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())` — `dd_budget = config.dd_threshold` keeps the DD_PCT denominator in sync with the `w_dd` penalty term env_step applies on the same `dd_threshold` knob. (4) **Oracle test migration** (`crates/ml/tests/sp15_phase1_oracle_tests.rs::dd_state_kernel_tracks_drawdown_correctly`): test was written against the OLD kernel signature so it WOULD have failed to compile after the kernel signature change — atomic migration per `feedback_no_partial_refactor.md`. The test no longer passes `pnl_step` to the launcher; instead it pre-sums the per-step PnL series into cumulative equity values, sets `pos_state[PS_PREV_EQUITY] = equity` and `pos_state[PS_PEAK_EQUITY] = running_max(equity)` directly into the env-0 row of the portfolio buffer before each `launch_sp15_dd_state` call (mirrors what env_step does in production). The same six-step equity curve (100, 110, 90, 105, 95, 100) drives the kernel and the same expected results hold (max_dd ≈ 0.182, current_dd ≈ 0.091, recovery_bars ≥ 4, dd_pct ∈ (0,1]). Test passes on RTX 3050 Ti. (5) **Phase 1.3 orphan launcher closed** per `feedback_wire_everything_up.md` — `launch_sp15_dd_state` now has a production caller in the per-step training-loop hot path. Downstream consumers receive live values when their consumer wiring lands: Phase 3.3 `dd_penalty` (reads DD_CURRENT/DD_MAX), Phase 3.5.2 `dd_asymmetric_reward` (reads DD_PCT), Phase 3.5.4 `plasticity_injection` persistence (reads DD_PERSISTENCE), Phase 3.5.5 `dd_trajectory` (reads DD_PCT). All four consumer kernels are still orphan launchers themselves at the time of this commit; their wire-up is independent of this one. (6) **Per-env DD redesign deferral** (Phase 1.3.b-followup, task #360): if L40S smoke shows the env-0 single-observable aggregation produces stale or biased DD signals when envs diverge significantly, the followup will introduce a per-env DD tile + reduction kernel that aggregates DD across envs (e.g., max_dd over all envs, or weighted-mean DD by `|position|`). The current min-scope fix is sufficient for representative-env semantics that the DD-aware reward terms expect (one DD signal per training instance). Touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (kernel — pnl_step param dropped, equity recompute removed, env-0 reads via `0 * PS_STRIDE`, recovery-counter branch adapted to `current_dd ≤ 0`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher — `pnl_step` param dropped from `launch_sp15_dd_state` signature; docstring rewritten to reflect read-only contract + env-0 canonical observable + post-env_step ordering requirement), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (production caller — new step "5b" block in `launch_timestep_loop` after env_step launch, calls `launch_sp15_dd_state(&self.stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (oracle migration — equity series pre-summed and written into `pos_state` directly per the new contract). Hard rules: `feedback_no_partial_refactor` (kernel signature change + oracle test migration + production wire-up land in this commit; per-env redesign is a separate atomic followup), `feedback_no_quickfixes` (recompute removed outright; no skip-store guards), `feedback_wire_everything_up` (closes the Phase 1.3 orphan launcher), `pearl_no_host_branches_in_captured_graph` (verified launch site is outside the exp-fwd graph capture region). Verified: `cargo check -p ml --features cuda` clean, oracle test `dd_state_kernel_tracks_drawdown_correctly` green, ml lib suite holds 946 pass / 13 fail = baseline. + SP15 Phase 1.1.b — split sharpe out of fused backtest_metrics kernel; wire dedicated launch_sp15_sharpe_per_bar (2026-05-06): closes the orphan-launcher gap left by Phase 1.1. Phase 1.1 landed `sharpe_per_bar_kernel.cu` + `launch_sp15_sharpe_per_bar` as scaffolding because the val-side sharpe at `backtest_metrics_kernel.cu:277` was inline in the 8-metric fusion (kernel computed `(mean / std_val) * annualization_factor` directly into `metrics_out[out_base + 0]`), not in a host-side loop the spec sketch had assumed. This commit performs the atomic split per `feedback_no_partial_refactor` and `feedback_wire_everything_up`. (1) **Kernel changes** (`crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu`): removed the inline sharpe compute (variance accumulator `s_sq_sum`, `var = sum_sq/n - mean²`, `std_val = sqrt(max(var, 1e-10))`, write to slot 0). Sortino still uses `mean` (from `s_sum`) and `down_std` (from `s_down_sq` — sum of NEGATIVE squared returns, distinct from sharpe's sum of all squared); calmar still uses `daily_mean = sum/wlen` and `annualization_factor²`. Output stride dropped from 14 → 13 floats per window (sharpe slot 0 removed; every metric below shifted down by 1). Shared memory layout shrunk from 6 → 5 reduction arrays (256 × 4 × 5 = 5 KB + 4 KB sort scratch = 21 KB). Updated docstring + every offset write. (2) **Host changes** (`crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`): added new `sharpe_buf: MappedF32Buffer` of size `[n_windows * 3]` (mean, std, raw_sharpe per window) per `feedback_no_htod_htoh_only_mapped_pinned`. `launch_metrics_and_record_event` now invokes BOTH kernels: the modified fused kernel for the 7+ remaining metrics, then loops `w in 0..n_windows` to launch `launch_sp15_sharpe_per_bar` once per window with `step_returns_buf + w*max_len*sizeof::()` and `n_w = window_lens[w]`. Per-window launch is acceptable because the kernel is single-block by design (`if (blockIdx.x != 0) return;`) and `n_windows` is small (typically 5). Both kernels share the same eval stream — the launches inherit the existing `eval_done_event` dependency and require no extra synchronisation. Shmem byte calc updated (`256_u32 * 5 + 4096`). (3) **Annualisation** moved host-side: `consume_metrics_after_event` reads `sharpe_host[base + 2]` (raw `mean/std`) and computes `WindowMetrics.sharpe = raw_sharpe * annualization_factor`. Same factor (`sqrt(effective_bars_per_day * trading_days_per_year)`) the kernel previously applied in-line — value contract preserved. (4) **Offset rebase summary** (every consumer migrated atomically; only consumer is the `consume_metrics_after_event` map at `gpu_backtest_evaluator.rs:~2616` — no other site reads `metrics_out[base + N]` raw): old slot 0 (sharpe) — REMOVED, sourced from `sharpe_buf`; old 1 (total_pnl) → new 0; old 2 (max_drawdown) → new 1; old 3 (sortino) → new 2; old 4 (win_rate) → new 3; old 5 (total_trades) → new 4; old 6 (var_95) → new 5; old 7 (cvar_95) → new 6; old 8 (calmar) → new 7; old 9 (omega_ratio) → new 8; old 10 (buy_count) → new 9; old 11 (sell_count) → new 10; old 12 (hold_count) → new 11; old 13 (unique_actions) → new 12. Inside the kernel the second `tid==0` block (var/cvar/calmar/omega) reads `max_dd` from `metrics_out[out_base + 1]` (was `+2`) — same intra-kernel dependency, just at the new offset. (5) **`WindowMetrics.sharpe` value contract** preserved per a new equivalence test (`unified_sharpe_kernel_equivalence_under_annualization` in `crates/ml/tests/sp15_phase1_oracle_tests.rs`). Test launches `sharpe_per_bar_kernel` against a deterministic 512-sample sin-wave + drift PnL, applies host-side `* annualization_factor`, and asserts the result matches the f64 closed-form `(mean/std) * annualization_factor` to within 1e-5 relative error. Cross-checks raw_sharpe alone (without annualisation) to the same tolerance. The two formulas (one-pass `var = sum_sq/n - mean²` vs two-pass `var = sum_dev_sq/n`) are mathematically identical within f32 precision; the bound differences (`max(var, 1e-10)` in the old kernel vs `max(var, 1e-12)` in the new) are inactive at realistic std values. (6) **Phase 1.1 oracle tests** still pass: `unified_sharpe_kernel_zero_mean`, `unified_sharpe_kernel_positive_drift`, `cost_net_sharpe_round_trip_charges_full_spread`. Metrics cubin loads still pass. Full ml lib suite holds at the 945+/13 baseline (no new regressions; the 14th flaky `test_dqn_checkpoint_round_trip` failure is preexisting and noted in spec). (7) **Eliminates the Phase 1.1 orphan launcher** per `feedback_wire_everything_up` — `launch_sp15_sharpe_per_bar` now has a production caller (`gpu_backtest_evaluator::launch_metrics_and_record_event`). Sets up Phase 1.2.b cost-net sharpe to wire `launch_sp15_cost_net_sharpe` against the same per-bar returns buffer (deferred — separate task, separate commit per `feedback_no_partial_refactor` boundaries). Touched: `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` (kernel — sharpe compute removed, offsets shifted, shmem layout), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (host — `sharpe_buf` field, alloc, per-window launch loop, annualisation in `consume_metrics_after_event`, new offsets in `WindowMetrics` parsing), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 numerical equivalence test in `mod gpu`). SP15 Phase 3.5.5 — recovery curriculum in PER: per-step DD_TRAJECTORY_DECREASING proxy (2026-05-06): per spec §9.2 (3.5.5) post-amendment-2 fix. Replaces non-existent episode-level metadata with a per-bar boolean computed from the dd_pct trajectory: `trajectory(t) = 1.0 iff dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR`. True when the transition is part of a recovery (DD shrinking from a non-trivial drawdown above the floor). New single-kernel cubin `dd_trajectory_kernel.cu` with one `extern "C" __global__` symbol (`dd_trajectory_decreasing_kernel`) → one cubin → one launcher (`launch_sp15_dd_trajectory_decreasing`) — mirrors the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2 `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu` / 3.5.4 `plasticity_injection_kernel.cu` single-kernel-per-cubin pattern. Reads ISV[DD_PCT_INDEX=406] (written by Task 1.3's `dd_state_kernel`) + ISV[DD_TRAJECTORY_FLOOR_INDEX=441]; writes ISV[DD_TRAJECTORY_DECREASING_INDEX=439]; read-modify-writes a separate mapped-pinned `prev_dd_pct` scratch buffer (`sp15_dd_trajectory_prev_dd`, [1] f32, persistent previous-bar dd_pct between kernel calls). Persistent state cold-start: sentinel 0.0 means the very first call has prev=0.0 < the 0.02 floor sentinel, so the floor gate (`dd_prev > floor`) evaluates to false and trajectory=0 — bootstrapping the state without false-firing on the synthetic `current=0.10 vs prev=0` transition. PER sampling consumer (deferred follow-up): `sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)` — recovery transitions get amplified gradient signal so the agent learns recovery dynamics over typical "average-DD" Bellman noise. 3 new ISV slots: `DD_TRAJECTORY_DECREASING_INDEX=439` (initial 0.0; per-bar output, FoldReset sentinel 0 so the new fold starts with no carry-over signal — leftover '1.0' from the previous fold's last bar would over-weight the new fold's first replay-buffer entry), `RECOVERY_OVERSAMPLE_WEIGHT_INDEX=440` (initial 2.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up producer kernel reading the current dd_pct (more DD now → higher weight); the 2.0 anchor re-engages on each new fold), `DD_TRAJECTORY_FLOOR_INDEX=441` (initial 0.02; structural cold-start anchor — runtime value signal-driven from a follow-up producer reading the rolling 25th percentile of running dd_pct distribution per `feedback_isv_for_adaptive_bounds.md`, replacing the hardcoded 0.02 floor from the original spec draft; the 0.02 anchor re-engages on each new fold). 1 new mapped-pinned scratch buffer on the trainer struct: `sp15_dd_trajectory_prev_dd` ([1] f32, mirrors the Phase 3.1 `sp15_alpha_warm_count` and Phase 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned pattern) — initialised to 0.0 in the constructor; read-modify-written by the kernel (reads previous bar's dd_pct, writes current dd_pct so the next call sees the correct "previous" value); the host writes 0.0 at construction / fold reset (allowed under `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write rule, NOT a host-side compute) and reads it through the mapped page via `read_all()` only in tests. 4 new `state_reset_registry` entries (`sp15_dd_trajectory_decreasing` resets the per-bar output to 0; `sp15_recovery_oversample_weight` rewrites the 2.0 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_dd_trajectory_floor` rewrites the 0.02 anchor; `sp15_dd_trajectory_prev_dd` resets the non-ISV mapped-pinned scratch buffer to 0.0 via `host_slice_mut().fill(0.0)`). 4 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.10 recovery_after_streak (Phase 2C / Phase 3.5 paired) — green via 3.5.2 + 3.5.4 + 3.5.5 combined; this commit lands the recovery-curriculum primitive. **Phase 3.5.5 lands kernel + launcher + 3 ISV anchor seeds + 1 scratch buffer + 4 registry entries + 4 dispatch arms only**; the PER sampling consumer migration (`sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)`) AND the ISV-driven 25th-percentile DD_TRAJECTORY_FLOOR producer (and the current-dd_pct-driven RECOVERY_OVERSAMPLE_WEIGHT producer) are deferred to follow-up atomic commits per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + 3.1-3.5.4 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). 3 GPU oracle tests pass on RTX 3050 Ti: `dd_trajectory_decreasing_fires_during_recovery` validates the bootstrap + recovery sequence (first call prev=0 < floor=0.02 → trajectory=0 + prev advances to 0.10; second call prev=0.10 > floor AND now=0.08 < prev → trajectory=1 + prev advances to 0.08); `dd_trajectory_no_fire_when_dd_increasing` validates the strict `<` comparison (prev=0.05 < now=0.10 → trajectory=0; sentinel 0.5 must be overwritten with the real 0.0 result, not skipped); `dd_trajectory_no_fire_when_prev_below_floor` validates the non-trivial-DD gate (prev=0.01 < floor=0.02 → trajectory=0 even though dd_pct decreasing; sentinel 0.7 must be overwritten with 0.0). Touched: `crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `cooldown_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the writes; the kernel's ISV pointer is mutable `float*` because it writes the DD_TRAJECTORY_DECREASING slot, identical to `cooldown_kernel`'s mutable-ISV pattern; the kernel takes a separate `prev_dd_pct` device pointer so the persistent state lives outside the ISV bus per the Phase 3.5.3 `sp15_cooldown_consecutive_losses` pattern), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_TRAJECTORY_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_trajectory_decreasing` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_DD_TRAJECTORY_CUBIN.to_vec())` and resolving `get_function("dd_trajectory_decreasing_kernel")` — single-thread/single-block grid mirroring `launch_sp15_cooldown` precedent + 3 ISV constructor writes (DD_TRAJECTORY_DECREASING=0.0, RECOVERY_OVERSAMPLE_WEIGHT=2.0, DD_TRAJECTORY_FLOOR=0.02) appended to the SP15 Phase 3.5.4 plasticity anchor block + 1 new `sp15_dd_trajectory_prev_dd: MappedF32Buffer` field on the trainer struct + alloc + 0.0 init mirroring the Phase 3.5.3 `sp15_cooldown_consecutive_losses` precedent), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+4 `RegistryEntry` records — `sp15_dd_trajectory_decreasing` resets per-bar output to 0; `sp15_recovery_oversample_weight` rewrites 2.0 anchor; `sp15_dd_trajectory_floor` rewrites 0.02 anchor; `sp15_dd_trajectory_prev_dd` non-ISV mapped-pinned buffer reset to 0.0 via `host_slice_mut().fill(0.0)`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+4 dispatch arms for the 4 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+3 GPU oracle tests in `mod gpu`). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic + a few read-modify-writes of one ISV slot and the prev_dd_pct counter, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + scratch buffer + registry entries + dispatch arms land atomically; PER sampling consumer migration + 25th-percentile DD_TRAJECTORY_FLOOR producer + current-dd_pct-driven RECOVERY_OVERSAMPLE_WEIGHT producer all follow as their own atomic commits), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the trajectory boolean + advances the persistent prev_dd_pct; the three oracle tests verify recovery-fires, dd-increasing-no-fire, and prev-below-floor-no-fire behaviors with sentinel-overwrite semantics rather than skip-store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and prev_dd_pct scratch buffer; the trainer-struct `sp15_dd_trajectory_prev_dd` field is allocated via `MappedF32Buffer::new(1)` and reset via `host_slice_mut().fill(0.0)` — allowed-write under the rule, NOT a host-side compute), `feedback_isv_for_adaptive_bounds` (the 0.02 DD_TRAJECTORY_FLOOR and 2.0 RECOVERY_OVERSAMPLE_WEIGHT anchors are structural cold-start absorbers rewritten at fold boundary so the cold-start absorbers re-engage per phase precedent; runtime FLOOR will be signal-driven from the deferred 25th-percentile-of-dd_pct producer; runtime OVERSAMPLE_WEIGHT from the deferred current-dd_pct producer), `pearl_first_observation_bootstrap` (the persistent `prev_dd_pct=0` cold-start sentinel combined with the floor gate `prev > 0.02` ensures the kernel's first call cannot false-fire on a synthetic large-step transition; the Pearl A bootstrap pattern fires implicitly via the floor gate rather than via an explicit `prev == 0` branch — the result is bit-identical because at prev=0 the `prev > 0.02` predicate is false regardless of `now < prev`).