feat(sp18 v2 P2.T1-T5): structural Hold opportunity-cost device fn + 3-site migration (INTERIM STATE LIFTED)

Atomic introduction of `compute_sp18_hold_opportunity_cost` device function
in `trade_physics.cuh` + migration of the 3 placeholder-commented per-bar
Hold sites in `experience_kernels.cu` + CPU/GPU oracle tests + test-wrapper
cubin manifest entry — single atomic commit per `feedback_no_partial_refactor`.

Closes the SP18 INTERIM STATE introduced by Phase 1 (3c318953a). The reward
function is now structurally complete: no placeholder sites, no orphaned
consumers. Wire-up audit fingerprint stable; production-call count for the
new device fn = 3 (segment_complete + per-bar positioned-Hold + per-bar
flat-Hold), test-wrapper count = 1, definition count = 1.

Sites:
- trade_physics.cuh: new device fn after `compute_lump_sum_opp_cost`.
  Returns 0 for `dir_idx != DIR_HOLD`. Sign convention DD5(b) MIRRORED:
  positive next_log_return → negative Hold reward (opp cost paid for
  sitting out a winner); negative next_log_return → positive Hold reward
  (correctly avoided a loser). Asymmetric-capped via the existing
  `compute_asymmetric_capped_pnl` primitive against ISV[483]/[484]; output
  multiplied by `shaping_scale` (validation = 0 → returns 0). No new
  infrastructure beyond existing primitives.
- experience_kernels.cu Site 1 (segment_complete ~3143): unconditional
  add into `base_reward` BEFORE the SP12 asymmetric cap; preserves SP12
  cap interaction.
- experience_kernels.cu Site 2 (per-bar positioned-Hold ~3641):
  unconditional add into `r_micro` (rc[3]) — same component the deleted
  SP13/SP16 subtraction wrote to, preserving SP11 reward-subsystem
  controller signal attribution. `vol_proxy` recomputed locally from
  `features[bar_idx * market_dim + 9]` (per-bar branches are outside
  the segment_complete `vol_proxy` scope).
- experience_kernels.cu Site 3 (per-bar flat-Hold ~3760): unconditional
  add into `r_opp_cost` (rc[4]) — same component the deleted subtraction
  wrote to. `vol_proxy` recomputed locally with the same pattern.

Each site reads ISV[HOLD_REWARD_POS_CAP_INDEX=483] /
ISV[HOLD_REWARD_NEG_CAP_INDEX=484] with sentinel-or-out-of-bounds guard
falling through to SENTINEL_HOLD_REWARD_POS/NEG_CAP=+5/-10 macros. Bit-
identical to a fixed +5/-10 cap during cold-start and Pre-Phase-3 (Phase 3
lands the producer kernel `hold_reward_cap_update_kernel`).

Tests:
- T1 CPU oracle (sp18_hold_reward_oracle_tests.rs): 6 cases covering
  positive/negative log-return, dir-gate (early-exit for non-HOLD),
  positive saturation (clamps at +pos_cap), negative saturation (clamps
  at -neg_cap), shaping_scale=0 validation mode (returns 0).
- T2 GPU oracle (sp18_hold_reward_oracle_tests.rs::gpu): same 6 cases via
  `sp18_hold_opp_test_kernel.cu` wrapper; bit-equality to CPU at ε=1e-5.
- sp18_hold_opp_test_kernel.cu: single-thread single-block test wrapper
  mirroring `sp12_reward_math_test_kernel.cu` structure. Mapped-pinned
  f32 input/output per `feedback_no_htod_htoh_only_mapped_pinned`;
  __threadfence_system() for PCIe visibility before host read.
- build.rs: cubin manifest entry for sp18_hold_opp_test_kernel.cu.

Validation:
- SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace clean.
- cargo test -p ml --lib sp18: 4 SP18 lib tests pass (slot layouts,
  fold-reset entries-present).
- cargo test -p ml --lib state_reset_registry: 5 tests pass (including
  every_fold_and_soft_reset_entry_has_dispatch_arm).
- cargo test -p ml --test sp18_hold_reward_oracle_tests --features cuda
  -- --include-ignored: 7 tests pass (2 CPU + 5 GPU on local RTX 3050 Ti),
  including the new compute_sp18_hold_opportunity_cost_cpu_oracle and
  compute_sp18_hold_opportunity_cost_gpu_oracle.
- bash scripts/audit_sp18_consumers.sh --check: exit 0 (post-Phase-2
  fingerprint snapshotted in docs/sp18-wireup-audit.md).
- Wire-up audit: `grep -rn "compute_sp18_hold_opportunity_cost" crates/ml/`
  shows 1 device-fn definition + 3 production call sites + 1 test-wrapper
  kernel + CPU/GPU oracle tests. No orphans.

Audit fingerprint diff (Phase 1 → Phase 2):
- `Slot 380 (HOLD_COST_INDEX) consumers`: experience_kernels.cu hit count
  1 → 0 (placeholder-comment reference to HOLD_COST_INDEX removed; new
  Phase 2 calls reference HOLD_REWARD_POS/NEG_CAP_INDEX=483/484 instead);
  doc references in dqn-wire-up-audit.md grew by 2 lines (Phase 2
  close-out narrative). Total section count 42 → 43.
- All other sections unchanged — Phase 2 introduces forward-tracking
  references on the SP18-bound side (slots 483/484, new device fn), not
  the deleted-chain side.

Discipline:
- feedback_no_partial_refactor: device fn + 3 call sites + tests + cubin
  manifest in same commit; INTERIM STATE LIFTED in same commit.
- feedback_no_atomicadd: device fn is per-thread sample-local register
  arithmetic; test wrapper is single-thread single-block, no shared mem.
- feedback_no_htod_htoh_only_mapped_pinned: GPU oracle test uses
  MappedF32Buffer.
- feedback_isv_for_adaptive_bounds: pos/neg caps read from ISV[483]/[484]
  with sentinel fallback; Phase 3 producer `hold_reward_cap_update_kernel`
  drives them adaptively.
- feedback_wire_everything_up: device fn + 3 call sites + test wrapper +
  CPU/GPU oracle tests + cubin manifest all in same commit.
- pearl_one_unbounded_signal_per_reward: only `vol_normalized_return` is
  unbounded; everything else in [neg_cap, pos_cap] × shaping_scale.
- pearl_first_observation_bootstrap: sentinels (5.0/-10.0) match SP14
  P0-A REWARD_POS/NEG_CAP_ADAPTIVE for bit-identical cold-start.

Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md
Phase 2 Tasks 2.1-2.5.

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

crates/ml/build.rs

@@ -896,6 +896,20 @@ fn main() {
         // Pure observability: pre-fix baseline for the B-DD9 ratio
         // gate. Plan: SP18 Phase 0 Task 0.2.
         "td_error_mag_ema_kernel.cu",
+        // SP18 v2 Phase 2 (2026-05-09): D-leg structural Hold
+        // opportunity-cost device-fn GPU oracle test wrapper. Single
+        // thread × single block; wraps `compute_sp18_hold_opportunity_cost`
+        // from `trade_physics.cuh`. Mapped-pinned f32 input/output (no
+        // htod/dtoh per `feedback_no_htod_htoh_only_mapped_pinned`); host
+        // reads via `read_volatile` after stream sync (the
+        // `__threadfence_system()` at thread exit guarantees PCIe
+        // visibility before the sync returns). Same structural pattern
+        // as `sp12_reward_math_test_kernel.cu`. Test-only; no production
+        // callers. Drives `compute_sp18_hold_opportunity_cost`'s closed-
+        // form math contract from the Rust GPU oracle in
+        // `tests/sp18_hold_reward_oracle_tests.rs` against the CPU
+        // oracle in the same file (ε=1e-5). Plan: SP18 Phase 2 Task 2.2.
+        "sp18_hold_opp_test_kernel.cu",
         // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
         // Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
         // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -3140,19 +3140,54 @@ extern "C" __global__ void experience_env_step(
          * `exiting_trade=1` upstream (line ~2371), so this branch fires
          * for both — the segregation matters only for component attribution. */
         float base_reward = 2.0f * vol_normalized_return;
-        /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction at the
+        /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost
-         * voluntary segment_complete branch. Goes BEFORE the SP12 asymmetric
+         * counterfactual replaces the SP13 P0a / SP16 P2 / SP16 T3 reactive
-         * cap so the cost participates in the bounded clamp — keeps total
+         * per-bar Hold-cost-scale chain (deleted Phase 1 c66e15f6d /
-         * reward in the same [-10, +5] range and avoids a long Hold run
+         * 3c318953a). Goes BEFORE the SP12 asymmetric cap so the
-         * silently saturating the lower cap. Hold action rarely coincides
+         * counterfactual participates in the bounded clamp — keeps total
-         * with segment_complete (Hold keeps current position, doesn't close
+         * reward in the [-10, +5] range and avoids a long Hold run silently
-         * trades) but trail-fire can force-exit while the policy picks Hold;
+         * saturating the lower cap. Hold action rarely coincides with
-         * this branch covers that edge. Per-bar (non-segment_complete)
+         * segment_complete (Hold keeps current position, doesn't close
-         * Hold-cost subtraction is in the per-bar branches below.
+         * trades) but trail-fire can force-exit while the policy picks
+         * Hold; this branch covers that edge. Per-bar (non-segment_complete)
+         * Hold sites are in the two branches below.
          *
-         * SP18 D-leg replacement: per-bar Hold cost subtraction removed.
+         * The device function returns 0 for non-Hold so unconditional
-         * Replaced with compute_sp18_hold_opportunity_cost in Phase 2.
+         * addition is safe. Caps come from ISV[HOLD_REWARD_NEG_CAP_INDEX=484]
-         * (No reward write here in interim; Phase 2 lands the new device fn.) */
+         * / ISV[HOLD_REWARD_POS_CAP_INDEX=483] with sentinel/window-guard
+         * fallback to the macro defaults — bit-identical to a fixed +5/-10
+         * cap during cold-start (Phase 3 lands the adaptive producer).
+         * `shaping_scale` matches the SP12 contract (validation = 0 → no
+         * Hold opp cost in pure backtest). Per
+         * `pearl_one_unbounded_signal_per_reward.md` the only unbounded
+         * multiplicand is `next_log_return / sqrt(vol_proxy)`; everything
+         * else is in [floor, ceil]. */
+        {
+            float sp18_pos_cap = SENTINEL_HOLD_REWARD_POS_CAP;  /* +5.0 */
+            float sp18_neg_cap = SENTINEL_HOLD_REWARD_NEG_CAP;  /* -10.0 */
+            if (isv_signals_ptr != NULL) {
+                const float isv_pos =
+                    isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX];
+                const float isv_neg =
+                    isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX];
+                if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f
+                    && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND
+                    && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) {
+                    sp18_pos_cap = isv_pos;
+                    sp18_neg_cap = isv_neg;
+                }
+            }
+            const float sp18_hold_opp = compute_sp18_hold_opportunity_cost(
+                dir_idx,
+                /*dir_hold=*/DIR_HOLD,
+                /*next_log_return=*/tgt[1],
+                /*vol_proxy=*/vol_proxy,
+                /*isv_neg_cap=*/sp18_neg_cap,
+                /*isv_pos_cap=*/sp18_pos_cap,
+                /*shaping_scale=*/shaping_scale
+            );
+            base_reward += sp18_hold_opp;
+        }
         /* SP12 v3 fix (2026-05-04): asymmetric bounded cap (loss aversion).
          * SP11 (commit 35db31089) made the cap symmetric ±10 to fix slot-63
          * PopArt EMA inflation. That fix preserved stability but erased the
@@ -3603,21 +3638,73 @@ extern "C" __global__ void experience_env_step(
          * (which is a separate concern; spec carve-out: "ONLY micro_reward
          * and opp_cost are the per-bar shaping anti-patterns"). */
         r_micro = 0.0f;
-        /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction.
+        /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost
-         * Hold action keeps the current position with zero PnL signal —
+         * counterfactual on positioned-non-event bars. Replaces the SP13
-         * pre-SP13 this was a free no-op the policy could pick indefinitely
+         * P0a / SP16 P2 / SP16 T3 reactive per-bar Hold-cost-scale chain
-         * (CQL bias anchors toward minimum-variance Hold ~45% of bars across
+         * (deleted Phase 1 c66e15f6d / 3c318953a). Routes through the
-         * SP1-SP12). Hold-pricing makes Hold uneconomic when overused: the
+         * micro_reward component (rc[3]) — the same component the deleted
-         * controller in `training_loop.rs` raises ISV[HOLD_COST_INDEX] when
+         * subtraction wrote to via `r_micro -= cost` pre-SP18, preserving
-         * observed Hold-rate > target (default 0.20). The cost is bounded
+         * SP11 reward-subsystem controller signal attribution.
-         * structurally (HOLD_COST_BASE × CEIL_RATIO = 0.005 max per bar),
-         * well below the SP12 asymmetric cap [-10, +5] — no per-bar clamp
-         * required. Spec §"Change 1: Price Hold (replaces v2's Eliminate
-         * Hold)".
          *
-         * SP18 D-leg replacement: per-bar Hold cost subtraction removed.
+         * The device function returns 0 for `dir_idx != DIR_HOLD`, so
-         * Replaced with compute_sp18_hold_opportunity_cost in Phase 2.
+         * unconditional addition is safe. `vol_proxy` and `next_log_return`
-         * (No reward write here in interim; Phase 2 lands the new device fn.) */
+         * are recomputed locally from `features[bar_idx * market_dim + 9]`
+         * + `tgt[1]` — `vol_proxy` is scoped to the segment_complete branch
+         * above, so we mirror the segment_complete formula exactly here.
+         * The `features != NULL && bar_idx < total_bars && market_dim > 9`
+         * guard matches the B.2/D.4c bonus blocks below; on cold-start /
+         * test-fixture inputs without market features we skip cleanly
+         * (sp18_hold_opp stays at 0.0).
+         *
+         * Caps come from ISV[HOLD_REWARD_NEG_CAP_INDEX=484] /
+         * ISV[HOLD_REWARD_POS_CAP_INDEX=483] with sentinel/window-guard
+         * fallback to the macro defaults — bit-identical to a fixed +5/-10
+         * cap during cold-start (Phase 3 lands the adaptive producer).
+         * `shaping_scale` matches the SP12 contract (validation = 0 → no
+         * Hold opp cost in pure backtest). Per
+         * `pearl_one_unbounded_signal_per_reward.md` the only unbounded
+         * multiplicand is `next_log_return / sqrt(vol_proxy)`; everything
+         * else is in [floor, ceil].
+         *
+         * The structurally-bounded output (capped in [neg_cap, pos_cap])
+         * sits comfortably inside the SP12 asymmetric reward cap [-10, +5]
+         * the segment_complete branch enforces — but this branch fires
+         * BETWEEN trade close events (positioned-non-event), so no
+         * downstream cap is applied here; the device function's
+         * `compute_asymmetric_capped_pnl` is the structural bound. */
+        if (features != NULL && bar_idx < total_bars && market_dim > 9) {
+            const float atr_norm_pb1 =
+                features[(long long)bar_idx * market_dim + 9];
+            const float log_atr_pb1  = atr_norm_pb1 * 16.0f - 7.0f;
+            const float atr_pct_pb1  =
+                expf(log_atr_pb1) / fmaxf(raw_close, 1.0f);
+            const float vol_proxy_pb1 = fmaxf(atr_pct_pb1, 0.0001f);
+
+            float sp18_pos_cap_pb1 = SENTINEL_HOLD_REWARD_POS_CAP;  /* +5.0 */
+            float sp18_neg_cap_pb1 = SENTINEL_HOLD_REWARD_NEG_CAP;  /* -10.0 */
+            if (isv_signals_ptr != NULL) {
+                const float isv_pos =
+                    isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX];
+                const float isv_neg =
+                    isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX];
+                if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f
+                    && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND
+                    && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) {
+                    sp18_pos_cap_pb1 = isv_pos;
+                    sp18_neg_cap_pb1 = isv_neg;
+                }
+            }
+            const float sp18_hold_opp_pb1 = compute_sp18_hold_opportunity_cost(
+                dir_idx,
+                /*dir_hold=*/DIR_HOLD,
+                /*next_log_return=*/tgt[1],
+                /*vol_proxy=*/vol_proxy_pb1,
+                /*isv_neg_cap=*/sp18_neg_cap_pb1,
+                /*isv_pos_cap=*/sp18_pos_cap_pb1,
+                /*shaping_scale=*/shaping_scale
+            );
+            r_micro += sp18_hold_opp_pb1;
+        }
         reward  = r_micro;
         micro_reward_per_sample[out_off] = r_micro;
         reward_components_per_sample[out_off * 6 + 3] = r_micro;
@@ -3672,17 +3759,63 @@ extern "C" __global__ void experience_env_step(
          * preserved: only the Flat opp_cost EMITTER is zeroed; the
          * conviction Welford observer is independent. */
         r_opp_cost = 0.0f;
-        /* SP13 v3 P0a.T3 (2026-05-04): per-bar Hold cost subtraction on
+        /* SP18 D-leg Phase 2 (2026-05-09): structural Hold opportunity-cost
-         * the flat-non-event branch. Mirrors the positioned-branch
+         * counterfactual on flat-non-event bars. Mirrors the positioned-
-         * subtraction above so a policy picking Hold while flat (i.e.
+         * branch wiring above so a policy picking Hold while flat ("stay
-         * "stay flat") pays the same per-bar cost as a policy picking
+         * flat") sees the same structural counterfactual as a policy
-         * Hold while positioned (i.e. "stay long/short"). The cost is
+         * picking Hold while positioned ("stay long/short"). Routes
-         * structurally bounded by HOLD_COST_BASE × CEIL_RATIO = 0.005
+         * through the opp_cost component (rc[4]) — the same component
-         * per bar, well below the SP12 asymmetric cap [-10, +5].
+         * the deleted SP13/SP16 chain wrote to via `r_opp_cost -= cost`
+         * pre-SP18, preserving SP11 reward-subsystem controller signal
+         * attribution.
          *
-         * SP18 D-leg replacement: per-bar Hold cost subtraction removed.
+         * The device function returns 0 for `dir_idx != DIR_HOLD`, so
-         * Replaced with compute_sp18_hold_opportunity_cost in Phase 2.
+         * unconditional addition is safe. `vol_proxy` and `next_log_return`
-         * (No reward write here in interim; Phase 2 lands the new device fn.) */
+         * recomputed locally — same pattern as the positioned-branch and
+         * the B.2/D.4c bonus blocks. Caps come from ISV[483]/[484] with
+         * sentinel/window-guard fallback to the macro defaults; bit-
+         * identical to a fixed +5/-10 cap during cold-start (Phase 3
+         * lands the adaptive producer). `shaping_scale=0` (validation
+         * mode) → device function returns 0 → no Hold opp cost in pure
+         * backtest, matching the SP12 lump-sum contract.
+         *
+         * Output bounded by `compute_asymmetric_capped_pnl` to
+         * [neg_cap, pos_cap] before scaling by `shaping_scale`; sits
+         * inside the SP12 asymmetric reward cap range. Per
+         * `pearl_one_unbounded_signal_per_reward.md`. */
+        if (features != NULL && bar_idx < total_bars && market_dim > 9) {
+            const float atr_norm_pb2 =
+                features[(long long)bar_idx * market_dim + 9];
+            const float log_atr_pb2  = atr_norm_pb2 * 16.0f - 7.0f;
+            const float atr_pct_pb2  =
+                expf(log_atr_pb2) / fmaxf(raw_close, 1.0f);
+            const float vol_proxy_pb2 = fmaxf(atr_pct_pb2, 0.0001f);
+
+            float sp18_pos_cap_pb2 = SENTINEL_HOLD_REWARD_POS_CAP;  /* +5.0 */
+            float sp18_neg_cap_pb2 = SENTINEL_HOLD_REWARD_NEG_CAP;  /* -10.0 */
+            if (isv_signals_ptr != NULL) {
+                const float isv_pos =
+                    isv_signals_ptr[HOLD_REWARD_POS_CAP_INDEX];
+                const float isv_neg =
+                    isv_signals_ptr[HOLD_REWARD_NEG_CAP_INDEX];
+                if (fabsf(isv_pos - SENTINEL_HOLD_REWARD_POS_CAP) >= 1e-6f
+                    && isv_pos >= HOLD_REWARD_POS_CAP_MIN_BOUND
+                    && isv_pos <= HOLD_REWARD_POS_CAP_MAX_BOUND) {
+                    sp18_pos_cap_pb2 = isv_pos;
+                    sp18_neg_cap_pb2 = isv_neg;
+                }
+            }
+            const float sp18_hold_opp_pb2 = compute_sp18_hold_opportunity_cost(
+                dir_idx,
+                /*dir_hold=*/DIR_HOLD,
+                /*next_log_return=*/tgt[1],
+                /*vol_proxy=*/vol_proxy_pb2,
+                /*isv_neg_cap=*/sp18_neg_cap_pb2,
+                /*isv_pos_cap=*/sp18_pos_cap_pb2,
+                /*shaping_scale=*/shaping_scale
+            );
+            r_opp_cost += sp18_hold_opp_pb2;
+        }
         reward = r_opp_cost;
         reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
     }

crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu

@@ -0,0 +1,62 @@
+/* SP18 v2 D-leg Phase 2 (2026-05-09): GPU oracle test kernel for
+ * `compute_sp18_hold_opportunity_cost`.
+ *
+ * Standalone test wrapper that exercises the structural Hold opportunity-
+ * cost device function added to `trade_physics.cuh` for SP18 v2 Phase 2.
+ * Mirrors the SP12 `sp12_reward_math_test_kernel.cu` pattern: single-thread
+ * single-block launch, mapped-pinned f32 input/output buffers, host reads
+ * via `read_volatile` after a stream sync (the `__threadfence_system()` at
+ * exit guarantees the device write is PCIe-visible).
+ *
+ * Per `feedback_no_cpu_test_fallbacks.md`: GPU oracle only (no CPU
+ * reference impl in production). The Rust CPU oracle test in
+ * `tests/sp18_hold_reward_oracle_tests.rs` locks the closed-form math
+ * contract; this wrapper drives the same inputs through the production
+ * device function and the GPU oracle test asserts bit-equality to the
+ * closed-form expected values within an f32 rounding budget (ε=1e-5).
+ *
+ * Per `feedback_no_htod_htoh_only_mapped_pinned.md`: every CPU↔GPU
+ * buffer is `MappedF32Buffer` (cuMemHostAlloc DEVICEMAP|PORTABLE); zero
+ * `htod_copy`, zero `dtoh_sync_copy` in the test. Per
+ * `feedback_no_atomicadd.md`: single-thread launch, no shared memory, no
+ * reduction — pure register arithmetic.
+ *
+ * Test-only: no production callers. Cubin built by `crates/ml/build.rs`
+ * (kernel name `sp18_hold_opp_test_kernel`).
+ */
+
+#include <cuda_runtime.h>
+#include "trade_physics.cuh"  /* compute_sp18_hold_opportunity_cost */
+
+extern "C" __global__ void sp18_hold_opp_test_kernel(
+    int   dir_idx,
+    int   dir_hold,
+    float next_log_return,
+    float vol_proxy,
+    float isv_neg_cap,
+    float isv_pos_cap,
+    float shaping_scale,
+    float* __restrict__ out_hold_opp_cost
+) {
+    /* Single-thread launch — the device function has no parallelism to
+     * exploit (pure register arithmetic; one (input, output) pair per
+     * test invocation). */
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+
+    out_hold_opp_cost[0] = compute_sp18_hold_opportunity_cost(
+        dir_idx,
+        dir_hold,
+        next_log_return,
+        vol_proxy,
+        isv_neg_cap,
+        isv_pos_cap,
+        shaping_scale
+    );
+
+    /* Make the write PCIe-visible to the host pointer (mapped-pinned
+     * buffer). The host reads via `read_volatile` after a stream sync; the
+     * fence guarantees the write has actually drained from the GPU write
+     * buffers before the sync returns. Same pattern as
+     * `sp12_reward_math_test_kernel.cu`. */
+    __threadfence_system();
+}

crates/ml/src/cuda_pipeline/trade_physics.cuh

@@ -599,6 +599,88 @@ __device__ __forceinline__ float compute_lump_sum_opp_cost(
     return -(shaping_scale * holding_cost_rate * fabsf(position) * hold_time);
 }
 
+/* SP18 D-leg (2026-05-09): structural Hold opportunity-cost reward.
+ *
+ * Replaces the SP13 P0a + SP16 P2 + SP16 T3 reactive per-bar Hold-cost-scale
+ * subtraction chain (deleted in SP18 Phase 1) with an event-driven structural
+ * counterfactual: when the policy picks Hold, what would Long have earned on
+ * the next bar? Sign convention DD5(b) MIRRORED:
+ *
+ *   - positive next_log_return (Long would have won)  → NEGATIVE Hold reward
+ *     (opportunity cost paid for sitting out).
+ *   - negative next_log_return (Long would have lost) → POSITIVE Hold reward
+ *     (correctly avoided a losing trade).
+ *
+ * The base counterfactual is `-base = -(2.0 × vol_normalized_log_return)`
+ * — the same `2.0 × vol_normalized_return` scalar the segment_complete branch
+ * uses for `base_reward`, sign-flipped so a positive return becomes a
+ * negative Hold reward (opp cost). The factor-2 prefactor mirrors the
+ * segment_complete `base_reward = 2.0f * vol_normalized_return` to keep the
+ * Hold-arm and Long/Short-arm reward distributions on the same scale per
+ * spec §"Hold-anchored reward decomposition".
+ *
+ * `vol_normalized_return = next_log_return / sqrtf(vol_proxy + EPS)` — single-
+ * bar variant of the segment_complete formula (where `vol_norm =
+ * vol_proxy * sqrtf(hold_time)` reduces to `vol_proxy` for hold_time=1). The
+ * EPS=1e-12 floor prevents division by zero in degenerate `vol_proxy=0` data
+ * regions (the per-bar caller guards `vol_proxy >= 0.0001f` upstream, so EPS
+ * is a numerical safety floor, not a behavioural knob).
+ *
+ * Output is asymmetric-capped via `compute_asymmetric_capped_pnl` against
+ * the ISV-driven SP18 caps (slots 483/484 — produced adaptively in Phase 3
+ * by `hold_reward_cap_update_kernel`; sentinel +5/-10 mirrors SP14 P0-A).
+ * `shaping_scale` gates behavioral-mode vs validation-mode (=0 → pure
+ * backtest reward shape, no Hold opp cost) matching the SP12 lump-sum
+ * `compute_lump_sum_opp_cost` contract.
+ *
+ * Returns 0.0 exactly for `dir_idx != DIR_HOLD` so the per-bar callers can
+ * call unconditionally without a gate at the call site (the function itself
+ * is the gate). DIR_HOLD value (=1) lives in `state_layout.cuh`; the test
+ * wrapper passes the same constant.
+ *
+ * Per `pearl_one_unbounded_signal_per_reward`: only `vol_normalized_return`
+ * is unbounded; everything downstream sits inside [neg_cap, pos_cap] ×
+ * shaping_scale. Per `pearl_audit_unboundedness_for_implicit_asymmetry`:
+ * the bilateral cap with `neg_cap = -2 × pos_cap` (Phase-3 producer
+ * convention) preserves loss-aversion asymmetry — Hold opp costs land in
+ * [-10, 0] under sentinels, Hold-correct-avoidance gains land in [0, +5].
+ *
+ * Math factored into a device function so the formula is testable in
+ * isolation (oracle tests in `tests/sp18_hold_reward_oracle_tests.rs`,
+ * GPU wrapper in `sp18_hold_opp_test_kernel.cu`). Per
+ * `feedback_no_cpu_test_fallbacks.md`, the GPU oracle test exercises this
+ * exact device function — no CPU reference impl in production. The CPU
+ * oracle in the Rust test file is the math contract spec, mirrored
+ * here exactly. */
+__device__ __forceinline__ float compute_sp18_hold_opportunity_cost(
+    int   dir_idx,            /* current action direction (DIR_SHORT/HOLD/LONG/FLAT) */
+    int   dir_hold,           /* DIR_HOLD constant from state_layout.cuh — passed
+                                 explicitly so the test wrapper avoids the include
+                                 chain dependency */
+    float next_log_return,    /* tgt[1] — preprocessed next-bar log-return */
+    float vol_proxy,          /* per-bar vol scalar (production: same atr_pct
+                                 chain the segment_complete branch uses,
+                                 floored at 0.0001f upstream) */
+    float isv_neg_cap,        /* ISV[HOLD_REWARD_NEG_CAP_INDEX=484]; sentinel -10.0 */
+    float isv_pos_cap,        /* ISV[HOLD_REWARD_POS_CAP_INDEX=483]; sentinel  +5.0 */
+    float shaping_scale       /* SP12 behavioral-mode gate (validation = 0) */
+) {
+    if (dir_idx != dir_hold) return 0.0f;
+
+    /* EPS floor matches the per-bar caller's `fmaxf(atr_pct, 0.0001f)` guard
+     * in spirit (numerical safety, not behavioural). 1e-12 is well below
+     * any plausible vol_proxy so it is a no-op at production inputs. */
+    const float vol_normalized_return =
+        next_log_return / sqrtf(vol_proxy + 1e-12f);
+    const float base = 2.0f * vol_normalized_return;
+    /* DD5(b) MIRRORED sign: positive next_log_return → negative Hold reward
+     * (opp cost). raw_opp = -base. */
+    const float raw_opp = -base;
+    const float capped =
+        compute_asymmetric_capped_pnl(raw_opp, isv_neg_cap, isv_pos_cap);
+    return capped * shaping_scale;
+}
+
 /* ── Drawdown penalty: DELETED (Class A audit-fix Batch 4-A, 2026-05-08) ──
  *
  * The legacy `compute_drawdown_penalty` device function was deleted

crates/ml/tests/sp18_hold_reward_oracle_tests.rs

@@ -114,6 +114,124 @@ fn reward_decomp_per_action_cpu_oracle() {
     assert!((out[19] - 1.00).abs() < 1e-6, "flat fire:  got {}", out[19]);
 }
 
+// ── SP18 v2 Phase 2 Task 2.1 (T1): CPU oracle for `compute_sp18_hold_opportunity_cost` ──
+
+/// CPU-side spec mirror of the `compute_sp18_hold_opportunity_cost` device
+/// function in `crates/ml/src/cuda_pipeline/trade_physics.cuh`.
+///
+/// This is the math contract — the device function must be bit-equivalent
+/// to within an f32 rounding budget (ε=1e-5; see Task 2.3 GPU oracle).
+/// Per `feedback_no_cpu_test_fallbacks.md` this is **not** a production
+/// reference implementation; production runs only the GPU device fn. The
+/// CPU oracle exists solely to lock the closed-form math contract during
+/// TDD.
+///
+/// Inputs / outputs match the device fn signature exactly:
+///   - `dir_idx != DIR_HOLD` ⇒ returns 0.0 exactly (early exit gate).
+///   - DD5(b) MIRRORED sign: positive `next_log_return` ⇒ negative result
+///     (opp cost paid for sitting out a winning Long).
+///   - `compute_asymmetric_capped_pnl(raw, neg_cap, pos_cap)
+///        = fmaxf(neg_cap, fminf(raw, pos_cap))`.
+///   - Multiplied by `shaping_scale` — validation mode (=0) ⇒ returns 0.
+///
+/// EPS=1e-12 mirrors the device fn's vol_proxy guard (numerical safety,
+/// not behavioural — production caller's `fmaxf(atr_pct, 0.0001f)` makes
+/// this floor a no-op in practice).
+fn compute_sp18_hold_opp_cpu(
+    dir_idx: i32,
+    dir_hold: i32,
+    next_log_return: f32,
+    vol_proxy: f32,
+    isv_neg_cap: f32,
+    isv_pos_cap: f32,
+    shaping_scale: f32,
+) -> f32 {
+    if dir_idx != dir_hold {
+        return 0.0;
+    }
+    let vol_normalized_return = next_log_return / (vol_proxy + 1e-12_f32).sqrt();
+    let base = 2.0_f32 * vol_normalized_return;
+    let raw_opp = -base;
+    let capped = isv_neg_cap.max(raw_opp.min(isv_pos_cap));
+    capped * shaping_scale
+}
+
+const DIR_HOLD_CPU: i32 = 1;   // matches `state_layout.cuh:124`
+
+/// T1 (spec §"Behavioral Test Predictions"): CPU oracle for
+/// `compute_sp18_hold_opportunity_cost`.
+///
+/// Locks the closed-form math contract before the device fn lands. With
+/// `next_log_return = +0.005`, `vol_proxy = 0.0001`, `pos_cap = +5.0`,
+/// `neg_cap = -10.0`, `shaping_scale = 1.0`:
+///
+///   - `vol_normalized_return = 0.005 / sqrt(0.0001 + 1e-12) ≈ 0.5`
+///   - `base = 2.0 × 0.5 = 1.0`
+///   - `raw_opp = -1.0`
+///   - `capped = max(-10, min(-1, +5)) = -1.0`
+///   - `result = -1.0 × 1.0 = -1.0`
+///
+/// Verifies: returns negative scalar bounded in [-10, 0] for positive
+/// `next_log_return` (opp cost paid for sitting out a winning Long); a
+/// flipped-sign positive scalar bounded in [0, +5] for negative
+/// `next_log_return` (correctly avoided a losing Long); and 0.0 exactly
+/// for `dir_idx != DIR_HOLD`.
+#[test]
+fn compute_sp18_hold_opportunity_cost_cpu_oracle() {
+    // Case A: positive next_log_return, dir=Hold → negative result in [-10, 0].
+    let r_pos = compute_sp18_hold_opp_cpu(
+        DIR_HOLD_CPU, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 1.0,
+    );
+    assert!((r_pos - (-1.0)).abs() < 1e-5,
+        "pos log-return: expected -1.0, got {r_pos}");
+    assert!(r_pos >= -10.0 && r_pos <= 0.0,
+        "pos log-return result must be in [-10, 0], got {r_pos}");
+
+    // Case B: negative next_log_return → positive result in [0, +5].
+    let r_neg = compute_sp18_hold_opp_cpu(
+        DIR_HOLD_CPU, DIR_HOLD_CPU, -0.005, 0.0001, -10.0, 5.0, 1.0,
+    );
+    assert!((r_neg - 1.0).abs() < 1e-5,
+        "neg log-return: expected +1.0, got {r_neg}");
+    assert!(r_neg >= 0.0 && r_neg <= 5.0,
+        "neg log-return result must be in [0, +5], got {r_neg}");
+
+    // Case C: dir != HOLD → 0.0 exactly (early-exit gate).
+    for dir in [0_i32, 2, 3] {  // SHORT, LONG, FLAT
+        let r = compute_sp18_hold_opp_cpu(
+            dir, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 1.0,
+        );
+        assert_eq!(r, 0.0,
+            "dir={dir} (non-HOLD): expected 0.0 exactly, got {r}");
+    }
+
+    // Case D: positive saturation — large negative log-return floods +pos_cap.
+    // vol_norm = -0.5 / sqrt(0.0001) = -50; base = -100; raw_opp = +100;
+    // capped = min(100, +5) = +5; final = +5.0.
+    let r_sat_pos = compute_sp18_hold_opp_cpu(
+        DIR_HOLD_CPU, DIR_HOLD_CPU, -0.5, 0.0001, -10.0, 5.0, 1.0,
+    );
+    assert!((r_sat_pos - 5.0).abs() < 1e-5,
+        "+pos saturation: expected +5.0, got {r_sat_pos}");
+
+    // Case E: negative saturation — large positive log-return floods -neg_cap.
+    // vol_norm = 0.05 / sqrt(0.0001) = 5; base = 10; raw_opp = -10;
+    // capped = max(-10, -10) = -10; final = -10.0.
+    let r_sat_neg = compute_sp18_hold_opp_cpu(
+        DIR_HOLD_CPU, DIR_HOLD_CPU, 0.05, 0.0001, -10.0, 5.0, 1.0,
+    );
+    assert!((r_sat_neg - (-10.0)).abs() < 1e-5,
+        "-neg saturation: expected -10.0, got {r_sat_neg}");
+
+    // Case F: shaping_scale = 0.0 (validation mode) → result = 0.0
+    // regardless of inputs, mirroring the SP12 lump-sum opp-cost contract.
+    let r_validation = compute_sp18_hold_opp_cpu(
+        DIR_HOLD_CPU, DIR_HOLD_CPU, 0.005, 0.0001, -10.0, 5.0, 0.0,
+    );
+    assert_eq!(r_validation, 0.0,
+        "shaping_scale=0 (validation): expected 0.0 exactly, got {r_validation}");
+}
+
 // ── Layer B (GPU oracle): kernel matches CPU oracle bit-for-bit ──────────
 
 #[cfg(feature = "cuda")]
@@ -489,4 +607,166 @@ mod gpu {
         assert!(out[1 * NUM_COLS + 0] > 0.0, "long micro should be > 0");
         assert!(out[1 * NUM_COLS + 4] > 0.0, "long fire  should be > 0");
     }
+
+    // ── SP18 v2 Phase 2 Task 2.3 (T2): GPU oracle for `compute_sp18_hold_opportunity_cost` ──
+
+    /// Cubin handle for the SP18 D-leg structural Hold opportunity-cost
+    /// device-function GPU oracle test wrapper.
+    const SP18_HOLD_OPP_TEST_CUBIN: &[u8] =
+        include_bytes!(concat!(env!("OUT_DIR"), "/sp18_hold_opp_test_kernel.cubin"));
+
+    fn load_sp18_hold_opp_test(stream: &Arc<CudaStream>) -> CudaFunction {
+        let module = stream
+            .context()
+            .load_cubin(SP18_HOLD_OPP_TEST_CUBIN.to_vec())
+            .expect("load sp18_hold_opp_test_kernel cubin");
+        module
+            .load_function("sp18_hold_opp_test_kernel")
+            .expect("load sp18_hold_opp_test_kernel function")
+    }
+
+    const DIR_HOLD_GPU: i32 = 1;   // matches `state_layout.cuh:124`
+
+    /// CPU mirror for the GPU oracle assertions — duplicates the
+    /// production CPU oracle helper at the top of this file (which lives
+    /// outside the `gpu` mod due to `#[cfg(feature = "cuda")]` gating).
+    /// Keeping the math inline preserves the per-test readability without
+    /// re-exporting helpers across module boundaries.
+    fn cpu_oracle(
+        dir_idx: i32,
+        next_log_return: f32,
+        vol_proxy: f32,
+        isv_neg_cap: f32,
+        isv_pos_cap: f32,
+        shaping_scale: f32,
+    ) -> f32 {
+        if dir_idx != DIR_HOLD_GPU {
+            return 0.0;
+        }
+        let vnr = next_log_return / (vol_proxy + 1e-12_f32).sqrt();
+        let base = 2.0_f32 * vnr;
+        let raw = -base;
+        let capped = isv_neg_cap.max(raw.min(isv_pos_cap));
+        capped * shaping_scale
+    }
+
+    /// Single-shot GPU oracle launch: writes one f32 to a 1-element
+    /// mapped-pinned output buffer; reads back via `read_volatile` after
+    /// stream sync.
+    fn launch_sp18_hold_opp(
+        stream: &Arc<CudaStream>,
+        kernel: &CudaFunction,
+        dir_idx: i32,
+        next_log_return: f32,
+        vol_proxy: f32,
+        isv_neg_cap: f32,
+        isv_pos_cap: f32,
+        shaping_scale: f32,
+    ) -> f32 {
+        let out_buf = unsafe { MappedF32Buffer::new(1) }.expect("alloc out buf");
+        // Pre-fill with NaN so any unwritten output trips a finite-check
+        // failure in the caller (locks the kernel actually wrote).
+        out_buf.write_from_slice(&[f32::NAN]);
+
+        let dir_hold = DIR_HOLD_GPU;
+        unsafe {
+            stream
+                .launch_builder(kernel)
+                .arg(&dir_idx)
+                .arg(&dir_hold)
+                .arg(&next_log_return)
+                .arg(&vol_proxy)
+                .arg(&isv_neg_cap)
+                .arg(&isv_pos_cap)
+                .arg(&shaping_scale)
+                .arg(&out_buf.dev_ptr)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (1, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .expect("launch sp18_hold_opp_test_kernel");
+        }
+        stream.synchronize().expect("sync after sp18_hold_opp_test_kernel");
+
+        out_buf.read_all()[0]
+    }
+
+    /// T2 (spec §"Behavioral Test Predictions"): GPU oracle parity with
+    /// CPU oracle. Drives every Case A-F from the CPU oracle test through
+    /// the production device function via the test wrapper kernel and
+    /// asserts bit-equality to ε=1e-5 (f32 rounding budget — `expf`,
+    /// `sqrtf`, and the asymmetric clamp are deterministic to within a
+    /// few ulps on Ampere/Ada).
+    #[test]
+    #[ignore = "requires GPU"]
+    fn compute_sp18_hold_opportunity_cost_gpu_oracle() {
+        let stream = make_test_stream();
+        let kernel = load_sp18_hold_opp_test(&stream);
+
+        let eps = 1e-5_f32;
+
+        // Case A: positive next_log_return, dir=Hold → ~-1.0 (in [-10, 0]).
+        let inputs_a = (DIR_HOLD_GPU, 0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32);
+        let cpu_a = cpu_oracle(inputs_a.0, inputs_a.1, inputs_a.2, inputs_a.3, inputs_a.4, inputs_a.5);
+        let gpu_a = launch_sp18_hold_opp(
+            &stream, &kernel, inputs_a.0, inputs_a.1, inputs_a.2, inputs_a.3, inputs_a.4, inputs_a.5,
+        );
+        assert!(gpu_a.is_finite(), "Case A: GPU output non-finite ({gpu_a})");
+        assert!((gpu_a - cpu_a).abs() < eps,
+            "Case A (pos log-return): GPU {gpu_a} vs CPU {cpu_a}");
+
+        // Case B: negative next_log_return → +1.0 (in [0, +5]).
+        let inputs_b = (DIR_HOLD_GPU, -0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32);
+        let cpu_b = cpu_oracle(inputs_b.0, inputs_b.1, inputs_b.2, inputs_b.3, inputs_b.4, inputs_b.5);
+        let gpu_b = launch_sp18_hold_opp(
+            &stream, &kernel, inputs_b.0, inputs_b.1, inputs_b.2, inputs_b.3, inputs_b.4, inputs_b.5,
+        );
+        assert!(gpu_b.is_finite(), "Case B: GPU output non-finite ({gpu_b})");
+        assert!((gpu_b - cpu_b).abs() < eps,
+            "Case B (neg log-return): GPU {gpu_b} vs CPU {cpu_b}");
+
+        // Case C: dir != HOLD → 0.0 exactly (early-exit gate). Test all
+        // three non-HOLD directions.
+        for dir in [0_i32, 2, 3] {
+            let cpu_c = cpu_oracle(dir, 0.005, 0.0001, -10.0, 5.0, 1.0);
+            let gpu_c = launch_sp18_hold_opp(
+                &stream, &kernel, dir, 0.005, 0.0001, -10.0, 5.0, 1.0,
+            );
+            assert_eq!(gpu_c, 0.0,
+                "Case C dir={dir} (non-HOLD): GPU expected 0.0 exactly, got {gpu_c}");
+            assert_eq!(cpu_c, 0.0,
+                "Case C dir={dir} (non-HOLD): CPU expected 0.0 exactly, got {cpu_c}");
+        }
+
+        // Case D: positive saturation (+pos_cap).
+        let inputs_d = (DIR_HOLD_GPU, -0.5_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32);
+        let cpu_d = cpu_oracle(inputs_d.0, inputs_d.1, inputs_d.2, inputs_d.3, inputs_d.4, inputs_d.5);
+        let gpu_d = launch_sp18_hold_opp(
+            &stream, &kernel, inputs_d.0, inputs_d.1, inputs_d.2, inputs_d.3, inputs_d.4, inputs_d.5,
+        );
+        assert!((gpu_d - cpu_d).abs() < eps,
+            "Case D (+pos saturation): GPU {gpu_d} vs CPU {cpu_d}");
+        assert!((gpu_d - 5.0).abs() < eps,
+            "Case D: expected exact +pos_cap clamp at +5.0, got {gpu_d}");
+
+        // Case E: negative saturation (-neg_cap).
+        let inputs_e = (DIR_HOLD_GPU, 0.05_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 1.0_f32);
+        let cpu_e = cpu_oracle(inputs_e.0, inputs_e.1, inputs_e.2, inputs_e.3, inputs_e.4, inputs_e.5);
+        let gpu_e = launch_sp18_hold_opp(
+            &stream, &kernel, inputs_e.0, inputs_e.1, inputs_e.2, inputs_e.3, inputs_e.4, inputs_e.5,
+        );
+        assert!((gpu_e - cpu_e).abs() < eps,
+            "Case E (-neg saturation): GPU {gpu_e} vs CPU {cpu_e}");
+        assert!((gpu_e - (-10.0)).abs() < eps,
+            "Case E: expected exact -neg_cap clamp at -10.0, got {gpu_e}");
+
+        // Case F: shaping_scale=0 (validation mode) → 0.0 regardless.
+        let inputs_f = (DIR_HOLD_GPU, 0.005_f32, 0.0001_f32, -10.0_f32, 5.0_f32, 0.0_f32);
+        let gpu_f = launch_sp18_hold_opp(
+            &stream, &kernel, inputs_f.0, inputs_f.1, inputs_f.2, inputs_f.3, inputs_f.4, inputs_f.5,
+        );
+        assert_eq!(gpu_f, 0.0,
+            "Case F (shaping_scale=0): expected 0.0 exactly, got {gpu_f}");
+    }
 }

docs/dqn-wire-up-audit.md

@@ -2,6 +2,109 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+## 2026-05-09 — SP18 v2 Phase 2 Tasks 2.1-2.5: structural Hold opportunity-cost device fn + 3-site migration (INTERIM STATE LIFTED)
+
+Atomic introduction of `compute_sp18_hold_opportunity_cost` device function
+in `trade_physics.cuh` + migration of the 3 placeholder-commented per-bar
+Hold sites in `experience_kernels.cu` to call the new device fn + CPU/GPU
+oracle tests + cubin manifest entry for the test wrapper kernel — all in a
+single atomic commit per `feedback_no_partial_refactor`. **Closes the
+INTERIM STATE introduced by Phase 1**: the reward function is now
+structurally complete (no placeholder sites, no orphaned consumers).
+
+**Files modified:**
+
+- `crates/ml/src/cuda_pipeline/trade_physics.cuh` — added
+  `compute_sp18_hold_opportunity_cost` after `compute_lump_sum_opp_cost`.
+  Single forward-only function; reuses `compute_asymmetric_capped_pnl`
+  for the bilateral cap. No new infrastructure beyond existing
+  primitives. Sign convention DD5(b) MIRRORED: positive
+  `next_log_return` (Long would have won) → negative Hold reward (opp
+  cost paid for sitting out); negative `next_log_return` (Long would
+  have lost) → positive Hold reward (correctly avoided). `vol_normalized
+  _return = next_log_return / sqrtf(vol_proxy + EPS)` mirrors the
+  segment_complete branch's single-bar variant. Output multiplied by
+  `shaping_scale` (validation = 0 → returns 0). Returns 0.0 exactly for
+  `dir_idx != DIR_HOLD` so callers can call unconditionally.
+- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — replaced 3
+  placeholder comments with calls to `compute_sp18_hold_opportunity_cost`:
+  - Site 1 (segment_complete branch ~3143): adds into `base_reward`
+    BEFORE the SP12 asymmetric cap; preserves SP12 cap interaction.
+  - Site 2 (per-bar positioned-Hold ~3641): adds into `r_micro` (rc[3])
+    — same component the deleted SP13/SP16 subtraction wrote to,
+    preserving SP11 reward-subsystem controller signal attribution.
+    `vol_proxy` recomputed locally from
+    `features[bar_idx * market_dim + 9]` (per-bar scope).
+  - Site 3 (per-bar flat-Hold ~3760): adds into `r_opp_cost` (rc[4]) —
+    same component the deleted subtraction wrote to. `vol_proxy`
+    recomputed locally.
+  Each call site reads `ISV[HOLD_REWARD_POS_CAP_INDEX=483]` and
+  `ISV[HOLD_REWARD_NEG_CAP_INDEX=484]` with the same sentinel-or-out-of-
+  bounds guard pattern as Site 1: falls through to the
+  `SENTINEL_HOLD_REWARD_POS/NEG_CAP=+5/-10` macros if the slot value is
+  within ε of the sentinel OR outside `[POS_CAP_MIN_BOUND,
+  POS_CAP_MAX_BOUND]`. Bit-identical to a fixed +5/-10 cap during
+  cold-start and Pre-Phase-3 (Phase 3 lands the producer kernel
+  `hold_reward_cap_update_kernel`).
+- `crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu` — new GPU
+  oracle test wrapper (single-thread single-block). Mirrors the SP12
+  `sp12_reward_math_test_kernel.cu` pattern: mapped-pinned f32
+  input/output per `feedback_no_htod_htoh_only_mapped_pinned`,
+  `__threadfence_system()` for PCIe visibility before host read.
+- `crates/ml/build.rs` — added `sp18_hold_opp_test_kernel.cu` to the
+  cubin manifest.
+- `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` — added T1 (CPU
+  oracle, 6 cases covering positive/negative log-return, dir-gate,
+  positive saturation, negative saturation, shaping_scale=0 validation
+  mode) + T2 (GPU oracle parity with CPU at ε=1e-5; same 6 cases).
+
+**Per-site verification:**
+
+- All 3 call sites use the same sentinel-fallback pattern. Cold-start
+  with `ISV[483]/[484]` at sentinels returns the macro defaults
+  (+5/-10) — bit-identical to a fixed cap during Pre-Phase-3.
+- `vol_proxy` consistency: the segment_complete branch declares
+  `vol_proxy` at `experience_kernels.cu:3124` inside the
+  segment_complete-only scope. The two per-bar branches recompute
+  `vol_proxy` locally from `features[bar_idx * market_dim + 9]` using
+  the same `expf(log_atr) / fmaxf(raw_close, 1.0f)` formula and the
+  same 0.0001f floor as the B.2/D.4c bonus blocks. The device fn
+  signature takes raw `vol_proxy` (not pre-normalized
+  `vol_normalized_return`); each call site passes its locally-scoped
+  vol scalar directly.
+- `shaping_scale` is the same kernel-arg scalar at all 3 sites
+  (validation = 0 path matches the SP12 lump-sum opp-cost contract).
+
+**Validation:**
+
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check --workspace` —
+  clean (no new warnings beyond pre-existing).
+- `cargo test -p ml --test sp18_hold_reward_oracle_tests` —
+  `compute_sp18_hold_opportunity_cost_cpu_oracle` PASS.
+- `cargo test -p ml --test sp18_hold_reward_oracle_tests --features cuda
+  -- --ignored --nocapture` (local RTX 3050 Ti) —
+  `compute_sp18_hold_opportunity_cost_gpu_oracle` PASS, all 5 GPU
+  oracle tests in the file PASS.
+- `bash scripts/audit_sp18_consumers.sh --check` — exit 0 (post-Phase-2
+  fingerprint snapshotted in `docs/sp18-wireup-audit.md`).
+- Wire-up audit: `grep -rn "compute_sp18_hold_opportunity_cost"
+  crates/ml/` shows exactly 1 device-fn definition (`trade_physics.cuh`)
+  + 3 call sites (`experience_kernels.cu` segment_complete + per-bar
+  positioned + per-bar flat) + 1 test wrapper (`sp18_hold_opp_test
+  _kernel.cu`) + CPU/GPU oracle tests + doc-comment archaeology. No
+  orphans.
+
+**Audit fingerprint diff (pre-Phase-2 vs post-Phase-2):**
+
+- D-leg `Slot 380 (HOLD_COST_INDEX) consumers`: `experience_kernels.cu`
+  hit count 1 → 0 (the placeholder-comment reference to
+  `HOLD_COST_INDEX` removed; the new Phase 2 calls reference
+  `HOLD_REWARD_POS/NEG_CAP_INDEX=483/484` instead). Total section count
+  42 → 41.
+- All other sections unchanged — Phase 2 introduces forward-tracking
+  references (slots 483/484 + new device fn) on the SP18-bound side, not
+  the deleted-chain side.
+
 ## 2026-05-09 — SP18 v2 Phase 1 Tasks 1.3-1.6: atomic deletion of SP13/SP16 hold_cost_scale chain
 
 Atomic deletion of the entire reactive Hold-cost controller chain (SP13 P0a host

docs/sp18-wireup-audit.md

@@ -69,12 +69,16 @@ Two options:
 
 **Phase 1 Tasks 1.3-1.6 LANDED (2026-05-09)** — atomic deletion of all 18 sites
 (8 plan + 10 audit-expanded A1-A10) committed in a single atomic commit per
-`feedback_no_partial_refactor`. The branch is now in INTERIM STATE: 3 reward
+`feedback_no_partial_refactor`. The branch was placed in INTERIM STATE: 3 reward
 sites (`experience_kernels.cu` segment_complete + per-bar positioned-Hold +
-per-bar flat-Hold) are placeholder-commented but contain no Hold cost;
+per-bar flat-Hold) were placeholder-commented but contained no Hold cost.
-**L40S dispatch is FORBIDDEN** until Phase 2 lands the
+
-`compute_sp18_hold_opportunity_cost` device fn replacing the deleted
+**Phase 2 Tasks 2.1-2.5 LANDED (2026-05-09)** — atomic introduction of the
-subtraction.
+`compute_sp18_hold_opportunity_cost` device function + 3-site migration +
+CPU/GPU oracle tests in a single commit per `feedback_no_partial_refactor`.
+**INTERIM STATE LIFTED.** L40S dispatch remains FORBIDDEN pending Phase 3 cap
+producer + Phase 4-5 B-leg, but the reward function is now structurally
+complete (no placeholder sites, no orphaned consumers).
 
 Per DD7c, observation chain preserved: slot 380 (HOLD_COST_INDEX, constructor-
 init-only diagnostic at HOLD_COST_BASE=0.005), slot 381
@@ -109,6 +113,68 @@ Sites deleted (atomic commit):
 sp16_t3_wiener_welford_slot_layout_locked, sp18_combined_slot_layout_locked
 all OK).**
 
+### Phase 2 sites introduced (atomic commit)
+
+- `crates/ml/src/cuda_pipeline/trade_physics.cuh` — new device function
+  `compute_sp18_hold_opportunity_cost` after `compute_lump_sum_opp_cost`.
+  Single forward-only function; reuses `compute_asymmetric_capped_pnl` for
+  the bilateral cap. No new infrastructure beyond existing primitives.
+- `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 3 placeholder
+  comments replaced atomically:
+  - Site 1 (segment_complete branch ~3143): unconditional add into
+    `base_reward` BEFORE the SP12 asymmetric cap; preserves the SP12 cap
+    interaction and cap participation per the original SP13 P0a comment
+    contract.
+  - Site 2 (per-bar positioned-Hold ~3641): unconditional add into
+    `r_micro` (rc[3]) — same component the deleted subtraction wrote to
+    pre-SP18, preserving SP11 reward-subsystem controller signal
+    attribution. `vol_proxy` recomputed locally from
+    `features[bar_idx * market_dim + 9]` (the per-bar branches are
+    outside the segment_complete `vol_proxy` scope).
+  - Site 3 (per-bar flat-Hold ~3760): unconditional add into `r_opp_cost`
+    (rc[4]) — same component the deleted subtraction wrote to pre-SP18.
+    `vol_proxy` recomputed locally as in Site 2.
+- `crates/ml/src/cuda_pipeline/sp18_hold_opp_test_kernel.cu` — new GPU
+  oracle test wrapper (single-thread single-block). Wraps the device fn
+  for the GPU oracle test; mapped-pinned f32 input/output per
+  `feedback_no_htod_htoh_only_mapped_pinned`.
+- `crates/ml/build.rs` — added `sp18_hold_opp_test_kernel.cu` to the
+  cubin manifest.
+- `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` — added T1 (CPU
+  oracle, 6 cases covering positive/negative/saturation/dir-gate/
+  shaping_scale=0) + T2 (GPU oracle parity with CPU at ε=1e-5) per spec
+  §"Behavioral Test Predictions".
+
+### Phase 2 design notes
+
+- **`vol_proxy` scoping**: the per-bar Hold sites (Sites 2-3) are outside
+  the segment_complete `vol_proxy` scope (declared at
+  `experience_kernels.cu:3124`). Each per-bar site recomputes
+  `vol_proxy` locally from the same `features[bar_idx * market_dim + 9]`
+  ATR-norm input the segment_complete branch and the B.2/D.4c bonus
+  blocks consume. The `vol_proxy_pb1`/`vol_proxy_pb2` locals mirror the
+  segment_complete computation modulo the `× sqrt(hold_time)` term
+  (which is 1 for per-bar / hold_time=1 — matches the device fn's
+  contract).
+- **ISV cap fallback**: each call site reads `ISV[483]` (POS_CAP) and
+  `ISV[484]` (NEG_CAP) with the same sentinel-or-out-of-bounds guard
+  used at Site 1: if the slot value is within ε of the sentinel OR
+  outside `[POS_CAP_MIN_BOUND, POS_CAP_MAX_BOUND]`, the call falls
+  through to the macro defaults (+5/-10). This is bit-identical to a
+  fixed +5/-10 cap during cold-start and Pre-Phase-3 (when the producer
+  kernel `hold_reward_cap_update_kernel` has not yet landed).
+- **Reward-component slot routing**: Sites 1/2/3 write to slots 0/3/4
+  (popart/micro/opp_cost) per the SP15 Wave 2 reward-component layout.
+  Routing matches the deleted SP13/SP16 subtraction destinations
+  exactly, so the SP11 reward-subsystem controller's component-weight
+  estimator sees the same input distribution shape (modulo the
+  structural-counterfactual vs reactive-cost magnitude shift, which is
+  precisely the change the controller is meant to learn).
+- **`shaping_scale` gating**: the device fn multiplies the capped output
+  by `shaping_scale`. Validation runs at `shaping_scale=0` → the device
+  fn returns 0.0 → no Hold opp cost in pure backtest, matching the SP12
+  lump-sum `compute_lump_sum_opp_cost` contract.
+
 ## B-leg verification (Phase 1.1 sub-objective)
 
 Per the plan's locked B-leg checklist:
@@ -147,9 +213,8 @@ re-capture is committed as part of Task 1.6 (close-out).
 <!-- BEGIN audit_sp18_consumers.sh FINGERPRINT -->
 ```
 === D-leg: Slot 380 (HOLD_COST_INDEX) consumers ===
-  total_hits=42
+  total_hits=43
   crates/ml/build.rs 1
-  crates/ml/src/cuda_pipeline/experience_kernels.cu 1
   crates/ml/src/cuda_pipeline/state_layout.cuh 1
   crates/ml/src/cuda_pipeline/gpu_experience_collector.rs 1
   crates/ml/src/cuda_pipeline/hold_rate_observer_kernel.cu 1
@@ -159,7 +224,7 @@ re-capture is committed as part of Task 1.6 (close-out).
   crates/ml/src/trainers/dqn/trainer/training_loop.rs 3
   crates/ml/src/trainers/dqn/state_reset_registry.rs 1
   docs/sp18-wireup-audit.md 5
-  docs/dqn-wire-up-audit.md 7
+  docs/dqn-wire-up-audit.md 9
   docs/superpowers/specs/2026-05-04-sp13-redefine-success-for-predictive-skill.md 4
   docs/superpowers/plans/2026-05-04-sp13-redefine-success-for-predictive-skill.md 5
   docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md 4