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

feat(sp15-p3.5.4): plasticity injection trigger + warm-up tracker (weight-reset deferred)

Browse Source 
Per spec §9.2 (3.5.4) post-amendment-2 fix. TWO-STEP recovery:
  1. Fire when DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD AND
     PLASTICITY_FIRED_THIS_FOLD == 0 → set fired flag, set warm-bars
     counter to M_warm (default 200). [DEFERRED: reset last 10% of
     advantage-head weights to Kaiming-He init via cuRAND.]
  2. Per-bar warm-bars decrement; action-selection layer (consumer wiring
     follow-up) reads max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_
     REMAINING) and forces Hold while > 0.

Weight-reset DEFERRED to Phase 3.5.4.b — kernel signature plumbed
(advantage_head_weights + n_weights) but no-op via (void) cast.
Documented in audit doc.

3 ISV slots: 436 PLASTICITY_FIRED_THIS_FOLD (debounce flag, resets at
fold boundary to re-arm next fold), 437 PLASTICITY_PERSISTENCE_THRESHOLD
(initial 100.0 sentinel; ISV-tracked from running mean of dd_persistence
in follow-up), 438 PLASTICITY_WARM_BARS_REMAINING (counter, OR-gates
with cooldown).

3 fold-reset registry entries + dispatch arms.

Three GPU oracle tests pass: fires-when-persistence-exceeds-threshold
(warm_bars [198, 200] post-fire-and-decrement), debounced-within-fold
(no re-fire when fired=1; warm decrements 50→49), 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.

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

35
crates/ml/build.rs
View File

@@ -1019,6 +1019,41 @@ fn main() {
// 1.1-1.5 + 3.1-3.5.2 precedent — kernel + launcher verify in
// isolation first via the GPU oracle tests below).
"cooldown_kernel.cu",
// SP15 Phase 3.5.4 (2026-05-06): plasticity injection — TWO-STEP
// recovery (Flat + cooldown) with Kaiming-He reset of last 10%
// of advantage-head weights. K=DD_PERSISTENCE threshold is
// ISV-tracked (slot 437); initial sentinel 100.0 hardcoded in
// the trainer constructor (ISV-driven from running mean of
// dd_persistence_bars is a documented follow-up sub-task per
// `feedback_isv_for_adaptive_bounds.md`). Single source file
// with one `extern "C" __global__` symbol
// (`plasticity_injection_kernel`) → one cubin → one launcher
// (`launch_sp15_plasticity_injection`). 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` single-kernel-per-cubin pattern.
// Reads ISV[DD_PERSISTENCE_INDEX=404] +
// ISV[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX=437]; read-modify-
// writes ISV[PLASTICITY_FIRED_THIS_FOLD_INDEX=436] +
// ISV[PLASTICITY_WARM_BARS_REMAINING_INDEX=438]. Per spec §9.2
// (3.5.4) post-amendment-2 fix: fire when DD_PERSISTENCE
// exceeds threshold AND PLASTICITY_FIRED_THIS_FOLD == 0 →
// set fired flag, set warm-bars counter to M_warm (default
// 200), [DEFERRED: reset last 10% of advantage-head weights];
// per-bar warm-bars decrement; consumer-wiring follow-up reads
// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)
// and forces Hold while > 0. Phase 3.5.4 lands kernel +
// launcher + 3 ISV anchor seeds + 3 registry entries + 3
// dispatch arms only; the action-selection two-step Flat +
// cooldown wiring AND the actual Kaiming-He weight-reset are
// both deferred to a Phase 3.5.4.b follow-up atomic commit per
// `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 +
// 3.1-3.5.3 precedent — kernel + launcher verify in isolation
// first via the GPU oracle tests below). The kernel signature
// accepts the advantage_head_weights pointer + n_weights param
// so the follow-up commit is purely additive on the kernel
// body, not on the launcher contract.
"plasticity_injection_kernel.cu",
];
// ALL kernels get common header (BF16 types + wrappers)

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

@@ -1712,6 +1712,95 @@ pub fn launch_sp15_cooldown(
Ok(())
}
/// SP15 Phase 3.5.4 (2026-05-06): cubin slot for the plasticity
/// injection trigger + warm-up tracker kernel. TWO-STEP recovery per
/// spec §9.2 (3.5.4) post-amendment-2 fix:
/// 1. Fire when `DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD`
/// AND `PLASTICITY_FIRED_THIS_FOLD == 0` → set fired flag, set
/// warm-bars counter to M_warm (default 200), [DEFERRED: reset
/// last 10% of advantage-head weights to Kaiming-He init].
/// 2. Per-bar warm-bars decrement; consumer-wiring follow-up reads
/// `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)`
/// and forces Hold while > 0 (the OR-gate combining cooldown +
/// plasticity warm-up is the action-selection two-step).
///
/// 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`
/// single-kernel-per-cubin pattern.
///
/// Phase 3.5.4 lands kernel + launcher + 3 ISV constructor seeds + 3
/// state-reset registry entries + 3 dispatch arms only; the action-
/// selection two-step Flat + cooldown wiring AND the actual Kaiming-
/// He weight-reset are both deferred to a Phase 3.5.4.b follow-up
/// atomic commit per `feedback_no_partial_refactor.md` (matching
/// Phase 1.1-1.5 + 3.1-3.5.3 precedent — kernel + launcher verify in
/// isolation first via the GPU oracle tests below). The kernel
/// signature accepts the advantage_head_weights pointer + n_weights
/// param so the follow-up commit is purely additive on the kernel
/// body, not on the launcher contract.
pub static SP15_PLASTICITY_INJECTION_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/plasticity_injection_kernel.cubin"));
/// SP15 Phase 3.5.4 (2026-05-06): launcher for the plasticity
/// injection trigger + warm-up tracker kernel. Free function (matches
/// `launch_sp15_dd_penalty` / `launch_sp15_regret_signal` /
/// `launch_sp15_hold_floor` / `launch_sp15_dd_asymmetric_reward` /
/// `launch_sp15_cooldown` 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
/// 404 is read for DD_PERSISTENCE; slot 437 is read for
/// PLASTICITY_PERSISTENCE_THRESHOLD; slots 436 + 438 are read-modify-
/// written for PLASTICITY_FIRED_THIS_FOLD + PLASTICITY_WARM_BARS_
/// REMAINING). `advantage_head_weights` MUST point to at least
/// `n_weights` writable f32 slots — but the weight-reset is DEFERRED
/// to Phase 3.5.4.b so the kernel body currently does NOT touch the
/// buffer; the parameter is wired through the launcher contract so
/// the follow-up commit is purely additive on the kernel body.
/// `m_warm` is the warm-up bar count to set on fire (default 200 per
/// spec §9.2 (3.5.4)). Single thread, single block (mirrors
/// `cooldown_kernel` / `dd_asymmetric_reward_kernel` / `hold_floor_kernel`
/// — no reduction).
pub fn launch_sp15_plasticity_injection(
stream: &Arc<CudaStream>,
isv: cudarc::driver::sys::CUdeviceptr,
advantage_head_weights: cudarc::driver::sys::CUdeviceptr,
n_weights: i32,
m_warm: f32,
) -> Result<(), MLError> {
let module = stream
.context()
.load_cubin(SP15_PLASTICITY_INJECTION_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!(
"load sp15_plasticity_injection cubin: {e}"
)))?;
let kernel = module
.load_function("plasticity_injection_kernel")
.map_err(|e| MLError::ModelError(format!(
"load plasticity_injection_kernel function: {e}"
)))?;
unsafe {
stream
.launch_builder(&kernel)
.arg(&isv)
.arg(&advantage_head_weights)
.arg(&n_weights)
.arg(&m_warm)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
shared_mem_bytes: 0,
})
.map_err(|e| MLError::ModelError(format!(
"launch plasticity_injection_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
@@ -20960,6 +21049,65 @@ impl GpuDqnTrainer {
*sig_ptr.add(COOLDOWN_BARS_REMAINING_INDEX) = 0.0_f32;
*sig_ptr.add(MEDIAN_STREAK_LENGTH_INDEX) = 0.0_f32;
// SP15 Phase 3.5.4 (2026-05-06): plasticity injection
// anchors (per spec §9.2 (3.5.4) post-amendment-2 fix).
// TWO-STEP recovery — fire when
// `DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD`
// AND `PLASTICITY_FIRED_THIS_FOLD == 0` → set fired
// flag, set warm-bars counter to M_warm (default 200),
// [DEFERRED: reset last 10% of advantage-head weights
// to Kaiming-He init]. Per-bar warm-bars decrement;
// consumer-wiring follow-up reads
// max(COOLDOWN_BARS_REMAINING,
// PLASTICITY_WARM_BARS_REMAINING) and forces Hold
// while > 0 (the OR-gate combining cooldown +
// plasticity warm-up is the action-selection two-step).
//
// PLASTICITY_FIRED_THIS_FOLD = 0.0 (debounce flag —
// fires at most once
// per fold; reset
// to 0.0 at fold
// boundary re-arms
// the gate for the
// next fold)
// PLASTICITY_PERSISTENCE_THRESHOLD = 100.0 (Invariant-1
// anchor —
// ISV-driven
// from running
// mean of
// dd_persistence_bars
// is a
// follow-up
// sub-task per
// `feedback_isv_for_adaptive_bounds.md`;
// the 100.0
// anchor
// re-engages
// on each
// new fold)
// PLASTICITY_WARM_BARS_REMAINING = 0.0 (cold-start: no
// active warm-up;
// FoldReset
// sentinel 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)
use crate::cuda_pipeline::sp15_isv_slots::{
PLASTICITY_FIRED_THIS_FOLD_INDEX,
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX,
PLASTICITY_WARM_BARS_REMAINING_INDEX,
};
*sig_ptr.add(PLASTICITY_FIRED_THIS_FOLD_INDEX) = 0.0_f32;
*sig_ptr.add(PLASTICITY_PERSISTENCE_THRESHOLD_INDEX) = 100.0_f32;
*sig_ptr.add(PLASTICITY_WARM_BARS_REMAINING_INDEX) = 0.0_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

125
crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu Normal file
View File

@@ -0,0 +1,125 @@
// crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu
//
// SP15 Phase 3.5.4 — plasticity injection trigger + warm-up tracker.
//
// TWO-STEP recovery per spec §9.2 (3.5.4) post-amendment-2 fix:
// 1. On fire: set PLASTICITY_FIRED_THIS_FOLD = 1.0, set
// PLASTICITY_WARM_BARS_REMAINING = M_warm, [DEFERRED: reset last
// 10% of advantage-head weights to Kaiming-He init (gain=√2 for
// ReLU/GLU activations) via cuRAND]. The action-selection layer
// (consumer-wiring follow-up) reads the
// PLASTICITY_FIRED_THIS_FOLD == 1 flag transition AND emits Flat
// that bar.
// 2. Per-bar decrement of PLASTICITY_WARM_BARS_REMAINING; the
// consumer-wiring follow-up uses
// `effective_cooldown = max(COOLDOWN_BARS_REMAINING,
// PLASTICITY_WARM_BARS_REMAINING)` to force Hold while either
// counter > 0.
//
// **Weight-reset is DEFERRED** to a Phase 3.5.4.b follow-up. This
// commit lands:
// - Fire/debounce trigger logic (PLASTICITY_FIRED_THIS_FOLD gates
// re-fire within the same fold; reset to 0.0 at fold boundary
// re-arms the next fold).
// - Per-bar warm-bars decrement (every kernel call, regardless of
// fire).
// - The advantage_head_weights pointer + n_weights param wired
// through the launcher signature (kernel accepts but does not
// modify them yet — same `(void)foo` no-op pattern used elsewhere
// for deferred parameter plumbing).
//
// Why deferred: weight-reset requires Kaiming-He init via cuRAND, and
// accurate targeting of "last 10% of advantage-head weights" requires
// plumbing specific param-buffer pointers (mag/ord/urg head Linear
// layer slices) that aren't available at this kernel's call site
// without trainer-internal access. Establishing the kernel signature +
// flag-tracking now keeps the follow-up commit purely additive on the
// weight-reset side, mirroring the Phase 1.1-1.5 + 3.1-3.5.3 precedent
// of kernel + launcher landing first and consumer wiring following
// atomically.
//
// Slot reads/writes (literal indices to keep this kernel self-contained;
// the canonical names live in `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`
// and the constructor + state_reset_registry pin those names to these
// indices via compile-time `assert_eq!` regression checks):
// ISV[404] DD_PERSISTENCE (read)
// ISV[437] PLASTICITY_PERSISTENCE_THRESHOLD (read)
// ISV[436] PLASTICITY_FIRED_THIS_FOLD (read-modify-write)
// ISV[438] PLASTICITY_WARM_BARS_REMAINING (read-modify-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;
// consumer wiring (action-selection two-step Flat+cooldown
// integration via `effective_cooldown =
// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)`)
// deferred to a follow-up atomic commit per the established Phase
// 1.1-1.5 / 3.1-3.5.3 precedent. Weight-reset is also deferred
// explicitly — the launcher signature accepts the pointer + count
// so the follow-up is purely additive on the kernel body, not on
// the launcher contract.
// * `feedback_no_stubs` — kernel returns the actual updated state
// (flag transition + warm-bars decrement); the deferred weight-
// reset is a no-op on the same call rather than a stubbed
// placeholder return. The three oracle tests verify fire-on-
// threshold-crossing, debounce-within-fold, and below-threshold-
// no-fire behaviors.
// * `feedback_isv_for_adaptive_bounds` — PLASTICITY_PERSISTENCE_
// THRESHOLD is ISV-tracked (slot 437); the producer driving it
// from the running mean of dd_persistence_bars is deferred but
// the kernel reads it rather than a hardcoded literal so the
// follow-up swap is a no-op at the consumer.
//
// Single-thread, single-block (mirrors `cooldown_kernel` /
// `dd_asymmetric_reward_kernel` / `hold_floor_kernel` /
// `regret_signal_kernel` / `dd_penalty_kernel` — no reduction).
extern "C" __global__ void plasticity_injection_kernel(
float* __restrict__ isv,
float* __restrict__ advantage_head_weights, /* deferred: weight-reset NO-OP for now */
int n_weights,
float m_warm /* default 200 bars */
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
const float persistence = isv[/*DD_PERSISTENCE_INDEX*/ 404];
const float threshold = isv[/*PLASTICITY_PERSISTENCE_THRESHOLD_INDEX*/ 437];
const float fired = isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436];
float warm = isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438];
// Trigger: persistence exceeds threshold AND not yet fired this
// fold. PLASTICITY_FIRED_THIS_FOLD is reset to 0.0 at fold boundary
// by the `sp15_plasticity_fired_this_fold` registry-arm dispatch,
// re-arming the gate for the next fold. The `< 0.5f` half-bit
// tolerance check matches the existing `cooldown_remaining <= 0.0f`
// float-equality-with-tolerance pattern in `cooldown_kernel`.
if (persistence > threshold && fired < 0.5f) {
isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436] = 1.0f;
warm = m_warm;
// Weight-reset DEFERRED to Phase 3.5.4.b — kernel param plumbed
// through but no-op. Follow-up: Kaiming-He reset of last 10% of
// advantage-head weights via cuRAND with gain=√2 for ReLU/GLU
// activations. The `(void)foo` cast pattern matches the Rust
// `_ = foo` convention used elsewhere in this codebase for
// deferred parameter plumbing.
(void)advantage_head_weights;
(void)n_weights;
}
// Per-bar decrement of warm-bars (every call, regardless of fire).
// The trigger above runs first so the same call that fires the gate
// also immediately consumes one bar of warm-up — bias downward by 1
// is acceptable per the test's [198, 200] tolerance and avoids the
// bookkeeping branch needed to skip the decrement on the trip step
// (mirrors `cooldown_kernel`'s identical trigger-then-decrement
// sequence).
if (warm > 0.0f) {
warm -= 1.0f;
}
isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438] = warm;
__threadfence_system();
}

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

@@ -1296,6 +1296,40 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.sp15_cooldown_consecutive_losses [1] f32 mapped-pinned scratch buffer — SP15 Phase 3.5.3 (2026-05-06) consecutive-loss streak counter for the cooldown gate (per spec §9.2 (3.5.3) post-amendment-2 fix). Initialised to 0.0 in the constructor; read-modify-written by `cooldown_kernel`: `+= 1.0f` on trade-close losses, `= 0.0f` on trade-close wins (streak broken), untouched on non-trade-close bars. FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)` (mirrors the Phase 3.1 `sp15_alpha_warm_count` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary) — 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. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.3).",
},
// SP15 Phase 3.5.4 (2026-05-06): plasticity injection trigger
// + warm-up tracker anchors (per spec §9.2 (3.5.4)
// post-amendment-2 fix). TWO-STEP recovery — fire when
// DD_PERSISTENCE > threshold AND fired_flag == 0 → set
// fired flag, set warm-bars counter to M_warm (default
// 200), [DEFERRED: Kaiming-He weight reset]; per-bar
// warm-bars decrement; consumer-wiring follow-up reads
// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)
// and forces Hold while > 0. PLASTICITY_FIRED_THIS_FOLD
// resets to 0.0 at fold boundary — this is the load-
// bearing semantics for the "once per fold" debounce: the
// gate fires at most once per fold, then re-arms on the
// next fold. PLASTICITY_PERSISTENCE_THRESHOLD is a
// structural cold-start anchor per
// `feedback_isv_for_adaptive_bounds.md` (runtime value
// signal-driven from a follow-up producer reading the
// running mean of dd_persistence_bars). PLASTICITY_WARM_
// BARS_REMAINING is counter state — FoldReset sentinel 0
// so the new fold starts with no active warm-up.
RegistryEntry {
name: "sp15_plasticity_fired_this_fold",
category: ResetCategory::FoldReset,
description: "ISV[PLASTICITY_FIRED_THIS_FOLD_INDEX=436] — SP15 Phase 3.5.4 (2026-05-06) plasticity injection debounce flag. Set to 1.0 by `plasticity_injection_kernel` when DD_PERSISTENCE crosses the threshold (gate fires at most once per fold). FoldReset sentinel 0.0 — the load-bearing semantics for the 'once per fold' debounce: resetting to 0 at fold boundary re-arms the gate so it can fire again on the next fold's first persistence-threshold crossing. Without this reset, after the gate fires in fold N it would stay armed forever and never fire again across folds N+1..M, defeating the per-fold recovery semantic. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.4).",
},
RegistryEntry {
name: "sp15_plasticity_persistence_threshold",
category: ResetCategory::FoldReset,
description: "ISV[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX=437] — SP15 Phase 3.5.4 (2026-05-06) plasticity injection trigger threshold (number of consecutive bars in drawdown that triggers the gate). FoldReset rewrites the structural cold-start anchor 100.0 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from a follow-up producer reading the running mean of dd_persistence_bars (so the threshold tracks the typical drawdown duration distribution); the 100.0 anchor re-engages on each new fold so the cold-start absorber holds the threshold at a finite scale until the producer accumulates observations. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §9.2 (3.5.4).",
},
RegistryEntry {
name: "sp15_plasticity_warm_bars_remaining",
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).",
},
// 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

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

@@ -8378,6 +8378,40 @@ impl DQNTrainer {
fused.trainer_mut().sp15_cooldown_consecutive_losses.host_slice_mut().fill(0.0);
}
}
// SP15 Phase 3.5.4 (2026-05-06): plasticity injection
// anchors (per spec §9.2 (3.5.4) post-amendment-2 fix).
// TWO-STEP recovery — fire when DD_PERSISTENCE > threshold
// AND fired_flag == 0 → set fired flag, set warm-bars
// counter to M_warm (default 200), [DEFERRED: Kaiming-He
// weight reset]; per-bar warm-bars decrement;
// consumer-wiring follow-up reads
// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)
// and forces Hold while > 0. PLASTICITY_FIRED_THIS_FOLD
// resets to 0.0 at fold boundary — load-bearing for the
// "once per fold" debounce semantic (gate re-arms each
// fold). PLASTICITY_PERSISTENCE_THRESHOLD is a structural
// cold-start anchor (100.0); ISV-driven from running mean
// of dd_persistence_bars is a follow-up sub-task.
// PLASTICITY_WARM_BARS_REMAINING is counter state — FoldReset
// sentinel 0 so the new fold starts with no active warm-up.
"sp15_plasticity_fired_this_fold" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::PLASTICITY_FIRED_THIS_FOLD_INDEX;
fused.trainer().write_isv_signal_at(PLASTICITY_FIRED_THIS_FOLD_INDEX, 0.0);
}
}
"sp15_plasticity_persistence_threshold" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::PLASTICITY_PERSISTENCE_THRESHOLD_INDEX;
fused.trainer().write_isv_signal_at(PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, 100.0);
}
}
"sp15_plasticity_warm_bars_remaining" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp15_isv_slots::PLASTICITY_WARM_BARS_REMAINING_INDEX;
fused.trainer().write_isv_signal_at(PLASTICITY_WARM_BARS_REMAINING_INDEX, 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

@@ -1394,6 +1394,186 @@ mod gpu {
"consec_losses = {consec}, expected 2 (non-trade-close bars must not touch streak counter)"
);
}
/// SP15 Phase 3.5.4 — plasticity injection fires when
/// `DD_PERSISTENCE` crosses the threshold AND
/// `PLASTICITY_FIRED_THIS_FOLD == 0`. Sets the fired flag and the
/// warm-bars counter to M_warm. Per spec §9.2 (3.5.4) post-amendment-2
/// fix: TWO-STEP recovery — kernel sets fired + warm-bars counter
/// (this commit); action-selection layer (consumer wiring follow-up)
/// reads the fired flag transition AND emits Flat that bar, then
/// forces Hold while
/// `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING) > 0`.
/// Weight-reset deferred to Phase 3.5.4.b — kernel param plumbed
/// through but no-op for now (the test passes a non-zero
/// advantage_head_weights buffer to verify the launcher accepts it
/// without modifying the contents).
#[test]
#[ignore = "requires GPU"]
fn plasticity_fires_when_persistence_exceeds_threshold() {
use ml::cuda_pipeline::sp15_isv_slots::{
DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_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_PERSISTENCE_INDEX] = 150.0; // exceeds threshold
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0; // not yet fired
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
isv_buf.write_from_slice(&isv_init);
// Empty advantage-head weights buffer for now (weight-reset
// deferred to Phase 3.5.4.b — kernel takes the pointer but
// does not modify contents).
let weights_buf = unsafe { MappedF32Buffer::new(64) }
.expect("alloc MappedF32Buffer for advantage_head_weights");
weights_buf.write_from_slice(&vec![1.0f32; 64]);
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
&stream,
isv_buf.dev_ptr,
weights_buf.dev_ptr,
/*n_weights*/ 64,
/*m_warm*/ 200.0,
)
.expect("launch plasticity_injection_kernel (fires)");
stream
.synchronize()
.expect("synchronize after plasticity_injection_kernel launch (fires)");
let isv_host = isv_buf.read_all();
let fired = isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX];
assert!(
(fired - 1.0).abs() < 1e-9,
"PLASTICITY_FIRED_THIS_FOLD = {fired}, expected 1.0"
);
// After fire: warm = m_warm = 200 set, then per-bar decrement
// takes warm to 199 in the same call. Acceptable range
// [198, 200] mirrors the cooldown_kernel trigger-then-decrement
// tolerance.
let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX];
assert!(
(198.0..=200.0).contains(&warm),
"PLASTICITY_WARM_BARS_REMAINING = {warm}, expected ~199-200"
);
}
/// SP15 Phase 3.5.4 — plasticity injection does NOT re-fire within
/// the same fold (debounced by `PLASTICITY_FIRED_THIS_FOLD == 1`).
/// Per spec §9.2 (3.5.4) post-amendment-2 fix: the gate fires at
/// most once per fold; PLASTICITY_FIRED_THIS_FOLD resets to 0.0 at
/// fold boundary re-arms the gate for the next fold (handled by
/// `sp15_plasticity_fired_this_fold` registry-arm dispatch). Mid-
/// warm-up state: `fired = 1`, `warm_bars = 50`. After one call
/// with `DD_PERSISTENCE = 200` (way over threshold): fired stays
/// 1 (debounced), warm_bars decrements 50 → 49.
#[test]
#[ignore = "requires GPU"]
fn plasticity_debounced_within_fold() {
use ml::cuda_pipeline::sp15_isv_slots::{
DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_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_PERSISTENCE_INDEX] = 200.0; // way over threshold
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 1.0; // ALREADY fired this fold
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 50.0; // mid-warm-up
isv_buf.write_from_slice(&isv_init);
let weights_buf = unsafe { MappedF32Buffer::new(64) }
.expect("alloc MappedF32Buffer for advantage_head_weights");
weights_buf.write_from_slice(&vec![1.0f32; 64]);
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
&stream,
isv_buf.dev_ptr,
weights_buf.dev_ptr,
64,
200.0,
)
.expect("launch plasticity_injection_kernel (debounced)");
stream
.synchronize()
.expect("synchronize after plasticity_injection_kernel launch (debounced)");
let isv_host = isv_buf.read_all();
// fired stays 1 (already fired, debounced); warm_bars decrements 50 → 49.
assert!(
(isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9,
"PLASTICITY_FIRED_THIS_FOLD = {}, expected 1.0 (debounced — must NOT re-fire)",
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX]
);
let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX];
assert!(
(warm - 49.0).abs() < 1e-9,
"PLASTICITY_WARM_BARS_REMAINING = {warm}, expected 49 (decremented from 50, no re-fire)"
);
}
/// SP15 Phase 3.5.4 — plasticity injection does NOT fire when
/// `DD_PERSISTENCE` is below the threshold. `warm_bars` stays at 0
/// (no decrement when already 0 per the kernel's `if (warm > 0.0f)`
/// guard). Sentinel non-zero output verification: starting with
/// `fired = 0, warm = 0` and `persistence < threshold`, both flags
/// remain at 0 after the launch.
#[test]
#[ignore = "requires GPU"]
fn plasticity_no_fire_below_threshold() {
use ml::cuda_pipeline::sp15_isv_slots::{
DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX,
PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_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_PERSISTENCE_INDEX] = 50.0; // BELOW threshold
isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 100.0;
isv_init[PLASTICITY_FIRED_THIS_FOLD_INDEX] = 0.0;
isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0;
isv_buf.write_from_slice(&isv_init);
let weights_buf = unsafe { MappedF32Buffer::new(64) }
.expect("alloc MappedF32Buffer for advantage_head_weights");
weights_buf.write_from_slice(&vec![1.0f32; 64]);
ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection(
&stream,
isv_buf.dev_ptr,
weights_buf.dev_ptr,
64,
200.0,
)
.expect("launch plasticity_injection_kernel (no-fire)");
stream
.synchronize()
.expect("synchronize after plasticity_injection_kernel launch (no-fire)");
let isv_host = isv_buf.read_all();
assert!(
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX].abs() < 1e-9,
"PLASTICITY_FIRED_THIS_FOLD = {}, expected 0.0 (below-threshold must not fire)",
isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX]
);
assert!(
isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX].abs() < 1e-9,
"PLASTICITY_WARM_BARS_REMAINING = {}, expected 0.0 (no fire → no warm-up set; the `if (warm > 0)` guard prevents underflow when already 0)",
isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX]
);
}
}
/// 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