build(ml-alpha): multi-cubin build.rs + placeholder kernels + local MappedF32Buffer

Cargo.toml: drops gbdt; adds memmap2 + approx; keeps ml-core only
(cannot depend on ml: would cycle since ml depends on ml-alpha for
the Mamba2 gate baseline).

build.rs: compiles 7 cubins (mamba2_alpha + 6 new placeholders)
with -O3 --use_fast_math --ftz --fmad. Skips kernels whose source
isn't present yet so partial check-ins work. Every env::var paired
with rerun-if-env-changed per the canonical build pearl.

src/pinned_mem.rs: local copy of MappedF32Buffer (mirrors
ml::cuda_pipeline::mapped_pinned::MappedF32Buffer). Drives the only
permitted CPU<->GPU path per feedback_no_htod_htoh_only_mapped_pinned.
Eventually the move-to-ml-core refactor will deduplicate; out of
scope for the Phase A branch.

Addendum: updates the import path to ml_alpha::pinned_mem.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

crates/ml-alpha/Cargo.toml

@@ -3,20 +3,13 @@ name = "ml-alpha"
 version.workspace = true
 edition.workspace = true
 rust-version.workspace = true
-description = "FoxhuntQ-Δ Phase 1a — minimal alpha-only crate for cheapest-cost falsification of bar-resolution signal hypothesis"
+description = "CfC perception + multi-horizon alpha heads (Phase A foundation for CfC+PPO greenfield)"
 publish = false
 
-# Phase 1a discipline: minimal dependency footprint to keep iteration fast.
-# This crate exists to answer ONE question: does supervised alpha at the
-# imbalance-bar resolution exceed validation accuracy > 0.52 (binary direction)
-# on a purged walk-forward held-out fold? If yes, FoxhuntQ-Δ proceeds.
-# If no, the bar-resolution hypothesis from `project_bar_resolution_is_actual_architecture`
-# is confirmed and FoxhuntQ-Δ does not proceed.
-#
-# DESIGN INVARIANT: depends only on ml-core for GPU primitives + std + cudarc.
-# NO ml, NO ml-dqn, NO ml-supervised. This keeps the cold-compile under ~2 min
-# vs ~12 min for the main ml crate. If we ever NEED something from ml, refactor
-# it up to ml-core first.
+# Phase A scope: snapshot-level CfC trunk + 5 multi-horizon BCE heads,
+# pre-compiled CUDA cubins, GPU-resident weights. The hard validation gate
+# (CfC vs Mamba2 AUC at every horizon) sits at the end of Phase A. Phase B
+# (PPO) extends this crate; live integration lives in trading_agent_service.
 
 [features]
 default = ["cuda"]
@@ -24,6 +17,9 @@ cuda = []
 
 [dependencies]
 ml-core = { workspace = true }
+# NOTE: cannot depend on `ml` (would cycle — ml depends on ml-alpha for the
+# mamba2_block gate reference). MappedF32Buffer is duplicated locally in
+# `src/pinned_mem.rs`; eventually we should move mapped-pinned to ml-core.
 cudarc = { version = "0.19", default-features = false, features = ["driver", "cublas", "dynamic-linking", "std", "cuda-version-from-build-system", "f16"] }
 
 # Standard async + error + logging plumbing
@@ -36,22 +32,16 @@ serde_json.workspace = true
 thiserror.workspace = true
 clap = { workspace = true, features = ["derive"] }
 
-# Arrow IPC for fxcache reading (replaces custom binary mmap; eliminates the
-# off-by-8 schema-mismatch bug class by reading dims/version from the file's
-# embedded schema metadata rather than relying on compile-time constants).
+# Arrow IPC for fxcache reading (gate reference path).
 arrow = { workspace = true, features = ["ipc"] }
 
-# Deterministic RNG for shuffling, splits, weighted subsampling
+# Deterministic RNG for weight init, shuffling.
 rand = { workspace = true }
 rand_chacha = "0.3"
 
-# Pure-Rust gradient boosted decision trees, used as the GBM baseline for
-# Phase 1a falsification (per Grinsztajn et al. NeurIPS 2022, GBM is the
-# canonical baseline for ~74-dim engineered tabular features, not MLP).
-# Pure Rust = no system libLightGBM/libxgboost install. If results show
-# marginal signal worth chasing, can upgrade to `lightgbm3` (real LightGBM)
-# in Phase 1b.
-gbdt = "0.1.3"
+# Phase A data path: mmap predecoded MBP-10 sidecars.
+memmap2 = { workspace = true }
 
 [dev-dependencies]
-tempfile = "3.10"
+tempfile = { workspace = true }
+approx = { workspace = true }

crates/ml-alpha/build.rs

@@ -1,107 +1,104 @@
-//! ml-alpha build.rs — precompile the Mamba2 temporal SSM scan kernel.
+//! Pre-compile all ml-alpha CUDA kernels into arch-specific cubins.
 //!
-//! Compiles `crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu` (the only
-//! kernel ml-alpha needs from the parent ml crate's library, for Phase 1d.1's
-//! GPU-resident stateful encoder) into an arch-specific cubin in OUT_DIR.
-//!
-//! Discipline:
-//! - **rerun-if-env-changed=CUDA_COMPUTE_CAP** is mandatory: cargo target PVCs
-//!   are shared across L40S (sm_89) and H100 (sm_90) builds; without this,
-//!   stale cubins surface as `CUDA_ERROR_NO_BINARY_FOR_GPU` at runtime.
-//! - The kernel source has zero `#include`s, so no header prepending is
-//!   needed (unlike ml-core's autograd kernels).
-//! - On non-CUDA builds (`CARGO_FEATURE_CUDA` unset) we silently skip; the
-//!   library still compiles, only the Mamba2Block::new() runtime path fails
-//!   with a missing-cubin error.
+//! Per `feedback_no_nvrtc.md`: no runtime kernel compilation.
+//! Per `pearl_build_rs_rerun_if_env_changed.md`: every `std::env::var`
+//! is paired with `cargo:rerun-if-env-changed`.
 
 use std::path::{Path, PathBuf};
 use std::process::Command;
 
+const KERNELS: &[&str] = &[
+    "mamba2_alpha_kernel",          // gate reference (Phase A->Mamba2 baseline)
+    "snap_feature_assemble",
+    "cfc_step",
+    "multi_horizon_heads",
+    "projection",
+    "bce_loss_multi_horizon",
+    "adamw_step",
+];
+
 fn main() {
     println!("cargo:rerun-if-changed=build.rs");
 
+    println!("cargo:rerun-if-env-changed=CARGO_FEATURE_CUDA");
     if std::env::var("CARGO_FEATURE_CUDA").is_err() {
+        eprintln!("  ml-alpha: cuda feature disabled, skipping kernel build");
         return;
     }
 
-    let out_dir = PathBuf::from(std::env::var("OUT_DIR").unwrap());
-
-    // Kernel source lives in this crate's `cuda/` directory — purpose-built
-    // simplified scan for the alpha supervised path (no ISV, no
-    // temporal_weight, no NULL-pointer branches that the DQN trainer's
-    // shared kernel carries).
-    let kernel_src_path = Path::new("cuda/mamba2_alpha_kernel.cu")
-        .canonicalize()
-        .unwrap_or_else(|_| {
-            let workspace = std::env::var("CARGO_MANIFEST_DIR")
-                .map(PathBuf::from)
-                .unwrap_or_default();
-            workspace
-                .join("cuda/mamba2_alpha_kernel.cu")
-                .canonicalize()
-                .expect("Cannot locate cuda/mamba2_alpha_kernel.cu")
-        });
-    println!("cargo:rerun-if-changed={}", kernel_src_path.display());
-
     println!("cargo:rerun-if-env-changed=CUDA_COMPUTE_CAP");
-    let cuda_compute_cap = std::env::var("CUDA_COMPUTE_CAP").unwrap_or_else(|_| "80".to_string());
-    let arch = format!("sm_{cuda_compute_cap}");
+    println!("cargo:rerun-if-env-changed=CUDA_HOME");
+    let cap = std::env::var("CUDA_COMPUTE_CAP").unwrap_or_else(|_| "80".to_string());
+    let arch = format!("sm_{cap}");
 
     let nvcc = match find_nvcc() {
         Some(p) => p,
         None => {
-            eprintln!("  ml-alpha: nvcc not found, skipping mamba2 kernel precompilation");
-            eprintln!("            Install CUDA toolkit or set CUDA_HOME for GPU builds");
+            eprintln!("  ml-alpha: nvcc not found, skipping kernel build (set CUDA_HOME or install CUDA toolkit)");
             return;
         }
     };
 
-    let cubin_path = out_dir.join("mamba2_alpha_kernel.cubin");
+    let out = PathBuf::from(std::env::var("OUT_DIR").expect("OUT_DIR not set by cargo"));
 
-    let status = Command::new(&nvcc)
+    for k in KERNELS {
+        let src = PathBuf::from(format!("cuda/{k}.cu"));
+        if !src.exists() {
+            eprintln!("  ml-alpha: skipping {k} — source not yet present");
+            continue;
+        }
+        println!("cargo:rerun-if-changed={}", src.display());
+        let cubin = out.join(format!("{k}.cubin"));
+        compile(&nvcc, &src, &cubin, &arch);
+    }
+}
+
+fn compile(nvcc: &Path, src: &Path, cubin: &Path, arch: &str) {
+    let status = Command::new(nvcc)
         .args([
             "-cubin",
             &format!("-arch={arch}"),
             "-O3",
+            "--use_fast_math",
             "--ftz=true",
             "--fmad=true",
-            "--prec-div=true",
-            "--prec-sqrt=true",
             "-o",
-            cubin_path.to_str().unwrap(),
-            kernel_src_path.to_str().unwrap(),
+            cubin.to_str().unwrap(),
+            src.to_str().unwrap(),
         ])
-        .status();
-
-    match status {
-        Ok(s) if s.success() => {
-            eprintln!(
-                "  ml-alpha: Compiled mamba2_alpha_kernel.cu -> mamba2_alpha_kernel.cubin ({arch})"
-            );
-        }
-        Ok(s) => panic!(
-            "ml-alpha: nvcc failed to compile mamba2_alpha_kernel.cu (exit={})",
-            s.code().unwrap_or(-1)
-        ),
-        Err(e) => panic!("ml-alpha: nvcc invocation error: {e}"),
+        .status()
+        .unwrap_or_else(|e| panic!("nvcc spawn failed for {}: {e}", src.display()));
+    if !status.success() {
+        panic!(
+            "nvcc failed for {} (exit {})",
+            src.display(),
+            status.code().unwrap_or(-1)
+        );
     }
+    eprintln!(
+        "  ml-alpha: compiled {} -> {} ({arch})",
+        src.display(),
+        cubin.display()
+    );
 }
 
 fn find_nvcc() -> Option<PathBuf> {
     if let Ok(home) = std::env::var("CUDA_HOME") {
-        let nvcc = PathBuf::from(home).join("bin/nvcc");
-        if nvcc.exists() {
-            return Some(nvcc);
-        }
-    }
-    for path in &["/usr/local/cuda/bin/nvcc", "/usr/bin/nvcc"] {
-        let p = PathBuf::from(path);
+        let p = PathBuf::from(home).join("bin/nvcc");
         if p.exists() {
             return Some(p);
         }
     }
-    match Command::new("nvcc").arg("--version").output() {
-        Ok(output) if output.status.success() => Some(PathBuf::from("nvcc")),
-        _ => None,
+    for cand in ["/usr/local/cuda/bin/nvcc", "/usr/bin/nvcc"] {
+        let p = PathBuf::from(cand);
+        if p.exists() {
+            return Some(p);
+        }
     }
+    Command::new("nvcc")
+        .arg("--version")
+        .output()
+        .ok()
+        .filter(|o| o.status.success())
+        .map(|_| PathBuf::from("nvcc"))
 }

crates/ml-alpha/cuda/adamw_step.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void adamw_step_stub() {}

crates/ml-alpha/cuda/bce_loss_multi_horizon.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void bce_loss_multi_horizon_stub() {}

crates/ml-alpha/cuda/cfc_step.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void cfc_step_stub() {}

crates/ml-alpha/cuda/multi_horizon_heads.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void multi_horizon_heads_stub() {}

crates/ml-alpha/cuda/projection.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void projection_stub() {}

crates/ml-alpha/cuda/snap_feature_assemble.cu

@@ -0,0 +1,2 @@
+// placeholder — real implementation in the task adding this kernel
+extern "C" __global__ void snap_feature_assemble_stub() {}

crates/ml-alpha/src/lib.rs

@@ -20,6 +20,7 @@ pub mod cfc;
 pub mod heads;
 pub mod isv;
 pub mod pinned;
+pub mod pinned_mem;
 pub mod trainer;
 pub mod data;
 pub mod gate;

crates/ml-alpha/src/pinned_mem.rs

@@ -0,0 +1,110 @@
+//! Mapped-pinned f32 buffer (CPU+GPU visible via `cuMemHostAlloc(DEVICEMAP)`).
+//!
+//! This is the **only** permitted CPU↔GPU path per
+//! `feedback_no_htod_htoh_only_mapped_pinned.md`. Mirror of
+//! `crates/ml/src/cuda_pipeline/mapped_pinned.rs::MappedF32Buffer` —
+//! duplicated here because adding `ml-alpha -> ml` would cycle (ml
+//! depends on ml-alpha for the Mamba2 gate baseline). Move both to
+//! `ml-core::mapped_pinned` to deduplicate (out of scope for Phase A).
+
+#![allow(unsafe_code)]
+
+use std::ffi::c_void;
+use std::mem::MaybeUninit;
+
+/// CPU+GPU visible buffer of `f32`s allocated via
+/// `cuMemHostAlloc(DEVICEMAP|PORTABLE)`. The kernel writes via `dev_ptr`
+/// (with `__threadfence_system()` in the kernel); the CPU reads via
+/// `host_ptr` after stream synchronisation.
+#[allow(missing_debug_implementations)]
+pub struct MappedF32Buffer {
+    pub host_ptr: *mut f32,
+    pub dev_ptr: cudarc::driver::sys::CUdeviceptr,
+    pub len: usize,
+}
+
+// Safety: only accessed from the thread that constructed it (which holds
+// the active CUDA context). The host pointer is valid for the life of the
+// allocation.
+unsafe impl Send for MappedF32Buffer {}
+unsafe impl Sync for MappedF32Buffer {}
+
+impl MappedF32Buffer {
+    /// Allocate `len` f32s of mapped-pinned memory.
+    ///
+    /// # Safety
+    /// Caller must ensure a CUDA context is active on the current thread.
+    pub unsafe fn new(len: usize) -> Result<Self, String> {
+        let num_bytes = len * std::mem::size_of::<f32>();
+        let flags = cudarc::driver::sys::CU_MEMHOSTALLOC_DEVICEMAP
+            | cudarc::driver::sys::CU_MEMHOSTALLOC_PORTABLE;
+
+        let host_ptr = cudarc::driver::result::malloc_host(num_bytes, flags)
+            .map_err(|e| format!("MappedF32Buffer alloc ({len} f32): {e}"))? as *mut f32;
+
+        // Zero-init so a stale read returns a safe default (0.0).
+        std::ptr::write_bytes(host_ptr, 0, len);
+
+        let mut dev_ptr_raw = MaybeUninit::uninit();
+        cudarc::driver::sys::cuMemHostGetDevicePointer_v2(
+            dev_ptr_raw.as_mut_ptr(),
+            host_ptr as *mut c_void,
+            0,
+        )
+        .result()
+        .map_err(|e| format!("cuMemHostGetDevicePointer (f32 buf): {e}"))?;
+
+        Ok(Self {
+            host_ptr,
+            dev_ptr: dev_ptr_raw.assume_init(),
+            len,
+        })
+    }
+
+    /// Read all `len` entries via `read_volatile`. Caller must have
+    /// synchronised the producing stream first.
+    pub fn read_all(&self) -> Vec<f32> {
+        let mut out = Vec::with_capacity(self.len);
+        unsafe {
+            for i in 0..self.len {
+                out.push(std::ptr::read_volatile(self.host_ptr.add(i)));
+            }
+        }
+        out
+    }
+
+    /// Write CPU-side data into the buffer via `host_ptr`. The kernel
+    /// reading via `dev_ptr` sees the data after a stream sync barrier
+    /// (mapped-pinned coherence).
+    pub fn write_from_slice(&self, slice: &[f32]) {
+        assert!(
+            slice.len() <= self.len,
+            "MappedF32Buffer write overflow: have {} have {}",
+            slice.len(),
+            self.len
+        );
+        unsafe {
+            for (i, &v) in slice.iter().enumerate() {
+                std::ptr::write_volatile(self.host_ptr.add(i), v);
+            }
+        }
+    }
+
+    /// Mutable host slice over the mapped pages.
+    ///
+    /// # Safety
+    /// `host_ptr` is valid for `self.len` f32s by construction; the caller
+    /// must not invalidate the allocation while the slice is live and
+    /// must respect unique-mutable-borrow semantics.
+    pub fn host_slice_mut(&mut self) -> &mut [f32] {
+        unsafe { std::slice::from_raw_parts_mut(self.host_ptr, self.len) }
+    }
+}
+
+impl Drop for MappedF32Buffer {
+    fn drop(&mut self) {
+        unsafe {
+            let _ = cudarc::driver::result::free_host(self.host_ptr as *mut c_void);
+        }
+    }
+}

docs/superpowers/plans/2026-05-16-ml-alpha-phase-a-api-addendum.md

@@ -22,7 +22,7 @@ use cudarc::driver::{
     CudaContext, CudaStream, CudaModule, CudaFunction, CudaSlice,
     DevicePtr, DevicePtrMut, LaunchConfig, PushKernelArg, DriverError,
 };
-use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer;
+use crate::pinned_mem::MappedF32Buffer;  // local copy — see pinned_mem.rs header
 use ml_core::device::MlDevice;
 ```