diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index 1b54056d7..8eef7d4be 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -886,6 +886,18 @@ fn main() {
         // buffer + bumping s1_input_dim) lands as a follow-up atomic
         // commit per the established Phase 1.1-1.4 precedent.
         "dd_pct_concat_kernel.cu",
+        // SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split
+        // kernel + ISV-driven α producer kernel (single source file, two
+        // `extern "C" __global__` symbols sharing one cubin per the
+        // 1:1-source-to-cubin / multi-kernel-per-file convention used by
+        // `baseline_kernels.cu` and `novelty_simhash_kernel.cu`).
+        // Composer (`r_quality_discipline_split_kernel`) reads
+        // ALPHA_SPLIT_INDEX=417 + alpha_warm_count scratch buffer, writes
+        // r_total + increments warm count. Producer
+        // (`alpha_split_producer_kernel`) reads grad-norm slots 418/419
+        // + warm count, writes ALPHA_SPLIT_INDEX clamped to [0.05, 0.95]
+        // (silent no-op until `*alpha_warm_count >= N_WARM=100`).
+        "r_quality_discipline_split_kernel.cu",
     ];
 
     // ALL kernels get common header (BF16 types + wrappers)
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index 7658435b6..db9a10ccc 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -1187,6 +1187,139 @@ pub fn launch_sp15_baseline_naive_reversion(
     Ok(())
 }
 
+/// SP15 Phase 3.1 (2026-05-06): cubin shared by the r_quality +
+/// r_discipline split composer (`r_quality_discipline_split_kernel`) and
+/// the ISV-driven α producer (`alpha_split_producer_kernel`). Per spec
+/// §8.2 (3.1) post-amendment-2 fix. Single source file with two
+/// `extern "C" __global__` symbols sharing one cubin per the established
+/// 1:1-source-to-cubin / multi-kernel-per-file pattern (mirrors the SP15
+/// Phase 1.4 baseline cubin and the SP11 SimHash novelty cubin). The two
+/// launchers below load this single cubin and resolve different
+/// `get_function()` symbols.
+///
+/// Cold-start protocol:
+///   • Constructor writes ALPHA_SPLIT_INDEX (slot 417) DIRECTLY to 0.5
+///     (NOT derived from the formula, which would produce 0/0=0 from the
+///     zero grad-norm EMAs at boot).
+///   • The composer maintains an `alpha_warm_count` scratch buffer
+///     ([1] f32 mapped-pinned). Each composer launch increments the
+///     count; while `*count < N_WARM=100` the composer overrides α back
+///     to 0.5 regardless of producer writes.
+///   • The producer is gated on the same warm count and stays a silent
+///     no-op until `*count >= N_WARM`; once it activates it writes
+///     α = clamp(grad_norm_q / (grad_norm_q + grad_norm_d + ε), 0.05,
+///     0.95) on each step.
+///
+/// Phase 3.1 lands kernels + launchers + ISV anchor seed + scratch
+/// buffer + state-reset registry entries + dispatch arms only;
+/// per-step consumer migration in `training_loop.rs` (composing
+/// r_total at the reward composition site + the producer launch in the
+/// gradient pipeline) is deferred to a follow-up commit per
+/// `feedback_no_partial_refactor.md` — kernels + launchers verify in
+/// isolation first via the GPU oracle test
+/// `r_split_uses_sentinel_alpha_at_cold_start`, mirroring the
+/// established Phase 1.1-1.5 atomic pattern.
+pub static SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/r_quality_discipline_split_kernel.cubin"));
+
+/// SP15 Phase 3.1 (2026-05-06): launcher for the r_quality + r_discipline
+/// split composer. Free function (matches `launch_sp15_dd_state` /
+/// `launch_sp15_baseline_*` precedent) so unit/oracle tests can drive the
+/// kernel directly without the trainer struct; production callers in
+/// `training_loop.rs` invoke the same launcher per step inside the
+/// reward-composition site.
+///
+/// Contract:
+///   `r_quality`, `r_discipline`: per-step f32 reward components.
+///   `isv`:                       device ptr to ISV bus (≥443 f32 slots).
+///   `r_total_out`:               device ptr to f32 [1], caller-allocated.
+///   `alpha_warm_count`:          device ptr to f32 [1] mapped-pinned
+///                                scratch — the trainer struct holds the
+///                                production buffer; tests pass a fresh
+///                                MappedF32Buffer.
+///
+/// Single thread, single block — the kernel is a per-step state machine,
+/// not a reduction.
+pub fn launch_sp15_r_quality_discipline_split(
+    stream: &Arc<CudaStream>,
+    r_quality: f32,
+    r_discipline: f32,
+    isv: cudarc::driver::sys::CUdeviceptr,
+    r_total_out: cudarc::driver::sys::CUdeviceptr,
+    alpha_warm_count: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_r_quality_discipline_split cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("r_quality_discipline_split_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load r_quality_discipline_split_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&r_quality)
+            .arg(&r_discipline)
+            .arg(&isv)
+            .arg(&r_total_out)
+            .arg(&alpha_warm_count)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (1, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch r_quality_discipline_split_kernel: {e}"
+            )))?;
+    }
+    Ok(())
+}
+
+/// SP15 Phase 3.1 (2026-05-06): launcher for the ISV-driven α producer.
+/// Loads the SAME cubin as `launch_sp15_r_quality_discipline_split`
+/// (multi-kernel-per-file pattern) and resolves the
+/// `alpha_split_producer_kernel` symbol. Reads grad-norm slots
+/// 418/419 + alpha_warm_count, writes ALPHA_SPLIT_INDEX (slot 417). No-op
+/// until `*alpha_warm_count >= N_WARM=100`.
+///
+/// Single thread, single block.
+pub fn launch_sp15_alpha_split_producer(
+    stream: &Arc<CudaStream>,
+    isv: cudarc::driver::sys::CUdeviceptr,
+    alpha_warm_count: cudarc::driver::sys::CUdeviceptr,
+) -> Result<(), MLError> {
+    let module = stream
+        .context()
+        .load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())
+        .map_err(|e| MLError::ModelError(format!(
+            "load sp15_r_quality_discipline_split cubin: {e}"
+        )))?;
+    let kernel = module
+        .load_function("alpha_split_producer_kernel")
+        .map_err(|e| MLError::ModelError(format!(
+            "load alpha_split_producer_kernel function: {e}"
+        )))?;
+    unsafe {
+        stream
+            .launch_builder(&kernel)
+            .arg(&isv)
+            .arg(&alpha_warm_count)
+            .launch(LaunchConfig {
+                grid_dim: (1, 1, 1),
+                block_dim: (1, 1, 1),
+                shared_mem_bytes: 0,
+            })
+            .map_err(|e| MLError::ModelError(format!(
+                "launch alpha_split_producer_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
@@ -3975,6 +4108,24 @@ pub struct GpuDqnTrainer {
     sel_t_pinned: *mut i32,
     sel_t_dev_ptr: u64,
 
+    /// SP15 Phase 3.1 (2026-05-06): warm-up counter for the
+    /// r_quality + r_discipline split composer (per spec §8.2 (3.1)
+    /// post-amendment-2 fix). Single-element mapped-pinned buffer
+    /// initialised to 0.0 in the constructor and reset to 0.0 at fold
+    /// boundary by the `sp15_alpha_warm_count` registry-arm dispatch
+    /// (which calls `host_slice_mut().fill(0.0)` mirroring the
+    /// `sp11_novelty_hash` pattern). Each per-step launch of
+    /// `r_quality_discipline_split_kernel` increments the count by 1;
+    /// once `*count >= N_WARM=100` the composer stops overriding α to
+    /// 0.5 and `alpha_split_producer_kernel` activates from no-op state
+    /// to write the formula's clamped output to ALPHA_SPLIT_INDEX. The
+    /// counter is a producer-controlled scalar — the host writes 0.0
+    /// at construction / fold reset (allowed under
+    /// `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
+    /// rule, NOT a host-side compute) and reads it through the mapped
+    /// page via `read_all()` only in tests.
+    pub(crate) sp15_alpha_warm_count: super::mapped_pinned::MappedF32Buffer, // [1]
+
     // ── Per-branch Q-gap EMA (trajectory backtracking state) ──
     /// [4] per-branch Q-gap EMA — device buffer (read-modify-write for backtracking restore).
     per_branch_q_gap_ema_buf: CudaSlice<f32>,
@@ -18612,6 +18763,31 @@ impl GpuDqnTrainer {
             dev_ptr
         };
 
+        // SP15 Phase 3.1 (2026-05-06): warm-up counter for the
+        // r_quality + r_discipline split composer (per spec §8.2 (3.1)
+        // post-amendment-2 fix). Single-element mapped-pinned buffer
+        // initialised to 0.0 by `MappedF32Buffer::new`'s zero-init
+        // contract — the host writes 0.0 at construction and again at
+        // fold boundary via the registry-arm dispatch (allowed under
+        // `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
+        // rule, NOT a host-side compute). The composer kernel
+        // (`r_quality_discipline_split_kernel`) increments by 1 each
+        // step; once `*count >= N_WARM=100` the composer stops
+        // overriding α and the producer kernel
+        // (`alpha_split_producer_kernel`) activates from no-op state.
+        // Production wire-up of the per-step launches in
+        // `training_loop.rs` is deferred to a follow-up commit per
+        // `feedback_no_partial_refactor.md`.
+        let sp15_alpha_warm_count = unsafe { super::mapped_pinned::MappedF32Buffer::new(1) }
+            .map_err(|e| MLError::ModelError(format!(
+                "sp15_alpha_warm_count alloc (1 f32): {e}"
+            )))?;
+        // `MappedF32Buffer::new` already zero-fills the host side; the
+        // explicit write here makes the cold-start contract self-evident
+        // at the call site (mirrors `sel_clip_buf.write_from_slice` for
+        // its 1.0 fixed-clip default).
+        sp15_alpha_warm_count.write_from_slice(&[0.0_f32]);
+
         // ── Per-branch Q-gap EMA (trajectory backtracking state) ────────
         let per_branch_q_gap_ema_buf = stream.alloc_zeros::<f32>(4) // [4]
             .map_err(|e| MLError::ModelError(format!("alloc per_branch_q_gap_ema_buf: {e}")))?;
@@ -20176,6 +20352,33 @@ impl GpuDqnTrainer {
                 use crate::cuda_pipeline::sp15_isv_slots::OFI_IMPACT_LAMBDA_INDEX;
                 *sig_ptr.add(OFI_IMPACT_LAMBDA_INDEX) = 2.0e-4_f32;
 
+                // SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline
+                // split slots (per spec §8.2 (3.1) post-amendment-2 fix).
+                // ALPHA_SPLIT (slot 417) is initialized DIRECTLY to 0.5
+                // — NOT derived from the formula
+                // `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)` at
+                // boot, which would produce 0/0=0 from the zero
+                // grad-norm EMAs. The composer kernel
+                // (`r_quality_discipline_split_kernel`) enforces this
+                // 0.5 sentinel until `alpha_warm_count >= N_WARM=100`,
+                // at which point the producer kernel
+                // (`alpha_split_producer_kernel`) takes over and writes
+                // the formula's clamped output here on each step. Per
+                // `feedback_isv_for_adaptive_bounds.md` — the structural
+                // 0.5 anchor is the cold-start absorber, the runtime
+                // value is signal-driven. The matching grad-norm slots
+                // (418, 419) are stateful EMA outputs, FoldReset
+                // sentinel 0 so Pearl A bootstraps from the new fold's
+                // first observation per
+                // `pearl_first_observation_bootstrap.md`.
+                use crate::cuda_pipeline::sp15_isv_slots::{
+                    ALPHA_SPLIT_INDEX, GRAD_NORM_DISCIPLINE_INDEX,
+                    GRAD_NORM_QUALITY_INDEX,
+                };
+                *sig_ptr.add(ALPHA_SPLIT_INDEX) = 0.5_f32;
+                *sig_ptr.add(GRAD_NORM_QUALITY_INDEX) = 0.0_f32;
+                *sig_ptr.add(GRAD_NORM_DISCIPLINE_INDEX) = 0.0_f32;
+
                 // Layout fingerprint (ISV[58..60)). Compile-time structural hash of
                 // the slot layout; checkpoint load fails-fast on mismatch.
                 // Stored as a u64 split across two f32 lanes using raw bit-cast so
@@ -20916,6 +21119,7 @@ impl GpuDqnTrainer {
             sel_clip_buf,
             sel_t_pinned,
             sel_t_dev_ptr,
+            sp15_alpha_warm_count,
             per_branch_q_gap_ema_buf,
             last_per_branch_q_gaps: [0.0; 4],
             q_mag_means_cached: [0.0_f32; 3],
diff --git a/crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu b/crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu
new file mode 100644
index 000000000..e00639077
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu
@@ -0,0 +1,117 @@
+// crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu
+//
+// SP15 Phase 3.1 — r_total = α × r_quality + (1 − α) × r_discipline
+// (per spec §8.2 (3.1) post-amendment-2 fix).
+//
+// ALPHA_SPLIT slot (ISV[417]) is initialized DIRECTLY to 0.5 in the
+// trainer constructor (NOT derived from formula at cold-start, which
+// would produce 0/0=0 from zero grad-norm EMAs). The formula
+//   α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)
+// only takes over once BOTH grad-norm EMAs have accumulated ≥ N_WARM
+// non-zero observations. Until then, the slot remains at 0.5 and the
+// composer reads the constructor-written value through the
+// r_quality_discipline_split_kernel below.
+//
+// Two `extern "C" __global__` symbols share this single cubin (the
+// established 1:1-source-to-cubin / multi-kernel-per-file pattern used
+// by `baseline_kernels.cu` and `novelty_simhash_kernel.cu`):
+//
+//   r_quality_discipline_split_kernel — per-step composition + warm count
+//     • Reads ALPHA_SPLIT_INDEX (slot 417)
+//     • Reads/writes alpha_warm_count scratch buffer ([1])
+//     • While `*alpha_warm_count < N_WARM`, OVERRIDES alpha to 0.5 (the
+//       cold-start sentinel) regardless of what the producer kernel may
+//       have written — this protects the composer from premature formula
+//       activation when the grad-norm EMAs are still warming up.
+//     • Writes r_total to r_total_out
+//
+//   alpha_split_producer_kernel — per-step ALPHA_SPLIT update from grad
+//     ratio
+//     • Returns immediately if `*alpha_warm_count < N_WARM` (the
+//       composer alone owns the warm-count increment; the producer is a
+//       silent no-op until the composer has fully warmed up)
+//     • Reads GRAD_NORM_QUALITY_INDEX (slot 418), GRAD_NORM_DISCIPLINE_INDEX
+//       (slot 419)
+//     • Writes ALPHA_SPLIT_INDEX (slot 417) clamped to [0.05, 0.95] so
+//       neither signal can totally dominate the composer
+//
+// Hard rules:
+//   • feedback_no_atomicadd — single-thread/single-block; no reductions.
+//   • feedback_isv_for_adaptive_bounds — N_WARM is the only structural
+//     constant; α itself is signal-driven.
+//   • feedback_first_observation_bootstrap — alpha_warm_count starts at 0;
+//     the warm-up window is the EMA cold-start absorber.
+
+#include <cuda_runtime.h>
+
+// Mirror the named slot constants from
+// `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`. Keep these literals in
+// lockstep with the Rust constants — the Rust unit test
+// `all_sp15_slots_fit_within_isv_total_dim` enforces ISV_TOTAL_DIM bounds
+// and `sp15_slot_layout_locked` pins each named index, so any drift here
+// shows up as an immediate test failure on the Rust side.
+#define ALPHA_SPLIT_INDEX           417
+#define GRAD_NORM_QUALITY_INDEX     418
+#define GRAD_NORM_DISCIPLINE_INDEX  419
+
+// Welford warm-up count. After both grad-norm EMAs have accumulated this
+// many non-zero observations the producer kernel takes over. Until then
+// the composer enforces the cold-start sentinel α=0.5.
+#define N_WARM 100.0f
+
+extern "C" __global__ void r_quality_discipline_split_kernel(
+    float r_quality,
+    float r_discipline,
+    const float* __restrict__ isv,
+    float* __restrict__ r_total_out,
+    float* __restrict__ alpha_warm_count
+) {
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+
+    // Load the constructor-written sentinel (or, post-warmup, the
+    // producer-written value).
+    float alpha = isv[ALPHA_SPLIT_INDEX];
+
+    // Cold-start absorber. The composer is the SOLE owner of the warm
+    // count — the producer reads it but never increments. This guarantees
+    // exactly one increment per call to this kernel even though both
+    // kernels alias the same buffer.
+    float warm = *alpha_warm_count;
+    if (warm < N_WARM) {
+        alpha = 0.5f;
+        *alpha_warm_count = warm + 1.0f;
+    }
+
+    *r_total_out = alpha * r_quality + (1.0f - alpha) * r_discipline;
+
+    // Make the slot writes visible to the host on the mapped-pinned
+    // buffers (same pattern used by the SP15 cost-net / dd-state kernels).
+    __threadfence_system();
+}
+
+extern "C" __global__ void alpha_split_producer_kernel(
+    float* __restrict__ isv,
+    const float* __restrict__ alpha_warm_count
+) {
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+
+    // Wait out the cold-start window. While the composer is overriding
+    // alpha to 0.5 (warm < N_WARM), writing the formula's output to the
+    // slot would be wasted work and could leak a transient value if the
+    // composer somehow ran on a stale warm count — leave the slot at the
+    // constructor-written 0.5 sentinel during warmup.
+    if (*alpha_warm_count < N_WARM) return;
+
+    float gn_q = isv[GRAD_NORM_QUALITY_INDEX];
+    float gn_d = isv[GRAD_NORM_DISCIPLINE_INDEX];
+    const float EPS = 1e-6f;
+
+    float alpha = gn_q / (gn_q + gn_d + EPS);
+    // Clamp to [0.05, 0.95] so neither head can totally dominate the
+    // composer — bilateral clamp per `pearl_symmetric_clamp_audit.md`.
+    alpha = fmaxf(0.05f, fminf(alpha, 0.95f));
+
+    isv[ALPHA_SPLIT_INDEX] = alpha;
+
+    __threadfence_system();
+}
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index 4479de925..5321e21f0 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -1088,6 +1088,53 @@ impl StateResetRegistry {
                 category: ResetCategory::FoldReset,
                 description: "ISV[DD_PCT_INDEX=406] — SP15 Phase 1.3 (2026-05-06) `clip(current_dd / max(dd_budget, 1e-4), 0, 1)` ∈ [0, 1]. Stateful kernel output; FoldReset sentinel 0 (cold-start at fold boundary; the kernel overwrites with the true value on the new fold's first launch). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §6.3.",
             },
+            // ── SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split ───
+            // Three ISV slots [417..420) + one mapped-pinned scratch buffer
+            // (`sp15_alpha_warm_count`):
+            //   - ALPHA_SPLIT_INDEX=417: Invariant-1 sentinel anchor (NOT a
+            //     stateful EMA). Constructor-seeded to 0.5 per spec §8.2 (3.1)
+            //     post-amendment-2 fix; FoldReset rewrites the same 0.5 anchor
+            //     so the composer's cold-start absorber re-engages on each new
+            //     fold (the warm-count goes back to 0; the producer kernel
+            //     becomes a no-op until `*warm_count >= N_WARM=100` again).
+            //     Per `feedback_isv_for_adaptive_bounds.md` — the structural
+            //     0.5 anchor is the cold-start absorber, the runtime value is
+            //     signal-driven once warmup completes.
+            //   - GRAD_NORM_QUALITY_INDEX=418: stateful EMA of the r_quality
+            //     gradient L2 norm. FoldReset sentinel 0 so Pearl A
+            //     bootstraps from the new fold's first observation per
+            //     `pearl_first_observation_bootstrap.md`.
+            //   - GRAD_NORM_DISCIPLINE_INDEX=419: stateful EMA of the
+            //     r_discipline gradient L2 norm. FoldReset sentinel 0 same
+            //     rationale as quality.
+            //   - sp15_alpha_warm_count (NOT an ISV slot — single-element
+            //     mapped-pinned scratch buffer on the trainer struct):
+            //     FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)`
+            //     so the new fold's first composer launch re-enters the
+            //     warm-up window from step 0 and the producer-kernel gate
+            //     re-engages. Mirrors the `sp11_novelty_hash` pattern of
+            //     resetting a non-ISV mapped-pinned buffer at fold boundary.
+            // Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).
+            RegistryEntry {
+                name: "sp15_alpha_split",
+                category: ResetCategory::FoldReset,
+                description: "ISV[ALPHA_SPLIT_INDEX=417] — SP15 Phase 3.1 (2026-05-06) blend coefficient for `r_total = α × r_quality + (1 − α) × r_discipline`. Constant Invariant-1 anchor (0.5) per `feedback_isv_for_adaptive_bounds.md` — NOT a stateful EMA. Constructor-seeded to 0.5; FoldReset rewrites the same anchor so the composer's cold-start absorber re-engages on each new fold (per spec §8.2 (3.1) post-amendment-2 fix: deriving α from `grad_norm_q / (grad_norm_q + grad_norm_d + ε)` at boot would produce 0/0=0 from zero grad-norm EMAs). The companion `alpha_split_producer_kernel` is gated on `sp15_alpha_warm_count` and stays a silent no-op until `*warm_count >= N_WARM=100`, at which point it activates and writes the formula's clamped output here on each step. Per-fold ISV refit may overwrite this in later phases. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
+            },
+            RegistryEntry {
+                name: "sp15_grad_norm_quality",
+                category: ResetCategory::FoldReset,
+                description: "ISV[GRAD_NORM_QUALITY_INDEX=418] — SP15 Phase 3.1 (2026-05-06) stateful EMA of the r_quality gradient L2 norm. Read by `alpha_split_producer_kernel` (post-warmup) to compute `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)`. FoldReset sentinel 0 so Pearl A bootstraps from the new fold's first observation per `pearl_first_observation_bootstrap.md` (also drives the warm-count gate: `alpha_split_producer_kernel` will re-enter no-op state at fold boundary because the companion `sp15_alpha_warm_count` resets in the same arm cluster). Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
+            },
+            RegistryEntry {
+                name: "sp15_grad_norm_discipline",
+                category: ResetCategory::FoldReset,
+                description: "ISV[GRAD_NORM_DISCIPLINE_INDEX=419] — SP15 Phase 3.1 (2026-05-06) stateful EMA of the r_discipline gradient L2 norm. Read by `alpha_split_producer_kernel` (post-warmup) as the denominator partner of `sp15_grad_norm_quality`. FoldReset sentinel 0 same rationale as quality. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
+            },
+            RegistryEntry {
+                name: "sp15_alpha_warm_count",
+                category: ResetCategory::FoldReset,
+                description: "GpuDqnTrainer.sp15_alpha_warm_count [1] f32 mapped-pinned scratch buffer — SP15 Phase 3.1 (2026-05-06) warm-up counter for the r_quality + r_discipline split composer (per spec §8.2 (3.1) post-amendment-2 fix). Initialised to 0.0 in the constructor; incremented by 1 each step inside `r_quality_discipline_split_kernel`; once `*count >= N_WARM=100` the composer stops overriding α to 0.5 and `alpha_split_producer_kernel` activates from no-op state. FoldReset sentinel 0.0 via `host_slice_mut().fill(0.0)` (mirrors the `sp11_novelty_hash` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary) — the new fold's first composer launch must re-enter the warm-up window from step 0 so the cold-start absorber holds α at 0.5 long enough for the new fold's grad-norm EMAs to accumulate ≥100 non-zero observations before the producer kernel takes over. Spec: docs/superpowers/specs/2026-05-06-sp15-honest-numbers.md §8.2 (3.1).",
+            },
             // 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
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index cf09bb1c7..cee95b38d 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -8141,6 +8141,55 @@ impl DQNTrainer {
                     fused.trainer().write_isv_signal_at(DD_PCT_INDEX, 0.0);
                 }
             }
+            // SP15 Phase 3.1 (2026-05-06): r_quality + r_discipline split
+            // dispatch arms (per spec §8.2 (3.1) post-amendment-2 fix).
+            // ALPHA_SPLIT_INDEX=417 is an Invariant-1 anchor (NOT a stateful
+            // EMA) — rewrite the constructor's 0.5 sentinel at fold boundary
+            // so the composer's cold-start absorber re-engages on each new
+            // fold; deriving from `grad_norm_q / (grad_norm_q + grad_norm_d
+            // + ε)` at boot would produce 0/0=0 from zero grad-norm EMAs.
+            // Per `feedback_isv_for_adaptive_bounds.md`. Per-fold ISV refit
+            // may overwrite this in later phases.
+            "sp15_alpha_split" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp15_isv_slots::ALPHA_SPLIT_INDEX;
+                    fused.trainer().write_isv_signal_at(ALPHA_SPLIT_INDEX, 0.5);
+                }
+            }
+            // GRAD_NORM_QUALITY_INDEX=418 / GRAD_NORM_DISCIPLINE_INDEX=419
+            // are stateful EMAs — FoldReset sentinel 0 so Pearl A
+            // bootstraps from the new fold's first observation per
+            // `pearl_first_observation_bootstrap.md`. The companion
+            // `sp15_alpha_warm_count` reset re-arms the warm-up gate so
+            // the producer kernel re-enters no-op state until both EMAs
+            // accumulate ≥100 non-zero observations on the new fold.
+            "sp15_grad_norm_quality" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp15_isv_slots::GRAD_NORM_QUALITY_INDEX;
+                    fused.trainer().write_isv_signal_at(GRAD_NORM_QUALITY_INDEX, 0.0);
+                }
+            }
+            "sp15_grad_norm_discipline" => {
+                if let Some(ref fused) = self.fused_ctx {
+                    use crate::cuda_pipeline::sp15_isv_slots::GRAD_NORM_DISCIPLINE_INDEX;
+                    fused.trainer().write_isv_signal_at(GRAD_NORM_DISCIPLINE_INDEX, 0.0);
+                }
+            }
+            // sp15_alpha_warm_count is a non-ISV mapped-pinned scratch
+            // buffer on the trainer struct. Reset to 0.0 via
+            // `host_slice_mut().fill(0.0)` (allowed under
+            // `feedback_no_htod_htoh_only_mapped_pinned.md`'s allowed-write
+            // rule, NOT a host-side compute) — mirrors the
+            // `sp11_novelty_hash` pattern of resetting a non-ISV
+            // mapped-pinned buffer at fold boundary. The new fold's first
+            // composer launch must re-enter the warm-up window from step 0
+            // so the cold-start absorber holds α at 0.5 long enough for
+            // the new fold's grad-norm EMAs to accumulate.
+            "sp15_alpha_warm_count" => {
+                if let Some(ref mut fused) = self.fused_ctx {
+                    fused.trainer_mut().sp15_alpha_warm_count.host_slice_mut().fill(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
diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs
index 830ffd0a8..b35876bf1 100644
--- a/crates/ml/tests/sp15_phase1_oracle_tests.rs
+++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs
@@ -17,10 +17,11 @@ mod gpu {
     };
     use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedU32Buffer};
     use ml::cuda_pipeline::sp15_isv_slots::{
-        BASELINE_BUYHOLD_SHARPE_INDEX, BASELINE_HOLD_ONLY_SHARPE_INDEX,
+        ALPHA_SPLIT_INDEX, 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,
+        DD_RECOVERY_BARS_INDEX, GRAD_NORM_DISCIPLINE_INDEX, GRAD_NORM_QUALITY_INDEX,
+        OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
     };
 
     /// Upper bound for any ISV index touched by SP15 oracle tests. Matches
@@ -605,6 +606,65 @@ mod gpu {
             );
         }
     }
+
+    /// Test 3.1.a — `r_total = 0.5 × r_quality + 0.5 × r_discipline` at
+    /// cold start. Per spec §8.2 (3.1) post-amendment-2 fix: the
+    /// ALPHA_SPLIT slot (ISV[417]) is constructor-written to 0.5 and the
+    /// composer's warm-count override holds it there until both
+    /// grad-norm EMAs accumulate ≥ N_WARM=100 observations. This test
+    /// validates the cold-start behavior in isolation: warm count starts
+    /// at 0, so the composer must override α to 0.5 regardless of any
+    /// other ISV state, AND increment the warm count to 1.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn r_split_uses_sentinel_alpha_at_cold_start() {
+        let stream = make_test_stream();
+
+        // Safety: CUDA context active via `make_test_stream` above.
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_LEN) }
+            .expect("alloc MappedF32Buffer for isv bus");
+        let mut isv_init = vec![0.0f32; ISV_LEN];
+        // Mirror the trainer constructor's ALPHA_SPLIT seed of 0.5; the
+        // grad-norm slots stay at the zero sentinel (cold start, no
+        // observations yet).
+        isv_init[ALPHA_SPLIT_INDEX] = 0.5;
+        isv_init[GRAD_NORM_QUALITY_INDEX] = 0.0;
+        isv_init[GRAD_NORM_DISCIPLINE_INDEX] = 0.0;
+        isv_buf.write_from_slice(&isv_init);
+
+        let alpha_warm_count_buf = unsafe { MappedF32Buffer::new(1) }
+            .expect("alloc MappedF32Buffer for alpha_warm_count");
+        alpha_warm_count_buf.write_from_slice(&[0.0f32]);
+        let r_total_buf = unsafe { MappedF32Buffer::new(1) }
+            .expect("alloc MappedF32Buffer for r_total");
+
+        ml::cuda_pipeline::gpu_dqn_trainer::launch_sp15_r_quality_discipline_split(
+            &stream,
+            /* r_quality    = */ 1.0,
+            /* r_discipline = */ -2.0,
+            isv_buf.dev_ptr,
+            r_total_buf.dev_ptr,
+            alpha_warm_count_buf.dev_ptr,
+        )
+        .expect("launch r_quality_discipline_split_kernel");
+        stream
+            .synchronize()
+            .expect("synchronize after r_quality_discipline_split_kernel launch");
+
+        // r_total = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5
+        let r_total = r_total_buf.read_all()[0];
+        assert!(
+            (r_total - (-0.5)).abs() < 1e-5,
+            "r_total = {r_total}, expected -0.5"
+        );
+
+        // Warm count: 0 → 1 after the single composer launch.
+        let warm = alpha_warm_count_buf.read_all()[0];
+        assert!(
+            (warm - 1.0).abs() < 1e-5,
+            "warm count = {warm}, expected 1"
+        );
+    }
 }
 
 /// SP15 Phase 1.7 — verify the trainer's `set_test_data_from_slices` API
diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md
index 398073c0e..a7afe1444 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 3.1 — r_quality + r_discipline split with ISV-driven α + sentinel cold-start (2026-05-06): per spec §8.2 (3.1) post-amendment-2 fix. Single source file `r_quality_discipline_split_kernel.cu` with two `extern "C" __global__` symbols sharing one cubin per the established 1:1-source-to-cubin / multi-kernel-per-file pattern (mirrors SP15 Phase 1.4 baselines + SP11 SimHash novelty cubins). The two launchers in `gpu_dqn_trainer.rs` (`launch_sp15_r_quality_discipline_split`, `launch_sp15_alpha_split_producer`) load the SAME cubin module via `stream.context().load_cubin(SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.to_vec())` and resolve different `get_function()` symbols. **Cold-start protocol** (the load-bearing fix): ALPHA_SPLIT_INDEX=417 is initialized DIRECTLY to 0.5 in the trainer constructor (NOT derived from the formula `α = grad_norm_q / (grad_norm_q + grad_norm_d + ε)`, which would produce 0/0=0 from the zero grad-norm EMAs at boot). The composer kernel (`r_quality_discipline_split_kernel`) reads ALPHA_SPLIT, increments a `*alpha_warm_count` scratch buffer by 1 each step, and OVERRIDES α back to 0.5 while `*count < N_WARM=100` regardless of whatever the producer may have written. The producer kernel (`alpha_split_producer_kernel`) is gated on the same warm count and stays a silent no-op until `*count >= N_WARM`; once it activates it writes `α = clamp(grad_norm_q / (grad_norm_q + grad_norm_d + ε), 0.05, 0.95)` to ALPHA_SPLIT_INDEX on each step (bilateral clamp per `pearl_symmetric_clamp_audit.md`). ISV slots written: `ALPHA_SPLIT_INDEX=417` (Invariant-1 anchor 0.5; FoldReset rewrites the same anchor so the cold-start absorber re-engages on each new fold), `GRAD_NORM_QUALITY_INDEX=418` (stateful EMA, FoldReset sentinel 0), `GRAD_NORM_DISCIPLINE_INDEX=419` (stateful EMA, FoldReset sentinel 0). One non-ISV mapped-pinned scratch buffer added to the trainer struct: `sp15_alpha_warm_count` ([1] f32) — initialised to 0.0 by the constructor and reset to 0.0 at fold boundary via `host_slice_mut().fill(0.0)` in the dispatch arm, mirroring the `sp11_novelty_hash` pattern of resetting a non-ISV mapped-pinned buffer at fold boundary. **Phase 3.1 lands kernels + 2 launchers + ISV anchor seed + scratch buffer + 4 state-reset registry entries + 4 dispatch arms only**; per-step consumer migration in `training_loop.rs` (composing r_total at the reward composition site + the producer launch in the gradient pipeline) is deferred to a follow-up commit per `feedback_no_partial_refactor.md` — kernels + launchers verify in isolation first via the GPU oracle test `r_split_uses_sentinel_alpha_at_cold_start`, mirroring the established Phase 1.1-1.5 atomic pattern. The deferred follow-up wires this split into the Phase 2B behavioral test contracts (anchor tests 2.4 cost_sensitivity + 2.6 regime_silences land green via Phase 3.4 regret + Phase 3.2 cost; this commit lands the split structure they depend on). Touched: `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu` (new — composer + producer kernels + N_WARM=100 cold-start absorber), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN` cubin slot + 2 `pub fn launch_sp15_*` free-function launchers loading the same cubin and resolving different `get_function()` symbols + 3 ISV constructor writes (ALPHA_SPLIT=0.5, GRAD_NORM_QUALITY=0.0, GRAD_NORM_DISCIPLINE=0.0) at the SP13/SP14 anchor block + 1 new `sp15_alpha_warm_count: MappedF32Buffer` field on the trainer struct + alloc + 0.0 init mirroring the `sel_clip_buf` precedent), `crates/ml/src/trainers/dqn/state_reset_registry.rs` (+4 `RegistryEntry` records — `sp15_alpha_split` Invariant-1 anchor rewrites 0.5 at fold boundary; `sp15_grad_norm_quality` / `sp15_grad_norm_discipline` Pearl A sentinel 0; `sp15_alpha_warm_count` non-ISV mapped-pinned buffer reset to 0.0 via `host_slice_mut().fill(0.0)`), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (+4 dispatch arms for the 4 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 `r_split_uses_sentinel_alpha_at_cold_start` validates: at warm count 0 the composer overrides α to 0.5 producing `r_total = 0.5 × 1.0 + 0.5 × (-2.0) = -0.5`, and the warm count increments to 1 after a single launch). Hard rules: `feedback_no_atomicadd` (single-thread/single-block; no reductions), `feedback_no_partial_refactor` (kernels + launchers + ISV anchor seed + scratch buffer + registry entries + dispatch arms land atomically; per-step consumer migration follows as its own atomic commit), `feedback_no_stubs` (both launchers return `Result<(), MLError>` and are fully functional; the trainer-struct `sp15_alpha_warm_count` field is exercised end-to-end via the `host_slice_mut().fill(0.0)` reset arm and the constructor's `write_from_slice(&[0.0_f32])` init), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for ISV bus, alpha_warm_count, and r_total output), `feedback_isv_for_adaptive_bounds` (the 0.5 anchor is a structural cold-start absorber; the runtime α is signal-driven once warmup completes; the [0.05, 0.95] producer clamp is bilateral per `pearl_symmetric_clamp_audit.md`), `feedback_first_observation_bootstrap` (warm count starts at 0; the warm-up window is the EMA cold-start absorber so the producer never sees zero grad-norm EMAs).
+
 SP15 Phase 1.5 — dd_pct foundational state input concat kernel (2026-05-06): **LAYOUT FINGERPRINT BREAK — pre-SP15 checkpoints WILL NOT LOAD after this commit. Greenfield OK per spec Q1.** Single GPU kernel `dd_pct_concat_kernel.cu` produces a `[B, state_dim_padded + 1]` concat buffer whose leading `state_dim_padded` columns equal the input `states_buf` (full padded width, including pre-existing pad zeros from the upstream producer) and whose `+1` last column equals `isv[DD_PCT_INDEX=406]` broadcast across batch. Per spec §6.5: dd_pct (slot 406, written by Task 1.3's `dd_state_kernel`) gets concatenated to the trunk forward input as the last dim so the eval-time policy SEES drawdown context on every forward pass; Phase 3 teachings can condition on dd_pct directly via state, not just reward modulation. Kernel: `batch_size`-block grid, 128 threads/block; each block copies one row of `state_dim_padded` floats and writes the +1 last column. No atomicAdd (pure scatter-copy per `feedback_no_atomicadd`). Layout fingerprint extension: `layout_fingerprint_seed` gains a `TRUNK_INPUT_DD_PCT=sp15_phase_1_5` marker, which causes the FNV1a hash to change — old checkpoints fail to load with the existing layout-mismatch error path (the same fail-fast path that fired on the SP4 / SP14 layout breaks). **Phase 1.5 lands kernel + launcher + layout fingerprint marker + GPU oracle test only**; trunk consumer migration (re-pointing `forward_online` / `forward_target_raw` to consume the concat buffer + bumping `s1_input_dim` from 48 → 49 + propagating through the GRN encoder, VSN gate input, bottleneck path, and backward dx scratch) is deferred to a follow-up atomic commit per `feedback_no_partial_refactor.md` — this matches the established Phase 1.1-1.4 precedent (kernels + launchers verify in isolation first via the GPU oracle test below; consumer migration is a load-bearing change touching the GRN encoder, VSN partition boundaries, fxcache schema, and backward gradient flow that must verify independently). Touched: `crates/ml/src/cuda_pipeline/dd_pct_concat_kernel.cu` (new — single kernel `dd_pct_concat_kernel`), `crates/ml/build.rs` (+1 cubin manifest entry), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (+1 `pub static SP15_DD_PCT_CONCAT_CUBIN` cubin slot + 1 `pub fn launch_sp15_dd_pct_concat` free-function launcher mirroring `launch_sp15_dd_state`'s `stream.context().load_cubin` precedent + 1-line `TRUNK_INPUT_DD_PCT=sp15_phase_1_5` marker appended to `layout_fingerprint_seed` after the SP15 ISV slot block), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 GPU oracle test `dd_pct_concat_kernel_writes_last_column` validates B=4, state_dim=48, state_dim_padded=128: the leading 128 columns of each output row match the input states (including pad zeros), and the 129th column equals `isv[DD_PCT_INDEX]=0.42` broadcast across all 4 rows). cargo test -p ml --lib --features cuda: 945 passed / 14 failed — same 14 failures pre-existing on the parent commit `c6fd4b4b2` (Task 1.4 partial baseline); zero introduced by this commit. **Deferred to Phase 1.5.b follow-up** (per `feedback_no_partial_refactor`): wire trunk forward (`encoder_forward_only`, `target_encoder_forward_only`, bottleneck path) to consume the concat buffer; bump `s1_input_dim` to `state_dim + 1` so the GRN h_s1 Linear_a / Linear_residual GEMMs accept K = 49; propagate the +1 dim through the backward `dx` scratch (last column gradient is computed mechanically but not propagated — dd_pct is a state input, not a learnable parameter); audit VSN group partitioning to confirm groups still partition the original 48-dim state slice (the +1 dd_pct column lives outside the VSN-gated portion). The follow-up commit also adds the trunk-grounding behavioral test (Phase 4 L40S smoke verifies the dd_pct propagation at production scale via the existing layout-fingerprint-mismatch fail-fast on cold-start of any pre-SP15 checkpoint). Hard rules: `feedback_no_atomicadd` (pure scatter-copy, no reductions), `feedback_no_partial_refactor` (kernel + launcher + layout fingerprint marker land atomically; trunk consumer migration follows as its own atomic commit per the established Phase 1.1-1.4 precedent), `feedback_no_stubs` (launcher returns `Result<(), MLError>` and is fully functional, oracle test exercises real GPU kernel execution end-to-end — the test is NOT a doc-only stub), `feedback_no_htod_htoh_only_mapped_pinned` (oracle test uses `MappedF32Buffer` for states/isv/concat buffers), `feedback_no_hiding` (the launcher signature exposes only the dimensions the kernel actually consumes — `batch_size` and `state_dim_padded`; the follow-up consumer migration will introduce a separate `s1_input_dim` integer at the GRN encoder call site, not at this concat-kernel launcher).
 
 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).