diff --git a/crates/ml-alpha/Cargo.toml b/crates/ml-alpha/Cargo.toml index bfa8442eb..ae65cef2c 100644 --- a/crates/ml-alpha/Cargo.toml +++ b/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 } diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index b0dcab0a2..48019c32c 100644 --- a/crates/ml-alpha/build.rs +++ b/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 { 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")) } diff --git a/crates/ml-alpha/cuda/adamw_step.cu b/crates/ml-alpha/cuda/adamw_step.cu new file mode 100644 index 000000000..7cee1bc75 --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu b/crates/ml-alpha/cuda/bce_loss_multi_horizon.cu new file mode 100644 index 000000000..5ce435dd2 --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/cuda/cfc_step.cu b/crates/ml-alpha/cuda/cfc_step.cu new file mode 100644 index 000000000..11508d365 --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/cuda/multi_horizon_heads.cu b/crates/ml-alpha/cuda/multi_horizon_heads.cu new file mode 100644 index 000000000..9b543b682 --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/cuda/projection.cu b/crates/ml-alpha/cuda/projection.cu new file mode 100644 index 000000000..263b2ad25 --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/cuda/snap_feature_assemble.cu b/crates/ml-alpha/cuda/snap_feature_assemble.cu new file mode 100644 index 000000000..fb98190cb --- /dev/null +++ b/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() {} diff --git a/crates/ml-alpha/src/lib.rs b/crates/ml-alpha/src/lib.rs index 751e0f78e..5860df9a8 100644 --- a/crates/ml-alpha/src/lib.rs +++ b/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; diff --git a/crates/ml-alpha/src/pinned_mem.rs b/crates/ml-alpha/src/pinned_mem.rs new file mode 100644 index 000000000..b67d5f546 --- /dev/null +++ b/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 { + let num_bytes = len * std::mem::size_of::(); + 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 { + 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); + } + } +} diff --git a/docs/superpowers/plans/2026-05-16-ml-alpha-phase-a-api-addendum.md b/docs/superpowers/plans/2026-05-16-ml-alpha-phase-a-api-addendum.md index 60e4b130e..2c74b1915 100644 --- a/docs/superpowers/plans/2026-05-16-ml-alpha-phase-a-api-addendum.md +++ b/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; ```