diff --git a/crates/ml/build.rs b/crates/ml/build.rs index b696846e2..51a832c26 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1036,20 +1036,24 @@ fn main() { // (3.5.4) post-amendment-2 fix: fire when DD_PERSISTENCE // exceeds threshold AND PLASTICITY_FIRED_THIS_FOLD == 0 → // set fired flag, set warm-bars counter to M_warm (default - // 200), [DEFERRED: reset last 10% of advantage-head weights]; - // per-bar warm-bars decrement; consumer-wiring follow-up reads - // max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING) - // and forces Hold while > 0. Phase 3.5.4 lands kernel + - // launcher + 3 ISV anchor seeds + 3 registry entries + 3 - // dispatch arms only; the action-selection two-step Flat + - // cooldown wiring AND the actual Kaiming-He weight-reset are - // both deferred to a Phase 3.5.4.b follow-up atomic commit per - // `feedback_no_partial_refactor.md` (matching Phase 1.1-1.5 + - // 3.1-3.5.3 precedent — kernel + launcher verify in isolation - // first via the GPU oracle tests below). The kernel signature - // accepts the advantage_head_weights pointer + n_weights param - // so the follow-up commit is purely additive on the kernel - // body, not on the launcher contract. + // 200), AND reset last 10% of advantage-head weights to + // Kaiming-He init (`Normal(0, sqrt(2/fan_in))` sampled via + // cuRAND `curand_normal()`); per-bar warm-bars decrement; + // consumer-wiring follow-up reads max(COOLDOWN_BARS_REMAINING, + // PLASTICITY_WARM_BARS_REMAINING) and forces Hold while > 0. + // Phase 3.5.4 (2026-05-06) landed kernel + launcher + 3 ISV + // anchor seeds + 3 registry entries + 3 dispatch arms with + // the weight-reset deferred (kernel accepted + // advantage_head_weights + n_weights but no-op'd via (void) + // cast); Wave 4.2 / Phase 3.5.4.b (2026-05-07) lands the + // actual cuRAND Kaiming-He reset, extending the launcher + // signature with `fan_in: i32` and `seed: u64`. The kernel + // includes `` for device-side + // `curand_init`/`curand_normal` (inlined into the cubin by + // nvcc — no host-side cuRAND linker dependency required). + // The action-selection two-step Flat + cooldown consumer + // wiring remains a separate follow-up (Phase 3.5.4.c) per + // `feedback_no_partial_refactor.md`. "plasticity_injection_kernel.cu", // SP15 Phase 3.5.5 (2026-05-06): recovery curriculum in PER — // per-step DD_TRAJECTORY_DECREASING proxy. Replaces non-existent diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 2649ad4e1..7122ad002 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1705,13 +1705,16 @@ pub fn launch_sp15_cooldown( Ok(()) } -/// SP15 Phase 3.5.4 (2026-05-06): cubin slot for the plasticity -/// injection trigger + warm-up tracker kernel. TWO-STEP recovery per -/// spec §9.2 (3.5.4) post-amendment-2 fix: +/// SP15 Phase 3.5.4 (2026-05-06) + Wave 4.2 / Phase 3.5.4.b +/// (2026-05-07): cubin slot for the plasticity injection trigger + +/// warm-up tracker + cuRAND Kaiming-He weight-reset kernel. TWO-STEP +/// recovery per spec §9.2 (3.5.4) post-amendment-2 fix: /// 1. Fire when `DD_PERSISTENCE > PLASTICITY_PERSISTENCE_THRESHOLD` /// AND `PLASTICITY_FIRED_THIS_FOLD == 0` → set fired flag, set -/// warm-bars counter to M_warm (default 200), [DEFERRED: reset -/// last 10% of advantage-head weights to Kaiming-He init]. +/// warm-bars counter to M_warm (default 200), AND reset the +/// trailing 10% of advantage-head weights to Kaiming-He init +/// (`Normal(0, sqrt(2/fan_in))` sampled via cuRAND +/// `curand_normal()`; gain=√2 for ReLU/GLU activations). /// 2. Per-bar warm-bars decrement; consumer-wiring follow-up reads /// `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)` /// and forces Hold while > 0 (the OR-gate combining cooldown + @@ -1722,47 +1725,65 @@ pub fn launch_sp15_cooldown( /// `dd_asymmetric_reward_kernel.cu` / 3.5.3 `cooldown_kernel.cu` /// single-kernel-per-cubin pattern. /// -/// Phase 3.5.4 lands kernel + launcher + 3 ISV constructor seeds + 3 -/// state-reset registry entries + 3 dispatch arms only; the action- -/// selection two-step Flat + cooldown wiring AND the actual Kaiming- -/// He weight-reset are both deferred to a Phase 3.5.4.b follow-up -/// atomic commit per `feedback_no_partial_refactor.md` (matching -/// Phase 1.1-1.5 + 3.1-3.5.3 precedent — kernel + launcher verify in -/// isolation first via the GPU oracle tests below). The kernel -/// signature accepts the advantage_head_weights pointer + n_weights -/// param so the follow-up commit is purely additive on the kernel -/// body, not on the launcher contract. +/// Phase 3.5.4 landed kernel + launcher + 3 ISV constructor seeds + +/// 3 state-reset registry entries + 3 dispatch arms with the weight- +/// reset deferred (kernel accepted `advantage_head_weights + +/// n_weights` but no-op'd via `(void)` cast). Wave 4.2 / Phase +/// 3.5.4.b lands the actual cuRAND Kaiming-He reset, extending the +/// launcher signature with `fan_in: i32` and `seed: u64` and +/// promoting the kernel grid from single-thread/single-block to a +/// 256-thread per-element grid for the weight write. The action- +/// selection two-step Flat + cooldown consumer wiring remains a +/// separate follow-up (Phase 3.5.4.c) per `feedback_no_partial_refactor.md`. pub static SP15_PLASTICITY_INJECTION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/plasticity_injection_kernel.cubin")); -/// SP15 Phase 3.5.4 (2026-05-06): launcher for the plasticity -/// injection trigger + warm-up tracker kernel. Free function (matches -/// `launch_sp15_dd_penalty` / `launch_sp15_regret_signal` / -/// `launch_sp15_hold_floor` / `launch_sp15_dd_asymmetric_reward` / -/// `launch_sp15_cooldown` precedent — single kernel, single cubin, -/// single launcher) so unit/oracle tests can drive the kernel -/// directly without the trainer struct; production callers invoke -/// the same launcher. +/// SP15 Phase 3.5.4 (2026-05-06) + Wave 4.2 / Phase 3.5.4.b +/// (2026-05-07): launcher for the plasticity injection trigger + +/// warm-up tracker + cuRAND Kaiming-He weight-reset kernel. Free +/// function (matches `launch_sp15_dd_penalty` / +/// `launch_sp15_regret_signal` / `launch_sp15_hold_floor` / +/// `launch_sp15_dd_asymmetric_reward` / `launch_sp15_cooldown` +/// precedent — single kernel, single cubin, single launcher) so +/// unit/oracle tests can drive the kernel directly without the +/// trainer struct; production callers invoke the same launcher. /// /// `isv` MUST be the ISV bus (≥ `SP15_SLOT_END=443` f32 slots — slot /// 404 is read for DD_PERSISTENCE; slot 437 is read for /// PLASTICITY_PERSISTENCE_THRESHOLD; slots 436 + 438 are read-modify- /// written for PLASTICITY_FIRED_THIS_FOLD + PLASTICITY_WARM_BARS_ /// REMAINING). `advantage_head_weights` MUST point to at least -/// `n_weights` writable f32 slots — but the weight-reset is DEFERRED -/// to Phase 3.5.4.b so the kernel body currently does NOT touch the -/// buffer; the parameter is wired through the launcher contract so -/// the follow-up commit is purely additive on the kernel body. +/// `n_weights` writable f32 slots — Wave 4.2 lands the actual +/// Kaiming-He reset of the trailing `n_weights / 10` slots when the +/// kernel fires. +/// /// `m_warm` is the warm-up bar count to set on fire (default 200 per -/// spec §9.2 (3.5.4)). Single thread, single block (mirrors -/// `cooldown_kernel` / `dd_asymmetric_reward_kernel` / `hold_floor_kernel` -/// — no reduction). +/// spec §9.2 (3.5.4)). `fan_in` is the input dimension of the +/// advantage head's last Linear layer (the Kaiming-He gain +/// denominator — `stddev = sqrt(2 / fan_in)`); production callers +/// derive it from the trainer's network config. `seed` is a +/// per-call deterministic value (e.g. derived from a trainer RNG +/// state or `SystemTime::now()` nanos) — the kernel re-initialises +/// per-thread `curandState`s from `(seed, sequence=tid, offset=0)` +/// every launch; the same `(seed, tid)` pair always yields the same +/// sample. +/// +/// Grid: `[((n_weights / 10 + 255) / 256), 1, 1]` × `[256, 1, 1]`. +/// Block 0 / thread 0 owns the trigger + warm-bars decrement (single- +/// thread, mirrors `cooldown_kernel` / `dd_asymmetric_reward_kernel` +/// / `hold_floor_kernel`); the remaining threads race-free write the +/// per-element Kaiming-He samples (each thread writes a unique +/// `advantage_head_weights[reset_start + tid]`, no atomic, no +/// reduction). When `n_weights / 10 == 0` (oracle edge case) the +/// grid degenerates to `[1, 1, 1]` so the trigger still runs. pub fn launch_sp15_plasticity_injection( stream: &Arc, isv: cudarc::driver::sys::CUdeviceptr, advantage_head_weights: cudarc::driver::sys::CUdeviceptr, n_weights: i32, m_warm: f32, + fan_in: i32, + seed: u64, ) -> Result<(), MLError> { let module = stream .context() @@ -1775,6 +1796,13 @@ pub fn launch_sp15_plasticity_injection( .map_err(|e| MLError::ModelError(format!( "load plasticity_injection_kernel function: {e}" )))?; + // Grid sized so every weight in the trailing 10% gets a thread. + // `block_dim = 256` matches the codebase's prevailing 1-D + // element-wise launch shape. `grid_dim` always ≥ 1 so block 0 + // thread 0 always runs the trigger, even when `n_weights / 10 == 0` + // (oracle edge case where the test passes a tiny weight buffer). + let reset_count = (n_weights / 10).max(0); + let grid_x = ((reset_count + 255) / 256).max(1) as u32; unsafe { stream .launch_builder(&kernel) @@ -1782,9 +1810,11 @@ pub fn launch_sp15_plasticity_injection( .arg(&advantage_head_weights) .arg(&n_weights) .arg(&m_warm) + .arg(&fan_in) + .arg(&seed) .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (1, 1, 1), + grid_dim: (grid_x, 1, 1), + block_dim: (256, 1, 1), shared_mem_bytes: 0, }) .map_err(|e| MLError::ModelError(format!( diff --git a/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu index 04549a804..ff96c93bc 100644 --- a/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu +++ b/crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu @@ -1,13 +1,14 @@ // crates/ml/src/cuda_pipeline/plasticity_injection_kernel.cu // // SP15 Phase 3.5.4 — plasticity injection trigger + warm-up tracker. +// SP15 Wave 4.2 / Phase 3.5.4.b — cuRAND Kaiming-He weight reset (now wired). // // TWO-STEP recovery per spec §9.2 (3.5.4) post-amendment-2 fix: // 1. On fire: set PLASTICITY_FIRED_THIS_FOLD = 1.0, set -// PLASTICITY_WARM_BARS_REMAINING = M_warm, [DEFERRED: reset last -// 10% of advantage-head weights to Kaiming-He init (gain=√2 for -// ReLU/GLU activations) via cuRAND]. The action-selection layer -// (consumer-wiring follow-up) reads the +// PLASTICITY_WARM_BARS_REMAINING = M_warm, AND reset the last 10% +// of advantage-head weights to Kaiming-He init (gain=√2 for +// ReLU/GLU activations) via cuRAND `curand_normal()`. The +// action-selection layer (consumer-wiring follow-up) reads the // PLASTICITY_FIRED_THIS_FOLD == 1 flag transition AND emits Flat // that bar. // 2. Per-bar decrement of PLASTICITY_WARM_BARS_REMAINING; the @@ -16,27 +17,35 @@ // PLASTICITY_WARM_BARS_REMAINING)` to force Hold while either // counter > 0. // -// **Weight-reset is DEFERRED** to a Phase 3.5.4.b follow-up. This -// commit lands: -// - Fire/debounce trigger logic (PLASTICITY_FIRED_THIS_FOLD gates -// re-fire within the same fold; reset to 0.0 at fold boundary -// re-arms the next fold). -// - Per-bar warm-bars decrement (every kernel call, regardless of -// fire). -// - The advantage_head_weights pointer + n_weights param wired -// through the launcher signature (kernel accepts but does not -// modify them yet — same `(void)foo` no-op pattern used elsewhere -// for deferred parameter plumbing). +// Wave 4.2 lands the previously-deferred Kaiming-He weight reset: // -// Why deferred: weight-reset requires Kaiming-He init via cuRAND, and -// accurate targeting of "last 10% of advantage-head weights" requires -// plumbing specific param-buffer pointers (mag/ord/urg head Linear -// layer slices) that aren't available at this kernel's call site -// without trainer-internal access. Establishing the kernel signature + -// flag-tracking now keeps the follow-up commit purely additive on the -// weight-reset side, mirroring the Phase 1.1-1.5 + 3.1-3.5.3 precedent -// of kernel + launcher landing first and consumer wiring following -// atomically. +// - Reset region: the last 10% of `advantage_head_weights` +// (`reset_count = n_weights / 10`, addressed via offset +// `advantage_head_weights + n_weights - reset_count`). Resetting +// the trailing slice keeps the upstream feature-extraction trunk +// intact while re-initialising the head's last layer where the +// drawdown-correlated gradients have presumably saturated. This +// matches the spec §9.2 (3.5.4) "last 10%" prescription. +// - Init formula: Kaiming-He for ReLU/GLU activations, +// `weights[i] ~ Normal(0, sqrt(2 / fan_in))` (gain = sqrt(2)). +// `fan_in` is supplied by the launcher (host computes it from the +// trainer's network config; for oracle tests the test sets it to +// a representative value). +// - Per-thread cuRAND state via `curand_init(seed, tid, 0, &state)` +// + `curand_normal(&state)` × stddev. `seed` is a host-passed +// `unsigned long long`; the launcher derives it deterministically +// from a trainer RNG state (or `SystemTime` nanos in the oracle +// tests). The same (seed, tid) pair always produces the same +// sample, satisfying determinism for the oracle test's mean/std +// assertions. +// - Grid: multi-block, `[((reset_count + 255) / 256), 1, 1]` × +// `[256, 1, 1]`. Block 0 / thread 0 still runs the trigger + +// warm-bars decrement (the only ISV-write code path); all other +// threads independently re-evaluate the same fire condition from +// the same ISV source-of-truth (ISV state is read at kernel-launch +// time, the trigger condition is a pure function of those reads, +// so every thread computes the same `fire_now` without any +// cross-block synchronisation). // // Slot reads/writes (literal indices to keep this kernel self-contained; // the canonical names live in `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` @@ -44,50 +53,60 @@ // indices via compile-time `assert_eq!` regression checks): // ISV[404] DD_PERSISTENCE (read) // ISV[437] PLASTICITY_PERSISTENCE_THRESHOLD (read) -// ISV[436] PLASTICITY_FIRED_THIS_FOLD (read-modify-write) -// ISV[438] PLASTICITY_WARM_BARS_REMAINING (read-modify-write) +// ISV[436] PLASTICITY_FIRED_THIS_FOLD (read-modify-write — block 0 thread 0 only) +// ISV[438] PLASTICITY_WARM_BARS_REMAINING (read-modify-write — block 0 thread 0 only) // // Hard rules: // -// * `feedback_no_atomicadd` — single-thread, single-block, pure -// scalar arithmetic; no reductions, no `atomicAdd`. -// * `feedback_no_partial_refactor` — kernel + launcher land first; -// consumer wiring (action-selection two-step Flat+cooldown -// integration via `effective_cooldown = -// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)`) -// deferred to a follow-up atomic commit per the established Phase -// 1.1-1.5 / 3.1-3.5.3 precedent. Weight-reset is also deferred -// explicitly — the launcher signature accepts the pointer + count -// so the follow-up is purely additive on the kernel body, not on -// the launcher contract. -// * `feedback_no_stubs` — kernel returns the actual updated state -// (flag transition + warm-bars decrement); the deferred weight- -// reset is a no-op on the same call rather than a stubbed -// placeholder return. The three oracle tests verify fire-on- -// threshold-crossing, debounce-within-fold, and below-threshold- -// no-fire behaviors. +// * `feedback_no_atomicadd` — weight-reset is per-thread independent +// (each thread writes a unique element of `advantage_head_weights`); +// no reductions, no `atomicAdd`. The trigger logic remains +// single-thread (block 0 / thread 0) so its ISV writes are +// race-free without atomics. +// * `feedback_no_partial_refactor` — kernel + launcher signature +// change + 4 oracle tests (3 existing migrated + 1 new) land +// atomically in this commit (Wave 4.2); consumer wiring (action +// selection forces Hold via `effective_cooldown = +// max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING)` +// remains the separate follow-up tracked under Phase 3.5.4.c. +// * `feedback_no_stubs` — kernel performs the actual Kaiming-He +// reset; no `(void)` casts remain. Wave 4.2 closes the Phase +// 3.5.4-deferred weight-reset NO-OP per +// `feedback_wire_everything_up.md`. // * `feedback_isv_for_adaptive_bounds` — PLASTICITY_PERSISTENCE_ // THRESHOLD is ISV-tracked (slot 437); the producer driving it // from the running mean of dd_persistence_bars is deferred but // the kernel reads it rather than a hardcoded literal so the // follow-up swap is a no-op at the consumer. -// -// Single-thread, single-block (mirrors `cooldown_kernel` / -// `dd_asymmetric_reward_kernel` / `hold_floor_kernel` / -// `regret_signal_kernel` / `dd_penalty_kernel` — no reduction). +// * `feedback_no_cpu_compute_strict` — cuRAND init + sampling are +// GPU-resident (`curand_init` / `curand_normal` are device-only +// functions inlined by nvcc into the cubin from +// ``). +// * `pearl_no_host_branches_in_captured_graph` — `seed` is a kernel +// argument (a `u64` value, not a host-side branch); the launcher +// itself can be invoked from inside a graph capture region +// because the kernel-launch shape (grid/block) is fixed at the +// call (a function of `n_weights / 10` which is also a launch +// argument, not a host-conditional). + +#include extern "C" __global__ void plasticity_injection_kernel( float* __restrict__ isv, - float* __restrict__ advantage_head_weights, /* deferred: weight-reset NO-OP for now */ + float* __restrict__ advantage_head_weights, int n_weights, - float m_warm /* default 200 bars */ + float m_warm, /* default 200 bars */ + int fan_in, /* Kaiming gain denominator */ + unsigned long long seed /* cuRAND seed */ ) { - if (threadIdx.x != 0 || blockIdx.x != 0) return; - + // Every thread independently reads the ISV state to determine + // `fire_now`. The trigger condition is a pure function of those + // reads, so all threads converge on the same answer without any + // cross-block synchronisation. The same reads also drive block 0 + // thread 0's ISV-write path below. const float persistence = isv[/*DD_PERSISTENCE_INDEX*/ 404]; const float threshold = isv[/*PLASTICITY_PERSISTENCE_THRESHOLD_INDEX*/ 437]; const float fired = isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436]; - float warm = isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438]; // Trigger: persistence exceeds threshold AND not yet fired this // fold. PLASTICITY_FIRED_THIS_FOLD is reset to 0.0 at fold boundary @@ -95,31 +114,58 @@ extern "C" __global__ void plasticity_injection_kernel( // re-arming the gate for the next fold. The `< 0.5f` half-bit // tolerance check matches the existing `cooldown_remaining <= 0.0f` // float-equality-with-tolerance pattern in `cooldown_kernel`. - if (persistence > threshold && fired < 0.5f) { - isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436] = 1.0f; - warm = m_warm; + const bool fire_now = (persistence > threshold) && (fired < 0.5f); - // Weight-reset DEFERRED to Phase 3.5.4.b — kernel param plumbed - // through but no-op. Follow-up: Kaiming-He reset of last 10% of - // advantage-head weights via cuRAND with gain=√2 for ReLU/GLU - // activations. The `(void)foo` cast pattern matches the Rust - // `_ = foo` convention used elsewhere in this codebase for - // deferred parameter plumbing. - (void)advantage_head_weights; - (void)n_weights; + // Block 0 / thread 0 owns the ISV writes (fired flag + warm-bars + // counter). Single-thread to mirror the cooldown_kernel / + // dd_asymmetric_reward_kernel / hold_floor_kernel / regret_signal + // / dd_penalty_kernel pattern — no reduction, no atomic. + if (threadIdx.x == 0 && blockIdx.x == 0) { + float warm = isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438]; + + if (fire_now) { + isv[/*PLASTICITY_FIRED_THIS_FOLD_INDEX*/ 436] = 1.0f; + warm = m_warm; + } + + // Per-bar decrement of warm-bars (every call, regardless of + // fire). The trigger above runs first so the same call that + // fires the gate also immediately consumes one bar of warm-up + // — bias downward by 1 is acceptable per the test's [198, 200] + // tolerance and avoids the bookkeeping branch needed to skip + // the decrement on the trip step (mirrors `cooldown_kernel`'s + // identical trigger-then-decrement sequence). + if (warm > 0.0f) { + warm -= 1.0f; + } + isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438] = warm; + + __threadfence_system(); } - // Per-bar decrement of warm-bars (every call, regardless of fire). - // The trigger above runs first so the same call that fires the gate - // also immediately consumes one bar of warm-up — bias downward by 1 - // is acceptable per the test's [198, 200] tolerance and avoids the - // bookkeeping branch needed to skip the decrement on the trip step - // (mirrors `cooldown_kernel`'s identical trigger-then-decrement - // sequence). - if (warm > 0.0f) { - warm -= 1.0f; - } - isv[/*PLASTICITY_WARM_BARS_REMAINING_INDEX*/ 438] = warm; + // Weight reset: last 10% of advantage-head weights via Kaiming-He + // sampling. Each thread writes a unique element so no atomic / + // reduction is required. Threads beyond `reset_count` early-out; + // when `fire_now == false` every thread early-outs and the kernel + // is a pure trigger/decrement (the previous Phase 3.5.4 behaviour). + if (!fire_now) return; - __threadfence_system(); + const int reset_count = n_weights / 10; + if (reset_count <= 0) return; + + const int reset_start = n_weights - reset_count; + const int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid >= reset_count) return; + + // Per-thread cuRAND state: (seed, sequence=tid, offset=0). Same + // (seed, tid) pair always produces the same sample, satisfying + // determinism for oracle-test mean/std assertions. `fan_in` is the + // input dimension of the advantage head's last Linear layer; for + // ReLU/GLU activations the Kaiming-He gain is sqrt(2). + curandState state; + curand_init(seed, (unsigned long long)tid, 0ULL, &state); + + const float gain = sqrtf(2.0f); + const float stddev = gain / sqrtf((float)fan_in); + advantage_head_weights[reset_start + tid] = curand_normal(&state) * stddev; } diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 67c21bc0d..5da01669c 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -1141,10 +1141,12 @@ mod gpu { /// reads the fired flag transition AND emits Flat that bar, then /// forces Hold while /// `max(COOLDOWN_BARS_REMAINING, PLASTICITY_WARM_BARS_REMAINING) > 0`. - /// Weight-reset deferred to Phase 3.5.4.b — kernel param plumbed - /// through but no-op for now (the test passes a non-zero - /// advantage_head_weights buffer to verify the launcher accepts it - /// without modifying the contents). + /// Wave 4.2 / Phase 3.5.4.b: kernel now also performs the + /// Kaiming-He weight-reset; this trigger-only test passes the + /// extra `fan_in` + `seed` launcher params and asserts the + /// trigger behaviour (fired flag + warm-bars). The dedicated + /// `plasticity_injection_kernel_resets_last_10pct_kaiming_he` + /// test covers the weight-reset numerics. #[test] #[ignore = "requires GPU"] fn plasticity_fires_when_persistence_exceeds_threshold() { @@ -1164,9 +1166,10 @@ mod gpu { isv_init[PLASTICITY_WARM_BARS_REMAINING_INDEX] = 0.0; isv_buf.write_from_slice(&isv_init); - // Empty advantage-head weights buffer for now (weight-reset - // deferred to Phase 3.5.4.b — kernel takes the pointer but - // does not modify contents). + // Wave 4.2 weight-reset is exercised in detail by the + // dedicated Kaiming-He oracle below; here we just provide a + // 64-element buffer so the launcher can run (n_weights / 10 = + // 6 trailing slots get re-sampled). let weights_buf = unsafe { MappedF32Buffer::new(64) } .expect("alloc MappedF32Buffer for advantage_head_weights"); weights_buf.write_from_slice(&vec![1.0f32; 64]); @@ -1177,6 +1180,8 @@ mod gpu { weights_buf.dev_ptr, /*n_weights*/ 64, /*m_warm*/ 200.0, + /*fan_in*/ 256, + /*seed*/ 0xCAFE_BABE_DEAD_BEEF_u64, ) .expect("launch plasticity_injection_kernel (fires)"); stream @@ -1238,6 +1243,8 @@ mod gpu { weights_buf.dev_ptr, 64, 200.0, + /*fan_in*/ 256, + /*seed*/ 0xDEAD_BEEF_F00D_CAFE_u64, ) .expect("launch plasticity_injection_kernel (debounced)"); stream @@ -1256,6 +1263,17 @@ mod gpu { (warm - 49.0).abs() < 1e-9, "PLASTICITY_WARM_BARS_REMAINING = {warm}, expected 49 (decremented from 50, no re-fire)" ); + + // Wave 4.2: when debounced, the kernel returns early before + // the weight-reset region, so all 64 weights must remain + // exactly 1.0 (no Kaiming-He sample written). + let weights_host = weights_buf.read_all(); + for (i, w) in weights_host.iter().enumerate() { + assert!( + (w - 1.0).abs() < 1e-9, + "weights[{i}] = {w}, expected 1.0 (debounced — kernel must skip weight-reset)" + ); + } } /// SP15 Phase 3.5.4 — plasticity injection does NOT fire when @@ -1293,6 +1311,8 @@ mod gpu { weights_buf.dev_ptr, 64, 200.0, + /*fan_in*/ 256, + /*seed*/ 0xFEED_FACE_BAAD_F00D_u64, ) .expect("launch plasticity_injection_kernel (no-fire)"); stream @@ -1310,6 +1330,189 @@ mod gpu { "PLASTICITY_WARM_BARS_REMAINING = {}, expected 0.0 (no fire → no warm-up set; the `if (warm > 0)` guard prevents underflow when already 0)", isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX] ); + + // Wave 4.2: when no-fire, the kernel returns early before the + // weight-reset region, so all 64 weights must remain exactly + // 1.0 (no Kaiming-He sample written). + let weights_host = weights_buf.read_all(); + for (i, w) in weights_host.iter().enumerate() { + assert!( + (w - 1.0).abs() < 1e-9, + "weights[{i}] = {w}, expected 1.0 (no-fire — kernel must skip weight-reset)" + ); + } + } + + /// SP15 Wave 4.2 / Phase 3.5.4.b — when the plasticity injection + /// gate fires, the trailing 10% of `advantage_head_weights` must + /// be re-initialised to Kaiming-He samples + /// (`Normal(0, sqrt(2/fan_in))`, gain=√2 for ReLU/GLU + /// activations) drawn via cuRAND `curand_normal()`. The first 90% + /// of the buffer must remain bit-identical to its pre-launch + /// value (1.0) — the kernel writes element `[reset_start + tid]` + /// only. + /// + /// Distribution checks (with `n_weights = 1280` → `reset_count = + /// 128`, `fan_in = 256`, expected `stddev = sqrt(2/256) ≈ + /// 0.08839`): + /// - Per-element diff from 1.0: every reset slot must have + /// changed (probability of an exact-1.0 sample is essentially + /// zero — Kaiming stddev ≈ 0.088 means 1.0 is ~11.3 σ out, and + /// anyway the kernel writes the *sample*, not the original). + /// - Mean of reset region: |mean| < 0.05 (sampling-noise + /// tolerance for n=128 from a Normal with stddev≈0.088 — the + /// true SE of the mean is `0.088 / sqrt(128) ≈ 0.0078`, so + /// `0.05` is a comfortable ~6σ envelope around 0). + /// - Std of reset region: within ±20% of `sqrt(2/fan_in)` + /// (n=128 sample std has ~6% relative SE; ±20% is comfortable + /// headroom). + /// - ISV[PLASTICITY_FIRED_THIS_FOLD_INDEX] flipped 0 → 1. + /// - ISV[PLASTICITY_WARM_BARS_REMAINING_INDEX] = m_warm − 1 + /// (per-bar decrement runs in the same launch as the trigger + /// per the cooldown_kernel trigger-then-decrement convention). + /// + /// Determinism: a fixed seed (`42`) drives `curand_init(seed, + /// tid, 0, &state)` per thread; re-running this test always + /// produces identical samples. + #[test] + #[ignore = "requires GPU"] + fn plasticity_injection_kernel_resets_last_10pct_kaiming_he() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_PERSISTENCE_INDEX, PLASTICITY_FIRED_THIS_FOLD_INDEX, + PLASTICITY_PERSISTENCE_THRESHOLD_INDEX, PLASTICITY_WARM_BARS_REMAINING_INDEX, + }; + + let stream = make_test_stream(); + + // ISV: ensure the trigger fires — large persistence, small + // threshold, fired flag clear, warm-bars 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[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); + + // Weight buffer: n_weights = 1280, reset_count = 128 (10%). + // Pre-fill every slot with 1.0 so we can detect both + // "unchanged" (first 90%) and "changed" (last 10%) by exact + // float comparison + sample-statistic checks. + const N_WEIGHTS: usize = 1280; + const FAN_IN: i32 = 256; + const SEED: u64 = 42; + const M_WARM: f32 = 200.0; + let weights_buf = unsafe { MappedF32Buffer::new(N_WEIGHTS) } + .expect("alloc MappedF32Buffer for advantage_head_weights"); + weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS]); + + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &stream, + isv_buf.dev_ptr, + weights_buf.dev_ptr, + N_WEIGHTS as i32, + M_WARM, + FAN_IN, + SEED, + ) + .expect("launch plasticity_injection_kernel (kaiming-he)"); + stream + .synchronize() + .expect("synchronize after plasticity_injection_kernel launch (kaiming-he)"); + + // ISV side-effects: fired flipped to 1, warm = M_warm - 1. + let isv_host = isv_buf.read_all(); + assert!( + (isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX] - 1.0).abs() < 1e-9, + "PLASTICITY_FIRED_THIS_FOLD = {}, expected 1.0 after fire", + isv_host[PLASTICITY_FIRED_THIS_FOLD_INDEX] + ); + let warm = isv_host[PLASTICITY_WARM_BARS_REMAINING_INDEX]; + assert!( + (warm - (M_WARM - 1.0)).abs() < 1e-6, + "PLASTICITY_WARM_BARS_REMAINING = {warm}, expected {} (M_warm - 1, per-bar decrement runs same call)", + M_WARM - 1.0 + ); + + // Weight-reset side-effects. + let weights_host = weights_buf.read_all(); + let reset_count: usize = N_WEIGHTS / 10; // 128 + let reset_start: usize = N_WEIGHTS - reset_count; // 1152 + + // First 90% must be bit-identical to the pre-launch value. + for (i, w) in weights_host[..reset_start].iter().enumerate() { + assert!( + (w - 1.0).abs() < 1e-9, + "weights[{i}] = {w}, expected 1.0 (outside reset region — kernel must not touch)" + ); + } + + // Last 10%: every slot must have moved off 1.0 (a Kaiming + // stddev≈0.088 makes the sentinel 1.0 ~11σ out, so an exact + // 1.0 sample is effectively zero-probability). + let resets = &weights_host[reset_start..]; + for (i, w) in resets.iter().enumerate() { + assert!( + (w - 1.0).abs() > 1e-6, + "weights[{}] = {w}, expected resampled value ≠ 1.0", + reset_start + i + ); + } + + // Mean of resets ≈ 0 within sampling-noise envelope. + let n = resets.len() as f64; + let mean: f64 = resets.iter().map(|&w| w as f64).sum::() / n; + assert!( + mean.abs() < 0.05, + "Kaiming-reset region mean = {mean:.6}, expected |mean| < 0.05 (n={n}, true SE ~ 0.0078)" + ); + + // Std of resets ≈ sqrt(2/fan_in) within ±20%. + let expected_std = (2.0_f64 / FAN_IN as f64).sqrt(); + let var: f64 = resets + .iter() + .map(|&w| (w as f64 - mean).powi(2)) + .sum::() + / (n - 1.0); + let std = var.sqrt(); + assert!( + (std - expected_std).abs() / expected_std < 0.20, + "Kaiming-reset region std = {std:.6}, expected {expected_std:.6} ± 20%" + ); + + // Determinism re-check: same (seed, tid) on a fresh buffer + // produces bit-identical samples (catches accidental + // entropy sources like clock-derived seeds inside the + // kernel). Re-launch with the same seed against a fresh + // 1.0-filled buffer — but note ISV state must also be reset + // (fired=0) so the trigger re-fires. + weights_buf.write_from_slice(&vec![1.0f32; N_WEIGHTS]); + 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); + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_plasticity_injection( + &stream, + isv_buf.dev_ptr, + weights_buf.dev_ptr, + N_WEIGHTS as i32, + M_WARM, + FAN_IN, + SEED, // same seed + ) + .expect("launch plasticity_injection_kernel (determinism re-check)"); + stream + .synchronize() + .expect("synchronize after plasticity_injection_kernel launch (determinism)"); + let resets_2 = weights_buf.read_all(); + for i in 0..reset_count { + let a = resets[i]; + let b = resets_2[reset_start + i]; + assert!( + (a - b).abs() < 1e-9, + "non-deterministic Kaiming sample at tid={i}: first={a}, second={b}" + ); + } } /// SP15 Phase 3.5.5 — DD_TRAJECTORY_DECREASING fires when dd_pct diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c9dd407a9..31829759a 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 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.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). SP15 Wave 4.1b — consumer migration: s1_input_dim 102→103, GRN w_s1 reshape, 4 forward + 3 backward sites wired (2026-05-07): atomic consumer migration that flips production callers of `bn_tanh_concat_kernel` over to Wave 4.1a's `bn_tanh_concat_dd_kernel`, eliminates the documented Wave 4.1a transient orphan, and bumps `s1_input_dim` from 102 to 103 across the entire trunk forward + backward path. Per `feedback_no_partial_refactor`: every consumer of the formula `bottleneck_dim + (STATE_DIM − market_dim)` migrates atomically — no parallel paths, no feature flag toggling old vs new dim. Per `feedback_wire_everything_up`: the Wave 4.1a launcher `launch_sp15_bn_concat_dd` and its underlying `bn_tanh_concat_dd_kernel` symbol now have production callers in this commit; the legacy `bn_tanh_concat_kernel` field is deleted from the trainer struct alongside its loader, eliminating the dead handle. (1) **`s1_input_dim` formula bump from `bn_dim + portfolio_dim` to `bn_dim + portfolio_dim + 1`** at four sites: `compute_param_sizes` (gpu_dqn_trainer.rs ~line 4002, the structural source of truth that drives GRN w_a_h_s1[0] and w_residual_h_s1[4] tensor sizes via `cfg.shared_h1 * s1_input_dim`), the trainer constructor's CublasGemmSet creation site (~19720 — feeds the cuBLAS forward gemm cache shape keys), `xavier_init_params_buf` (~29821 — drives Xavier-uniform init's (fan_in=s1_input_dim, fan_out=shared_h1) tuple for the new dd_pct column), and the matching site in `gpu_experience_collector.rs` (~1148, where the experience collector builds its own `CublasGemmSet` for the inference forward path); plus the backward gemm cache site in `batched_backward.rs` (~303 — the dW/dX shape keys for h_s1's Linear_a / Linear_residual now key on K=103 instead of 102). The cuBLAS gemm cache is built once at construction and re-keyed automatically — no manual invalidation needed because the cache is a fresh `HashMap` per `CublasGemmSet::new`. (2) **GRN `w_a_h_s1[0]` and `w_residual_h_s1[4]` reshape from `[shared_h1, 102]` to `[shared_h1, 103]`** flows naturally from the s1_input_dim bump: `compute_param_sizes`'s array entries `cfg.shared_h1 * s1_input_dim` at indices [0] and [4] absorb the +1; `xavier_init_params_buf`'s `fan_dims[0..95]` core_fan tuples `(cfg.shared_h1, s1_input_dim)` at indices [0] and [4] absorb the +1 in the (fan_out, fan_in) drive for Xavier-uniform init — the new column at index 102 (the dd_pct slot) gets standard Xavier-uniform initialisation alongside the original 102 columns. Mirrors the SP14 EGF aux→Q wire pattern that bumped `w_b0fc` to `[adv_h, shared_h2 + 1]`, except SP14 explicitly zeroed the appended column post-Xavier (because aux_softmax_diff is bidirectional ±1 which collides with Xavier's zero-mean assumption); dd_pct is non-negative bounded ∈ [0, 1] (Wave 1.3 contract), so symmetric Xavier-uniform init is the correct semantic — small-magnitude gradients flow from day 0, the network learns to attend to dd_pct via SGD without artificial cold-start zero. The flat parameter buffer auto-grows by `shared_h1 × 2 × 4 = shared_h1 × 8` bytes (2 tensors × 1 column × 4 B/f32) — each padded by `align4` so the actual byte growth tracks. (3) **`bn_concat_dim()` accessor bumped by +1** (gpu_dqn_trainer.rs ~10624) — TLOB backward in `fused_training.rs:1610` reads this to compute its row-stride into `bn_d_concat_buf`. The TLOB offset arithmetic `bottleneck_dim + (TLOB_START − market_dim)` is unchanged (it indexes into the portfolio slice, which is stride-stable), but the row stride is now 103 instead of 102. The change keeps TLOB's gradient slicing aligned with the new bn_d_concat layout. (4) **5 `concat_dim` local-variable bumps** at all sites computing `bn_dim + portfolio_dim` directly: 1 buffer-allocation site (gpu_dqn_trainer.rs:20068, the constructor's `bn_concat_buf` / `tg_bn_concat_buf` / `on_next_bn_concat_buf` / `bn_d_concat_buf` allocator — all 4 buffers grow by `B × 1 × 4` bytes from the `+1` column), 3 forward sites (DDQN ~27135, online ~27431, target ~27641), 2 backward sites (`aux_bottleneck_vsn_backward_dispatch` ~28522 and the main bottleneck backward in `apply_iqn_trunk_gradient`-equivalent ~29022), plus the matching site in `gpu_experience_collector.rs` (~3820). The `concat_dim` local in each site is the row stride passed to the kernel; the kernel's read-window for bn_tanh / portfolio is unchanged ([0..bn_dim) for tanh, [bn_dim..bn_dim+portfolio_dim) for portfolio passthrough), so the dd_pct column at index `bn_dim + portfolio_dim` is the new write column on the forward path and the new ignored column on the backward path. (5) **3 forward-site migrations from `self.bn_tanh_concat_kernel` raw `launch_builder` to `launch_sp15_bn_concat_dd` free function** (DDQN ~27158, online ~27467, target ~27666) — the new launcher takes the additional `isv: CUdeviceptr` arg (= `self.isv_signals_dev_ptr`); the kernel reads slot 406 (DD_PCT_INDEX) on-device and broadcasts the scalar across batch as the (concat_dd_dim − 1)th column. The free-function call replaces the inline 8-arg `launch_builder` block with a 10-arg launcher call, and the cubin module is loaded by the launcher itself (cubin pointer cache is process-lifetime — no per-call cubin reload cost; the loader's `Arc` is reused via cudarc's internal symbol cache). (6) **gpu_experience_collector.rs forward kernel migration** — the experience collector's bn_tanh_concat launch site (~3853) is the fourth forward call site listed in the Wave 4.1a audit comment block (the trainer has 3, the collector has 1 — together they form the "4 launch sites" the comment refers to). The collector loads `bn_tanh_concat_dd_kernel` instead of `bn_tanh_concat_f32_kernel` (both kernels have all-f32 I/O; the new kernel adds the ISV arg + appended dd_pct column). The collector's `isv_signals_dev_ptr` field is set by the trainer via `set_isv_signals_ptr` from `training_loop.rs:859` BEFORE the first epoch's forward graph capture, so the captured graph sees a non-null bus pointer at replay time. Per the one-step-lag semantic (mirrors mag_concat): step k's bn_tanh_concat_dd reads ISV[406] which holds the dd_pct value `launch_sp15_dd_state` wrote at step k−1 (the dd_state launch happens AFTER the captured forward graph in the collector's per-step env loop, line ~4546). At step 0 the slot is sentinel 0 per `state_reset_registry` — bootstrap-correct, not a stub. (7) **GRN backward sites — no kernel signature change required**: the 3 GRN backward call sites in gpu_dqn_trainer.rs (main backward chain ~11036 in `apply_iqn_trunk_gradient`, ensemble backward chain ~11336, CQL backward chain ~12006) all flow through `BatchedForward::encoder_backward_chain` which uses `self.s1_input_dim` for the dX shape — bumped automatically by step (1)'s formula change. The dX writes a 103-column gradient into `bn_d_concat_buf`; the dd_pct column at index 102 carries a gradient that has no learnable input source (dd_pct is read from ISV — a constant from the kernel's perspective). The `vsn_d_gated_state_portfolio_pad_kernel` (line 28546 / 29091) reads only [bn_dim..bn_dim+portfolio_dim) so it correctly skips the dd_pct gradient column; the `bn_tanh_backward_kernel` (line 28519 / 29027) reads only [0..bn_dim) so it also correctly skips. The dd_pct column's gradient is silently discarded — harmless because no parameter chain-rules through it (the chain is broken at the ISV bus boundary, exactly as designed by Wave 1.3.b's "ISV is a producer/consumer bus, not a learnable signal" contract). (8) **mag_concat / OFI concat audit verdict — DECOUPLED**: the `mag_concat_buf` shape is `[B, shared_h2 + branch_0_size]` (fixed direction-conditioning width = 4 per production config), and `ord_concat_buf` / `urg_concat_buf` are `[B, shared_h2 + 3]` (fixed OFI-conditioning width = 3). Neither uses `s1_input_dim` — they widen `shared_h2` (the second shared-trunk hidden size, post-encoder), not the trunk INPUT dim. The bump from 102→103 is structurally invisible to mag_concat / OFI, and no migration is required. Confirmed by exhaustive grep over `crates/ml/src/cuda_pipeline/`. (9) **Legacy `bn_tanh_concat_kernel` trainer field DELETED** per `feedback_no_legacy_aliases`: the `CudaFunction` field on the trainer struct, the corresponding loader call in `compile_training_kernels` (~31551), the entry in the 44-tuple return type (~31437), the entry in the 44-element destructuring at the constructor (~17533), and the field assignment in the `Self {}` literal (~22358) are all removed in lockstep. The function tuple shrinks to 43 elements. The kernel symbol itself remains in the cubin source (`dqn_utility_kernels.cu:771`) for the SP15 oracle parity tests in `crates/ml/tests/sp15_phase1_oracle_tests.rs` — but no production code path resolves it anymore. (10) **Test migration**: `test_gpu_backtest_evaluator_state_dim_calculation` (gpu_backtest_evaluator.rs:3315) was a pre-existing failure in the 946/13 baseline — it asserted `STATE_DIM == 96` while the canonical constant in `ml_core::state_layout` is 128 (the codebase evolved 48→112→128 as the OFI / TLOB / MTF blocks expanded). Per `feedback_trust_code_not_docs`, the test was migrated to assert 128; the docstring is updated to record the 96-dim layout's deprecation date and refer maintainers to `ml_core::state_layout` as the source of truth. The fix is in this commit because the same `feedback_trust_code_not_docs` correction applies to both Wave 4.1a's "state_dim 48→49" terminology fix and this test — landing them together preserves audit-trail continuity. (11) **CUDA Graph capture interaction**: per `pearl_no_host_branches_in_captured_graph`, the new kernel launches all happen via `launch_sp15_bn_concat_dd` (a free function with no host branches inside the launcher itself — the only Rust-side branch is the `if self.config.bottleneck_dim > 0` check in the call-site, which is a host-side compile-time-equivalent guard that resolves once before graph capture begins). The `module.load_cubin` + `module.load_function` calls inside the launcher would normally be a graph-capture concern, but cudarc's internal cubin cache makes them no-ops on second + subsequent capture (the `Arc` is cached per-context). On first capture, the cubin is already resident from the constructor's eager load, so the launcher's load calls hit the cache and do not introduce graph-replay-incompatible operations. (12) **xavier_init forward-wire log line preserved**: `info!(total_params = total, s1_input_dim, "Xavier-initialized params_buf and target_params_buf")` (~30157) automatically reflects the bumped value (e.g. `s1_input_dim = 103`) — operators reading the log post-deploy see the new dim immediately, no log-format change needed. Touched: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (s1_input_dim formula bump at compute_param_sizes + constructor + xavier_init + bn_concat_dim accessor; concat_dim local-var bump at 1 alloc + 3 forward + 2 backward sites; 3 forward-call migrations to `launch_sp15_bn_concat_dd`; legacy field + loader + tuple-element + destructuring + assignment removal — 5 mechanical sites for the 1 dead field), `crates/ml/src/cuda_pipeline/batched_backward.rs` (s1_input_dim formula bump in CublasBackwardSet::new for the backward gemm cache), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (s1_input_dim formula bump + concat_dim local-var bump + bn_tanh_concat_fn loader symbol change to bn_tanh_concat_dd_kernel + exp_bn_concat allocation +1 column + launch_builder arg list extended with isv_signals_dev_ptr), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (test migration: state_dim_calculation 96→128 + docstring rationale update). Hard rules: `feedback_no_partial_refactor` (every consumer of `s1_input_dim` and `bn_concat_buf` row-stride migrates in this commit; zero parallel paths; the trainer's legacy `bn_tanh_concat_kernel` field + its loader + its tuple slot + its destructuring + its assignment all delete in lockstep), `feedback_wire_everything_up` (Wave 4.1a's `launch_sp15_bn_concat_dd` and `bn_tanh_concat_dd_kernel` symbol — both ORPHAN at end of Wave 4.1a — both have production callers in this commit; the orphan window opened in `a8da1cb9c` and closes here), `feedback_no_legacy_aliases` (the trainer's `bn_tanh_concat_kernel` field is the production-path alias that this commit deletes; the kernel source itself is retained ONLY for oracle parity tests, not as a deprecated wrapper around the new kernel), `feedback_no_cpu_compute_strict` (no CPU compute introduced — every dim-bump propagates through GPU kernel launches; xavier_init still constructs on CPU per the established offline init pattern, which is exempted from the no-CPU-compute rule because it runs once at construction before any forward pass), `feedback_isv_for_adaptive_bounds` (s1_input_dim is structural, not adaptive — no ISV slot added; dd_pct itself is ISV-driven via `dd_state_kernel` per Wave 1.3.b's existing contract), `feedback_no_stubs` (every dim bump is a real propagation; no zero-init shortcut for the new column — Xavier-uniform init handles it), `feedback_no_quickfixes` (the `state_dim_calculation` test was migrated, not deleted, and the migration carries a docstring rationale that records WHY the constant changed — future maintainers can audit the 48→112→128 evolution chain through the audit doc + state_layout.rs comments), `feedback_no_functionality_removal` (the legacy bn_tanh_concat kernel symbol stays in the cubin for oracle tests; only the dead production field is removed), `pearl_no_host_branches_in_captured_graph` (kernel launches go through `launch_sp15_bn_concat_dd` whose only Rust-side branch is in the call-site's `if bottleneck_dim > 0` guard — pre-graph-capture host-side, not in-capture). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 pre-existing unrelated warnings, all `unused import: DevicePtrMut`); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored bn_concat dd_pct --nocapture` 1 of 1 oracle test green (`bn_tanh_concat_dd_kernel_writes_dd_pct_column` — Wave 4.1a's standalone parity test still passes, confirming the kernel-side contract is unchanged); `cargo test -p ml --features cuda --lib` HELD-OR-IMPROVED the 946 pass / 13 fail Wave 4.1a baseline (the migrated `test_gpu_backtest_evaluator_state_dim_calculation` flips from FAIL to PASS, bringing the count to ≥947 pass / ≤12 fail post-Wave-4.1b — actual count documented in the commit message body).