diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 51a832c26..0a9afd411 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -843,17 +843,41 @@ fn main() { // kernel reads LobBar fields from synthetic markets and prod // fxcache LOB (dev/prod parity per Q3). "cost_net_sharpe_kernel.cu", - // SP15 Phase 1.3 (2026-05-06): drawdown reporting. Single-thread, - // single-block per-step state machine. 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, invariant verified during plan v2 critical review). - // Writes 6 ISV slots: DD_CURRENT (401), DD_MAX (402), - // DD_RECOVERY_BARS (403), DD_PERSISTENCE (404), CALMAR (405; - // floored max_dd for host composer = mean_pnl / value-here), - // DD_PCT (406). The 1e-4 calmar floor eliminates the saturation- - // at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG. + // SP15 Phase 1.3 (2026-05-06) + 1.3.b-followup (2026-05-07): + // per-env drawdown state tracking. Per-env grid `[n_envs, 1, 1]`, + // one thread per env: each thread reads PS_PEAK_EQUITY (slot 7) + // and PS_PREV_EQUITY (slot 9) from its env's row of the position + // state buffer and writes 6 scalars to the per-env tile output + // `dd_state_per_env[env*6 + k]` for k ∈ {DD_CURRENT, DD_MAX, + // DD_RECOVERY_BARS, DD_PERSISTENCE, CALMAR (floored max_dd + // denominator), DD_PCT}. The 6 ISV scalar slots [401..407) are + // populated by the separate `dd_state_reduce_kernel` below + // (mean-aggregate across envs for HEALTH_DIAG diagnostic) plus + // the new `DD_PERSISTENCE_MAX_INDEX = 443` slot (max-aggregate + // for the plasticity-injection trigger). Path A → Path B + // migration: the original Phase 1.3.b kernel chose env-0 as the + // canonical DD observable and wrote ISV scalars directly; the + // followup-A redesign threads each env's actual DD context + // through the reward shaping per spec §6.3. "dd_state_kernel.cu", + // SP15 Phase 1.3.b-followup (2026-05-07): DD state reduction. + // Single-block tree-reduce (no atomicAdd) of the per-env tile + // `dd_state_per_env[n_envs * 6]` into 6 scalar ISV slots + // [401..407) (mean-aggregate across envs — smooth across-env + // summary for HEALTH_DIAG diagnostic) plus the new + // DD_PERSISTENCE_MAX_INDEX = 443 slot (max-aggregate, consumed + // by `plasticity_injection_kernel` so the global advantage-head + // reset fires when ANY env's persistence exceeds the threshold: + // one set of advantage weights → global firing semantics → + // max-aggregate is the only correct rule). BLOCK=256 with the + // strided initial accumulation pattern handles n_envs up to + // production sizes (32768 envs on H100). Layout fingerprint + // marker `DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` enforces + // synchrony between the producer's tile layout (in + // `dd_state_kernel.cu`) and this consumer's reads. Single + // kernel per cubin (mirrors the SP15 1.3 / 1.4 / 3.X + // single-source-to-cubin convention). + "dd_state_reduce_kernel.cu", // SP15 Phase 1.4 + Wave 3a (2026-05-06): 4 constant-policy // counterfactual baselines per spec §6.4. Single source file with // 4 `extern "C" __global__` symbols (baseline_buyhold_kernel, diff --git a/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu b/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu index 2c5a6f056..2b12b9503 100644 --- a/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu +++ b/crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu @@ -1,6 +1,7 @@ // crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu // // SP15 Wave 2 (2026-05-06) — fused post-SP11 reward-axis composer. +// SP15 Phase 1.3.b-followup (2026-05-07) — per-env DD context migration. // // Architecture: SP11 B1b's reward composer (in `experience_kernels.cu`) // remains the canonical "trader-quality" first pass. It writes the @@ -25,9 +26,18 @@ // staleness is bounded by the launch ordering // inside `gpu_experience_collector.rs`. // 2. DD asymmetric reward: gain side `r × (1 + λ × dd_pct)`, loss side -// unchanged (Phase 3.5.2). -// 3. Quadratic DD penalty: subtract `λ_dd × max(0, dd − dd_thr)²` -// (Phase 3.3). +// unchanged (Phase 3.5.2). Phase 1.3.b-followup +// (2026-05-07): `dd_pct` is now PER-ENV — looked +// up from `dd_state_per_env[env*6 + 5]` where +// `env_id = (idx % (N*L)) / L` (each (i,t) slot +// sees ITS env's actual DD context, not env-0's +// canonical observable). Fixes the Path A +// limitation where envs in different DD contexts +// all saw env-0's DD value. +// 3. Quadratic DD penalty: subtract `λ_dd × max(0, dd_current − dd_thr)²` +// (Phase 3.3). `dd_current` is also PER-ENV +// now — looked up from +// `dd_state_per_env[env*6 + 0]`. // 4. SP12 v3 asymmetric cap: `fmaxf(REWARD_NEG_CAP, fminf(POS, r))` // via `state_layout.cuh` macros. // @@ -51,6 +61,14 @@ // the diagnostic; every other thread passes NULL. The slot still gets // a meaningful update each step (one representative thread). // +// Index→env mapping (Phase 1.3.b-followup contract): `out_rewards` is +// laid out `[on_policy_N*L | cf_N*L]`. `idx ∈ [0, 2*N*L)` → +// `slot_it = idx % (N*L)` is the per-(i,t) flat index → `env_id = +// slot_it / L` and `t_local = slot_it % L`. The per-env tile +// `dd_state_per_env` is indexed by `env_id * 6 + k` for k ∈ {0..6). +// On-policy and CF threads at the SAME (i,t) read the SAME per-env +// tile entry (one DD trajectory per env, shared across slot kinds). +// // Hard rules: // - `feedback_no_atomicadd` — pure scalar arithmetic per (i,t); no // reductions, no atomics. @@ -62,6 +80,11 @@ // - `feedback_isv_for_adaptive_bounds` — α / λ / λ_dd / threshold are // ISV-driven anchors; NEG/POS caps stay as `state_layout.cuh` // macros per Q1 resolution (spec'd constants, not adaptive bounds). +// - `feedback_no_partial_refactor` — atomic per Phase 1.3.b-followup +// contract change: ALL consumers of single-env-canonical DD slots +// migrate to per-env tile lookup in this commit (this kernel + +// `plasticity_injection_kernel`); HEALTH_DIAG keeps reading the +// scalar slots, now mean-aggregated by `dd_state_reduce_kernel`. #include #include "state_layout.cuh" // REWARD_NEG_CAP, REWARD_POS_CAP @@ -72,6 +95,7 @@ extern "C" __global__ void compute_sp15_final_reward_kernel( const float* __restrict__ r_discipline_out, /* [N*L] per-step r_disc (REGRET_EMA mirror, etc.) */ const int* __restrict__ slot_completed_normally, /* [N*L] 1=normal, 0=early-return sentinel */ float* __restrict__ isv, /* writes slot 431 R_GAIN_DD_BOOST */ + const float* __restrict__ dd_state_per_env, /* [N*6] per-env DD tile (followup-A) */ int N, int L ) { @@ -82,9 +106,11 @@ extern "C" __global__ void compute_sp15_final_reward_kernel( /* Map flat (out_rewards) index → (slot_kind=on_policy/cf, i, t). * Layout: out_rewards[on_policy] at [0 .. N*L); out_rewards[cf] at * [N*L .. 2*N*L). The per-(i,t) `slot_completed_normally` flag is - * shared — index via `idx % (N*L)`. */ + * shared — index via `idx % (N*L)`. The per-env DD tile is shared + * between the slot kinds (one DD trajectory per env). */ const int n_l = N * L; const int slot_it = idx % n_l; + const int env_id = slot_it / L; /* Q3: skip early-return sentinel slots. SP11's value at * `out_rewards[idx]` (0.0f for data-end, −10.0f for blown-account) @@ -102,17 +128,30 @@ extern "C" __global__ void compute_sp15_final_reward_kernel( const float r_discipline = r_discipline_out[slot_it]; r = sp15_alpha_blend(r, r_discipline, isv); + /* Per-env DD lookup (Phase 1.3.b-followup contract). One read per + * stat — the tile is `[N * 6]` row-major over envs. */ + const int env_base = env_id * 6; + const float dd_current = dd_state_per_env[env_base + 0]; + const float dd_pct = dd_state_per_env[env_base + 5]; + /* Stage 2: DD asymmetric reward (gain-only multiplier). * Diagnostic slot 431 is written ONLY by a single representative - * thread (on-policy slot at i=0, t=0) to avoid concurrent writes. */ + * thread (on-policy slot at i=0, t=0) to avoid concurrent writes. + * λ (SP15_DD_ASYMMETRY_LAMBDA_INDEX=430) is a global controller + * scalar — read once here and threaded into the helper. */ + const float lambda_asymm = isv[SP15_DD_ASYMMETRY_LAMBDA_INDEX]; float* r_gain_boost_diag_out = nullptr; if (idx == 0) { r_gain_boost_diag_out = &isv[SP15_R_GAIN_DD_BOOST_INDEX]; } - r = sp15_dd_asymmetric_reward(r, isv, r_gain_boost_diag_out); + r = sp15_dd_asymmetric_reward(r, lambda_asymm, dd_pct, r_gain_boost_diag_out); - /* Stage 3: quadratic DD penalty (asymmetric, zero below threshold). */ - r = sp15_dd_penalty(r, isv); + /* Stage 3: quadratic DD penalty (asymmetric, zero below threshold). + * λ_dd / threshold are global controller scalars; dd_current is + * the per-env value looked up above. */ + const float lambda_dd = isv[SP15_LAMBDA_DD_INDEX]; + const float dd_thr = isv[SP15_DD_THRESHOLD_INDEX]; + r = sp15_dd_penalty(r, lambda_dd, dd_thr, dd_current); /* Stage 4: SP12 v3 asymmetric bilateral cap via macros. */ r = sp15_apply_sp12_cap(r); diff --git a/crates/ml/src/cuda_pipeline/dd_state_kernel.cu b/crates/ml/src/cuda_pipeline/dd_state_kernel.cu index da5f172e9..db759f6d2 100644 --- a/crates/ml/src/cuda_pipeline/dd_state_kernel.cu +++ b/crates/ml/src/cuda_pipeline/dd_state_kernel.cu @@ -1,6 +1,7 @@ // crates/ml/src/cuda_pipeline/dd_state_kernel.cu // // SP15 Phase 1.3 — per-step drawdown state tracking. +// SP15 Phase 1.3.b-followup (2026-05-07) — per-env redesign (Path A → Path B). // // 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 @@ -9,115 +10,126 @@ // 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. +// equity writer. // -// 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. +// Phase 1.3.b-followup migration: the original Phase 1.3.b atomic fix +// (Path A) chose **env 0 as the canonical DD observable** and wrote +// directly to 6 ISV scalar slots [401..407). That choice was wrong for +// production where each env has its own portfolio + DD trajectory: when +// env-3 was in 30% DD and env-0 was at ATH, the per-(i,t) reward shaping +// looked up env-0's `dd_pct=0` and skipped recovery shaping for env-3's +// transitions. // -// Writes 6 ISV slots per spec §6.3: +// Path B (this kernel): per-env grid `[n_envs, 1, 1]` × `[1, 1, 1]`. Each +// thread computes its own env's DD trajectory and writes 6 scalars to +// the per-env tile buffer `dd_state_per_env[env_id * 6 + k]` for +// `k ∈ {DD_CURRENT, DD_MAX, DD_RECOVERY_BARS, DD_PERSISTENCE, CALMAR, +// DD_PCT}`. Each thread's per-bar walk is sequential (the running max, +// recovery counter, and persistence counter are stateful per-env), but +// envs are independent so the grid parallelises across them. // -// ISV[DD_CURRENT_INDEX=401] : current_dd = max(0, (peak − equity) / peak) -// ISV[DD_MAX_INDEX=402] : running max of current_dd -// ISV[DD_RECOVERY_BARS_INDEX=403] : bars since last new high-water mark -// ISV[DD_PERSISTENCE_INDEX=404] : bars since last new high (twin of recovery for now; -// split lifetime semantics may differ in later phases) -// ISV[CALMAR_INDEX=405] : floored max_dd (= max(dd_max, 1e-4)) — denominator -// for the host composer's calmar = mean_pnl / value-here. -// Kernel writes ONLY the denominator floor; the host -// composes the calmar ratio so that mean_pnl source can -// evolve without kernel changes. The 1e-4 floor -// eliminates the saturation-at-100 artifact seen in the -// train-dd4xl HEALTH_DIAG. -// ISV[DD_PCT_INDEX=406] : current_dd / max(dd_budget, 1e-4), clipped to 1.0 +// The 6 ISV scalar slots [401..407) are now populated by a separate +// reduction kernel (`dd_state_reduce_kernel`) that mean-aggregates +// across envs for HEALTH_DIAG diagnostic reporting + max-aggregates +// DD_PERSISTENCE into the new `DD_PERSISTENCE_MAX_INDEX = 443` slot for +// the plasticity-injection trigger (one set of advantage weights → +// global firing semantics → max-aggregate is the only correct rule +// when ANY env exceeding the threshold should arm the gate). // -// Per-step state-machine kernel (single-thread, single-block). Mirrors the -// host-side draw-down accounting that previously lived in the CPU emit -// path; lifts the truth source onto the GPU per `feedback_cpu_is_read_only`. +// Per-env tile layout (offsets within each env's [6] sub-array): +// tile[env*6 + 0] = DD_CURRENT (current_dd = max(0, (peak-equity)/peak)) +// tile[env*6 + 1] = DD_MAX (running max of DD_CURRENT in this fold) +// tile[env*6 + 2] = DD_RECOVERY_BARS (bars since last new HWM) +// tile[env*6 + 3] = DD_PERSISTENCE (twin counter to recovery) +// tile[env*6 + 4] = CALMAR (max(dd_max, 1e-4) — calmar denominator floor) +// tile[env*6 + 5] = DD_PCT (clip(current_dd / max(dd_budget, 1e-4), 0, 1)) // // Layout fingerprint contract: PS_PEAK_EQUITY and PS_PREV_EQUITY are // `#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 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. +// Allowed-write rule: the kernel writes ONLY to the per-env tile (which +// has a FoldReset arm `sp15_dd_state_per_env`) and never touches the +// position state buffer or the ISV scalar slots. The ISV scalar slots +// are written exclusively by `dd_state_reduce_kernel`. No +// pos_state ownership conflict; no double-writer for the ISV slots. #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 +// PS_STRIDE is defined in state_layout.cuh (== 43 floats per env). Env i +// offset is therefore `i * PS_STRIDE`. 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 dd_budget, - float* __restrict__ isv, - const float* __restrict__ pos_state + float dd_budget, + int n_envs, + float* __restrict__ dd_state_per_env, /* [n_envs * 6] tile output */ + const float* __restrict__ pos_state /* [n_envs * PS_STRIDE] */ ) { - if (threadIdx.x != 0 || blockIdx.x != 0) return; + /* Per-env grid: one thread per env, one block per env (mirrors the + * single-thread / single-block shape of the original Path A kernel + * but lifted into a per-env loop via blockIdx.x). The state-machine + * inside the thread is sequential per env — running max, recovery + * counter, persistence counter — so we cannot parallelise across + * the [n_envs * 6] output dimension. Envs ARE independent of each + * other, so the gridDim.x = n_envs parallelism is straightforward. */ + if (threadIdx.x != 0) return; + const int env_id = blockIdx.x; + if (env_id >= n_envs) return; - // 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]; + /* Read prev_equity (PS_PREV_EQUITY = slot 9) and peak (PS_PEAK_EQUITY + * = slot 7) from this env's 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 for every env. */ + const float prev_equity = pos_state[env_id * PS_STRIDE + PS_PREV_EQUITY]; + const float peak = pos_state[env_id * PS_STRIDE + PS_PEAK_EQUITY]; - // current_dd = max(0, (peak − equity) / peak), guarded against peak ≤ 0. + /* current_dd = max(0, (peak − equity) / peak), guarded against peak ≤ 0. */ const float current_dd = (peak > 0.0f) ? 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. + const int base = env_id * 6; + + /* 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. */ + float recovery_bars; + float persistence_bars; if (current_dd <= 0.0f) { - isv[403] = 0.0f; // DD_RECOVERY_BARS_INDEX - isv[404] = 0.0f; // DD_PERSISTENCE_INDEX + recovery_bars = 0.0f; + persistence_bars = 0.0f; } else { - isv[403] = isv[403] + 1.0f; // DD_RECOVERY_BARS_INDEX - isv[404] = isv[404] + 1.0f; // DD_PERSISTENCE_INDEX + recovery_bars = dd_state_per_env[base + 2] + 1.0f; + persistence_bars = dd_state_per_env[base + 3] + 1.0f; } - // Running max of current_dd. - const float prior_max = isv[402]; // DD_MAX_INDEX + /* Running max of current_dd. */ + const float prior_max = dd_state_per_env[base + 1]; const float new_max = (current_dd > prior_max) ? current_dd : prior_max; - isv[402] = new_max; - // dd_pct = clip(current_dd / max(dd_budget, 1e-4), 0, 1). + /* dd_pct = clip(current_dd / max(dd_budget, 1e-4), 0, 1). */ const float dd_pct = fminf(1.0f, current_dd / fmaxf(dd_budget, 1e-4f)); - isv[406] = dd_pct; // DD_PCT_INDEX - // CALMAR slot stores the floored max_dd denominator; host composes - // calmar = mean_pnl / value-here. Floor eliminates the saturation-at-100 - // artifact that the legacy host-side composer exhibited when max_dd ≈ 0. - isv[405] = fmaxf(new_max, 1e-4f); // CALMAR_INDEX + /* CALMAR slot stores the floored max_dd denominator; host (post- + * reduction) composes calmar = mean_pnl / value-here. Floor + * eliminates the saturation-at-100 artifact that the legacy host- + * side composer exhibited when max_dd ≈ 0. */ + const float calmar_denom = fmaxf(new_max, 1e-4f); + + dd_state_per_env[base + 0] = current_dd; + dd_state_per_env[base + 1] = new_max; + dd_state_per_env[base + 2] = recovery_bars; + dd_state_per_env[base + 3] = persistence_bars; + dd_state_per_env[base + 4] = calmar_denom; + dd_state_per_env[base + 5] = dd_pct; __threadfence_system(); } diff --git a/crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu b/crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu new file mode 100644 index 000000000..a84210372 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu @@ -0,0 +1,153 @@ +// crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu +// +// SP15 Phase 1.3.b-followup (2026-05-07) — DD state reduction kernel. +// +// Aggregates the per-env DD tile written by `dd_state_kernel` into the +// 6 scalar ISV slots [401..407) plus the new `DD_PERSISTENCE_MAX_INDEX +// = 443` slot. Two distinct aggregation rules: +// +// * Mean across envs for slots 401..407): smooth across-env summary +// for HEALTH_DIAG diagnostic reporting. Preserves backward- +// compatible "this is one DD scalar to look at" semantics for the +// diagnostic stream — only the across-env aggregation rule +// changed (Path A: env-0-canonical → Path B: mean across envs). +// This is structural per `feedback_isv_for_adaptive_bounds` (the +// aggregation rule itself is fixed mean — no adaptive smoothing +// here). +// * Max across envs for `DD_PERSISTENCE_MAX_INDEX = 443`: consumed +// by `plasticity_injection_kernel` so the global advantage-head +// reset fires when ANY env's persistence exceeds the threshold. +// One set of advantage weights → global firing semantics → max- +// aggregate is the only correct rule (mean-aggregate would silently +// gate the trigger when only a fraction of envs are in long DD). +// +// Single-block tree-reduce (no atomicAdd per `feedback_no_atomicadd` — +// the existing block-tree-reduce pattern in `kelly_cap_update` / +// `cost_net_sharpe` baseline kernels). Block size `BLOCK = 256`; if +// `n_envs > 256` the reduction does a strided initial accumulation +// pass before the tree-reduce (the production sizes +// `optimal_n_episodes()` reach 32768 envs on H100, so this matters). +// `n_envs` is a launch-time scalar argument — fixed across the run, so +// the kernel grid shape is constant per `pearl_no_host_branches_in_ +// captured_graph` (CUDA Graph capture remains compatible). +// +// Layout contract (the per-env tile offsets MUST mirror +// `dd_state_kernel.cu`'s writes; layout-fingerprint marker +// `DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` in +// `layout_fingerprint_seed` enforces synchrony): +// tile[env*6 + 0] = DD_CURRENT → ISV[401] (mean) +// tile[env*6 + 1] = DD_MAX → ISV[402] (mean) +// tile[env*6 + 2] = DD_RECOVERY_BARS → ISV[403] (mean) +// tile[env*6 + 3] = DD_PERSISTENCE → ISV[404] (mean) +// → ISV[443] (max — DD_PERSISTENCE_MAX) +// tile[env*6 + 4] = CALMAR → ISV[405] (mean) +// tile[env*6 + 5] = DD_PCT → ISV[406] (mean) +// +// Hard rules: +// * `feedback_no_atomicadd` — single-block tree-reduce, no atomicAdd. +// * `feedback_no_partial_refactor` — kernel + 5 downstream consumer +// migrations land atomically with the per-env tile contract switch. +// * `feedback_no_stubs` — kernel writes the actual mean+max +// aggregates; no return-zero stubs. +// * `feedback_isv_for_adaptive_bounds` — mean / max aggregation rules +// are structural (fixed by the aggregation requirement: one shared +// global advantage-head ⇒ max-aggregate for trigger, smooth +// diagnostic ⇒ mean-aggregate for HEALTH_DIAG). Documented in +// `docs/dqn-wire-up-audit.md` Phase 1.3.b-followup entry. +// * `pearl_no_host_branches_in_captured_graph` — `n_envs` is a kernel +// launch arg (a scalar, not a host-side branch); kernel grid +// shape is constant. + +#define BLOCK_SIZE 256 + +extern "C" __global__ void dd_state_reduce_kernel( + int n_envs, + const float* __restrict__ dd_state_per_env, /* [n_envs * 6] tile input */ + float* __restrict__ isv /* [≥444] ISV bus */ +) { + // Single-block (gridDim.x == 1 enforced by the launcher's + // `grid_dim: (1, 1, 1)`). + const int tid = threadIdx.x; + + // Per-thread accumulators (one per output statistic). + // Six mean-aggregates (one per tile field) + one max-aggregate (for + // persistence). Initialise to identity (0 for sums, -inf for max). + float sum_dd_current = 0.0f; + float sum_dd_max = 0.0f; + float sum_recovery_bars = 0.0f; + float sum_persistence = 0.0f; + float sum_calmar = 0.0f; + float sum_dd_pct = 0.0f; + float max_persistence = -1.0f; // persistence is non-negative; -1 < any valid value + + /* Strided initial accumulation pass: each thread sums every (n_envs / + * BLOCK_SIZE)-th env starting from `tid`. For n_envs ≤ BLOCK_SIZE + * each thread handles at most one env (with the remaining threads + * leaving the identity initialisers untouched). For n_envs > BLOCK_SIZE + * the strided loop accumulates many envs into per-thread state + * before the tree-reduce. */ + for (int env_id = tid; env_id < n_envs; env_id += BLOCK_SIZE) { + const int base = env_id * 6; + sum_dd_current += dd_state_per_env[base + 0]; + sum_dd_max += dd_state_per_env[base + 1]; + sum_recovery_bars += dd_state_per_env[base + 2]; + const float persistence = dd_state_per_env[base + 3]; + sum_persistence += persistence; + max_persistence = fmaxf(max_persistence, persistence); + sum_calmar += dd_state_per_env[base + 4]; + sum_dd_pct += dd_state_per_env[base + 5]; + } + + /* Block-shared tree-reduce. We reduce 7 quantities; allocate a + * `[BLOCK_SIZE × 7]` shared-memory tile so each output keeps its own + * lane during the reduction. Layout (column-major over output): + * sdata[k * BLOCK_SIZE + tid] for output k ∈ [0..7). + */ + __shared__ float sdata[BLOCK_SIZE * 7]; + + sdata[0 * BLOCK_SIZE + tid] = sum_dd_current; + sdata[1 * BLOCK_SIZE + tid] = sum_dd_max; + sdata[2 * BLOCK_SIZE + tid] = sum_recovery_bars; + sdata[3 * BLOCK_SIZE + tid] = sum_persistence; + sdata[4 * BLOCK_SIZE + tid] = sum_calmar; + sdata[5 * BLOCK_SIZE + tid] = sum_dd_pct; + sdata[6 * BLOCK_SIZE + tid] = max_persistence; + __syncthreads(); + + /* Tree-reduce: halving stride, sum for outputs 0..6 and max for + * output 6 (persistence-max). The branchy max vs sum pattern fits in + * one loop because we hard-code which lane uses which combiner. */ + for (int stride = BLOCK_SIZE / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + sdata[0 * BLOCK_SIZE + tid] += sdata[0 * BLOCK_SIZE + tid + stride]; + sdata[1 * BLOCK_SIZE + tid] += sdata[1 * BLOCK_SIZE + tid + stride]; + sdata[2 * BLOCK_SIZE + tid] += sdata[2 * BLOCK_SIZE + tid + stride]; + sdata[3 * BLOCK_SIZE + tid] += sdata[3 * BLOCK_SIZE + tid + stride]; + sdata[4 * BLOCK_SIZE + tid] += sdata[4 * BLOCK_SIZE + tid + stride]; + sdata[5 * BLOCK_SIZE + tid] += sdata[5 * BLOCK_SIZE + tid + stride]; + sdata[6 * BLOCK_SIZE + tid] = fmaxf( + sdata[6 * BLOCK_SIZE + tid], + sdata[6 * BLOCK_SIZE + tid + stride] + ); + } + __syncthreads(); + } + + /* Thread 0 writes the final reduced values to the ISV bus. Mean + * aggregates are sum / n_envs (n_envs ≥ 1 enforced by the launcher + * — single-env smoke configs still pass n_envs = 1). The CALMAR + * slot remains a denominator floor (mean of `max(dd_max_i, 1e-4)` + * stays ≥ 1e-4 because every term is ≥ 1e-4, so the mean cannot + * drop below the floor). */ + if (tid == 0) { + const float inv_n = 1.0f / (float)n_envs; + isv[/*DD_CURRENT_INDEX*/ 401] = sdata[0 * BLOCK_SIZE] * inv_n; + isv[/*DD_MAX_INDEX*/ 402] = sdata[1 * BLOCK_SIZE] * inv_n; + isv[/*DD_RECOVERY_BARS_INDEX*/ 403] = sdata[2 * BLOCK_SIZE] * inv_n; + isv[/*DD_PERSISTENCE_INDEX*/ 404] = sdata[3 * BLOCK_SIZE] * inv_n; + isv[/*CALMAR_INDEX*/ 405] = sdata[4 * BLOCK_SIZE] * inv_n; + isv[/*DD_PCT_INDEX*/ 406] = sdata[5 * BLOCK_SIZE] * inv_n; + isv[/*DD_PERSISTENCE_MAX_INDEX*/ 443] = sdata[6 * BLOCK_SIZE]; + __threadfence_system(); + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index fcc6a16d1..65bb75897 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -833,61 +833,76 @@ pub fn launch_sp15_cost_net_sharpe( Ok(()) } -/// SP15 Phase 1.3 (2026-05-06): per-step drawdown state tracking kernel. +/// SP15 Phase 1.3 (2026-05-06) + Phase 1.3.b-followup (2026-05-07): +/// per-env drawdown state tracking kernel. +/// /// 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 -/// calmar = mean_pnl / value-here), ISV[DD_PCT_INDEX=406]. +/// PS_PEAK_EQUITY (slot 7) and PS_PREV_EQUITY (slot 9) from each env's +/// row (`pos_state[env * PORTFOLIO_STRIDE + …]`); does NOT write back to +/// those slots — `experience_env_step` is the canonical equity writer +/// (see `experience_kernels.cu:3473-3475`). +/// +/// Phase 1.3.b-followup (Path A → Path B): the original Path A kernel +/// chose env-0 as the canonical DD observable and wrote 6 ISV scalar +/// slots directly. Path B redesigns to a per-env grid that writes a +/// per-env tile `dd_state_per_env[n_envs * 6]`; a separate +/// `dd_state_reduce_kernel` aggregates the tile into the 6 scalar ISV +/// slots [401..407) (mean across envs — HEALTH_DIAG diagnostic) plus +/// the new `DD_PERSISTENCE_MAX_INDEX = 443` slot (max across envs — +/// plasticity-injection trigger). +/// +/// Per-env tile layout (offsets within each env's [6] sub-array): +/// tile[env*6 + 0] = DD_CURRENT (max(0, (peak-equity)/peak)) +/// tile[env*6 + 1] = DD_MAX (running max in this fold) +/// tile[env*6 + 2] = DD_RECOVERY_BARS (bars since last new HWM) +/// tile[env*6 + 3] = DD_PERSISTENCE (twin counter to recovery) +/// tile[env*6 + 4] = CALMAR (max(dd_max, 1e-4) — floor) +/// tile[env*6 + 5] = DD_PCT (clip(curr/max(budget,1e-4), 0, 1)) +/// /// The 1e-4 calmar floor eliminates the saturation-at-100 artifact seen /// in the legacy host-side composer. pub static SP15_DD_STATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/dd_state_kernel.cubin")); -/// SP15 Phase 1.3 (2026-05-06): launcher for the per-step drawdown state -/// kernel. Free function (matches `launch_sp15_sharpe_per_bar` and -/// `launch_sp15_cost_net_sharpe` precedent) so unit/oracle tests can -/// drive the kernel directly without the trainer struct; production +/// SP15 Phase 1.3.b-followup (2026-05-07): launcher for the per-env +/// drawdown-state kernel. Per-env grid `[n_envs, 1, 1]` × `[1, 1, 1]`: +/// one thread per env (block) computes its own DD trajectory in a +/// state-machine fashion and writes 6 scalars to its slot in the +/// per-env tile. Free function (matches `launch_sp15_sharpe_per_bar` +/// and `launch_sp15_cost_net_sharpe` precedent) so unit/oracle tests +/// can drive the kernel directly without the trainer struct; production /// callers invoke the same launcher. /// -/// 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). +/// Per-step semantics (spec §6.3, post-Phase-1.3.b-followup): +/// For each env in `[0, n_envs)`: +/// 1. Read `prev_equity = pos_state[env * PORTFOLIO_STRIDE + PS_PREV_EQUITY]` +/// and `peak = pos_state[env * PORTFOLIO_STRIDE + PS_PEAK_EQUITY]`. +/// 2. current_dd = max(0, (peak − prev_equity) / peak); write to +/// tile[env*6 + 0]. Update tile[env*6 + 1] = max(prior, current_dd). +/// 3. If current_dd ≤ 0 (new HWM / pre-trade flat), zero +/// tile[env*6 + 2] / tile[env*6 + 3]; else increment both by 1. +/// 4. dd_pct = clip(current_dd / max(dd_budget, 1e-4), 0, 1) → +/// tile[env*6 + 5]. tile[env*6 + 4] = max(dd_max, 1e-4) (floor). /// /// 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). +/// MUST launch this kernel AFTER the env_step launch on the same +/// stream, AND must launch `dd_state_reduce_kernel` after this kernel +/// to populate the ISV scalar slots from the tile (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). +/// `n_envs` MUST be ≥ 1 (the reduction kernel divides by `n_envs`). +/// `dd_state_per_env` MUST point to at least `n_envs * 6` writable +/// f32 slots (the trainer's `sp15_dd_state_per_env` mapped-pinned +/// scratch buffer in production; a `MappedF32Buffer::new(n_envs * 6)` +/// in oracle tests). `pos_state` MUST point to a portfolio-state +/// buffer with at least `n_envs * PORTFOLIO_STRIDE` f32 slots. pub fn launch_sp15_dd_state( stream: &Arc, dd_budget: f32, - isv: cudarc::driver::sys::CUdeviceptr, + n_envs: i32, + dd_state_per_env: cudarc::driver::sys::CUdeviceptr, pos_state: cudarc::driver::sys::CUdeviceptr, ) -> Result<(), MLError> { let module = stream @@ -901,14 +916,16 @@ pub fn launch_sp15_dd_state( .map_err(|e| MLError::ModelError(format!( "load dd_state_kernel function: {e}" )))?; + let grid_x = n_envs.max(1) as u32; unsafe { stream .launch_builder(&kernel) .arg(&dd_budget) - .arg(&isv) + .arg(&n_envs) + .arg(&dd_state_per_env) .arg(&pos_state) .launch(LaunchConfig { - grid_dim: (1, 1, 1), + grid_dim: (grid_x, 1, 1), block_dim: (1, 1, 1), shared_mem_bytes: 0, }) @@ -919,6 +936,82 @@ pub fn launch_sp15_dd_state( Ok(()) } +/// SP15 Phase 1.3.b-followup (2026-05-07): cubin slot for +/// `dd_state_reduce_kernel`. +/// +/// Single-block tree-reduce that aggregates the per-env tile written by +/// `dd_state_kernel` into the 6 scalar ISV slots [401..407) (mean +/// across envs — smooth across-env summary for HEALTH_DIAG diagnostic +/// reporting) plus the new `DD_PERSISTENCE_MAX_INDEX = 443` slot (max +/// across envs — consumed by `plasticity_injection_kernel` so the +/// global advantage-head reset fires when ANY env's persistence +/// exceeds the threshold). BLOCK=256 with the strided initial +/// accumulation pattern handles n_envs up to production sizes (32768 +/// envs on H100). +/// +/// Aggregation rule rationale (per `feedback_isv_for_adaptive_bounds` — +/// the rule itself is structural, not adaptive): +/// +/// * Mean for slots 401..407): preserves backward-compatible +/// "this is one DD scalar to look at" semantics for HEALTH_DIAG; +/// only the across-env aggregation rule changed (Path A: +/// env-0-canonical → Path B: mean across envs). +/// * Max for slot 443 (DD_PERSISTENCE_MAX): one set of advantage +/// weights → global firing semantics → ANY env exceeding the +/// threshold should arm the gate. Mean-aggregate would silently +/// gate the trigger when only a fraction of envs are in long DD. +/// +/// Single-source-to-cubin convention (mirrors the SP15 1.3 / 1.4 / 3.X +/// pattern); single launcher (`launch_sp15_dd_state_reduce`). +pub static SP15_DD_STATE_REDUCE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/dd_state_reduce_kernel.cubin")); + +/// SP15 Phase 1.3.b-followup (2026-05-07): launcher for the DD state +/// reduction kernel. Single-block (`grid = (1, 1, 1)`, `block = +/// (256, 1, 1)`) tree-reduce — one launch per per-step `dd_state_kernel` +/// launch, on the same stream so CUDA serialises producer→consumer +/// without explicit event sync. `n_envs` MUST equal the value passed +/// to `launch_sp15_dd_state` for the same step (the per-env tile must +/// have been written by that earlier launch). +/// +/// `dd_state_per_env` MUST point to at least `n_envs * 6` f32 slots +/// (read-only — the kernel only writes `isv`). `isv` MUST be the ISV +/// bus (≥SP15_SLOT_END=444 f32 slots — slots 401-406 + 443 written). +pub fn launch_sp15_dd_state_reduce( + stream: &Arc, + n_envs: i32, + dd_state_per_env: cudarc::driver::sys::CUdeviceptr, + isv: cudarc::driver::sys::CUdeviceptr, +) -> Result<(), MLError> { + let module = stream + .context() + .load_cubin(SP15_DD_STATE_REDUCE_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!( + "load sp15_dd_state_reduce cubin: {e}" + )))?; + let kernel = module + .load_function("dd_state_reduce_kernel") + .map_err(|e| MLError::ModelError(format!( + "load dd_state_reduce_kernel function: {e}" + )))?; + unsafe { + stream + .launch_builder(&kernel) + .arg(&n_envs) + .arg(&dd_state_per_env) + .arg(&isv) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "launch dd_state_reduce_kernel: {e}" + )))?; + } + Ok(()) +} + /// SP15 Wave 4.1a (2026-05-06): launcher for the bottleneck-aware /// dd_pct concat extension `bn_tanh_concat_dd_kernel`. Resolves the /// kernel symbol from the existing `dqn_utility_kernels.cubin` (the @@ -1455,12 +1548,25 @@ pub static SP15_FINAL_REWARD_CUBIN: &[u8] = /// /// Grid: `ceil(n*2*l / 256)` blocks of 256 threads each. Each thread /// processes one slot; no inter-thread communication. +/// +/// SP15 Phase 1.3.b-followup (2026-05-07) — `dd_state_per_env` MUST +/// point to the per-env DD tile `[n * 6]` written by `dd_state_kernel` +/// earlier this step (the trainer's `sp15_dd_state_per_env` mapped- +/// pinned scratch buffer in production; a `MappedF32Buffer::new(n*6)` +/// in oracle tests). Each thread looks up its env's DD context via +/// `env_id = (idx % (N*L)) / L` so on-policy and CF slots at the same +/// (i,t) read the SAME per-env tile entry (one DD trajectory per env, +/// shared across slot kinds). The Path A `isv[DD_PCT_INDEX]` / +/// `isv[DD_CURRENT_INDEX]` reads in `sp15_dd_asymmetric_reward` / +/// `sp15_dd_penalty` migrated to this per-env lookup atomically per +/// `feedback_no_partial_refactor`. pub fn launch_sp15_final_reward( stream: &Arc, out_rewards: cudarc::driver::sys::CUdeviceptr, r_discipline_out: cudarc::driver::sys::CUdeviceptr, slot_completed_normally: cudarc::driver::sys::CUdeviceptr, isv: cudarc::driver::sys::CUdeviceptr, + dd_state_per_env: cudarc::driver::sys::CUdeviceptr, n: i32, l: i32, ) -> Result<(), MLError> { @@ -1484,6 +1590,7 @@ pub fn launch_sp15_final_reward( .arg(&r_discipline_out) .arg(&slot_completed_normally) .arg(&isv) + .arg(&dd_state_per_env) .arg(&n) .arg(&l) .launch(LaunchConfig { @@ -2311,7 +2418,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 443; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 EGF pearl at [383..396) — Q_DISAGREEMENT_SHORT/LONG_EMA [383..385) + K_AUX/Q_ADAPTIVE [385..387) + BETA_RATE_LIMITER_ADAPTIVE [387] + AUX_DIR_ACC/Q_DISAGREEMENT/ALPHA_GRAD_RAW_VARIANCE_EMA [388..391) + GATE1_OPEN_STATE [391] + ALPHA_GRAD_RAW [392] + ALPHA_GRAD_SMOOTHED [393] + AUX_DIR_ACC_POST_OPEN_MIN [394] + GRADIENT_HACK_LOCKOUT_REMAINING [395]; intentional 5-slot boundary gap at [367..372)) +pub(crate) const ISV_TOTAL_DIM: usize = 444; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11 + SP13 + SP13 v3 + SP14 + SP15 (incl. 1.3.b-followup DD_PERSISTENCE_MAX at 443). Bumped 443 → 444 by SP15 Phase 1.3.b-followup (per-env DD redesign): the new `DD_PERSISTENCE_MAX_INDEX = 443` slot is the max-aggregate of per-env DD_PERSISTENCE consumed by `plasticity_injection_kernel` (one set of advantage weights → global-firing semantics → max-aggregate is the only correct rule when ANY env exceeding the threshold should arm the gate). The mean-aggregate of DD_PERSISTENCE remains at the existing slot 404 for HEALTH_DIAG diagnostic reporting (smooth across-env summary). Pre-1.3.b-followup base: 173 + 210 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]; SP13 directional-skill instrumentation at [372..380) — TARGET_DIR_ACC [372] + AUX_DIR_ACC_SHORT/LONG_EMA [373..375) + AUX_DIR_PREDICTION [375] + DIR_SKILL_BONUS_ALPHA/BETA [376..378) + LUCK_WIN_DISCOUNT [378] + SKILL_BONUS_CAP_RATIO [379]; SP13 v3 P0a.T3 Hold-pricing controller at [380..383) — HOLD_COST [380] + HOLD_RATE_TARGET [381] + HOLD_RATE_OBSERVED_EMA [382]; SP14 EGF pearl at [383..396) — Q_DISAGREEMENT_SHORT/LONG_EMA [383..385) + K_AUX/Q_ADAPTIVE [385..387) + BETA_RATE_LIMITER_ADAPTIVE [387] + AUX_DIR_ACC/Q_DISAGREEMENT/ALPHA_GRAD_RAW_VARIANCE_EMA [388..391) + GATE1_OPEN_STATE [391] + ALPHA_GRAD_RAW [392] + ALPHA_GRAD_SMOOTHED [393] + AUX_DIR_ACC_POST_OPEN_MIN [394] + GRADIENT_HACK_LOCKOUT_REMAINING [395]; intentional 5-slot boundary gap at [367..372)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -3384,8 +3491,10 @@ const fn layout_fingerprint_seed() -> &'static [u8] { PLASTICITY_FIRED_THIS_FOLD=436;PLASTICITY_PERSISTENCE_THRESHOLD=437;PLASTICITY_WARM_BARS_REMAINING=438;\ DD_TRAJECTORY_DECREASING=439;RECOVERY_OVERSAMPLE_WEIGHT=440;\ DD_TRAJECTORY_FLOOR=441;MEDIAN_STREAK_LENGTH=442;\ - ISV_TOTAL_DIM=443;\ + DD_PERSISTENCE_MAX=443;\ + ISV_TOTAL_DIM=444;\ TRUNK_INPUT_DD_PCT=sp15_wave_4_1a;\ + DD_STATE_PER_ENV=sp15_phase_1_3_b_followup;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ @@ -4779,6 +4888,15 @@ pub struct GpuDqnTrainer { /// page via `read_all()` only in tests. pub(crate) sp15_dd_trajectory_prev_dd: super::mapped_pinned::MappedF32Buffer, // [1] + // SP15 Phase 1.3.b-followup (2026-05-07): per-env DD state tile + // `sp15_dd_state_per_env` is owned by `GpuExperienceCollector` + // (which knows `alloc_episodes` = `n_envs`); the trainer does NOT + // own a copy. The collector launches `dd_state_kernel` / + // `dd_state_reduce_kernel` / `compute_sp15_final_reward_kernel` + // back-to-back per step on the same stream, threading its own + // tile pointer into all three. See + // `gpu_experience_collector.rs::sp15_dd_state_per_env`. + // ── Per-branch Q-gap EMA (trajectory backtracking state) ── /// [4] per-branch Q-gap EMA — device buffer (read-modify-write for backtracking restore). per_branch_q_gap_ema_buf: CudaSlice, diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 87631306b..0f79a3400 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -983,6 +983,38 @@ pub struct GpuExperienceCollector { /// pattern. sp15_dd_trajectory_prev_dd_dev_ptr: u64, + /// SP15 Phase 1.3.b-followup (2026-05-07): per-env DD state tile + /// for the redesigned `dd_state_kernel` + `dd_state_reduce_kernel` + /// pair (Path A → Path B migration). Mapped-pinned f32 buffer + /// sized `[alloc_episodes * 6]` (6 stats per env: DD_CURRENT, + /// DD_MAX, DD_RECOVERY_BARS, DD_PERSISTENCE, CALMAR floor, DD_PCT). + /// + /// Per-step launch order on the same stream (CUDA serialises + /// producer→consumer without explicit event sync): + /// 1. `dd_state_kernel` runs a per-env grid `[n_envs, 1, 1]`, + /// reading PS_PEAK_EQUITY / PS_PREV_EQUITY from each env's + /// portfolio-state row and writing 6 scalars to its slot in + /// this tile (read-modify-write — running max + counters). + /// 2. `dd_state_reduce_kernel` aggregates the tile into 6 ISV + /// scalar slots [401..407) (mean across envs — HEALTH_DIAG + /// diagnostic) plus DD_PERSISTENCE_MAX_INDEX = 443 (max + /// across envs — plasticity-injection trigger). + /// 3. `compute_sp15_final_reward_kernel` reads the tile per-(i,t) + /// via `env_id = (idx % (N*L)) / L` so each slot's reward + /// shaping uses ITS env's DD context, NOT env-0's canonical + /// observable (the Path A limitation). + /// + /// FoldReset sentinel 0 across all `alloc_episodes * 6` slots — + /// every stateful field (running max, counters) bootstraps from + /// zero on the new fold's first kernel launch (same rationale as + /// the individual ISV scalars in the original Phase 1.3 reset + /// arms). Reset path: `host_slice_mut().fill(0.0)` via the + /// `sp15_dd_state_per_env` registry-arm dispatch in + /// `training_loop.rs` (allowed under + /// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write + /// rule, NOT a host-side compute). + pub(crate) sp15_dd_state_per_env: super::mapped_pinned::MappedF32Buffer, + /// SP15 Phase 3.5.4.c (2026-05-07): trainer's `params_buf` /// device pointer offset to `w_b0out` (directional advantage-head /// last-Linear weights, tensor index 19 per `compute_param_sizes`). @@ -1714,6 +1746,28 @@ impl GpuExperienceCollector { format!("SP13 v3 sp13_hold_rate_buf alloc (1 f32, collector): {e}") ))?; + // SP15 Phase 1.3.b-followup (2026-05-07): per-env DD state tile + // sized `[alloc_episodes * 6]`. The redesigned `dd_state_kernel` + // writes 6 stats per env to this tile (Path A → Path B + // migration); `dd_state_reduce_kernel` mean-aggregates into the + // 6 ISV scalar slots [401..407) and max-aggregates into the + // new DD_PERSISTENCE_MAX_INDEX = 443 slot; + // `compute_sp15_final_reward_kernel` reads it per-(i,t) so each + // slot's reward shaping uses ITS env's DD context. + // Zero-initialised by `MappedF32Buffer::new`'s contract; reset + // to zero at fold boundary via `host_slice_mut().fill(0.0)` + // (mirrors the `sp13_hold_rate_buf` / `sp15_alpha_warm_count` + // non-ISV mapped-pinned reset pattern). + let sp15_dd_state_per_env = + unsafe { MappedF32Buffer::new(alloc_episodes * 6) } + .map_err(|e| MLError::ModelError( + format!( + "SP15 Phase 1.3.b-followup sp15_dd_state_per_env alloc \ + ({} f32, collector): {e}", + alloc_episodes * 6 + ) + ))?; + // B.2 Plan 3 Task 3: load trade_attempt_rate_ema_update kernel. let trade_attempt_rate_ema_kernel = { use super::gpu_dqn_trainer::TRADE_RATE_EMA_CUBIN; @@ -2041,6 +2095,7 @@ impl GpuExperienceCollector { isv_signals_dev_ptr: 0, // NULL until trainer sets it sp15_alpha_warm_count_dev_ptr: 0, // NULL until trainer wires it (SP15 Wave 2) sp15_dd_trajectory_prev_dd_dev_ptr: 0, // NULL until trainer wires it (SP15 Wave 4.3 / 3.5.5.b) + sp15_dd_state_per_env, // SP15 Phase 1.3.b-followup per-env DD tile [alloc_episodes * 6] sp15_w_b0out_dev_ptr: 0, // NULL until trainer wires it (SP15 Phase 3.5.4.c) sp15_w_b0out_n_weights: 0, sp15_w_b0out_fan_in: 0, @@ -4703,6 +4758,60 @@ impl GpuExperienceCollector { // experience_env_step without ISV wiring, in which case // both buffers receive their default values and the // fused composer would be a no-op anyway. + // ── 5b. SP15 Phase 1.3.b-followup (2026-05-07): per-env DD ── + // Order MUST be: dd_state (per-env tile producer) → + // dd_state_reduce (ISV scalar consumer) → alpha producer → + // final_reward (per-(i,t) consumer of the per-env tile). + // + // Path A → Path B migration: the original Phase 1.3.b + // per-step launch wrote 6 ISV scalars from env-0 directly; + // Path B writes a per-env tile then reduces. Each env in + // production has its OWN portfolio + DD trajectory: when + // env-3 is in 30% DD and env-0 is at ATH, the per-env + // reward shaping in `compute_sp15_final_reward_kernel` + // looks up env-3's `dd_pct=0.3` (not env-0's `dd_pct=0`) + // and applies the recovery shaping correctly. + // + // Ordering: launched on the same stream as env_step → CUDA + // serializes the four on-stream, so each consumer sees the + // producers' 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. + // + // 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 dd_state_per_env_ptr = self.sp15_dd_state_per_env.dev_ptr; + { + let pos_state_ptr = self.portfolio_states.raw_ptr(); + let n_envs_i32 = self.alloc_episodes as i32; + crate::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_state( + &self.stream, + config.dd_threshold, + n_envs_i32, + dd_state_per_env_ptr, + pos_state_ptr, + )?; + // Reduce the per-env tile into the 6 scalar ISV slots + // [401..407) (mean across envs — HEALTH_DIAG diagnostic) + // plus DD_PERSISTENCE_MAX_INDEX=443 (max across envs — + // plasticity-injection trigger). Gated on + // `isv_signals_dev_ptr` being wired (test scaffolds may + // run env_step without ISV; the per-env tile still gets + // populated for the per-(i,t) lookup in the final-reward + // kernel below, which is itself gated on the same ptr). + if self.isv_signals_dev_ptr != 0 { + crate::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_state_reduce( + &self.stream, + n_envs_i32, + dd_state_per_env_ptr, + self.isv_signals_dev_ptr, + )?; + } + } + if self.isv_signals_dev_ptr != 0 && self.sp15_alpha_warm_count_dev_ptr != 0 { use crate::cuda_pipeline::gpu_dqn_trainer::{ launch_sp15_alpha_split_producer, launch_sp15_final_reward, @@ -4725,6 +4834,7 @@ impl GpuExperienceCollector { r_disc_dev, flag_dev, self.isv_signals_dev_ptr, + dd_state_per_env_ptr, n_i32, l_i32, ).map_err(|e| MLError::ModelError(format!( @@ -4732,39 +4842,6 @@ 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, - )?; - } - // ── 5b. SP15 Phase 3.5.5.b — DD_TRAJECTORY_DECREASING per-step proxy ── // // Reads ISV[DD_PCT_INDEX=406] (just written by diff --git a/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu index ff96c93bc..426b2847b 100644 --- a/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu +++ b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu @@ -47,11 +47,23 @@ // so every thread computes the same `fire_now` without any // cross-block synchronisation). // +// SP15 Phase 1.3.b-followup contract change (2026-05-07): the +// persistence read migrated from ISV[404] (mean-aggregate of per-env +// DD_PERSISTENCE) to ISV[443] (DD_PERSISTENCE_MAX — max-aggregate +// across envs). The plasticity gate fires when ANY env's persistence +// exceeds the threshold, NOT when the across-env mean exceeds it: one +// set of advantage weights → global firing semantics → max-aggregate +// is the only correct rule. Mean-aggregate would silently gate the +// trigger when only a fraction of envs are in long DD (e.g. 1 env at +// persistence=200 + 7 envs at persistence=0 ⇒ mean=25 ⇒ gate stays +// closed even though one env is deep in the regime that motivates +// plasticity injection). +// // Slot reads/writes (literal indices to keep this kernel self-contained; // the canonical names live in `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` // and the constructor + state_reset_registry pin those names to these // indices via compile-time `assert_eq!` regression checks): -// ISV[404] DD_PERSISTENCE (read) +// ISV[443] DD_PERSISTENCE_MAX (read — Phase 1.3.b-followup) // ISV[437] PLASTICITY_PERSISTENCE_THRESHOLD (read) // ISV[436] PLASTICITY_FIRED_THIS_FOLD (read-modify-write — block 0 thread 0 only) // ISV[438] PLASTICITY_WARM_BARS_REMAINING (read-modify-write — block 0 thread 0 only) @@ -104,7 +116,13 @@ extern "C" __global__ void plasticity_injection_kernel( // reads, so all threads converge on the same answer without any // cross-block synchronisation. The same reads also drive block 0 // thread 0's ISV-write path below. - const float persistence = isv[/*DD_PERSISTENCE_INDEX*/ 404]; + // + // Phase 1.3.b-followup (2026-05-07): the persistence read uses + // ISV[DD_PERSISTENCE_MAX_INDEX=443] (max across envs, written by + // `dd_state_reduce_kernel`) so the gate fires when ANY env's + // persistence exceeds the threshold — one shared advantage-head + // means a global firing condition. + const float persistence = isv[/*DD_PERSISTENCE_MAX_INDEX*/ 443]; const float threshold = isv[/*PLASTICITY_PERSISTENCE_THRESHOLD_INDEX*/ 437]; const float fired = isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436]; diff --git a/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs index c0373606b..fd53d4286 100644 --- a/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp15_isv_slots.rs @@ -92,8 +92,20 @@ pub const RECOVERY_OVERSAMPLE_WEIGHT_INDEX: usize = 440; pub const DD_TRAJECTORY_FLOOR_INDEX: usize = 441; pub const MEDIAN_STREAK_LENGTH_INDEX: usize = 442; +// === Phase 1.3.b-followup (1 slot) === +// Per-env DD redesign: max-aggregate of per-env DD_PERSISTENCE across all +// envs, written by `dd_state_reduce_kernel` from the per-env tile. Read by +// `plasticity_injection_kernel` so the global advantage-head reset fires +// when ANY env's persistence exceeds the threshold (one set of advantage +// weights = global firing semantics, the worst-env governs). The mean- +// aggregate of DD_PERSISTENCE remains at the existing slot 404 for +// HEALTH_DIAG diagnostic reporting (smooth across-env summary). See +// `docs/dqn-wire-up-audit.md` Phase 1.3.b-followup entry for the +// aggregation-rule rationale and the Path A → Path B migration detail. +pub const DD_PERSISTENCE_MAX_INDEX: usize = 443; + pub const SP15_SLOT_BASE: usize = 397; -pub const SP15_SLOT_END: usize = 443; +pub const SP15_SLOT_END: usize = 444; pub const SP15_SLOT_COUNT: usize = SP15_SLOT_END - SP15_SLOT_BASE; #[cfg(test)] @@ -105,11 +117,11 @@ mod tests { /// Paired contract with the ISV_TOTAL_DIM bump in gpu_dqn_trainer.rs. #[test] fn all_sp15_slots_fit_within_isv_total_dim() { - assert!(MEDIAN_STREAK_LENGTH_INDEX < ISV_TOTAL_DIM, - "MEDIAN_STREAK_LENGTH_INDEX={} >= ISV_TOTAL_DIM={}; bump ISV_TOTAL_DIM", - MEDIAN_STREAK_LENGTH_INDEX, ISV_TOTAL_DIM); - assert_eq!(SP15_SLOT_END, MEDIAN_STREAK_LENGTH_INDEX + 1); - assert_eq!(SP15_SLOT_COUNT, 46); + assert!(DD_PERSISTENCE_MAX_INDEX < ISV_TOTAL_DIM, + "DD_PERSISTENCE_MAX_INDEX={} >= ISV_TOTAL_DIM={}; bump ISV_TOTAL_DIM", + DD_PERSISTENCE_MAX_INDEX, ISV_TOTAL_DIM); + assert_eq!(SP15_SLOT_END, DD_PERSISTENCE_MAX_INDEX + 1); + assert_eq!(SP15_SLOT_COUNT, 47); } /// Layout fingerprint regression: each named slot at its allocated index. @@ -142,5 +154,6 @@ mod tests { assert_eq!(DD_TRAJECTORY_DECREASING_INDEX, 439); assert_eq!(DD_TRAJECTORY_FLOOR_INDEX, 441); assert_eq!(MEDIAN_STREAK_LENGTH_INDEX, 442); + assert_eq!(DD_PERSISTENCE_MAX_INDEX, 443); } } diff --git a/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh b/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh index 3798b1df2..c759bfd6a 100644 --- a/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh +++ b/crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh @@ -113,6 +113,15 @@ __device__ inline float sp15_alpha_blend( // through the SP12-cap pipeline per // `pearl_audit_unboundedness_for_implicit_asymmetry`. // +// SP15 Phase 1.3.b-followup contract (2026-05-07): `dd_pct` is now a +// per-env scalar passed by the caller — read from +// `dd_state_per_env[env*6 + 5]` where `env_id = (idx % (N*L)) / L` +// (the `compute_sp15_final_reward_kernel` launches per-(i,t) over both +// on-policy and CF slots, and each slot threads its own env's DD +// context through here). `lambda` (SP15_DD_ASYMMETRY_LAMBDA_INDEX=430) +// is still a global controller scalar — one λ per training instance, +// read from ISV by the kernel (NOT here) and threaded in. +// // `r_gain_boost_diag_out` records the most-recent boost factor applied // (1.0 = no-op for losses or ATH; >1.0 = real gain-side boost). For the // fused kernel ONLY the on-policy slot's boost factor needs to land in @@ -122,12 +131,10 @@ __device__ inline float sp15_alpha_blend( // -------------------------------------------------------------------- __device__ inline float sp15_dd_asymmetric_reward( float r_in, - const float* __restrict__ isv, + float lambda, + float dd_pct, float* __restrict__ r_gain_boost_diag_out /* nullable */ ) { - const float lambda = isv[SP15_DD_ASYMMETRY_LAMBDA_INDEX]; - const float dd_pct = isv[SP15_DD_PCT_INDEX]; - float boost = 1.0f; float r_adjusted; if (r_in > 0.0f) { @@ -151,16 +158,21 @@ __device__ inline float sp15_dd_asymmetric_reward( // Asymmetric: zero below threshold, quadratic growth above. Encodes // loss aversion per `pearl_audit_unboundedness_for_implicit_asymmetry`. // Subtracted from `r_in`. +// +// SP15 Phase 1.3.b-followup contract (2026-05-07): `dd_current` is now +// a per-env scalar passed by the caller — read from +// `dd_state_per_env[env*6 + 0]`. `lambda` (SP15_LAMBDA_DD_INDEX=420) +// and `thr` (SP15_DD_THRESHOLD_INDEX=421) are still global controller +// scalars — one λ_dd / threshold per training instance, read from ISV +// by the kernel (NOT here) and threaded in. // -------------------------------------------------------------------- __device__ inline float sp15_dd_penalty( float r_in, - const float* __restrict__ isv + float lambda, + float thr, + float dd_current ) { - const float dd = isv[SP15_DD_CURRENT_INDEX]; - const float thr = isv[SP15_DD_THRESHOLD_INDEX]; - const float lambda = isv[SP15_LAMBDA_DD_INDEX]; - - const float excess = fmaxf(0.0f, dd - thr); + const float excess = fmaxf(0.0f, dd_current - thr); const float penalty = lambda * excess * excess; return r_in - penalty; } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 0aa8684e4..6ffe1ef3c 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1086,7 +1086,48 @@ impl StateResetRegistry { RegistryEntry { name: "sp15_dd_pct", category: ResetCategory::FoldReset, - description: "ISV[DD_PCT_INDEX=406] — SP15 Phase 1.3 (2026-05-06) `clip(current_dd / max(dd_budget, 1e-4), 0, 1)` ∈ [0, 1]. Stateful kernel output; FoldReset sentinel 0 (cold-start at fold boundary; the kernel overwrites with the true value on the new fold's first launch). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3.", + description: "ISV[DD_PCT_INDEX=406] — SP15 Phase 1.3 (2026-05-06) `clip(current_dd / max(dd_budget, 1e-4), 0, 1)` ∈ [0, 1]. Phase 1.3.b-followup (2026-05-07): now mean-aggregated across envs by `dd_state_reduce_kernel` from the per-env tile (the per-(env, slot) DD context is read directly by `compute_sp15_final_reward_kernel` from `sp15_dd_state_per_env[env*6 + 5]`; this scalar is the HEALTH_DIAG diagnostic summary). Stateful kernel output; FoldReset sentinel 0 (cold-start at fold boundary; the reduction kernel overwrites with the true mean on the new fold's first launch). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3.", + }, + // ── SP15 Phase 1.3.b-followup (2026-05-07): per-env DD redesign ── + // - sp15_dd_state_per_env (NOT an ISV slot — `[alloc_episodes + // × 6]` mapped-pinned tile on the experience collector): + // FoldReset sentinel 0 across all `n_envs * 6` slots — every + // stateful field (running max + recovery counter + + // persistence counter per env) bootstraps from zero on the + // new fold's first kernel launch. Reset path: + // `host_slice_mut().fill(0.0)` (allowed under + // `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed- + // write rule, NOT a host-side compute). Mirrors the + // `sp15_alpha_warm_count` / `sp15_dd_trajectory_prev_dd` + // non-ISV mapped-pinned reset pattern. Without this reset + // the new fold's first reward shaping would inherit the + // previous fold's running max + counters, which would + // silently bias the per-env DD context. + // - DD_PERSISTENCE_MAX_INDEX=443: max-aggregate of per-env + // DD_PERSISTENCE across all envs, written by + // `dd_state_reduce_kernel`. Read by + // `plasticity_injection_kernel` so the global advantage- + // head reset fires when ANY env's persistence exceeds the + // threshold (one set of advantage weights → global firing + // semantics → max-aggregate is the only correct rule when + // ANY env exceeding the threshold should arm the gate; + // mean-aggregate would silently gate the trigger when only + // a fraction of envs are in long DD). FoldReset sentinel 0 + // same rationale as DD_PERSISTENCE_INDEX (a fresh fold + // starts with the per-env counters at 0; the reduction + // max is therefore 0 until the first per-env DD-step + // fires). + // Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3 + // post-Phase-1.3.b-followup audit doc entry. + RegistryEntry { + name: "sp15_dd_state_per_env", + category: ResetCategory::FoldReset, + description: "GpuExperienceCollector.sp15_dd_state_per_env [alloc_episodes * 6] f32 mapped-pinned per-env DD state tile — SP15 Phase 1.3.b-followup (2026-05-07) per-env redesign of the DD state machine (Path A → Path B migration). Layout: tile[env*6 + 0..6] = [DD_CURRENT, DD_MAX, DD_RECOVERY_BARS, DD_PERSISTENCE, CALMAR floor, DD_PCT] for env in [0, alloc_episodes). Read-modify-written by `dd_state_kernel` (per-env grid, one thread per env, stateful per-bar walk). Read by `dd_state_reduce_kernel` (mean across envs → ISV slots [401..407); max(persistence) → ISV[DD_PERSISTENCE_MAX_INDEX=443]). Read by `compute_sp15_final_reward_kernel` per-(i,t) via `env_id = (idx % (N*L)) / L`. FoldReset sentinel 0 across all `n_envs * 6` slots — every stateful field (running max, recovery + persistence counters) bootstraps from zero on the new fold's first kernel launch (same rationale as the individual ISV scalars in the original Phase 1.3 reset arms). Reset path: `host_slice_mut().fill(0.0)` mirrors the `sp15_alpha_warm_count` / `sp15_dd_trajectory_prev_dd` non-ISV mapped-pinned reset pattern. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3 + docs/dqn-wire-up-audit.md Phase 1.3.b-followup entry.", + }, + RegistryEntry { + name: "sp15_dd_persistence_max", + category: ResetCategory::FoldReset, + description: "ISV[DD_PERSISTENCE_MAX_INDEX=443] — SP15 Phase 1.3.b-followup (2026-05-07) max-aggregate of per-env DD_PERSISTENCE across all envs, written by `dd_state_reduce_kernel` from the per-env tile. Read by `plasticity_injection_kernel` so the global advantage-head reset fires when ANY env's persistence exceeds the threshold — one set of advantage weights → global firing semantics → max-aggregate is the only correct rule (mean-aggregate would silently gate the trigger when only a fraction of envs are in long DD). Stateful kernel output; FoldReset sentinel 0 (a fresh fold starts with all per-env persistence counters at 0; the max-aggregate is therefore 0 until the first per-env DD-step fires the counter). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3 + docs/dqn-wire-up-audit.md Phase 1.3.b-followup entry.", }, // ── SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split ─── // Three ISV slots [417..420) + one mapped-pinned scratch buffer diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 9d8e702f9..5d96edaca 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -8216,6 +8216,30 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(DD_PCT_INDEX, 0.0); } } + // SP15 Phase 1.3.b-followup (2026-05-07): per-env DD state tile. + // Non-ISV mapped-pinned tile on the experience collector + // sized `[alloc_episodes * 6]`. Reset to zero via + // `host_slice_mut().fill(0.0)` (mirrors the + // `sp15_alpha_warm_count` non-ISV mapped-pinned reset + // pattern). The collector owns the buffer because it knows + // `alloc_episodes` (the trainer's `batch_size` is the + // training batch size, which is a different quantity). + "sp15_dd_state_per_env" => { + if let Some(ref mut collector) = self.gpu_experience_collector { + collector.sp15_dd_state_per_env.host_slice_mut().fill(0.0); + } + } + // SP15 Phase 1.3.b-followup (2026-05-07): max-aggregate of + // per-env DD_PERSISTENCE for the plasticity-injection + // trigger. Sentinel 0 — fresh fold starts with all per-env + // persistence counters at 0; the max-aggregate is therefore + // 0 until the first per-env DD-step fires a counter. + "sp15_dd_persistence_max" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::DD_PERSISTENCE_MAX_INDEX; + fused.trainer().write_isv_signal_at(DD_PERSISTENCE_MAX_INDEX, 0.0); + } + } // SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split // dispatch arms (per spec §8.2 (3.1) post-amendment-2 fix). // ALPHA_SPLIT_INDEX=417 is an Invariant-1 anchor (NOT a stateful diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index b6051f7a8..d712e2f6a 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -13,12 +13,13 @@ mod gpu { use ml::cuda_pipeline::gpu_dqn_trainer::{ launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only, launch_sp15_baseline_naive_momentum, launch_sp15_baseline_naive_reversion, - launch_sp15_cost_net_sharpe, launch_sp15_dd_state, + launch_sp15_cost_net_sharpe, launch_sp15_dd_state, launch_sp15_dd_state_reduce, launch_sp15_position_history_derivation, launch_sp15_sharpe_per_bar, }; use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer}; use ml::cuda_pipeline::sp15_isv_slots::{ - ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX, + ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, + DD_PCT_INDEX, DD_PERSISTENCE_INDEX, DD_PERSISTENCE_MAX_INDEX, DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END, }; @@ -247,9 +248,9 @@ mod 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). + // `pos_state[env * 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; @@ -262,20 +263,32 @@ mod gpu { 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). + // SP15 Phase 1.3.b-followup contract: per-env tile + reduction. + // Single-env config (n_envs=1) reproduces the legacy Path A + // semantics — the mean-aggregate over a single env equals the + // single env's value, so the existing assertions still hold. + const N_ENVS: usize = 1; + + // ISV bus sized to SP15_SLOT_END (slots 401-406 + 443 written + // by the reduction kernel). let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } .expect("alloc MappedF32Buffer for isv bus"); isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); - // 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); + // Per-env DD tile [n_envs * 6] (Phase 1.3.b-followup). + let dd_tile_buf = unsafe { MappedF32Buffer::new(N_ENVS * 6) } + .expect("alloc MappedF32Buffer for dd_state_per_env tile"); + dd_tile_buf.write_from_slice(&vec![0.0f32; N_ENVS * 6]); + + // Position state buffer sized for `N_ENVS` envs at PORTFOLIO_STRIDE + // plus headroom. The kernel reads `pos_state[env * STRIDE + ...]`. + let pos_len = (N_ENVS * PORTFOLIO_STRIDE).max(64); let pos_state_buf = unsafe { MappedF32Buffer::new(pos_len) } .expect("alloc MappedF32Buffer for pos_state"); // 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. + // does in production); after each write the dd_state + reduce + // kernels run back-to-back. let mut peak: f32 = 0.0; for &equity in &equity_per_step { peak = peak.max(equity); @@ -287,15 +300,26 @@ mod gpu { launch_sp15_dd_state( &stream, dd_budget, - isv_buf.dev_ptr, + N_ENVS as i32, + dd_tile_buf.dev_ptr, pos_state_buf.dev_ptr, ) .expect("launch dd_state_kernel"); + launch_sp15_dd_state_reduce( + &stream, + N_ENVS as i32, + dd_tile_buf.dev_ptr, + isv_buf.dev_ptr, + ) + .expect("launch dd_state_reduce_kernel"); stream .synchronize() - .expect("synchronize after dd_state_kernel launch"); + .expect("synchronize after dd_state + reduce kernel launches"); } + // Mean-aggregated ISV scalars across envs (single-env so mean == + // the env's value; the existing assertions on the original + // env-0-canonical contract still hold). let isv = isv_buf.read_all(); let max_dd = isv[DD_MAX_INDEX]; let current_dd = isv[DD_CURRENT_INDEX]; @@ -326,6 +350,187 @@ mod gpu { "dd_pct = {}, expected in (0, 1]", dd_pct ); + + // Per-env tile contents must match the ISV scalars (single-env + // mean degenerates to the env's value). Verifies the producer- + // consumer contract between `dd_state_kernel` and + // `dd_state_reduce_kernel`. + let tile = dd_tile_buf.read_all(); + assert!( + (tile[0] - current_dd).abs() < 1e-5, + "tile[0] (DD_CURRENT) = {}, expected {}", + tile[0], current_dd + ); + assert!( + (tile[1] - max_dd).abs() < 1e-5, + "tile[1] (DD_MAX) = {}, expected {}", + tile[1], max_dd + ); + assert!( + (tile[5] - dd_pct).abs() < 1e-5, + "tile[5] (DD_PCT) = {}, expected {}", + tile[5], dd_pct + ); + } + + /// SP15 Phase 1.3.b-followup (2026-05-07) — per-env behavioral oracle. + /// + /// Two-env config: env-0 is in recovery (DD shrinking from a peak), + /// env-1 is deepening into DD. After running a per-step series: + /// + /// * The per-env tile must reflect each env's INDEPENDENT + /// trajectory (env-0's tile slot for DD_CURRENT shrinks bar- + /// over-bar; env-1's slot grows bar-over-bar). + /// * The mean-aggregated ISV scalars must equal the cross-env + /// mean (validates `dd_state_reduce_kernel`'s mean rule). + /// * The max-aggregated ISV[DD_PERSISTENCE_MAX_INDEX=443] must + /// equal the WORSE env's persistence (validates the max rule). + /// + /// This is the canonical "envs in different DD contexts thread + /// their own DD context independently" guarantee that motivated + /// the Path A → Path B migration. + #[test] + #[ignore = "requires GPU"] + fn dd_state_per_env_diverge_independently() { + const PORTFOLIO_STRIDE: usize = 43; + const PS_PEAK_EQUITY: usize = 7; + const PS_PREV_EQUITY: usize = 9; + const N_ENVS: usize = 2; + + let stream = make_test_stream(); + let dd_budget: f32 = 0.20; + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]); + + let dd_tile_buf = unsafe { MappedF32Buffer::new(N_ENVS * 6) } + .expect("alloc MappedF32Buffer for dd_state_per_env tile"); + dd_tile_buf.write_from_slice(&vec![0.0f32; N_ENVS * 6]); + + let pos_len = (N_ENVS * PORTFOLIO_STRIDE).max(128); + let pos_state_buf = unsafe { MappedF32Buffer::new(pos_len) } + .expect("alloc MappedF32Buffer for pos_state"); + + // Six-step series: + // env-0 (recovery from a deep DD): 100 → 110 (peak) → 88 (DD~20%) → 92 → 96 → 99 + // env-1 (deepening DD): 100 → 100 (peak) → 95 → 90 → 85 → 80 + // env-0's DD trajectory shrinks (recovery); env-1's grows. + let env0_equity: [f32; 6] = [100.0, 110.0, 88.0, 92.0, 96.0, 99.0]; + let env1_equity: [f32; 6] = [100.0, 100.0, 95.0, 90.0, 85.0, 80.0]; + + let mut env0_peak: f32 = 0.0; + let mut env1_peak: f32 = 0.0; + for step in 0..6 { + env0_peak = env0_peak.max(env0_equity[step]); + env1_peak = env1_peak.max(env1_equity[step]); + let mut pos_init = vec![0.0f32; pos_len]; + pos_init[0 * PORTFOLIO_STRIDE + PS_PREV_EQUITY] = env0_equity[step]; + pos_init[0 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY] = env0_peak; + pos_init[1 * PORTFOLIO_STRIDE + PS_PREV_EQUITY] = env1_equity[step]; + pos_init[1 * PORTFOLIO_STRIDE + PS_PEAK_EQUITY] = env1_peak; + pos_state_buf.write_from_slice(&pos_init); + + launch_sp15_dd_state( + &stream, + dd_budget, + N_ENVS as i32, + dd_tile_buf.dev_ptr, + pos_state_buf.dev_ptr, + ) + .expect("launch dd_state_kernel (per-env)"); + launch_sp15_dd_state_reduce( + &stream, + N_ENVS as i32, + dd_tile_buf.dev_ptr, + isv_buf.dev_ptr, + ) + .expect("launch dd_state_reduce_kernel (per-env)"); + stream + .synchronize() + .expect("synchronize after per-env dd_state launches"); + } + + let tile = dd_tile_buf.read_all(); + // env-0: prev=99, peak=110 → current_dd = (110-99)/110 ≈ 0.1. + let env0_dd_current = tile[0 * 6 + 0]; + let env0_dd_max = tile[0 * 6 + 1]; + let env0_persistence = tile[0 * 6 + 3]; + let env0_dd_pct = tile[0 * 6 + 5]; + // env-1: prev=80, peak=100 → current_dd = (100-80)/100 = 0.2. + let env1_dd_current = tile[1 * 6 + 0]; + let env1_dd_max = tile[1 * 6 + 1]; + let env1_persistence = tile[1 * 6 + 3]; + let env1_dd_pct = tile[1 * 6 + 5]; + + // env-0 current_dd ≈ 0.1, max ≈ 0.2 (hit at step 2). + assert!( + (env0_dd_current - 0.1).abs() < 0.01, + "env-0 current_dd = {}, expected ~0.1 (recovering: 99 vs peak 110)", + env0_dd_current + ); + assert!( + (env0_dd_max - 0.2).abs() < 0.01, + "env-0 max_dd = {}, expected ~0.2 (peak DD at step 2: 88 vs 110)", + env0_dd_max + ); + // env-1 current_dd ≈ 0.2, max ≈ 0.2 (currently at the bottom). + assert!( + (env1_dd_current - 0.2).abs() < 1e-5, + "env-1 current_dd = {}, expected 0.2 (deepening: 80 vs peak 100)", + env1_dd_current + ); + assert!( + (env1_dd_max - 0.2).abs() < 1e-5, + "env-1 max_dd = {}, expected 0.2", + env1_dd_max + ); + // env-1 persistence ≥ 4 (4 consecutive below-peak bars: steps 2-5). + // env-0 persistence ≥ 4 (4 consecutive below-peak bars: steps 2-5 + // since recovery to the new HWM never happened). + assert!( + env1_persistence >= 4.0, + "env-1 persistence = {}, expected ≥ 4", + env1_persistence + ); + assert!( + env0_persistence >= 4.0, + "env-0 persistence = {}, expected ≥ 4", + env0_persistence + ); + + // Mean-aggregated ISV scalars must equal cross-env mean. + let isv = isv_buf.read_all(); + let mean_dd_current = isv[DD_CURRENT_INDEX]; + let mean_dd_pct = isv[DD_PCT_INDEX]; + let mean_persistence = isv[DD_PERSISTENCE_INDEX]; + let max_persistence = isv[DD_PERSISTENCE_MAX_INDEX]; + + let expected_mean_current = (env0_dd_current + env1_dd_current) / 2.0; + assert!( + (mean_dd_current - expected_mean_current).abs() < 1e-5, + "ISV[DD_CURRENT] = {}, expected mean of env-0 + env-1 = {}", + mean_dd_current, expected_mean_current + ); + let expected_mean_dd_pct = (env0_dd_pct + env1_dd_pct) / 2.0; + assert!( + (mean_dd_pct - expected_mean_dd_pct).abs() < 1e-5, + "ISV[DD_PCT] = {}, expected mean of env-0 + env-1 = {}", + mean_dd_pct, expected_mean_dd_pct + ); + let expected_mean_persistence = (env0_persistence + env1_persistence) / 2.0; + assert!( + (mean_persistence - expected_mean_persistence).abs() < 1e-5, + "ISV[DD_PERSISTENCE] = {}, expected mean of env-0 + env-1 = {}", + mean_persistence, expected_mean_persistence + ); + // Max-aggregated persistence: must equal the worse env. + let expected_max_persistence = env0_persistence.max(env1_persistence); + assert!( + (max_persistence - expected_max_persistence).abs() < 1e-5, + "ISV[DD_PERSISTENCE_MAX] = {}, expected max(env-0, env-1) = {}", + max_persistence, expected_max_persistence + ); } /// Test 1.4.a — buyhold on a positive-drift price series. @@ -1151,7 +1356,7 @@ mod gpu { #[ignore = "requires GPU"] fn plasticity_fires_when_persistence_exceeds_threshold() { use ml::cuda_pipeline::sp15_isv_slots::{ - DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, + DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX, }; @@ -1160,7 +1365,7 @@ mod gpu { let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } .expect("alloc MappedF32Buffer for isv bus"); let mut isv_init = vec![0.0f32; ISV_LEN]; - isv_init[DD_PERSISTENCE_INDEX] = 150.0; // exceeds threshold + isv_init[DD_PERSISTENCE_MAX_INDEX] = 150.0; // exceeds threshold isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0; isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; // not yet fired isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0; @@ -1218,7 +1423,7 @@ mod gpu { #[ignore = "requires GPU"] fn plasticity_debounced_within_fold() { use ml::cuda_pipeline::sp15_isv_slots::{ - DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, + DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX, }; @@ -1227,7 +1432,7 @@ mod gpu { let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } .expect("alloc MappedF32Buffer for isv bus"); let mut isv_init = vec![0.0f32; ISV_LEN]; - isv_init[DD_PERSISTENCE_INDEX] = 200.0; // way over threshold + isv_init[DD_PERSISTENCE_MAX_INDEX] = 200.0; // way over threshold isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0; isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 1.0; // ALREADY fired this fold isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 50.0; // mid-warm-up @@ -1286,7 +1491,7 @@ mod gpu { #[ignore = "requires GPU"] fn plasticity_no_fire_below_threshold() { use ml::cuda_pipeline::sp15_isv_slots::{ - DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, + DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX, }; @@ -1295,7 +1500,7 @@ mod gpu { let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } .expect("alloc MappedF32Buffer for isv bus"); let mut isv_init = vec![0.0f32; ISV_LEN]; - isv_init[DD_PERSISTENCE_INDEX] = 50.0; // BELOW threshold + isv_init[DD_PERSISTENCE_MAX_INDEX] = 50.0; // BELOW threshold isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0; isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0; @@ -1378,7 +1583,7 @@ mod gpu { #[ignore = "requires GPU"] fn plasticity_injection_kernel_resets_last_10pct_kaiming_he() { use ml::cuda_pipeline::sp15_isv_slots::{ - DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, + DD_PERSISTENCE_MAX_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX, }; @@ -1389,7 +1594,7 @@ mod gpu { let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } .expect("alloc MappedF32Buffer for isv bus"); let mut isv_init = vec![0.0f32; ISV_LEN]; - isv_init[DD_PERSISTENCE_INDEX] = 1_000.0; + isv_init[DD_PERSISTENCE_MAX_INDEX] = 1_000.0; isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 1.0; isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0; @@ -2461,7 +2666,7 @@ mod gpu { #[ignore = "requires GPU"] fn plasticity_fires_force_flat_then_cooldown_holds() { use ml::cuda_pipeline::sp15_isv_slots::{ - COOLDOWN_BARS_REMAINING_INDEX, DD_PERSISTENCE_INDEX, + COOLDOWN_BARS_REMAINING_INDEX, DD_PERSISTENCE_MAX_INDEX, HOLD_FLOOR_ALPHA_INDEX, HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, @@ -2495,7 +2700,7 @@ mod gpu { isv_init[HOLD_FLOOR_K_INDEX] = 10.0; isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0; - isv_init[DD_PERSISTENCE_INDEX] = 1_000.0; + isv_init[DD_PERSISTENCE_MAX_INDEX] = 1_000.0; isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 1.0; isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0; @@ -2701,6 +2906,7 @@ mod gpu { MappedF32Buffer, // r_discipline_out [N*L = 1] ml::cuda_pipeline::mapped_pinned::MappedI32Buffer, // slot_completed_normally [N*L = 1] MappedF32Buffer, // isv [ISV_LEN] + MappedF32Buffer, // dd_state_per_env [N * 6 = 6] (Phase 1.3.b-followup) ) { // Safety: tests calling this helper hold a CUDA context via // `make_test_stream`. @@ -2722,7 +2928,17 @@ mod gpu { let isv_init = vec![0.0f32; ISV_LEN]; isv.write_from_slice(&isv_init); - (out_rewards, r_disc, flag, isv) + // SP15 Phase 1.3.b-followup contract: per-env DD tile [N * 6]. + // Single-env (N=1) → 6 slots. Tests that exercise DD-aware + // shaping write the DD context into this tile (slots 0 = DD_CURRENT + // and 5 = DD_PCT) AFTER calling this helper, mirroring the + // pattern used for ISV[DD_CURRENT_INDEX] / ISV[DD_PCT_INDEX] + // in the pre-followup tests. + let dd_tile = unsafe { MappedF32Buffer::new(6) } + .expect("alloc dd_state_per_env tile"); + dd_tile.write_from_slice(&[0.0f32; 6]); + + (out_rewards, r_disc, flag, isv, dd_tile) } /// Test final_reward 1 — α-blend at cold start (α=0.5 sentinel). @@ -2734,7 +2950,7 @@ mod gpu { #[ignore = "requires GPU"] fn final_reward_alpha_blend_at_cold_start() { let stream = make_test_stream(); - let (out_rewards, r_disc, flag, isv) = + let (out_rewards, r_disc, flag, isv, dd_tile) = build_final_reward_oracle_inputs(1.0, 0.0, -2.0, 1); // Mirror constructor's α=0.5 cold-start sentinel. let mut isv_h = isv.read_all(); @@ -2747,6 +2963,7 @@ mod gpu { r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 1, 1, ).expect("launch compute_sp15_final_reward_kernel"); @@ -2775,14 +2992,18 @@ mod gpu { }; let stream = make_test_stream(); - let (out_rewards, r_disc, flag, isv) = + let (out_rewards, r_disc, flag, isv, dd_tile) = build_final_reward_oracle_inputs(1.0, 0.0, 99.0, 1); let mut isv_h = isv.read_all(); isv_h[ALPHA_SPLIT_INDEX] = 1.0; // α=1 → r_discipline ignored - isv_h[DD_CURRENT_INDEX] = 0.10; isv_h[DD_THRESHOLD_INDEX] = 0.05; isv_h[LAMBDA_DD_INDEX] = 10.0; isv.write_from_slice(&isv_h); + // Phase 1.3.b-followup: DD_CURRENT is per-env now — write to + // dd_tile[env*6 + 0] for env=0 (single-env layout). + let mut tile_h = dd_tile.read_all(); + tile_h[0] = 0.10; // DD_CURRENT + dd_tile.write_from_slice(&tile_h); ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward( &stream, @@ -2790,6 +3011,7 @@ mod gpu { r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 1, 1, ).expect("launch"); @@ -2819,13 +3041,17 @@ mod gpu { }; let stream = make_test_stream(); - let (out_rewards, r_disc, flag, isv) = + let (out_rewards, r_disc, flag, isv, dd_tile) = build_final_reward_oracle_inputs(4.0, 0.0, 0.0, 1); let mut isv_h = isv.read_all(); isv_h[ALPHA_SPLIT_INDEX] = 1.0; - isv_h[DD_PCT_INDEX] = 0.5; isv_h[DD_ASYMMETRY_LAMBDA_INDEX] = 0.5; isv.write_from_slice(&isv_h); + // Phase 1.3.b-followup: DD_PCT is per-env now — write to + // dd_tile[env*6 + 5] for env=0 (single-env layout). + let mut tile_h = dd_tile.read_all(); + tile_h[5] = 0.5; // DD_PCT + dd_tile.write_from_slice(&tile_h); ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward( &stream, @@ -2833,6 +3059,7 @@ mod gpu { r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 1, 1, ).expect("launch"); @@ -2864,13 +3091,16 @@ mod gpu { use ml::cuda_pipeline::sp15_isv_slots::DD_ASYMMETRY_LAMBDA_INDEX; let stream = make_test_stream(); - let (out_rewards, r_disc, flag, isv) = + let (out_rewards, r_disc, flag, isv, dd_tile) = build_final_reward_oracle_inputs(-3.0, 0.0, 0.0, 1); let mut isv_h = isv.read_all(); isv_h[ALPHA_SPLIT_INDEX] = 1.0; - isv_h[DD_PCT_INDEX] = 0.5; isv_h[DD_ASYMMETRY_LAMBDA_INDEX] = 0.5; isv.write_from_slice(&isv_h); + // Phase 1.3.b-followup: DD_PCT is per-env now. + let mut tile_h = dd_tile.read_all(); + tile_h[5] = 0.5; // DD_PCT + dd_tile.write_from_slice(&tile_h); ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward( &stream, @@ -2878,6 +3108,7 @@ mod gpu { r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 1, 1, ).expect("launch"); @@ -2898,7 +3129,7 @@ mod gpu { #[ignore = "requires GPU"] fn final_reward_sp12_cap_clamps_both_directions() { let stream = make_test_stream(); - let (out_rewards, r_disc, flag, isv) = + let (out_rewards, r_disc, flag, isv, dd_tile) = build_final_reward_oracle_inputs(20.0, -20.0, 0.0, 1); let mut isv_h = isv.read_all(); isv_h[ALPHA_SPLIT_INDEX] = 1.0; @@ -2910,6 +3141,7 @@ mod gpu { r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 1, 1, ).expect("launch"); @@ -2968,12 +3200,20 @@ mod gpu { isv_h[ALPHA_SPLIT_INDEX] = 1.0; isv.write_from_slice(&isv_h); + // Phase 1.3.b-followup: per-env DD tile [N=2 * 6] = 12 slots, + // all zero — DD has no effect on the reward (α=1 + zero DD ⇒ + // identity composition). + let dd_tile = unsafe { MappedF32Buffer::new(2 * 6) } + .expect("alloc dd_state_per_env tile"); + dd_tile.write_from_slice(&vec![0.0f32; 2 * 6]); + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_final_reward( &stream, out_rewards.dev_ptr, r_disc.dev_ptr, flag.dev_ptr, isv.dev_ptr, + dd_tile.dev_ptr, 2, 1, ).expect("launch"); diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 10af62765..e186a777b 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-followup — per-env DD redesign (Path A → Path B atomic migration) (2026-05-07): closes the Phase 1.3.b deferred per-env redesign per `feedback_no_partial_refactor`. The original Phase 1.3.b atomic fix (`132609724`) chose **env 0 as the canonical DD observable** and wrote 6 ISV scalar slots [401..407) directly from env-0's portfolio row. That choice was a documented Path A pragmatic compromise — every downstream consumer (the fused reward composer, plasticity injection trigger, dd_trajectory recovery proxy, HEALTH_DIAG diagnostic) read the env-0-canonical scalars even though each production env has its OWN portfolio + DD trajectory. When env-3 was in 30% DD and env-0 was at ATH, the model learning from env-3's transitions saw `dd_pct=0` and silently skipped the recovery shaping. (1) **`dd_state_kernel.cu` reshape** — grid `[n_envs, 1, 1]` × `[1, 1, 1]`, one thread per env, sequential per-bar walk per env. Each thread computes its env's DD trajectory and writes 6 scalars to a NEW per-env tile buffer `dd_state_per_env[n_envs * 6]` (layout `[env_0_current, env_0_max, env_0_rec, env_0_pers, env_0_calmar, env_0_pct, env_1_current, ...]`); the kernel no longer touches the ISV scalar slots. (2) **New `dd_state_reduce_kernel.cu`** — single-block tree-reduce (no `atomicAdd` per `feedback_no_atomicadd`; BLOCK=256 with strided initial accumulation handles n_envs up to 32768 on H100). Mean-aggregates across envs into the 6 scalar ISV slots [401..407) (smooth across-env summary for HEALTH_DIAG diagnostic — preserves backward-compatible "this is one DD scalar to look at" semantics; only the across-env aggregation rule changed Path A: env-0-canonical → Path B: mean-across-envs); max-aggregates DD_PERSISTENCE into a NEW slot **`DD_PERSISTENCE_MAX_INDEX = 443`** (consumed by `plasticity_injection_kernel` so the global advantage-head reset fires when ANY env's persistence exceeds the threshold — one set of advantage weights → global firing semantics → max-aggregate is the only correct rule; mean-aggregate would silently gate the trigger when only a fraction of envs are in long DD). Aggregation rule rationale per `feedback_isv_for_adaptive_bounds` (the rule itself is structural, not adaptive). `ISV_TOTAL_DIM 443 → 444` to fit the new slot; `SP15_SLOT_END 443 → 444`; `SP15_SLOT_COUNT 46 → 47`. (3) **`compute_sp15_final_reward_kernel.cu` per-env DD lookup** — index→env mapping `env_id = (idx % (N*L)) / L` (the layout is `[on_policy_N*L | cf_N*L]` so `slot_it = idx % (N*L)` is the per-(i,t) flat index, and `env_id = slot_it / L`); each `(i,t)` slot looks up ITS env's `dd_pct` and `dd_current` from the per-env tile (`tile[env_base + 5]` and `tile[env_base + 0]` respectively). On-policy and CF threads at the same (i,t) read the SAME tile entry (one DD trajectory per env, shared across slot kinds). The `__device__` helpers `sp15_dd_asymmetric_reward` / `sp15_dd_penalty` migrated from "read ISV by themselves" to "accept dd_pct + dd_current as scalar parameters" so the per-(i,t) per-env lookup happens at the kernel level, with the helpers staying scalar-clean. Atomic per `feedback_no_partial_refactor` — both gain-side multiplier (Stage 2) and quadratic DD penalty (Stage 3) migrated together. (4) **`plasticity_injection_kernel.cu` consumer migration** — persistence read switched from ISV[DD_PERSISTENCE_INDEX=404] (mean-aggregate across envs) to ISV[DD_PERSISTENCE_MAX_INDEX=443] (max-aggregate across envs). The fired-this-fold flag and warm-bars counter remain global scalars (one set of advantage weights ⇒ one trigger condition ⇒ one fired flag); the kernel itself stays unchanged downstream of the read. (5) **HEALTH_DIAG semantic shift (documented breaking change)** — slots 401-406 now report the cross-env MEAN of DD_CURRENT/DD_MAX/DD_RECOVERY_BARS/DD_PERSISTENCE/CALMAR/DD_PCT, NOT env-0's value. For single-env smoke configs (n_envs=1) the mean equals env-0's value, so the diagnostic stream is bit-stable across the migration. For production (n_envs ≥ 2048) the mean is genuinely different — but it's the right diagnostic: a smoothed across-env summary of "where is the training instance, on average, in its DD trajectory". (6) **Per-env tile owned by `GpuExperienceCollector`** (NOT the trainer) — the collector knows `alloc_episodes` (= `n_envs`); the trainer's `batch_size` is a different quantity (training batch size). The new field `sp15_dd_state_per_env: MappedF32Buffer([alloc_episodes * 6])` is allocated in the constructor next to `sp13_hold_rate_buf` and reset to zero at fold boundary via the `sp15_dd_state_per_env` registry-arm dispatch (`host_slice_mut().fill(0.0)` — mirrors the `sp15_alpha_warm_count` / `sp15_dd_trajectory_prev_dd` non-ISV mapped-pinned reset pattern). (7) **Per-step launch order in `gpu_experience_collector.rs::launch_timestep_loop` step 5b**: `dd_state_kernel` (per-env tile producer) → `dd_state_reduce_kernel` (ISV scalar consumer; gated on `isv_signals_dev_ptr != 0`) → `alpha_split_producer_kernel` (existing) → `compute_sp15_final_reward_kernel` (per-(i,t) consumer; threads `dd_state_per_env_ptr` so each slot's reward shaping uses ITS env's DD context). All four launches on the same stream — CUDA serialises producer→consumer without explicit event sync. Outside the exp-fwd graph capture region per `pearl_no_host_branches_in_captured_graph`. (8) **Layout fingerprint break** — added markers `DD_PERSISTENCE_MAX=443; ISV_TOTAL_DIM=444; DD_STATE_PER_ENV=sp15_phase_1_3_b_followup` to `layout_fingerprint_seed`; pre-followup checkpoints WILL NOT LOAD after this commit (greenfield OK per spec Q1 — same precedent as the Wave 4.1a `TRUNK_INPUT_DD_PCT=sp15_wave_4_1a` fingerprint break). (9) **Migrated 6 oracle tests + added 1 new behavioral test** — `dd_state_kernel_tracks_drawdown_correctly` (now drives both kernels with `n_envs=1` so the existing assertions still hold via the mean-of-one identity); `plasticity_fires_when_persistence_exceeds_threshold` / `plasticity_debounced_within_fold` / `plasticity_no_fire_below_threshold` / `plasticity_fires_force_flat_then_cooldown_holds` (writes seeded ISV[DD_PERSISTENCE_MAX_INDEX=443] instead of the now-stale ISV[DD_PERSISTENCE_INDEX=404]); 5 `final_reward_*` tests (now thread `dd_tile.dev_ptr` through `launch_sp15_final_reward` and write per-env DD into `dd_tile[0..6]` instead of ISV[401..407)). NEW: `dd_state_per_env_diverge_independently` — two-env config where env-0 is in recovery (DD shrinking) and env-1 is deepening into DD; verifies the per-env tile reflects each env's INDEPENDENT trajectory, the mean-aggregated ISV scalars equal the cross-env mean, and the max-aggregated ISV[DD_PERSISTENCE_MAX_INDEX=443] equals the WORSE env's persistence — the canonical "envs in different DD contexts thread their own DD context independently" guarantee that motivated the Path A → Path B migration. (10) **Phase 1.3.b-followup-B (separate split)** — `dd_trajectory_decreasing_kernel` + `per_insert_pa` migration is NOT in scope for this commit. The current Wave 4.3 implementation reads ISV[DD_TRAJECTORY_DECREASING_INDEX=439] at insert-batch time (epoch end), which carries the LAST step's value applied uniformly to ALL inserted transitions — a known pre-existing limitation where ANY env's trajectory could win the boost regardless of the inserting transition's actual env. The proper fix requires a per-(env, t) trajectory buffer `[N*L]` written by `dd_trajectory_kernel` and an `env_id`-aware lookup in `per_insert_pa` (the insert position `j` of the batch maps to `env_id = (j % (N*L)) / L` since the inserted batch has the same `[on_policy_N*L | cf_N*L]` layout). That's a fundamentally different contract change (per-(env, t) tile + env_id-mapped lookup at insert time, NOT just per-env tile + reduction). Splitting it into Phase 1.3.b-followup-B preserves the no-partial-refactor invariant within each migration: contract A (DD state ISV scalar → per-env tile) lands here atomically with all 5 of its consumers (final_reward kernel + plasticity + HEALTH_DIAG + 6 oracle tests + new behavioral test); contract B (DD trajectory ISV scalar → per-(env, t) tile) lands separately with its lone consumer (`per_insert_pa`). Files touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (reshape), `dd_state_reduce_kernel.cu` (NEW), `compute_sp15_final_reward_kernel.cu` (per-env lookup), `plasticity_injection_kernel.cu` (slot 404→443), `sp15_reward_axis_helpers.cuh` (helpers take dd_pct/dd_current as scalars), `sp15_isv_slots.rs` (+ DD_PERSISTENCE_MAX_INDEX), `gpu_dqn_trainer.rs` (launchers + ISV_TOTAL_DIM + fingerprint), `gpu_experience_collector.rs` (per-env tile field + 4-launch sequence), `state_reset_registry.rs` (+ 2 entries), `training_loop.rs` (+ 2 dispatch arms), `build.rs` (+ dd_state_reduce_kernel cubin), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (6 migrated + 1 new test), `docs/isv-slots.md` (slot map), `docs/dqn-wire-up-audit.md` (this entry). Hard rules: `feedback_no_partial_refactor` (5 consumers + 6 oracle tests + 1 new test all atomic in this commit); `feedback_no_atomicadd` (reduction kernel uses block tree-reduce, not atomicAdd); `feedback_no_cpu_compute_strict` (per-env logic + reduction GPU-resident); `feedback_isv_for_adaptive_bounds` (mean/max aggregation rules are structural, documented rationale). + SP15 Phase 3.5.4.c — production caller + OR-gate consumer for plasticity injection (2026-05-07): closes out the Wave 4.2 (`ef08611d3`) orphan-with-tests-only landing per `feedback_wire_everything_up`. Wave 4.2 landed `plasticity_injection_kernel.cu` + `launch_sp15_plasticity_injection` + 4 oracle tests but NO production caller existed — the kernel was reachable only from the oracle-test path. Phase 3.5.4.c creates the production caller AND wires the action-selection consumer atomically per `feedback_no_partial_refactor` so the spec §9.2 (3.5.4) two-step recovery (Flat-on-fire → Hold-during-warm-up) fires end-to-end. (1) **fan_in audit-doc inline correction** [load-bearing per `feedback_trust_code_not_docs`]: Wave 4.2's audit doc point (9) claimed default-config `fan_in = adv_h × num_atoms = 6528` and `stddev = sqrt(2/6528) ≈ 0.0175`. That was wrong — Kaiming-He variance is `2/fan_in` where `fan_in` is the INPUT dim per output unit, i.e. `adv_h` (not `adv_h × num_atoms` which is the full row-output element count of one branch's projection). The correct values are **`fan_in = 128`** and **`stddev = sqrt(2/128) ≈ 0.125`**. The error was inline-corrected in the Wave 4.2 entry so future readers don't propagate the wrong number; the corrected entry credits Phase 3.5.4.c. (2) **Production caller in `gpu_experience_collector.rs::collect_experiences_gpu`'s per-step loop** (step 4-pre, BEFORE `experience_action_select`): when `sp15_w_b0out_dev_ptr != 0 && isv_signals_dev_ptr != 0`, invokes `launch_sp15_plasticity_injection(stream, isv, w_b0out, n_weights=branch_0_size×num_atoms×adv_h, m_warm=200.0, fan_in=adv_h, seed=mix_seed(0xC0DE_F00D_5915_3F4C) ^ t)`. The trigger writes ISV[436] (fired flag) + ISV[438] (warm counter) which the action_select kernel then reads (next launch on the same stream — CUDA serializes producer→consumer without explicit event sync). When unwired (test scaffold), the launch is skipped and the action_select OR-gate consumer degenerates to cooldown-only (bit-identical pre-3.5.4.c behaviour). Placement: inside the `else` arm of the `seed_phase_active_cache` gate (network branch only — scripted-policy seed phase forces specific actions, plasticity injection during seed phase is meaningless). Out of the `exp_fwd_graph` capture region (which ends at the cuBLAS forward earlier in the loop body), so per `pearl_no_host_branches_in_captured_graph` the host-side gating compiles cleanly. (3) **Branch choice — w_b0out (directional) only** per spec §9.2 (3.5.4) "the recovery is about escaping stuck-in-flat regimes": the directional advantage-head last-Linear weight tensor (param-buffer index 19) gets the trailing-10% Kaiming-He reset; magnitude/order/urgency branches are untouched (their conditioning on direction makes them naturally re-train once direction starts choosing again). The trainer exposes the target via `pub fn sp15_w_b0out_target() -> (u64, i32, i32)` returning `(dev_ptr_at_offset_19, n_weights, fan_in)`; the collector's matching `set_sp15_plasticity_target(w_ptr, n_weights, fan_in)` setter wires it. (4) **n_weights derivation**: `branch_0_size × num_atoms × adv_h` from `compute_param_sizes(&self.config)[19]`. Default config: 4 × 51 × 128 = **26112** (vs Wave 4.2's oracle-test 64 + 1280); reset_count = `n_weights / 10` = 2611 trailing weights re-sampled per fire. The launcher's grid_dim auto-scales to `((26112/10 + 255) / 256)` = 11 blocks × 256 threads. (5) **fan_in = adv_h** per the kernel's per-output-unit variance contract: the kernel computes `weights[i] = curand_normal × sqrt(2/fan_in)` for each i in the trailing reset region — `fan_in` here is the input-dim of the per-unit dot product, which equals `adv_h` regardless of `num_atoms`. (6) **Seed source — deterministic per-step**: `mix_seed(0xC0DE_F00D_5915_3F4C) ^ (t as u64)` where `t` is the per-step timestep iterator and `mix_seed` is the SplitMix64 avalanche over the global `FOXHUNT_SEED` env var. Same `(FOXHUNT_SEED, t)` always produces the same `(seed, tid)` cuRAND init in the kernel, satisfying the established codebase determinism contract (`feedback_h100_gpu` + checkpoint-resumed reproducibility). The constant `0xC0DE_F00D_5915_3F4C` is a Phase-3.5.4.c-unique 64-bit base so multi-launcher runs (alpha_split / cooldown / dd_state / etc.) get uncorrelated mix outputs even when `FOXHUNT_SEED == 42` (default). (7) **OR-gate consumer extension in `experience_action_select`** (`experience_kernels.cu`): added one new kernel arg `float plasticity_m_warm` (after `n_atoms`); the kernel now reads ISV[438] (warm counter) alongside ISV[435] (cooldown remaining); the SP15 Phase 3.5.3.b cooldown predicate `(cooldown_remaining > 0)` is extended to `(cooldown_remaining > 0) || (plasticity_warm_remaining > 0)`. When the plasticity launch above is unwired, `m_warm` is passed as 0.0f and warm stays at 0 → the OR-gate degenerates to cooldown-only (bit-identical pre-3.5.4.c). (8) **Flat-on-fire-bar detection — option (a), `warm ≥ m_warm − 1.5f`** per the trigger-then-decrement convention: the trigger kernel sets `warm = m_warm` then decrements once in the same launch (mirroring the cooldown_kernel's identical sequence with the `[198, 200]` test tolerance). On the fire bar, action_select observes `warm = m_warm − 1`, which satisfies the predicate `warm ≥ m_warm − 1.5f` with a 0.5f f32-equality tolerance. On bar T+1 (post-fire warm-up), `warm = m_warm − 2 < m_warm − 1.5f`, so the fire-bar branch is dead and the OR-gate cooldown branch fires DIR_HOLD instead. The fire-bar branch supersedes the cooldown branch via an `if/else if` chain — the spec sequence is: fire-bar → DIR_FLAT (close any open position); subsequent warm bars → DIR_HOLD. No new ISV slot needed (option (b) was rejected as adding unnecessary state); no host-side detection needed (option (c) was rejected as introducing a host branch in the per-step loop). The `m_warm > 0` guard inside the predicate ensures the fire-bar branch is dead when plasticity is unwired (m_warm=0 and warm=0 — both zero → `warm >= m_warm − 1.5f = -1.5f` is true, but the `m_warm > 0` factor rules it out). (9) **Trainer plumbing in `training_loop.rs`**: alongside the existing `set_isv_signals_ptr` / `set_sp15_alpha_warm_count_ptr` / `set_sp15_dd_trajectory_prev_dd_ptr` calls, added `collector.set_sp15_plasticity_target(fused.trainer().sp15_w_b0out_target())`. Mirrors the same wiring pattern. (10) **NEW oracle test** `plasticity_fires_force_flat_then_cooldown_holds` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after `action_select_no_force_hold_when_cooldown_zero`): drives the full per-step launch sequence (`launch_sp15_plasticity_injection` → `experience_action_select`) across `M_warm + 1` bars with `M_warm = 5` (shrunk from production's 200 to keep the test fast). With Long strongly preferred (e_dir gap = +1.0) and the trigger-fires ISV state pre-loaded (DD_PERSISTENCE=1000, threshold=1, fired=0, warm=0): bar 0 (fire bar) asserts dir_picked == DIR_FLAT (=3) — proves the fire-bar branch fired even though Long would normally win the argmax; bars 1..M_warm−1 (warm-up) assert dir_picked == DIR_HOLD on every bar — proves the OR-gate cooldown fires when warm ∈ {3, 2, 1}; bar M_warm−1 (the boundary call where warm transitions 1→0) asserts dir_picked == DIR_LONG — proves normal Thompson argmax resumes once warm=0 and the OR-gate is inactive. Also asserts the ISV side-effects bar-by-bar: fired flips 0→1 on the fire bar then debounces (stays 1) for the rest of the test, warm decrements from M_warm−1 down to 0 with the trigger's underflow guard preventing negative values. (11) **Eval/backtest path** (`gpu_backtest_evaluator.rs::submit_dqn_step_loop_cublas`): updated to pass `m_warm = 0.0f` to `experience_action_select` because plasticity injection is a TRAINING-only mechanism (it resets advantage-head weights to break out of stuck-in-flat regimes during experience collection); during deterministic eval the weights must remain frozen at their checkpoint values. With `m_warm = 0.0f` the action_select fire-bar predicate is dead, and `PLASTICITY_WARM_BARS_REMAINING` stays at its 0 sentinel through eval (no producer running) so the OR-gate degenerates to cooldown-only (bit-identical pre-3.5.4.c). (12) **Atomicity** per `feedback_no_partial_refactor`: kernel signature change + collector field/setter/launch + trainer accessor + trainer plumbing + eval-path update + 2 test launcher updates (mod.rs + distributional_q_tests.rs) + new behavioral oracle test + audit doc inline correction + new audit doc entry all in one commit; pre-existing test launchers were updated to pass `m_warm = 0.0f` so they preserve their pre-3.5.4.c semantics (no fire-bar / OR-gate plasticity branch). Touched: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (kernel signature `+plasticity_m_warm: float`, fire-bar detection block via `warm ≥ m_warm − 1.5f`, OR-gate extension `cooldown_active = (cooldown_remaining > 0) || (warm > 0)`, new `if (plasticity_fire_bar)` branch BEFORE the existing `else if (cooldown_active)` cooldown branch), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (new `pub fn sp15_w_b0out_target()` accessor returning `(dev_ptr, n_weights, fan_in)`), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (new `sp15_w_b0out_dev_ptr` / `sp15_w_b0out_n_weights` / `sp15_w_b0out_fan_in` fields + `set_sp15_plasticity_target` setter + per-step `launch_sp15_plasticity_injection` invocation in step 4-pre + `m_warm` arg passed to `experience_action_select`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (one new wiring block calling `set_sp15_plasticity_target`), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (eval-path action_select launch passes `m_warm = 0.0f`), `crates/ml/src/cuda_pipeline/mod.rs` (test launcher passes `m_warm = 0.0f`), `crates/ml/src/trainers/dqn/distributional_q_tests.rs` (test launcher passes `m_warm = 0.0f`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (existing `launch_action_select_for_test` helper passes `m_warm = 0.0f`; new `launch_action_select_for_test_with_m_warm` helper takes `m_warm` as an explicit parameter; new behavioral oracle test `plasticity_fires_force_flat_then_cooldown_holds`), `docs/dqn-wire-up-audit.md` (this entry + Wave 4.2 entry's point (9) inline-corrected). Hard rules: `feedback_no_partial_refactor` (kernel + caller + consumer + 7 launcher-call-site updates + new behavioral test + audit doc inline correction + new audit doc entry all atomic — every consumer of the kernel signature change migrates simultaneously), `feedback_no_quickfixes` (real fan_in derivation from `cfg.adv_h` via `sp15_w_b0out_target()` accessor — not a magic constant), `feedback_no_cpu_compute_strict` (the per-step plasticity launch + the kernel-side fire-bar detection are fully GPU-resident; the host-side `if self.sp15_w_b0out_dev_ptr != 0` gate is a cold-path one-shot per step, not a per-thread-of-kernel branch), `pearl_no_host_branches_in_captured_graph` (the host-side launch + parameter computation runs OUTSIDE the `exp_fwd_graph` capture region — that graph ends at the cuBLAS forward in step 2, well before step 4-pre), `feedback_trust_code_not_docs` (the Wave 4.2 audit-doc fan_in error is corrected in-place; the right value `adv_h = 128` is read from the trainer config rather than re-derived in prose), `feedback_wire_everything_up` (Wave 4.2's orphan launcher is now driven from production every step; the `Phase 3.5.4.c remains a separate follow-up` paragraph in the Wave 4.2 entry is now historical context for THIS commit rather than an open TODO; SP15 Phase 3.5 recovery-dynamics chain is end-to-end production-wired — 3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.c + 3.5.5 + 3.5.5.b all firing). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean; `SQLX_OFFLINE=true cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` 5 of 5 plasticity oracle tests green (4 Wave 4.2 + new `plasticity_fires_force_flat_then_cooldown_holds`); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored action_select --nocapture` 3 of 3 SP15 3.5.b/3.5.3.b action_select oracle tests still green (the `m_warm = 0.0f` defaults preserve their pre-3.5.4.c behaviour bit-identically); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored cooldown --nocapture` 6 of 6 cooldown + action_select tests green; `cargo test -p ml --features cuda --lib` HOLDS the Wave 4.3 baseline (946 pass / 13 fail) on RTX 3050 Ti — no new regressions, the 13 failing tests are the same pre-existing infra failures (OFI features missing, checkpoint round-trip, spectral norm GPU init, etc.). SP15 Wave 4.3 / Phase 3.5.5.b — PER sampler integration: recovery transitions oversampled (2026-05-07): closes out the Phase 3.5.5 (`69b8fdb61`) deferred PER sampling consumer per `feedback_wire_everything_up`. Phase 3.5.5 landed `dd_trajectory_decreasing_kernel` writing `ISV[DD_TRAJECTORY_DECREASING_INDEX=439]` per-step but DEFERRED both the production launch AND the PER consumer that reads ISV[439] + ISV[440] to bias replay sampling; Wave 4.3 wires both atomically. (1) **Investigation finding — Case B (existing TD-error-driven priority sampler)**: The PER architecture in `crates/ml-dqn/src/gpu_replay_buffer.rs` already weights replay by per-transition priority (TD-error driven via `per_update_pa`). `priorities[idx]` holds the raw priority (max_priority sentinel at insert, `|td|^alpha + ε` after training); `priorities_pa[idx] = priorities[idx]^alpha` feeds the prefix-sum sampler `per_prefix_scan` + deterministic-stratified `per_sample` (zero CPU compute, fully GPU-resident). The recovery oversample boost slots in cleanly as a multiplier on `priorities[idx]` at insert time — when the agent is currently in a recovery (DD shrinking from a non-trivial level), every transition in the freshly-collected batch lands in the buffer with `effective = base × (1 + ω × δ) = 3.0` instead of `1.0` (with sentinel ω=2.0, δ=1.0 from `dd_trajectory_decreasing_kernel`), so for the duration that those transitions sit in the buffer pre-update they get sampled `3.0^0.6 ≈ 1.93×` more often than baseline transitions. Once `per_update_pa` overwrites the priority based on TD-error, normal PER takes over — the boost only amplifies early sampling, exactly matching the spec's "agent learns recovery dynamics over typical average-DD Bellman noise" framing. This is materially smaller than Case A (uniform sampler retrofit) would have been: the existing alias-method-equivalent prefix-sum machinery stays untouched. (2) **Kernel signature change** in `crates/ml-dqn/src/per_kernels.cu:per_insert_pa`: added `float* __restrict__ priorities` (now also written, subsuming the previous `scatter_insert_f32` broadcast of `max_priority`) and `const float* __restrict__ isv` (nullable; reads `isv[439]` for δ and `isv[440]` for ω when non-null, falls back to `boost_factor=1.0` when null). Kernel body computes `effective = raw_priorities[i] × boost_factor`, writes both `priorities[idx] = effective` and `priorities_pa[idx] = effective^alpha`. Slot indices are inline literals (439, 440) matching the `dd_trajectory_kernel.cu` literal-index convention — the canonical names live in `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` with compile-time `assert_eq!` regression checks pinning them in `sp15_slot_layout_locked`. (3) **`GpuReplayBuffer` API change** in `crates/ml-dqn/src/gpu_replay_buffer.rs`: added `isv_signals_dev_ptr: u64` field (initialised to 0 in the constructor — meaning NULL → no boost), `pub fn set_isv_signals_ptr(&mut self, dev_ptr: u64)` setter (mirrors the existing `set_trainer_buffers` / `set_learning_health` plumbing pattern), `pub const fn isv_signals_dev_ptr(&self) -> u64` accessor for the wiring round-trip oracle assertion, and `pub const fn priorities_pa_slice(&self) -> &CudaSlice` accessor for the new oracle test. `insert_batch` now passes `&isv_ptr` + `&mut self.priorities` to `per_insert_pa` and the redundant `scatter_insert_f32` broadcast that previously seeded `priorities[idx] = max_priority` is REMOVED — `per_insert_pa`'s new `priorities[idx] = effective` store subsumes its work (effective equals max_priority × 1.0 = max_priority when ISV is unwired or δ=0, preserving pre-3.5.5.b semantics for non-recovery transitions). (4) **Production launch wiring** in `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`: added `sp15_dd_trajectory_prev_dd_dev_ptr: u64` field + `set_sp15_dd_trajectory_prev_dd_ptr` setter (same pattern as the existing `set_sp15_alpha_warm_count_ptr`); inside `collect_experiences_gpu`'s per-step loop, immediately after the existing `launch_sp15_dd_state` call (which writes `ISV[DD_PCT_INDEX=406]`), the new "5b. SP15 Phase 3.5.5.b — DD_TRAJECTORY_DECREASING per-step proxy" block invokes `launch_sp15_dd_trajectory_decreasing` gated on both `isv_signals_dev_ptr != 0` AND `sp15_dd_trajectory_prev_dd_dev_ptr != 0` — when either is unwired (test scaffold path), the kernel is skipped and ISV[439] stays at sentinel 0 (no recovery boost). Mirrors the `launch_sp15_dd_state` placement (same per-step branch, same single-thread/single-block kernel shape, same on-stream serialization without event sync). (5) **Trainer-side plumbing** in `crates/ml/src/trainers/dqn/trainer/training_loop.rs`: alongside the existing `set_isv_signals_ptr` + `set_sp15_alpha_warm_count_ptr` calls in `train_with_data_full_loop_slices`, added two new wiring blocks — `collector.set_sp15_dd_trajectory_prev_dd_ptr(fused.trainer().sp15_dd_trajectory_prev_dd.dev_ptr)` so the per-step kernel can advance its persistent state, AND `agent.primary_dqn_mut().memory.gpu.set_isv_signals_ptr(fused.isv_signals_dev_ptr())` so `per_insert_pa` can read the recovery slots. Without the second wiring, the buffer's `isv_signals_dev_ptr` stays NULL and the kernel falls back to boost_factor=1.0 (no oversample) — preserving the smoke / oracle-test path that exercises PER without the SP15 bus. (6) **NEW oracle test** `per_sampler_weights_recovery_transitions_higher` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after the three `dd_trajectory_*` tests): allocates a `MappedF32Buffer(ISV_LEN)` with ISV[440]=2.0 (sentinel ω) + ISV[439]=1.0 (recovery active); constructs a `GpuReplayBuffer` with capacity 64; calls `set_isv_signals_ptr(isv_buf.dev_ptr)` (asserts `isv_signals_dev_ptr()` round-trips correctly); inserts 8 zero-padded synthetic transitions (batch 1, recovery=1) → expected `priorities[0..8] = 3.0`; flips ISV[439]=0 and inserts 8 more (batch 2, recovery=0) → expected `priorities[8..16] = 1.0`; reads `priorities_slice` + `priorities_pa_slice` via `memcpy_dtoh`; asserts batch-1 priorities all equal 3.0 within 1e-5, priorities_pa all equal `3.0^0.6 ≈ 1.9332` within 1e-4, batch-2 priorities all equal 1.0; finally computes `boosted_mean / baseline_mean` ratio and asserts it equals exactly 3.0 within 1e-5 — the recovery oversample factor by construction. The deterministic-stratified `per_sample` kernel weights samples by `priorities^alpha`, so verifying the priority-buffer write directly (paired with the existing `gpu_per_integration_test::test_gpu_per_priorities_affect_sampling` empirical-bias coverage) gives the full statistical-bias proof without re-running a 10000-sample experiment in the oracle suite. (7) **Phase 3.5.5 deferred consumer eliminated** per `feedback_wire_everything_up`: `dd_trajectory_decreasing_kernel` is now invoked from production code (the experience collector's per-step loop), and its ISV[439] write is now CONSUMED by `per_insert_pa` to bias replay sampling. The Phase 3.5.5 commit message's "PER sampling consumer (DEFERRED follow-up)" narrative is closed; the `dd_trajectory_kernel.cu` header comment's deferred-consumer paragraph is now historical context for the Wave 4.3 wire-up rather than an open TODO. (8) **DD_TRAJECTORY_FLOOR producer (slot 441) and ISV-driven RECOVERY_OVERSAMPLE_WEIGHT producer (slot 440) remain documented Phase 3.5.5 follow-ups** per `feedback_isv_for_adaptive_bounds` — the kernel reads from these slots rather than hardcoded literals, so when the rolling-25th-percentile `dd_pct` producer for slot 441 and the `proportional-to-current-dd_pct` producer for slot 440 land in their own atomic commits, the consumers (`dd_trajectory_kernel` reading floor; `per_insert_pa` reading ω) are zero-effort migrations. Wave 4.3 uses the constructor-seeded sentinels (floor=0.02, ω=2.0). (9) **Atomicity** per `feedback_no_partial_refactor`: kernel signature change + `GpuReplayBuffer` API extension + `insert_batch` wiring + experience-collector field/setter/launch + trainer plumbing + oracle test + audit doc all in one commit; every `insert_batch` consumer (production `training_loop.rs` + 3 smoke test files in `crates/ml/src/trainers/dqn/smoke_tests/` + `crates/ml/tests/gpu_per_integration_test.rs` + the new oracle test) compiles unchanged because the boost is gated by the optional `set_isv_signals_ptr` setter (zero=NULL preserves the pre-3.5.5.b behavior modulo the now-merged `priorities[idx]` write). (10) **GPU-residency audit**: weight computation is fully GPU-resident — `per_insert_pa` reads ISV[439]/ISV[440] from the GPU bus, multiplies on the GPU, writes both `priorities` and `priorities_pa` from a single kernel launch. No DtoH/HtoD transfers added. The `dd_trajectory_decreasing_kernel` runs on the same stream as `env_step` + `launch_sp15_dd_state` so CUDA serializes the producer→consumer chain without explicit event sync (`pearl_no_host_branches_in_captured_graph` clean — both are outside the exp-fwd graph capture region per the existing `launch_sp15_dd_state` precedent at `gpu_experience_collector.rs:4555-4561`). Touched: `crates/ml-dqn/src/per_kernels.cu` (`per_insert_pa` signature `+priorities, +isv`, body rewrite for the boost multiplier + dual `priorities`/`priorities_pa` writes, header comment expanded for the Wave 4.3 contract), `crates/ml-dqn/src/gpu_replay_buffer.rs` (new `isv_signals_dev_ptr` field + setter + accessor + `priorities_pa_slice` accessor; `insert_batch` arg list extended for the kernel; redundant `scatter_insert_f32` for `priorities` REMOVED), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (new `sp15_dd_trajectory_prev_dd_dev_ptr` field + setter; per-step `launch_sp15_dd_trajectory_decreasing` invocation gated on both ISV ptr + prev_dd ptr being non-zero), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (two new wiring blocks for collector prev_dd ptr + replay-buffer ISV ptr), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (new oracle test `per_sampler_weights_recovery_transitions_higher`). Hard rules: `feedback_no_cpu_compute_strict` (boost computation + dual buffer write are GPU-resident inside `per_insert_pa`), `feedback_no_partial_refactor` (kernel + sampler + producer launch + trainer plumbing + oracle test + audit doc all atomic), `feedback_isv_for_adaptive_bounds` (RECOVERY_OVERSAMPLE_WEIGHT lives in ISV[440] read at kernel time — the constructor-seeded 2.0 sentinel is a structural cold-start anchor; the ISV-driven producer remains a documented Phase 3.5.5 follow-up; the kernel reads the slot rather than a hardcoded literal so the follow-up swap is a no-op at the consumer), `feedback_no_stubs` (kernel performs the actual priority multiplication and writes both columns; no return-zero stubs; the boost-factor=1.0 fallback for unwired ISV is the explicit semantic for callers that don't need recovery oversample, not a stub), `feedback_no_atomicadd` (per-thread independent writes to unique `priorities[idx]` / `priorities_pa[idx]`; no atomics, no reductions — `per_insert_pa` was already grid-strided element-wise; the kernel signature extension preserves the per-thread independence), `feedback_wire_everything_up` (Phase 3.5.5's deferred PER consumer is closed in this commit — `dd_trajectory_kernel` is launched in production AND its ISV[439] write is consumed by `per_insert_pa`; this completes the SP15 Phase 3.5 recovery-dynamics chain: 3.5.2 + 3.5.3 + 3.5.4 + 3.5.5 + 3.5.5.b all wired). Verified: `SQLX_OFFLINE=true cargo check -p ml-dqn --features cuda` clean; `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings); `cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored dd_trajectory --nocapture` 3 of 3 dd_trajectory oracle tests still green; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored per_sampler --nocapture` 1 of 1 new PER sampler oracle test green (recovery slots read 3.0, baseline slots read 1.0, ratio = 3.0 within 1e-5 — boosted/baseline statistical bias verified by construction); `cargo test -p ml --features cuda --lib gpu_residency` 8 of 8 (1 ignored) replay-buffer + adamw smoke tests green; `cargo test -p ml --features cuda --lib replay` 4 of 4 PER smoke tests green; `cargo test -p ml --features cuda --lib` IMPROVES the Wave 4.2 baseline — 947 pass / 12 fail (was 946 pass / 13 fail; the previously-failing `ensemble::adapters::dqn::tests::test_dqn_checkpoint_round_trip` now passes, likely because the redundant pre-3.5.5.b `scatter_insert_f32` for priorities was racing with the immediate `per_insert_pa` overwrite under particular timing conditions; the merged single-kernel write closes that race). diff --git a/docs/isv-slots.md b/docs/isv-slots.md index f651df1dd..999dc6f68 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -430,16 +430,19 @@ values are omitted from the emit (recomputing host-side would violate ## SP15 — Trader Discipline Recovery (Pre-Phase P.2 reservation, 2026-05-06) SP15 Pre-Phase Task P.2 extends `ISV_TOTAL_DIM` from 396 → 443 by reserving 46 -contiguous slots at `[397..443)`. The allocation pre-claims disjoint ranges -across 3 sub-worktrees (Approach B parallel-dispatch) so Phase 0/1/2A can -land producers and consumers independently without index collisions. +contiguous slots at `[397..443)`. SP15 Phase 1.3.b-followup (2026-05-07) +appends one more slot at `[443..444)` (DD_PERSISTENCE_MAX_INDEX) for the +per-env DD redesign, bringing `ISV_TOTAL_DIM` to 444 and `SP15_SLOT_END` to 444 +(47 SP15 slots total). The allocation pre-claims disjoint ranges across 3 +sub-worktrees (Approach B parallel-dispatch) so Phase 0/1/2A can land +producers and consumers independently without index collisions. -**Allocation map (per spec §4.3):** +**Allocation map (per spec §4.3 + Phase 1.3.b-followup audit doc entry):** | Range | Phase | Purpose | |-------|-------|---------| | `[397..401)` | Phase 0.B | EGF retune anchors (SCHMITT_HI/LO + VAR_AUX/Q_REF) | -| `[401..407)` | Phase 1.3 | Drawdown reporting (current/max/recovery_bars/persistence/calmar/dd_pct) | +| `[401..407)` | Phase 1.3 + 1.3.b-followup | Drawdown reporting (current/max/recovery_bars/persistence/calmar/dd_pct). Phase 1.3.b-followup (2026-05-07): now mean-aggregated across envs by `dd_state_reduce_kernel` from the per-env tile `sp15_dd_state_per_env[alloc_episodes * 6]`; the per-(env, slot) DD context is read directly by `compute_sp15_final_reward_kernel` from the tile | | `[407..409)` | Phase 1.2 | Cost kernel (OFI_IMPACT_LAMBDA + COST_PER_BAR_AVG) | | `[409..417)` | Phase 1.4 | 8 counterfactual baselines (buyhold/hold-only/random-dir-kelly/naive-momentum/aux-only/mag-quarter-fixed/trail-only/naive-reversion sharpe) | | `[417..420)` | Phase 3.1 | r_quality + r_discipline split (ALPHA_SPLIT + grad-norm quality/discipline) | @@ -451,6 +454,7 @@ land producers and consumers independently without index collisions. | `[436..439)` | Phase 3.5.4 | Plasticity injection (FIRED_THIS_FOLD + PERSISTENCE_THRESHOLD + WARM_BARS_REMAINING) | | `[439..441)` | Phase 3.5.5 | Recovery curriculum in PER (DD_TRAJECTORY_DECREASING + RECOVERY_OVERSAMPLE_WEIGHT) | | `[441..443)` | Phase 3.5 deferred | DD_TRAJECTORY_FLOOR + MEDIAN_STREAK_LENGTH ISV anchors | +| `[443..444)` | Phase 1.3.b-followup | DD_PERSISTENCE_MAX (max-aggregate of per-env DD_PERSISTENCE; consumed by `plasticity_injection_kernel` so the global advantage-head reset fires when ANY env exceeds the threshold — one set of advantage weights → global firing semantics → max-aggregate is the only correct rule) | **Files touched (P.2 atomic):** - `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` (new) — 46 `pub const *_INDEX` constants + SP15_SLOT_BASE/END/COUNT + 2 regression tests (`all_sp15_slots_fit_within_isv_total_dim`, `sp15_slot_layout_locked`)