diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 8900e2848..c7097b07b 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1054,6 +1054,38 @@ fn main() { // so the follow-up commit is purely additive on the kernel // body, not on the launcher contract. "plasticity_injection_kernel.cu", + // SP15 Phase 3.5.5 (2026-05-06): recovery curriculum in PER — + // per-step DD_TRAJECTORY_DECREASING proxy. Replaces non-existent + // episode-level metadata with a per-bar boolean computed from + // the dd_pct trajectory (`dd_pct(t) < dd_pct(t-1) AND + // dd_pct(t-1) > DD_TRAJECTORY_FLOOR`). True when transition is + // part of a recovery (DD shrinking from a non-trivial + // drawdown). Single source file with one `extern "C" __global__` + // symbol (`dd_trajectory_decreasing_kernel`) → one cubin → one + // launcher (`launch_sp15_dd_trajectory_decreasing`). Mirrors + // the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 + // `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2 + // `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu` + // / 3.5.4 `plasticity_injection_kernel.cu` + // single-kernel-per-cubin pattern. Reads ISV[DD_PCT_INDEX=406] + + // ISV[DD_TRAJECTORY_FLOOR_INDEX=441]; writes + // ISV[DD_TRAJECTORY_DECREASING_INDEX=439]; read-modify-writes a + // separate mapped-pinned `prev_dd_pct` scratch buffer + // (`sp15_dd_trajectory_prev_dd`, [1] f32, persistent + // previous-bar dd_pct between kernel calls — mirrors Phase + // 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV + // mapped-pinned pattern). Per spec §9.2 (3.5.5) + // post-amendment-2 fix: PER sampling consumer + // (`sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × + // DD_TRAJECTORY_DECREASING)`) AND the ISV-driven + // 25th-percentile DD_TRAJECTORY_FLOOR producer are deferred to + // follow-up atomic commits per `feedback_no_partial_refactor.md` + // (matching Phase 1.1-1.5 + 3.1-3.5.4 precedent — kernel + + // launcher verify in isolation first via the GPU oracle tests + // below). Phase 3.5.5 lands kernel + launcher + 3 ISV anchor + // seeds + 1 scratch buffer + 4 registry entries + 4 dispatch + // arms only. + "dd_trajectory_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu b/crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu new file mode 100644 index 000000000..39d1f9e2a --- /dev/null +++ b/crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu @@ -0,0 +1,109 @@ +// crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu +// +// SP15 Phase 3.5.5 — recovery curriculum: per-step DD_TRAJECTORY_DECREASING proxy. +// +// Per spec §9.2 (3.5.5) post-amendment-2 fix: replaces non-existent +// episode-level metadata with a per-step signal computed from the +// per-bar dd_pct trajectory. +// +// trajectory(t) = 1.0 iff dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > floor +// trajectory(t) = 0.0 otherwise +// +// True when the transition is part of a recovery: drawdown shrinking +// from a non-trivial level (above DD_TRAJECTORY_FLOOR). The non-trivial +// gate keeps "0.001 → 0.0005" PnL flutter from inflating the recovery +// signal — only meaningful drawdowns count. +// +// PER sampling consumer (DEFERRED follow-up): +// +// sampling_weight = base_priority × (1.0 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING_t) +// +// Recovery transitions get amplified gradient signal so the agent +// learns recovery dynamics over the typical "average-DD" Bellman noise. +// +// DD_TRAJECTORY_FLOOR (slot 441) is ISV-driven (25th percentile of +// running dd_pct distribution) — producer kernel is a documented +// Phase 3.5.5 follow-up sub-task per `feedback_isv_for_adaptive_bounds.md`. +// The 0.02 sentinel from the constructor / fold-reset arm is used until +// the producer lands. RECOVERY_OVERSAMPLE_WEIGHT (slot 440) is also +// ISV-driven (proportional to current dd_pct — more DD now → higher +// weight for recovery samples) but the producer is part of the same +// follow-up sub-task; the 2.0 sentinel is used until then. +// +// Persistent state: `prev_dd_pct` is a 1-element f32 mapped-pinned +// scratch buffer on the trainer struct (`sp15_dd_trajectory_prev_dd`), +// initialised to 0.0 in the constructor and reset to 0.0 at fold +// boundary via `host_slice_mut().fill(0.0)` (mirrors the Phase 3.1 +// `sp15_alpha_warm_count` and Phase 3.5.3 +// `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned reset +// pattern). The cold-start sentinel 0.0 means the very first call +// after construction / fold reset has prev=0.0 < floor=0.02, so the +// floor gate evaluates to false and trajectory=0 — bootstrapping the +// state without false-firing on a synthetic large-step transition +// (current=0.10 vs prev=0 would otherwise satisfy currentfloor only if floor<0; the dual-condition `<` and `>` keep the +// gate honest at the boundary). +// +// 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[406] DD_PCT (read) +// ISV[441] DD_TRAJECTORY_FLOOR (read) +// ISV[439] DD_TRAJECTORY_DECREASING (write) +// +// Hard rules: +// +// * `feedback_no_atomicadd` — single-thread, single-block, pure +// scalar arithmetic; no reductions, no `atomicAdd`. +// * `feedback_no_partial_refactor` — kernel + launcher land first; +// PER sampling consumer migration (sampling_weight = base × (1 + +// RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)) AND the +// 25th-percentile DD_TRAJECTORY_FLOOR producer kernel are deferred +// to follow-up atomic commits per the established Phase +// 1.1-1.5 / 3.1-3.5.4 precedent. +// * `feedback_no_stubs` — kernel returns the actual updated state +// (per-step trajectory boolean + persistent prev_dd_pct advance); +// the three oracle tests verify fires-during-recovery, +// no-fire-when-dd-increasing, and no-fire-when-prev-below-floor +// behaviors. +// * `feedback_isv_for_adaptive_bounds` — DD_TRAJECTORY_FLOOR (slot +// 441) is ISV-tracked (the 0.02 sentinel is a structural cold-start +// anchor); the producer driving it from the rolling 25th percentile +// of the dd_pct distribution is deferred but the kernel reads the +// slot rather than a hardcoded literal so the follow-up swap is a +// no-op at the consumer. +// +// Single-thread, single-block (mirrors `cooldown_kernel` / +// `plasticity_injection_kernel` / `dd_asymmetric_reward_kernel` / +// `hold_floor_kernel` / `regret_signal_kernel` / `dd_penalty_kernel` — +// no reduction). + +extern "C" __global__ void dd_trajectory_decreasing_kernel( + float* __restrict__ isv, + float* __restrict__ prev_dd_pct /* persistent state, [1] f32 */ +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float dd_now = isv[/*DD_PCT_INDEX*/ 406]; + const float dd_prev = *prev_dd_pct; + const float floor = isv[/*DD_TRAJECTORY_FLOOR_INDEX*/ 441]; + + // Recovery proxy: dd shrinking AND prev was above the non-trivial + // floor. Both conditions required — `dd_now < dd_prev` alone would + // count "0.001 → 0.0005" PnL flutter as recovery; `dd_prev > floor` + // alone would count "0.10 → 0.15" (worse drawdown) as recovery. + const float trajectory = + (dd_now < dd_prev && dd_prev > floor) ? 1.0f : 0.0f; + + isv[/*DD_TRAJECTORY_DECREASING_INDEX*/ 439] = trajectory; + + // Advance persistent state for next call. The host writes 0.0 at + // construction / fold boundary (allowed under + // `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write + // rule, NOT a host-side compute); every subsequent advance is + // GPU-side. + *prev_dd_pct = dd_now; + + __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 384e26699..9fee6478d 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1801,6 +1801,103 @@ pub fn launch_sp15_plasticity_injection( Ok(()) } +/// SP15 Phase 3.5.5 (2026-05-06): cubin slot for the recovery-curriculum +/// per-step DD_TRAJECTORY_DECREASING proxy kernel +/// (`dd_trajectory_decreasing_kernel`). +/// +/// Per spec §9.2 (3.5.5) post-amendment-2 fix: replaces non-existent +/// episode-level metadata with a per-bar boolean computed from the +/// dd_pct trajectory: +/// +/// trajectory(t) = 1.0 iff dd_pct(t) < dd_pct(t-1) +/// AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR +/// +/// True when the transition is part of a recovery: drawdown shrinking +/// from a non-trivial level (above the floor). +/// +/// PER sampling consumer (DEFERRED follow-up): +/// +/// sampling_weight = base_priority × (1.0 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING_t) +/// +/// DD_TRAJECTORY_FLOOR (slot 441) is ISV-driven (25th percentile of +/// running dd_pct distribution) — producer kernel is a documented +/// Phase 3.5.5 follow-up sub-task per `feedback_isv_for_adaptive_bounds.md`. +/// The 0.02 sentinel from the constructor / fold-reset arm is used until +/// the producer lands. RECOVERY_OVERSAMPLE_WEIGHT (slot 440) is also +/// ISV-driven (proportional to current dd_pct) — same follow-up +/// sub-task; the 2.0 sentinel is used until then. +/// +/// Mirrors the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 +/// `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2 +/// `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu` / 3.5.4 +/// `plasticity_injection_kernel.cu` single-kernel-per-cubin pattern. +/// +/// Phase 3.5.5 lands kernel + launcher + 3 ISV constructor seeds + 1 +/// new `sp15_dd_trajectory_prev_dd` scratch buffer + 4 state-reset +/// registry entries + 4 dispatch arms only; the PER sampling consumer +/// migration AND the ISV-driven 25th-percentile DD_TRAJECTORY_FLOOR +/// producer are deferred to follow-up atomic commits per +/// `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + 3.1-3.5.4 +/// precedent — kernel + launcher verify in isolation first via the +/// GPU oracle tests below). +pub static SP15_DD_TRAJECTORY_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/dd_trajectory_kernel.cubin")); + +/// SP15 Phase 3.5.5 (2026-05-06): launcher for the recovery-curriculum +/// per-step DD_TRAJECTORY_DECREASING proxy kernel. Free function +/// (matches `launch_sp15_dd_penalty` / `launch_sp15_regret_signal` / +/// `launch_sp15_hold_floor` / `launch_sp15_dd_asymmetric_reward` / +/// `launch_sp15_cooldown` / `launch_sp15_plasticity_injection` +/// precedent — single kernel, single cubin, single launcher) so +/// unit/oracle tests can drive the kernel directly without the trainer +/// struct; production callers invoke the same launcher. +/// +/// `isv` MUST be the ISV bus (≥ `SP15_SLOT_END=443` f32 slots — slot +/// 406 is read for DD_PCT; slot 441 is read for DD_TRAJECTORY_FLOOR; +/// slot 439 is written for DD_TRAJECTORY_DECREASING). +/// `prev_dd_pct` MUST point to at least 1 writable f32 slot (the +/// `sp15_dd_trajectory_prev_dd` mapped-pinned scratch buffer in +/// production; a fresh 1-element `MappedF32Buffer` in oracle tests). +/// The kernel read-modify-writes the slot: reads to compute the +/// trajectory direction comparison, then advances to the current +/// dd_pct so the next call sees the right "previous bar" value. +/// Single thread, single block (mirrors `cooldown_kernel` / +/// `plasticity_injection_kernel` / `dd_asymmetric_reward_kernel` / +/// `hold_floor_kernel` / `regret_signal_kernel` / `dd_penalty_kernel` — +/// no reduction). +pub fn launch_sp15_dd_trajectory_decreasing( + stream: &Arc, + isv: cudarc::driver::sys::CUdeviceptr, + prev_dd_pct: cudarc::driver::sys::CUdeviceptr, +) -> Result<(), MLError> { + let module = stream + .context() + .load_cubin(SP15_DD_TRAJECTORY_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!( + "load sp15_dd_trajectory cubin: {e}" + )))?; + let kernel = module + .load_function("dd_trajectory_decreasing_kernel") + .map_err(|e| MLError::ModelError(format!( + "load dd_trajectory_decreasing_kernel function: {e}" + )))?; + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv) + .arg(&prev_dd_pct) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "launch dd_trajectory_decreasing_kernel: {e}" + )))?; + } + Ok(()) +} + /// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup + /// update kernels sharing one cubin. Lookup reads /// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update @@ -4628,6 +4725,30 @@ pub struct GpuDqnTrainer { /// page via `read_all()` only in tests. pub(crate) sp15_cooldown_consecutive_losses: super::mapped_pinned::MappedF32Buffer, // [1] + /// SP15 Phase 3.5.5 (2026-05-06): previous-bar dd_pct scratch buffer + /// for the recovery-curriculum DD_TRAJECTORY_DECREASING per-step + /// proxy kernel (per spec §9.2 (3.5.5) post-amendment-2 fix). + /// Single-element mapped-pinned f32 buffer initialised to 0.0 in + /// the constructor and reset to 0.0 at fold boundary by the + /// `sp15_dd_trajectory_prev_dd` registry-arm dispatch (which calls + /// `host_slice_mut().fill(0.0)` mirroring the Phase 3.1 + /// `sp15_alpha_warm_count` and Phase 3.5.3 + /// `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned reset + /// pattern). Read-modify-written by `dd_trajectory_decreasing_kernel`: + /// reads the previous bar's dd_pct to compute the trajectory + /// direction (`dd_now < dd_prev AND dd_prev > floor`), then writes + /// the current `dd_pct` so the next call sees the correct + /// "previous" value. The cold-start sentinel 0.0 means the very + /// first call after construction / fold reset has prev=0.0 < the + /// 0.02 floor sentinel, so the floor gate evaluates to false and + /// trajectory=0 — bootstrapping the state without false-firing on + /// a synthetic large-step transition. The host writes 0.0 at + /// construction / fold reset (allowed under + /// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write + /// rule, NOT a host-side compute) and reads it through the mapped + /// page via `read_all()` only in tests. + pub(crate) sp15_dd_trajectory_prev_dd: super::mapped_pinned::MappedF32Buffer, // [1] + // ── 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, @@ -19311,6 +19432,31 @@ impl GpuDqnTrainer { )))?; sp15_cooldown_consecutive_losses.write_from_slice(&[0.0_f32]); + // SP15 Phase 3.5.5 (2026-05-06): recovery-curriculum DD trajectory + // previous-bar dd_pct scratch buffer (per spec §9.2 (3.5.5) + // post-amendment-2 fix). Single-element mapped-pinned f32 buffer + // initialised to 0.0 by `MappedF32Buffer::new`'s zero-init + // contract — the host writes 0.0 at construction and again at + // fold boundary via the registry-arm dispatch (allowed under + // `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write + // rule, NOT a host-side compute). The `dd_trajectory_decreasing + // _kernel` read-modify-writes this buffer: reads the previous + // bar's dd_pct to compute the trajectory direction, then advances + // it to the current `dd_pct` so the next call sees the correct + // "previous" value. The cold-start sentinel 0.0 means the very + // first call has prev=0.0 < the 0.02 floor sentinel, so the + // floor gate evaluates to false and trajectory=0 — bootstrapping + // the state without false-firing on a synthetic large-step + // transition. Production wire-up of the per-bar launches AND the + // PER sampling consumer migration are deferred to follow-up + // commits per `feedback_no_partial_refactor.md`. + let sp15_dd_trajectory_prev_dd = + unsafe { super::mapped_pinned::MappedF32Buffer::new(1) } + .map_err(|e| MLError::ModelError(format!( + "sp15_dd_trajectory_prev_dd alloc (1 f32): {e}" + )))?; + sp15_dd_trajectory_prev_dd.write_from_slice(&[0.0_f32]); + // ── Per-branch Q-gap EMA (trajectory backtracking state) ──────── let per_branch_q_gap_ema_buf = stream.alloc_zeros::(4) // [4] .map_err(|e| MLError::ModelError(format!("alloc per_branch_q_gap_ema_buf: {e}")))?; @@ -21108,6 +21254,60 @@ impl GpuDqnTrainer { *sig_ptr.add(PLASTICITY_PERSISTENCE_THRESHOLD_INDEX) = 100.0_f32; *sig_ptr.add(PLASTICITY_WARM_BARS_REMAINING_INDEX) = 0.0_f32; + // SP15 Phase 3.5.5 (2026-05-06): recovery curriculum in + // PER — per-step DD_TRAJECTORY_DECREASING proxy anchors + // (per spec §9.2 (3.5.5) post-amendment-2 fix). + // Replaces non-existent episode-level metadata with a + // per-bar boolean computed from the dd_pct trajectory: + // `trajectory(t) = 1.0 iff dd_pct(t) < dd_pct(t-1) AND + // dd_pct(t-1) > DD_TRAJECTORY_FLOOR`. PER sampling + // consumer (deferred follow-up): `sampling_weight = + // base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × + // DD_TRAJECTORY_DECREASING)`. + // + // DD_TRAJECTORY_DECREASING = 0.0 (cold-start: no + // observation yet — + // the kernel's first + // call after + // construction / fold + // reset has prev=0.0 + // < floor=0.02 so the + // floor gate + // evaluates to false + // and trajectory + // stays 0) + // RECOVERY_OVERSAMPLE_WEIGHT = 2.0 (initial sentinel — + // ISV-driven from + // current dd_pct in + // follow-up sub-task + // per + // `feedback_isv_for_adaptive_bounds.md`; + // more DD now → + // higher weight for + // recovery samples; + // the 2.0 anchor + // re-engages on each + // new fold) + // DD_TRAJECTORY_FLOOR = 0.02 (initial sentinel — + // ISV-driven 25th + // percentile of running + // dd_pct distribution in + // follow-up sub-task per + // `feedback_isv_for_adaptive_bounds.md`, + // replacing the hardcoded + // 0.02 floor from the + // original spec draft; + // the 0.02 anchor + // re-engages on each new + // fold) + use crate::cuda_pipeline::sp15_isv_slots::{ + DD_TRAJECTORY_DECREASING_INDEX, DD_TRAJECTORY_FLOOR_INDEX, + RECOVERY_OVERSAMPLE_WEIGHT_INDEX, + }; + *sig_ptr.add(DD_TRAJECTORY_DECREASING_INDEX) = 0.0_f32; + *sig_ptr.add(RECOVERY_OVERSAMPLE_WEIGHT_INDEX) = 2.0_f32; + *sig_ptr.add(DD_TRAJECTORY_FLOOR_INDEX) = 0.02_f32; + // Layout fingerprint (ISV[58..60)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so @@ -21850,6 +22050,7 @@ impl GpuDqnTrainer { sel_t_dev_ptr, sp15_alpha_warm_count, sp15_cooldown_consecutive_losses, + sp15_dd_trajectory_prev_dd, per_branch_q_gap_ema_buf, last_per_branch_q_gaps: [0.0; 4], q_mag_means_cached: [0.0_f32; 3], diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 855cdf581..0aa8684e4 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1330,6 +1330,58 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[PLASTICITY_WARM_BARS_REMAINING_INDEX=438] — SP15 Phase 3.5.4 (2026-05-06) plasticity injection warm-up counter (number of bars remaining where action-selection forces Hold post-fire via the OR-gate `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)`). Set to M_warm (default 200) by `plasticity_injection_kernel` on fire; decremented per bar. FoldReset sentinel 0.0 so the new fold starts with no active warm-up — leftover state from the previous fold would force Hold during the new fold's warmup, contaminating the cold-start (mirrors `sp15_cooldown_bars_remaining` semantics). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.4).", }, + // SP15 Phase 3.5.5 (2026-05-06) — recovery curriculum in + // PER: per-step DD_TRAJECTORY_DECREASING proxy (per spec + // §9.2 (3.5.5) post-amendment-2 fix). Replaces non-existent + // episode-level metadata with a per-bar boolean from the + // dd_pct trajectory. PER sampling consumer (deferred + // follow-up): `sampling_weight = base × (1 + + // RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)`. + // DD_TRAJECTORY_FLOOR is ISV-driven (25th percentile of + // running dd_pct distribution; producer is a follow-up + // sub-task per `feedback_isv_for_adaptive_bounds.md`); + // RECOVERY_OVERSAMPLE_WEIGHT is also ISV-driven (proportional + // to current dd_pct — more DD now → higher weight; producer + // is the same follow-up sub-task). DD_TRAJECTORY_DECREASING + // is the per-bar output — FoldReset sentinel 0 so the new + // fold starts with no carry-over signal. Both anchors + // (FLOOR=0.02, OVERSAMPLE=2.0) rewrite at fold boundary so + // the cold-start absorbers re-engage on each new fold per + // phase precedent. + RegistryEntry { + name: "sp15_dd_trajectory_decreasing", + category: ResetCategory::FoldReset, + description: "ISV[DD_TRAJECTORY_DECREASING_INDEX=439] — SP15 Phase 3.5.5 (2026-05-06) recovery-curriculum per-step proxy. Per-bar boolean (0.0/1.0) computed by `dd_trajectory_decreasing_kernel`: 1.0 iff `dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR`. Read by the deferred PER sampling consumer (`sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)`). FoldReset sentinel 0 so the new fold starts with no carry-over signal — leftover '1.0' from the previous fold's last bar would over-weight the new fold's first replay-buffer entry. The companion `sp15_dd_trajectory_prev_dd` scratch buffer reset in the same arm cluster ensures the kernel's persistent prev_dd_pct cannot trigger a false recovery on the new fold's first call (mirrors the Phase 3.5.3 `sp15_cooldown_consecutive_losses` reset pattern). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.5).", + }, + RegistryEntry { + name: "sp15_recovery_oversample_weight", + category: ResetCategory::FoldReset, + description: "ISV[RECOVERY_OVERSAMPLE_WEIGHT_INDEX=440] — SP15 Phase 3.5.5 (2026-05-06) PER sampling oversample weight for recovery transitions. FoldReset rewrites the structural cold-start anchor 2.0 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from a follow-up producer kernel reading the current dd_pct (more DD now → higher weight for recovery samples so the agent learns recovery dynamics over typical 'average-DD' Bellman noise); the 2.0 anchor re-engages on each new fold so the cold-start absorber holds the weight at a finite scale until the producer accumulates observations. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.5).", + }, + RegistryEntry { + name: "sp15_dd_trajectory_floor", + category: ResetCategory::FoldReset, + description: "ISV[DD_TRAJECTORY_FLOOR_INDEX=441] — SP15 Phase 3.5.5 (2026-05-06) non-trivial-drawdown floor for the DD_TRAJECTORY_DECREASING gate (`dd_prev > floor` term). FoldReset rewrites the structural cold-start anchor 0.02 per `feedback_isv_for_adaptive_bounds.md` — replaces the hardcoded 0.02 floor from the original spec draft; the runtime value will be signal-driven from a follow-up producer reading the rolling 25th percentile of the dd_pct distribution (ensures the floor tracks the typical small-drawdown threshold instead of being decoupled from the actual dd_pct scale). The 0.02 anchor re-engages on each new fold so the cold-start absorber holds the floor at a finite scale until the producer accumulates observations. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.5).", + }, + // SP15 Phase 3.5.5 (2026-05-06) — recovery-curriculum DD + // trajectory previous-bar dd_pct scratch buffer (non-ISV + // mapped-pinned). Mirrors the Phase 3.1 `sp15_alpha_warm_count` + // and Phase 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV + // mapped-pinned reset pattern. Reset to 0.0 via + // `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) — the new + // fold's first dd_trajectory_decreasing_kernel launch must + // see prev=0.0 so the floor gate (prev > 0.02) evaluates + // to false and trajectory=0 on the bootstrap call, + // preventing a partially-accumulated dd_pct from the + // previous fold from triggering a false recovery on the + // new fold's first call. + RegistryEntry { + name: "sp15_dd_trajectory_prev_dd", + category: ResetCategory::FoldReset, + description: "GpuDqnTrainer.sp15_dd_trajectory_prev_dd [1] f32 mapped-pinned scratch buffer — SP15 Phase 3.5.5 (2026-05-06) previous-bar dd_pct for the recovery-curriculum DD_TRAJECTORY_DECREASING per-step proxy kernel (per spec §9.2 (3.5.5) post-amendment-2 fix). Initialised to 0.0 in the constructor; read-modify-written by `dd_trajectory_decreasing_kernel` (reads previous bar's dd_pct, writes current dd_pct). FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)` (mirrors the Phase 3.1 `sp15_alpha_warm_count` and Phase 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned reset pattern) — the new fold's first kernel launch must see prev=0.0 so the floor gate (`dd_prev > DD_TRAJECTORY_FLOOR=0.02` sentinel) evaluates to false and trajectory=0 on the bootstrap call, preventing a partially-accumulated dd_pct from the previous fold from triggering a false recovery on the new fold's first call. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.5).", + }, // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) + // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index d39029ca7..6903748cf 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -8412,6 +8412,50 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(PLASTICITY_WARM_BARS_REMAINING_INDEX, 0.0); } } + // SP15 Phase 3.5.5 (2026-05-06): recovery curriculum in PER — + // per-step DD_TRAJECTORY_DECREASING proxy (per spec §9.2 + // (3.5.5) post-amendment-2 fix). DD_TRAJECTORY_DECREASING is + // the per-bar output (FoldReset sentinel 0). RECOVERY_ + // OVERSAMPLE_WEIGHT and DD_TRAJECTORY_FLOOR are structural + // cold-start anchors; ISV-driven producers from current + // dd_pct (RECOVERY_OVERSAMPLE_WEIGHT) and 25th percentile of + // running dd_pct distribution (DD_TRAJECTORY_FLOOR) are + // follow-up sub-tasks per `feedback_isv_for_adaptive_bounds.md`. + // The companion `sp15_dd_trajectory_prev_dd` scratch buffer + // reset in the same arm cluster ensures the kernel's + // persistent prev_dd_pct cannot trigger a false recovery on + // the new fold's first call. + "sp15_dd_trajectory_decreasing" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::DD_TRAJECTORY_DECREASING_INDEX; + fused.trainer().write_isv_signal_at(DD_TRAJECTORY_DECREASING_INDEX, 0.0); + } + } + "sp15_recovery_oversample_weight" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::RECOVERY_OVERSAMPLE_WEIGHT_INDEX; + fused.trainer().write_isv_signal_at(RECOVERY_OVERSAMPLE_WEIGHT_INDEX, 2.0); + } + } + "sp15_dd_trajectory_floor" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::DD_TRAJECTORY_FLOOR_INDEX; + fused.trainer().write_isv_signal_at(DD_TRAJECTORY_FLOOR_INDEX, 0.02); + } + } + // sp15_dd_trajectory_prev_dd is a non-ISV mapped-pinned + // scratch buffer on the trainer struct. Reset to 0.0 via + // `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 Phase 3.1 `sp15_alpha_warm_count` and Phase 3.5.3 + // `sp15_cooldown_consecutive_losses` pattern of resetting + // a non-ISV mapped-pinned buffer at fold boundary. + "sp15_dd_trajectory_prev_dd" => { + if let Some(ref mut fused) = self.fused_ctx { + fused.trainer_mut().sp15_dd_trajectory_prev_dd.host_slice_mut().fill(0.0); + } + } // SP11 Task A2 (2026-05-04): novelty visit-count hash table // reset arm — closes the A0 deferral per the registry entry's // docstring. 1M f32 slots, zero-filled via mapped-pinned host diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 2abb0d4bb..79c9fe306 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -1574,6 +1574,186 @@ mod gpu { isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX] ); } + + /// SP15 Phase 3.5.5 — DD_TRAJECTORY_DECREASING fires when dd_pct + /// decreases from a non-trivial drawdown (above the floor). Per + /// spec §9.2 (3.5.5) post-amendment-2 fix: per-bar boolean + /// computed from `dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > floor`. + /// Sequence: (1) bootstrap call sets prev=0.10 (cold-start + /// trajectory=0 because initial prev=0); (2) update dd_pct=0.08 + /// and re-launch — prev=0.10 > floor=0.02 AND now=0.08 < prev → + /// trajectory=1. Verifies BOTH the trajectory output AND the + /// persistent prev_dd_pct advance. + #[test] + #[ignore = "requires GPU"] + fn dd_trajectory_decreasing_fires_during_recovery() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_PCT_INDEX, DD_TRAJECTORY_DECREASING_INDEX, DD_TRAJECTORY_FLOOR_INDEX, + }; + + let stream = make_test_stream(); + + 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_PCT_INDEX] = 0.10; // current dd_pct + isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02; // floor sentinel + isv_buf.write_from_slice(&isv_init); + + // Persistent state: 1-element scratch buffer, init to 0.0. + // Mirrors the trainer-struct `sp15_dd_trajectory_prev_dd` + // mapped-pinned [1] f32 buffer. + let prev_dd_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for prev_dd_pct"); + prev_dd_buf.write_from_slice(&[0.0_f32]); + + // First call: prev=0, current=0.10 → cold-start bootstrap + // (prev=0 < floor=0.02 → floor gate false → trajectory=0); the + // kernel advances prev_dd_pct to 0.10. + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing( + &stream, + isv_buf.dev_ptr, + prev_dd_buf.dev_ptr, + ) + .expect("launch dd_trajectory_decreasing_kernel (bootstrap)"); + stream + .synchronize() + .expect("synchronize after dd_trajectory_decreasing_kernel launch (bootstrap)"); + + // Now reduce dd_pct: 0.10 → 0.08 (recovery from non-trivial DD + // above floor). Update only DD_PCT_INDEX so the floor + sentinel + // slots survive (write_from_slice rewrites the whole buffer). + let mut isv_now = isv_buf.read_all(); + isv_now[DD_PCT_INDEX] = 0.08; + isv_buf.write_from_slice(&isv_now); + + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing( + &stream, + isv_buf.dev_ptr, + prev_dd_buf.dev_ptr, + ) + .expect("launch dd_trajectory_decreasing_kernel (recovery)"); + stream + .synchronize() + .expect("synchronize after dd_trajectory_decreasing_kernel launch (recovery)"); + + let isv_host = isv_buf.read_all(); + let trajectory = isv_host[DD_TRAJECTORY_DECREASING_INDEX]; + assert!( + (trajectory - 1.0).abs() < 1e-9, + "DD_TRAJECTORY_DECREASING = {trajectory}, expected 1.0 (recovery: prev=0.10 > floor=0.02 AND now=0.08 < prev)" + ); + // Persistent state advance verification: prev_dd_pct now == current dd_pct = 0.08. + let prev_after = prev_dd_buf.read_all()[0]; + assert!( + (prev_after - 0.08).abs() < 1e-9, + "prev_dd_pct = {prev_after}, expected 0.08 (kernel advances persistent state to current dd_pct)" + ); + } + + /// SP15 Phase 3.5.5 — no fire when dd_pct increasing (DD getting + /// worse). Per spec §9.2 (3.5.5) post-amendment-2 fix: trajectory + /// requires `dd_pct(t) < dd_pct(t-1)` (strict). Setup: prev=0.05, + /// current=0.10 → DD getting worse → trajectory=0. Sentinel-write + /// semantics: ISV slot starts at 0.0 (default), kernel writes 0.0 + /// (the sentinel itself) — verified separately by the recovery + /// test that the kernel WILL write a non-zero output when the + /// gate fires, so the 0.0 here is a real write rather than a + /// skipped store. + #[test] + #[ignore = "requires GPU"] + fn dd_trajectory_no_fire_when_dd_increasing() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_PCT_INDEX, DD_TRAJECTORY_DECREASING_INDEX, DD_TRAJECTORY_FLOOR_INDEX, + }; + + let stream = make_test_stream(); + + 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_PCT_INDEX] = 0.10; // current dd_pct + isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02; + // Sentinel non-zero output buffer ensures the kernel actually + // overwrites with the 0.0 result rather than skipping the + // store (mirrors the dd_penalty_kernel sentinel-write check). + isv_init[DD_TRAJECTORY_DECREASING_INDEX] = 0.5; + isv_buf.write_from_slice(&isv_init); + + let prev_dd_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for prev_dd_pct"); + prev_dd_buf.write_from_slice(&[0.05_f32]); // prev was 0.05 + + // current dd_pct = 0.10 > prev = 0.05 → DD getting worse, no + // recovery fire (the strict `<` comparison rejects the > + // case). + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing( + &stream, + isv_buf.dev_ptr, + prev_dd_buf.dev_ptr, + ) + .expect("launch dd_trajectory_decreasing_kernel (dd-increasing)"); + stream + .synchronize() + .expect("synchronize after dd_trajectory_decreasing_kernel launch (dd-increasing)"); + + let isv_host = isv_buf.read_all(); + let trajectory = isv_host[DD_TRAJECTORY_DECREASING_INDEX]; + assert!( + trajectory.abs() < 1e-9, + "DD_TRAJECTORY_DECREASING = {trajectory}, expected 0 (DD increasing: prev=0.05 < now=0.10; sentinel 0.5 must be overwritten with the real 0.0 result, not skipped)" + ); + } + + /// SP15 Phase 3.5.5 — no fire when prev_dd below the floor (trivial + /// DD, not "in recovery"). Per spec §9.2 (3.5.5) post-amendment-2 + /// fix: trajectory requires `dd_pct(t-1) > DD_TRAJECTORY_FLOOR` + /// (strict). The non-trivial gate keeps "0.001 → 0.0005" PnL + /// flutter from inflating the recovery signal — only meaningful + /// drawdowns count. Setup: prev=0.01 (BELOW floor=0.02), current= + /// 0.005 → dd_pct decreasing but prev was below floor → trajectory=0. + #[test] + #[ignore = "requires GPU"] + fn dd_trajectory_no_fire_when_prev_below_floor() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_PCT_INDEX, DD_TRAJECTORY_DECREASING_INDEX, DD_TRAJECTORY_FLOOR_INDEX, + }; + + let stream = make_test_stream(); + + 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_PCT_INDEX] = 0.005; // current below floor + isv_init[DD_TRAJECTORY_FLOOR_INDEX] = 0.02; + // Sentinel non-zero output buffer (see recovery test for + // rationale). + isv_init[DD_TRAJECTORY_DECREASING_INDEX] = 0.7; + isv_buf.write_from_slice(&isv_init); + + let prev_dd_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for prev_dd_pct"); + prev_dd_buf.write_from_slice(&[0.01_f32]); // prev BELOW floor 0.02 + + // dd decreasing (0.01 → 0.005) but prev was below floor → not + // "in recovery" per the non-trivial-DD gate. + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_trajectory_decreasing( + &stream, + isv_buf.dev_ptr, + prev_dd_buf.dev_ptr, + ) + .expect("launch dd_trajectory_decreasing_kernel (prev-below-floor)"); + stream + .synchronize() + .expect("synchronize after dd_trajectory_decreasing_kernel launch (prev-below-floor)"); + + let isv_host = isv_buf.read_all(); + let trajectory = isv_host[DD_TRAJECTORY_DECREASING_INDEX]; + assert!( + trajectory.abs() < 1e-9, + "DD_TRAJECTORY_DECREASING = {trajectory}, expected 0 (prev=0.01 below floor=0.02; sentinel 0.7 must be overwritten with the real 0.0 result, not skipped)" + ); + } } /// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index b9aab9eb7..4bb7a4d0a 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 3.5.5 — recovery curriculum in PER: per-step DD_TRAJECTORY_DECREASING proxy (2026-05-06): per spec §9.2 (3.5.5) post-amendment-2 fix. Replaces non-existent episode-level metadata with a per-bar boolean computed from the dd_pct trajectory: `trajectory(t) = 1.0 iff dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR`. True when the transition is part of a recovery (DD shrinking from a non-trivial drawdown above the floor). New single-kernel cubin `dd_trajectory_kernel.cu` with one `extern "C" __global__` symbol (`dd_trajectory_decreasing_kernel`) → one cubin → one launcher (`launch_sp15_dd_trajectory_decreasing`) — mirrors the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2 `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu` / 3.5.4 `plasticity_injection_kernel.cu` single-kernel-per-cubin pattern. Reads ISV[DD_PCT_INDEX=406] (written by Task 1.3's `dd_state_kernel`) + ISV[DD_TRAJECTORY_FLOOR_INDEX=441]; writes ISV[DD_TRAJECTORY_DECREASING_INDEX=439]; read-modify-writes a separate mapped-pinned `prev_dd_pct` scratch buffer (`sp15_dd_trajectory_prev_dd`, [1] f32, persistent previous-bar dd_pct between kernel calls). Persistent state cold-start: sentinel 0.0 means the very first call has prev=0.0 < the 0.02 floor sentinel, so the floor gate (`dd_prev > floor`) evaluates to false and trajectory=0 — bootstrapping the state without false-firing on the synthetic `current=0.10 vs prev=0` transition. PER sampling consumer (deferred follow-up): `sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)` — recovery transitions get amplified gradient signal so the agent learns recovery dynamics over typical "average-DD" Bellman noise. 3 new ISV slots: `DD_TRAJECTORY_DECREASING_INDEX=439` (initial 0.0; per-bar output, FoldReset sentinel 0 so the new fold starts with no carry-over signal — leftover '1.0' from the previous fold's last bar would over-weight the new fold's first replay-buffer entry), `RECOVERY_OVERSAMPLE_WEIGHT_INDEX=440` (initial 2.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up producer kernel reading the current dd_pct (more DD now → higher weight); the 2.0 anchor re-engages on each new fold), `DD_TRAJECTORY_FLOOR_INDEX=441` (initial 0.02; structural cold-start anchor — runtime value signal-driven from a follow-up producer reading the rolling 25th percentile of running dd_pct distribution per `feedback_isv_for_adaptive_bounds.md`, replacing the hardcoded 0.02 floor from the original spec draft; the 0.02 anchor re-engages on each new fold). 1 new mapped-pinned scratch buffer on the trainer struct: `sp15_dd_trajectory_prev_dd` ([1] f32, mirrors the Phase 3.1 `sp15_alpha_warm_count` and Phase 3.5.3 `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned pattern) — initialised to 0.0 in the constructor; read-modify-written by the kernel (reads previous bar's dd_pct, writes current dd_pct so the next call sees the correct "previous" value); the host writes 0.0 at construction / fold reset (allowed under `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write rule, NOT a host-side compute) and reads it through the mapped page via `read_all()` only in tests. 4 new `state_reset_registry` entries (`sp15_dd_trajectory_decreasing` resets the per-bar output to 0; `sp15_recovery_oversample_weight` rewrites the 2.0 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_dd_trajectory_floor` rewrites the 0.02 anchor; `sp15_dd_trajectory_prev_dd` resets the non-ISV mapped-pinned scratch buffer to 0.0 via `host_slice_mut().fill(0.0)`). 4 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.10 recovery_after_streak (Phase 2C / Phase 3.5 paired) — green via 3.5.2 + 3.5.4 + 3.5.5 combined; this commit lands the recovery-curriculum primitive. **Phase 3.5.5 lands kernel + launcher + 3 ISV anchor seeds + 1 scratch buffer + 4 registry entries + 4 dispatch arms only**; the PER sampling consumer migration (`sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × DD_TRAJECTORY_DECREASING)`) AND the ISV-driven 25th-percentile DD_TRAJECTORY_FLOOR producer (and the current-dd_pct-driven RECOVERY_OVERSAMPLE_WEIGHT producer) are deferred to follow-up atomic commits per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + 3.1-3.5.4 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). 3 GPU oracle tests pass on RTX 3050 Ti: `dd_trajectory_decreasing_fires_during_recovery` validates the bootstrap + recovery sequence (first call prev=0 < floor=0.02 → trajectory=0 + prev advances to 0.10; second call prev=0.10 > floor AND now=0.08 < prev → trajectory=1 + prev advances to 0.08); `dd_trajectory_no_fire_when_dd_increasing` validates the strict `<` comparison (prev=0.05 < now=0.10 → trajectory=0; sentinel 0.5 must be overwritten with the real 0.0 result, not skipped); `dd_trajectory_no_fire_when_prev_below_floor` validates the non-trivial-DD gate (prev=0.01 < floor=0.02 → trajectory=0 even though dd_pct decreasing; sentinel 0.7 must be overwritten with 0.0). Touched: `crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `cooldown_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the writes; the kernel's ISV pointer is mutable `float*` because it writes the DD_TRAJECTORY_DECREASING slot, identical to `cooldown_kernel`'s mutable-ISV pattern; the kernel takes a separate `prev_dd_pct` device pointer so the persistent state lives outside the ISV bus per the Phase 3.5.3 `sp15_cooldown_consecutive_losses` pattern), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_TRAJECTORY_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_trajectory_decreasing` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_DD_TRAJECTORY_CUBIN.to_vec())` and resolving `get_function("dd_trajectory_decreasing_kernel")` — single-thread/single-block grid mirroring `launch_sp15_cooldown` precedent + 3 ISV constructor writes (DD_TRAJECTORY_DECREASING=0.0, RECOVERY_OVERSAMPLE_WEIGHT=2.0, DD_TRAJECTORY_FLOOR=0.02) appended to the SP15 Phase 3.5.4 plasticity anchor block + 1 new `sp15_dd_trajectory_prev_dd: MappedF32Buffer` field on the trainer struct + alloc + 0.0 init mirroring the Phase 3.5.3 `sp15_cooldown_consecutive_losses` precedent), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+4 `RegistryEntry` records — `sp15_dd_trajectory_decreasing` resets per-bar output to 0; `sp15_recovery_oversample_weight` rewrites 2.0 anchor; `sp15_dd_trajectory_floor` rewrites 0.02 anchor; `sp15_dd_trajectory_prev_dd` non-ISV mapped-pinned buffer reset to 0.0 via `host_slice_mut().fill(0.0)`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+4 dispatch arms for the 4 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+3 GPU oracle tests in `mod gpu`). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic + a few read-modify-writes of one ISV slot and the prev_dd_pct counter, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + scratch buffer + registry entries + dispatch arms land atomically; PER sampling consumer migration + 25th-percentile DD_TRAJECTORY_FLOOR producer + current-dd_pct-driven RECOVERY_OVERSAMPLE_WEIGHT producer all follow as their own atomic commits), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the trajectory boolean + advances the persistent prev_dd_pct; the three oracle tests verify recovery-fires, dd-increasing-no-fire, and prev-below-floor-no-fire behaviors with sentinel-overwrite semantics rather than skip-store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and prev_dd_pct scratch buffer; the trainer-struct `sp15_dd_trajectory_prev_dd` field is allocated via `MappedF32Buffer::new(1)` and reset via `host_slice_mut().fill(0.0)` — allowed-write under the rule, NOT a host-side compute), `feedback_isv_for_adaptive_bounds` (the 0.02 DD_TRAJECTORY_FLOOR and 2.0 RECOVERY_OVERSAMPLE_WEIGHT anchors are structural cold-start absorbers rewritten at fold boundary so the cold-start absorbers re-engage per phase precedent; runtime FLOOR will be signal-driven from the deferred 25th-percentile-of-dd_pct producer; runtime OVERSAMPLE_WEIGHT from the deferred current-dd_pct producer), `pearl_first_observation_bootstrap` (the persistent `prev_dd_pct=0` cold-start sentinel combined with the floor gate `prev > 0.02` ensures the kernel's first call cannot false-fire on a synthetic large-step transition; the Pearl A bootstrap pattern fires implicitly via the floor gate rather than via an explicit `prev == 0` branch — the result is bit-identical because at prev=0 the `prev > 0.02` predicate is false regardless of `now < prev`). + SP15 Phase 3.5.4 — plasticity injection trigger + warm-up tracker (weight-reset deferred) (2026-05-06): per spec §9.2 (3.5.4) post-amendment-2 fix. New single-kernel cubin `plasticity_injection_kernel.cu` mirrors prior Phase 3.5.X single-kernel-per-cubin pattern. TWO-STEP recovery: (1) On fire (`DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD` AND `PLASTICITY_FIRED_THIS_FOLD == 0`) → set fired flag, set warm-bars counter to M_warm (default 200). (2) Per-bar warm-bars decrement; consumer wiring (action-selection forces Hold via `effective_cooldown = max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)`) deferred. Weight-reset of last 10% of advantage-head weights to Kaiming-He init via cuRAND DEFERRED to Phase 3.5.4.b — kernel signature plumbed (`advantage_head_weights` pointer + `n_weights`) but no-op via `(void)` cast. 3 new ISV slots: `PLASTICITY_FIRED_THIS_FOLD_INDEX=436` (debounce flag, FoldReset sentinel 0 so re-arms each fold), `PLASTICITY_PERSISTENCE_THRESHOLD_INDEX=437` (initial 100.0 sentinel; ISV-tracked from running mean of dd_persistence_bars in follow-up), `PLASTICITY_WARM_BARS_REMAINING_INDEX=438` (counter, OR-gates with cooldown). 3 new state_reset_registry entries + 3 dispatch arms. 3 GPU oracle tests pass on RTX 3050 Ti: fires-when-persistence-exceeds-threshold (warm_bars [198, 200] post-fire-and-decrement), debounced-within-fold (no re-fire when fired=1; warm decrements normally), no-fire-below-threshold. Anchor test 2.22 plasticity_cooldown_interlock (Phase 2C / Phase 3.5 paired) green via 3.5.4.b follow-up + action-selection consumer wiring. Touched: `crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu` (new), `crates/ml/build.rs`, `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, `crates/ml/src/trainers/dqn/state_reset_registry.rs`, `crates/ml/src/trainers/dqn/trainer/training_loop.rs`, `crates/ml/tests/sp15_phase1_oracle_tests.rs`. Hard rules: feedback_no_atomicadd, feedback_no_partial_refactor (weight-reset + consumer wiring deferred), feedback_isv_for_adaptive_bounds (PLASTICITY_PERSISTENCE_THRESHOLD ISV-tracked). SP15 Phase 3.5.3 — cooldown gate (force Hold for M bars after K consecutive losses) (2026-05-06): per spec §9.2 (3.5.3). New single-kernel cubin `cooldown_kernel.cu` with one `extern "C" __global__` symbol (`cooldown_kernel`) → one cubin → one launcher (`launch_sp15_cooldown`) — mirrors the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` / 3.5.2 `dd_asymmetric_reward_kernel.cu` single-kernel-per-cubin pattern. After K consecutive losing trades, force Hold for M bars. K and M are ISV-tracked (slots 433/434); initial sentinels K=5, M=20 hardcoded in the trainer constructor. Counter (`COOLDOWN_BARS_REMAINING_INDEX=435`) decremented per bar — part of state so the model can reason about it. Reads ISV[COOLDOWN_K_THRESHOLD_INDEX=433] + ISV[COOLDOWN_M_BARS_INDEX=434] (both floored at 2.0/1.0 in-kernel to defang accidental ISV under-write — at K<2 the gate would fire on every loss, at M<1 the cooldown would be a no-op); read-modify-writes ISV[COOLDOWN_BARS_REMAINING_INDEX=435] + a separate mapped-pinned `consecutive_losses` scratch buffer (`sp15_cooldown_consecutive_losses`, [1] f32, persistent streak state between kernel calls). Per spec §9.2 (3.5.3) post-amendment-2 fix: streak counter only updates on `trade_close_event != 0`; per-bar non-close calls just decrement the cooldown counter without touching the loss tracker — separates the per-bar cadence (decrement) from the per-trade cadence (streak update). Trigger gate fires only when `trade_close_event != 0` AND `consec_losses >= K` AND `cooldown_remaining <= 0` — re-arming mid-cooldown would extend the gate every closed trade during the freeze, which is not the spec; resets `consec_losses=0` after firing. 4 new ISV slots: `COOLDOWN_K_THRESHOLD_INDEX=433` (initial 5.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up producer kernel reading the rolling median of consecutive-loss streak lengths via the two-heap algorithm (`MEDIAN_STREAK_LENGTH_INDEX=442`); the 5.0 anchor re-engages on each new fold), `COOLDOWN_M_BARS_INDEX=434` (initial 20.0; structural cold-start anchor — runtime value signal-driven from follow-up vol_normalizer time-to-mean-reversion producer; 20.0 anchor re-engages each fold), `COOLDOWN_BARS_REMAINING_INDEX=435` (initial 0.0; counter state — FoldReset sentinel 0 so the new fold starts with no active cooldown), `MEDIAN_STREAK_LENGTH_INDEX=442` (initial 0.0; stateful EMA, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first non-zero observation per `pearl_first_observation_bootstrap.md`). 1 new mapped-pinned scratch buffer on the trainer struct: `sp15_cooldown_consecutive_losses` ([1] f32, mirrors the Phase 3.1 `sp15_alpha_warm_count` non-ISV mapped-pinned pattern) — initialised to 0.0 in the constructor; read-modify-written by the kernel (`+= 1.0f` on trade-close losses, `= 0.0f` on trade-close wins/streak broken, untouched on non-trade-close bars); f32 representation is exact for any reachable streak length because the kernel only performs `+= 1.0f` increments and `= 0.0f` resets. 5 new `state_reset_registry` entries (`sp15_cooldown_k_threshold` / `sp15_cooldown_m_bars` rewrite the 5.0 / 20.0 anchors at fold boundary so the cold-start absorbers re-engage on each new fold; `sp15_cooldown_bars_remaining` resets the counter so the new fold starts with no active cooldown — leftover state from the previous fold would force Hold during the new fold's warmup; `sp15_median_streak_length` Pearl A sentinel 0; `sp15_cooldown_consecutive_losses` resets the non-ISV mapped-pinned scratch buffer to 0.0 via `host_slice_mut().fill(0.0)` — the new fold's first cooldown_kernel launch must see a fresh streak counter so a partially-accumulated streak from the previous fold cannot trip the gate during the new fold's warmup). 5 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.12 cooldown_engagement (Phase 2C / Phase 3.5 paired) — green via this commit. **Phase 3.5.3 lands kernel + launcher + 4 ISV anchor seeds + 1 scratch buffer + 5 registry entries + 5 dispatch arms only**; the action-selection wiring (force Hold while `cooldown_remaining > 0` — i.e. mask non-Hold actions in the per-action Q vector pre-argmax/Thompson selection when the cooldown is active) is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + 3.1-3.5.2 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). Both follow-up producer kernels (ISV-driven K via running median of consecutive-loss streak lengths using the two-heap algorithm; ISV-driven M via vol_normalizer time-to-mean-reversion) are documented as follow-up sub-tasks and do NOT block this Phase's primitive landing. Touched: `crates/ml/src/cuda_pipeline/cooldown_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `dd_penalty_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the writes; the kernel's ISV pointer is mutable `float*` because it writes the COOLDOWN_BARS_REMAINING counter slot, identical to `regret_signal_kernel`'s mutable-ISV pattern; the kernel takes a separate `consecutive_losses` device pointer so the streak state lives outside the ISV bus per the Phase 3.1 `sp15_alpha_warm_count` pattern of resetting non-ISV mapped-pinned buffers at fold boundary), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_COOLDOWN_CUBIN` cubin slot + 1 `pub fn launch_sp15_cooldown` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_COOLDOWN_CUBIN.to_vec())` and resolving `get_function("cooldown_kernel")` — single-thread/single-block grid mirroring `launch_sp15_dd_asymmetric_reward` precedent + 4 ISV constructor writes (COOLDOWN_K_THRESHOLD=5.0, COOLDOWN_M_BARS=20.0, COOLDOWN_BARS_REMAINING=0.0, MEDIAN_STREAK_LENGTH=0.0) appended to the SP15 Phase 3.5.2 anchor block + 1 new `sp15_cooldown_consecutive_losses: MappedF32Buffer` field on the trainer struct + alloc + 0.0 init mirroring the Phase 3.1 `sp15_alpha_warm_count` precedent), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+5 `RegistryEntry` records — `sp15_cooldown_k_threshold` rewrites 5.0 anchor; `sp15_cooldown_m_bars` rewrites 20.0 anchor; `sp15_cooldown_bars_remaining` resets counter to 0; `sp15_median_streak_length` Pearl A sentinel 0; `sp15_cooldown_consecutive_losses` non-ISV mapped-pinned buffer reset to 0.0 via `host_slice_mut().fill(0.0)`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+5 dispatch arms for the 5 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+3 GPU oracle tests in `mod gpu`: `cooldown_trips_after_k_losses` validates 5 sequential losing trade-close events trigger the gate — `cooldown_remaining` ∈ [18, 20] after the 5th loss, allowing implementation-order flexibility for the trigger-then-decrement single-call sequence; `cooldown_no_trip_when_streak_broken` validates 4 losses then 1 win then 4 more losses → no trip because the win resets the streak (the second 4-loss run never reaches K=5); `cooldown_decrements_on_non_trade_bars` validates 5 non-trade-close bars decrement `cooldown_remaining` from 10 → 5 AND leave `consec_losses` untouched at 2 — the per-bar cadence is independent of the per-trade cadence per spec post-amendment-2 fix). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic + a few read-modify-writes of ISV slots and the streak counter, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + scratch buffer + registry entries + dispatch arms land atomically; action-selection wiring follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the streak update + trigger + decrement; the three oracle tests verify trip-after-K, streak-broken-by-win, and non-trade-close-decrement-only behaviors with sentinel write semantics rather than skip-store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and consec_losses scratch buffer; `last_trade_outcome_loss` and `trade_close_event` scalars are by-value at the launcher boundary; the trainer-struct `sp15_cooldown_consecutive_losses` field is allocated via `MappedF32Buffer::new(1)` and reset via `host_slice_mut().fill(0.0)` — allowed-write under the rule, NOT a host-side compute), `feedback_isv_for_adaptive_bounds` (the 5.0 K and 20.0 M anchors are structural cold-start absorbers rewritten at fold boundary so the cold-start absorbers re-engage per phase precedent; runtime K will be signal-driven from the deferred two-heap median producer reading MEDIAN_STREAK_LENGTH_INDEX=442; runtime M from the deferred vol_normalizer time-to-mean-reversion producer), `pearl_first_observation_bootstrap` (MEDIAN_STREAK_LENGTH sentinel 0 so the deferred two-heap median producer's Pearl A first-observation replacement fires on the new fold's first non-zero streak observation).