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

feat(sp15-p3.5.5): recovery curriculum — per-step DD_TRAJECTORY_DECREASING proxy

Browse Source 
Per spec §9.2 (3.5.5) post-amendment-2: replaces non-existent
episode-level metadata with per-bar signal that fires when
dd_pct(t) < dd_pct(t-1) AND dd_pct(t-1) > DD_TRAJECTORY_FLOOR — i.e.
transition is part of a recovery from non-trivial DD.

PER sampler (Phase 3.5.5.b follow-up) will read this and weight:
  sampling_weight = base × (1 + RECOVERY_OVERSAMPLE_WEIGHT × signal)
so recovery transitions get amplified gradient signal, completing the
downward-spiral break-out chain (3.5.2 reward asymmetry → 3.5.3
cooldown gate → 3.5.4 plasticity → 3.5.5 PER recovery curriculum).

3 ISV slots: 439 DD_TRAJECTORY_DECREASING, 440 RECOVERY_OVERSAMPLE_
WEIGHT (2.0 sentinel; ISV-driven from current dd_pct in follow-up),
441 DD_TRAJECTORY_FLOOR (0.02 sentinel; ISV-driven 25th percentile
of running dd_pct distribution in Phase 3.5.5.c follow-up per
feedback_isv_for_adaptive_bounds).

New sp15_dd_trajectory_prev_dd MappedF32Buffer (size 1) tracks
prev_dd across kernel calls — mirrors Task 3.5.3 sp15_cooldown_
consecutive_losses non-ISV mapped-pinned scratch pattern.

4 fold-reset registry entries + dispatch arms (3 ISV + 1 scratch).

Per established Phase precedent: kernel + launcher land first; PER
sampler integration and 25th-percentile floor producer are purely
additive follow-ups per feedback_no_partial_refactor.

Anchor test 2.10 recovery_after_streak (Phase 2C / Phase 3.5 paired) —
fully green via 3.5.2 + 3.5.4 + 3.5.5 combined.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 18:19:53 +02:00
parent e0e0abfb28
commit 69b8fdb61a
7 changed files with 620 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

32
crates/ml/build.rs
View File

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

109
crates/ml/src/cuda_pipeline/dd_trajectory_kernel.cu Normal file
View File

@@ -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 current<prev with
// prev>floor 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();
}

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

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

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

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

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

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

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

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

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long