diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 1677f1b3b..73a151da1 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -316,6 +316,17 @@ fn main() {
         // the same .cuh header directly; the test wrapper has no production
         // callers.
         "sp4_histogram_p99_test_kernel.cu",
+        // SP12 v3 (2026-05-04): standalone test kernel exercised only by
+        // `tests/sp12_reward_math_tests.rs`. Wraps the three reward
+        // composition device functions added to `trade_physics.cuh`
+        // (compute_asymmetric_capped_pnl, compute_min_hold_penalty,
+        // compute_lump_sum_opp_cost) so the Rust GPU oracle tests can
+        // drive each formula on synthetic inputs with analytically-known
+        // expected outputs per `feedback_no_cpu_test_fallbacks`. Producer
+        // call site is in `experience_kernels.cu::experience_env_step`
+        // segment_complete branch; the test wrapper has no production
+        // callers.
+        "sp12_reward_math_test_kernel.cu",
         // SP4 Layer A Task A5 (2026-04-30): first end-to-end SP4 producer.
         // Single-block kernel reads `denoise_target_q_buf`, computes
         // p99(|target_q|) via `sp4_histogram_p99<256>`, writes the step
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 697e81a07..4c65cbe94 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -2796,8 +2796,14 @@ extern "C" __global__ void experience_env_step(
          *
          * Constants live in state_layout.cuh as Invariant-1 numerical
          * anchors (Phase 1). Phase 2 (deferred) lifts the ratio to an
-         * ISV-driven controller per feedback_isv_for_adaptive_bounds. */
-        float capped_pnl = fmaxf(REWARD_NEG_CAP, fminf(base_reward, REWARD_POS_CAP));
+         * ISV-driven controller per feedback_isv_for_adaptive_bounds.
+         *
+         * Math factored into `compute_asymmetric_capped_pnl` in
+         * `trade_physics.cuh` so the formula is testable in isolation
+         * (oracle tests in `tests/sp12_reward_math_tests.rs`). */
+        float capped_pnl = compute_asymmetric_capped_pnl(
+            base_reward, REWARD_NEG_CAP, REWARD_POS_CAP
+        );
         if (trail_triggered) {
             r_trail = capped_pnl;
         } else {
@@ -2822,15 +2828,23 @@ extern "C" __global__ void experience_env_step(
              *
              * Applied BEFORE `r_popart = capped_pnl` so the penalty flows
              * through the popart component (which the SP11 controller
-             * weighs) — keeps mag-ratio canary semantics consistent. */
-            float adjusted_pnl = capped_pnl;
-            if (segment_hold_time < min_hold_target) {
-                const float deficit = min_hold_target - segment_hold_time;
-                const float soft_factor = deficit
-                    / fmaxf(deficit + min_hold_temperature, 1e-6f);
-                const float min_hold_penalty = min_hold_penalty_max * soft_factor;
-                adjusted_pnl -= shaping_scale * min_hold_penalty;
-            }
+             * weighs) — keeps mag-ratio canary semantics consistent.
+             *
+             * Math factored into `compute_min_hold_penalty` in
+             * `trade_physics.cuh` so the formula is testable in isolation
+             * (oracle tests in `tests/sp12_reward_math_tests.rs`). The
+             * device function returns the penalty MAGNITUDE (positive) and
+             * subsumes the `hold_time >= target` early-exit; preserving the
+             * outer `if (segment_hold_time < min_hold_target)` guard would
+             * be redundant since the function already returns 0 in that
+             * case and `capped_pnl - shaping_scale * 0 == capped_pnl`. */
+            const float min_hold_penalty = compute_min_hold_penalty(
+                segment_hold_time,
+                min_hold_target,
+                min_hold_temperature,
+                min_hold_penalty_max
+            );
+            float adjusted_pnl = capped_pnl - shaping_scale * min_hold_penalty;
             r_popart = adjusted_pnl;
             /* SP11 Fix 39 B1b fix-up (2026-05-04): mirror r_popart to the
              * dedicated per-bar buffer feeding ISV[POPART_COMPONENT_MAG_EMA_
@@ -2885,12 +2899,16 @@ extern "C" __global__ void experience_env_step(
          * pre-helper hold count (set to `saved_hold_time` at line ~2493).
          * `shaping_scale` gates the term off in pure-validation mode (=0)
          * matching the spec's "step_returns are pure per-bar P&L"
-         * convention for backtest reward shape. */
-        const float lump_opp_cost = shaping_scale
-                                  * holding_cost_rate
-                                  * fabsf(pre_trade_position)
-                                  * segment_hold_time;
-        r_opp_cost = -lump_opp_cost;
+         * convention for backtest reward shape.
+         *
+         * Math factored into `compute_lump_sum_opp_cost` in
+         * `trade_physics.cuh` so the formula is testable in isolation
+         * (oracle tests in `tests/sp12_reward_math_tests.rs`). The device
+         * function unconditionally returns the (negative) reward magnitude;
+         * the segment-complete gating happens at this call site. */
+        r_opp_cost = compute_lump_sum_opp_cost(
+            shaping_scale, holding_cost_rate, pre_trade_position, segment_hold_time
+        );
         reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
 
         /* ── Layer 3: CEA counterfactual loop REMOVED (4-branch: direction(4) replaces exposure(9)).
diff --git a/crates/ml/src/cuda_pipeline/sp12_reward_math_test_kernel.cu b/crates/ml/src/cuda_pipeline/sp12_reward_math_test_kernel.cu
new file mode 100644
index 000000000..d69592f2c
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/sp12_reward_math_test_kernel.cu
@@ -0,0 +1,65 @@
+/* SP12 v3 reward-math GPU oracle test kernel.
+ *
+ * Standalone test wrapper that exercises the three device-inline reward
+ * composition functions added to `trade_physics.cuh` for SP12 v3:
+ *   - compute_asymmetric_capped_pnl
+ *   - compute_min_hold_penalty
+ *   - compute_lump_sum_opp_cost
+ *
+ * One thread per kernel launch is sufficient — these functions are pure
+ * register-only arithmetic with no shared state. The wrapper mirrors the
+ * production kernel's call sites bit-for-bit so the tests verify the
+ * exact math the production reward chain executes (no CPU reference impl
+ * per `feedback_no_cpu_test_fallbacks`).
+ *
+ * Outputs go to mapped-pinned buffers (cuMemHostAlloc DEVICEMAP|PORTABLE)
+ * with a `__threadfence_system()` before thread exit so the host can read
+ * via `read_volatile` after a stream sync. Same pattern as
+ * `thompson_test_kernel.cu` and `sp4_histogram_p99_test_kernel.cu`.
+ *
+ * Test-only: no production callers. Cubin built by `crates/ml/build.rs`
+ * (kernel name `sp12_reward_math_test_kernel`).
+ */
+
+#include <cuda_runtime.h>
+#include "trade_physics.cuh"  /* compute_asymmetric_capped_pnl, compute_min_hold_penalty, compute_lump_sum_opp_cost */
+
+extern "C" __global__ void sp12_reward_math_test_kernel(
+    /* ── asymmetric cap inputs ── */
+    const float* __restrict__ in_base_reward,
+    float neg_cap,
+    float pos_cap,
+    float* __restrict__ out_capped_pnl,
+    /* ── min-hold penalty inputs ── */
+    const float* __restrict__ in_hold_time_mh,
+    float min_hold_target,
+    float min_hold_temperature,
+    float min_hold_penalty_max,
+    float* __restrict__ out_min_hold_penalty,
+    /* ── lump-sum opp_cost inputs ── */
+    float shaping_scale,
+    float holding_cost_rate,
+    const float* __restrict__ in_position,
+    const float* __restrict__ in_hold_time_oc,
+    float* __restrict__ out_opp_cost
+) {
+    /* Single-thread launch — the wrapper has no parallelism to exploit
+     * (each test invocation drives one (input, expected-output) pair). */
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+
+    out_capped_pnl[0] = compute_asymmetric_capped_pnl(
+        in_base_reward[0], neg_cap, pos_cap
+    );
+    out_min_hold_penalty[0] = compute_min_hold_penalty(
+        in_hold_time_mh[0], min_hold_target, min_hold_temperature, min_hold_penalty_max
+    );
+    out_opp_cost[0] = compute_lump_sum_opp_cost(
+        shaping_scale, holding_cost_rate, in_position[0], in_hold_time_oc[0]
+    );
+
+    /* Make all three writes PCIe-visible to host pointers (mapped-pinned
+     * buffers). The host reads via `read_volatile` after a stream sync;
+     * the fence guarantees the writes have actually drained from the GPU
+     * write buffers before the sync returns. */
+    __threadfence_system();
+}
diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh
index 098f69082..30022e05c 100644
--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh
+++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh
@@ -454,6 +454,86 @@ __device__ __forceinline__ void record_kelly_trade_outcome(
     }
 }
 
+/* ── SP12 v3 (2026-05-04): per-trade event-driven reward composition ──────
+ * Three small device functions extracted from `experience_kernels.cu` so the
+ * math is testable in isolation. Production `experience_env_step` calls these
+ * inside the `segment_complete` branch (around line ~2800). The Rust unit
+ * tests in `tests/sp12_reward_math_tests.rs` exercise them via a single-thread
+ * wrapper kernel (`sp12_reward_math_test_kernel.cu`) — synthetic inputs,
+ * analytical expected outputs, no CPU reference impl per
+ * `feedback_no_cpu_test_fallbacks`.
+ *
+ * Constants live in `state_layout.cuh` (REWARD_NEG_CAP, REWARD_POS_CAP,
+ * MIN_HOLD_TARGET, MIN_HOLD_PENALTY_MAX, MIN_HOLD_TEMPERATURE_*). The device
+ * functions take the bounds/parameters as arguments rather than reading the
+ * macros directly so the test wrapper can drive the same code path with
+ * arbitrary values without #define gymnastics.
+ * ──────────────────────────────────────────────────────────────────────── */
+
+/* SP12 v3 Change 1: asymmetric bounded reward cap (prospect-theory loss
+ * aversion). With production constants `neg_cap=-10, pos_cap=+5` the 2:1
+ * ratio matches the Kahneman/Tversky meta-analysis estimate of human loss
+ * aversion (~2.0-2.25). Restores selectivity that SP11's symmetric ±10 cap
+ * (commit 35db31089) erased while preserving the bilateral bound that
+ * fixed slot-63 PopArt EMA inflation. Per
+ * `pearl_audit_unboundedness_for_implicit_asymmetry`. */
+__device__ __forceinline__ float compute_asymmetric_capped_pnl(
+    float base_reward,
+    float neg_cap,   /* production: REWARD_NEG_CAP = -10.0f */
+    float pos_cap    /* production: REWARD_POS_CAP =  +5.0f */
+) {
+    return fmaxf(neg_cap, fminf(base_reward, pos_cap));
+}
+
+/* SP12 v3 Change 2: min-hold soft penalty for voluntary trade exits.
+ * Returns the penalty MAGNITUDE (positive). Caller subtracts it from the
+ * reward (production multiplies by `shaping_scale` at the call site).
+ *
+ * Soft saturation: `factor = deficit / (deficit + T)` — bounded [0, 1],
+ * smooth (no cliff), monotone increasing in deficit, monotone decreasing
+ * in T. When `hold_time >= min_hold_target` the deficit is non-positive
+ * and the function returns 0 — gating handled here so the call site
+ * stays a single expression.
+ *
+ * Numerical guard: production uses `fmaxf(deficit + T, 1e-6f)` to defend
+ * against simultaneous (deficit ≤ 0, T ≤ 0) malformed configs. The early
+ * `hold_time >= target` exit makes the guard unreachable for sensible
+ * inputs but the bound mirrors production exactly. */
+__device__ __forceinline__ float compute_min_hold_penalty(
+    float hold_time,
+    float min_hold_target,
+    float min_hold_temperature,
+    float min_hold_penalty_max
+) {
+    if (hold_time >= min_hold_target) return 0.0f;
+    const float deficit     = min_hold_target - hold_time;
+    const float soft_factor = deficit / fmaxf(deficit + min_hold_temperature, 1e-6f);
+    return min_hold_penalty_max * soft_factor;
+}
+
+/* SP12 v3 Change 3: lump-sum carrying-cost at trade close.
+ * Returns the (negative) reward magnitude. Production wires the result into
+ * `r_opp_cost` (rc[4]) only on `segment_complete` bars; per-bar callers
+ * MUST gate before invocation (the device function itself does NOT check
+ * exiting_trade — that gating is a call-site invariant of
+ * `experience_env_step`'s segment_complete branch).
+ *
+ * Per `pearl_event_driven_reward_density_alignment`: preserves the
+ * carrying-cost economic concept (k × |position| × hold_time) while moving
+ * its delivery from per-bar density-positive shaping to a single event-
+ * driven payment at trade close. */
+__device__ __forceinline__ float compute_lump_sum_opp_cost(
+    float shaping_scale,
+    float holding_cost_rate,
+    float position,
+    float hold_time
+) {
+    /* Parens preserve the production multiplication order
+     * (`-(scale * rate * |pos| * hold)`) so the refactor is bit-equivalent
+     * under f32 rounding to the inline computation it replaces. */
+    return -(shaping_scale * holding_cost_rate * fabsf(position) * hold_time);
+}
+
 /* ── Drawdown penalty: smooth ramp from threshold to floor ───────────── */
 /* Returns a penalty in [-5*w_dd, 0]. Applied every step so the model    */
 /* learns to reduce position size DURING drawdown.                       */
diff --git a/crates/ml/tests/sp12_reward_math_tests.rs b/crates/ml/tests/sp12_reward_math_tests.rs
new file mode 100644
index 000000000..c134fee55
--- /dev/null
+++ b/crates/ml/tests/sp12_reward_math_tests.rs
@@ -0,0 +1,443 @@
+#![allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory.
+
+//! SP12 v3 reward-math GPU oracle tests.
+//!
+//! Validates the three reward composition device functions added to
+//! `trade_physics.cuh` for SP12 v3 (commit 17cfbb250):
+//!   1. `compute_asymmetric_capped_pnl`  — prospect-theory asymmetric cap
+//!   2. `compute_min_hold_penalty`        — soft-saturation patience penalty
+//!   3. `compute_lump_sum_opp_cost`       — lump-sum carrying-cost on exit
+//!
+//! Each test drives a synthetic input through the production device
+//! function (via the standalone wrapper `sp12_reward_math_test_kernel.cu`)
+//! and verifies the GPU-computed result against an analytically-known
+//! expected value. Per `feedback_no_cpu_test_fallbacks`: GPU oracle only,
+//! no CPU reference impl. Per `feedback_no_htod_htoh_only_mapped_pinned`:
+//! every CPU↔GPU buffer is a `MappedF32Buffer` (cuMemHostAlloc with
+//! DEVICEMAP|PORTABLE); zero `htod_copy`, zero `dtoh_sync_copy`.
+//!
+//! All tests are `#[ignore = "requires GPU"]`-gated to match every other
+//! GPU oracle test in this crate (sp4/sp5/sp11). Run on a GPU host:
+//!
+//!     SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
+//!       cargo test -p ml --test sp12_reward_math_tests --features cuda \
+//!         -- --ignored --nocapture
+//!
+//! The wrapper kernel is single-thread / single-block; each test launch
+//! drives one (input → output) triple in parallel (the kernel writes all
+//! three outputs in one invocation), so the test helper returns a tuple
+//! `(capped_pnl, min_hold_penalty, opp_cost)` and individual assertions
+//! pick the relevant slot.
+
+#![cfg(feature = "cuda")]
+
+use std::sync::Arc;
+
+use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
+use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
+
+/// Test-only cubin built by `crates/ml/build.rs`.
+const SP12_REWARD_MATH_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/sp12_reward_math_test_kernel.cubin"));
+
+/// Resolve a CUDA stream against device 0. Mirrors `make_test_stream` in
+/// every other oracle test in this crate (`sp4_producer_unit_tests.rs`,
+/// `sp11_producer_unit_tests.rs`, `distributional_q_tests.rs`).
+fn make_test_stream() -> Arc<CudaStream> {
+    let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
+    ctx.default_stream()
+}
+
+/// Load the test kernel function from the embedded cubin.
+fn load_test_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
+    let module = stream
+        .context()
+        .load_cubin(SP12_REWARD_MATH_CUBIN.to_vec())
+        .expect("load sp12_reward_math_test_kernel cubin");
+    module
+        .load_function("sp12_reward_math_test_kernel")
+        .expect("load sp12_reward_math_test_kernel function")
+}
+
+/// Inputs for one launch of the SP12 reward-math wrapper kernel.
+///
+/// Carries the inputs each device function consumes. The wrapper kernel
+/// computes all three results in one invocation; tests project the
+/// returned tuple's relevant slot.
+#[derive(Clone, Copy, Debug)]
+struct Sp12RewardInputs {
+    /* asymmetric-cap inputs */
+    base_reward: f32,
+    neg_cap:     f32,
+    pos_cap:     f32,
+    /* min-hold-penalty inputs */
+    hold_time_mh:         f32,
+    min_hold_target:      f32,
+    min_hold_temperature: f32,
+    min_hold_penalty_max: f32,
+    /* lump-sum opp-cost inputs */
+    shaping_scale:     f32,
+    holding_cost_rate: f32,
+    position:          f32,
+    hold_time_oc:      f32,
+}
+
+impl Sp12RewardInputs {
+    /// Production constants (from `state_layout.cuh`) for the asymmetric cap.
+    /// `min_hold_*` and opp-cost params get neutral defaults so a test that
+    /// only cares about the asymmetric cap doesn't pay attention to the
+    /// other two outputs.
+    fn defaults() -> Self {
+        Self {
+            base_reward: 0.0,
+            neg_cap:    -10.0,
+            pos_cap:     5.0,
+            hold_time_mh:         30.0, // == target → penalty 0
+            min_hold_target:      30.0,
+            min_hold_temperature: 10.0,
+            min_hold_penalty_max:  3.0,
+            shaping_scale:     1.0,
+            holding_cost_rate: 0.001,
+            position:          0.0,    // → opp_cost 0
+            hold_time_oc:      0.0,
+        }
+    }
+}
+
+/// Returns `(capped_pnl, min_hold_penalty, opp_cost)` from one wrapper
+/// kernel launch. Single-thread single-block — the device functions are
+/// pure register arithmetic, no parallelism to exploit.
+fn launch_sp12_reward_math(
+    stream: &Arc<CudaStream>,
+    f: &CudaFunction,
+    inputs: Sp12RewardInputs,
+) -> (f32, f32, f32) {
+    // Safety: CUDA context active on this thread (resolved through the
+    // stream's context above). All CPU↔GPU buffers are mapped-pinned.
+    let in_base_reward = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc in_base_reward (1 f32)");
+    in_base_reward.write_from_slice(&[inputs.base_reward]);
+
+    let in_hold_time_mh = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc in_hold_time_mh (1 f32)");
+    in_hold_time_mh.write_from_slice(&[inputs.hold_time_mh]);
+
+    let in_position = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc in_position (1 f32)");
+    in_position.write_from_slice(&[inputs.position]);
+
+    let in_hold_time_oc = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc in_hold_time_oc (1 f32)");
+    in_hold_time_oc.write_from_slice(&[inputs.hold_time_oc]);
+
+    let out_capped_pnl = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc out_capped_pnl (1 f32)");
+    let out_min_hold_penalty = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc out_min_hold_penalty (1 f32)");
+    let out_opp_cost = unsafe { MappedF32Buffer::new(1) }
+        .expect("alloc out_opp_cost (1 f32)");
+
+    let in_base_reward_dev = in_base_reward.dev_ptr;
+    let in_hold_time_mh_dev = in_hold_time_mh.dev_ptr;
+    let in_position_dev = in_position.dev_ptr;
+    let in_hold_time_oc_dev = in_hold_time_oc.dev_ptr;
+    let out_capped_pnl_dev = out_capped_pnl.dev_ptr;
+    let out_min_hold_penalty_dev = out_min_hold_penalty.dev_ptr;
+    let out_opp_cost_dev = out_opp_cost.dev_ptr;
+
+    unsafe {
+        stream
+            .launch_builder(f)
+            // asymmetric cap
+            .arg(&in_base_reward_dev)
+            .arg(&inputs.neg_cap)
+            .arg(&inputs.pos_cap)
+            .arg(&out_capped_pnl_dev)
+            // min-hold penalty
+            .arg(&in_hold_time_mh_dev)
+            .arg(&inputs.min_hold_target)
+            .arg(&inputs.min_hold_temperature)
+            .arg(&inputs.min_hold_penalty_max)
+            .arg(&out_min_hold_penalty_dev)
+            // lump-sum opp_cost
+            .arg(&inputs.shaping_scale)
+            .arg(&inputs.holding_cost_rate)
+            .arg(&in_position_dev)
+            .arg(&in_hold_time_oc_dev)
+            .arg(&out_opp_cost_dev)
+            .launch(LaunchConfig {
+                grid_dim:         (1, 1, 1),
+                block_dim:        (1, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .expect("launch sp12_reward_math_test_kernel");
+    }
+    stream.synchronize().expect("sync after sp12_reward_math launch");
+
+    (
+        out_capped_pnl.read_all()[0],
+        out_min_hold_penalty.read_all()[0],
+        out_opp_cost.read_all()[0],
+    )
+}
+
+// ─── Asymmetric capped P&L (Change 1 — prospect-theory loss aversion) ──────
+
+/// Above the positive cap → clipped to `pos_cap`.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_asymmetric_cap_clips_above_pos_cap() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.base_reward = 20.0;
+    let (capped, _, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (capped - 5.0).abs() < 1e-5,
+        "+20 should clip to pos_cap=+5, got {capped}"
+    );
+}
+
+/// Below the negative cap → clipped to `neg_cap`.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_asymmetric_cap_clips_below_neg_cap() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.base_reward = -20.0;
+    let (capped, _, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (capped - (-10.0)).abs() < 1e-5,
+        "-20 should clip to neg_cap=-10, got {capped}"
+    );
+}
+
+/// Inside the band (zero) → identity passthrough.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_asymmetric_cap_passes_through_zero() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.base_reward = 0.0;
+    let (capped, _, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(capped.abs() < 1e-5, "0 should pass through, got {capped}");
+}
+
+/// Exactly at `pos_cap` → boundary identity.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_asymmetric_cap_at_pos_cap() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.base_reward = 5.0;
+    let (capped, _, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (capped - 5.0).abs() < 1e-5,
+        "+5 (== pos_cap) should pass through, got {capped}"
+    );
+}
+
+/// Exactly at `neg_cap` → boundary identity.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_asymmetric_cap_at_neg_cap() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.base_reward = -10.0;
+    let (capped, _, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (capped - (-10.0)).abs() < 1e-5,
+        "-10 (== neg_cap) should pass through, got {capped}"
+    );
+}
+
+// ─── Min-hold soft penalty (Change 2 — patience curriculum) ────────────────
+
+/// `hold_time == target` → deficit 0, soft factor 0, penalty 0.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_min_hold_penalty_at_target_is_zero() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.hold_time_mh         = 30.0;
+    inputs.min_hold_target      = 30.0;
+    inputs.min_hold_temperature = 10.0;
+    inputs.min_hold_penalty_max = 3.0;
+    let (_, penalty, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        penalty.abs() < 1e-5,
+        "hold_time == target should give 0 penalty, got {penalty}"
+    );
+}
+
+/// `hold_time = 0`, target = 30, T = 10 → deficit 30, factor 30/40 = 0.75,
+/// penalty = 3 × 0.75 = 2.25.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_min_hold_penalty_zero_hold_time_max_target() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.hold_time_mh         = 0.0;
+    inputs.min_hold_target      = 30.0;
+    inputs.min_hold_temperature = 10.0;
+    inputs.min_hold_penalty_max = 3.0;
+    let (_, penalty, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (penalty - 2.25).abs() < 1e-5,
+        "hold=0,target=30,T=10,max=3 should give penalty=2.25, got {penalty}"
+    );
+}
+
+/// `hold_time = 15`, target = 30, T = 10 → deficit 15, factor 15/25 = 0.6,
+/// penalty = 3 × 0.6 = 1.8.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_min_hold_penalty_half_way_deficit() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.hold_time_mh         = 15.0;
+    inputs.min_hold_target      = 30.0;
+    inputs.min_hold_temperature = 10.0;
+    inputs.min_hold_penalty_max = 3.0;
+    let (_, penalty, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (penalty - 1.8).abs() < 1e-5,
+        "hold=15,target=30,T=10,max=3 should give penalty=1.8, got {penalty}"
+    );
+}
+
+/// `hold_time > target` → no penalty (early-exit branch in device fn).
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_min_hold_penalty_past_target_no_penalty() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.hold_time_mh         = 50.0;
+    inputs.min_hold_target      = 30.0;
+    inputs.min_hold_temperature = 10.0;
+    inputs.min_hold_penalty_max = 3.0;
+    let (_, penalty, _) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        penalty.abs() < 1e-5,
+        "hold > target should give 0 penalty, got {penalty}"
+    );
+}
+
+/// Higher temperature smooths the curve: same deficit gives a SMALLER
+/// penalty under T=20 than under T=10. With deficit=30 (hold=0,target=30):
+///   T=10 → factor=30/40=0.75 → penalty=2.25
+///   T=20 → factor=30/50=0.60 → penalty=1.80
+/// Verifies `compute_min_hold_penalty` is monotone-decreasing in T at
+/// fixed deficit (the curriculum-anneal direction).
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_min_hold_penalty_temperature_smooths_transition() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+
+    let mut inputs_t10 = Sp12RewardInputs::defaults();
+    inputs_t10.hold_time_mh         = 0.0;
+    inputs_t10.min_hold_target      = 30.0;
+    inputs_t10.min_hold_temperature = 10.0;
+    inputs_t10.min_hold_penalty_max = 3.0;
+    let (_, penalty_t10, _) = launch_sp12_reward_math(&stream, &f, inputs_t10);
+
+    let mut inputs_t20 = Sp12RewardInputs::defaults();
+    inputs_t20.hold_time_mh         = 0.0;
+    inputs_t20.min_hold_target      = 30.0;
+    inputs_t20.min_hold_temperature = 20.0;
+    inputs_t20.min_hold_penalty_max = 3.0;
+    let (_, penalty_t20, _) = launch_sp12_reward_math(&stream, &f, inputs_t20);
+
+    assert!(
+        penalty_t10 > penalty_t20,
+        "higher T should give lower penalty for same deficit; T=10 → {penalty_t10}, T=20 → {penalty_t20}"
+    );
+    assert!(
+        (penalty_t20 - 1.8).abs() < 1e-5,
+        "hold=0,target=30,T=20,max=3 should give penalty=1.8, got {penalty_t20}"
+    );
+}
+
+// ─── Lump-sum opportunity cost (Change 3 — event-driven density) ───────────
+
+/// Standard exit: `scale=1, rate=0.001, position=0.5, hold_time=20` →
+/// `r_opp_cost = -0.001 × 0.5 × 20 = -0.01`.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_lump_sum_opp_cost_at_exit() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.shaping_scale     = 1.0;
+    inputs.holding_cost_rate = 0.001;
+    inputs.position          = 0.5;
+    inputs.hold_time_oc      = 20.0;
+    let (_, _, opp_cost) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (opp_cost - (-0.01)).abs() < 1e-6,
+        "expected -0.01 (= -1 * 0.001 * 0.5 * 20), got {opp_cost}"
+    );
+}
+
+/// Zero position → zero cost regardless of `hold_time`.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_lump_sum_opp_cost_zero_position() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.shaping_scale     = 1.0;
+    inputs.holding_cost_rate = 0.001;
+    inputs.position          = 0.0;
+    inputs.hold_time_oc      = 20.0;
+    let (_, _, opp_cost) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        opp_cost.abs() < 1e-6,
+        "zero position should give zero cost, got {opp_cost}"
+    );
+}
+
+/// Zero hold_time → zero cost regardless of `position`.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_lump_sum_opp_cost_zero_hold_time() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.shaping_scale     = 1.0;
+    inputs.holding_cost_rate = 0.001;
+    inputs.position          = 0.5;
+    inputs.hold_time_oc      = 0.0;
+    let (_, _, opp_cost) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        opp_cost.abs() < 1e-6,
+        "zero hold_time should give zero cost, got {opp_cost}"
+    );
+}
+
+/// Negative position uses `|position|`: same magnitude as positive case.
+#[test]
+#[ignore = "requires GPU"]
+fn sp12_lump_sum_opp_cost_negative_position_uses_abs() {
+    let stream = make_test_stream();
+    let f = load_test_kernel(&stream);
+    let mut inputs = Sp12RewardInputs::defaults();
+    inputs.shaping_scale     = 1.0;
+    inputs.holding_cost_rate = 0.001;
+    inputs.position          = -0.5;
+    inputs.hold_time_oc      = 20.0;
+    let (_, _, opp_cost) = launch_sp12_reward_math(&stream, &f, inputs);
+    assert!(
+        (opp_cost - (-0.01)).abs() < 1e-6,
+        "|-0.5| should match +0.5 magnitude (-0.01), got {opp_cost}"
+    );
+}
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index e625dd666..d97b79cb5 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -18,6 +18,8 @@ SP12 v3 — per-trade event-driven reward composition (2026-05-04): three archit
 
 **Backtest-kernel cap site investigation**: spec §design-1 mentioned `crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu` as the matching cap location. Investigation confirmed `backtest_plan_kernel.cu` lines 167/171 contain `fmaxf(-2.0f, fminf(unrealized / (plan_profit × equity + 1e-6f), 2.0f))` for `pisv[PLAN_ISV_PNL_VS_TARGET]` and `pisv[PLAN_ISV_PNL_VS_STOP]` — these are STATE FEATURE normalizations (the policy reads them as input features), not reward caps. Same pattern mirrored in `experience_kernels.cu` lines 852-856. The state-feature representations need to remain symmetric so the policy's perception of "way over target" vs "way under stop" is balanced. Asymmetrizing those would be a different change with different rationale; not in scope for SP12. The reward cap is exclusively at `experience_kernels.cu:2758`, the only site touched by Change 1.
 
+SP12 v3 follow-up — GPU oracle test scaffold for the 3 reward math changes (2026-05-04): the architectural commit `17cfbb250` deferred the GPU oracle test scaffold; this follow-up lands it without behavior change. Three reward composition formulas refactored from inline expressions in `experience_kernels.cu` into device-inline functions in `trade_physics.cuh`: `compute_asymmetric_capped_pnl(base_reward, neg_cap, pos_cap)`, `compute_min_hold_penalty(hold_time, target, T, max)`, `compute_lump_sum_opp_cost(scale, rate, position, hold_time)`. The `experience_env_step` `segment_complete` branch now calls these helpers in place of the inline math. Bit-exact preservation: the lump-sum opp_cost helper wraps the multiplication in parens to match the production `-(scale * rate * |pos| * hold)` order under f32 rounding; the min-hold helper subsumes the production `if (hold_time < target)` early-exit (returns 0 when at/past target — `capped_pnl - shaping_scale * 0 == capped_pnl`). New test kernel `crates/ml/src/cuda_pipeline/sp12_reward_math_test_kernel.cu` exposes the three device functions for GPU oracle testing (single-thread / single-block; outputs to mapped-pinned buffers with `__threadfence_system()` per the standard pattern in `thompson_test_kernel.cu` and `sp4_histogram_p99_test_kernel.cu`); cubin registered in `build.rs::kernels_with_common`. New test module `crates/ml/tests/sp12_reward_math_tests.rs` covers all three formulas with 14 GPU oracle tests: 5 for asymmetric cap (above/below caps, zero, exact boundaries), 5 for min-hold soft penalty (at target, hold=0, mid-deficit, past target, temperature-monotonicity), 4 for lump-sum opp_cost (standard exit, zero position, zero hold_time, |position|). Per `feedback_no_cpu_test_fallbacks` (GPU oracle only) + `feedback_no_htod_htoh_only_mapped_pinned` (every CPU↔GPU buffer is `MappedF32Buffer`); `#[ignore = "requires GPU"]`-gated to match the existing sp4/sp5/sp11 producer-test convention. Verified on local RTX 3050 Ti (sm_86): 14 passed, 0 failed (3.32s). Touched: `crates/ml/build.rs` (+1 cubin entry), `crates/ml/src/cuda_pipeline/trade_physics.cuh` (+3 `__device__ __forceinline__` functions), `crates/ml/src/cuda_pipeline/experience_kernels.cu` (3 inline math sites refactored to device-fn calls; comments cross-reference the test file), `crates/ml/src/cuda_pipeline/sp12_reward_math_test_kernel.cu` (new), `crates/ml/tests/sp12_reward_math_tests.rs` (new).
+
 SP6 clean-compile — wire-or-delete all 13 ml + 2 ml-dqn warnings (2026-05-01): W1 SEMANTIC: Pearl 2's `iqn_branch[b]` (per-branch IQN budget, `[f32; 4]` returned by `compute_adaptive_budgets`) was destructured but never consumed — both parallel and sequential Pearl 5 IQN paths used `iqn_trunk / 4.0` (the mean), silently averaging out Pearl 2's per-branch differentiation. Fixed in both paths by moving `let iqn_budget_per_branch = iqn_branch[branch_idx] / 4.0_f32` inside the `for branch_idx in 0..4` loop so each pass uses its branch-specific budget. `iqn_trunk` is now only referenced in comments → renamed to `_iqn_trunk` in the destructure. W2 IMPORT DELETE: `ATOM_NUM_ATOMS_BASE` and `NOISY_SIGMA_BASE` removed from `fused_training.rs:44` — consumers are in Layer B's `atoms_update_kernel.cu` and `experience_kernels.cu`, not `fused_training.rs`; imports were speculative additions from an earlier draft. W3 DELETE: `v_blocks` at `gpu_iql_trainer.rs:756` — computed `((b+255)/256)` but both downstream kernel launches used `v_blocks2` (the 2× variant); stale dead code. W4 IMPORT DELETE: `OrderRouter` removed from `ml-dqn/src/dqn.rs:25` — import only, never used in the file body. W5 IMPORT DELETE: `get_snapshots_for_timestamp` removed from `data_loading.rs:312` — imported inside an `if !mbp10_data_dir.is_empty()` block but the function was never called; `get_trades_for_bar` and `OFICalculator` ARE used. W6 DELETE: `update_target_networks` method in `ml-dqn/src/dqn.rs:1621` — 42-line orphan with docstring claiming it's "shared implementation used by train_step and apply_accumulated_gradients" but `grep` showed zero call sites; real training path uses the fused CUDA EMA kernel on flat buffers. Cascaded: `convergence_half_life` import (now unused) deleted from line 14. W7-W12 FILE-LEVEL `#![allow(unsafe_code)]`: `cublas_algo_deterministic.rs` and `sp4_wiener_ema.rs` — workspace `Cargo.toml` has `unsafe_code = "warn"` which fires for every `unsafe impl / unsafe fn`. Both files require unsafe for cuBLAS-Lt handle passing and CUDA kernel launches; added file-level attr matching the established pattern in `mapped_pinned.rs`, `gpu_iqn_head.rs`, `gpu_weights.rs`, etc. W13 same treatment in `sp4_wiener_ema.rs` (covered by the same file-level attr). W14-W15 VISIBILITY SCOPE: `ControllerPrevValues` and `ControllerFireCounts` in `trainer/mod.rs:172,197` changed from `pub struct` to `pub(crate) struct` — both are consumed only within the `ml` crate; `pub` was unreachable from outside. Per `feedback_no_hiding`: zero `#[allow(...)]` suppressions added at item level; 2 file-level `#![allow(unsafe_code)]` follow established crate pattern. cargo check (ml + ml-dqn) + cargo build --release + cargo test --lib sp5/sp4/state_reset_registry all clean (0 warnings; 13/13 pass).
 
 SP5 Layer A bug-fix — add reset_named_state dispatch arms for 21 SP5 entries (2026-05-02): L40S smoke at SP5 Layer B HEAD (commit `3ad5e011b`) failed in 9.4s at the first fold-reset boundary with `StateResetRegistry reset dispatch: unknown name 'sp5_atom_v_center'`. Every Layer A task added `RegistryEntry` blocks to `state_reset_registry.rs` but **none added the matching dispatch arms in `reset_named_state`** — the runtime validator catches this on first fold boundary, crashing all 3 folds before training runs (0/3 checkpoints saved). Local unit tests verified slot layouts but never exercised the fold-reset dispatch path, so the gap was masked through Layer A close-out. Added 21 dispatch arms covering all SP5 Layer A FoldReset entries: `sp5_atom_v_center/v_half/headroom/clip_rate` (Pearl 1), `sp5_branch_entropy/q_var_per_branch` (shared signal), `sp5_wiener_state` (no-op — `reset_sp4_wiener_state` already bulk-zeros the full 543-float buffer including the SP5 block at offsets [213..543)), `sp5_noisy_sigma/sigma_fraction` (Pearl 3), `sp5_budget_c51/iqn/cql/ens/flatness` (Pearl 2), `sp5_adam_beta1/beta2/eps` (Pearl 4), `sp5_grad_prev_buf` (Pearl 4 aux — calls a new `reset_sp5_grad_prev_buf` bulk-zero method on `GpuDqnTrainer`, mirrors `reset_sp4_clamp_engage_counters` pattern), `sp5_iqn_tau` (Pearl 5), `sp5_trail_dist_per_dir` (Pearl 8), `sp5_atom_num_atoms` (Pearl 1-ext). **Pearl 6 (slots 280..286) intentionally absent** — those slots are cross-fold-persistent (the whole point of A6 per `project_magnitude_eval_collapse_kelly_capped`). Each ISV-range arm zeros its slot range to fire Pearl A's first-observation sentinel on the new fold's first producer launch. Touched: `cuda_pipeline/gpu_dqn_trainer.rs` (+1 method `reset_sp5_grad_prev_buf`), `trainers/dqn/trainer/training_loop.rs` (+21 dispatch arms before the `_ =>` fallback). cargo check + cargo test --lib (sp4+sp5+state_reset_registry: 13/13 pass) both clean. Smoke re-deployment pending.