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

feat(sp15-p1.3.b-followup): per-env DD redesign — Path A env-0-canonical → Path B per-env tile + reduction

Browse Source 
Closes the Phase 1.3.b deferred per-env redesign per
feedback_no_partial_refactor. Path A (env-0-canonical, commit 132609724)
made every downstream consumer read env-0's DD context; production envs
each have their own DD trajectory, so when env-3 was at 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. Path B threads
each env's actual DD context through the reward shaping atomically:

(1) dd_state_kernel reshape — grid [n_envs, 1, 1], one thread per env,
    writes 6 scalars per env to a new per-env tile dd_state_per_env
    [n_envs * 6]. No more ISV scalar writes.
(2) NEW dd_state_reduce_kernel — single-block tree-reduce (no
    atomicAdd; BLOCK=256 with strided initial pass for n_envs up to
    32768 on H100). Mean-aggregates per-env tile → 6 scalar ISV slots
    [401..407) for HEALTH_DIAG diagnostic. Max-aggregates DD_PERSISTENCE
    → new slot DD_PERSISTENCE_MAX_INDEX=443 for plasticity injection
    trigger (one shared advantage-head ⇒ global firing ⇒ max-aggregate
    is the only correct rule). ISV_TOTAL_DIM 443→444; SP15_SLOT_END
    443→444; SP15_SLOT_COUNT 46→47.
(3) compute_sp15_final_reward_kernel migration — per-(i,t) per-env DD
    lookup via env_id = (idx % (N*L)) / L. Helpers sp15_dd_asymmetric_reward
    and sp15_dd_penalty migrated to take dd_pct + dd_current as scalar
    parameters (the kernel reads from the per-env tile, threads scalars
    in). On-policy + CF threads at the same (i,t) read the SAME tile
    entry (one DD trajectory per env, shared across slot kinds).
(4) plasticity_injection_kernel migration — persistence read switched
    from ISV[404] (mean) to ISV[443] (max). One set of advantage
    weights ⇒ ANY env exceeding the threshold should arm the gate.
(5) 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. Reset to zero via the
    sp15_dd_state_per_env registry-arm dispatch.
(6) Per-step launch order: dd_state → dd_state_reduce →
    alpha_split_producer → final_reward, all on the same stream
    (CUDA serialises producer→consumer without explicit event sync).
(7) HEALTH_DIAG semantic shift (documented breaking change): slots
    401-406 now report cross-env mean, not env-0 value. For n_envs=1
    smoke configs the mean equals env-0's value (bit-stable migration).
(8) Layout fingerprint break: added markers DD_PERSISTENCE_MAX=443;
    ISV_TOTAL_DIM=444; DD_STATE_PER_ENV=sp15_phase_1_3_b_followup.
    Pre-followup checkpoints will not load (greenfield OK per spec Q1).
(9) 6 oracle tests migrated + 1 NEW behavioral test
    `dd_state_per_env_diverge_independently` — two-env config (env-0
    in recovery, env-1 deepening) verifies independent trajectories
    + mean-aggregate + max-aggregate semantics.

Phase 1.3.b-followup-B (separate split): dd_trajectory_decreasing_kernel
+ per_insert_pa migration is NOT in scope. The current Wave 4.3 reads
ISV[439] at insert-batch time (epoch end) — applied uniformly to ALL
inserted transitions (a known pre-existing limitation). Proper fix
requires per-(env, t) trajectory buffer [N*L] + env_id-aware lookup
in per_insert_pa via env_id = (j % (N*L)) / L. That's a different
contract change; splitting preserves no-partial-refactor within each
migration.

Atomic per feedback_no_partial_refactor: all 5 consumers of single-env-
canonical DD slots (final_reward kernel + plasticity kernel +
HEALTH_DIAG diagnostic + 6 oracle tests + new behavioral test) migrate
to per-env tile lookup in this commit; the new DD_PERSISTENCE_MAX
ISV slot lands with its sole consumer (plasticity).

Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean;
cargo check -p ml --features cuda --tests clean;
CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests
--features cuda -- --ignored: 17/17 oracle tests green (2 dd_state
incl. new per-env behavioral + 5 plasticity + 3 dd_trajectory + 6
final_reward + 1 per_sampler); cargo test -p ml --features cuda --lib:
947 pass / 12 fail HOLDS the 483cef454 baseline (no new regressions).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-07 09:35:58 +02:00
parent 483cef454c
commit 5b394f1035
14 changed files with 1004 additions and 227 deletions
Download Patch File Download Diff File Expand all files Collapse all files

44
crates/ml/build.rs
View File

@@ -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,

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

@@ -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 <cuda_runtime.h>
#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);

166
crates/ml/src/cuda_pipeline/dd_state_kernel.cu
View File

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

153
crates/ml/src/cuda_pipeline/dd_state_reduce_kernel.cu Normal file
View File

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

204
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -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<CudaStream>,
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<CudaStream>,
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<CudaStream>,
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<f32>,

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

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

22
crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu
View File

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

25
crates/ml/src/cuda_pipeline/sp15_isv_slots.rs
View File

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

32
crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh
View File

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

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

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

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

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

304
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

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