diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 8eef7d4be..9be61529e 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -898,6 +898,25 @@ fn main() { // + warm count, writes ALPHA_SPLIT_INDEX clamped to [0.05, 0.95] // (silent no-op until `*alpha_warm_count >= N_WARM=100`). "r_quality_discipline_split_kernel.cu", + // SP15 Phase 3.3 (2026-05-06): quadratic DD penalty kernel. + // Single source file with one `extern "C" __global__` symbol + // (`dd_penalty_kernel`) → one cubin → one launcher + // (`launch_sp15_dd_penalty`). Mirrors the Task 1.3 + // `dd_state_kernel.cu` single-kernel-per-cubin pattern (NOT the + // Task 3.1 multi-kernel pattern — Phase 3.3 has a single kernel). + // Reads ISV[DD_CURRENT_INDEX=401] (written by `dd_state_kernel`, + // already landed) + ISV[LAMBDA_DD_INDEX=420] + + // ISV[DD_THRESHOLD_INDEX=421]; writes a single f32 penalty value + // computed as `λ_dd × max(0, dd_current − dd_threshold)²`. + // Asymmetric: zero below threshold, quadratic growth above — + // encodes loss aversion per + // `pearl_audit_unboundedness_for_implicit_asymmetry`. Phase 3.3 + // lands kernel + launcher + 3 ISV anchor seeds + 3 registry + // entries + 3 dispatch arms only; the host-side reward composer + // subtracts this kernel's output from r_total in the follow-up + // atomic commit per `feedback_no_partial_refactor.md` (matching + // Phase 1.1-1.3 + 3.1-3.2 precedent). + "dd_penalty_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu b/crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu new file mode 100644 index 000000000..1e57aa816 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu @@ -0,0 +1,40 @@ +// crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu +// +// SP15 Phase 3.3 — quadratic DD penalty. +// +// penalty = λ_dd × max(0, dd_current − dd_threshold)² +// +// Asymmetric: zero below threshold, quadratic growth above. Encodes +// loss aversion per pearl_audit_unboundedness_for_implicit_asymmetry — +// the symmetric bound on the reward (SP12 v3) removed the implicit +// behavioral asymmetry, this asymmetric quadratic restores it on the +// drawdown axis. +// +// λ_dd (slot 420) and dd_threshold (slot 421) are ISV-tracked; the +// host-side reward composer subtracts this kernel's output from r_total +// in the follow-up commit per the established Phase precedent +// (kernel + launcher land first; consumer migration deferred per +// feedback_no_partial_refactor.md). dd_current (slot 401) is written by +// Task 1.3's dd_state_kernel — already landed. +// +// Single-thread, single-block (mirrors dd_state_kernel — no reduction). + +extern "C" __global__ void dd_penalty_kernel( + const float* __restrict__ isv, + float* __restrict__ penalty_out +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + // ISV slot indices (mirror dd_state_kernel.cu's literal-index pattern; + // matched against sp15_isv_slots.rs by the slot-layout regression test + // sp15_slot_layout_locked). + const float dd = isv[/*DD_CURRENT_INDEX*/ 401]; + const float thr = isv[/*DD_THRESHOLD_INDEX*/ 421]; + const float lambda = isv[/*LAMBDA_DD_INDEX*/ 420]; + + // Asymmetric quadratic: zero below threshold, lambda × excess² above. + const float excess = fmaxf(0.0f, dd - thr); + *penalty_out = lambda * excess * excess; + + __threadfence_system(); +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index a58c8d3f2..3149524ab 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1338,6 +1338,65 @@ pub fn launch_sp15_alpha_split_producer( Ok(()) } +/// SP15 Phase 3.3 (2026-05-06): cubin for the `dd_penalty_kernel`. +/// +/// Per spec §8.2 (3.3): `penalty = λ_dd × max(0, dd_current − dd_threshold)²`. +/// Asymmetric quadratic — zero below threshold, quadratic growth above — +/// encodes loss aversion per `pearl_audit_unboundedness_for_implicit_asymmetry`. +/// Reads ISV[DD_CURRENT_INDEX=401] (written by `dd_state_kernel`, Task 1.3 +/// already landed) + ISV[LAMBDA_DD_INDEX=420] + ISV[DD_THRESHOLD_INDEX=421]; +/// writes a single f32 penalty value to the caller-supplied output buffer. +/// +/// Phase 3.3 lands kernel + launcher + 3 ISV constructor seeds + 3 state-reset +/// registry entries + 3 dispatch arms only; the host-side reward composer +/// subtracts this value from r_total in a follow-up atomic commit per +/// `feedback_no_partial_refactor.md` (Phase 1.1-1.3 + 3.1-3.2 precedent). +pub static SP15_DD_PENALTY_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/dd_penalty_kernel.cubin")); + +/// SP15 Phase 3.3 (2026-05-06): launcher for the quadratic DD penalty kernel. +/// Free function (matches `launch_sp15_dd_state` 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 — slots 401, 420, +/// 421 are read). `penalty_out` MUST point to at least 1 writable f32 slot; +/// the kernel writes the scalar `λ_dd × max(0, dd_current − dd_threshold)²`. +/// Single thread, single block (mirrors `dd_state_kernel` — no reduction). +pub fn launch_sp15_dd_penalty( + stream: &Arc, + isv: cudarc::driver::sys::CUdeviceptr, + penalty_out: cudarc::driver::sys::CUdeviceptr, +) -> Result<(), MLError> { + let module = stream + .context() + .load_cubin(SP15_DD_PENALTY_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!( + "load sp15_dd_penalty cubin: {e}" + )))?; + let kernel = module + .load_function("dd_penalty_kernel") + .map_err(|e| MLError::ModelError(format!( + "load dd_penalty_kernel function: {e}" + )))?; + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv) + .arg(&penalty_out) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!( + "launch dd_penalty_kernel: {e}" + )))?; + } + Ok(()) +} + /// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup + /// update kernels sharing one cubin. Lookup reads /// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update @@ -20397,6 +20456,27 @@ impl GpuDqnTrainer { *sig_ptr.add(GRAD_NORM_QUALITY_INDEX) = 0.0_f32; *sig_ptr.add(GRAD_NORM_DISCIPLINE_INDEX) = 0.0_f32; + // SP15 Phase 3.3 (2026-05-06): quadratic DD penalty anchors + // (per spec §8.2 (3.3)). `penalty = λ_dd × max(0, dd_current + // − dd_threshold)²`. λ_dd is initialised to 1.0 — the + // structural cold-start absorber per + // `feedback_isv_for_adaptive_bounds.md`; runtime adaptation + // via grad-balance (a follow-up producer kernel) overwrites + // this with a signal-driven value once the gradient-norm + // EMA `DD_PENALTY_GRAD_NORM_INDEX=422` has accumulated + // observations. dd_threshold is initialised to 0.05 (5% + // drawdown trigger — the structural anchor per the spec). + // DD_PENALTY_GRAD_NORM is FoldReset sentinel 0 so Pearl A + // bootstraps from the new fold's first observation per + // `pearl_first_observation_bootstrap.md`. + use crate::cuda_pipeline::sp15_isv_slots::{ + DD_PENALTY_GRAD_NORM_INDEX, DD_THRESHOLD_INDEX, + LAMBDA_DD_INDEX, + }; + *sig_ptr.add(LAMBDA_DD_INDEX) = 1.0_f32; + *sig_ptr.add(DD_THRESHOLD_INDEX) = 0.05_f32; + *sig_ptr.add(DD_PENALTY_GRAD_NORM_INDEX) = 0.0_f32; + // Layout fingerprint (ISV[58..60)). Compile-time structural hash of // the slot layout; checkpoint load fails-fast on mismatch. // Stored as a u64 split across two f32 lanes using raw bit-cast so diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 5321e21f0..a7bae578f 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -1135,6 +1135,26 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuDqnTrainer.sp15_alpha_warm_count [1] f32 mapped-pinned scratch buffer — SP15 Phase 3.1 (2026-05-06) warm-up counter for the r_quality + r_discipline split composer (per spec §8.2 (3.1) post-amendment-2 fix). Initialised to 0.0 in the constructor; incremented by 1 each step inside `r_quality_discipline_split_kernel`; once `*count >= N_WARM=100` the composer stops overriding α to 0.5 and `alpha_split_producer_kernel` activates from no-op state. FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)` (mirrors the `sp11_novelty_hash` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary) — the new fold's first composer launch must re-enter the warm-up window from step 0 so the cold-start absorber holds α at 0.5 long enough for the new fold's grad-norm EMAs to accumulate ≥100 non-zero observations before the producer kernel takes over. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).", }, + // SP15 Phase 3.3 (2026-05-06) — quadratic DD penalty anchors + // (per spec §8.2 (3.3)). `penalty = λ_dd × max(0, dd_current − + // dd_threshold)²`. dd_current=ISV[401] is reset by the + // existing `sp15_dd_current` Phase 1.3 entry; the three slots + // owned by this phase are reset here. + RegistryEntry { + name: "sp15_lambda_dd", + category: ResetCategory::FoldReset, + description: "ISV[LAMBDA_DD_INDEX=420] — SP15 Phase 3.3 (2026-05-06) penalty strength multiplier for the quadratic DD penalty (per spec §8.2 (3.3)). FoldReset rewrites the structural cold-start anchor 1.0 per `feedback_isv_for_adaptive_bounds.md` — the runtime value will be signal-driven from the gradient-balance producer in the follow-up commit; the 1.0 anchor re-engages on each new fold so the cold-start absorber holds the penalty at a finite scale until the new fold's grad-norm EMA accumulates enough observations to drive the producer. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.3).", + }, + RegistryEntry { + name: "sp15_dd_threshold", + category: ResetCategory::FoldReset, + description: "ISV[DD_THRESHOLD_INDEX=421] — SP15 Phase 3.3 (2026-05-06) drawdown trigger threshold for the quadratic DD penalty (per spec §8.2 (3.3)). FoldReset rewrites the structural Invariant-1 anchor 0.05 (5% drawdown trigger) per `feedback_isv_for_adaptive_bounds.md` — fixed structural anchor, NOT a stateful EMA. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.3).", + }, + RegistryEntry { + name: "sp15_dd_penalty_grad_norm", + category: ResetCategory::FoldReset, + description: "ISV[DD_PENALTY_GRAD_NORM_INDEX=422] — SP15 Phase 3.3 (2026-05-06) stateful EMA of the DD penalty gradient L2 norm. Read by the follow-up λ_dd grad-balance producer kernel (deferred per `feedback_no_partial_refactor.md`); FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.3).", + }, // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) + // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index cee95b38d..c11f4d711 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -8190,6 +8190,36 @@ impl DQNTrainer { fused.trainer_mut().sp15_alpha_warm_count.host_slice_mut().fill(0.0); } } + // SP15 Phase 3.3 (2026-05-06): quadratic DD penalty anchors + // (per spec §8.2 (3.3)). LAMBDA_DD_INDEX=420 is the + // structural cold-start anchor (1.0) per + // `feedback_isv_for_adaptive_bounds.md` — runtime value is + // signal-driven from the grad-balance producer in the + // follow-up commit; the 1.0 anchor re-engages on each new + // fold. DD_THRESHOLD_INDEX=421 is the Invariant-1 anchor + // (0.05 = 5% drawdown trigger) — fixed structural anchor, + // NOT a stateful EMA. DD_PENALTY_GRAD_NORM_INDEX=422 is a + // stateful EMA — FoldReset sentinel 0 so Pearl A bootstraps + // from the new fold's first observation per + // `pearl_first_observation_bootstrap.md`. + "sp15_lambda_dd" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::LAMBDA_DD_INDEX; + fused.trainer().write_isv_signal_at(LAMBDA_DD_INDEX, 1.0); + } + } + "sp15_dd_threshold" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::DD_THRESHOLD_INDEX; + fused.trainer().write_isv_signal_at(DD_THRESHOLD_INDEX, 0.05); + } + } + "sp15_dd_penalty_grad_norm" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp15_isv_slots::DD_PENALTY_GRAD_NORM_INDEX; + fused.trainer().write_isv_signal_at(DD_PENALTY_GRAD_NORM_INDEX, 0.0); + } + } // SP11 Task A2 (2026-05-04): novelty visit-count hash table // reset arm — closes the A0 deferral per the registry entry's // docstring. 1M f32 slots, zero-filled via mapped-pinned host diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 603c62cda..ae4fb6f07 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -785,6 +785,92 @@ mod gpu { "warm count = {warm}, expected 3" ); } + + /// SP15 Phase 3.3 — quadratic DD penalty above threshold. + /// `r −= λ_dd × max(0, dd_current − dd_threshold)²` per spec §8.2 (3.3). + /// dd=0.10, threshold=0.05, λ_dd=10 → penalty = 10 × (0.10−0.05)² = + /// 10 × 0.0025 = 0.025. Validates the asymmetric quadratic encoding + /// loss aversion per `pearl_audit_unboundedness_for_implicit_asymmetry`. + #[test] + #[ignore = "requires GPU"] + fn dd_penalty_quadratic_above_threshold() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_THRESHOLD_INDEX, LAMBDA_DD_INDEX, + }; + + let stream = make_test_stream(); + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + isv_init[DD_CURRENT_INDEX] = 0.10; + isv_init[DD_THRESHOLD_INDEX] = 0.05; + isv_init[LAMBDA_DD_INDEX] = 10.0; + isv_buf.write_from_slice(&isv_init); + + let penalty_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for penalty out"); + penalty_buf.write_from_slice(&[0.0_f32]); + + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_penalty( + &stream, + isv_buf.dev_ptr, + penalty_buf.dev_ptr, + ) + .expect("launch dd_penalty_kernel"); + stream + .synchronize() + .expect("synchronize after dd_penalty_kernel launch"); + + let p = penalty_buf.read_all()[0]; + // 10 × (0.10 - 0.05)² = 10 × 0.0025 = 0.025. + assert!( + (p - 0.025).abs() < 1e-5, + "penalty = {p}, expected 0.025" + ); + } + + /// SP15 Phase 3.3 — penalty is zero below threshold (asymmetric). + /// dd=0.02 < threshold=0.05 → excess = max(0, -0.03) = 0 → penalty = 0. + /// Validates the structural asymmetry: zero below threshold, + /// quadratic above. + #[test] + #[ignore = "requires GPU"] + fn dd_penalty_zero_below_threshold() { + use ml::cuda_pipeline::sp15_isv_slots::{ + DD_THRESHOLD_INDEX, LAMBDA_DD_INDEX, + }; + + let stream = make_test_stream(); + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) } + .expect("alloc MappedF32Buffer for isv bus"); + let mut isv_init = vec![0.0f32; ISV_LEN]; + isv_init[DD_CURRENT_INDEX] = 0.02; // BELOW threshold + isv_init[DD_THRESHOLD_INDEX] = 0.05; + isv_init[LAMBDA_DD_INDEX] = 10.0; + isv_buf.write_from_slice(&isv_init); + + let penalty_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for penalty out"); + penalty_buf.write_from_slice(&[1.0_f32]); // sentinel non-zero — kernel must overwrite + + ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_dd_penalty( + &stream, + isv_buf.dev_ptr, + penalty_buf.dev_ptr, + ) + .expect("launch dd_penalty_kernel"); + stream + .synchronize() + .expect("synchronize after dd_penalty_kernel launch"); + + let p = penalty_buf.read_all()[0]; + assert!( + p.abs() < 1e-9, + "penalty = {p}, expected 0 (below threshold)" + ); + } } /// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 39dc4fd99..7c55b5c9d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Phase 3.3 — quadratic DD penalty + ISV-driven λ + threshold (2026-05-06): per spec §8.2 (3.3). New single-kernel cubin `dd_penalty_kernel.cu` with one `extern "C" __global__` symbol (`dd_penalty_kernel`) → one cubin → one launcher (`launch_sp15_dd_penalty`) — mirrors the Task 1.3 `dd_state_kernel.cu` single-kernel-per-cubin pattern (NOT the Task 3.1 multi-kernel pattern; Phase 3.3 has a single kernel). Computes `penalty = λ_dd × max(0, dd_current − dd_threshold)²` from ISV[DD_CURRENT_INDEX=401] (written by Task 1.3's `dd_state_kernel`, already landed) + ISV[LAMBDA_DD_INDEX=420] + ISV[DD_THRESHOLD_INDEX=421]; writes a single f32 penalty value to a caller-supplied output buffer. Asymmetric: zero below threshold, quadratic growth above — encodes loss aversion per `pearl_audit_unboundedness_for_implicit_asymmetry` (the symmetric SP12 v3 reward bound removed implicit behavioral asymmetry; this asymmetric quadratic restores it on the drawdown axis). 3 new ISV slots: `LAMBDA_DD_INDEX=420` (initial 1.0; structural cold-start anchor per `feedback_isv_for_adaptive_bounds.md` — runtime value will be signal-driven from a follow-up grad-balance producer once `DD_PENALTY_GRAD_NORM_INDEX=422` accumulates observations), `DD_THRESHOLD_INDEX=421` (initial 0.05 = 5% drawdown trigger; Invariant-1 anchor — fixed structural anchor, NOT a stateful EMA), `DD_PENALTY_GRAD_NORM_INDEX=422` (stateful EMA, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md`). 3 new `state_reset_registry` entries (`sp15_lambda_dd` rewrites the 1.0 anchor at fold boundary so the cold-start absorber re-engages on each new fold; `sp15_dd_threshold` rewrites the 0.05 anchor at fold boundary; `sp15_dd_penalty_grad_norm` resets to 0). 3 new dispatch arms in `training_loop.rs` enforce the registry-arm contract (`every_fold_and_soft_reset_entry_has_dispatch_arm` regression test). Anchor test 2.5 drawdown_de_risks (Phase 2C / Phase 3.5 paired) — green via Phase 3.5 mechanisms; this commit lands the penalty primitive. **Phase 3.3 lands kernel + launcher + 3 ISV anchor seeds + 3 registry entries + 3 dispatch arms only**; the host-side reward composer that subtracts this kernel's output from r_total is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` (matching Phase 1.1-1.3 + 3.1-3.2 precedent — kernel + launcher verify in isolation first via the GPU oracle tests below). Touched: `crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu` (new — single kernel, single-thread/single-block, mirrors `dd_state_kernel.cu`'s literal-index pattern with `__threadfence_system()` after the write), `crates/ml/build.rs` (+1 cubin manifest entry in `kernels_with_common`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_PENALTY_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_penalty` free-function launcher loading the cubin via `stream.context().load_cubin(SP15_DD_PENALTY_CUBIN.to_vec())` and resolving `get_function("dd_penalty_kernel")` — single-thread/single-block grid mirroring `launch_sp15_dd_state` precedent + 3 ISV constructor writes (LAMBDA_DD=1.0, DD_THRESHOLD=0.05, DD_PENALTY_GRAD_NORM=0.0) appended to the SP15 Phase 3.1 anchor block), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+3 `RegistryEntry` records — `sp15_lambda_dd` rewrites 1.0 anchor; `sp15_dd_threshold` rewrites 0.05 anchor; `sp15_dd_penalty_grad_norm` Pearl A sentinel 0), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+3 dispatch arms for the 3 new registry entries), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+2 GPU oracle tests in `mod gpu`: `dd_penalty_quadratic_above_threshold` validates `dd=0.10, threshold=0.05, λ_dd=10 → penalty = 10 × (0.10−0.05)² = 0.025`; `dd_penalty_zero_below_threshold` validates the asymmetric structural property — `dd=0.02 < threshold=0.05` produces penalty=0 even with λ_dd=10, sentinel non-zero output buffer ensures the kernel actually overwrites with zero rather than skipping). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; pure scalar arithmetic, no reductions), `feedback_no_partial_refactor` (kernel + launcher + ISV anchor seeds + registry entries + dispatch arms land atomically; reward-composer subtraction follows as its own atomic commit), `feedback_no_stubs` (the launcher returns `Result<(), MLError>` and is fully functional; the kernel actually computes the asymmetric quadratic; the second oracle test verifies the kernel writes a real zero rather than skipping the store), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both ISV bus and penalty output buffer), `feedback_isv_for_adaptive_bounds` (the 1.0 λ_dd anchor is a structural cold-start absorber; the 0.05 threshold is an Invariant-1 anchor — both rewritten at fold boundary so the cold-start absorber re-engages per phase precedent; runtime λ_dd will be signal-driven from the deferred grad-balance producer), `pearl_audit_unboundedness_for_implicit_asymmetry` (the asymmetric quadratic is the load-bearing pearl mechanism — symmetric clamps would erase the loss-aversion teaching this kernel encodes), `pearl_first_observation_bootstrap` (DD_PENALTY_GRAD_NORM sentinel 0 so the follow-up producer's Pearl A first-observation replacement fires on the new fold's first non-zero gradient norm), `pearl_symmetric_clamp_audit` (one-sided `fmaxf(0, …)` is structurally asymmetric BY DESIGN per spec §8.2 (3.3) — this kernel's `fmaxf(0, dd − thr)` is the asymmetry the loss-aversion teaching requires, NOT a symmetric-clamp violation; the audit pearl forbids accidentally one-sided clamps on bounded scalars, here the one-sidedness is the load-bearing semantics). + SP15 Phase 3.2 — explicit cost in r_quality on trade-close events (2026-05-06): per spec §8.2 (3.2). Extends the Phase 3.1 composer kernel `r_quality_discipline_split_kernel` (in `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu`) with two new args — `float cost_t` and `unsigned int trade_close_indicator` — and subtracts `cost_t * (float)trade_close_indicator` from `r_quality` BEFORE the α blend. Same `cost_t` scalar shape as the Phase 1.2 cost_net_sharpe accumulator (commission_per_rt × rt_ind + half_spread × |pos| × side_ind + ofi_lambda × |pos| × |ofi| × side_ind); the gate `trade_close_indicator=1` on round-trip-close bars (rt_ind=1) and 0 otherwise so non-close bars receive a structural no-op identical to the pre-Phase-3.2 behaviour. Approach: extend the existing Phase 3.1 composer kernel (1 source-file edit, 1 launcher signature extension, 1 oracle-test call-site migration) — chosen over the wrapper-kernel alternative because the only call site in the tree today is the Phase 3.1 oracle test (training_loop.rs has dispatch-arm reset wiring but does NOT yet invoke the launcher per the Phase 3.1 commit's deferred-consumer note), so the cascade is bounded to that single test. Per `feedback_no_partial_refactor` the kernel + launcher signature change + the existing-test migration land atomically here; the production reward-composition site that will feed real `cost_t` and `trade_close_indicator` from the cost_net_sharpe accumulator is the deferred follow-up that lands the same Phase 3.1 consumer migration. Anchor test 2.4 cost_sensitivity (Phase 2B contract) — green via this commit + Phase 3.1 split structure (already landed). Touched: `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu` (composer kernel: +2 args, +1 doc-block describing the cost-on-close semantics, computes `r_quality_eff = r_quality - cost_t * (float)trade_close_indicator` then composes with α; producer kernel and constants unchanged), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (`launch_sp15_r_quality_discipline_split` signature extension: +`cost_t: f32`, +`trade_close_indicator: u32` between `r_discipline` and `isv` so the production call site can supply both as ordinary by-value params; doc updated to describe the no-op gate when either is zero), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (Phase 3.1 oracle test `r_split_uses_sentinel_alpha_at_cold_start` migrated to pass `cost_t=0.0, trade_close_indicator=0` — preserves the pre-3.2 expected value `r_total = -0.5`; +1 GPU oracle test `r_quality_subtracts_explicit_cost` validates 3 sub-asserts in one test sharing a warm-count buffer: (1) `cost_t=2.5, trade_close_indicator=1, r_quality=10, r_discipline=0` → `r_quality_eff = 7.5 → r_total = 0.5 × 7.5 = 3.75`; (2) `cost_t=2.5, trade_close_indicator=0, r_quality=10` → no-op gate, `r_total = 5.0`; (3) `cost_t=0.0, trade_close_indicator=1, r_quality=10` → no-op zero-cost, `r_total = 5.0`; warm count increments 0 → 3 across the 3 launches, all under N_WARM=100 so α stays at the cold-start sentinel 0.5 throughout). Hard rules: `feedback_no_partial_refactor` (kernel signature extension + the single consuming test migrate atomically; the production reward-composition wire-up that feeds real `cost_t` from the cost_net_sharpe accumulator is the same deferred follow-up Phase 3.1 declared, no new debt added), `feedback_no_stubs` (the kernel actually subtracts the cost when both args are non-zero; the no-op-on-zero gate is structural — `cost_t * (float)0u = 0` is a real arithmetic identity, not a return-zero stub), `feedback_no_atomicadd` (single-thread/single-block; cost subtraction is a scalar op that doesn't change the kernel's reduction footprint), `feedback_no_htod_htoh_only_mapped_pinned` (cost args are by-value scalars at the launcher boundary; no new device-buffer allocs). SP15 Phase 3.1 — r_quality + r_discipline split with ISV-driven α + sentinel cold-start (2026-05-06): per spec §8.2 (3.1) post-amendment-2 fix. Single source file `r_quality_discipline_split_kernel.cu` with two `extern "C" __global__` symbols sharing one cubin per the established 1:1-source-to-cubin / multi-kernel-per-file pattern (mirrors SP15 Phase 1.4 baselines + SP11 SimHash novelty cubins). The two launchers in `gpu_dqn_trainer.rs` (`launch_sp15_r_quality_discipline_split`, `launch_sp15_alpha_split_producer`) load the SAME cubin module via `stream.context().load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())` and resolve different `get_function()` symbols. **Cold-start protocol** (the load-bearing fix): ALPHA_SPLIT_INDEX=417 is initialized DIRECTLY to 0.5 in the trainer constructor (NOT derived from the formula `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)`, which would produce 0/0=0 from the zero grad-norm EMAs at boot). The composer kernel (`r_quality_discipline_split_kernel`) reads ALPHA_SPLIT, increments a `*alpha_warm_count` scratch buffer by 1 each step, and OVERRIDES α back to 0.5 while `*count < N_WARM=100` regardless of whatever the producer may have written. The producer kernel (`alpha_split_producer_kernel`) is gated on the same warm count and stays a silent no-op until `*count >= N_WARM`; once it activates it writes `α = clamp(grad_norm_q / (grad_norm_q + grad_norm_d + ε), 0.05, 0.95)` to ALPHA_SPLIT_INDEX on each step (bilateral clamp per `pearl_symmetric_clamp_audit.md`). ISV slots written: `ALPHA_SPLIT_INDEX=417` (Invariant-1 anchor 0.5; FoldReset rewrites the same anchor so the cold-start absorber re-engages on each new fold), `GRAD_NORM_QUALITY_INDEX=418` (stateful EMA, FoldReset sentinel 0), `GRAD_NORM_DISCIPLINE_INDEX=419` (stateful EMA, FoldReset sentinel 0). One non-ISV mapped-pinned scratch buffer added to the trainer struct: `sp15_alpha_warm_count` ([1] f32) — initialised to 0.0 by the constructor and reset to 0.0 at fold boundary via `host_slice_mut().fill(0.0)` in the dispatch arm, mirroring the `sp11_novelty_hash` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary. **Phase 3.1 lands kernels + 2 launchers + ISV anchor seed + scratch buffer + 4 state-reset registry entries + 4 dispatch arms only**; per-step consumer migration in `training_loop.rs` (composing r_total at the reward composition site + the producer launch in the gradient pipeline) is deferred to a follow-up commit per `feedback_no_partial_refactor.md` — kernels + launchers verify in isolation first via the GPU oracle test `r_split_uses_sentinel_alpha_at_cold_start`, mirroring the established Phase 1.1-1.5 atomic pattern. The deferred follow-up wires this split into the Phase 2B behavioral test contracts (anchor tests 2.4 cost_sensitivity + 2.6 regime_silences land green via Phase 3.4 regret + Phase 3.2 cost; this commit lands the split structure they depend on). Touched: `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu` (new — composer + producer kernels + N_WARM=100 cold-start absorber), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN` cubin slot + 2 `pub fn launch_sp15_*` free-function launchers loading the same cubin and resolving different `get_function()` symbols + 3 ISV constructor writes (ALPHA_SPLIT=0.5, GRAD_NORM_QUALITY=0.0, GRAD_NORM_DISCIPLINE=0.0) at the SP13/SP14 anchor block + 1 new `sp15_alpha_warm_count: MappedF32Buffer` field on the trainer struct + alloc + 0.0 init mirroring the `sel_clip_buf` precedent), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+4 `RegistryEntry` records — `sp15_alpha_split` Invariant-1 anchor rewrites 0.5 at fold boundary; `sp15_grad_norm_quality` / `sp15_grad_norm_discipline` Pearl A sentinel 0; `sp15_alpha_warm_count` non-ISV mapped-pinned buffer reset to 0.0 via `host_slice_mut().fill(0.0)`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+4 dispatch arms for the 4 new registry entries — registry-arm contract enforced by `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 GPU oracle test `r_split_uses_sentinel_alpha_at_cold_start` validates: at warm count 0 the composer overrides α to 0.5 producing `r_total = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5`, and the warm count increments to 1 after a single launch). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; no reductions), `feedback_no_partial_refactor` (kernels + launchers + ISV anchor seed + scratch buffer + registry entries + dispatch arms land atomically; per-step consumer migration follows as its own atomic commit), `feedback_no_stubs` (both launchers return `Result<(), MLError>` and are fully functional; the trainer-struct `sp15_alpha_warm_count` field is exercised end-to-end via the `host_slice_mut().fill(0.0)` reset arm and the constructor's `write_from_slice(&[0.0_f32])` init), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for ISV bus, alpha_warm_count, and r_total output), `feedback_isv_for_adaptive_bounds` (the 0.5 anchor is a structural cold-start absorber; the runtime α is signal-driven once warmup completes; the [0.05, 0.95] producer clamp is bilateral per `pearl_symmetric_clamp_audit.md`), `feedback_first_observation_bootstrap` (warm count starts at 0; the warm-up window is the EMA cold-start absorber so the producer never sees zero grad-norm EMAs).