diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 7280a4156..5cc31f855 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -854,6 +854,25 @@ fn main() {
         // DD_PCT (406). The 1e-4 calmar floor eliminates the saturation-
         // at-100 artifact previously seen in the train-dd4xl HEALTH_DIAG.
         "dd_state_kernel.cu",
+        // SP15 Phase 1.4 partial (2026-05-06): 4 constant-policy
+        // counterfactual baselines per spec §6.4. Single source file with
+        // 4 `extern "C" __global__` symbols (baseline_buyhold_kernel,
+        // baseline_hold_only_kernel, baseline_naive_momentum_kernel,
+        // baseline_naive_reversion_kernel) sharing one cubin per the 1:1
+        // source-to-cubin convention — the four launchers in
+        // gpu_dqn_trainer.rs all load the same cubin module and resolve
+        // different `get_function()` symbols. Each kernel: single-block
+        // BLOCK=256 with a templated 2-pass shared-memory tree-reduce
+        // (no atomicAdd). Reads price/half_spread/ofi arrays + commission
+        // scalar; computes constant-policy or last-bar-return-based
+        // actions; writes sharpe_net to its own ISV slot (409 buyhold,
+        // 410 hold_only, 412 naive_momentum, 416 naive_reversion).
+        // Trunk-shared baselines (random_dir_kelly slot 411, aux_only
+        // slot 413, mag_quarter_fixed slot 414, trail_only slot 415) are
+        // out of scope for this commit — they need partial-policy-
+        // forward access from the main eval pass and are deferred to
+        // Task 1.4.b.
+        "baseline_kernels.cu",
     ];
 
     // ALL kernels get common header (BF16 types + wrappers)
diff --git a/crates/ml/src/cuda_pipeline/baseline_kernels.cu b/crates/ml/src/cuda_pipeline/baseline_kernels.cu
new file mode 100644
index 000000000..01ca24293
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/baseline_kernels.cu
@@ -0,0 +1,225 @@
+// crates/ml/src/cuda_pipeline/baseline_kernels.cu
+//
+// SP15 Phase 1.4 (partial) — 4 constant-policy counterfactual baselines.
+// Per spec §6.4. Pure-CUDA kernels (no trunk-share). Each kernel reads
+// the price/spread/ofi arrays + a commission scalar, computes the
+// baseline's per-bar PnL stream from a constant or last-bar-return-based
+// action policy, then runs a 2-pass shared-memory tree-reduce mean/std
+// (no atomicAdd per `feedback_no_atomicadd`) and writes sharpe_net to
+// its allocated ISV slot.
+//
+// Single source file — multiple `extern "C" __global__` symbols share
+// one cubin per the build.rs 1:1-source-to-cubin convention. The four
+// launchers in `gpu_dqn_trainer.rs` load this single cubin and resolve
+// different `get_function()` symbols, mirroring the multi-kernel-per-
+// file pattern from SP11 (`novelty_simhash_kernel.cu`).
+//
+// ISV slots written (allocated in `sp15_isv_slots.rs`, spec §4.3):
+//   slot 409 — BASELINE_BUYHOLD_SHARPE_INDEX
+//   slot 410 — BASELINE_HOLD_ONLY_SHARPE_INDEX
+//   slot 412 — BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX
+//   slot 416 — BASELINE_NAIVE_REVERSION_SHARPE_INDEX
+//
+// Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413,
+// mag_quarter_fixed slot 414, trail_only slot 415) are out of scope
+// for this commit — they need partial-policy-forward access from the
+// main eval pass. Those slots stay at sentinel 0.0 until follow-up
+// Task 1.4.b lands them atomically with the eval-side wire-up.
+//
+// Cost semantics (matches `cost_net_sharpe_kernel.cu` per-side rules
+// from spec §6.2): a "trade event" is any bar where the action changes
+// from the prior bar (curr_action != prev_action). On a trade event
+// the bar pays half the round-trip commission plus a half-spread cost
+// times |position|; for these constant-policy baselines |position|=1
+// when in the market. The OFI-impact term is included with weight 0
+// (i.e. the baseline does not pay OFI impact) — the test harness pins
+// `ofi[]` to 0 anyway, so this is moot in oracle tests; in production
+// the comparator-net-of-cost is fair because the cost model matches
+// the model-under-test (both pay full RT commission + entry/exit
+// spread, neither pays OFI impact for this baseline class).
+
+#define BUYHOLD_SHARPE_SLOT      409
+#define HOLD_ONLY_SHARPE_SLOT    410
+#define MOMENTUM_SHARPE_SLOT     412
+#define REVERSION_SHARPE_SLOT    416
+
+// Templated helper: 2-pass shared-memory tree-reduce of mean and std of
+// the per-bar PnL stream produced by `ActionFn`. Returns sharpe to thread
+// 0 only — caller stores it to the appropriate ISV slot.
+//
+// Template parameter `ActionFn` is a CUDA functor:
+//   __device__ int operator()(const float* prices, int i, int n) const;
+//   returns +1 (long), 0 (flat/hold), -1 (short) for bar `i`.
+template <typename ActionFn>
+__device__ float compute_baseline_sharpe(
+    const float* __restrict__ prices,
+    const float* __restrict__ half_spread,
+    const float* __restrict__ /* ofi */,  // unused for this baseline class (lambda=0)
+    int                       n,
+    float                     commission_per_rt,
+    ActionFn                  action_fn
+) {
+    const int tid   = (int)threadIdx.x;
+    const int BLOCK = (int)blockDim.x;
+    __shared__ float reduce_buf[256];
+    __shared__ float mean_shared;
+
+    // Pass 1: per-bar PnL contribution (price diff × prev_action) net of
+    // a half-RT commission + half-spread cost on each action change.
+    // Bars indexed 1..n-1 (need a prior price for the diff).
+    float local_pnl = 0.0f;
+    for (int i = tid + 1; i < n; i += BLOCK) {
+        const int   prev_action = action_fn(prices, i - 1, n);
+        const int   curr_action = action_fn(prices, i, n);
+        const float gross       = (float)prev_action * (prices[i] - prices[i - 1]);
+        float       bar_pnl     = gross;
+        if (curr_action != prev_action) {
+            // Trade event at bar i: pay half RT commission + half-spread
+            // × |position|. |position|=1 for these baselines whenever
+            // they hold a position; |position|=0 for hold_only (which
+            // never changes action so this branch is never taken).
+            const float pos_abs = 1.0f;
+            bar_pnl -= 0.5f * commission_per_rt + half_spread[i] * pos_abs;
+        }
+        local_pnl += bar_pnl;
+    }
+    reduce_buf[tid] = local_pnl;
+    __syncthreads();
+    for (int s = BLOCK / 2; s > 0; s >>= 1) {
+        if (tid < s) reduce_buf[tid] += reduce_buf[tid + s];
+        __syncthreads();
+    }
+    if (tid == 0) {
+        mean_shared = (n > 1) ? reduce_buf[0] / (float)(n - 1) : 0.0f;
+    }
+    __syncthreads();
+
+    // Pass 2: variance about mean → std → sharpe.
+    float local_sq = 0.0f;
+    for (int i = tid + 1; i < n; i += BLOCK) {
+        const int   prev_action = action_fn(prices, i - 1, n);
+        const int   curr_action = action_fn(prices, i, n);
+        const float gross       = (float)prev_action * (prices[i] - prices[i - 1]);
+        float       bar_pnl     = gross;
+        if (curr_action != prev_action) {
+            const float pos_abs = 1.0f;
+            bar_pnl -= 0.5f * commission_per_rt + half_spread[i] * pos_abs;
+        }
+        const float d = bar_pnl - mean_shared;
+        local_sq += d * d;
+    }
+    reduce_buf[tid] = local_sq;
+    __syncthreads();
+    for (int s = BLOCK / 2; s > 0; s >>= 1) {
+        if (tid < s) reduce_buf[tid] += reduce_buf[tid + s];
+        __syncthreads();
+    }
+    if (tid == 0) {
+        const float var = (n > 1) ? reduce_buf[0] / (float)(n - 1) : 0.0f;
+        const float std = sqrtf(fmaxf(var, 1e-12f));
+        return (std > 0.0f) ? mean_shared / std : 0.0f;
+    }
+    return 0.0f;
+}
+
+// CUDA functor structs — one per baseline policy.
+struct BuyholdAction {
+    __device__ int operator()(const float* /*prices*/, int /*i*/, int /*n*/) const {
+        return 1;  // always Long position=+1
+    }
+};
+struct HoldOnlyAction {
+    __device__ int operator()(const float* /*prices*/, int /*i*/, int /*n*/) const {
+        return 0;  // always Hold; never enter the market
+    }
+};
+struct MomentumAction {
+    __device__ int operator()(const float* prices, int i, int /*n*/) const {
+        if (i == 0) return 0;  // no prior bar → no signal
+        const float diff = prices[i] - prices[i - 1];
+        if (diff > 0.0f) return  1;
+        if (diff < 0.0f) return -1;
+        return 0;
+    }
+};
+struct ReversionAction {
+    __device__ int operator()(const float* prices, int i, int /*n*/) const {
+        if (i == 0) return 0;
+        const float diff = prices[i] - prices[i - 1];
+        if (diff > 0.0f) return -1;
+        if (diff < 0.0f) return  1;
+        return 0;
+    }
+};
+
+extern "C" __global__ void baseline_buyhold_kernel(
+    const float* __restrict__ prices,
+    const float* __restrict__ half_spread,
+    const float* __restrict__ ofi,
+    int                       n,
+    float                     commission_per_rt,
+    float* __restrict__       isv
+) {
+    if (blockIdx.x != 0) return;
+    const float sharpe = compute_baseline_sharpe(
+        prices, half_spread, ofi, n, commission_per_rt, BuyholdAction()
+    );
+    if (threadIdx.x == 0) {
+        isv[BUYHOLD_SHARPE_SLOT] = sharpe;
+        __threadfence_system();
+    }
+}
+
+extern "C" __global__ void baseline_hold_only_kernel(
+    const float* __restrict__ prices,
+    const float* __restrict__ half_spread,
+    const float* __restrict__ ofi,
+    int                       n,
+    float                     commission_per_rt,
+    float* __restrict__       isv
+) {
+    if (blockIdx.x != 0) return;
+    const float sharpe = compute_baseline_sharpe(
+        prices, half_spread, ofi, n, commission_per_rt, HoldOnlyAction()
+    );
+    if (threadIdx.x == 0) {
+        isv[HOLD_ONLY_SHARPE_SLOT] = sharpe;
+        __threadfence_system();
+    }
+}
+
+extern "C" __global__ void baseline_naive_momentum_kernel(
+    const float* __restrict__ prices,
+    const float* __restrict__ half_spread,
+    const float* __restrict__ ofi,
+    int                       n,
+    float                     commission_per_rt,
+    float* __restrict__       isv
+) {
+    if (blockIdx.x != 0) return;
+    const float sharpe = compute_baseline_sharpe(
+        prices, half_spread, ofi, n, commission_per_rt, MomentumAction()
+    );
+    if (threadIdx.x == 0) {
+        isv[MOMENTUM_SHARPE_SLOT] = sharpe;
+        __threadfence_system();
+    }
+}
+
+extern "C" __global__ void baseline_naive_reversion_kernel(
+    const float* __restrict__ prices,
+    const float* __restrict__ half_spread,
+    const float* __restrict__ ofi,
+    int                       n,
+    float                     commission_per_rt,
+    float* __restrict__       isv
+) {
+    if (blockIdx.x != 0) return;
+    const float sharpe = compute_baseline_sharpe(
+        prices, half_spread, ofi, n, commission_per_rt, ReversionAction()
+    );
+    if (threadIdx.x == 0) {
+        isv[REVERSION_SHARPE_SLOT] = sharpe;
+        __threadfence_system();
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 9b4b07872..afa3f69ad 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -903,6 +903,213 @@ pub fn launch_sp15_dd_state(
     Ok(())
 }
 
+/// SP15 Phase 1.4 partial (2026-05-06): cubin shared by all 4 constant-
+/// policy counterfactual baselines (`baseline_buyhold_kernel`,
+/// `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`,
+/// `baseline_naive_reversion_kernel`). Per spec §6.4. Each baseline
+/// computes its sharpe_net by running a constant-policy or last-bar-
+/// return-based action stream through a 2-pass shared-memory tree-reduce
+/// (single-block BLOCK=256, no atomicAdd) and writing to its own ISV
+/// slot. The four launchers below all load this single cubin and resolve
+/// different `get_function()` symbols — multi-kernel-per-file pattern
+/// per `novelty_simhash_kernel.cu` precedent.
+///
+/// Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413,
+/// mag_quarter_fixed slot 414, trail_only slot 415) are out of scope for
+/// this commit — they need partial-policy-forward access from the main
+/// eval pass and are deferred to Task 1.4.b. Their slots stay at sentinel
+/// 0.0 in the meantime.
+pub static SP15_BASELINE_KERNELS_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/baseline_kernels.cubin"));
+
+/// SP15 Phase 1.4 partial (2026-05-06): launcher for the buyhold
+/// counterfactual baseline. Free function (matches `launch_sp15_dd_state`
+/// precedent) so unit/oracle tests can drive the kernel directly without
+/// the trainer struct; production callers (eval-pass baseline launches)
+/// invoke the same launcher.
+///
+/// Policy: position=+1 every bar (no exits). Writes sharpe_net to
+/// `ISV[BASELINE_BUYHOLD_SHARPE_INDEX=409]`. `prices`, `half_spread`,
+/// `ofi`, `isv` are device f32 pointers; `n` is bar count. `isv` MUST
+/// be the ISV bus (≥443 f32 slots).
+pub fn launch_sp15_baseline_buyhold(
+    stream: &Arc<CudaStream>,
+    prices: cudarc::driver::sys::CUdeviceptr,
+    half_spread: cudarc::driver::sys::CUdeviceptr,
+    ofi: cudarc::driver::sys::CUdeviceptr,
+    n: i32,
+    commission_per_rt: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_baseline_kernels cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("baseline_buyhold_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load baseline_buyhold_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&prices)
+            .arg(&half_spread)
+            .arg(&ofi)
+            .arg(&n)
+            .arg(&commission_per_rt)
+            .arg(&isv)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch baseline_buyhold_kernel: {e}"
+            )))?;
+    }
+    Ok(())
+}
+
+/// SP15 Phase 1.4 partial (2026-05-06): launcher for the hold-only
+/// counterfactual baseline. Policy: action=Hold every bar (no entry, no
+/// exit). Writes sharpe_net to `ISV[BASELINE_HOLD_ONLY_SHARPE_INDEX=410]`.
+/// On a strict no-trade trajectory `mean(pnl)=0` and `std(pnl)=0` so
+/// the kernel's `(std > 0) ? mean / std : 0` guard emits 0, which is
+/// the spec-mandated sentinel for an undefined sharpe.
+pub fn launch_sp15_baseline_hold_only(
+    stream: &Arc<CudaStream>,
+    prices: cudarc::driver::sys::CUdeviceptr,
+    half_spread: cudarc::driver::sys::CUdeviceptr,
+    ofi: cudarc::driver::sys::CUdeviceptr,
+    n: i32,
+    commission_per_rt: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_baseline_kernels cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("baseline_hold_only_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load baseline_hold_only_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&prices)
+            .arg(&half_spread)
+            .arg(&ofi)
+            .arg(&n)
+            .arg(&commission_per_rt)
+            .arg(&isv)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch baseline_hold_only_kernel: {e}"
+            )))?;
+    }
+    Ok(())
+}
+
+/// SP15 Phase 1.4 partial (2026-05-06): launcher for the naive-momentum
+/// counterfactual baseline. Policy: `direction = sign(prices[i] −
+/// prices[i-1])` with |position|=1. Writes sharpe_net to
+/// `ISV[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX=412]`.
+pub fn launch_sp15_baseline_naive_momentum(
+    stream: &Arc<CudaStream>,
+    prices: cudarc::driver::sys::CUdeviceptr,
+    half_spread: cudarc::driver::sys::CUdeviceptr,
+    ofi: cudarc::driver::sys::CUdeviceptr,
+    n: i32,
+    commission_per_rt: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_baseline_kernels cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("baseline_naive_momentum_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load baseline_naive_momentum_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&prices)
+            .arg(&half_spread)
+            .arg(&ofi)
+            .arg(&n)
+            .arg(&commission_per_rt)
+            .arg(&isv)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch baseline_naive_momentum_kernel: {e}"
+            )))?;
+    }
+    Ok(())
+}
+
+/// SP15 Phase 1.4 partial (2026-05-06): launcher for the naive-reversion
+/// counterfactual baseline. Policy: `direction = -sign(prices[i] −
+/// prices[i-1])` with |position|=1 (mirror of naive_momentum). Writes
+/// sharpe_net to `ISV[BASELINE_NAIVE_REVERSION_SHARPE_INDEX=416]`.
+pub fn launch_sp15_baseline_naive_reversion(
+    stream: &Arc<CudaStream>,
+    prices: cudarc::driver::sys::CUdeviceptr,
+    half_spread: cudarc::driver::sys::CUdeviceptr,
+    ofi: cudarc::driver::sys::CUdeviceptr,
+    n: i32,
+    commission_per_rt: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_BASELINE_KERNELS_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_baseline_kernels cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("baseline_naive_reversion_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load baseline_naive_reversion_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&prices)
+            .arg(&half_spread)
+            .arg(&ofi)
+            .arg(&n)
+            .arg(&commission_per_rt)
+            .arg(&isv)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (256, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch baseline_naive_reversion_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
diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs
index aa7eed965..c72e11756 100644
--- a/crates/ml/tests/sp15_phase1_oracle_tests.rs
+++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs
@@ -11,10 +11,14 @@ mod gpu {
 
     use cudarc::driver::{CudaContext, CudaStream};
     use ml::cuda_pipeline::gpu_dqn_trainer::{
+        launch_sp15_baseline_buyhold, launch_sp15_baseline_hold_only,
+        launch_sp15_baseline_naive_momentum, launch_sp15_baseline_naive_reversion,
         launch_sp15_cost_net_sharpe, launch_sp15_dd_state, launch_sp15_sharpe_per_bar,
     };
     use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer};
     use ml::cuda_pipeline::sp15_isv_slots::{
+        BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX,
+        BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX, BASELINE_NAIVE_REVERSION_SHARPE_INDEX,
         COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX,
         DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
     };
@@ -294,4 +298,212 @@ mod gpu {
             dd_pct
         );
     }
+
+    /// Test 1.4.a — buyhold on a positive-drift price series.
+    ///
+    /// Build a synthetic price walk with mean-per-bar return +0.5 and
+    /// alternating ±1 noise (std≈1 per bar). Buyhold with position=+1
+    /// every bar collects the per-bar return + noise, so the sharpe
+    /// (mean/std of per-bar PnL) should be near 0.5 (drift / noise).
+    /// Bounds chosen to absorb the n=1000 sample-size noise: lower 0.2
+    /// rules out a sharpe-near-zero degenerate baseline; upper 1.0
+    /// catches accidental zero-cost or reduction bugs that would
+    /// inflate the sharpe.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn baseline_buyhold_positive_on_drift() {
+        let stream = make_test_stream();
+
+        // Synthetic +drift price series: per-bar diff = 0.5 + alternating
+        // ±1 noise → mean(diff)=0.5, std(diff)=1 → sharpe≈0.5 for
+        // buyhold (position=+1 every bar).
+        let n = 1000usize;
+        let mut prices = vec![4500.0f32; n];
+        for i in 1..n {
+            let dz: f32 = if i % 2 == 0 { 1.0 } else { -1.0 };
+            prices[i] = prices[i - 1] + 0.5 + dz;
+        }
+        let half_spread = vec![0.0f32; n]; // zero spread for clean test
+        let ofi = vec![0.0f32; n];
+
+        // Safety: CUDA context active via `make_test_stream` above.
+        let prices_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for prices");
+        prices_buf.write_from_slice(&prices);
+        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 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);
+
+        launch_sp15_baseline_buyhold(
+            &stream,
+            prices_buf.dev_ptr,
+            half_spread_buf.dev_ptr,
+            ofi_buf.dev_ptr,
+            n as i32,
+            /* commission_per_rt = */ 0.0,
+            isv_buf.dev_ptr,
+        )
+        .expect("launch baseline_buyhold_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after baseline_buyhold_kernel launch");
+
+        let isv = isv_buf.read_all();
+        let sharpe = isv[BASELINE_BUYHOLD_SHARPE_INDEX];
+        assert!(
+            sharpe > 0.2,
+            "buyhold sharpe = {}, expected > 0.2 on +drift series",
+            sharpe
+        );
+        assert!(
+            sharpe < 1.0,
+            "buyhold sharpe = {}, expected < 1.0 (sanity bound)",
+            sharpe
+        );
+    }
+
+    /// Test 1.4.b — hold_only on any price series emits sharpe = 0.
+    ///
+    /// HoldOnly never enters the market, so per-bar PnL is identically
+    /// zero on every bar. The kernel's reduction yields mean=0 std=0
+    /// and the `(std > 0) ? mean / std : 0` guard returns 0, the
+    /// spec-mandated sentinel for an undefined sharpe (no trades).
+    #[test]
+    #[ignore = "requires GPU"]
+    fn baseline_hold_only_emits_zero() {
+        let stream = make_test_stream();
+        let n = 1000usize;
+        let prices = vec![4500.0f32; n];
+        let half_spread = vec![0.0f32; n];
+        let ofi = vec![0.0f32; n];
+
+        // Safety: CUDA context active via `make_test_stream` above.
+        let prices_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for prices");
+        prices_buf.write_from_slice(&prices);
+        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 isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
+            .expect("alloc MappedF32Buffer for isv bus");
+        isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]);
+
+        launch_sp15_baseline_hold_only(
+            &stream,
+            prices_buf.dev_ptr,
+            half_spread_buf.dev_ptr,
+            ofi_buf.dev_ptr,
+            n as i32,
+            /* commission_per_rt = */ 0.0,
+            isv_buf.dev_ptr,
+        )
+        .expect("launch baseline_hold_only_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after baseline_hold_only_kernel launch");
+
+        let isv = isv_buf.read_all();
+        let sharpe = isv[BASELINE_HOLD_ONLY_SHARPE_INDEX];
+        assert!(
+            sharpe.abs() < 1e-5,
+            "hold_only sharpe = {}, expected ~0 (no trades)",
+            sharpe
+        );
+    }
+
+    /// Test 1.4.d + 1.4.h — symmetry: naive_momentum and naive_reversion
+    /// are sign-flipped policies, so on the same price series their
+    /// sharpe values should sum to ~0. The series is deliberately
+    /// mean-reverting (reversion outperforms) but the symmetry test is
+    /// invariant to which side wins.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn baseline_momentum_reversion_symmetry() {
+        let stream = make_test_stream();
+
+        // Mean-reverting series: AR(1) with negative coefficient anchors
+        // prices around 4500. On this series naive_momentum picks the
+        // wrong direction and naive_reversion picks the right one — the
+        // sum near zero is the sharpe-symmetry invariant we test.
+        let n = 1000usize;
+        let mut prices = vec![4500.0f32; n];
+        for i in 1..n {
+            let pull = (4500.0 - prices[i - 1]) * 0.3; // strong reversion
+            let dz: f32 = if i % 2 == 0 { 1.0 } else { -1.0 };
+            prices[i] = prices[i - 1] + pull + dz;
+        }
+        let half_spread = vec![0.0f32; n];
+        let ofi = vec![0.0f32; n];
+
+        // Safety: CUDA context active via `make_test_stream` above.
+        let prices_buf = unsafe { MappedF32Buffer::new(n) }
+            .expect("alloc MappedF32Buffer for prices");
+        prices_buf.write_from_slice(&prices);
+        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 isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
+            .expect("alloc MappedF32Buffer for isv bus");
+        isv_buf.write_from_slice(&vec![0.0f32; ISV_LEN]);
+
+        launch_sp15_baseline_naive_momentum(
+            &stream,
+            prices_buf.dev_ptr,
+            half_spread_buf.dev_ptr,
+            ofi_buf.dev_ptr,
+            n as i32,
+            0.0,
+            isv_buf.dev_ptr,
+        )
+        .expect("launch baseline_naive_momentum_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after baseline_naive_momentum_kernel launch");
+
+        launch_sp15_baseline_naive_reversion(
+            &stream,
+            prices_buf.dev_ptr,
+            half_spread_buf.dev_ptr,
+            ofi_buf.dev_ptr,
+            n as i32,
+            0.0,
+            isv_buf.dev_ptr,
+        )
+        .expect("launch baseline_naive_reversion_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after baseline_naive_reversion_kernel launch");
+
+        let isv = isv_buf.read_all();
+        let momentum_sharpe = isv[BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX];
+        let reversion_sharpe = isv[BASELINE_NAIVE_REVERSION_SHARPE_INDEX];
+
+        // Sign-flipped policies: their sum should be near zero (anti-
+        // correlated picks on the same per-bar return stream). Bound 0.5
+        // absorbs the asymmetry from cost charges on action-change bars
+        // (each policy hits trade events on different bars).
+        assert!(
+            (momentum_sharpe + reversion_sharpe).abs() < 0.5,
+            "momentum + reversion = {} + {} = {}, expected ~0 (anti-correlated)",
+            momentum_sharpe,
+            reversion_sharpe,
+            momentum_sharpe + reversion_sharpe
+        );
+    }
 }
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 2feb00aeb..02a782812 100644
--- a/docs/dqn-wire-up-audit.md
+++ b/docs/dqn-wire-up-audit.md
@@ -2,6 +2,8 @@
 
 **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
 
+SP15 Phase 1.4 partial — 4 constant-policy counterfactual baselines (2026-05-06): single source file `baseline_kernels.cu` with 4 `extern "C" __global__` symbols (`baseline_buyhold_kernel`, `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`, `baseline_naive_reversion_kernel`) sharing one cubin per the build.rs 1:1-source-to-cubin convention with multiple kernels per file (mirroring the SP11 `novelty_simhash_kernel.cu` precedent — lookup + update kernels in one cubin). Each kernel: single-block BLOCK=256 with a templated `compute_baseline_sharpe<ActionFn>` device-helper that runs a 2-pass shared-memory tree-reduce of mean/std of the per-bar PnL stream produced by a CUDA functor (`BuyholdAction` / `HoldOnlyAction` / `MomentumAction` / `ReversionAction`) — no atomicAdd per `feedback_no_atomicadd`. Cost semantics match `cost_net_sharpe_kernel.cu` (per-side rules from spec §6.2): a "trade event" is any bar where `curr_action != prev_action`; on a trade event the bar pays `0.5 × commission_per_rt + half_spread[i] × |pos|` (|pos|=1 for these baselines whenever they hold a position; HoldOnly never changes action so the branch is never taken). OFI-impact term is omitted (lambda=0 implicit) — matches the test harness's `ofi[]=0` and the constant-policy class is explicitly not paying the production model's OFI-impact charge. ISV slots written: `BASELINE_BUYHOLD_SHARPE_INDEX=409`, `BASELINE_HOLD_ONLY_SHARPE_INDEX=410`, `BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX=412`, `BASELINE_NAIVE_REVERSION_SHARPE_INDEX=416`. **Trunk-shared baselines (random_dir_kelly slot 411, aux_only slot 413, mag_quarter_fixed slot 414, trail_only slot 415) are out of scope for this commit** — they need partial-policy-forward access from the main eval pass and are deferred to Task 1.4.b; their slots stay at sentinel 0.0 in the meantime per `pearl_first_observation_bootstrap.md`. **Phase 1.4 partial lands kernels + 4 free-function launchers only**; consumer wire-up (per-eval-pass launches in `gpu_backtest_evaluator.rs` and HEALTH_DIAG `baseline_deltas` emit) is deferred to a follow-up commit per `feedback_no_partial_refactor.md` — kernel + launchers verify in isolation first via the 3 GPU oracle tests below, mirroring the Phase 1.1 + 1.2 + 1.3 atomic pattern. No state-reset registry entries are added: these slots are read-only outputs of an idempotent per-eval-pass kernel (each launch fully overwrites its own slot from the input price/spread arrays — no fold-stateful EMA); the slot-0 sentinel from constructor zero-fill is the correct cold-start value because the consumer reads each slot only after a freshly-completed launch in the same eval pass. Touched: `crates/ml/src/cuda_pipeline/baseline_kernels.cu` (new — 4 kernels + templated helper + 4 functor structs), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_BASELINE_KERNELS_CUBIN` cubin slot shared by all 4 launchers + 4 `pub fn launch_sp15_baseline_*` free-function launchers all loading the same cubin and resolving different `get_function()` symbols — mirrors `launch_sp15_dd_state`'s `stream.context().load_cubin` precedent), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+3 GPU oracle tests inside the existing `mod gpu` block: `baseline_buyhold_positive_on_drift` validates buyhold sharpe ∈ (0.2, 1.0) on a synthetic +0.5/bar drift series with ±1 noise (sharpe≈0.5 expected); `baseline_hold_only_emits_zero` validates `|sharpe| < 1e-5` because no-trade trajectory has mean=std=0 and the kernel's `(std > 0) ? mean / std : 0` guard returns 0 (the spec-mandated sentinel for an undefined sharpe); `baseline_momentum_reversion_symmetry` validates `|momentum_sharpe + reversion_sharpe| < 0.5` on a deliberately mean-reverting AR(1) series — the sign-flipped policies anti-correlate, the bound 0.5 absorbs cost-asymmetry from action-change bars hitting on different bars). All 3 tests pass on local RTX 3050 Ti (sm_86) in 1.93s. cargo test -p ml --lib --features cuda: 946 passed / 13 failed — same 13 failures pre-existing on the parent commit `9e8460248` (Task 1.3 baseline); zero introduced by this commit. **Deviation from task spec**: per-eval-pass production launches + HEALTH_DIAG `baseline_deltas` emit (steps 7+8 of the task brief) are NOT included this commit — the `feedback_no_partial_refactor.md` precedent set by Tasks 1.1 + 1.2 + 1.3 (all explicitly defer consumer migration as its own atomic follow-up) takes precedence over the brief's wire-up steps; consumer wiring is a load-bearing change that touches the eval pass and HEALTH_DIAG schema and must verify in isolation first. Hard rules: `feedback_no_atomicadd` (block-tree-reduce only via shared-memory `reduce_buf[256]`), `feedback_no_partial_refactor` (kernels + launchers land atomically; consumer wire-up follows as its own atomic commit), `feedback_no_stubs` (all 4 launchers return `Result<(), MLError>` and are fully functional, not placeholders — the trunk-shared baselines are explicitly out-of-scope, NOT stubbed), `feedback_no_htod_htoh_only_mapped_pinned` (oracle tests use `MappedF32Buffer` for prices/half_spread/ofi/isv buffers), `feedback_no_hiding` (the 4 deferred trunk-shared baselines are documented as out-of-scope here, not silently zero-stubbed — slots 411/413/414/415 remain at constructor-zero sentinel until Task 1.4.b lands them with the eval-side wire-up).
+
 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.3 — drawdown reporting kernel (2026-05-06): single GPU kernel `dd_state_kernel.cu` per-step state machine (single-thread, single-block) that lifts the truth source for current/max drawdown, recovery-bar counter, persistence counter, calmar denominator, and dd_pct from the legacy host-side composer onto the GPU. Reads existing `PS_PEAK_EQUITY` (slot 7) and `PS_PREV_EQUITY` (slot 9) from the position state buffer (canonical equity tracking — NO new ISV equity slot, invariant verified during plan v2 critical review per spec §6.3). Writes 6 ISV slots: `DD_CURRENT_INDEX=401` (per-step `max(0, (peak − equity) / peak)`), `DD_MAX_INDEX=402` (running max within fold), `DD_RECOVERY_BARS_INDEX=403` (bars since last new high-water mark; resets in-flight when a fresh peak fires), `DD_PERSISTENCE_INDEX=404` (twin counter to DD_RECOVERY_BARS in this phase; split-lifetime semantics may diverge in a later phase), `CALMAR_INDEX=405` (floored max-DD denominator `max(dd_max, 1e-4)` for the host composer's `calmar = mean_pnl / value-here` — the 1e-4 floor eliminates the saturation-at-100 artifact previously seen in train-dd4xl HEALTH_DIAG when max-DD ≈ 0), `DD_PCT_INDEX=406` (`clip(current_dd / max(dd_budget, 1e-4), 0, 1)`). **Phase 1.3 lands kernel + free-function launcher + state-reset registry entries + dispatch arms only**; per-step production wiring in `training_loop.rs` and the host-side calmar composer + HEALTH_DIAG emit are deferred to a follow-up commit per `feedback_no_partial_refactor.md` — the kernel + registry contract land atomically and the kernel is verified-working in isolation first via the GPU oracle test below, mirroring the Phase 1.1 + 1.2 atomic pattern. Touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (new), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_STATE_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_state` free-function launcher mirroring `launch_sp15_cost_net_sharpe`'s `stream.context().load_cubin` precedent — single-thread/single-block grid since the kernel is a per-step state machine, not a reduction), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+6 `RegistryEntry` records — `sp15_dd_current` / `sp15_dd_max` / `sp15_dd_recovery_bars` / `sp15_dd_persistence` / `sp15_dd_pct` are all FoldReset sentinel 0 stateful kernel outputs per `pearl_first_observation_bootstrap.md`; `sp15_calmar` is FoldReset sentinel `1e-4` — the SAME floor value the kernel writes — so the very first calmar division at fold start uses the floor rather than ±inf which would re-trip the saturation guard), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+6 dispatch arms for the 6 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 `dd_state_kernel_tracks_drawdown_correctly` validates a 6-step synthetic equity curve [+100, +10, −20, +15, −10, +5] → peak 110 at step 2, max-DD ≈ 0.182 at step 3 ((110−90)/110), never recovers, current_dd ≈ 0.091 at step 6 with recovery_bars=4, dd_pct ≈ 0.45 ∈ (0, 1] for `dd_budget=0.20`). **Deviation from task spec**: per-step production wire-up + HEALTH_DIAG emit (steps 7+8 of the task brief) are NOT included this commit — the `feedback_no_partial_refactor.md` precedent set by Tasks 1.1 + 1.2 (both explicitly defer consumer migration as its own atomic follow-up) takes precedence over the brief's wire-up steps; consumer migration is a load-bearing change and the kernel must verify in isolation before integration. Hard rules: `feedback_no_atomicadd` (no reductions; pure per-step state-machine), `feedback_no_partial_refactor` (kernel + launcher + 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 both ISV bus and pos_state buffer), `feedback_isv_for_adaptive_bounds` (calmar floor 1e-4 is a constant Invariant-1 anchor — fixed structural protection, NOT an adaptive EMA), `feedback_cpu_is_read_only` (truth source for drawdown state lifted from legacy host-side composer onto the GPU; calmar division remains host-side composition reading the kernel's denominator).