feat(sp15-p3.3): quadratic DD penalty + ISV-driven λ + threshold

Per spec §8.2 (3.3). 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.

3 ISV slots: 420 LAMBDA_DD (initial 1.0; ISV-tracked from grad-balance
in follow-up), 421 DD_THRESHOLD (initial 0.05 = 5% drawdown trigger),
422 DD_PENALTY_GRAD_NORM (initial 0.0).

3 fold-reset registry entries + dispatch arms.

Per established Phase precedent: kernel + launcher land first; reward
composition site (subtract penalty from r_total) deferred to follow-up
commit per feedback_no_partial_refactor.

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.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

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)

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();
+}

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<CudaStream>,
+    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

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

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

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