From 483cef454c639650ebdfd28ae274bb63022d9c4f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 7 May 2026 08:58:17 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp15-p3.5.4.c):=20production=20caller=20+?= =?UTF-8?q?=20OR-gate=20consumer=20for=20plasticity=20injection=20?= =?UTF-8?q?=E2=80=94=20closes=203.5.4=20end-to-end?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Wave 4.2 (ef08611d3) landed cuRAND Kaiming-He weight reset as orphan- with-tests-only. Phase 3.5.4.c creates the production caller and the action-selection consumer atomically per spec lines 4346-4448. Production caller in gpu_experience_collector.rs step 4-pre (BEFORE experience_action_select): - fan_in = cfg.adv_h = 128 (CORRECTED from Wave 4.2's audit-doc error claiming 6528 = adv_h × num_atoms; that's total weight count of one branch's projection, not per-unit input dim. Per feedback_trust_code_not_docs the audit-doc error is also fixed inline in this commit.) - n_weights = branch_0_size × num_atoms × adv_h = 4 × 51 × 128 = 26112 (default config). Resets last 10% = 2611 directional advantage-head weights via cuRAND curand_normal × sqrt(2/128) ≈ 0.125 stddev (was wrongly documented as 0.0175). - Branch: w_b0out (directional) only — plasticity is about escaping stuck-in-flat regimes; targeted at directional choice. - Seed: mix_seed(Phase-3.5.4.c-unique base) ^ t per-step deterministic source; same (FOXHUNT_SEED, t) always yields the same Kaiming-He samples per kernel thread. - Trainer exposes target via new pub fn sp15_w_b0out_target() returning (dev_ptr, n_weights, fan_in); collector consumes via new set_sp15_plasticity_target(); training_loop wires the two. OR-gate consumer in experience_action_select: - New kernel arg float plasticity_m_warm threaded into the cooldown-mask code path. - Wave 1.B's cooldown_active = (cooldown_remaining > 0) extended to cooldown_active = (cooldown_remaining > 0) || (plasticity_warm_remaining > 0). - Flat-on-fire-bar detection — option (a), warm ≥ m_warm − 1.5f per the trigger-then-decrement convention. The kernel sees warm = m_warm − 1 on the fire bar; subsequent warm-up bars observe warm ≤ m_warm − 2. New if (plasticity_fire_bar) branch BEFORE the existing else if (cooldown_active) branch — the fire-bar gets DIR_FLAT, subsequent warm bars get DIR_HOLD via the OR-gate cooldown. Two-step recovery semantics: - Bar T (fire): plasticity launches → ISV[436]=1, ISV[438]= m_warm then decrements to m_warm−1. action_select detects fire-bar → dir_idx = DIR_FLAT. - Bars T+1..T+M_warm−1 (warm): action_select OR-gate forces dir_idx = DIR_HOLD. - Bar T+M_warm: warm transitions to 0, OR-gate inactive, normal Thompson/argmax resumes. When the production caller is unwired (test scaffold path): plasticity_m_warm passes 0.0f, the kernel's fire-bar predicate is dead (warm == 0 too), and the OR-gate degenerates to cooldown-only — bit-identical to the pre-3.5.4.c behaviour. Existing 3 SP15 3.5.b/3.5.3.b action_select oracle tests pass unchanged at m_warm = 0.0f. New behavioral oracle test plasticity_fires_force_flat_then_cooldown_ holds verifies the full sequence on RTX 3050 Ti with M_warm = 5 (shrunk from production's 200 for test runtime): bar 0 → DIR_FLAT, bars 1-3 → DIR_HOLD, bar 4 (warm boundary) → DIR_LONG. ISV side- effects (fired flips 0→1 then debounces, warm decrements with underflow guard) verified bar-by-bar. Eval/backtest path passes m_warm = 0.0f because plasticity is a training-only mechanism (during deterministic eval the weights must remain frozen at their checkpoint values). Atomic per feedback_no_partial_refactor: kernel + caller + consumer + trainer plumbing + 4 launcher-call-site updates (1 production collector + 1 eval-path + 2 test scaffolds) + new behavioral test + audit-doc fan_in inline correction + new audit-doc entry all in this commit. No fallback. SP15 Phase 3.5 recovery-dynamics chain is now end-to-end production-wired (3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.c + 3.5.5 + 3.5.5.b all firing). Verified: cargo check -p ml --features cuda clean; cargo check -p ml --features cuda --tests clean; CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored --nocapture 34 of 34 SP15 oracle tests green (5 plasticity + 3 action_select + 3 cooldown + 23 others); cargo test -p ml --features cuda --lib HOLDS the Wave 4.3 baseline (946 pass / 13 fail) on RTX 3050 Ti — no new regressions. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 97 +++++- .../cuda_pipeline/gpu_backtest_evaluator.rs | 11 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 37 ++ .../cuda_pipeline/gpu_experience_collector.rs | 158 +++++++++ crates/ml/src/cuda_pipeline/mod.rs | 3 + .../trainers/dqn/distributional_q_tests.rs | 3 + .../src/trainers/dqn/trainer/training_loop.rs | 26 ++ crates/ml/tests/sp15_phase1_oracle_tests.rs | 329 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 4 +- 9 files changed, 664 insertions(+), 4 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index c5d4957bc..fc9f2f61b 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1061,7 +1061,34 @@ extern "C" __global__ void experience_action_select( const float* __restrict__ b_logits_dir, /* [N, b0_size, n_atoms] direction-branch RAW C51 logits */ const float* __restrict__ per_sample_support, /* [N, b0_size, 3] stride-12 per-direction [v_min, v_max, delta_z] */ const float* __restrict__ atom_positions, /* [b0_size, n_atoms] adaptive non-linear positions; NULL = linear */ - int n_atoms /* C51 atom count */ + int n_atoms, /* C51 atom count */ + /* SP15 Phase 3.5.4.c (2026-05-07) — plasticity-injection consumer + * args. The plasticity-injection kernel runs BEFORE this kernel on + * the same stream and writes ISV[PLASTICITY_FIRED_THIS_FOLD=436] + + * ISV[PLASTICITY_WARM_BARS_REMAINING=438]. This kernel reads those + * slots to: + * (a) detect the FIRE BAR — the trigger writes warm = m_warm + * then immediately decrements to (m_warm − 1). On the fire + * bar this kernel observes warm ≥ (m_warm − 1.5f). On that + * bar dir_idx is forced to DIR_FLAT (close any open + * position — step 1 of the spec §9.2 (3.5.4) two-step + * recovery). + * (b) extend the SP15 Phase 3.5.3.b cooldown OR-gate from + * (cooldown_remaining > 0) to + * (max(cooldown_remaining, plasticity_warm_remaining) > 0); + * on bars T+1..T+M (post-fire warm-up) plasticity_warm_ + * remaining > 0 and dir_idx is forced to DIR_HOLD (step 2 of + * the two-step). The pre-existing 3.5.3.b code path that + * reads ISV[435] is preserved bit-identically when + * plasticity is unwired (m_warm == 0): warm stays 0, so the + * OR-gate degenerates to (cooldown_remaining > 0). + * + * `m_warm` is the warm-up bar count the plasticity-injection + * launcher passes to its trigger (default 200 per spec). Pass + * 0.0f to disable the fire-bar detection (test path / production + * code path that hasn't wired the plasticity launch yet — when + * m_warm == 0, warm == 0 too, so neither branch fires). */ + float plasticity_m_warm ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -1296,6 +1323,15 @@ extern "C" __global__ void experience_action_select( float q_eff_dir[4]; float hold_floor = 0.0f; int cooldown_active = 0; + /* SP15 Phase 3.5.4.c (2026-05-07) — plasticity-injection two-step + * recovery state. The trigger kernel writes ISV[438] = m_warm on + * the fire bar (then immediately decrements by 1 in the same + * launch), and decrements per bar thereafter. This kernel reads + * the warm counter and m_warm to discriminate the fire bar from + * subsequent warm-up bars; the fire-bar gets DIR_FLAT (step 1), + * subsequent warm-up bars get DIR_HOLD via the OR-gate (step 2). */ + int plasticity_fire_bar = 0; + int plasticity_warm_active = 0; if (isv_signals_ptr != NULL && b0_size <= 4) { /* Direction-policy entropy. Stable softmax over e_dir[0..b0_size). */ float ed_max = e_dir[0]; @@ -1325,7 +1361,40 @@ extern "C" __global__ void experience_action_select( /* Cooldown counter — ISV[COOLDOWN_BARS_REMAINING=435]. f32 * representation of an integer counter; > 0 means cooldown * still active. */ - cooldown_active = (isv_signals_ptr[/*COOLDOWN_BARS_REMAINING_INDEX*/ 435] > 0.0f); + const float cooldown_remaining = + isv_signals_ptr[/*COOLDOWN_BARS_REMAINING_INDEX*/ 435]; + + /* SP15 Phase 3.5.4.c — plasticity warm counter + + * fire-bar detection. The trigger kernel sets warm = m_warm + * on fire and decrements once in the same launch, so the + * fire bar observes warm == m_warm − 1 (use a 0.5f tolerance + * for the f32 equality compare per the cooldown_kernel + * trigger-then-decrement convention). When plasticity is + * unwired (m_warm == 0), warm == 0 too — both predicates are + * false and behaviour reduces to the pre-3.5.4.c (cooldown- + * only) path bit-identically. + * + * Per pearl_no_host_branches_in_captured_graph: the kernel- + * side detection avoids any host read of ISV between the + * plasticity launch and this action_select launch (both run + * outside the exp_fwd_graph capture region — exp_fwd_graph + * ends at the cuBLAS forward in gpu_experience_collector.rs + * before step 4 — but we keep the detection on-GPU + * regardless to preserve graph-friendly ordering invariants + * for any future graph-capture extension). */ + const float plasticity_warm_remaining = + isv_signals_ptr[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438]; + plasticity_warm_active = (plasticity_warm_remaining > 0.0f); + plasticity_fire_bar = (plasticity_m_warm > 0.0f) && + (plasticity_warm_remaining + >= plasticity_m_warm - 1.5f); + + /* OR-gate per spec §9.2 (3.5.4) post-amendment-2: the cooldown + * mask fires when EITHER counter is active. Pre-3.5.4.c + * behaviour preserved when plasticity is unwired (warm == 0 + * → only the cooldown branch can flip the gate). */ + cooldown_active = (cooldown_remaining > 0.0f) + || plasticity_warm_active; } for (int d = 0; d < b0_size; d++) { q_eff_dir[d] = e_dir[d]; @@ -1353,10 +1422,32 @@ extern "C" __global__ void experience_action_select( if (!(thompson_temp > 0.5f)) thompson_temp = 0.5f; if (thompson_temp > 2.0f) thompson_temp = 2.0f; - if (cooldown_active && DIR_HOLD < b0_size) { + if (plasticity_fire_bar && DIR_FLAT < b0_size) { + /* SP15 Phase 3.5.4.c — fire-bar Flat injection (step 1 of the + * spec §9.2 (3.5.4) two-step recovery). On the bar where the + * plasticity-injection trigger fired this fold (detected via + * `plasticity_warm_remaining ≥ m_warm − 1.5f` per the + * trigger-then-decrement convention) we force dir_idx = + * DIR_FLAT to close any open position. The OR-gate cooldown + * below would have forced DIR_HOLD on the same bar (warm > 0); + * the fire-bar branch supersedes it because Flat-this-bar is + * the directed exit, not Hold-this-bar. + * + * Per spec the warm counter then runs out over bars + * T+1..T+M during which the OR-gate forces DIR_HOLD (the + * `else if (cooldown_active)` branch below). DIR_FLAT = 3 + * by `state_layout.cuh`. Skips Pass 2 + RNG draws same as + * the cooldown branch — reproducibility-at-fixed-seed + * preserved because plasticity_fire_bar is a deterministic + * function of the kernel's ISV reads. */ + dir_idx = DIR_FLAT; + } else if (cooldown_active && DIR_HOLD < b0_size) { /* SP15 Phase 3.5.3.b — cooldown mask: when * `ISV[COOLDOWN_BARS_REMAINING=435] > 0`, force the picked * direction to Hold per spec §9.2 (3.5.3) post-amendment-2 fix. + * SP15 Phase 3.5.4.c extends the predicate to the OR of + * COOLDOWN_BARS_REMAINING and PLASTICITY_WARM_BARS_REMAINING + * so the post-fire warm-up window also forces Hold. * Hard short-circuit (skip Pass 2 entirely) — sidesteps the * temperature-blend numerics where a finite-sentinel-on-non-Hold * approach would still let a pure-Thompson (τ=1) sample dominate diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index d46a96784..7277b9842 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -1994,6 +1994,17 @@ impl GpuBacktestEvaluator { .arg(&per_sample_support_ptr) // per_sample_support: [N, 4, 3] stride-12 mapped pinned (trainer-owned) .arg(&atom_positions_ptr) // atom_positions: [b0, n_atoms] adaptive C51 positions (trainer-owned) .arg(&n_atoms_i32) // n_atoms: C51 atom count (i32 to match `int` ABI) + // SP15 Phase 3.5.4.c (2026-05-07): plasticity_m_warm — + // pass 0.0f during eval/backtest. Plasticity injection + // is a TRAINING-only mechanism (it resets advantage-head + // weights to break out of stuck-in-flat regimes during + // experience collection); during deterministic eval the + // weights must remain frozen at their checkpoint values. + // m_warm = 0.0f makes the action_select fire-bar + // predicate dead, and PLASTICITY_WARM_BARS_REMAINING + // stays at its 0 sentinel through eval (no producer + // running) so the OR-gate degenerates to cooldown-only. + .arg(&0.0f32) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "chunked action_select chunk_start={chunk_start}: {e}" diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 7122ad002..fcc6a16d1 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -10157,6 +10157,43 @@ impl GpuDqnTrainer { self.target_params_buf.raw_ptr() } + /// SP15 Phase 3.5.4.c (2026-05-07): return the + /// `(w_b0out_dev_ptr, n_weights, fan_in)` triple for the + /// directional advantage-head last-Linear weight tensor (param- + /// buffer tensor index 19). The collector's + /// `set_sp15_plasticity_target()` consumes this triple to wire + /// the per-step `launch_sp15_plasticity_injection` against the + /// directional branch only. + /// + /// Layout (per `compute_param_sizes` and the row/col-dim pairs + /// at `gpu_dqn_trainer.rs:29946`): + /// `w_b0out: [branch_0_size × num_atoms, adv_h]` row-major + /// (output_dim rows × input_dim columns). + /// `n_weights = branch_0_size × num_atoms × adv_h` + /// (default config: 4 × 51 × 128 = 26112). + /// `fan_in = adv_h` (= input dim of the per-unit weight + /// vector — Kaiming-He gain `sqrt(2/fan_in)` is per-output- + /// unit variance; default config: 128). + /// + /// The Wave 4.2 audit doc previously documented `fan_in = + /// adv_h × num_atoms = 6528`. That value is the total weight + /// count of one row of branch outputs (`num_atoms` atoms per + /// direction × `adv_h` inputs), not the input dim per unit; the + /// Kaiming-He variance derivation uses input dim per unit, which + /// is `adv_h`. The Phase 3.5.4.c audit-doc inline correction + /// restores the right value (`128`) and the right stddev + /// (`sqrt(2/128) ≈ 0.125`). + pub fn sp15_w_b0out_target(&self) -> (u64, i32, i32) { + let param_sizes = compute_param_sizes(&self.config); + let byte_offset = padded_byte_offset(¶m_sizes, 19); + let dev_ptr = self.params_buf.raw_ptr() + byte_offset; + let n_weights = (self.config.branch_0_size + * self.config.num_atoms + * self.config.adv_h) as i32; + let fan_in = self.config.adv_h as i32; + (dev_ptr, n_weights, fan_in) + } + /// Backward gradient w.r.t. h_s2 (trunk activation) — f32. pub fn bw_d_h_s2_buf(&self) -> &CudaSlice { &self.bw_d_h_s2 diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index dea8f47cf..87631306b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -983,6 +983,48 @@ pub struct GpuExperienceCollector { /// pattern. sp15_dd_trajectory_prev_dd_dev_ptr: u64, + /// SP15 Phase 3.5.4.c (2026-05-07): trainer's `params_buf` + /// device pointer offset to `w_b0out` (directional advantage-head + /// last-Linear weights, tensor index 19 per `compute_param_sizes`). + /// Set via `set_sp15_plasticity_target()` after trainer + /// construction. 0 = NULL → the per-step plasticity-injection + /// launch in step 4-pre is skipped AND `m_warm` is passed as 0.0f + /// to `experience_action_select` so the OR-gate consumer + /// degenerates to cooldown-only (bit-identical pre-3.5.4.c + /// behaviour). Production sets this in `training_loop.rs` + /// immediately after `set_sp15_dd_trajectory_prev_dd_ptr` — + /// mirrors the same wiring pattern. + /// + /// Plasticity is injected into the directional branch only — + /// per spec §9.2 (3.5.4) the recovery is about escaping stuck- + /// in-flat regimes (a directional decision problem), so resetting + /// the trailing 10% of the directional output projection is + /// targeted at the exact failure mode. mag/order/urgency branches + /// are untouched — their conditioning on direction makes them + /// naturally re-train once direction starts choosing again. + sp15_w_b0out_dev_ptr: u64, + /// SP15 Phase 3.5.4.c: total weight count for `w_b0out` + /// (= `branch_0_size × num_atoms × adv_h`; default config + /// 4 × 51 × 128 = 26112). Computed once at wiring time from + /// the trainer's network config; passed as the `n_weights` + /// arg to `launch_sp15_plasticity_injection` so the kernel's + /// `reset_count = n_weights / 10` (~2611 default config) + /// addresses the trailing-10%-slice within the directional + /// output projection. + sp15_w_b0out_n_weights: i32, + /// SP15 Phase 3.5.4.c: Kaiming-He fan_in for `w_b0out`. Per + /// `compute_param_sizes` line 4074 the layout is + /// `[branch_0_size × num_atoms, adv_h]` (row-major output_dim × + /// input_dim) so fan_in = adv_h (default config 128). The Wave + /// 4.2 audit doc previously documented fan_in = + /// `adv_h × num_atoms = 6528` — that was incorrect (total + /// weight count, not input dim per unit) and the Phase 3.5.4.c + /// audit-doc inline correction restores the right value. With + /// the correct fan_in the Kaiming-He stddev is + /// `sqrt(2 / 128) ≈ 0.125` (vs the audit doc's earlier + /// erroneous `sqrt(2/6528) ≈ 0.0175`). + sp15_w_b0out_fan_in: i32, + /// SP4 Layer A Task A13.5 (2026-05-01); GPU-Pearls refactor (2026-05-01) /// — cross-boundary Pearls A+D wiring for /// `launch_reward_component_ema_inplace`. The producer kernel writes @@ -1999,6 +2041,9 @@ impl GpuExperienceCollector { isv_signals_dev_ptr: 0, // NULL until trainer sets it sp15_alpha_warm_count_dev_ptr: 0, // NULL until trainer wires it (SP15 Wave 2) sp15_dd_trajectory_prev_dd_dev_ptr: 0, // NULL until trainer wires it (SP15 Wave 4.3 / 3.5.5.b) + sp15_w_b0out_dev_ptr: 0, // NULL until trainer wires it (SP15 Phase 3.5.4.c) + sp15_w_b0out_n_weights: 0, + sp15_w_b0out_fan_in: 0, // SP4 Task A13.5: Pearls A+D buffers for reward_component_ema. // All NULL until FusedTrainingCtx wires them post-collector // construction (mirrors A14/A15 Pearl C wiring path). 2026-05-01 @@ -2100,6 +2145,42 @@ impl GpuExperienceCollector { self.sp15_dd_trajectory_prev_dd_dev_ptr = dev_ptr; } + /// SP15 Phase 3.5.4.c (2026-05-07): set the trainer's `params_buf` + /// device pointer offset to `w_b0out` (directional advantage-head + /// last-Linear weights, tensor index 19 per `compute_param_sizes`) + /// + the matching `n_weights` (= `branch_0_size × num_atoms × + /// adv_h`; default config 4 × 51 × 128 = 26112) + Kaiming-He + /// `fan_in` (= `adv_h`; default 128). When all three are non-zero + /// alongside `isv_signals_dev_ptr`, the per-step + /// `launch_sp15_plasticity_injection` runs in step 4-pre (BEFORE + /// `experience_action_select`); the trigger writes + /// `ISV[PLASTICITY_FIRED_THIS_FOLD=436]` + `ISV[PLASTICITY_WARM_BARS_ + /// REMAINING=438]`, the action_select kernel's OR-gate consumer + /// reads both slots + `m_warm` (defaulted 200.0 per spec §9.2 + /// (3.5.4)) to force `dir_idx = DIR_FLAT` on the fire bar and + /// `dir_idx = DIR_HOLD` on subsequent warm-up bars. Pass 0 for any + /// of the three to disable (test scaffold path; the OR-gate + /// degenerates to cooldown-only — bit-identical to the + /// pre-3.5.4.c path). + /// + /// The trainer derives `n_weights` and `fan_in` from + /// `compute_param_sizes(&self.config)[19]` and `self.config.adv_h` + /// respectively (per the param-buffer layout at + /// `gpu_dqn_trainer.rs:4074` and the Kaiming-He fan-dim layout at + /// `gpu_dqn_trainer.rs:29946` — the row-dim `branch_0_size × + /// num_atoms` is the OUTPUT count, the col-dim `adv_h` is the + /// INPUT count = fan_in for the per-unit Kaiming-He variance). + pub fn set_sp15_plasticity_target( + &mut self, + w_b0out_dev_ptr: u64, + n_weights: i32, + fan_in: i32, + ) { + self.sp15_w_b0out_dev_ptr = w_b0out_dev_ptr; + self.sp15_w_b0out_n_weights = n_weights; + self.sp15_w_b0out_fan_in = fan_in; + } + /// Set trade plan params device pointer for plan activation/enforcement. /// The env_step kernel reads plan params to copy into portfolio state on /// Flat→Positioned transitions. Pass 0 to disable (NULL = no plan). @@ -4199,6 +4280,70 @@ impl GpuExperienceCollector { let na_i32 = self.num_atoms as i32; let support_dev_ptr = self.per_sample_support_buf.dev_ptr; let null_atom_positions = 0u64; // NULL = use linear atom spacing + + // ── 4-pre. SP15 Phase 3.5.4.c (2026-05-07): plasticity-injection + // production launch — runs BEFORE experience_action_select on the + // same stream so the trigger's ISV writes (slots 436 / 438) are + // serialized-visible to the action_select kernel's reads. Closes + // out the Wave 4.2 (`ef08611d3`) orphan-with-tests-only landing + // per `feedback_wire_everything_up`. + // + // Gated on `sp15_w_b0out_dev_ptr != 0`: the trainer wires this + // dev_ptr via `set_sp15_plasticity_target` after construction + // (mirrors `set_sp15_alpha_warm_count_ptr` / `set_sp15_dd_ + // trajectory_prev_dd_ptr`). Until wired, the launch is skipped + // and the ISV slots stay at their constructor sentinels (warm = 0, + // fired = 0); the action_select kernel's OR-gate consumer then + // degenerates to cooldown-only (bit-identical to the + // pre-3.5.4.c path). + // + // `plasticity_m_warm_arg` is the warm-up bar count (default 200 + // per spec §9.2 (3.5.4)) — the trigger writes it to ISV[438] on + // fire, action_select reads it back to detect the fire bar + // (warm ≥ m_warm − 1.5f per the kernel's trigger-then-decrement + // convention). Pass 0.0f when plasticity is unwired so the fire- + // bar predicate is dead in the kernel. + // + // `plasticity_seed` derives from the global FOXHUNT_SEED + // (`mix_seed`) XOR'd with the per-step timestep `t` so the + // cuRAND init in the trigger kernel produces deterministic-but- + // distinct samples each fire (multi-fold runs that fire on + // different timesteps see different Kaiming-He samples). Same + // (FOXHUNT_SEED, t) pair always yields the same reset. + // + // Out of the exp_fwd_graph capture region (which ends at the + // cuBLAS forward earlier in this loop body), so per + // `pearl_no_host_branches_in_captured_graph` the host-side + // gating compiles cleanly. + let plasticity_m_warm_arg: f32 = + if self.sp15_w_b0out_dev_ptr != 0 && self.isv_signals_dev_ptr != 0 { + const M_WARM: f32 = 200.0; // spec §9.2 (3.5.4) default + let plasticity_seed = crate::cuda_pipeline::mix_seed( + /* SP15 Phase 3.5.4.c historic base — arbitrary + * 64-bit constant unique to this launch site so + * multi-launcher runs (alpha_split / cooldown / + * etc.) get uncorrelated mix outputs even when + * FOXHUNT_SEED defaults to 42. */ + 0xC0DE_F00D_5915_3F4Cu64 + ) ^ (t as u64); + crate::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &self.stream, + self.isv_signals_dev_ptr, + self.sp15_w_b0out_dev_ptr, + self.sp15_w_b0out_n_weights, + M_WARM, + self.sp15_w_b0out_fan_in, + plasticity_seed, + ).map_err(|e| MLError::ModelError(format!( + "sp15 plasticity_injection t={t}: {e}" + )))?; + M_WARM + } else { + // Plasticity unwired — pass m_warm = 0.0 so the + // action_select fire-bar branch is inactive. + 0.0f32 + }; + unsafe { self.stream .launch_builder(&self.action_select_kernel) @@ -4233,6 +4378,19 @@ impl GpuExperienceCollector { .arg(&support_dev_ptr) // per_sample_support: [N, 4, 3] stride-12 mapped pinned .arg(&null_atom_positions) // atom_positions: NULL → linear support .arg(&na_i32) // n_atoms: C51 atom count (i32 to match `int` ABI) + // SP15 Phase 3.5.4.c (2026-05-07): plasticity-injection + // m_warm. When the per-step plasticity launch above + // wired this collector's `sp15_w_b0out_dev_ptr`, the + // m_warm passed to the trigger flows into the + // consumer here so the action_select kernel can + // detect the fire bar (warm ≥ m_warm − 1.5f) and + // force DIR_FLAT, then DIR_HOLD on subsequent + // warm-up bars via the OR-gate. When unwired + // (`sp15_w_b0out_dev_ptr == 0`), m_warm = 0.0f + // disables the fire-bar branch (warm == 0 too → + // OR-gate degenerates to cooldown-only — bit- + // identical to the pre-3.5.4.c path). + .arg(&plasticity_m_warm_arg) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_action_select t={t}: {e}" diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 3a6d3e288..1a4ef0ccd 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -1369,6 +1369,9 @@ mod tests { .arg(&support_buf) // per_sample_support [batch, b0, 3] .arg(&null_ptr) // atom_positions = NULL → linear support .arg(&n_atoms) // n_atoms + // SP15 Phase 3.5.4.c: plasticity_m_warm — 0.0f for this + // diagnostic test (no plasticity wiring in this test scaffold). + .arg(&0.0f32) .launch(cfg) .expect("launch experience_action_select"); } diff --git a/crates/ml/src/trainers/dqn/distributional_q_tests.rs b/crates/ml/src/trainers/dqn/distributional_q_tests.rs index b16edd96f..c23e88f85 100644 --- a/crates/ml/src/trainers/dqn/distributional_q_tests.rs +++ b/crates/ml/src/trainers/dqn/distributional_q_tests.rs @@ -1025,6 +1025,9 @@ impl ProdActionSelectFixture { .arg(&self.per_sample_support) .arg(&null_ptr) // atom_positions = NULL → linear support .arg(&self.n_atoms) + // SP15 Phase 3.5.4.c: plasticity_m_warm — 0.0f for this + // distributional Q test (no plasticity wiring needed). + .arg(&0.0f32) .launch(cfg) .expect("launch experience_action_select"); } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index ef952af0b..9d8e702f9 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -893,6 +893,32 @@ impl DQNTrainer { } } + // SP15 Phase 3.5.4.c (2026-05-07): wire the trainer's + // `params_buf` device pointer offset to `w_b0out` (directional + // advantage-head last-Linear weights, tensor index 19 per + // `compute_param_sizes`) + `n_weights` (= `branch_0_size × + // num_atoms × adv_h`; default config 4 × 51 × 128 = 26112) + + // Kaiming-He `fan_in` (= `adv_h`; default 128) into the + // collector so the per-step `launch_sp15_plasticity_injection` + // can run inside `collect_experiences_gpu` BEFORE + // `experience_action_select`. Closes out the Wave 4.2 + // (`ef08611d3`) orphan-with-tests-only landing per + // `feedback_wire_everything_up`. The plasticity trigger writes + // ISV[PLASTICITY_FIRED_THIS_FOLD=436] + + // ISV[PLASTICITY_WARM_BARS_REMAINING=438] which the + // action_select kernel then reads to force DIR_FLAT on the + // fire bar and DIR_HOLD on subsequent warm-up bars (OR-gate + // extension of the SP15 Phase 3.5.3.b cooldown mask). + // Mirrors the `set_sp15_dd_trajectory_prev_dd_ptr` wiring + // pattern. + if let Some(ref fused) = self.fused_ctx { + let (w_ptr, n_weights, fan_in) = + fused.trainer().sp15_w_b0out_target(); + if let Some(ref mut collector) = self.gpu_experience_collector { + collector.set_sp15_plasticity_target(w_ptr, n_weights, fan_in); + } + } + // SP15 Wave 4.3 / Phase 3.5.5.b (2026-05-07): wire the // trainer's ISV bus dev_ptr into the GPU PER replay buffer // so `per_insert_pa` can read diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index ac24763a9..b6051f7a8 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -2123,6 +2123,84 @@ mod gpu { .arg(per_sample_support) .arg(&null_ptr) // atom_positions = NULL → linear support .arg(&n_atoms) + // SP15 Phase 3.5.4.c (2026-05-07): plasticity_m_warm. + // Default 0.0f → fire-bar predicate dead. The new + // `plasticity_fires_force_flat_then_cooldown_holds` test + // overrides this via `launch_action_select_for_test_with_m_warm`. + .arg(&0.0f32) + .launch(cfg) + .expect("launch experience_action_select"); + } + stream.synchronize().expect("sync after action_select"); + stream.clone_dtoh(actions_buf).expect("readback actions") + } + + /// SP15 Phase 3.5.4.c (2026-05-07): variant of `launch_action_select_for_test` + /// that takes an explicit `m_warm` so the new behavioral test + /// `plasticity_fires_force_flat_then_cooldown_holds` can drive the + /// kernel's fire-bar / OR-gate consumer paths. All other args + /// match the m_warm-defaulted helper. + #[allow(clippy::too_many_arguments)] + fn launch_action_select_for_test_with_m_warm( + stream: &Arc, + batch: usize, + n_atoms: i32, + q_values: &CudaSlice, + b_logits_dir: &CudaSlice, + per_sample_support: &CudaSlice, + actions_buf: &mut CudaSlice, + q_gaps_buf: &mut CudaSlice, + intent_buf: &mut CudaSlice, + conv_buf: &mut CudaSlice, + mag_conv_buf: &mut CudaSlice, + isv_dev_ptr: u64, + m_warm: f32, + ) -> Vec { + let kernel = load_action_select(stream); + let blocks = ((batch as u32).div_ceil(256)).max(1); + let cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + let n_i32: i32 = batch as i32; + let null_ptr: u64 = 0; + let eps_start: f32 = 0.0; + let eps_end: f32 = 0.0; + unsafe { + stream + .launch_builder(kernel) + .arg(q_values) + .arg(&mut *actions_buf) + .arg(&mut *q_gaps_buf) + .arg(&eps_start) + .arg(&eps_end) + .arg(&0_i32) // current_epoch + .arg(&1_i32) // total_epochs + .arg(&n_i32) + .arg(&AS_B0) + .arg(&AS_B1) + .arg(&AS_B2) + .arg(&AS_B3) + .arg(&0.0_f32) // q_gap_threshold + .arg(&null_ptr) + .arg(&0_i32) + .arg(&0.0_f32) + .arg(&1.0_f32) + .arg(&1.0_f32) + .arg(&1.0_f32) + .arg(&0_i32) + .arg(&null_ptr) + .arg(&isv_dev_ptr) + .arg(&0_i32) + .arg(&mut *intent_buf) + .arg(&mut *conv_buf) + .arg(&mut *mag_conv_buf) + .arg(b_logits_dir) + .arg(per_sample_support) + .arg(&null_ptr) + .arg(&n_atoms) + .arg(&m_warm) .launch(cfg) .expect("launch experience_action_select"); } @@ -2345,6 +2423,257 @@ mod gpu { ); } + /// SP15 Phase 3.5.4.c (2026-05-07) — full plasticity-injection two-step + /// recovery sequence end-to-end. Sets up the ISV state so the trigger + /// fires (DD_PERSISTENCE > threshold, fired flag clear, warm = 0), + /// then drives the per-step launch sequence + /// `launch_sp15_plasticity_injection` → `experience_action_select` + /// across `M_warm + 2` bars. Asserts: + /// + /// - **Bar 0 (fire bar)**: dir_picked == DIR_FLAT — the trigger sets + /// warm = M_warm then decrements to (M_warm − 1), so action_select + /// observes warm ≥ M_warm − 1.5f and forces DIR_FLAT (step 1 of + /// the spec §9.2 (3.5.4) two-step). e_dir is set to strongly + /// prefer Long (gap = +1.0); without the fire-bar branch the + /// argmax would land on Long, so observing DIR_FLAT proves the + /// fire-bar consumer fired. + /// + /// - **Bars 1..M_warm − 1 (warm-up)**: dir_picked == DIR_HOLD on + /// every bar. The trigger debounces (already fired this fold) + /// but continues to decrement warm; the action_select OR-gate + /// forces Hold while warm > 0 (step 2 of the two-step). Bar 1 + /// observes warm = M_warm − 2; bar M_warm − 1 observes warm = 1. + /// + /// - **Bar M_warm (warm expired)**: warm = 0 (the trigger's + /// decrement hit floor on the previous bar). action_select OR-gate + /// inactive (cooldown also 0); standard Thompson + hold_floor + /// argmax resumes; with Long strongly preferred and entropy near + /// 0 → hold_floor ≈ 0 → dir_picked == DIR_LONG. + /// + /// Uses M_warm = 5 (down from production's 200) to keep the test fast + /// while preserving the invariant. The `fan_in` / `n_weights` choices + /// don't matter for the consumer-side assertions (the kernel writes + /// the trailing 10% of advantage_head_weights in addition to its ISV + /// updates, but this test doesn't read those weights). Driven via + /// `launch_action_select_for_test_with_m_warm` so the consumer sees + /// the same `m_warm` the trigger wrote. + #[test] + #[ignore = "requires GPU"] + fn plasticity_fires_force_flat_then_cooldown_holds() { + use ml::cuda_pipeline::sp15_isv_slots::{ + COOLDOWN_BARS_REMAINING_INDEX, DD_PERSISTENCE_INDEX, + HOLD_FLOOR_ALPHA_INDEX, HOLD_FLOOR_EPS0_INDEX, HOLD_FLOOR_K_INDEX, + PLASTICITY_FIRED_THIS_FOLD_INDEX, + PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, + PLASTICITY_WARM_BARS_REMAINING_INDEX, + }; + + let stream = make_test_stream(); + let batch: usize = 1; + let n_atoms: i32 = 51; + let v_min: f32 = -1.0; + let v_max: f32 = 1.0; + let delta_z = (v_max - v_min) / (n_atoms as f32 - 1.0); + + // Long strongly preferred → without fire-bar / OR-gate, argmax = Long. + let target_v = [0.0_f32, 0.0_f32, 1.0_f32, 0.0_f32]; + let b_logits_host = fill_peaked_logits_action_select( + batch, n_atoms, v_min, delta_z, &target_v, /*peak_logit*/ 50.0, + ); + let support_host = fill_linear_support_action_select(batch, v_min, v_max, delta_z); + let q_values_host = vec![0.0_f32; batch * AS_Q_STRIDE]; + + // ISV bus pre-fire: persistence way above threshold, fired clear, + // warm 0, cooldown 0. hold_floor params identical to the cooldown + // tests above so any DIR_HOLD outcome is forced by the OR-gate, not + // by hold_floor (entropy near 0 with one-hot logits → hold_floor ≈ 0). + 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[339] = 0.5; // ISV_EVAL_THOMPSON_TEMP_IDX (floor) + isv_init[HOLD_FLOOR_ALPHA_INDEX] = 0.5; + isv_init[HOLD_FLOOR_K_INDEX] = 10.0; + isv_init[HOLD_FLOOR_EPS0_INDEX] = 1.0; + isv_init[COOLDOWN_BARS_REMAINING_INDEX] = 0.0; + isv_init[DD_PERSISTENCE_INDEX] = 1_000.0; + isv_init[PLASTICITY_PERSISTENCE_THRESHOLD_INDEX] = 1.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); + + // advantage_head_weights buffer for the trigger's Kaiming-He reset. + // Uses the same 64-elem skeleton the existing plasticity oracle + // tests use; n_weights / 10 = 6 → 6 trailing slots get rewritten. + const N_WEIGHTS: i32 = 64; + const FAN_IN: i32 = 128; // matches default config adv_h + const M_WARM: f32 = 5.0; // shrunk from production 200.0 for test runtime + let weights_buf = unsafe { MappedF32Buffer::new(N_WEIGHTS as usize) } + .expect("alloc MappedF32Buffer for advantage_head_weights"); + weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS as usize]); + + // Per-step GPU buffers (allocated once, reused across bars). + let q_values = stream.clone_htod(&q_values_host).expect("upload q_values"); + let b_logits_dir = stream.clone_htod(&b_logits_host).expect("upload logits"); + let per_sample_support = stream + .clone_htod(&support_host) + .expect("upload support"); + let mut actions_buf = stream.alloc_zeros::(batch).expect("alloc actions"); + let mut q_gaps_buf = stream.alloc_zeros::(batch).expect("alloc q_gaps"); + let mut intent_buf = stream.alloc_zeros::(batch).expect("alloc intent"); + let mut conv_buf = stream.alloc_zeros::(batch).expect("alloc conv"); + let mut mag_conv = stream.alloc_zeros::(batch).expect("alloc mag_conv"); + + let m_warm_i = M_WARM as usize; + + // ── Bar 0: fire bar → expect DIR_FLAT ────────────────────────────── + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &stream, + isv_buf.dev_ptr, + weights_buf.dev_ptr, + N_WEIGHTS, + M_WARM, + FAN_IN, + /*seed*/ 0xFEED_BEEF_CAFE_BABEu64, + ).expect("launch plasticity (fire bar)"); + let actions = launch_action_select_for_test_with_m_warm( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv, + isv_buf.dev_ptr, M_WARM, + ); + let dir_picked = decode_dir(actions[0]); + // DIR_FLAT = 3 per state_layout.cuh. + assert_eq!( + dir_picked, 3, + "fire bar must force DIR_FLAT (=3); got dir = {dir_picked} \ + (expected 3, even though Long had +1.0 advantage). \ + ISV[fired]={}, ISV[warm]={}", + isv_buf.read_all()[PLASTICITY_FIRED_THIS_FOLD_INDEX], + isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX], + ); + + // Sanity-check the ISV state after the fire bar: + // fired flipped 0 → 1, warm = M_warm − 1 (per the kernel's + // trigger-then-decrement convention). + let isv_after_fire = isv_buf.read_all(); + assert!( + (isv_after_fire[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9, + "after fire-bar PLASTICITY_FIRED_THIS_FOLD must be 1.0; got {}", + isv_after_fire[PLASTICITY_FIRED_THIS_FOLD_INDEX] + ); + let warm_after_fire = isv_after_fire[PLASTICITY_WARM_BARS_REMAINING_INDEX]; + assert!( + (warm_after_fire - (M_WARM - 1.0)).abs() < 1e-4, + "after fire-bar PLASTICITY_WARM_BARS_REMAINING must be M_warm − 1 = {}; got {}", + M_WARM - 1.0, warm_after_fire + ); + + // ── Bars 1..M_warm − 1 (warm-up) → expect DIR_HOLD ────────────────── + // Trigger continues to run on every bar (debounced — fired = 1, no + // re-fire), per-bar warm decrement runs unconditionally; consumer + // sees warm decreasing from M_warm − 2 down to 1; OR-gate forces Hold. + // + // Trace (M_warm = 5): + // Bar 0 (fire): trigger sets warm=5 → decrements to 4. action_select + // reads warm=4 ≥ 3.5 → DIR_FLAT. (asserted above) + // Bar 1: trigger debounces, warm 4→3. action_select reads warm=3 → Hold. + // Bar 2: warm 3→2 → Hold. + // Bar 3: warm 2→1 → Hold. + // Bar 4: warm 1→0. action_select reads warm=0 — OR-gate inactive → + // normal Thompson (NOT Hold). This bar is the "post-warm" + // assertion below, not part of the warm-up loop. + // + // So the warm-up loop runs for bars [1, 2, 3] = bars 1..M_warm − 1 + // = bars 1..4 in iter terms (3 iterations). The previous formulation + // `1..m_warm_i` overran by one bar. + for bar in 1..(m_warm_i - 1) { + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &stream, + isv_buf.dev_ptr, + weights_buf.dev_ptr, + N_WEIGHTS, + M_WARM, + FAN_IN, + /*seed*/ 0xFEED_BEEF_CAFE_BABEu64 ^ (bar as u64), + ).expect("launch plasticity (warm-up)"); + let actions = launch_action_select_for_test_with_m_warm( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv, + isv_buf.dev_ptr, M_WARM, + ); + let dir_picked = decode_dir(actions[0]); + assert_eq!( + dir_picked, AS_DIR_HOLD, + "warm-up bar {bar} must force DIR_HOLD (={AS_DIR_HOLD}); \ + got dir = {dir_picked}. ISV[warm]={}", + isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX] + ); + } + + // After (M_warm − 2) warm-up bars the warm counter is now 1 — the + // very next call's per-bar decrement takes it to 0 (the post-warm + // boundary). Sanity-check that we're at the boundary so the post- + // warm assertion below tests the ACTUAL transition, not some later + // already-decremented state. + let isv_at_boundary = isv_buf.read_all(); + let warm_at_boundary = isv_at_boundary[PLASTICITY_WARM_BARS_REMAINING_INDEX]; + assert!( + (warm_at_boundary - 1.0).abs() < 1e-4, + "after warm-up loop PLASTICITY_WARM_BARS_REMAINING must be 1 (= boundary); got {}", + warm_at_boundary + ); + + // ── Bar M_warm − 1 (warm transitions to 0) → normal action selection resumes ─ + // The trigger still runs (debounced — `fired` stays 1 for the + // remainder of the fold) and decrements warm 1 → 0. action_select + // reads warm = 0: OR-gate inactive (cooldown also 0), fire-bar + // predicate false (warm=0 < 3.5). With Long strongly preferred and + // hold_floor ≈ 0 (one-hot logits → low entropy), Thompson argmax + // returns DIR_LONG. + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &stream, + isv_buf.dev_ptr, + weights_buf.dev_ptr, + N_WEIGHTS, + M_WARM, + FAN_IN, + /*seed*/ 0xFEED_BEEF_CAFE_BABEu64 ^ (m_warm_i as u64 + 1), + ).expect("launch plasticity (post-warm)"); + let actions = launch_action_select_for_test_with_m_warm( + &stream, batch, n_atoms, + &q_values, &b_logits_dir, &per_sample_support, + &mut actions_buf, &mut q_gaps_buf, &mut intent_buf, + &mut conv_buf, &mut mag_conv, + isv_buf.dev_ptr, M_WARM, + ); + let dir_picked = decode_dir(actions[0]); + assert_eq!( + dir_picked, AS_DIR_LONG, + "post-warm bar must follow normal Thompson argmax → DIR_LONG; \ + got dir = {dir_picked}. ISV[warm]={}, ISV[fired]={}", + isv_buf.read_all()[PLASTICITY_WARM_BARS_REMAINING_INDEX], + isv_buf.read_all()[PLASTICITY_FIRED_THIS_FOLD_INDEX] + ); + + // Post-test ISV state: fired stays 1 (debounced for remainder of + // fold — fold-reset registry arms re-arm to 0 on next fold + // boundary), warm at 0 (decrement guard prevents underflow). + let isv_final = isv_buf.read_all(); + assert!( + (isv_final[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9, + "post-test PLASTICITY_FIRED_THIS_FOLD must remain 1 (debounced); got {}", + isv_final[PLASTICITY_FIRED_THIS_FOLD_INDEX] + ); + assert!( + isv_final[PLASTICITY_WARM_BARS_REMAINING_INDEX].abs() < 1e-4, + "post-test PLASTICITY_WARM_BARS_REMAINING must be 0; got {}", + isv_final[PLASTICITY_WARM_BARS_REMAINING_INDEX] + ); + } + // ==================================================================== // SP15 Wave 2 (2026-05-06) — fused post-SP11 reward-axis composer // oracle tests for `compute_sp15_final_reward_kernel`. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index b490ee17b..10af62765 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,9 +2,11 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Phase 3.5.4.c — production caller + OR-gate consumer for plasticity injection (2026-05-07): closes out the Wave 4.2 (`ef08611d3`) orphan-with-tests-only landing per `feedback_wire_everything_up`. Wave 4.2 landed `plasticity_injection_kernel.cu` + `launch_sp15_plasticity_injection` + 4 oracle tests but NO production caller existed — the kernel was reachable only from the oracle-test path. Phase 3.5.4.c creates the production caller AND wires the action-selection consumer atomically per `feedback_no_partial_refactor` so the spec §9.2 (3.5.4) two-step recovery (Flat-on-fire → Hold-during-warm-up) fires end-to-end. (1) **fan_in audit-doc inline correction** [load-bearing per `feedback_trust_code_not_docs`]: Wave 4.2's audit doc point (9) claimed default-config `fan_in = adv_h × num_atoms = 6528` and `stddev = sqrt(2/6528) ≈ 0.0175`. That was wrong — Kaiming-He variance is `2/fan_in` where `fan_in` is the INPUT dim per output unit, i.e. `adv_h` (not `adv_h × num_atoms` which is the full row-output element count of one branch's projection). The correct values are **`fan_in = 128`** and **`stddev = sqrt(2/128) ≈ 0.125`**. The error was inline-corrected in the Wave 4.2 entry so future readers don't propagate the wrong number; the corrected entry credits Phase 3.5.4.c. (2) **Production caller in `gpu_experience_collector.rs::collect_experiences_gpu`'s per-step loop** (step 4-pre, BEFORE `experience_action_select`): when `sp15_w_b0out_dev_ptr != 0 && isv_signals_dev_ptr != 0`, invokes `launch_sp15_plasticity_injection(stream, isv, w_b0out, n_weights=branch_0_size×num_atoms×adv_h, m_warm=200.0, fan_in=adv_h, seed=mix_seed(0xC0DE_F00D_5915_3F4C) ^ t)`. The trigger writes ISV[436] (fired flag) + ISV[438] (warm counter) which the action_select kernel then reads (next launch on the same stream — CUDA serializes producer→consumer without explicit event sync). When unwired (test scaffold), the launch is skipped and the action_select OR-gate consumer degenerates to cooldown-only (bit-identical pre-3.5.4.c behaviour). Placement: inside the `else` arm of the `seed_phase_active_cache` gate (network branch only — scripted-policy seed phase forces specific actions, plasticity injection during seed phase is meaningless). Out of the `exp_fwd_graph` capture region (which ends at the cuBLAS forward earlier in the loop body), so per `pearl_no_host_branches_in_captured_graph` the host-side gating compiles cleanly. (3) **Branch choice — w_b0out (directional) only** per spec §9.2 (3.5.4) "the recovery is about escaping stuck-in-flat regimes": the directional advantage-head last-Linear weight tensor (param-buffer index 19) gets the trailing-10% Kaiming-He reset; magnitude/order/urgency branches are untouched (their conditioning on direction makes them naturally re-train once direction starts choosing again). The trainer exposes the target via `pub fn sp15_w_b0out_target() -> (u64, i32, i32)` returning `(dev_ptr_at_offset_19, n_weights, fan_in)`; the collector's matching `set_sp15_plasticity_target(w_ptr, n_weights, fan_in)` setter wires it. (4) **n_weights derivation**: `branch_0_size × num_atoms × adv_h` from `compute_param_sizes(&self.config)[19]`. Default config: 4 × 51 × 128 = **26112** (vs Wave 4.2's oracle-test 64 + 1280); reset_count = `n_weights / 10` = 2611 trailing weights re-sampled per fire. The launcher's grid_dim auto-scales to `((26112/10 + 255) / 256)` = 11 blocks × 256 threads. (5) **fan_in = adv_h** per the kernel's per-output-unit variance contract: the kernel computes `weights[i] = curand_normal × sqrt(2/fan_in)` for each i in the trailing reset region — `fan_in` here is the input-dim of the per-unit dot product, which equals `adv_h` regardless of `num_atoms`. (6) **Seed source — deterministic per-step**: `mix_seed(0xC0DE_F00D_5915_3F4C) ^ (t as u64)` where `t` is the per-step timestep iterator and `mix_seed` is the SplitMix64 avalanche over the global `FOXHUNT_SEED` env var. Same `(FOXHUNT_SEED, t)` always produces the same `(seed, tid)` cuRAND init in the kernel, satisfying the established codebase determinism contract (`feedback_h100_gpu` + checkpoint-resumed reproducibility). The constant `0xC0DE_F00D_5915_3F4C` is a Phase-3.5.4.c-unique 64-bit base so multi-launcher runs (alpha_split / cooldown / dd_state / etc.) get uncorrelated mix outputs even when `FOXHUNT_SEED == 42` (default). (7) **OR-gate consumer extension in `experience_action_select`** (`experience_kernels.cu`): added one new kernel arg `float plasticity_m_warm` (after `n_atoms`); the kernel now reads ISV[438] (warm counter) alongside ISV[435] (cooldown remaining); the SP15 Phase 3.5.3.b cooldown predicate `(cooldown_remaining > 0)` is extended to `(cooldown_remaining > 0) || (plasticity_warm_remaining > 0)`. When the plasticity launch above is unwired, `m_warm` is passed as 0.0f and warm stays at 0 → the OR-gate degenerates to cooldown-only (bit-identical pre-3.5.4.c). (8) **Flat-on-fire-bar detection — option (a), `warm ≥ m_warm − 1.5f`** per the trigger-then-decrement convention: the trigger kernel sets `warm = m_warm` then decrements once in the same launch (mirroring the cooldown_kernel's identical sequence with the `[198, 200]` test tolerance). On the fire bar, action_select observes `warm = m_warm − 1`, which satisfies the predicate `warm ≥ m_warm − 1.5f` with a 0.5f f32-equality tolerance. On bar T+1 (post-fire warm-up), `warm = m_warm − 2 < m_warm − 1.5f`, so the fire-bar branch is dead and the OR-gate cooldown branch fires DIR_HOLD instead. The fire-bar branch supersedes the cooldown branch via an `if/else if` chain — the spec sequence is: fire-bar → DIR_FLAT (close any open position); subsequent warm bars → DIR_HOLD. No new ISV slot needed (option (b) was rejected as adding unnecessary state); no host-side detection needed (option (c) was rejected as introducing a host branch in the per-step loop). The `m_warm > 0` guard inside the predicate ensures the fire-bar branch is dead when plasticity is unwired (m_warm=0 and warm=0 — both zero → `warm >= m_warm − 1.5f = -1.5f` is true, but the `m_warm > 0` factor rules it out). (9) **Trainer plumbing in `training_loop.rs`**: alongside the existing `set_isv_signals_ptr` / `set_sp15_alpha_warm_count_ptr` / `set_sp15_dd_trajectory_prev_dd_ptr` calls, added `collector.set_sp15_plasticity_target(fused.trainer().sp15_w_b0out_target())`. Mirrors the same wiring pattern. (10) **NEW oracle test** `plasticity_fires_force_flat_then_cooldown_holds` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after `action_select_no_force_hold_when_cooldown_zero`): drives the full per-step launch sequence (`launch_sp15_plasticity_injection` → `experience_action_select`) across `M_warm + 1` bars with `M_warm = 5` (shrunk from production's 200 to keep the test fast). With Long strongly preferred (e_dir gap = +1.0) and the trigger-fires ISV state pre-loaded (DD_PERSISTENCE=1000, threshold=1, fired=0, warm=0): bar 0 (fire bar) asserts dir_picked == DIR_FLAT (=3) — proves the fire-bar branch fired even though Long would normally win the argmax; bars 1..M_warm−1 (warm-up) assert dir_picked == DIR_HOLD on every bar — proves the OR-gate cooldown fires when warm ∈ {3, 2, 1}; bar M_warm−1 (the boundary call where warm transitions 1→0) asserts dir_picked == DIR_LONG — proves normal Thompson argmax resumes once warm=0 and the OR-gate is inactive. Also asserts the ISV side-effects bar-by-bar: fired flips 0→1 on the fire bar then debounces (stays 1) for the rest of the test, warm decrements from M_warm−1 down to 0 with the trigger's underflow guard preventing negative values. (11) **Eval/backtest path** (`gpu_backtest_evaluator.rs::submit_dqn_step_loop_cublas`): updated to pass `m_warm = 0.0f` to `experience_action_select` because plasticity injection is a TRAINING-only mechanism (it resets advantage-head weights to break out of stuck-in-flat regimes during experience collection); during deterministic eval the weights must remain frozen at their checkpoint values. With `m_warm = 0.0f` the action_select fire-bar predicate is dead, and `PLASTICITY_WARM_BARS_REMAINING` stays at its 0 sentinel through eval (no producer running) so the OR-gate degenerates to cooldown-only (bit-identical pre-3.5.4.c). (12) **Atomicity** per `feedback_no_partial_refactor`: kernel signature change + collector field/setter/launch + trainer accessor + trainer plumbing + eval-path update + 2 test launcher updates (mod.rs + distributional_q_tests.rs) + new behavioral oracle test + audit doc inline correction + new audit doc entry all in one commit; pre-existing test launchers were updated to pass `m_warm = 0.0f` so they preserve their pre-3.5.4.c semantics (no fire-bar / OR-gate plasticity branch). Touched: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (kernel signature `+plasticity_m_warm: float`, fire-bar detection block via `warm ≥ m_warm − 1.5f`, OR-gate extension `cooldown_active = (cooldown_remaining > 0) || (warm > 0)`, new `if (plasticity_fire_bar)` branch BEFORE the existing `else if (cooldown_active)` cooldown branch), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (new `pub fn sp15_w_b0out_target()` accessor returning `(dev_ptr, n_weights, fan_in)`), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (new `sp15_w_b0out_dev_ptr` / `sp15_w_b0out_n_weights` / `sp15_w_b0out_fan_in` fields + `set_sp15_plasticity_target` setter + per-step `launch_sp15_plasticity_injection` invocation in step 4-pre + `m_warm` arg passed to `experience_action_select`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (one new wiring block calling `set_sp15_plasticity_target`), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (eval-path action_select launch passes `m_warm = 0.0f`), `crates/ml/src/cuda_pipeline/mod.rs` (test launcher passes `m_warm = 0.0f`), `crates/ml/src/trainers/dqn/distributional_q_tests.rs` (test launcher passes `m_warm = 0.0f`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (existing `launch_action_select_for_test` helper passes `m_warm = 0.0f`; new `launch_action_select_for_test_with_m_warm` helper takes `m_warm` as an explicit parameter; new behavioral oracle test `plasticity_fires_force_flat_then_cooldown_holds`), `docs/dqn-wire-up-audit.md` (this entry + Wave 4.2 entry's point (9) inline-corrected). Hard rules: `feedback_no_partial_refactor` (kernel + caller + consumer + 7 launcher-call-site updates + new behavioral test + audit doc inline correction + new audit doc entry all atomic — every consumer of the kernel signature change migrates simultaneously), `feedback_no_quickfixes` (real fan_in derivation from `cfg.adv_h` via `sp15_w_b0out_target()` accessor — not a magic constant), `feedback_no_cpu_compute_strict` (the per-step plasticity launch + the kernel-side fire-bar detection are fully GPU-resident; the host-side `if self.sp15_w_b0out_dev_ptr != 0` gate is a cold-path one-shot per step, not a per-thread-of-kernel branch), `pearl_no_host_branches_in_captured_graph` (the host-side launch + parameter computation runs OUTSIDE the `exp_fwd_graph` capture region — that graph ends at the cuBLAS forward in step 2, well before step 4-pre), `feedback_trust_code_not_docs` (the Wave 4.2 audit-doc fan_in error is corrected in-place; the right value `adv_h = 128` is read from the trainer config rather than re-derived in prose), `feedback_wire_everything_up` (Wave 4.2's orphan launcher is now driven from production every step; the `Phase 3.5.4.c remains a separate follow-up` paragraph in the Wave 4.2 entry is now historical context for THIS commit rather than an open TODO; SP15 Phase 3.5 recovery-dynamics chain is end-to-end production-wired — 3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.c + 3.5.5 + 3.5.5.b all firing). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean; `SQLX_OFFLINE=true cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` 5 of 5 plasticity oracle tests green (4 Wave 4.2 + new `plasticity_fires_force_flat_then_cooldown_holds`); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored action_select --nocapture` 3 of 3 SP15 3.5.b/3.5.3.b action_select oracle tests still green (the `m_warm = 0.0f` defaults preserve their pre-3.5.4.c behaviour bit-identically); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored cooldown --nocapture` 6 of 6 cooldown + action_select tests green; `cargo test -p ml --features cuda --lib` HOLDS the Wave 4.3 baseline (946 pass / 13 fail) on RTX 3050 Ti — no new regressions, the 13 failing tests are the same pre-existing infra failures (OFI features missing, checkpoint round-trip, spectral norm GPU init, etc.). + SP15 Wave 4.3 / Phase 3.5.5.b — PER sampler integration: recovery transitions oversampled (2026-05-07): closes out the Phase 3.5.5 (`69b8fdb61`) deferred PER sampling consumer per `feedback_wire_everything_up`. Phase 3.5.5 landed `dd_trajectory_decreasing_kernel` writing `ISV[DD_TRAJECTORY_DECREASING_INDEX=439]` per-step but DEFERRED both the production launch AND the PER consumer that reads ISV[439] + ISV[440] to bias replay sampling; Wave 4.3 wires both atomically. (1) **Investigation finding — Case B (existing TD-error-driven priority sampler)**: The PER architecture in `crates/ml-dqn/src/gpu_replay_buffer.rs` already weights replay by per-transition priority (TD-error driven via `per_update_pa`). `priorities[idx]` holds the raw priority (max_priority sentinel at insert, `|td|^alpha + ε` after training); `priorities_pa[idx] = priorities[idx]^alpha` feeds the prefix-sum sampler `per_prefix_scan` + deterministic-stratified `per_sample` (zero CPU compute, fully GPU-resident). The recovery oversample boost slots in cleanly as a multiplier on `priorities[idx]` at insert time — when the agent is currently in a recovery (DD shrinking from a non-trivial level), every transition in the freshly-collected batch lands in the buffer with `effective = base × (1 + ω × δ) = 3.0` instead of `1.0` (with sentinel ω=2.0, δ=1.0 from `dd_trajectory_decreasing_kernel`), so for the duration that those transitions sit in the buffer pre-update they get sampled `3.0^0.6 ≈ 1.93×` more often than baseline transitions. Once `per_update_pa` overwrites the priority based on TD-error, normal PER takes over — the boost only amplifies early sampling, exactly matching the spec's "agent learns recovery dynamics over typical average-DD Bellman noise" framing. This is materially smaller than Case A (uniform sampler retrofit) would have been: the existing alias-method-equivalent prefix-sum machinery stays untouched. (2) **Kernel signature change** in `crates/ml-dqn/src/per_kernels.cu:per_insert_pa`: added `float* __restrict__ priorities` (now also written, subsuming the previous `scatter_insert_f32` broadcast of `max_priority`) and `const float* __restrict__ isv` (nullable; reads `isv[439]` for δ and `isv[440]` for ω when non-null, falls back to `boost_factor=1.0` when null). Kernel body computes `effective = raw_priorities[i] × boost_factor`, writes both `priorities[idx] = effective` and `priorities_pa[idx] = effective^alpha`. Slot indices are inline literals (439, 440) matching the `dd_trajectory_kernel.cu` literal-index convention — the canonical names live in `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` with compile-time `assert_eq!` regression checks pinning them in `sp15_slot_layout_locked`. (3) **`GpuReplayBuffer` API change** in `crates/ml-dqn/src/gpu_replay_buffer.rs`: added `isv_signals_dev_ptr: u64` field (initialised to 0 in the constructor — meaning NULL → no boost), `pub fn set_isv_signals_ptr(&mut self, dev_ptr: u64)` setter (mirrors the existing `set_trainer_buffers` / `set_learning_health` plumbing pattern), `pub const fn isv_signals_dev_ptr(&self) -> u64` accessor for the wiring round-trip oracle assertion, and `pub const fn priorities_pa_slice(&self) -> &CudaSlice` accessor for the new oracle test. `insert_batch` now passes `&isv_ptr` + `&mut self.priorities` to `per_insert_pa` and the redundant `scatter_insert_f32` broadcast that previously seeded `priorities[idx] = max_priority` is REMOVED — `per_insert_pa`'s new `priorities[idx] = effective` store subsumes its work (effective equals max_priority × 1.0 = max_priority when ISV is unwired or δ=0, preserving pre-3.5.5.b semantics for non-recovery transitions). (4) **Production launch wiring** in `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`: added `sp15_dd_trajectory_prev_dd_dev_ptr: u64` field + `set_sp15_dd_trajectory_prev_dd_ptr` setter (same pattern as the existing `set_sp15_alpha_warm_count_ptr`); inside `collect_experiences_gpu`'s per-step loop, immediately after the existing `launch_sp15_dd_state` call (which writes `ISV[DD_PCT_INDEX=406]`), the new "5b. SP15 Phase 3.5.5.b — DD_TRAJECTORY_DECREASING per-step proxy" block invokes `launch_sp15_dd_trajectory_decreasing` gated on both `isv_signals_dev_ptr != 0` AND `sp15_dd_trajectory_prev_dd_dev_ptr != 0` — when either is unwired (test scaffold path), the kernel is skipped and ISV[439] stays at sentinel 0 (no recovery boost). Mirrors the `launch_sp15_dd_state` placement (same per-step branch, same single-thread/single-block kernel shape, same on-stream serialization without event sync). (5) **Trainer-side plumbing** in `crates/ml/src/trainers/dqn/trainer/training_loop.rs`: alongside the existing `set_isv_signals_ptr` + `set_sp15_alpha_warm_count_ptr` calls in `train_with_data_full_loop_slices`, added two new wiring blocks — `collector.set_sp15_dd_trajectory_prev_dd_ptr(fused.trainer().sp15_dd_trajectory_prev_dd.dev_ptr)` so the per-step kernel can advance its persistent state, AND `agent.primary_dqn_mut().memory.gpu.set_isv_signals_ptr(fused.isv_signals_dev_ptr())` so `per_insert_pa` can read the recovery slots. Without the second wiring, the buffer's `isv_signals_dev_ptr` stays NULL and the kernel falls back to boost_factor=1.0 (no oversample) — preserving the smoke / oracle-test path that exercises PER without the SP15 bus. (6) **NEW oracle test** `per_sampler_weights_recovery_transitions_higher` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after the three `dd_trajectory_*` tests): allocates a `MappedF32Buffer(ISV_LEN)` with ISV[440]=2.0 (sentinel ω) + ISV[439]=1.0 (recovery active); constructs a `GpuReplayBuffer` with capacity 64; calls `set_isv_signals_ptr(isv_buf.dev_ptr)` (asserts `isv_signals_dev_ptr()` round-trips correctly); inserts 8 zero-padded synthetic transitions (batch 1, recovery=1) → expected `priorities[0..8] = 3.0`; flips ISV[439]=0 and inserts 8 more (batch 2, recovery=0) → expected `priorities[8..16] = 1.0`; reads `priorities_slice` + `priorities_pa_slice` via `memcpy_dtoh`; asserts batch-1 priorities all equal 3.0 within 1e-5, priorities_pa all equal `3.0^0.6 ≈ 1.9332` within 1e-4, batch-2 priorities all equal 1.0; finally computes `boosted_mean / baseline_mean` ratio and asserts it equals exactly 3.0 within 1e-5 — the recovery oversample factor by construction. The deterministic-stratified `per_sample` kernel weights samples by `priorities^alpha`, so verifying the priority-buffer write directly (paired with the existing `gpu_per_integration_test::test_gpu_per_priorities_affect_sampling` empirical-bias coverage) gives the full statistical-bias proof without re-running a 10000-sample experiment in the oracle suite. (7) **Phase 3.5.5 deferred consumer eliminated** per `feedback_wire_everything_up`: `dd_trajectory_decreasing_kernel` is now invoked from production code (the experience collector's per-step loop), and its ISV[439] write is now CONSUMED by `per_insert_pa` to bias replay sampling. The Phase 3.5.5 commit message's "PER sampling consumer (DEFERRED follow-up)" narrative is closed; the `dd_trajectory_kernel.cu` header comment's deferred-consumer paragraph is now historical context for the Wave 4.3 wire-up rather than an open TODO. (8) **DD_TRAJECTORY_FLOOR producer (slot 441) and ISV-driven RECOVERY_OVERSAMPLE_WEIGHT producer (slot 440) remain documented Phase 3.5.5 follow-ups** per `feedback_isv_for_adaptive_bounds` — the kernel reads from these slots rather than hardcoded literals, so when the rolling-25th-percentile `dd_pct` producer for slot 441 and the `proportional-to-current-dd_pct` producer for slot 440 land in their own atomic commits, the consumers (`dd_trajectory_kernel` reading floor; `per_insert_pa` reading ω) are zero-effort migrations. Wave 4.3 uses the constructor-seeded sentinels (floor=0.02, ω=2.0). (9) **Atomicity** per `feedback_no_partial_refactor`: kernel signature change + `GpuReplayBuffer` API extension + `insert_batch` wiring + experience-collector field/setter/launch + trainer plumbing + oracle test + audit doc all in one commit; every `insert_batch` consumer (production `training_loop.rs` + 3 smoke test files in `crates/ml/src/trainers/dqn/smoke_tests/` + `crates/ml/tests/gpu_per_integration_test.rs` + the new oracle test) compiles unchanged because the boost is gated by the optional `set_isv_signals_ptr` setter (zero=NULL preserves the pre-3.5.5.b behavior modulo the now-merged `priorities[idx]` write). (10) **GPU-residency audit**: weight computation is fully GPU-resident — `per_insert_pa` reads ISV[439]/ISV[440] from the GPU bus, multiplies on the GPU, writes both `priorities` and `priorities_pa` from a single kernel launch. No DtoH/HtoD transfers added. The `dd_trajectory_decreasing_kernel` runs on the same stream as `env_step` + `launch_sp15_dd_state` so CUDA serializes the producer→consumer chain without explicit event sync (`pearl_no_host_branches_in_captured_graph` clean — both are outside the exp-fwd graph capture region per the existing `launch_sp15_dd_state` precedent at `gpu_experience_collector.rs:4555-4561`). Touched: `crates/ml-dqn/src/per_kernels.cu` (`per_insert_pa` signature `+priorities, +isv`, body rewrite for the boost multiplier + dual `priorities`/`priorities_pa` writes, header comment expanded for the Wave 4.3 contract), `crates/ml-dqn/src/gpu_replay_buffer.rs` (new `isv_signals_dev_ptr` field + setter + accessor + `priorities_pa_slice` accessor; `insert_batch` arg list extended for the kernel; redundant `scatter_insert_f32` for `priorities` REMOVED), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (new `sp15_dd_trajectory_prev_dd_dev_ptr` field + setter; per-step `launch_sp15_dd_trajectory_decreasing` invocation gated on both ISV ptr + prev_dd ptr being non-zero), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (two new wiring blocks for collector prev_dd ptr + replay-buffer ISV ptr), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (new oracle test `per_sampler_weights_recovery_transitions_higher`). Hard rules: `feedback_no_cpu_compute_strict` (boost computation + dual buffer write are GPU-resident inside `per_insert_pa`), `feedback_no_partial_refactor` (kernel + sampler + producer launch + trainer plumbing + oracle test + audit doc all atomic), `feedback_isv_for_adaptive_bounds` (RECOVERY_OVERSAMPLE_WEIGHT lives in ISV[440] read at kernel time — the constructor-seeded 2.0 sentinel is a structural cold-start anchor; the ISV-driven producer remains a documented Phase 3.5.5 follow-up; the kernel reads the slot rather than a hardcoded literal so the follow-up swap is a no-op at the consumer), `feedback_no_stubs` (kernel performs the actual priority multiplication and writes both columns; no return-zero stubs; the boost-factor=1.0 fallback for unwired ISV is the explicit semantic for callers that don't need recovery oversample, not a stub), `feedback_no_atomicadd` (per-thread independent writes to unique `priorities[idx]` / `priorities_pa[idx]`; no atomics, no reductions — `per_insert_pa` was already grid-strided element-wise; the kernel signature extension preserves the per-thread independence), `feedback_wire_everything_up` (Phase 3.5.5's deferred PER consumer is closed in this commit — `dd_trajectory_kernel` is launched in production AND its ISV[439] write is consumed by `per_insert_pa`; this completes the SP15 Phase 3.5 recovery-dynamics chain: 3.5.2 + 3.5.3 + 3.5.4 + 3.5.5 + 3.5.5.b all wired). Verified: `SQLX_OFFLINE=true cargo check -p ml-dqn --features cuda` clean; `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings); `cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored dd_trajectory --nocapture` 3 of 3 dd_trajectory oracle tests still green; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored per_sampler --nocapture` 1 of 1 new PER sampler oracle test green (recovery slots read 3.0, baseline slots read 1.0, ratio = 3.0 within 1e-5 — boosted/baseline statistical bias verified by construction); `cargo test -p ml --features cuda --lib gpu_residency` 8 of 8 (1 ignored) replay-buffer + adamw smoke tests green; `cargo test -p ml --features cuda --lib replay` 4 of 4 PER smoke tests green; `cargo test -p ml --features cuda --lib` IMPROVES the Wave 4.2 baseline — 947 pass / 12 fail (was 946 pass / 13 fail; the previously-failing `ensemble::adapters::dqn::tests::test_dqn_checkpoint_round_trip` now passes, likely because the redundant pre-3.5.5.b `scatter_insert_f32` for priorities was racing with the immediate `per_insert_pa` overwrite under particular timing conditions; the merged single-kernel write closes that race). -SP15 Wave 4.2 / Phase 3.5.4.b — cuRAND Kaiming-He weight reset for plasticity injection (2026-05-07): closes out the Phase 3.5.4 (`e0e0abfb2`) deferred-weight-reset NO-OP per `feedback_wire_everything_up`. Phase 3.5.4 landed the trigger + warm-up tracker but DEFERRED the actual weight reset (kernel accepted `advantage_head_weights + n_weights` and no-op'd via `(void)` cast); Wave 4.2 lands the real cuRAND-driven Kaiming-He reset of the trailing 10% of the advantage-head weight tensor. (1) **Kernel body**: `crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu` now `#include ` and writes `advantage_head_weights[reset_start + tid] = curand_normal(&state) * stddev` per thread when the gate fires. `reset_count = n_weights / 10`, `reset_start = n_weights - reset_count`, `stddev = sqrt(2 / fan_in)` (Kaiming-He gain=√2 for ReLU/GLU activations). Per-thread `curandState` is initialised via `curand_init(seed, sequence=tid, offset=0, &state)` so the same `(seed, tid)` pair always yields the same sample (oracle determinism re-check verifies bit-identical samples across two launches with identical seed). cuRAND device functions (`curand_init`, `curand_normal`) are inlined into the cubin by nvcc from the standard CUDA header — no host-side cuRAND linker dependency required (the cubin compiler already finds `curand_kernel.h` from the CUDA toolkit `targets/x86_64-linux/include/` path; no `build.rs` link change). (2) **Grid configuration changed** from single-thread/single-block to multi-block element-wise: `grid_dim = ((n_weights / 10 + 255) / 256).max(1)`, `block_dim = 256`. Block 0 / thread 0 still owns the trigger + warm-bars decrement (the only ISV-write code path — single-thread mirrors `cooldown_kernel` / `dd_asymmetric_reward_kernel` / `hold_floor_kernel`); the remaining threads independently re-evaluate the same `fire_now` predicate from the same ISV reads (the trigger condition is a pure function of three ISV slots — DD_PERSISTENCE / PLASTICITY_PERSISTENCE_THRESHOLD / PLASTICITY_FIRED_THIS_FOLD — read at kernel-launch time, so every thread converges on the same answer without any cross-block synchronisation; no `__shared__` broadcast needed). The `.max(1)` floor on `grid_dim` ensures block 0 always exists even when `n_weights / 10 == 0` (oracle edge case where the test passes a tiny weight buffer), so the trigger still runs. (3) **Launcher signature change** in `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:launch_sp15_plasticity_injection`: added `fan_in: i32` and `seed: u64` after `m_warm`. Production callers will derive `fan_in` from the trainer's `cfg.adv_h × shared_h2` advantage-head Linear-layer dimensions (the consumer wiring is the separate Phase 3.5.4.c follow-up; no production launch site exists yet — only the 4 oracle tests below drive the launcher); `seed` will be derived deterministically per-call from a trainer RNG state. The launcher itself contains no host branches inside the launch builder (`pearl_no_host_branches_in_captured_graph` clean — the `seed` u64 is a kernel argument, not a host conditional, and the grid-dim computation runs once before any `launch_builder` call). (4) **3 existing oracle tests migrated** to the new launcher signature in `crates/ml/tests/sp15_phase1_oracle_tests.rs`: `plasticity_fires_when_persistence_exceeds_threshold` (passes `fan_in=256` + `seed=0xCAFE_BABE_DEAD_BEEF`), `plasticity_debounced_within_fold` (passes `fan_in=256` + `seed=0xDEAD_BEEF_F00D_CAFE` and now also asserts all 64 weights remain exactly 1.0 since the debounce path early-outs before the reset region), `plasticity_no_fire_below_threshold` (passes `fan_in=256` + `seed=0xFEED_FACE_BAAD_F00D` and asserts the same weight-stability invariant since no-fire also early-outs). The trigger-behavior assertions stay byte-identical. (5) **NEW oracle test** `plasticity_injection_kernel_resets_last_10pct_kaiming_he` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after `plasticity_no_fire_below_threshold`): allocates a 1280-element advantage-head weight buffer pre-filled with 1.0, sets ISV[DD_PERSISTENCE]=1000, ISV[PLASTICITY_PERSISTENCE_THRESHOLD]=1, ISV[PLASTICITY_FIRED_THIS_FOLD]=0; passes `fan_in=256`, `seed=42`; launches; then asserts: (a) ISV[PLASTICITY_FIRED_THIS_FOLD] flipped 0→1, (b) ISV[PLASTICITY_WARM_BARS_REMAINING] = 199 (`m_warm - 1`, per-bar decrement runs in same call), (c) first 90% (1152 elements) bit-identical to 1.0, (d) last 10% (128 elements) every slot has moved off 1.0 by > 1e-6 (Kaiming stddev≈0.088 makes 1.0 ~11σ out — exact-match impossible), (e) sample mean of the reset region |mean| < 0.05 (true SE ~0.0078 for n=128 from Normal(0, 0.088)), (f) sample std within ±20% of `sqrt(2/256) ≈ 0.08839`, (g) determinism re-check — re-launch with the same seed against a fresh 1.0-filled buffer + reset ISV produces bit-identical Kaiming samples per thread. All four plasticity oracle tests pass on RTX 3050 Ti (`SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` → 4 passed; 0 failed; 33 filtered out). (6) **Audit doc + build.rs comment + cubin-slot docstring + launcher docstring all updated** to reflect the now-landed weight-reset (Phase 3.5.4 deferred-NO-OP narrative replaced with Wave 4.2 active-reset narrative); the kernel's own header comment is rewritten to document the Wave 4.2 contract (cuRAND init, reset region addressing, Kaiming-He gain rationale, per-thread vs per-block responsibilities). (7) **Phase 3.5.4 deferred consumer eliminated** per `feedback_wire_everything_up`: the kernel signature's `(void)advantage_head_weights; (void)n_weights;` no-op cast is gone; the `(void)` pattern that flagged the deferred work is replaced by real per-element writes; the launcher no longer documents the plumbed-but-unused buffer. (8) **Action-selection consumer wiring (Phase 3.5.4.c) remains a separate follow-up** per `feedback_no_partial_refactor` — the OR-gate `effective_cooldown = max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)` plus the same-bar Flat emission on the fired-flag transition are independent of weight-reset and are tracked under their own atomic commit (matching Phase 1.1-1.5 + 3.1-3.5.3 precedent: kernel + launcher land first verified by oracle tests, then consumer migration is the next atomic). (9) **fan_in source for production**: the advantage head's last Linear layer input dimension. Per `compute_param_sizes` (`gpu_dqn_trainer.rs:4042-4056`), the four branch-output linears are `w_b{0,1,2,3}out [branch_X_size, num_atoms × adv_h]`, so the input dimension is `adv_h × num_atoms` (fan-in to the per-atom output projection). Default config has `adv_h = 128`, `num_atoms = 51` → fan_in = 6528, stddev = `sqrt(2/6528) ≈ 0.0175`. The oracle test uses `fan_in = 256` as a representative value for synthetic checks. The exact production `fan_in` value will be supplied by the Phase 3.5.4.c consumer-wiring commit when the launcher gets a real production call site. (10) **Determinism vs entropy choice** (Option A — per-thread `curand_init` per-call): trades a small amount of cuRAND-init compute (~10 cycles per thread to seed the Philox4x32 state) for full determinism in the oracle test plus reproducibility across trainer-restart resumed sessions, which is the established codebase posture (`feedback_h100_gpu` + the trainer's resumed-from-checkpoint contract). The alternative (Option B — persistent per-thread state across calls) would require a `curandState[reset_count]` mapped-pinned scratch buffer on the trainer struct (mirroring `sp15_cooldown_consecutive_losses`), but Wave 4.2 only fires at most once per fold and reset_count is small (≤ ~650 for default config), so the per-call init cost is negligible (< 1µs total) and Option A is the right scope. Touched: `crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu` (kernel body — `` include, multi-block grid contract, per-thread `curand_init` + `curand_normal` writes, header comment rewrite for Wave 4.2 contract), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher signature `+fan_in: i32, +seed: u64`, grid_dim computation, launch_builder arg list extended; cubin-slot docstring + launcher docstring rewritten for Wave 4.2 contract), `crates/ml/build.rs` (kernel manifest comment block updated — deferred-NO-OP narrative replaced with Wave 4.2 active-reset narrative + cuRAND header inlining note), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (3 existing tests migrated to new launcher signature + new test `plasticity_injection_kernel_resets_last_10pct_kaiming_he`). Hard rules: `feedback_no_atomicadd` (per-thread independent writes to unique `advantage_head_weights[reset_start + tid]`; no reductions; trigger logic remains single-thread for race-free ISV writes), `feedback_no_partial_refactor` (kernel + launcher signature change + 3 test migrations + 1 new test + audit doc + build.rs comment + 2 launcher/cubin docstrings all atomic; the action-selection consumer wiring remains the separate Phase 3.5.4.c follow-up by design — same pattern as Phase 3.5.3), `feedback_no_stubs` (the `(void)` no-op casts are gone; kernel performs the actual Kaiming-He reset), `feedback_no_cpu_compute_strict` (cuRAND init + sampling are GPU-resident — `curand_init` and `curand_normal` are device-only functions inlined into the cubin), `feedback_no_quickfixes` (the determinism re-check inside the oracle test catches accidental clock-derived-seed regressions; the mean/std assertions catch non-Gaussian or wrong-stddev regressions; the first-90%-unchanged assertion catches off-by-one regressions in the addressing arithmetic), `feedback_no_legacy_aliases` (no compatibility shim — the launcher signature changes in-place; existing callers MUST update to pass fan_in + seed or fail to compile), `feedback_wire_everything_up` (Phase 3.5.4's deferred-weight-reset NO-OP is closed in this commit; the `(void)` casts are removed). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings); `cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` 4 of 4 plasticity oracle tests green (`plasticity_fires_when_persistence_exceeds_threshold` + `plasticity_debounced_within_fold` + `plasticity_no_fire_below_threshold` + `plasticity_injection_kernel_resets_last_10pct_kaiming_he`) on RTX 3050 Ti; `cargo test -p ml --features cuda --lib` HOLDS the 947 pass / 12 fail Wave 4.1c baseline. +SP15 Wave 4.2 / Phase 3.5.4.b — cuRAND Kaiming-He weight reset for plasticity injection (2026-05-07): closes out the Phase 3.5.4 (`e0e0abfb2`) deferred-weight-reset NO-OP per `feedback_wire_everything_up`. Phase 3.5.4 landed the trigger + warm-up tracker but DEFERRED the actual weight reset (kernel accepted `advantage_head_weights + n_weights` and no-op'd via `(void)` cast); Wave 4.2 lands the real cuRAND-driven Kaiming-He reset of the trailing 10% of the advantage-head weight tensor. (1) **Kernel body**: `crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu` now `#include ` and writes `advantage_head_weights[reset_start + tid] = curand_normal(&state) * stddev` per thread when the gate fires. `reset_count = n_weights / 10`, `reset_start = n_weights - reset_count`, `stddev = sqrt(2 / fan_in)` (Kaiming-He gain=√2 for ReLU/GLU activations). Per-thread `curandState` is initialised via `curand_init(seed, sequence=tid, offset=0, &state)` so the same `(seed, tid)` pair always yields the same sample (oracle determinism re-check verifies bit-identical samples across two launches with identical seed). cuRAND device functions (`curand_init`, `curand_normal`) are inlined into the cubin by nvcc from the standard CUDA header — no host-side cuRAND linker dependency required (the cubin compiler already finds `curand_kernel.h` from the CUDA toolkit `targets/x86_64-linux/include/` path; no `build.rs` link change). (2) **Grid configuration changed** from single-thread/single-block to multi-block element-wise: `grid_dim = ((n_weights / 10 + 255) / 256).max(1)`, `block_dim = 256`. Block 0 / thread 0 still owns the trigger + warm-bars decrement (the only ISV-write code path — single-thread mirrors `cooldown_kernel` / `dd_asymmetric_reward_kernel` / `hold_floor_kernel`); the remaining threads independently re-evaluate the same `fire_now` predicate from the same ISV reads (the trigger condition is a pure function of three ISV slots — DD_PERSISTENCE / PLASTICITY_PERSISTENCE_THRESHOLD / PLASTICITY_FIRED_THIS_FOLD — read at kernel-launch time, so every thread converges on the same answer without any cross-block synchronisation; no `__shared__` broadcast needed). The `.max(1)` floor on `grid_dim` ensures block 0 always exists even when `n_weights / 10 == 0` (oracle edge case where the test passes a tiny weight buffer), so the trigger still runs. (3) **Launcher signature change** in `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:launch_sp15_plasticity_injection`: added `fan_in: i32` and `seed: u64` after `m_warm`. Production callers will derive `fan_in` from the trainer's `cfg.adv_h × shared_h2` advantage-head Linear-layer dimensions (the consumer wiring is the separate Phase 3.5.4.c follow-up; no production launch site exists yet — only the 4 oracle tests below drive the launcher); `seed` will be derived deterministically per-call from a trainer RNG state. The launcher itself contains no host branches inside the launch builder (`pearl_no_host_branches_in_captured_graph` clean — the `seed` u64 is a kernel argument, not a host conditional, and the grid-dim computation runs once before any `launch_builder` call). (4) **3 existing oracle tests migrated** to the new launcher signature in `crates/ml/tests/sp15_phase1_oracle_tests.rs`: `plasticity_fires_when_persistence_exceeds_threshold` (passes `fan_in=256` + `seed=0xCAFE_BABE_DEAD_BEEF`), `plasticity_debounced_within_fold` (passes `fan_in=256` + `seed=0xDEAD_BEEF_F00D_CAFE` and now also asserts all 64 weights remain exactly 1.0 since the debounce path early-outs before the reset region), `plasticity_no_fire_below_threshold` (passes `fan_in=256` + `seed=0xFEED_FACE_BAAD_F00D` and asserts the same weight-stability invariant since no-fire also early-outs). The trigger-behavior assertions stay byte-identical. (5) **NEW oracle test** `plasticity_injection_kernel_resets_last_10pct_kaiming_he` (`crates/ml/tests/sp15_phase1_oracle_tests.rs` after `plasticity_no_fire_below_threshold`): allocates a 1280-element advantage-head weight buffer pre-filled with 1.0, sets ISV[DD_PERSISTENCE]=1000, ISV[PLASTICITY_PERSISTENCE_THRESHOLD]=1, ISV[PLASTICITY_FIRED_THIS_FOLD]=0; passes `fan_in=256`, `seed=42`; launches; then asserts: (a) ISV[PLASTICITY_FIRED_THIS_FOLD] flipped 0→1, (b) ISV[PLASTICITY_WARM_BARS_REMAINING] = 199 (`m_warm - 1`, per-bar decrement runs in same call), (c) first 90% (1152 elements) bit-identical to 1.0, (d) last 10% (128 elements) every slot has moved off 1.0 by > 1e-6 (Kaiming stddev≈0.088 makes 1.0 ~11σ out — exact-match impossible), (e) sample mean of the reset region |mean| < 0.05 (true SE ~0.0078 for n=128 from Normal(0, 0.088)), (f) sample std within ±20% of `sqrt(2/256) ≈ 0.08839`, (g) determinism re-check — re-launch with the same seed against a fresh 1.0-filled buffer + reset ISV produces bit-identical Kaiming samples per thread. All four plasticity oracle tests pass on RTX 3050 Ti (`SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` → 4 passed; 0 failed; 33 filtered out). (6) **Audit doc + build.rs comment + cubin-slot docstring + launcher docstring all updated** to reflect the now-landed weight-reset (Phase 3.5.4 deferred-NO-OP narrative replaced with Wave 4.2 active-reset narrative); the kernel's own header comment is rewritten to document the Wave 4.2 contract (cuRAND init, reset region addressing, Kaiming-He gain rationale, per-thread vs per-block responsibilities). (7) **Phase 3.5.4 deferred consumer eliminated** per `feedback_wire_everything_up`: the kernel signature's `(void)advantage_head_weights; (void)n_weights;` no-op cast is gone; the `(void)` pattern that flagged the deferred work is replaced by real per-element writes; the launcher no longer documents the plumbed-but-unused buffer. (8) **Action-selection consumer wiring (Phase 3.5.4.c) remains a separate follow-up** per `feedback_no_partial_refactor` — the OR-gate `effective_cooldown = max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)` plus the same-bar Flat emission on the fired-flag transition are independent of weight-reset and are tracked under their own atomic commit (matching Phase 1.1-1.5 + 3.1-3.5.3 precedent: kernel + launcher land first verified by oracle tests, then consumer migration is the next atomic). (9) **fan_in source for production** [INLINE-CORRECTED 2026-05-07 by Phase 3.5.4.c per `feedback_trust_code_not_docs`]: the advantage head's last Linear layer input dimension. Per `compute_param_sizes` line 4074 the size is `branch_0_size × num_atoms × adv_h` (the row-major weight tensor's flat element count); per the matching Kaiming-He fan-dim layout at `gpu_dqn_trainer.rs:29946` the (rows, cols) shape is `[branch_0_size × num_atoms, adv_h]` (output_dim × input_dim). Kaiming-He variance is `2 / fan_in` where `fan_in` is the INPUT dim per output unit — i.e. `adv_h`, NOT `adv_h × num_atoms`. Default config has `adv_h = 128`, `num_atoms = 51` → **fan_in = 128**, stddev = **`sqrt(2/128) ≈ 0.125`**. The pre-Phase-3.5.4.c version of this entry erroneously claimed `fan_in = adv_h × num_atoms = 6528` and `stddev ≈ 0.0175`; that conflated the total weight count of one branch's output projection (`num_atoms` outputs × `adv_h` inputs per output) with the per-unit input dim. The Wave 4.2 oracle tests still pass with `fan_in = 256` (synthetic representative value) because the assertions are scale-relative (`std within ±20% of sqrt(2/fan_in)`), not absolute; only the production-`fan_in` documentation was wrong. Phase 3.5.4.c's `set_sp15_plasticity_target()` derives the correct `fan_in = self.config.adv_h` directly from the trainer config. (10) **Determinism vs entropy choice** (Option A — per-thread `curand_init` per-call): trades a small amount of cuRAND-init compute (~10 cycles per thread to seed the Philox4x32 state) for full determinism in the oracle test plus reproducibility across trainer-restart resumed sessions, which is the established codebase posture (`feedback_h100_gpu` + the trainer's resumed-from-checkpoint contract). The alternative (Option B — persistent per-thread state across calls) would require a `curandState[reset_count]` mapped-pinned scratch buffer on the trainer struct (mirroring `sp15_cooldown_consecutive_losses`), but Wave 4.2 only fires at most once per fold and reset_count is small (≤ ~650 for default config), so the per-call init cost is negligible (< 1µs total) and Option A is the right scope. Touched: `crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu` (kernel body — `` include, multi-block grid contract, per-thread `curand_init` + `curand_normal` writes, header comment rewrite for Wave 4.2 contract), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher signature `+fan_in: i32, +seed: u64`, grid_dim computation, launch_builder arg list extended; cubin-slot docstring + launcher docstring rewritten for Wave 4.2 contract), `crates/ml/build.rs` (kernel manifest comment block updated — deferred-NO-OP narrative replaced with Wave 4.2 active-reset narrative + cuRAND header inlining note), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (3 existing tests migrated to new launcher signature + new test `plasticity_injection_kernel_resets_last_10pct_kaiming_he`). Hard rules: `feedback_no_atomicadd` (per-thread independent writes to unique `advantage_head_weights[reset_start + tid]`; no reductions; trigger logic remains single-thread for race-free ISV writes), `feedback_no_partial_refactor` (kernel + launcher signature change + 3 test migrations + 1 new test + audit doc + build.rs comment + 2 launcher/cubin docstrings all atomic; the action-selection consumer wiring remains the separate Phase 3.5.4.c follow-up by design — same pattern as Phase 3.5.3), `feedback_no_stubs` (the `(void)` no-op casts are gone; kernel performs the actual Kaiming-He reset), `feedback_no_cpu_compute_strict` (cuRAND init + sampling are GPU-resident — `curand_init` and `curand_normal` are device-only functions inlined into the cubin), `feedback_no_quickfixes` (the determinism re-check inside the oracle test catches accidental clock-derived-seed regressions; the mean/std assertions catch non-Gaussian or wrong-stddev regressions; the first-90%-unchanged assertion catches off-by-one regressions in the addressing arithmetic), `feedback_no_legacy_aliases` (no compatibility shim — the launcher signature changes in-place; existing callers MUST update to pass fan_in + seed or fail to compile), `feedback_wire_everything_up` (Phase 3.5.4's deferred-weight-reset NO-OP is closed in this commit; the `(void)` casts are removed). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings); `cargo check -p ml --features cuda --tests` clean; `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored plasticity --nocapture` 4 of 4 plasticity oracle tests green (`plasticity_fires_when_persistence_exceeds_threshold` + `plasticity_debounced_within_fold` + `plasticity_no_fire_below_threshold` + `plasticity_injection_kernel_resets_last_10pct_kaiming_he`) on RTX 3050 Ti; `cargo test -p ml --features cuda --lib` HOLDS the 947 pass / 12 fail Wave 4.1c baseline. SP15 Wave 4.1c — behavioral KL test: dd_pct trunk integration shifts synthetic policy distribution (2026-05-07): closes out Wave 4.1 (Phase 1.5.b consumer migration). Wave 4.1a (`a8da1cb9c`) landed `bn_tanh_concat_dd_kernel`; Wave 4.1b (`eb9515e41`) wired `s1_input_dim` 102→103 through GRN reshape + 4 forward + 3 backward call sites. Wave 4.1c proves the wiring actually changes the policy: a synthetic GPU forward composes `launch_sp15_bn_concat_dd` (the production kernel that injects ISV[DD_PCT_INDEX=406] into the trunk input) with a single cuBLAS `cublasSgemm_v2` against random Xavier-init weights `W[proj_h=4, concat_dd_dim=103]`, then computes softmax + KL on the host post-readback. Two passes differing ONLY in `ISV[DD_PCT]` (0.0 at-ATH vs 0.10 in-DD) yield mean KL = 1.158e-4 (max KL = 3.005e-4) — well above the `1e-6` threshold (~100× headroom), confirming the new column-102 weights are non-zero and the dd_pct column flows from ISV → bn_concat[:,102] → SGEMM K-dim 103. (1) **New test `dd_pct_trunk_input_shifts_policy_distribution`** in module `sp15_wave_4_1c_behavioral` at `crates/ml/tests/sp15_phase1_oracle_tests.rs` — `#[ignore = "requires GPU"]` (matches the existing oracle-test gating). Imports `launch_sp15_bn_concat_dd` (Wave 4.1a launcher), `MappedF32Buffer` (mapped-pinned alloc per `feedback_no_htod_htoh_only_mapped_pinned`), `PerStreamCublasHandles::new` (the public cuBLAS handle factory at `shared_cublas_handle.rs:174`), `cublas_sys::cublasSgemm_v2` (raw FFI matching `gpu_tlob.rs:954`'s pattern), and the helper `kl_divergence` from Wave 4.1a's seeded `sp15_wave_4_1a_test_helpers` module via `super::sp15_wave_4_1a_test_helpers::kl_divergence`. (2) **Synthetic forward layout** — `synthetic_forward` helper composes the production bn_concat_dd kernel + a synthetic linear projection: cuBLAS column-major sgemm `op(A)=W^T [concat_dd_dim, proj_h]` × `op(B)=bn_concat [concat_dd_dim, B]` → `C = [proj_h, B]` column-major, m=proj_h, n=batch, k=concat_dd_dim. Same convention `gpu_tlob.rs:472-490` uses for the QKV projection. The post-output transpose to row-major `[B, proj_h]` happens on host, post-readback. (3) **NOT a `forward_trunk_for_test` per the seeded plan** — the seeded comment in Wave 4.1a (`sp15_phase1_oracle_tests.rs:2304-2319`) noted `forward_trunk_for_test` would require either (a) a public surface change on `DQNTrainer` exposing internal cuBLAS handles + GRN scratch buffers + weight pointers (the trainer's `fused_ctx` is `pub(crate)` and only initialised inside the training loop at `training_loop.rs:547` — `DQNTrainer::new` returns with `fused_ctx: None`), or (b) duplicating the trunk's cuBLAS setup in a test (≥200 lines of buffer + handle plumbing). Both options are architecturally heavier than the test's purpose justifies. Per the dispatch's "the test's purpose is 'non-zero KL proves the wire is connected' not 'verifies trained behavior'", the synthetic single-layer projection is the right scope: it exercises the new column-102 weights on the dd_pct value — exactly the path the real GRN's `Linear_a` first GEMM takes for `w_a_h_s1[:, 102]`. (4) **Random Xavier-uniform weight init** — bound = sqrt(6 / (fan_in + fan_out)) = sqrt(6 / 107) ≈ 0.237 for (concat_dd_dim=103, proj_h=4). Same formula `gpu_dqn_trainer.rs::xavier_init_params_buf` applies for the production GRN `w_a_h_s1` weights. Deterministic LCG seed=42 (Numerical Recipes constants) so the test is byte-reproducible across runs. (5) **What this test buys** — layout fingerprint + behavioral coupling. If a future migration breaks the wiring (e.g. column-102 Xavier init silently zeros, the post-concat buffer truncates back to 102 columns, the SGEMM K-dim falls back to 102, or `bn_tanh_concat_dd` stops writing the dd_pct column), `mean_kl` collapses to ~0 and the test fires with a diagnostic message that pinpoints the regression to the 4.1a/4.1b/4.1c chain. (6) **What this test does NOT verify** — the full GRN composition (ELU / GLU / LayerNorm / residual) propagating dd_pct through h_s2 + the branch advantage heads. That end-to-end behavior is exercised by the L40S smoke + production training runs. The test is a **pinpoint sanity check** for the trunk-input wiring, not an integration test of the full network. (7) **Numerical reasonableness** — the dd_pct contribution to the synthetic logits is `0.10 × W[h, 102]` for each output unit h. With Xavier W ~ U[-0.237, +0.237], the pre-softmax logit shift is O(0.10 × 0.237) ≈ 0.024 per unit. Through a 4-class softmax on small base logits this produces row-distribution shifts with KL ~ 1e-4 (observed: mean=1.158e-4, max=3.005e-4 across the 4-row batch). The threshold of 1e-6 sits 100× below the observed magnitude — large enough to reject pure numerical noise, small enough to pass on random-init weights without ramping up to the trained-network signal magnitude. (8) **Pearl applicability** — `pearl_canary_input_freshness_launch_order`: producers feeding a canary MUST launch before it. In the synthetic forward, `bn_tanh_concat_dd` (producer) writes the post-concat buffer; `cublasSgemm_v2` (consumer) reads it. Both submitted to the same stream → executes in order; explicit `stream.synchronize()` between launch and host readback per `feedback_no_htod_htoh_only_mapped_pinned`. Per `feedback_no_cpu_compute_strict`: the SGEMM is GPU-only (cuBLAS); the only host computation is the post-readback softmax + KL, which is post-output policy extraction (NOT trunk replication — the GRN composition stays on GPU in production code, untouched by this test). Per `feedback_isv_for_adaptive_bounds`: dd_pct is the adaptive observable being tested; this is a downstream behavioral check on the existing ISV producer chain (Wave 1.3.b's `dd_state_kernel`). (9) **Atomicity** — purely additive: no kernel changes, no production code changes, no public-API surface changes, no audit-doc changes other than this entry. Touched: `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 new module `sp15_wave_4_1c_behavioral` with 1 ignored test + 2 helper fns + 1 assertion block). (10) **Phase 1.5.b orphan launcher chain fully eliminated** per `feedback_wire_everything_up`: kernel landed (4.1a) → consumer migration (4.1b) → behavioral verification (4.1c) — three atomic commits, the 3a/3b/3c split-pattern matching Wave 3's a/b decomposition. The Wave 4.1a transient orphan window opened in `a8da1cb9c` (kernel) → closed in `eb9515e41` (consumer migration) → behavioral coverage added in this commit. Hard rules: `feedback_no_partial_refactor` (test addition is atomic with the audit-doc entry — no parallel half-test path), `feedback_wire_everything_up` (the seeded helpers `kl_divergence` + `minimal_trainer_for_tests` (Wave 4.1a) get a real consumer in this commit per the documented plan; `minimal_trainer_for_tests` was not used here because the seeded comment correctly identified that exposing the trunk forward via the trainer's surface is non-trivial — the helper remains for future cargo-cult tests that don't need GPU forward, e.g. weight-shape introspection), `feedback_no_cpu_compute_strict` (cuBLAS SGEMM is GPU; row_softmax + kl_divergence run on host post-readback as policy-extraction analysis, not as trunk-forward replacement — the actual trunk forward remains GPU-only in production), `feedback_no_stubs` (all helpers do real work; no NULL-skip; no return-zero shortcut), `feedback_no_quickfixes` (the test threshold is honest: 1e-6 is set 100× below the observed magnitude so a real wiring break makes the test fire — NOT silently passing), `feedback_isv_for_adaptive_bounds` (dd_pct is ISV-driven via slot 406 by Wave 1.3.b's existing contract; no new ISV slot added). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda --tests` clean (18 pre-existing unrelated warnings, no new warnings); `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored bn_concat dd_pct --nocapture` 2 of 2 oracle tests green (Wave 4.1a `bn_tanh_concat_dd_kernel_writes_dd_pct_column` + Wave 4.1c `dd_pct_trunk_input_shifts_policy_distribution`); `cargo test -p ml --features cuda --lib` HOLDS the 947 pass / 12 fail Wave 4.1b baseline (the new test is in the `--test` integration-test target, not the lib target — lib count unchanged as expected).