diff --git a/crates/ml/build.rs b/crates/ml/build.rs index b78b5998a..8900e2848 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -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) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 849c02775..384e26699 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, + 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 diff --git a/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu new file mode 100644 index 000000000..04549a804 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu @@ -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(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 2a5f4bf47..855cdf581 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -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 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f2d3f1488..d39029ca7 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -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 diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 70ecee498..2abb0d4bb 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index aa674948f..b9aab9eb7 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.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). SP15 Phase 3.5.2 — asymmetric reward under DD (gain-amplified by dd_pct, pre-SP12-cap) (2026-05-06): per spec §9.2 (3.5.2). New single-kernel cubin `dd_asymmetric_reward_kernel.cu` with one `extern "C" __global__` symbol (`dd_asymmetric_reward_kernel`) → one cubin → one launcher (`launch_sp15_dd_asymmetric_reward`) — mirrors the Phase 3.3 `dd_penalty_kernel.cu` / 3.4 `regret_signal_kernel.cu` / 3.5 `hold_floor_kernel.cu` single-kernel-per-cubin pattern. For gains: `r_adjusted = r × (1 + λ × dd_pct)`. For losses: unchanged. Reads ISV[DD_ASYMMETRY_LAMBDA_INDEX=430] + ISV[DD_PCT_INDEX=406] (written by Task 1.3's `dd_state_kernel`, already landed); writes the adjusted reward to a caller-supplied output buffer AND records the most-recent boost factor at ISV[R_GAIN_DD_BOOST_INDEX=431] (diagnostic-only — read by HEALTH_DIAG in the consumer-wiring follow-up; 1.0 ⇒ no-op (loss side, or dd_pct=0 at ATH); >1.0 ⇒ gain × DD-recovery boost actually fired). Multiplier applies BEFORE the SP12 v3 NEG/POS reward cap clamp — saturating the POS_CAP for big recovery trades is behaviorally correct per spec (encourages frequent small recoveries over rare large ones; the cap is the saturation point, the boost is the route there). Asymmetric: gain-only multiplier preserves loss aversion through the SP12-cap pipeline per `pearl_audit_unboundedness_for_implicit_asymmetry` (Phase 3.3 restored asymmetry on the drawdown-axis penalty side; 3.5.2 closes the recovery half of the same axis). 3 new ISV slots: `DD_ASYMMETRY_LAMBDA_INDEX=430` (initial 0.5; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up λ producer kernel that tracks the running variance of dd_pct (`DD_DIST_VAR_INDEX=432`); the 0.5 anchor re-engages on each new fold), `R_GAIN_DD_BOOST_INDEX=431` (initial 1.0 sentinel — diagnostic slot recording the most-recent boost factor; FoldReset rewrites the 1.0 sentinel so the first HEALTH_DIAG emission on a new fold (before any kernel launch) reads as 'no boost yet' rather than carrying the previous fold's last-recorded value), `DD_DIST_VAR_INDEX=432` (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`). 3 new `state_reset_registry` entries (`sp15_dd_asymmetry_lambda` rewrites the 0.5 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_r_gain_dd_boost` rewrites the 1.0 sentinel; `sp15_dd_dist_var` Pearl A sentinel 0). 3 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 the combined 3.5.2 + 3.5.3 + 3.5.4 + 3.5.5 mechanisms; this commit lands the asymmetric-reward primitive. **Phase 3.5.2 lands kernel + launcher + 3 ISV anchor seeds + 3 registry entries + 3 dispatch arms only**; the host-side reward composer that applies this multiplier BEFORE the SP12 cap is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + 3.1-3.5 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). The follow-up λ producer (running variance of dd_pct via the existing SP4 wiener_state companion buffer pattern) is also documented as a follow-up sub-task and does NOT block this Phase's primitive landing. Touched: `crates/ml/src/cuda_pipeline/dd_asymmetric_reward_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 diagnostic R_GAIN_DD_BOOST slot, contrasted with `dd_penalty_kernel`'s `const float*` read-only ISV; identical to `regret_signal_kernel`'s mutable-ISV 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_ASYMMETRIC_REWARD_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_asymmetric_reward` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_DD_ASYMMETRIC_REWARD_CUBIN.to_vec())` and resolving `get_function("dd_asymmetric_reward_kernel")` — single-thread/single-block grid mirroring `launch_sp15_hold_floor` precedent + 3 ISV constructor writes (DD_ASYMMETRY_LAMBDA=0.5, R_GAIN_DD_BOOST=1.0, DD_DIST_VAR=0.0) appended to the SP15 Phase 3.5 anchor block), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+3 `RegistryEntry` records — `sp15_dd_asymmetry_lambda` rewrites 0.5 anchor; `sp15_r_gain_dd_boost` rewrites 1.0 sentinel; `sp15_dd_dist_var` Pearl A sentinel 0), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+3 dispatch arms for the 3 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+3 GPU oracle tests in `mod gpu`: `dd_asymmetric_reward_gain_amplified_by_dd_pct` validates `r=4.0, λ=0.5, dd_pct=0.5 → r_adjusted = 4.0 × (1 + 0.25) = 5.0` AND that R_GAIN_DD_BOOST is updated to 1.25 (diagnostic-slot write semantics); `dd_asymmetric_reward_loss_unchanged` validates `r=-3.0` passes through unchanged regardless of λ/dd_pct AND R_GAIN_DD_BOOST=1.0 on loss path (sentinel non-zero output buffer ensures the kernel actually overwrites with -3.0 rather than skipping); `dd_asymmetric_reward_no_op_at_ath` validates dd_pct=0 collapses the gain multiplier to 1.0 (no boost when at ATH) AND R_GAIN_DD_BOOST=1.0). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic + one diagnostic write, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + registry entries + dispatch arms land atomically; reward-composer pre-SP12-cap multiplier follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the asymmetric multiplier AND writes the diagnostic boost factor; the three oracle tests verify gain-amplified, loss-unchanged, and no-op-at-ATH 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 r_adjusted output buffer; `r_in` scalar is by-value at the launcher boundary), `feedback_isv_for_adaptive_bounds` (the 0.5 λ anchor is a structural cold-start absorber rewritten at fold boundary so the cold-start absorber re-engages per phase precedent; runtime λ will be signal-driven from the deferred running-variance-of-dd_pct producer), `pearl_audit_unboundedness_for_implicit_asymmetry` (the gain-only multiplier IS the asymmetry — losses untouched preserves loss aversion through the SP12-cap pipeline; symmetric application would erase the recovery-axis loss-aversion teaching this kernel encodes), `pearl_symmetric_clamp_audit` (the `r_in > 0` branch is structurally asymmetric BY DESIGN per spec §9.2 (3.5.2); losses pass through unchanged so the asymmetric NEG/POS caps in SP12 v3 remain semantically aligned — NEG_CAP gates loss magnitude, POS_CAP gates boosted-gain magnitude — this is the load-bearing mechanism for the recovery-axis loss-aversion teaching, NOT a symmetric-clamp violation), `pearl_first_observation_bootstrap` (DD_DIST_VAR sentinel 0 so the deferred λ producer's Pearl A first-observation replacement fires on the new fold's first non-zero variance observation).