feat(sp15-p1.2): cost-net sharpe kernel — commission + spread + OFI-impact

Per spec §6.2 per-side semantics:
  cost_t = commission_per_rt × rt_ind[t]
         + half_spread[t] × |pos[t]| × side_ind[t]
         + ofi_lambda × |pos[t]| × |ofi[t]| × side_ind[t]

Commission charged at close; half-spread × |pos| at entry AND exit
(sums to one full spread per RT); OFI impact same per-side. Initial
λ=2.0e-4 in ISV[OFI_IMPACT_LAMBDA_INDEX=407] as Invariant-1 anchor
(constructor-write + FoldReset rewrite per feedback_isv_for_adaptive_bounds);
per-fold ISV refit may overwrite in later phases. Mean cost-per-bar
emitted to ISV[COST_PER_BAR_AVG_INDEX=408] (stateful kernel output;
FoldReset sentinel 0 + Pearl A first-observation bootstrap).

Same kernel reads LobBar fields from synthetic markets (Phase 2A) and
real fxcache LOB (prod) — dev/prod parity per Q3.

Phase 1.2 lands kernel + launcher + anchor seed + 2 registry entries
+ 2 dispatch arms only; consumer migration deferred to a follow-up
commit per feedback_no_partial_refactor (mirrors Phase 1.1 atomic
pattern). Oracle test cost_net_sharpe_round_trip_charges_full_spread
validates single round-trip cost = 2.00 / 10 bars = 0.20/bar with
mean_pnl_net = 0.80 on 1.69s RTX 3050 Ti (sm_86). Zero regressions
introduced (946 passed / 13 pre-existing failures, same as Task 1.1).

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

crates/ml/build.rs

@@ -833,6 +833,16 @@ fn main() {
         // (val × sqrt(525600)) split. Annualisation is the host-side caller's
         // responsibility (sharpe_annualised = out[2] × sqrt(N_bars_per_year)).
         "sharpe_per_bar_kernel.cu",
+        // SP15 Phase 1.2 (2026-05-06): cost-net sharpe. Single-block,
+        // BLOCK=256 kernel computing mean/std/sharpe of (pnl - cost) per
+        // spec §6.2 per-side semantics: commission_per_rt × rt_ind +
+        // half_spread × |pos| × side_ind + ofi_lambda × |pos| × |ofi| ×
+        // side_ind. Charges full spread per RT (entry + exit). Reads
+        // ofi_lambda from ISV[OFI_IMPACT_LAMBDA_INDEX=407]; emits
+        // mean cost-per-bar to ISV[COST_PER_BAR_AVG_INDEX=408]. Same
+        // kernel reads LobBar fields from synthetic markets and prod
+        // fxcache LOB (dev/prod parity per Q3).
+        "cost_net_sharpe_kernel.cu",
     ];
 
     // ALL kernels get common header (BF16 types + wrappers)

crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu

@@ -0,0 +1,111 @@
+// crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu
+//
+// SP15 Phase 1.2 — cost-net sharpe.
+//
+// Per spec §6.2 per-side semantics:
+//   cost_t = commission_per_rt × rt_ind[t]
+//          + half_spread[t] × |position[t]| × side_ind[t]
+//          + ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]
+//
+// Commission is charged at close (rt_ind=1). Half-spread × |position| is
+// charged on side-events at entry AND exit (sums to one full spread per
+// round-trip). OFI impact same per-side. ofi_lambda is read from
+// ISV[OFI_IMPACT_LAMBDA_INDEX=407] (constant across the window;
+// constructor-seeded to 2e-4, FoldReset-rewritten to 2e-4 since this is an
+// Invariant-1 anchor, not a stateful EMA).
+//
+// Outputs:
+//   isv[COST_PER_BAR_AVG_INDEX=408] = sum(cost_t) / n
+//   out[0] = mean(pnl - cost_t)
+//   out[1] = std(pnl - cost_t)
+//   out[2] = sharpe_net = out[0] / max(out[1], 1e-12)
+//
+// Same kernel reads LobBar fields from synthetic markets (Phase 2A) and
+// real fxcache LOB (prod) — dev/prod parity per Q3.
+//
+// Single-block launch, BLOCK=256. No atomicAdd (per feedback_no_atomicadd):
+// reduction uses shared-memory tree-reduce only. `__threadfence_system()`
+// flushes mapped-pinned writes before the host reads via `read_volatile`.
+
+extern "C" __global__ void cost_net_sharpe_kernel(
+    const float*        __restrict__ pnl,
+    const float*        __restrict__ half_spread,
+    const float*        __restrict__ ofi,
+    const unsigned int* __restrict__ side_ind,
+    const unsigned int* __restrict__ rt_ind,
+    const float*        __restrict__ position,
+    int                 n,
+    float               commission_per_rt,
+    float*              __restrict__ isv,
+    float*              __restrict__ out
+) {
+    if (blockIdx.x != 0) return;
+
+    const int tid   = (int)threadIdx.x;
+    const int BLOCK = (int)blockDim.x;
+    __shared__ float reduce_a[256];
+    __shared__ float reduce_b[256];
+    __shared__ float mean_pnl_shared;
+    __shared__ float mean_cost_shared;
+
+    // Read OFI lambda once (slot 407 = OFI_IMPACT_LAMBDA_INDEX).
+    const float ofi_lambda = isv[407];
+
+    // Pass 1: parallel sums of pnl_t and cost_t.
+    float local_pnl  = 0.0f;
+    float local_cost = 0.0f;
+    for (int i = tid; i < n; i += BLOCK) {
+        local_pnl += pnl[i];
+        float c_comm   = commission_per_rt * (float)rt_ind[i];
+        float pos_abs  = fabsf(position[i]);
+        float c_spread = half_spread[i] * pos_abs * (float)side_ind[i];
+        float c_ofi    = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i];
+        local_cost += c_comm + c_spread + c_ofi;
+    }
+    reduce_a[tid] = local_pnl;
+    reduce_b[tid] = local_cost;
+    __syncthreads();
+    for (int s = BLOCK / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            reduce_a[tid] += reduce_a[tid + s];
+            reduce_b[tid] += reduce_b[tid + s];
+        }
+        __syncthreads();
+    }
+    if (tid == 0) {
+        mean_pnl_shared  = (n > 0) ? reduce_a[0] / (float)n : 0.0f;
+        mean_cost_shared = (n > 0) ? reduce_b[0] / (float)n : 0.0f;
+        // ISV slot 408 = COST_PER_BAR_AVG_INDEX (per sp15_isv_slots.rs).
+        isv[408] = mean_cost_shared;
+        __threadfence_system();
+    }
+    __syncthreads();
+
+    const float mean_net = mean_pnl_shared - mean_cost_shared;
+
+    // Pass 2: variance of (pnl_t - cost_t).
+    float local_sq = 0.0f;
+    for (int i = tid; i < n; i += BLOCK) {
+        float c_comm   = commission_per_rt * (float)rt_ind[i];
+        float pos_abs  = fabsf(position[i]);
+        float c_spread = half_spread[i] * pos_abs * (float)side_ind[i];
+        float c_ofi    = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i];
+        float net      = pnl[i] - (c_comm + c_spread + c_ofi);
+        float d        = net - mean_net;
+        local_sq += d * d;
+    }
+    reduce_a[tid] = local_sq;
+    __syncthreads();
+    for (int s = BLOCK / 2; s > 0; s >>= 1) {
+        if (tid < s) reduce_a[tid] += reduce_a[tid + s];
+        __syncthreads();
+    }
+    if (tid == 0) {
+        float var = (n > 0) ? reduce_a[0] / (float)n : 0.0f;
+        float std = sqrtf(fmaxf(var, 1e-12f));
+        out[0] = mean_net;
+        out[1] = std;
+        out[2] = (std > 0.0f) ? mean_net / std : 0.0f;
+        __threadfence_system();
+    }
+}

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -754,6 +754,82 @@ pub fn launch_sp15_sharpe_per_bar(
     Ok(())
 }
 
+/// SP15 Phase 1.2 (2026-05-06): cost-net sharpe cubin. Per spec §6.2
+/// per-side cost semantics: commission_per_rt charged at close (rt_ind=1);
+/// half-spread × |pos| charged at entry AND exit on side_ind events (sums
+/// to one full spread per round-trip); ofi_lambda × |pos| × |ofi| same
+/// per-side. Reads ofi_lambda from ISV[OFI_IMPACT_LAMBDA_INDEX=407];
+/// emits mean cost-per-bar to ISV[COST_PER_BAR_AVG_INDEX=408]. Same
+/// kernel reads LobBar fields from synthetic markets (Phase 2A) and real
+/// fxcache LOB (prod) — dev/prod parity per Q3.
+pub static SP15_COST_NET_SHARPE_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/cost_net_sharpe_kernel.cubin"));
+
+/// SP15 Phase 1.2 (2026-05-06): launcher for the cost-net sharpe kernel.
+/// Free function (matches `launch_sp15_sharpe_per_bar` precedent) so unit
+/// / oracle tests can drive the kernel directly without the trainer struct;
+/// production callers invoke the same launcher.
+///
+/// Per-side cost semantics (spec §6.2):
+///   `cost_t = commission_per_rt × rt_ind[t]`
+///          `+ half_spread[t] × |position[t]| × side_ind[t]`
+///          `+ ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]`
+///
+/// `pnl`, `half_spread`, `ofi`, `position`, `isv`, `out` are device f32
+/// pointers. `side_ind` and `rt_ind` are device u32 pointers. `out` MUST
+/// point to ≥3 f32 slots: `[mean_pnl_net, std, sharpe_net]`. `isv` MUST
+/// be the ISV bus (≥443 f32 slots) — slot 407 is read for `ofi_lambda`,
+/// slot 408 is written with `mean(cost_t)`.
+#[allow(clippy::too_many_arguments)]
+pub fn launch_sp15_cost_net_sharpe(
+    stream: &Arc<CudaStream>,
+    pnl: cudarc::driver::sys::CUdeviceptr,
+    half_spread: cudarc::driver::sys::CUdeviceptr,
+    ofi: cudarc::driver::sys::CUdeviceptr,
+    side_ind: cudarc::driver::sys::CUdeviceptr,
+    rt_ind: cudarc::driver::sys::CUdeviceptr,
+    position: cudarc::driver::sys::CUdeviceptr,
+    n: i32,
+    commission_per_rt: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+    out: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_COST_NET_SHARPE_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_cost_net_sharpe cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("cost_net_sharpe_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load cost_net_sharpe_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&pnl)
+            .arg(&half_spread)
+            .arg(&ofi)
+            .arg(&side_ind)
+            .arg(&rt_ind)
+            .arg(&position)
+            .arg(&n)
+            .arg(&commission_per_rt)
+            .arg(&isv)
+            .arg(&out)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch cost_net_sharpe_kernel: {e}"
+            )))?;
+    }
+    Ok(())
+}
+
 /// SP11 Fix 39 (2026-05-04, Task A2): SimHash novelty signal — lookup +
 /// update kernels sharing one cubin. Lookup reads
 /// `1/sqrt(1+count)` ∈ [0, 1] for each (state, action) bucket; update
@@ -19727,6 +19803,21 @@ impl GpuDqnTrainer {
                 *sig_ptr.add(HOLD_COST_INDEX)        = HOLD_COST_BASE;
                 *sig_ptr.add(HOLD_RATE_TARGET_INDEX) = HOLD_RATE_TARGET_DEFAULT;
 
+                // SP15 Phase 1.2 (2026-05-06): OFI-impact lambda
+                // (per spec §6.2). Constant Invariant-1 anchor — NOT
+                // a stateful EMA — so constructor-write per
+                // `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV
+                // refit may overwrite this in later phases; the
+                // FoldReset arm rewrites the same default to keep the
+                // cost kernel functional during cold-start of each
+                // new fold (must never reach sentinel 0, which would
+                // silently disable the OFI-impact term). The matching
+                // COST_PER_BAR_AVG_INDEX=408 is a stateful kernel
+                // output (FoldReset sentinel 0; first kernel launch
+                // writes the true mean).
+                use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX;
+                *sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4_f32;
+
                 // Layout fingerprint (ISV[58..60)). Compile-time structural hash of
                 // the slot layout; checkpoint load fails-fast on mismatch.
                 // Stored as a u64 split across two f32 lanes using raw bit-cast so

crates/ml/src/trainers/dqn/state_reset_registry.rs

@@ -1005,6 +1005,31 @@ impl StateResetRegistry {
                 category: ResetCategory::FoldReset,
                 description: "ISV[AUX_DIR_ACC_POST_OPEN_MIN_INDEX=394] — SP14 anti-gradient-hacking circuit breaker: running minimum of aux_dir_acc since Gate 1 last opened. Initialised to 1.0 (no minimum observed) when Gate 1 opens; updated each step post-open as `min = fminf(min, current_aux_dir_acc)`. If `open_min − current_aux_dir_acc > LOCKOUT_TRIGGER_DROP=0.05` while q_disagreement also rises by >LOCKOUT_TRIGGER_DIS_RISE=0.10, the lockout circuit fires (B.5) and writes LOCKOUT_EPOCHS=2.0 to GRADIENT_HACK_LOCKOUT_REMAINING_INDEX=395. FoldReset sentinel SENTINEL_POST_OPEN_MIN=1.0 (no minimum observed at fold start; the circuit breaker must accumulate post-open observations before it can fire). Produced by `alpha_grad_compute_kernel` (B.4). Spec: docs/superpowers/specs/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md.",
             },
+            // ── SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots ────────────
+            // Two ISV slots [407, 408]:
+            //   - OFI_IMPACT_LAMBDA_INDEX=407: Invariant-1 anchor (NOT a
+            //     stateful EMA). Constructor-seeded to 2.0e-4 per spec §6.2;
+            //     FoldReset rewrites the same anchor (NOT sentinel 0) so the
+            //     cost kernel functions correctly on the new fold's first
+            //     launch. Per-fold ISV refit may overwrite this in later
+            //     phases; the FoldReset arm restores the default to break
+            //     cold-start of each new fold (must never reach sentinel 0,
+            //     which would silently disable the OFI-impact term).
+            //   - COST_PER_BAR_AVG_INDEX=408: stateful kernel output
+            //     (mean(cost_t)). FoldReset sentinel 0; Pearl A bootstraps
+            //     from the new fold's first kernel launch per
+            //     `pearl_first_observation_bootstrap.md`.
+            // Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md.
+            RegistryEntry {
+                name: "sp15_ofi_impact_lambda",
+                category: ResetCategory::FoldReset,
+                description: "ISV[OFI_IMPACT_LAMBDA_INDEX=407] — SP15 Phase 1.2 (2026-05-06) OFI-impact lambda for the cost-net sharpe kernel (per spec §6.2). Per-side cost: `λ × |position| × |ofi| × side_ind`. Constant Invariant-1 anchor (2.0e-4) per `feedback_isv_for_adaptive_bounds.md` — NOT a stateful EMA. Constructor-seeded; FoldReset rewrites the same default so the cost kernel functions correctly on the new fold's first launch (must never reach sentinel 0, which would silently disable the OFI-impact term). Per-fold ISV refit may overwrite this in later phases; the reset path restores the default to break cold-start of each new fold. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.2.",
+            },
+            RegistryEntry {
+                name: "sp15_cost_per_bar_avg",
+                category: ResetCategory::FoldReset,
+                description: "ISV[COST_PER_BAR_AVG_INDEX=408] — SP15 Phase 1.2 (2026-05-06) mean cost-per-bar emitted by the cost-net sharpe kernel (`sum(cost_t) / n` over the sharpe window). Stateful kernel output (NOT cross-fold persistent); FoldReset sentinel 0 so the new fold's first `launch_sp15_cost_net_sharpe` writes the new fold's true mean. Per `pearl_first_observation_bootstrap.md`. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.2.",
+            },
             // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf
             // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) +
             // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3

crates/ml/src/trainers/dqn/trainer/training_loop.rs

@@ -8066,6 +8066,33 @@ impl DQNTrainer {
                     );
                 }
             }
+            // SP15 Phase 1.2 (2026-05-06): cost-net sharpe slots.
+            // OFI_IMPACT_LAMBDA_INDEX=407 is an Invariant-1 anchor (NOT
+            // a stateful EMA) — rewrite the constructor's value at fold
+            // boundary so the cost kernel functions correctly on the new
+            // fold's first launch (must never reach sentinel 0, which
+            // would silently disable the OFI-impact term). Per
+            // `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV refit
+            // may overwrite this in later phases.
+            "sp15_ofi_impact_lambda" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX;
+                    fused.trainer().write_isv_signal_at(
+                        OFI_IMPACT_LAMBDA_INDEX,
+                        2.0e-4,
+                    );
+                }
+            }
+            // COST_PER_BAR_AVG_INDEX=408 is a stateful kernel output —
+            // FoldReset sentinel 0 so Pearl A bootstraps from the new
+            // fold's first `launch_sp15_cost_net_sharpe` per
+            // `pearl_first_observation_bootstrap.md`.
+            "sp15_cost_per_bar_avg" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp15_isv_slots::COST_PER_BAR_AVG_INDEX;
+                    fused.trainer().write_isv_signal_at(COST_PER_BAR_AVG_INDEX, 0.0);
+                }
+            }
             // SP11 Task A2 (2026-05-04): novelty visit-count hash table
             // reset arm — closes the A0 deferral per the registry entry's
             // docstring. 1M f32 slots, zero-filled via mapped-pinned host

crates/ml/tests/sp15_phase1_oracle_tests.rs

@@ -10,8 +10,22 @@ mod gpu {
     use std::sync::Arc;
 
     use cudarc::driver::{CudaContext, CudaStream};
-    use ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_sharpe_per_bar;
-    use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
+    use ml::cuda_pipeline::gpu_dqn_trainer::{
+        launch_sp15_cost_net_sharpe, launch_sp15_sharpe_per_bar,
+    };
+    use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer};
+    use ml::cuda_pipeline::sp15_isv_slots::{
+        COST_PER_BAR_AVG_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
+    };
+
+    /// Upper bound for any ISV index touched by SP15 oracle tests. Matches
+    /// `gpu_dqn_trainer::ISV_TOTAL_DIM` (which is `pub(crate)` and not
+    /// reachable from integration tests). The `sp15_isv_slots` regression
+    /// test `all_sp15_slots_fit_within_isv_total_dim` enforces that
+    /// `SP15_SLOT_END <= ISV_TOTAL_DIM`, so allocating an ISV buffer of
+    /// length `SP15_SLOT_END` is large enough for every SP15 slot the
+    /// kernel reads or writes.
+    const ISV_LEN: usize = SP15_SLOT_END;
 
     /// Build a CUDA stream against device 0. Mirrors the helper in
     /// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`. If no
@@ -88,4 +102,114 @@ mod gpu {
         let out = out_buf.read_all();
         assert!((out[2] - 0.5).abs() < 0.05, "sharpe={}, expected ~0.5", out[2]);
     }
+
+    /// Test 1.2 — single round-trip oracle for the cost-net sharpe kernel.
+    ///
+    /// Setup: n=10 bars, position +1 contract held throughout, half_spread
+    /// 0.25, commission $1.50/RT, ofi=0 (zero out the OFI-impact term so
+    /// the spread + commission contributions are isolated).
+    /// Side events: bar 0 (entry, side_ind=1) and bar 5 (exit,
+    /// side_ind=1, rt_ind=1). Gross PnL: 10.0 booked at bar 5 (the
+    /// close-out bar).
+    ///
+    /// Expected per spec §6.2:
+    ///   total cost = 1.50 commission + 0.25 (entry half-spread)
+    ///              + 0.25 (exit half-spread) = 2.00
+    ///   cost_per_bar_avg = 2.00 / 10 = 0.20
+    ///   mean(pnl) = 10.0 / 10 = 1.0
+    ///   mean_pnl_net = 1.0 − 0.20 = 0.80
+    ///
+    /// Validates:
+    ///   • Per-side cost semantics (entry + exit = full spread per RT).
+    ///   • Commission charged at close (rt_ind=1 only on bar 5).
+    ///   • OFI-impact term scales with `|position|·|ofi|·side_ind`
+    ///     (here zero because `ofi=0`).
+    ///   • COST_PER_BAR_AVG_INDEX (slot 408) emitted to ISV.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn cost_net_sharpe_round_trip_charges_full_spread() {
+        let stream = make_test_stream();
+        let n = 10usize;
+
+        // Inputs.
+        let mut pnl = vec![0.0f32; n];
+        pnl[5] = 10.0;
+        let half_spread = vec![0.25f32; n];
+        let ofi = vec![0.0f32; n];
+        let position = vec![1.0f32; n];
+        let mut side_ind = vec![0u32; n];
+        side_ind[0] = 1;
+        side_ind[5] = 1;
+        let mut rt_ind = vec![0u32; n];
+        rt_ind[5] = 1;
+
+        // Safety: CUDA context active via `make_test_stream` above.
+        let pnl_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for pnl");
+        pnl_buf.write_from_slice(&pnl);
+        let half_spread_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for half_spread");
+        half_spread_buf.write_from_slice(&half_spread);
+        let ofi_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for ofi");
+        ofi_buf.write_from_slice(&ofi);
+        let position_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for position");
+        position_buf.write_from_slice(&position);
+        let side_ind_buf = unsafe { MappedU32Buffer::new(n) }
+            .expect("alloc MappedU32Buffer for side_ind");
+        side_ind_buf.write_from_slice(&side_ind);
+        let rt_ind_buf = unsafe { MappedU32Buffer::new(n) }
+            .expect("alloc MappedU32Buffer for rt_ind");
+        rt_ind_buf.write_from_slice(&rt_ind);
+
+        // ISV bus seeded with OFI_IMPACT_LAMBDA=0 (mirrors the spec's
+        // zero-OFI test variable). The constructor seeds slot 407 to 2e-4
+        // in production; here we override to 0 so the spread term is the
+        // sole non-commission cost contributor.
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
+            .expect("alloc MappedF32Buffer for isv bus");
+        let mut isv_init = vec![0.0f32; ISV_LEN];
+        isv_init[OFI_IMPACT_LAMBDA_INDEX] = 0.0;
+        isv_buf.write_from_slice(&isv_init);
+
+        let out_buf = unsafe { MappedF32Buffer::new(3) }
+            .expect("alloc MappedF32Buffer for cost-net sharpe output");
+
+        launch_sp15_cost_net_sharpe(
+            &stream,
+            pnl_buf.dev_ptr,
+            half_spread_buf.dev_ptr,
+            ofi_buf.dev_ptr,
+            side_ind_buf.dev_ptr,
+            rt_ind_buf.dev_ptr,
+            position_buf.dev_ptr,
+            n as i32,
+            /* commission_per_rt = */ 1.50,
+            isv_buf.dev_ptr,
+            out_buf.dev_ptr,
+        )
+        .expect("launch cost_net_sharpe_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after cost_net_sharpe_kernel launch");
+
+        let isv_host = isv_buf.read_all();
+        let cost_per_bar_avg = isv_host[COST_PER_BAR_AVG_INDEX];
+        // Total cost = 1.50 commission + 0.25 + 0.25 = 2.00 over 10 bars
+        // → cost_per_bar_avg = 0.20.
+        assert!(
+            (cost_per_bar_avg - 0.20).abs() < 1e-3,
+            "cost_per_bar_avg = {}, expected 0.20",
+            cost_per_bar_avg
+        );
+
+        let out = out_buf.read_all();
+        // mean_pnl_net = mean(pnl) − mean(cost) = 1.0 − 0.20 = 0.80.
+        assert!(
+            (out[0] - 0.80).abs() < 1e-2,
+            "mean_pnl_net = {}, expected 0.80",
+            out[0]
+        );
+    }
 }

docs/dqn-wire-up-audit.md

@@ -4,6 +4,8 @@
 
 SP15 Phase 1.1 — unified per-bar sharpe kernel (2026-05-06): single GPU kernel `sharpe_per_bar_kernel.cu` computes mean/std/sharpe via 2-pass shared-memory tree-reduce (single-block, BLOCK=256, no atomicAdd). Replaces the SP14-era split between `sharpe_ema` (per-batch EMA, train) and `sharpe_annualised` (val × sqrt(525600)). Same formula and same window definition for train and val; annualisation is the host-side caller's responsibility (`sharpe_annualised = out[2] × sqrt(N_bars_per_year)`). **Phase 1.1 lands kernel + free-function launcher only**; consumer migration in `metrics.rs` / `training_loop.rs` is deferred to a follow-up commit per `feedback_no_partial_refactor.md` (atomic consumer migration is a load-bearing change too risky to bundle with kernel introduction; the kernel is verified-working in isolation first via the two GPU oracle tests below). Touched: `crates/ml/src/cuda_pipeline/sharpe_per_bar_kernel.cu` (new), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_SHARPE_PER_BAR_CUBIN` cubin slot + 1 `pub fn launch_sp15_sharpe_per_bar` free-function launcher matching the `sp4_histogram_p99_test_kernel` cold-path standalone-launcher precedent), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+2 GPU oracle tests inside the existing `mod gpu` block: `unified_sharpe_kernel_zero_mean` validates mean=0/std=1/sharpe=0 on alternating ±1 PnL; `unified_sharpe_kernel_positive_drift` validates sharpe≈0.5 on `0.5 + alternating ±1` PnL). Both tests pass on local RTX 3050 Ti (sm_86) in 1.78s. cargo test -p ml --lib --features cuda: 946 passed / 13 failed — all 13 failures pre-existing on the parent commit (verified via `git stash` baseline run); zero introduced by this commit. Hard rules: `feedback_no_atomicadd` (block-tree-reduce only), `feedback_no_partial_refactor` (kernel + launcher land atomically; consumer migration follows as its own atomic commit), `feedback_no_stubs` (launcher returns `Result<(), MLError>` and is fully functional, not a placeholder), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for both input PnL and output `[mean, std, sharpe]`).
 
+SP15 Phase 1.2 — cost-net sharpe kernel (2026-05-06): single GPU kernel `cost_net_sharpe_kernel.cu` computes mean(pnl − cost) / std(pnl − cost) sharpe with per-side cost semantics per spec §6.2 (`cost_t = commission_per_rt × rt_ind[t] + half_spread[t] × |position[t]| × side_ind[t] + ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]`). Commission charged at close (rt_ind=1); half-spread × |pos| charged at entry AND exit (sums to one full spread per round-trip); OFI impact same per-side. Single-block BLOCK=256 2-pass shared-memory tree-reduce — no atomicAdd. Reads `ofi_lambda` from `ISV[OFI_IMPACT_LAMBDA_INDEX=407]`; emits mean cost-per-bar to `ISV[COST_PER_BAR_AVG_INDEX=408]`. Same kernel reads LobBar fields from synthetic markets (Phase 2A) and real fxcache LOB (prod) — dev/prod parity per Q3. **Phase 1.2 lands kernel + free-function launcher + ISV anchor seed + state-reset registry entries only**; consumer migration in metrics / val sites is deferred to a follow-up commit per `feedback_no_partial_refactor.md` (consumer migration is a load-bearing change; the kernel is verified-working in isolation first via the GPU oracle test below, mirroring the Phase 1.1 atomic pattern). Touched: `crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu` (new), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_COST_NET_SHARPE_CUBIN` cubin slot + 1 `pub fn launch_sp15_cost_net_sharpe` free-function launcher mirroring `launch_sp15_sharpe_per_bar`'s `stream.context().load_cubin` precedent + 1 constructor-write `*sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4` Invariant-1 anchor seed at the SP13/SP14 anchor block per `feedback_isv_for_adaptive_bounds.md`), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+2 `RegistryEntry` records: `sp15_ofi_impact_lambda` rewrites the constructor's 2.0e-4 default at fold boundary since this is an Invariant-1 anchor not a stateful EMA — must never reach sentinel 0 which would silently disable the OFI-impact term; `sp15_cost_per_bar_avg` is a stateful kernel output, FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first launch per `pearl_first_observation_bootstrap.md`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+2 dispatch arms for the 2 new registry entries — registry-arm contract enforced by `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 GPU oracle test `cost_net_sharpe_round_trip_charges_full_spread` validates a single round-trip on n=10 bars: 1 entry @ bar 0 + 1 exit @ bar 5, position +1 contract, half_spread 0.25, commission 1.50, ofi_lambda overridden to 0 in the test ISV bus to isolate spread+commission. Expected: `cost_per_bar_avg = (1.50 + 0.25 + 0.25) / 10 = 0.20` and `mean_pnl_net = 10/10 − 0.20 = 0.80`). Test passes on local RTX 3050 Ti (sm_86) in 1.69s. cargo test -p ml --lib --features cuda: 946 passed / 13 failed — same 13 failures pre-existing on the parent commit `3667cd1b0` (Task 1.1 baseline); zero introduced by this commit. State-reset registry tests: 4/4 pass including `every_fold_and_soft_reset_entry_has_dispatch_arm` validating both new entries have matching dispatch arms. Hard rules: `feedback_no_atomicadd` (block-tree-reduce only), `feedback_no_partial_refactor` (kernel + launcher + anchor-seed + registry entries + dispatch arms land atomically; consumer migration follows as its own atomic commit), `feedback_no_stubs` (launcher returns `Result<(), MLError>` and is fully functional), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for f32 inputs/output and `MappedU32Buffer` for u32 side/rt indicators), `feedback_isv_for_adaptive_bounds` (OFI lambda is a constant Invariant-1 anchor written at constructor + FoldReset, NOT a runtime EMA — per-fold ISV refit may overwrite this in later phases).
+
 SP12 v3 — per-trade event-driven reward composition (2026-05-04): three architectural changes in one atomic commit per spec `2026-05-04-sp12-quality-over-quantity-reward-shaping.md` (committed at `ecab584c3`). Empirical motivation: 50-epoch validation `train-multi-seed-pmbwn` on commit `6a259942e` showed sharpe-gaming (sharpe=80 held while PnL declined 30% over 8 epochs; 62% trade rate, 46% win rate stuck). **Three causes** addressed in one commit per `feedback_no_partial_refactor`:
 
 (1) **Asymmetric bounded cap** at `experience_kernels.cu:2758` — replaces SP11's symmetric ±10 cap (commit `35db31089`) with `fmaxf(REWARD_NEG_CAP=-10, fminf(base_reward, REWARD_POS_CAP=+5))`. The 2:1 ratio matches Kahneman/Tversky meta-analysis pegging human loss aversion at ~2.0-2.25 — calibration to a known psychological constant (Invariant-1 numerical anchor in `state_layout.cuh`), not a tuned multiplier. SP11's stability work preserved (cap still bilateral). Per `pearl_audit_unboundedness_for_implicit_asymmetry`: SP11 erased the pre-fix implicit downside-asymmetry by symmetrizing the bound; the model became risk-neutral and locked into HFT noise extraction. SP12 widens only the negative cap back toward the pre-SP11 unbounded regime.