diff --git a/Cargo.lock b/Cargo.lock index 36a9e18f4..9990fb4b4 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -111,6 +111,7 @@ name = "bloomery-engram" version = "0.1.0" dependencies = [ "bloomery-gguf", + "bloomery-os", "bloomery-threads", "libc", "memmap2", @@ -130,6 +131,7 @@ dependencies = [ name = "bloomery-gguf" version = "0.1.0" dependencies = [ + "bloomery-os", "bloomery-refset", "libc", "memmap2", @@ -324,6 +326,7 @@ dependencies = [ "bloomery-gguf", "bloomery-levers", "bloomery-models", + "bloomery-os", "bloomery-placement", "bloomery-qdot", "bloomery-refset", @@ -361,6 +364,13 @@ dependencies = [ "thiserror", ] +[[package]] +name = "bloomery-os" +version = "0.1.0" +dependencies = [ + "libc", +] + [[package]] name = "bloomery-placement" version = "0.1.0" @@ -432,6 +442,7 @@ name = "bloomery-threads" version = "0.1.0" dependencies = [ "bloomery-levers", + "bloomery-os", "libc", ] diff --git a/Cargo.toml b/Cargo.toml index f2898ca4a..d2a078b29 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -1,6 +1,6 @@ [workspace] resolver = "3" -members = ["crates/qdot", "crates/gguf", "crates/models", "crates/model", "crates/model-hf", "crates/placement", "crates/threads", "crates/levers", "crates/verbs", "crates/gpu", "crates/gpu-deepseek41", "crates/gpu-glm5next", "crates/gpu-mimo2", "crates/gpu-solar2", "crates/gpu-gemma-embedding2", "crates/gpu-gates", "crates/engram", "crates/engram-lab", "crates/sampler", "crates/tokenizer", "crates/serve", "crates/jinja", "crates/vision", "crates/gpu-vision", "crates/gpu-gemma4a", "crates/refset", "crates/runtime", "crates/app", "crates/decision", "crates/embed", "crates/hf", "crates/safetensors"] +members = ["crates/qdot", "crates/gguf", "crates/models", "crates/model", "crates/model-hf", "crates/placement", "crates/threads", "crates/levers", "crates/verbs", "crates/gpu", "crates/gpu-deepseek41", "crates/gpu-glm5next", "crates/gpu-mimo2", "crates/gpu-solar2", "crates/gpu-gemma-embedding2", "crates/gpu-gates", "crates/engram", "crates/engram-lab", "crates/sampler", "crates/tokenizer", "crates/serve", "crates/jinja", "crates/vision", "crates/gpu-vision", "crates/gpu-gemma4a", "crates/refset", "crates/runtime", "crates/app", "crates/decision", "crates/embed", "crates/hf", "crates/safetensors", "crates/os"] # 의도적으로 컴파일되지 않는 재현기다(MUL-7). 워크스페이스에 넣으면 --workspace가 영구히 빨갛다. exclude = ["crates/oxide-ice-unroll", "crates/oxide-ice-unroll/variants/unroll_smoke", "fuzz", "labs/hip-native"] @@ -16,6 +16,9 @@ cuda-core = { version = "0.4.0", git = "https://github.com/NVIDIA/cuda-rust.git" # one semver line unify into one version, so the gates cut the text handed to the driver with the loader's own code. oxide-artifacts = "0.2.1" libc = "0.2" +# The one owner of the operating-system calls the host chain makes: the crate +# that holds what Linux and macOS do differently. +bloomery-os = { path = "crates/os" } # The decision head's host sgemm (crates/decision); its own threading stays off. matrixmultiply = { version = "0.3.10", default-features = false, features = ["std"] } memmap2 = "0.9" diff --git a/crates/engram/Cargo.toml b/crates/engram/Cargo.toml index ededac692..356e75956 100644 --- a/crates/engram/Cargo.toml +++ b/crates/engram/Cargo.toml @@ -13,6 +13,7 @@ workspace = true [dependencies] bloomery-gguf = { path = "../gguf" } +bloomery-os.workspace = true bloomery-threads = { path = "../threads" } libc.workspace = true memmap2.workspace = true diff --git a/crates/engram/src/direct.rs b/crates/engram/src/direct.rs index ae58d96f8..beab6649c 100644 --- a/crates/engram/src/direct.rs +++ b/crates/engram/src/direct.rs @@ -33,15 +33,10 @@ pub const DIRECT_ALIGN: usize = 4096; /// are multiples of [`DIRECT_ALIGN`]; the kernel refuses any other with /// `EINVAL`. pub fn open_direct(path: &Path) -> Result { - use std::os::unix::fs::OpenOptionsExt; - std::fs::OpenOptions::new() - .read(true) - .custom_flags(libc::O_DIRECT) - .open(path) - .map_err(|source| EngramError::OpenDirect { - path: path.to_path_buf(), - source, - }) + os::open_direct(path).map_err(|source| EngramError::OpenDirect { + path: path.to_path_buf(), + source: source.into(), + }) } /// A file read without the page cache ([`DirectFile::open`]) or through it diff --git a/crates/engram/src/lib.rs b/crates/engram/src/lib.rs index d03175741..b2ae15886 100644 --- a/crates/engram/src/lib.rs +++ b/crates/engram/src/lib.rs @@ -540,21 +540,8 @@ impl Site { /// `posix_fadvise(DONTNEED)` over `at .. at + len` of `file`: the page-cache /// half of an eviction, after the mapping's own entries are dropped. fn fadvise_dontneed(file: &File, at: usize, len: usize) -> Result<(), std::io::Error> { - // SAFETY: the descriptor is `file`'s own, borrowed for the call, and the - // range is plain numbers: fadvise reads and writes no memory of ours. - let rc = unsafe { - libc::posix_fadvise( - file.as_raw_fd(), - fadvise_off(at), - fadvise_off(len), - libc::POSIX_FADV_DONTNEED, - ) - }; - if rc == 0 { - Ok(()) - } else { - Err(std::io::Error::from_raw_os_error(rc)) - } + os::fadvise_dontneed(file.as_raw_fd(), fadvise_off(at), fadvise_off(len)) + .map_err(std::io::Error::from) } /// `ids` as runs of adjacent rows, each `(first, last)` inclusive, sorted and @@ -699,7 +686,16 @@ impl TableBytes<'_> { /// [`Site::populate`]'s read: `POPULATE_READ` over each row's pages. fn populate(&self, ids: &[u32]) -> Result<(), EngramError> { - self.advise(ids, libc::MADV_POPULATE_READ, "madvise(POPULATE_READ)") + let page = u64::try_from(self.page).expect("a page size fits a u64"); + for &id in ids { + let (first, len) = self.pages(id)?; + let addr = first as *mut libc::c_void; + // SAFETY: the range is whole pages inside the mapping (`pages` + // rounds a row out to the mapping's own pages). + unsafe { os::populate_span(addr, len, page, || os::advise_populate(addr, len)) } + .map_err(|e| self.io("madvise(POPULATE_READ)", e.into()))?; + } + Ok(()) } /// [`Site::resident_rows`]' count: the rows of `ids` whose pages `mincore` @@ -748,14 +744,12 @@ impl TableBytes<'_> { /// when the page is in the page cache. fn mincore_chunk(&self, at: usize, len: usize) -> Result<[u8; MINCORE_PAGES], EngramError> { let mut vec = [0u8; MINCORE_PAGES]; - // SAFETY: `at .. at + len` is page-aligned and inside the mapping - // (`pages` rounds a row out to the mapping's own pages), and `vec` - // holds one byte for each of the at most `MINCORE_PAGES` pages of - // that range. - let rc = unsafe { libc::mincore(at as *mut libc::c_void, len, vec.as_mut_ptr()) }; - if rc != 0 { - return Err(self.io("mincore", std::io::Error::last_os_error())); - } + // `at .. at +len` is page-aligned and inside the mapping (`pages` + // rounds a row out to the mapping's own pages), and `vec` holds one + // byte for each of the at most `MINCORE_PAGES` pages of that range. + // SAFETY: the caller's contract above. + unsafe { os::mincore(at as *mut libc::c_void, len, &mut vec) } + .map_err(|e| self.io("mincore", e.into()))?; Ok(vec) } @@ -938,7 +932,8 @@ pub fn faults() -> Faults { /// This is the counter a helper-thread design needs: the process total cannot /// say which thread paid for a fault, and the whole claim of moving the read /// off the step thread is that the step thread stops taking them. Sample it on -/// the thread it describes — a thread cannot ask for another's. +/// the thread it describes — a thread cannot ask for another's. Linux only: a +/// kernel without `RUSAGE_THREAD` (macOS) refuses it, and that panics by name. pub fn faults_thread() -> Faults { faults_of(RUSAGE_THREAD) } @@ -946,11 +941,17 @@ pub fn faults_thread() -> Faults { fn faults_of(who: libc::c_int) -> Faults { // SAFETY: `usage` is a live, writable `rusage` for the duration of the call // and `getrusage` fills it or leaves it untouched. - let usage = unsafe { + let (rc, usage) = unsafe { let mut usage: libc::rusage = std::mem::zeroed(); - libc::getrusage(who, &mut usage); - usage + let rc = libc::getrusage(who, &mut usage); + (rc, usage) }; + assert_eq!( + rc, + 0, + "getrusage({who}) refused: {}", + std::io::Error::last_os_error() + ); Faults { major: usage.ru_majflt.max(0) as u64, minor: usage.ru_minflt.max(0) as u64, diff --git a/crates/engram/tests/engram.rs b/crates/engram/tests/engram.rs index 74d5e05b6..b834a8129 100644 --- a/crates/engram/tests/engram.rs +++ b/crates/engram/tests/engram.rs @@ -866,8 +866,15 @@ fn hw_engram_expert_stack_evicts_and_reads_back() { /// A direct handle reads the file's bytes at an aligned offset into an aligned /// buffer, carries the `O_DIRECT` flag, and a path that is not there is /// refused by name. +#[cfg_attr( + target_os = "linux", + ignore = "hw: the box's disk-backed target directory" +)] +#[cfg_attr( + not(target_os = "linux"), + ignore = "the Linux fact it reads: the open flags in /proc/self/fdinfo" +)] #[test] -#[ignore = "hw: the box's disk-backed target directory"] fn hw_engram_open_direct_reads_the_files_bytes() { use std::os::fd::AsRawFd; let scratch = Scratch::new("direct"); @@ -881,10 +888,11 @@ fn hw_engram_open_direct_reads_the_files_bytes() { .find_map(|l| l.strip_prefix("flags:")) .map(|v| u32::from_str_radix(v.trim(), 8).unwrap()) .expect("fdinfo names the open flags"); + let direct_flag = os::DIRECT_FLAG.expect("this platform has a direct-open flag"); assert_ne!( - flags & libc::O_DIRECT as u32, + flags & direct_flag, 0, - "the handle carries O_DIRECT: flags {flags:o}" + "the handle carries the direct flag: flags {flags:o}" ); let align = engram::DIRECT_ALIGN; diff --git a/crates/gguf/Cargo.toml b/crates/gguf/Cargo.toml index 11f58274d..c5d96c9b7 100644 --- a/crates/gguf/Cargo.toml +++ b/crates/gguf/Cargo.toml @@ -17,6 +17,8 @@ sha2.workspace = true thiserror.workspace = true [dev-dependencies] +# The lib test reads a thread's own page faults. +bloomery-os.workspace = true libc.workspace = true # The dequant oracle's sets open through their refset family; refset reads the GGUF files # through this crate, so the tests link the lib once and refset beside it. diff --git a/crates/gguf/src/lib.rs b/crates/gguf/src/lib.rs index 8f949148c..b8f95e0c4 100644 --- a/crates/gguf/src/lib.rs +++ b/crates/gguf/src/lib.rs @@ -1321,16 +1321,15 @@ mod split_tests { /// Page faults the calling thread has taken, minor and major. fn thread_faults() -> u64 { - // SAFETY: `rusage` is a C struct of integers; all-zero bytes are a value of it. - let mut usage: libc::rusage = unsafe { std::mem::zeroed() }; - // SAFETY: `usage` is a live, writable `rusage` for the duration of the call. - let rc = unsafe { libc::getrusage(libc::RUSAGE_THREAD, &mut usage) }; - assert_eq!(rc, 0, "getrusage(RUSAGE_THREAD)"); - u64::try_from(usage.ru_minflt + usage.ru_majflt).unwrap() + os::thread_faults().expect("getrusage(RUSAGE_THREAD)") } /// A file refused by the strict pass costs a populating open its header's /// pages, not its data's: the populate waits for the validation. + #[cfg_attr( + not(target_os = "linux"), + ignore = "the Linux fact it reads: a per-thread page-fault count" + )] #[test] fn a_refused_file_is_not_populated() { const DATA: usize = 64 << 20; diff --git a/crates/model/Cargo.toml b/crates/model/Cargo.toml index c08429616..d1956d291 100644 --- a/crates/model/Cargo.toml +++ b/crates/model/Cargo.toml @@ -15,6 +15,9 @@ workspace = true bloomery-gguf = { path = "../gguf" } bloomery-levers = { path = "../levers" } bloomery-models = { path = "../models" } +# The OS calls the placement path makes (the affinity mask, the page-cache +# advice, the resident-page count) are the one owner's. +bloomery-os.workspace = true # The placement planner, as a pure crate so its tests run on the Mac; this # crate re-exports it at `placement` and keeps the parts the Mac cannot run # (host_lock, workstation::host_available) beside the re-export. diff --git a/crates/model/src/fileio.rs b/crates/model/src/fileio.rs index 3a3a6c3a1..ab47a79c0 100644 --- a/crates/model/src/fileio.rs +++ b/crates/model/src/fileio.rs @@ -1,47 +1,15 @@ //! File-system calls std does not wrap. Each returns the bare `io::Error`; //! the caller names the path and the operation. -use std::ffi::CString; use std::io; -use std::os::unix::ffi::OsStrExt; use std::path::Path; -/// `p` as a C string; a path holding a NUL is `InvalidInput`. -fn c_path(p: &Path) -> io::Result { - CString::new(p.as_os_str().as_bytes()) - .map_err(|e| io::Error::new(io::ErrorKind::InvalidInput, e)) -} - /// Bytes free to this process on the filesystem that holds `dir`. pub fn free_bytes(dir: &Path) -> io::Result { - let c = c_path(dir)?; - let mut st = std::mem::MaybeUninit::::uninit(); - // SAFETY: `c` is a NUL-terminated path that lives across the call, and - // `st` is storage for one `statvfs`, which the call fills on success. - if unsafe { libc::statvfs(c.as_ptr(), st.as_mut_ptr()) } != 0 { - return Err(io::Error::last_os_error()); - } - // SAFETY: the call returned 0, so it wrote every field of `st`. - let st = unsafe { st.assume_init() }; - Ok(st.f_bavail.saturating_mul(st.f_frsize)) + os::free_bytes(dir).map_err(io::Error::from) } /// Rename `from` to `to`, refused with `AlreadyExists` when `to` exists. pub fn rename_noreplace(from: &Path, to: &Path) -> io::Result<()> { - let (a, b) = (c_path(from)?, c_path(to)?); - // SAFETY: both are NUL-terminated paths that live across the call. - let rc = unsafe { - libc::renameat2( - libc::AT_FDCWD, - a.as_ptr(), - libc::AT_FDCWD, - b.as_ptr(), - libc::RENAME_NOREPLACE, - ) - }; - if rc == 0 { - Ok(()) - } else { - Err(io::Error::last_os_error()) - } + os::rename_noreplace(from, to).map_err(io::Error::from) } diff --git a/crates/model/src/placement/host_lock.rs b/crates/model/src/placement/host_lock.rs index 3daa05218..42ef6d3ff 100644 --- a/crates/model/src/placement/host_lock.rs +++ b/crates/model/src/placement/host_lock.rs @@ -32,6 +32,7 @@ use std::thread; use std::time::{Duration, Instant}; use gguf::{Gguf, Split}; +use os::Populated; use super::{Device, ExpertList, Format, ModelTensor, PlacementError, Plan, Segment}; use crate::r8file::{R8Error, R8Source, Sidecar}; @@ -861,83 +862,55 @@ pub(crate) fn drop_pages( libc::off_t::try_from(v) .map_err(|_| PlacementError::Host(format!("{name} offset {v} passes off_t"))) }; - // SAFETY: the descriptor is this file's own and lives across the call; - // fadvise reads no memory of ours. - let rc = unsafe { - libc::posix_fadvise( - file.as_raw_fd(), - off(first)?, - off(end - first)?, - libc::POSIX_FADV_DONTNEED, - ) - }; - if rc != 0 { - return Err(PlacementError::Host(format!( - "posix_fadvise(DONTNEED) of {name} bytes {first}..{end}: {}", - io::Error::from_raw_os_error(rc) - ))); - } + os::fadvise_dontneed(file.as_raw_fd(), off(first)?, off(end - first)?).map_err(|e| { + PlacementError::Host(format!( + "posix_fadvise(DONTNEED) of {name} bytes {first}..{end}: {e}" + )) + })?; Ok(end - first) } -/// How [`populate_span`] faulted a span's pages in. -#[derive(Clone, Copy, Debug, PartialEq, Eq)] -enum Populated { - /// `MADV_POPULATE_READ`. - Advised, - /// A read of one byte a page: the kernel has no `MADV_POPULATE_READ`. - Touched, -} - /// The line a populate prints once when the kernel has no /// `MADV_POPULATE_READ`. +#[cfg(target_os = "linux")] const NO_POPULATE_READ: &str = "host populate: madvise(MADV_POPULATE_READ) is not in this kernel \ (EINVAL; it came in Linux 5.14), so the load reads one byte a page \ instead"; +/// The line a populate prints once when the platform has no +/// `MADV_POPULATE_READ`. +#[cfg(target_os = "macos")] +const NO_POPULATE_READ: &str = "host populate: this platform has no madvise(MADV_POPULATE_READ), \ + so the load reads one byte a page instead"; + /// `madvise(MADV_POPULATE_READ)` over `span`, a live sub-slice of a set's /// mappings: a shard's read-only file mapping, or the sidecar's, which is its /// file's mapping or a resident anonymous copy. -fn advise_populate(span: &[u8]) -> io::Result<()> { - // SAFETY: `span` is a live slice (the caller's contract above); - // MADV_POPULATE_READ faults its pages in on either kind of mapping and - // never writes them. - let rc = unsafe { - libc::madvise( - span.as_ptr().cast_mut().cast(), - span.len(), - libc::MADV_POPULATE_READ, - ) - }; - if rc == 0 { - Ok(()) - } else { - Err(io::Error::last_os_error()) - } +fn advise_populate(span: &[u8]) -> Result<(), os::OsError> { + // SAFETY: `span` is a live slice (the caller's contract above), so its + // pages are inside one live mapping. + unsafe { os::advise_populate(span.as_ptr().cast_mut().cast(), span.len()) } } /// Fault the pages of `span` in through `advise` (the madvise call, /// [`advise_populate`]); a kernel that does not know the advice (EINVAL, -/// before Linux 5.14) gets the same pages faulted by a read of one byte each -/// `page` bytes instead. Any other refusal is the walk's, by name. +/// before Linux 5.14) or a platform without it ([`os::OsError::Unsupported`]) +/// gets the same pages faulted by a read of one byte each `page` bytes +/// instead. Any other refusal is the walk's, by name. The error stays typed +/// through this seam: the fallback reads the variant and the errno. fn populate_span( span: &[u8], page: u64, - advise: impl Fn(&[u8]) -> io::Result<()>, + advise: impl Fn(&[u8]) -> Result<(), os::OsError>, ) -> Result { - match advise(span) { - Ok(()) => Ok(Populated::Advised), - Err(e) if e.raw_os_error() == Some(libc::EINVAL) => { - let step = usize::try_from(page).map_err(|_| format!("a page of {page} bytes"))?; - for at in (0..span.len()).step_by(step) { - // SAFETY: `at` is below `span.len()`, so the byte is inside - // the live slice; the read is volatile so that it is made. - let _ = unsafe { std::ptr::read_volatile(span.as_ptr().add(at)) }; - } - Ok(Populated::Touched) - } - Err(e) => Err(format!("madvise(MADV_POPULATE_READ): {e}")), + // SAFETY: `span` is a live slice (the caller's contract above), so its + // pages are inside one live mapping. + unsafe { + os::populate_span(span.as_ptr().cast_mut().cast(), span.len(), page, || { + advise(span) + }) } + .map_err(|e| format!("madvise(MADV_POPULATE_READ): {e}")) } /// What one chunk's thread hands back: every span its call succeeded on — @@ -1207,20 +1180,12 @@ fn host_pages( fn resident_pages(span: &[u8], page: u64) -> io::Result<(u64, u64)> { let n = (span.len() as u64).div_ceil(page); let mut vec = vec![0u8; usize::try_from(n).map_err(io::Error::other)?]; - // SAFETY: `span` starts on a page boundary inside one live mapping (a set - // run starts at a page multiple of the mapping's page-aligned start), and - // `vec` has one byte for each of its pages. mincore only reads the page - // tables; nothing is written through `addr`. - let rc = unsafe { - libc::mincore( - span.as_ptr().cast_mut().cast(), - span.len(), - vec.as_mut_ptr(), - ) - }; - if rc != 0 { - return Err(io::Error::last_os_error()); - } + // `span` starts on a page boundary inside one live mapping (a set run + // starts at a page multiple of the mapping's page-aligned start), and + // `vec` has one byte for each of its pages. + // SAFETY: the caller's contract above. + unsafe { os::mincore(span.as_ptr().cast_mut().cast(), span.len(), &mut vec) } + .map_err(io::Error::from)?; Ok((vec.iter().filter(|&&b| b & 1 != 0).count() as u64, n)) } @@ -1237,21 +1202,36 @@ mod tests { use gguf::write::{Layout, TensorDecl, Writer}; use gguf::{Split, Weights}; - /// A kernel without `MADV_POPULATE_READ` (EINVAL) gets the span's pages - /// read one byte each instead; an advice the kernel takes reads nothing - /// more, and any other refusal is named. The madvise call is the seam: the errno comes from - /// the test, not from the running kernel. + /// A kernel without `MADV_POPULATE_READ` (EINVAL) or a platform without + /// it (`Unsupported`) gets the span's pages read one byte each instead; an + /// advice the kernel takes reads nothing more, and any other refusal is + /// named. The madvise call is the seam: the error comes from the test, not + /// from the running kernel. #[test] fn a_kernel_without_populate_read_reads_a_byte_a_page() { let page = 4096u64; let span: Vec = (0..3 * 4096 + 7).map(|i| (i % 251) as u8).collect(); - let einval = |_: &[u8]| Err(std::io::Error::from_raw_os_error(libc::EINVAL)); + let einval = |_: &[u8]| { + Err(os::OsError::Io(std::io::Error::from_raw_os_error( + libc::EINVAL, + ))) + }; assert_eq!(populate_span(&span, page, einval), Ok(Populated::Touched)); + let absent = |_: &[u8]| { + Err(os::OsError::Unsupported { + call: "madvise(MADV_POPULATE_READ)", + }) + }; + assert_eq!(populate_span(&span, page, absent), Ok(Populated::Touched)); assert_eq!( populate_span(&span, page, |_: &[u8]| Ok(())), Ok(Populated::Advised) ); - let enomem = |_: &[u8]| Err(std::io::Error::from_raw_os_error(libc::ENOMEM)); + let enomem = |_: &[u8]| { + Err(os::OsError::Io(std::io::Error::from_raw_os_error( + libc::ENOMEM, + ))) + }; let e = populate_span(&span, page, enomem).unwrap_err(); assert!(e.starts_with("madvise(MADV_POPULATE_READ): "), "{e}"); assert_eq!( diff --git a/crates/model/src/r8file.rs b/crates/model/src/r8file.rs index 0fcc2c374..773fe1ff8 100644 --- a/crates/model/src/r8file.rs +++ b/crates/model/src/r8file.rs @@ -510,14 +510,8 @@ fn off(path: &Path, v: u64) -> Result { /// cache. fn sync_and_drop(file: &File, path: &Path, range: Range) -> Result<(), ModelError> { let (at, len) = (off(path, range.start)?, off(path, range.end - range.start)?); - let flags = libc::SYNC_FILE_RANGE_WAIT_BEFORE - | libc::SYNC_FILE_RANGE_WRITE - | libc::SYNC_FILE_RANGE_WAIT_AFTER; - // SAFETY: the descriptor is `file`'s own and lives across the call; the - // call reads no memory of ours. - if unsafe { libc::sync_file_range(file.as_raw_fd(), at, len, flags) } != 0 { - return Err(io_err(path, "sync_file_range", io::Error::last_os_error())); - } + os::sync_range(file.as_raw_fd(), at, len) + .map_err(|e| io_err(path, "sync_file_range", e.into()))?; drop_cached(file, path, range) } @@ -525,16 +519,8 @@ fn sync_and_drop(file: &File, path: &Path, range: Range) -> Result<(), Mode /// that lie wholly inside it and no process maps. fn drop_cached(file: &File, path: &Path, range: Range) -> Result<(), ModelError> { let (at, len) = (off(path, range.start)?, off(path, range.end - range.start)?); - // SAFETY: as in `sync_and_drop`: `file`'s own descriptor, no memory read. - let rc = unsafe { libc::posix_fadvise(file.as_raw_fd(), at, len, libc::POSIX_FADV_DONTNEED) }; - if rc != 0 { - return Err(io_err( - path, - "posix_fadvise(DONTNEED)", - io::Error::from_raw_os_error(rc), - )); - } - Ok(()) + os::fadvise_dontneed(file.as_raw_fd(), at, len) + .map_err(|e| io_err(path, "posix_fadvise(DONTNEED)", e.into())) } /// The finished `.part` becomes `out`: synced, still the file this run wrote diff --git a/crates/os/Cargo.toml b/crates/os/Cargo.toml new file mode 100644 index 000000000..02945785c --- /dev/null +++ b/crates/os/Cargo.toml @@ -0,0 +1,15 @@ +[package] +name = "bloomery-os" +version = "0.1.0" +edition = "2024" +publish = false +license = "MIT" + +[lib] +name = "os" + +[lints] +workspace = true + +[dependencies] +libc.workspace = true diff --git a/crates/os/src/lib.rs b/crates/os/src/lib.rs new file mode 100644 index 000000000..8320ebe7d --- /dev/null +++ b/crates/os/src/lib.rs @@ -0,0 +1,873 @@ +//! bloomery-os: the one owner of the operating-system calls the host chain +//! makes, so that what Linux and macOS do differently is written once. +//! +//! Every call here has a Linux body — the call the kernel has — and a macOS +//! body: the call that platform has, or a named refusal +//! ([`OsError::Unsupported`]) when it has none. A platform with neither is a +//! compile error, never a silent fallback: a caller that gets a refusal +//! handles it by name, as it already handles the kernel's. +//! +//! The Linux bodies are the ones the callers had before this crate existed, +//! moved here unchanged; the macOS bodies are the forms the Mac port map +//! names. A type the two platforms shape differently (a `mincore` vector, a +//! `statvfs` field) is settled here, so every caller sees one signature. + +use std::ffi::CString; +use std::fs::File; +use std::os::fd::RawFd; +use std::path::Path; + +#[cfg(not(any(target_os = "linux", target_os = "macos")))] +compile_error!("bloomery-os has a body for Linux and macOS only; add this platform's calls here"); + +/// A refusal by the platform: a call it does not have, or one the kernel +/// refused. +#[derive(Debug)] +pub enum OsError { + /// The platform has no such call; `call` names it. + Unsupported { call: &'static str }, + /// The kernel refused the call. + Io(std::io::Error), +} + +impl OsError { + /// Whether the call is one the platform does not have, or one the kernel + /// does not know: the two refusals the populate fallback stands for. + fn is_no_advice(&self) -> bool { + match self { + Self::Unsupported { .. } => true, + Self::Io(e) => e.raw_os_error() == Some(libc::EINVAL), + } + } +} + +impl std::fmt::Display for OsError { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + Self::Unsupported { call } => write!(f, "unsupported on this platform: {call}"), + Self::Io(e) => write!(f, "{e}"), + } + } +} + +impl std::error::Error for OsError { + fn source(&self) -> Option<&(dyn std::error::Error + 'static)> { + match self { + Self::Io(e) => Some(e), + Self::Unsupported { .. } => None, + } + } +} + +impl From for OsError { + fn from(e: std::io::Error) -> Self { + Self::Io(e) + } +} + +impl From for std::io::Error { + fn from(e: OsError) -> Self { + match e { + OsError::Unsupported { call } => std::io::Error::new( + std::io::ErrorKind::Unsupported, + format!("unsupported on this platform: {call}"), + ), + OsError::Io(e) => e, + } + } +} + +/// `p` as a C string; a path holding a NUL is refused. +fn c_path(p: &Path) -> Result { + use std::os::unix::ffi::OsStrExt; + CString::new(p.as_os_str().as_bytes()) + .map_err(|e| OsError::Io(std::io::Error::new(std::io::ErrorKind::InvalidInput, e))) +} + +// ------------------------------------------------------------------- affinity + +/// The cpus the calling thread may run on, ascending. +/// +/// Linux: the kernel's mask for the calling thread, read into a buffer that +/// grows until the mask fits. macOS: every logical cpu of the machine, for +/// the platform has no per-thread mask. +#[cfg(target_os = "linux")] +pub fn mask_cpus() -> Result, OsError> { + let mut words = 16usize; + loop { + let mut set = vec![0u64; words]; + let bytes = words * std::mem::size_of::(); + // SAFETY: `set` is `bytes` writable bytes, and `sched_getaffinity` + // writes at most that many. + let rc = unsafe { libc::sched_getaffinity(0, bytes, set.as_mut_ptr().cast()) }; + if rc == 0 { + return Ok(set + .iter() + .enumerate() + .flat_map(|(w, &bits)| { + (0..64u32) + .filter(move |b| bits >> b & 1 == 1) + .map(move |b| w as u32 * 64 + b) + }) + .collect()); + } + let e = std::io::Error::last_os_error(); + // EINVAL: the mask is wider than the buffer. + if e.raw_os_error() != Some(libc::EINVAL) || words >= 1 << 14 { + return Err(OsError::Io(e)); + } + words *= 2; + } +} + +/// The cpus the calling thread may run on, ascending. +/// +/// Linux: the kernel's mask for the calling thread, read into a buffer that +/// grows until the mask fits. macOS: every logical cpu of the machine, for +/// the platform has no per-thread mask. +#[cfg(target_os = "macos")] +pub fn mask_cpus() -> Result, OsError> { + let n = logical_cpus()?; + Ok((0..n).collect()) +} + +/// The machine's logical cpu count (`hw.logicalcpu`). +#[cfg(target_os = "macos")] +fn logical_cpus() -> Result { + let mut n: libc::c_int = 0; + let mut len = std::mem::size_of::(); + // SAFETY: `n` is a live `c_int` and `len` its size; the call writes at + // most that many bytes into it. + let rc = unsafe { + libc::sysctlbyname( + c"hw.logicalcpu".as_ptr(), + (&raw mut n).cast(), + (&raw mut len), + std::ptr::null_mut(), + 0, + ) + }; + if rc != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + u32::try_from(n).map_err(|_| { + OsError::Io(std::io::Error::new( + std::io::ErrorKind::InvalidData, + "hw.logicalcpu is not a cpu count", + )) + }) +} + +/// Set the calling thread's mask to `cpus`. +/// +/// Linux: the kernel's mask for the calling thread. macOS: the platform has +/// no per-thread mask, so the call is refused by name. +#[cfg(target_os = "linux")] +pub fn set_mask(cpus: &[u32]) -> Result<(), OsError> { + let named = 8 * std::mem::size_of::(); + for &c in cpus { + if c as usize >= named { + return Err(OsError::Io(std::io::Error::new( + std::io::ErrorKind::InvalidInput, + format!("cpu {c} is past the {named} a cpu_set_t names"), + ))); + } + } + // SAFETY: `set` is a zeroed cpu_set_t only written through CPU_SET with + // indices below its own bit count, and sched_setaffinity reads it with + // the matching size. + let rc = unsafe { + let mut set: libc::cpu_set_t = std::mem::zeroed(); + for &c in cpus { + libc::CPU_SET(c as usize, &mut set); + } + libc::sched_setaffinity(0, std::mem::size_of::(), &set) + }; + if rc == 0 { + Ok(()) + } else { + Err(OsError::Io(std::io::Error::last_os_error())) + } +} + +/// Set the calling thread's mask to `cpus`. +/// +/// Linux: the kernel's mask for the calling thread. macOS: the platform has +/// no per-thread mask, so the call is refused by name. +#[cfg(target_os = "macos")] +pub fn set_mask(_cpus: &[u32]) -> Result<(), OsError> { + Err(OsError::Unsupported { + call: "sched_setaffinity", + }) +} + +/// Pin the calling thread to one logical cpu. +/// +/// Linux: the kernel's mask, set to that one cpu. macOS: the platform has no +/// per-thread mask, so the call is refused by name and the pool floats. +#[cfg(target_os = "linux")] +pub fn pin(cpu: u32) -> Result<(), OsError> { + set_mask(&[cpu]) +} + +/// Pin the calling thread to one logical cpu. +/// +/// Linux: the kernel's mask, set to that one cpu. macOS: the platform has no +/// per-thread mask, so the call is refused by name and the pool floats. +#[cfg(target_os = "macos")] +pub fn pin(_cpu: u32) -> Result<(), OsError> { + Err(OsError::Unsupported { + call: "sched_setaffinity", + }) +} + +/// Widen the calling thread's mask to every cpu the set can name. +/// +/// Linux: the kernel's mask, set to every cpu it names. macOS: the platform +/// has no per-thread mask to widen — a thread there runs on every cpu — so +/// the call has nothing to do and says so. +#[cfg(target_os = "linux")] +pub fn float_everywhere() -> Result<(), OsError> { + let max = u32::try_from(8 * std::mem::size_of::()) + .expect("a cpu_set_t names fewer cpus than a u32 holds"); + let cpus: Vec = (0..max).collect(); + set_mask(&cpus) +} + +/// Widen the calling thread's mask to every cpu the set can name. +/// +/// Linux: the kernel's mask, set to every cpu it names. macOS: the platform +/// has no per-thread mask to widen — a thread there runs on every cpu — so +/// the call has nothing to do and says so. +#[cfg(target_os = "macos")] +pub fn float_everywhere() -> Result<(), OsError> { + Ok(()) +} + +// -------------------------------------------------------------------- rusage + +/// The page faults the calling thread has taken, minor and major. +/// +/// Linux: the kernel's count for the calling thread. macOS: the platform +/// counts a process, not a thread, so the call is refused by name. +#[cfg(target_os = "linux")] +pub fn thread_faults() -> Result { + // SAFETY: `rusage` is a C struct of integers; all-zero bytes are a value of it. + let mut usage: libc::rusage = unsafe { std::mem::zeroed() }; + // SAFETY: `usage` is a live, writable `rusage` for the duration of the call. + let rc = unsafe { libc::getrusage(libc::RUSAGE_THREAD, &mut usage) }; + assert_eq!(rc, 0, "getrusage(RUSAGE_THREAD)"); + Ok(u64::try_from(usage.ru_minflt + usage.ru_majflt).expect("a fault count is not negative")) +} + +/// The page faults the calling thread has taken, minor and major. +/// +/// Linux: the kernel's count for the calling thread. macOS: the platform +/// counts a process, not a thread, so the call is refused by name. +#[cfg(target_os = "macos")] +pub fn thread_faults() -> Result { + Err(OsError::Unsupported { + call: "getrusage(RUSAGE_THREAD)", + }) +} + +// ---------------------------------------------------------------------- file + +/// The open flag a direct read carries, where the platform has one: Linux's +/// `O_DIRECT`. macOS turns the cache off with a descriptor command +/// (`F_NOCACHE`), not an open flag, so it has none. +#[cfg(target_os = "linux")] +pub const DIRECT_FLAG: Option = Some(libc::O_DIRECT as u32); + +/// The open flag a direct read carries, where the platform has one: Linux's +/// `O_DIRECT`. macOS turns the cache off with a descriptor command +/// (`F_NOCACHE`), not an open flag, so it has none. +#[cfg(target_os = "macos")] +pub const DIRECT_FLAG: Option = None; + +/// `path` opened for reads that bypass the page cache. +/// +/// Linux: the `O_DIRECT` flag. macOS: the flag has no counterpart, so the +/// file is opened and its cache turned off with `F_NOCACHE`. +#[cfg(target_os = "linux")] +pub fn open_direct(path: &Path) -> Result { + use std::os::unix::fs::OpenOptionsExt; + std::fs::OpenOptions::new() + .read(true) + .custom_flags(libc::O_DIRECT) + .open(path) + .map_err(OsError::Io) +} + +/// `path` opened for reads that bypass the page cache. +/// +/// Linux: the `O_DIRECT` flag. macOS: the flag has no counterpart, so the +/// file is opened and its cache turned off with `F_NOCACHE`. +#[cfg(target_os = "macos")] +pub fn open_direct(path: &Path) -> Result { + use std::os::fd::AsRawFd; + let file = std::fs::OpenOptions::new() + .read(true) + .open(path) + .map_err(OsError::Io)?; + // SAFETY: the descriptor is this file's own and lives across the call. + let rc = unsafe { libc::fcntl(file.as_raw_fd(), libc::F_NOCACHE, 1) }; + if rc != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + Ok(file) +} + +/// Rename `from` to `to`, refused when `to` exists. +/// +/// Linux: `renameat2` with `RENAME_NOREPLACE`. macOS: `renamex_np` with +/// `RENAME_EXCL`. +#[cfg(target_os = "linux")] +pub fn rename_noreplace(from: &Path, to: &Path) -> Result<(), OsError> { + let (a, b) = (c_path(from)?, c_path(to)?); + // SAFETY: both are NUL-terminated paths that live across the call. + let rc = unsafe { + libc::renameat2( + libc::AT_FDCWD, + a.as_ptr(), + libc::AT_FDCWD, + b.as_ptr(), + libc::RENAME_NOREPLACE, + ) + }; + if rc == 0 { + Ok(()) + } else { + Err(OsError::Io(std::io::Error::last_os_error())) + } +} + +/// Rename `from` to `to`, refused when `to` exists. +/// +/// Linux: `renameat2` with `RENAME_NOREPLACE`. macOS: `renamex_np` with +/// `RENAME_EXCL`. +#[cfg(target_os = "macos")] +pub fn rename_noreplace(from: &Path, to: &Path) -> Result<(), OsError> { + let (a, b) = (c_path(from)?, c_path(to)?); + // SAFETY: both are NUL-terminated paths that live across the call. + let rc = unsafe { libc::renamex_np(a.as_ptr(), b.as_ptr(), libc::RENAME_EXCL) }; + if rc == 0 { + Ok(()) + } else { + Err(OsError::Io(std::io::Error::last_os_error())) + } +} + +/// Bytes free to this process on the filesystem that holds `dir`. +/// +/// Linux: the `statvfs` fields as they are. macOS: the same fields, whose +/// widths that platform sets differently, widened to the same numbers. +#[cfg(target_os = "linux")] +pub fn free_bytes(dir: &Path) -> Result { + let c = c_path(dir)?; + let mut st = std::mem::MaybeUninit::::uninit(); + // SAFETY: `c` is a NUL-terminated path that lives across the call, and + // `st` is storage for one `statvfs`, which the call fills on success. + if unsafe { libc::statvfs(c.as_ptr(), st.as_mut_ptr()) } != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + // SAFETY: the call returned 0, so it wrote every field of `st`. + let st = unsafe { st.assume_init() }; + Ok(st.f_bavail.saturating_mul(st.f_frsize)) +} + +/// Bytes free to this process on the filesystem that holds `dir`. +/// +/// Linux: the `statvfs` fields as they are. macOS: the same fields, whose +/// widths that platform sets differently, widened to the same numbers. +#[cfg(target_os = "macos")] +pub fn free_bytes(dir: &Path) -> Result { + let c = c_path(dir)?; + let mut st = std::mem::MaybeUninit::::uninit(); + // SAFETY: as in the Linux body: `c` is a NUL-terminated path that lives + // across the call, and `st` is storage for one `statvfs`. + if unsafe { libc::statvfs(c.as_ptr(), st.as_mut_ptr()) } != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + // SAFETY: the call returned 0, so it wrote every field of `st`. + let st = unsafe { st.assume_init() }; + // `f_bavail` is a 32-bit count on macOS and `f_frsize` a 64-bit one: the + // same numbers Linux's fields hold, widened once. + Ok(u64::from(st.f_bavail).saturating_mul(st.f_frsize)) +} + +/// Drop the clean pages of `fd` over `at .. at + len` from the page cache. +/// +/// Linux: `posix_fadvise(DONTNEED)`. macOS: the platform has no such call, +/// so it is refused by name. +#[cfg(target_os = "linux")] +pub fn fadvise_dontneed(fd: RawFd, at: libc::off_t, len: libc::off_t) -> Result<(), OsError> { + // SAFETY: the descriptor is the caller's own and lives across the call; + // fadvise reads and writes no memory of ours. + let rc = unsafe { libc::posix_fadvise(fd, at, len, libc::POSIX_FADV_DONTNEED) }; + if rc == 0 { + Ok(()) + } else { + Err(OsError::Io(std::io::Error::from_raw_os_error(rc))) + } +} + +/// Drop the clean pages of `fd` over `at .. at + len` from the page cache. +/// +/// Linux: `posix_fadvise(DONTNEED)`. macOS: the platform has no such call, +/// so it is refused by name. +#[cfg(target_os = "macos")] +pub fn fadvise_dontneed(_fd: RawFd, _at: libc::off_t, _len: libc::off_t) -> Result<(), OsError> { + Err(OsError::Unsupported { + call: "posix_fadvise", + }) +} + +/// Write the bytes of `fd` over `at .. at + len` out to the drive. +/// +/// Linux: `sync_file_range` over the range. macOS: the platform has no +/// ranged write-out, so the whole file is synced with `F_FULLFSYNC`. +#[cfg(target_os = "linux")] +pub fn sync_range(fd: RawFd, at: libc::off_t, len: libc::off_t) -> Result<(), OsError> { + let flags = libc::SYNC_FILE_RANGE_WAIT_BEFORE + | libc::SYNC_FILE_RANGE_WRITE + | libc::SYNC_FILE_RANGE_WAIT_AFTER; + // SAFETY: the descriptor is the caller's own and lives across the call; the + // call reads no memory of ours. + if unsafe { libc::sync_file_range(fd, at, len, flags) } != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + Ok(()) +} + +/// Write the bytes of `fd` over `at .. at + len` out to the drive. +/// +/// Linux: `sync_file_range` over the range. macOS: the platform has no +/// ranged write-out, so the whole file is synced with `F_FULLFSYNC`. +#[cfg(target_os = "macos")] +pub fn sync_range(fd: RawFd, _at: libc::off_t, _len: libc::off_t) -> Result<(), OsError> { + // SAFETY: the descriptor is the caller's own and lives across the call. + let rc = unsafe { libc::fcntl(fd, libc::F_FULLFSYNC, 0) }; + if rc != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + Ok(()) +} + +// ------------------------------------------------------------------- memory + +/// How a span's pages were faulted in. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum Populated { + /// The platform's populate advice. + Advised, + /// A read of one byte a page. + Touched, +} + +/// The platform's populate advice: fault the pages of `addr .. addr + len` in. +/// +/// Linux: `madvise(MADV_POPULATE_READ)`. macOS: the platform has no such +/// advice, so the call is refused by name and the pages are read in instead. +/// +/// # Safety +/// `addr .. addr + len` must be whole pages inside one live mapping. +#[cfg(target_os = "linux")] +pub unsafe fn advise_populate(addr: *mut libc::c_void, len: usize) -> Result<(), OsError> { + // SAFETY: the caller's contract above; MADV_POPULATE_READ faults the + // pages in and never writes them. + let rc = unsafe { libc::madvise(addr, len, libc::MADV_POPULATE_READ) }; + if rc == 0 { + Ok(()) + } else { + Err(OsError::Io(std::io::Error::last_os_error())) + } +} + +/// The platform's populate advice: fault the pages of `addr .. addr + len` in. +/// +/// Linux: `madvise(MADV_POPULATE_READ)`. macOS: the platform has no such +/// advice, so the call is refused by name and the pages are read in instead. +/// +/// # Safety +/// `addr .. addr + len` must be whole pages inside one live mapping. +#[cfg(target_os = "macos")] +pub unsafe fn advise_populate(_addr: *mut libc::c_void, _len: usize) -> Result<(), OsError> { + Err(OsError::Unsupported { + call: "madvise(MADV_POPULATE_READ)", + }) +} + +/// Fault the pages of `addr .. addr + len` in. +/// +/// `advise` is the platform's populate advice call. A refusal that names the +/// advice as unknown (EINVAL) or absent ([`OsError::Unsupported`]) gets the +/// same pages faulted by a read of one byte each `page` bytes instead; any +/// other refusal is the caller's, by name. +/// +/// # Safety +/// `addr .. addr + len` must be inside one live mapping: the pages are read +/// one byte a page when the advice is not there. +pub unsafe fn populate_span( + addr: *mut libc::c_void, + len: usize, + page: u64, + advise: impl Fn() -> Result<(), OsError>, +) -> Result { + match advise() { + Ok(()) => Ok(Populated::Advised), + Err(e) if e.is_no_advice() => { + // SAFETY: the caller's contract above. + unsafe { touch(addr, len, page)? }; + Ok(Populated::Touched) + } + Err(e) => Err(e), + } +} + +/// The pages of `addr .. addr + len` faulted in by a read of one byte a page. +/// +/// # Safety +/// `addr .. addr + len` must be inside one live mapping. +unsafe fn touch(addr: *mut libc::c_void, len: usize, page: u64) -> Result<(), OsError> { + let step = usize::try_from(page).map_err(|_| { + OsError::Io(std::io::Error::new( + std::io::ErrorKind::InvalidInput, + format!("a page of {page} bytes"), + )) + })?; + for at in (0..len).step_by(step) { + // SAFETY: `at` is below `len`, so the byte is inside the caller's live + // range; the read is volatile so that it is made. + let _ = unsafe { std::ptr::read_volatile(addr.cast::().add(at)) }; + } + Ok(()) +} + +/// The kernel's `mincore` vector for `vec`: a byte a page. +#[cfg(target_os = "linux")] +fn mincore_vec(vec: &mut [u8]) -> *mut libc::c_uchar { + vec.as_mut_ptr() +} + +/// The kernel's `mincore` vector for `vec`: a `c_char` a page. +#[cfg(target_os = "macos")] +fn mincore_vec(vec: &mut [u8]) -> *mut libc::c_char { + vec.as_mut_ptr().cast() +} + +/// The address `mincore` takes: a mutable pointer on Linux, a constant one +/// on macOS. +#[cfg(target_os = "linux")] +fn mincore_addr(addr: *mut libc::c_void) -> *mut libc::c_void { + addr +} + +/// The address `mincore` takes: a mutable pointer on Linux, a constant one +/// on macOS. +#[cfg(target_os = "macos")] +fn mincore_addr(addr: *mut libc::c_void) -> *const libc::c_void { + addr.cast_const() +} + +/// The resident pages of `addr .. addr + len`, one byte a page in `vec`. +/// +/// Linux: the kernel's `mincore`, whose vector is a byte a page. macOS: the +/// same call, whose vector is a `c_char` a page; the shapes are settled here. +/// +/// # Safety +/// `addr .. addr + len` must be inside one live mapping, and `vec` must hold +/// one byte for each of the pages of that range. +pub unsafe fn mincore(addr: *mut libc::c_void, len: usize, vec: &mut [u8]) -> Result<(), OsError> { + // SAFETY: the caller's contract above; mincore only reads the page tables + // and writes nothing through `addr`. + let rc = unsafe { libc::mincore(mincore_addr(addr), len, mincore_vec(vec)) }; + if rc != 0 { + return Err(OsError::Io(std::io::Error::last_os_error())); + } + Ok(()) +} + +#[cfg(test)] +mod tests { + use super::{OsError, Populated}; + use std::path::PathBuf; + + /// A scratch file of this process's own making, removed when the test + /// ends. + fn scratch(name: &str) -> PathBuf { + let p = std::env::temp_dir().join(format!("bloomery-os-{name}-{}", std::process::id())); + let _ = std::fs::remove_file(&p); + p + } + + /// `len` bytes on a page boundary, zeroed: the mapping the populate + /// advice and the resident-page count take. + struct Page { + ptr: *mut u8, + len: usize, + layout: std::alloc::Layout, + } + + impl Page { + fn new(len: usize) -> Page { + let layout = + std::alloc::Layout::from_size_align(len, 4096).expect("a page-aligned layout"); + // SAFETY: `layout` has a non-zero size; the memory is zeroed below + // before anything reads it. + let ptr = unsafe { std::alloc::alloc(layout) }; + assert!(!ptr.is_null(), "the page-aligned memory is there"); + // SAFETY: `ptr` is `len` writable bytes of the allocation above. + unsafe { std::ptr::write_bytes(ptr, 0, len) }; + Page { ptr, len, layout } + } + + fn addr(&self) -> *mut std::ffi::c_void { + self.ptr.cast() + } + } + + impl Drop for Page { + fn drop(&mut self) { + // SAFETY: `ptr` came from `alloc` with this layout, and nothing + // uses it after the drop. + unsafe { std::alloc::dealloc(self.ptr, self.layout) }; + } + } + + /// The cpus the calling thread may run on: at least one, ascending. + #[test] + fn the_mask_names_the_cpus_the_thread_may_run_on() { + let cpus = super::mask_cpus().expect("the mask reads"); + assert!(!cpus.is_empty(), "the mask names a cpu"); + assert!( + cpus.windows(2).all(|w| w[0] < w[1]), + "the mask is ascending: {cpus:?}" + ); + } + + /// The mask the caller sets is the mask the thread then has; the platform + /// with no per-thread mask refuses the call by name. + #[test] + fn the_mask_is_the_one_the_caller_sets() { + #[cfg(target_os = "linux")] + { + let cpus = super::mask_cpus().expect("the mask reads"); + let one = &cpus[..1]; + super::set_mask(one).expect("the mask is set"); + assert_eq!(super::mask_cpus().expect("the mask reads back"), one); + } + #[cfg(target_os = "macos")] + { + let e = super::set_mask(&[0]).expect_err("the platform has no mask"); + assert!( + matches!(e, OsError::Unsupported { call } if call == "sched_setaffinity"), + "{e}" + ); + } + } + + /// A pin of a cpu the mask does not name is refused. + #[cfg(target_os = "linux")] + #[test] + fn a_pin_past_the_mask_is_refused() { + let e = super::pin(u32::MAX).expect_err("the cpu is past the mask"); + assert!( + matches!(&e, OsError::Io(e) if e.kind() == std::io::ErrorKind::InvalidInput), + "{e}" + ); + } + + /// A pin of a cpu the mask names holds the thread on it; the platform with + /// no per-thread mask refuses the pin by name. + #[test] + fn a_pin_holds_the_thread_on_its_cpu() { + #[cfg(target_os = "linux")] + { + let cpus = super::mask_cpus().expect("the mask reads"); + let cpu = cpus[0]; + super::pin(cpu).expect("the pin holds"); + assert_eq!(super::mask_cpus().expect("the mask reads back"), vec![cpu]); + } + #[cfg(target_os = "macos")] + { + let e = super::pin(0).expect_err("the platform has no mask"); + assert!( + matches!(e, OsError::Unsupported { call } if call == "sched_setaffinity"), + "{e}" + ); + } + } + + /// Floating widens a narrow mask back to the cpus the process may run on; + /// the platform with no mask has nothing to widen and says so. + #[test] + fn floating_widens_the_mask() { + #[cfg(target_os = "linux")] + { + let cpus = super::mask_cpus().expect("the mask reads"); + super::set_mask(&cpus[..1]).expect("the mask is set"); + super::float_everywhere().expect("the mask widens"); + let after = super::mask_cpus().expect("the mask reads back"); + assert!( + after.len() >= cpus.len(), + "the mask widened: {cpus:?} then {after:?}" + ); + } + #[cfg(target_os = "macos")] + super::float_everywhere().expect("there is no mask to widen"); + } + + /// The page faults of the calling thread are a count that does not go + /// back; the platform that counts a process, not a thread, refuses the + /// per-thread count by name. + #[test] + fn the_thread_faults_are_counted() { + #[cfg(target_os = "linux")] + { + let before = super::thread_faults().expect("the fault count reads"); + let after = super::thread_faults().expect("the fault count reads"); + assert!( + after >= before, + "the count does not go back: {before} then {after}" + ); + } + #[cfg(target_os = "macos")] + { + let e = super::thread_faults().expect_err("the platform counts a process"); + assert!( + matches!(e, OsError::Unsupported { call } if call == "getrusage(RUSAGE_THREAD)"), + "{e}" + ); + } + } + + /// A direct open of a file the process wrote is the handle of that file; + /// the flag the platform has for it is named. + #[test] + fn a_direct_open_is_the_handle_of_the_file() { + let p = scratch("direct"); + std::fs::write(&p, vec![7u8; 4096]).expect("the file is written"); + let file = super::open_direct(&p).expect("the file opens for direct reads"); + let len = file.metadata().expect("the file has a size").len(); + let _ = std::fs::remove_file(&p); + assert_eq!(len, 4096, "the handle is the file that was written"); + #[cfg(target_os = "linux")] + assert!(super::DIRECT_FLAG.is_some(), "the platform has the flag"); + #[cfg(target_os = "macos")] + assert!(super::DIRECT_FLAG.is_none(), "the platform has no flag"); + } + + /// A rename that replaces nothing moves the file; a rename onto an + /// existing file is refused. + #[test] + fn a_rename_onto_an_existing_file_is_refused() { + let from = scratch("rename-from"); + let to = scratch("rename-to"); + let there = scratch("rename-there"); + std::fs::write(&from, b"a").expect("the file is written"); + std::fs::write(&there, b"b").expect("the file is written"); + super::rename_noreplace(&from, &to).expect("the rename moves the file"); + assert!(to.exists(), "the file is at its new name"); + let e = super::rename_noreplace(&there, &to).expect_err("the target exists"); + assert!( + matches!(&e, OsError::Io(e) if e.raw_os_error() == Some(libc::EEXIST)), + "{e}" + ); + let _ = std::fs::remove_file(&to); + let _ = std::fs::remove_file(&there); + } + + /// The free bytes of a directory are its free blocks times their size, so + /// the count is a whole number of 512-byte sectors (every filesystem's + /// fragment size is one); a count of blocks without their size is not. + #[test] + fn the_free_bytes_of_a_directory_are_counted() { + let dir = std::env::temp_dir(); + let free = super::free_bytes(&dir).expect("the filesystem answers"); + assert!(free > 0, "the filesystem has room: {free}"); + assert_eq!( + free % 512, + 0, + "free_bytes {free} is not blocks times their size" + ); + } + + /// The page cache drops the clean pages of a file; the platform with no + /// `posix_fadvise` refuses the drop by name. + #[test] + fn the_page_cache_drops_a_files_clean_pages() { + let p = scratch("fadvise"); + std::fs::write(&p, b"bloomery").expect("the file is written"); + let file = std::fs::File::open(&p).expect("the file opens"); + let fd = { + use std::os::fd::AsRawFd; + file.as_raw_fd() + }; + let dropped = super::fadvise_dontneed(fd, 0, 8); + let _ = std::fs::remove_file(&p); + #[cfg(target_os = "linux")] + dropped.expect("the pages drop"); + #[cfg(target_os = "macos")] + { + let e = dropped.expect_err("the platform has no call"); + assert!( + matches!(e, OsError::Unsupported { call } if call == "posix_fadvise"), + "{e}" + ); + } + } + + /// A file's bytes are written out. + #[test] + fn a_files_bytes_are_written_out() { + let p = scratch("sync"); + std::fs::write(&p, b"bloomery").expect("the file is written"); + let file = std::fs::File::open(&p).expect("the file opens"); + let fd = { + use std::os::fd::AsRawFd; + file.as_raw_fd() + }; + super::sync_range(fd, 0, 8).expect("the bytes are written out"); + let _ = std::fs::remove_file(&p); + } + + /// The pages of a span are faulted in: by the platform's advice where it + /// has one, by a read of one byte a page where it has none. + #[test] + fn a_span_is_faulted_in() { + let span = Page::new(8192); + let addr = span.addr(); + // SAFETY: the span is the page-aligned allocation above, so its pages + // are inside one live mapping. + let got = unsafe { + super::populate_span(addr, span.len, 4096, || { + super::advise_populate(addr, span.len) + }) + } + .expect("the pages fault in"); + #[cfg(target_os = "linux")] + assert_eq!(got, Populated::Advised); + #[cfg(target_os = "macos")] + assert_eq!(got, Populated::Touched); + } + + /// The resident pages of a span are counted, a byte a page. + #[test] + fn a_span_s_resident_pages_are_counted() { + let span = Page::new(8192); + let mut vec = [0u8; 2]; + // SAFETY: the span is the page-aligned allocation above, so its pages + // are inside one live mapping, and `vec` holds a byte for each. + unsafe { super::mincore(span.addr(), span.len, &mut vec) }.expect("the pages are counted"); + // Linux answers for every page of the range; macOS answers for the + // first page of it, so the claim both make is that one. + #[cfg(target_os = "linux")] + assert!( + vec.iter().all(|b| b & 1 == 1), + "the pages are resident: {vec:?}" + ); + #[cfg(target_os = "macos")] + assert_eq!(vec[0] & 1, 1, "the first page is resident: {vec:?}"); + } +} diff --git a/crates/threads/Cargo.toml b/crates/threads/Cargo.toml index 99dbba90e..6f0067d8e 100644 --- a/crates/threads/Cargo.toml +++ b/crates/threads/Cargo.toml @@ -13,4 +13,5 @@ workspace = true [dependencies] bloomery-levers = { path = "../levers" } +bloomery-os.workspace = true libc.workspace = true diff --git a/crates/threads/src/helper.rs b/crates/threads/src/helper.rs index efa0fe72f..a5ba4a4f2 100644 --- a/crates/threads/src/helper.rs +++ b/crates/threads/src/helper.rs @@ -214,33 +214,13 @@ fn place(cpu: Option) -> Result, HelperError> { /// Widen the calling thread's affinity to every cpu the set can name; false /// when the kernel refuses. fn float_everywhere() -> bool { - let max = 8 * std::mem::size_of::(); - // SAFETY: `set` is a zeroed cpu_set_t only written through `CPU_SET` with - // indices below its own bit count, and `sched_setaffinity` reads it with - // the matching size. - unsafe { - let mut set: libc::cpu_set_t = std::mem::zeroed(); - for c in 0..max { - libc::CPU_SET(c, &mut set); - } - libc::sched_setaffinity(0, std::mem::size_of::(), &set) == 0 - } + os::float_everywhere().is_ok() } /// Pin the calling thread to one logical cpu; false when the kernel refuses /// (a container, a cpu outside the process's set). fn pin_to(cpu: usize) -> bool { - if cpu >= 8 * std::mem::size_of::() { - return false; - } - // SAFETY: `set` is a zeroed cpu_set_t only written through `CPU_SET` with - // an index inside it, and `sched_setaffinity` reads it with the matching - // size. - unsafe { - let mut set: libc::cpu_set_t = std::mem::zeroed(); - libc::CPU_SET(cpu, &mut set); - libc::sched_setaffinity(0, std::mem::size_of::(), &set) == 0 - } + u32::try_from(cpu).is_ok_and(|c| os::pin(c).is_ok()) } /// The cpus of a sysfs cpu list (`0,32`, `0-1`, `0-3,8`); a malformed entry diff --git a/crates/threads/src/lib.rs b/crates/threads/src/lib.rs index 6f527eeb5..c503ec3f4 100644 --- a/crates/threads/src/lib.rs +++ b/crates/threads/src/lib.rs @@ -744,35 +744,11 @@ pub fn chunk_bounds(n: usize, threads: usize, t: usize) -> (usize, usize) { // --------------------------------------------------------------------- mask -/// The cpus the calling thread may run on (`sched_getaffinity`), ascending. -/// The buffer grows until the kernel's mask fits, so a machine of more than -/// 1024 cpus reads whole. +/// The cpus the calling thread may run on, ascending: the kernel's mask for +/// the calling thread, read into a buffer that grows until the mask fits, so +/// a machine of more than 1024 cpus reads whole. fn mask_cpus() -> std::io::Result> { - let mut words = 16usize; - loop { - let mut set = vec![0u64; words]; - let bytes = words * std::mem::size_of::(); - // SAFETY: `set` is `bytes` writable bytes, and `sched_getaffinity` - // writes at most that many. - let rc = unsafe { libc::sched_getaffinity(0, bytes, set.as_mut_ptr().cast()) }; - if rc == 0 { - return Ok(set - .iter() - .enumerate() - .flat_map(|(w, &bits)| { - (0..64u32) - .filter(move |b| bits >> b & 1 == 1) - .map(move |b| w as u32 * 64 + b) - }) - .collect()); - } - let e = std::io::Error::last_os_error(); - // EINVAL: the mask is wider than the buffer. - if e.raw_os_error() != Some(libc::EINVAL) || words >= 1 << 14 { - return Err(e); - } - words *= 2; - } + os::mask_cpus().map_err(std::io::Error::from) } /// The pool's shape from this machine: its sysfs topology cut to the calling @@ -805,28 +781,10 @@ pub fn report() -> Result { /// Pin the calling thread to one logical cpu. Returns the kernel's refusal /// instead of panicking — a seccomp filter or a cpu outside the process's mask -/// cannot be pinned, and the pool must keep running, floating. +/// cannot be pinned, and the pool must keep running, floating. A platform +/// with no per-thread mask refuses the pin by name. fn pin(cpu: u32) -> std::io::Result<()> { - let named = 8 * std::mem::size_of::(); - if cpu as usize >= named { - return Err(std::io::Error::new( - std::io::ErrorKind::InvalidInput, - format!("cpu {cpu} is past the {named} a cpu_set_t names"), - )); - } - // SAFETY: `set` is a zeroed cpu_set_t only written through CPU_SET with - // an index inside it, and sched_setaffinity reads it with the matching - // size. - let rc = unsafe { - let mut set: libc::cpu_set_t = std::mem::zeroed(); - libc::CPU_SET(cpu as usize, &mut set); - libc::sched_setaffinity(0, std::mem::size_of::(), &set) - }; - if rc == 0 { - Ok(()) - } else { - Err(std::io::Error::last_os_error()) - } + os::pin(cpu).map_err(std::io::Error::from) } #[cfg(test)] @@ -835,19 +793,10 @@ mod tests { use std::sync::atomic::Ordering; /// Run `f` on a thread pinned to `cpus`, so the kernel's mask for it is - /// theirs: a test of the real `sched_getaffinity` read. + /// theirs: a test of the real mask read. fn on_cpus(cpus: Vec, f: impl FnOnce() -> T + Send + 'static) -> T { std::thread::spawn(move || { - // SAFETY: `set` is a zeroed cpu_set_t only written through - // CPU_SET, and sched_setaffinity reads it with the matching size. - let rc = unsafe { - let mut set: libc::cpu_set_t = std::mem::zeroed(); - for &c in &cpus { - libc::CPU_SET(c as usize, &mut set); - } - libc::sched_setaffinity(0, std::mem::size_of::(), &set) - }; - assert_eq!(rc, 0, "the test thread pins to {cpus:?}"); + os::set_mask(&cpus).unwrap_or_else(|e| panic!("the test thread pins to {cpus:?}: {e}")); f() }) .join() @@ -860,6 +809,10 @@ mod tests { /// thread of two the pool takes two cores or one with its sibling, and a /// thread count the lever gave stands over it. The unpinned process sees /// the whole mask, the topology whole and no cut. + #[cfg_attr( + not(target_os = "linux"), + ignore = "the Linux fact it reads: a thread's own affinity mask" + )] #[test] fn the_pool_plans_inside_the_threads_mask() { let whole = super::mask_cpus().expect("the mask reads"); @@ -900,6 +853,10 @@ mod tests { /// The built pool and the plan before it are one placement, and a pool /// built on a one-cpu thread pins its workers inside that mask. + #[cfg_attr( + not(target_os = "linux"), + ignore = "the Linux fact it reads: a thread's own affinity mask" + )] #[test] fn a_built_pool_is_the_plan() { let planned = super::plan_pool().expect("a plan"); @@ -934,8 +891,15 @@ mod tests { assert!(p.pin_failed()); assert!(p.plan().workers.is_empty()); let refused = super::pin(1 << 20).expect_err("no cpu_set_t holds that cpu"); - assert_eq!(refused.kind(), std::io::ErrorKind::InvalidInput); - assert!(refused.to_string().contains("1048576"), "{refused}"); + // Linux: the cpu is past the mask's size, and the refusal says so. + // A platform with no per-thread mask refuses the pin by name. + #[cfg(target_os = "linux")] + { + assert_eq!(refused.kind(), std::io::ErrorKind::InvalidInput); + assert!(refused.to_string().contains("1048576"), "{refused}"); + } + #[cfg(target_os = "macos")] + assert_eq!(refused.kind(), std::io::ErrorKind::Unsupported, "{refused}"); } /// A pool has a CCD map only once every participant is pinned: the diff --git a/crates/threads/tests/helper.rs b/crates/threads/tests/helper.rs index a35a5876e..a8511bff0 100644 --- a/crates/threads/tests/helper.rs +++ b/crates/threads/tests/helper.rs @@ -8,17 +8,11 @@ use threads::helper::{Placement, helpers, mask, sibling_of, spawn_helper}; /// Set the calling thread's mask to `cpus`. fn set_mask(cpus: &[usize]) { - // SAFETY: `set` is a zeroed cpu_set_t only written through `CPU_SET` with - // indices the test took from a mask of the same size, and - // `sched_setaffinity` reads it with the matching size. - let ok = unsafe { - let mut set: libc::cpu_set_t = std::mem::zeroed(); - for &c in cpus { - libc::CPU_SET(c, &mut set); - } - libc::sched_setaffinity(0, std::mem::size_of::(), &set) == 0 - }; - assert!(ok, "sched_setaffinity to {cpus:?}"); + let cpus: Vec = cpus + .iter() + .map(|&c| u32::try_from(c).expect("a cpu id fits a u32")) + .collect(); + os::set_mask(&cpus).unwrap_or_else(|e| panic!("the mask is set to {cpus:?}: {e}")); } /// The mask a helper named `name`, placed as `asked`, runs its body with when @@ -47,6 +41,10 @@ fn helper_mask(name: &str, opener: Vec, asked: Placement) -> (Vec, /// already, left as it is; on the cpu asked for; and on the SMT sibling of /// the cpu asked for, from an opener pinned to that cpu. A test thread whose own /// mask holds one cpu cannot tell the float from the pin and skips by name. +#[cfg_attr( + not(target_os = "linux"), + ignore = "the Linux fact it reads: a thread's own affinity mask" +)] #[test] fn a_helper_leaves_its_openers_one_cpu_mask() { let wide = mask().expect("the test thread's mask"); diff --git a/docs/contrib/mac.md b/docs/contrib/mac.md new file mode 100644 index 000000000..97398ae74 --- /dev/null +++ b/docs/contrib/mac.md @@ -0,0 +1,574 @@ +# Porting bloomery to Apple Silicon (Metal) + +This is the map and the design for running bloomery on an Apple Silicon Mac (aarch64-apple-darwin, Metal GPU, one +unified memory). It says what builds on a Mac today, which work is left, in what order, and what proves each step. It +changes no code. The target machine is the M5 Ultra, and the goal is the fastest decode and prefill on the M5 +generation, so the Metal backend (Part C) is the main part and the CPU path (Part B) is cut to what the Metal work needs. + +**How to read the numbers.** A claim about the code cites `path:line` at main `8d74220bb`. A number computed here is +marked `[derived]` with its inputs. A thing that was not seen is marked `[unverified]`. A hardware or toolchain fact +from the web carries its URL and one line copied from that page; a fact read from a local file names the file. Every +URL was opened in this round. "Measured" means a command run in this round on the one Mac that was available, an +Apple M1 Pro with 32 GiB (the machine facts are in section 2); every probe's log is under +`target/mac-portmap/` of the worktree (gitignored). No speed number is measured here: `AGENTS.md` says a timing comes +from a runner in `tools/ref/`, and none runs on a Mac. The only speed numbers are byte-over-bandwidth ceilings marked +`[derived]`. + +Related documents: [`dgx-spark.md`](dgx-spark.md) (the port map this one is modelled on; its Appendix A is the P0 site +map for aarch64 Linux), [`../research/tierengine.md`](../research/tierengine.md) (item (h) is unified memory on GB10), +`AGENTS.md` ("Derive first", "Never"). + +## 1. Where the port stands + +| Phase (issue #593) | Done when | State (2026-10-21) | Left | +|---|---|---|---| +| **M0** native host build | the host chain builds and tests on aarch64-apple-darwin; every Linux-only site is behind one platform module; on x86 Linux nothing moved (`just ptx-scan` equal) | 14 of 21 non-device crates already build and test natively (`tools/mac-check.sh test`: 1477 passed, 0 failed, 210 ignored in 52 test binaries); 7 do not, from 45 compile errors at 38 source lines in 10 files (Part A) | the platform module and its callers (Part E, M0) | +| **M1** CPU path | NEON for `qdot`; the CPU engine behind `serve::Engine` | `qdot` has a scalar mirror for every kernel, so it builds off x86 once `threads` builds; a NEON form of the two generic kernels was prototyped and is bit-identical to the scalar mirror on 2400 rows (Part B); no CPU engine implements `serve::Engine` | off the critical path to a fast Metal decode (Part E) | +| **M2** Metal backend | one model's decode set runs on Metal and a Mac gate holds it | nothing exists. The CUDA side has 376 `#[kernel]` fns and 90 `#[cuda_module]` blocks; a Qwen3-MoE decode step launches 15 distinct kernels (Part C) | the device seam, the kernels, the chain (Part E) | +| **M3** a Mac gate tier | gates run on a Mac; `AGENTS.md` "Never" says so | the static tier and the pure crates' tests run (`just mac-static`, `just mac-test`); no runner for a gate or a timing exists on a Mac (Part D) | Part D | + +The order that reaches a fast Metal decode soonest is M0, then M2, then M3; M1 moves out of the critical path (Part E +gives the reasons). + +## 2. The machines + +The measured machine is the only Mac used in this round; the M5 rows are Apple's own statements, with the line copied. +No M5 machine was available, so nothing about the Neural Accelerators is measured. + +| Fact | Value | Source | +|---|---|---| +| Chip, memory, cores | Apple M1 Pro; `hw.memsize` 34359738368 B (32 GiB); 10 cores, 8 performance + 2 efficiency (`hw.perflevel0.logicalcpu` 8, `hw.perflevel1.logicalcpu` 2); 16 GPU cores | measured: `sysctl`, `system_profiler SPDisplaysDataType`; `machine-probe.log` | +| OS, Metal | macOS 27.0 (26A428); Metal 4 supported; GPU family `apple7`, not `apple8`/`apple9` | measured: `sw_vers`, the Swift `supportsFamily` probe (`probes/dev.swift`, `machine-probe.log`) | +| Page size, cache line | 16384 B (`hw.pagesize`), 128 B (`hw.cachelinesize`); L2 12582912 B for the performance cores, 4194304 B for the efficiency cores | measured: `sysctl` | +| CPU ISA | `hw.optional.arm.FEAT_DotProd` 1, `FEAT_I8MM` 0, `FEAT_BF16` 0, `FEAT_SME` 0 | measured: `sysctl` | +| GPU memory limits | `hasUnifiedMemory` true; `recommendedMaxWorkingSetSize` 26800603136 B (0.78 of `hw.memsize` [derived]); `maxBufferLength` 20100448256 B; `maxThreadgroupMemoryLength` 32768 B; `iogpu.wired_limit_mb` 0 | measured: Swift `MTLDevice` probe, `sysctl iogpu.wired_limit_mb` | +| Missing tools | `flock`, `timeout`, `gtimeout`, `setsid` are not on this Mac | measured: `which` | +| Toolchain | the pinned `nightly-2026-08-28` (`rustc 1.100.0-nightly (e457a7b0d 2026-08-27)`) builds natively; `aarch64-apple-darwin` is tier 1 | measured; https://doc.rust-lang.org/nightly/rustc/platform-support.html : "aarch64-apple-darwin ARM64 macOS (11.0+, Big Sur+)" under "Tier 1 with Host Tools" | + +Memory bandwidth and the largest unified memory, by chip class (Apple's figures; "up to" is Apple's wording): + +| Class | Bandwidth | Largest memory | Source and copied line | +|---|---|---|---| +| M1 Pro | 200 GB/s | 32 GB | https://www.apple.com/newsroom/2021/10/introducing-m1-pro-and-m1-max-the-most-powerful-chips-apple-has-ever-built/ : "M1 Pro can be configured with up to 32GB of fast unified memory, with up to 200GB/s of memory bandwidth" | +| M1 Max | 400 GB/s | 64 GB | same page: "M1 Max also offers a higher-bandwidth on-chip fabric, and doubles the memory interface compared with M1 Pro for up to 400GB/s" | +| M2 Ultra | 800 GB/s | 192 GB | https://www.apple.com/newsroom/2023/06/apple-introduces-m2-ultra/ : "M2 Ultra features 800GB/s of system memory bandwidth, far greater than anything found in a PC." | +| M4 | 120 GB/s | 32 GB | https://www.apple.com/newsroom/2024/10/apple-introduces-m4-pro-and-m4-max/ : "M4 supports up to 32GB of unified memory and has higher memory bandwidth of 120GB/s." | +| M4 Pro | 273 GB/s | 64 GB | same page: "M4 Pro supports up to 64GB of fast unified memory and 273GB/s of memory bandwidth" | +| M4 Max | 546 GB/s | 128 GB | same page: "M4 Max supports up to 128GB of fast unified memory and up to 546GB/s of memory bandwidth" | +| M5 | 153 GB/s | not stated on the page read | https://www.apple.com/newsroom/2025/10/apple-unleashes-m5-the-next-big-leap-in-ai-performance-for-apple-silicon/ : "a nearly 30 percent increase in unified memory bandwidth to 153GB/s" | +| M5 Pro | 307 GB/s | 64 GB | https://www.apple.com/newsroom/2026/03/apple-debuts-m5-pro-and-m5-max-to-supercharge-the-most-demanding-pro-workflows/ : "M5 Pro supports up to 64GB of unified memory with higher unified memory bandwidth up to 307GB/s." | +| M5 Max | 614 GB/s | 128 GB | same page: "M5 Max supports up to 128GB of unified memory with higher unified memory bandwidth up to 614GB/s." | +| **M5 Ultra** | **1.2 TB/s** | **512 GB** | https://www.apple.com/newsroom/2026/08/apple-introduces-new-mac-studio-with-m5-max-and-m5-ultra/ : "Combined with up to 512GB of unified memory and 1.2TB/s of memory bandwidth, 50 percent higher than before, Mac Studio lets users run massive models entirely on device" | + +M5 Ultra GPU: same page, "Its up-to-80-core GPU, the most powerful Apple silicon GPU ever, brings Neural Accelerators to the +Ultra chip for the first time, enabling up to 4.3x the peak AI compute performance when compared to M3 Ultra." The same +page says "Mac Studio with 512GB of unified memory is coming in late October." The M5 Ultra's CPU core count, GPU clock, +peak TFLOPS and the Neural Accelerator's rate per data type are not stated on the pages read: `[unverified]`. + +Neural Accelerators (the M5 generation's GPU matrix hardware): https://www.apple.com/newsroom/2025/10/apple-unleashes-m5-the-next-big-leap-in-ai-performance-for-apple-silicon/ : +"The 10-core GPU features a dedicated Neural Accelerator in each core, delivering over 4x peak GPU compute compared to M4, +and over 6x peak GPU compute for AI performance compared to M1." and "Developers can also build solutions for their apps +by directly programming the Neural Accelerators using Tensor APIs in Metal 4." More in section C.4. + +## Part A. Native compile census + +**What was run.** For each of the 21 non-device crates (the 14 that `python3 tools/recipes.py pure-crates` selects and +the 7 it rejects without a device root), in dependency order, `cargo check -p --all-targets --keep-going --color never` +with the pinned toolchain's `bin` first on `PATH` and `CARGO_TARGET_DIR` the worktree's `target/`, no `--target`, no env +file: the way `tools/mac-check.sh` runs its native `test` (header `tools/mac-check.sh:9-14`). `--keep-going` is added so one +red target does not hide the others of a crate. One log a crate: `target/mac-portmap/check/.log` and `.rc`. The 10 +device crates (`bloomery-gpu*`, `bloomery-app`) were not built: `AGENTS.md` "Never" forbids plain cargo on them. The +excluded crates (`crates/oxide-ice-unroll`, `fuzz`, `labs/hip-native`, `Cargo.toml:4`) were not looked at. + +### A.1 Per crate (layer 1: the tree as it is) + +| Crate | rc | Own errors | Distinct causes | +|---|---|---|---| +| bloomery-levers, -jinja, -sampler, -verbs, -models, -refset, -runtime, -tokenizer, -vision, -embed, -placement, -hf, -decision, -serve | 0 | 0 | none (14 crates, all targets: lib, bins, tests, examples) | +| bloomery-gguf | 101 | 1 (lib test only; the lib and every other target build) | `libc::RUSAGE_THREAD` | +| bloomery-threads | 101 | 20: 16 in the lib, 4 more in the lib test | `libc::cpu_set_t`, `CPU_SET`, `sched_setaffinity`, `sched_getaffinity` | +| bloomery-qdot | 101 | 0: stops at `bloomery-threads` | the dependency's 16 | +| bloomery-engram | 101 | 0: stops at `bloomery-threads` | the dependency's 16 | +| bloomery-engram-lab | 101 | 0: stops at `bloomery-threads` | the dependency's 16 | +| bloomery-model | 101 | 0: stops at `bloomery-threads` | the dependency's 16 | +| bloomery-model-hf | 101 | 0: stops at `bloomery-threads` | the dependency's 16 | + +Layer 1 totals: 14 crates rc 0, 7 crates rc 101; 21 distinct errors (`gguf` 1, `threads` 20), 2 distinct causes. The five +crates behind `threads` show nothing of their own, so a second and a third layer were measured in a scratch copy of the +tree (outside the worktree, a scratch directory's `layer2` and `layer3`; the copies are not part of +the repository and the worktree is untouched): + +- **Layer 2**: `threads`' affinity calls replaced by stubs (a `shim` module returning refusal). Logs + `target/mac-portmap/check-layer2/`. `qdot` builds, rc 0 on every target (lib, 2 bins, 3 test files). `engram` has 5 errors in the lib, + `model` 14 in the lib (`engram-lab` and `model-hf` stop at those), and `threads`' own `tests/helper.rs` 4. +- **Layer 3**: also the `libc` items of layer 2 stubbed and the three type mismatches cast. Logs + `target/mac-portmap/check-layer3/`. `threads`, `qdot`, `engram-lab`, `model` and `model-hf` build on every target (rc 0), including + `model`'s bins and tests. `engram` has one error left in `tests/engram.rs:885`. + +So the host chain's compile surface for the Mac is **45 errors at 38 source lines in 10 files, 8 distinct causes** (the two type-shape breaks counted as one); with +those fixed, every non-device crate builds natively. The stubs were scratch only; `qdot` needed no change. + +### A.2 Per site + +`lib` is the crate's library; `test` is a `#[cfg(test)]` module or a `tests/` file. "macOS form" is the replacement the +site needs; a form marked `[unverified]` was not tried. + +| Site | Crate | Item | What it does | Kind | macOS form | +|---|---|---|---|---|---| +| `crates/threads/src/lib.rs:757` (`mask_cpus`, :750) | threads | `libc::sched_getaffinity` | reads the calling thread's allowed cpus; the pool cuts its topology by it | lib | none: macOS has no affinity mask; the mask is every online cpu (`hw.logicalcpu`) | +| `crates/threads/src/lib.rs:810-823` (`pin`, :809) | threads | `cpu_set_t`, `CPU_SET`, `sched_setaffinity` | pins a pool worker to one cpu; a refusal is returned, not fatal (:806-808) | lib | none: `thread_policy_set(THREAD_AFFINITY_POLICY)` returned `KERN_NOT_SUPPORTED` (46) on this M1 Pro (measured); return `Unsupported`, the pool floats. `pthread_set_qos_class_self_np(QOS_CLASS_USER_INTERACTIVE, 0)` returned 0 (measured) and is the only placement lever `[unverified]` for P versus E cores | +| `crates/threads/src/helper.rs:217-226` (`float_everywhere`, :216) | threads | same | widens the caller's mask to every cpu | lib | none (no mask to widen): `true` | +| `crates/threads/src/helper.rs:233-242` (`pin_to`, :232) | threads | same | pins the helper thread | lib | as `pin` | +| `crates/threads/src/lib.rs:844-848` (`on_cpus`, :839) | threads | same | a test thread pins itself to read its own mask | test | skip by name on macOS | +| `crates/threads/tests/helper.rs:15-19` (`set_mask`, :10) | threads | same | the helper gate sets the opener's mask | test | skip by name on macOS (the file says "Linux only", :3) | +| `crates/gguf/src/lib.rs:1322` (`thread_faults`, :1318) | gguf | `libc::RUSAGE_THREAD` | counts one thread's page faults for a populate test | test | `getrusage` on macOS takes only `RUSAGE_SELF` and `RUSAGE_CHILDREN` (`man 2 getrusage`, this Mac); a per-thread fault counter was not found `[unverified]`: count process faults in a single-test run | +| `crates/engram/src/direct.rs:39` (`open_direct`, :35) | engram | `libc::O_DIRECT` | opens a file to bypass the page cache | lib | `open`, then `fcntl(fd, F_NOCACHE, 1)` (`man 2 fcntl`: "F_NOCACHE Turns data caching off/on."); the alignment contract (`DIRECT_ALIGN`) is Linux's, whether `F_NOCACHE` needs it `[unverified]` | +| `crates/engram/tests/engram.rs:885` (:871) | engram | `libc::O_DIRECT` | a gate reads `/proc/self/fdinfo` flags | test | skip by name | +| `crates/engram/src/lib.rs:546-550` (`fadvise_dontneed`, :542) | engram | `posix_fadvise`, `POSIX_FADV_DONTNEED` | drops a file range from the page cache after an eviction | lib | no equivalent found; a no-op with a named note `[unverified]`. `F_NOCACHE` affects later reads only | +| `crates/engram/src/lib.rs:702` (`populate`, :701) | engram | `MADV_POPULATE_READ` | faults a row's pages in | lib | `man 2 madvise` on this Mac lists no `MADV_POPULATE`: read one byte a page (the Linux path already has this fallback for kernels older than 5.14, `crates/model/src/placement/host_lock.rs:889-937`) | +| `crates/engram/src/lib.rs:755` (`mincore_chunk`, :749) | engram | `libc::mincore` | counts resident pages | lib | the vector is `*mut c_char` here (`man 2 mincore`: `char *vec`): cast `.cast()`; the bit is `MINCORE_INCORE` | +| `crates/model/src/fileio.rs:29` (`free_bytes`, :19) | model | `statvfs` fields | free bytes on a filesystem | lib | `f_bavail` is `u32` and `f_frsize` `u64` on macOS: widen both | +| `crates/model/src/fileio.rs:37-42` (`rename_noreplace`, :33) | model | `renameat2`, `RENAME_NOREPLACE` | atomic rename that refuses an existing target | lib | `renamex_np(from, to, RENAME_EXCL)` (`man 2 rename`) | +| `crates/model/src/placement/host_lock.rs:867-871` (`drop_pages`, :818) | model | `posix_fadvise`, `POSIX_FADV_DONTNEED` | drops pages of a set's file | lib | as `fadvise_dontneed` | +| `crates/model/src/placement/host_lock.rs:909` (`advise_populate`, :901) | model | `MADV_POPULATE_READ` | faults a span in | lib | as `populate` (the `Populated::Touched` path exists, :889-937) | +| `crates/model/src/placement/host_lock.rs:1218` (`resident_pages`, :1207) | model | `libc::mincore` | resident page count | lib | as `mincore_chunk` | +| `crates/model/src/r8file.rs:512-517` (`sync_and_drop`, :510) | model | `sync_file_range`, `SYNC_FILE_RANGE_*` | writes a range out before dropping it | lib | `fcntl(F_FULLFSYNC)` after the writes (`man 2 fcntl`: "F_FULLFSYNC Does the same thing as fsync(2) then asks the …") | +| `crates/model/src/r8file.rs:528` (`drop_cached`, :525) | model | `posix_fadvise`, `POSIX_FADV_DONTNEED` | drops written pages | lib | as `fadvise_dontneed` | + +Sites that compile on a Mac and still do the wrong thing there are not compile errors, so the census cannot see them. +Two kinds were listed by reading the code: + +- **Linux files** (`/proc`, `/sys`): 40 lines in 12 files (a grep, `target/mac-portmap/proc-sys-sites.txt`; a few are message strings or `#[ignore]` reasons, not reads). In the + library and the load path: `crates/threads/src/topology.rs:83-129` and `helper.rs:177` (sysfs topology; with no sysfs + `detect_topology` returns one group of one cpu, `topology.rs:99-112`, so the pool's default is one thread `[derived]`); + `crates/model/src/placement/workstation.rs:35-90` (`/proc/meminfo`, `/proc/self/cgroup`, `/sys/fs/cgroup`: the host room). + In tools and tests only: `crates/model/src/bin/bench_v41_host.rs` (11 lines), `crates/engram-lab/src/bin/engram-rate.rs` + (6) and `src/lib.rs:51-58`, `crates/model/src/bin/bloomery-decode.rs:143`, `crates/model/src/ops.rs:6573`, + `crates/serve/tests/limits.rs:26`, `crates/threads/tests/pool.rs:208,351`. +- **Units and sizes**: `ru_maxrss` is KB on Linux and bytes on macOS (`man 2 getrusage`: "the maximum resident set size + utilized (in bytes)"), read at `crates/model-hf/src/bin/fixture.rs:387` and `crates/model-hf/tests/fixture.rs:852`. + `ROW_PAGE` is 4096 (`crates/placement/src/placement/workstation.rs:141`) and `host_available_read` refuses any other page + size before it reads anything, even with `BLOOMERY_HOST_ROOM` set (`crates/model/src/placement/workstation.rs:24-28`); this Mac's + page is 16384 (measured), so every placement plan on a Mac stops there until the page becomes a parameter. The row-room + arithmetic uses the constant too (`crates/placement/src/placement.rs:3055-3062`). The cache line is 128 B here (measured); + code that pads to 64 B for false sharing still works. + +### A.3 Against `recipes.py pure-crates` and the Spark port + +- `tools/recipes.py pure-crates` rejects exactly the 17 crates found: 10 for the device root and the 7 above, each with its + first Linux-only item. It rejects `bloomery-gguf` for `RUSAGE_THREAD` at `gguf/src/lib.rs:1322`; the census agrees that the + failure is in the lib test alone (`cargo test -p bloomery-gguf` does not build, the lib does). +- The rule's name list (`tools/recipes.py:5512-5525`) does not hold `libc::O_DIRECT`, and no name list can hold the two + type-shape breaks (the `mincore` vector, the `statvfs` field widths). Each sits in a crate the rule already rejects, so + today nothing is hidden; a crate fenced by name would still fail natively. +- The rule's `#[target_feature(` pattern (`tools/recipes.py:5514`) does not look at the feature: a NEON body under + `cfg(target_arch = "aarch64")` is true on the Mac and would reject `qdot` again (Part D). +- The Spark port fenced every x86 site (`dgx-spark.md` Appendix A, lines 1485-1520): `qdot` under `cfg(target_arch = + "x86_64")` (`crates/qdot/src/lib.rs:62`), the prefetch in `crates/model/src/ops.rs:5875-5893`, `deepseek2/attn.rs`, the bin and + the tests; those hold on the Mac (the census shows no x86 error). The Spark port left the libc items out because they are + POSIX on Linux ("not blockers on DGX OS `[unverified until a compile]`", `dgx-spark.md` Appendix A). They are the whole + of the Mac's census: 45 errors, none of them an x86 item. + +## Part C. The Metal backend (M2), design + +The decode step on this engine is bandwidth-bound on its weights and launch-bound on its chain; prefill is compute-bound. +The design below is cut to that: each choice states what it costs or gains, `[derived]` where it is computed. Where +the number needs a Mac measurement that cannot be derived, the term is named and left open (section C.5). + +### C.1 The device seam + +**What `crates/gpu` exposes, by kind** (the callers are the family crates `gpu-deepseek41`, `-glm5next`, `-mimo2`, `-solar2`, +`-vision`, `-gemma4a`, `-gemma-embedding2`, `gpu-gates` and `app`; counts are greps over those crates and `crates/gpu`, the +command is in `target/mac-portmap/surf.sh`; they are `[derived from code]`): + +| Kind | What the code uses | Count | Examples | +|---|---|---|---| +| Device memory | `cuda_core::DeviceBuffer`: `zeroed` 449, `from_host` 201; typed windows | 280 import lines | `crates/gpu/src/arch/qwen3moe/scratch.rs:90`, `crates/gpu/src/tensor.rs:41` (`window`), `crates/gpu/src/weights.rs:54` (`DevWeight`) | +| Pinned and mapped host memory | `PinnedHostBuffer` 16 sites; `cuMemHostAlloc` 22; `cuMemHostGetDevicePointer` 8; a mapped page shared with a tier node | 46 | `crates/gpu/src/checkpoint.rs:330`, `crates/gpu/src/graph.rs:491` (`MappedHost`), :544 | +| Copies and fills | `cuMemcpyHtoDAsync` 18, `DtoH` 10, `DtoD` 4, `Memcpy2D` 8, `cuMemsetD32Async` 4 | 44 | `crates/gpu/src/graph.rs` (`TapRows` copies are graph nodes) | +| Streams | `CudaStream` (the stream every launch takes), per-role streams | 250 import lines | `crates/gpu/src/lib.rs:2281` (`StreamRole`), :2290 (`role_stream`) | +| Events, host sync | `CudaEvent`, `.record` 37, `.wait` 37, `.synchronize()` 107, `cuLaunchHostFunc` 8, `cuStreamBatchMemOp` 8 | | `crates/gpu/src/checkpoint.rs:93`, `graph.rs:341` (`Branch`/`Forked`), :781-824 (`op_wait_geq`, `op_write`), :882 (`HostFlags`) | +| Graph capture and replay | `Gpu::capture`, `Graph::launch`, node introspection (`cuGraphGetNodes` 5, `cuGraphLaunch` 4) | | `crates/gpu/src/lib.rs:2776`, `crates/gpu/src/graph.rs:51,82,186,204` | +| Kernel launch | `#[cuda_module]` 90 blocks, `pub fn enqueue_*` 316, `prepare_*` per kernel; arguments are typed buffers and `u32`s | 316 | `crates/gpu/src/lib.rs:3021-3052` (`q4k_gemv_prefetch`/`q4k_gemv`/`q4k_gemv_mcol`) | +| Module loading | embedded device module loaded once per context | | `crates/gpu/src/lib.rs:2425` (`bundle_module`) | +| Fault word | one `u32` array per `Gpu`, kernels raise by atomic min | | `crates/gpu/src/fault.rs` (`ptx_asm!`), `crates/gpu/src/lib.rs:2549-2617` | +| Device query, memory room | `census`, `mem_info`, `allocation_granularity`, integrated flag | | `crates/gpu/src/lib.rs:2146`, :2692, :2705 | + +Device-side, the 376 `#[kernel]` fns use what NVIDIA exposes: `cuda_device::wmma` (`mma_m16n8k16` 22, `ldmatrix` 14), +`dotprod::dp4a_s32` (`crates/gpu/src/cores.rs:14,224`, `quant_sel.rs`, `kquant/walk.rs`, `q5.rs`, `head_argmax.rs`), +`async_copy` 8, `warp::shuffle_xor_f32`, `float::{fma_rn,mul_rn,add_rn}_f32` and `ptx_asm!` in 7 places in 5 files +(`fault.rs`, `cores.rs`, `flash_gqa_prefill.rs`, `linear/delta.rs`, `gemm/gemm32.rs`). Metal has no `dp4a`: `dot(char4,char4)` +does not compile, and `simdgroup_matrix` takes `half`, `bfloat` and `float` only (both measured on this M1 Pro, +`target/mac-portmap/msl-probe.log`; the spec line is https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf , MSL +4.1, section 2.4: "Metal supports the following SIMD-group matrix type names, where T is half, bfloat (in Metal 3.1 and later) or +float and Cols and Rows are 8"). An atomic min on a device `atomic_uint` compiles (measured), so the fault word moves. + +**The proposal.** Two cuts, in this order. + +1. **A device layer, one owner, selected at build time.** A new crate, `bloomery-device`, owns exactly the kinds above that + are resource and scheduling types: `Buffer`, `Pinned`, `Stream`, `Event`, `Graph` (capture, replay, node count), + `HostFlags` and `MappedHost`, the fault word, the device census. Its public API names no `cuda_core` and no Metal type. + The CUDA implementation wraps today's types without changing their behaviour; the Metal one is new. The family crates + import the layer's names instead of `cuda_core`'s (about 280 import lines, a rename). It is selected by a Cargo feature, + not by a trait object: the hot path is 12 L + 4 launches a token (C.2), and an indirect call per launch is a cost + with no gain, since a build holds one backend. This follows `AGENTS.md` "common code first": the owner of the resource + types is one place; no common owner grows a branch per backend, because the branch is the module boundary. +2. **The chain walk stays common; the launches go behind a small op set.** `crates/gpu/src/arch/qwen3moe/dispatch.rs` mixes the + order of the chain (`layer` :422, `mixer` :439, `ffn` :989: what runs after what, from the `LayerPlan`) with the launches + (which kernel an op takes by weight type). The order and the plans (`plan.rs`, `program.rs`, the arena sizing in + `scratch.rs`) do not name a backend except through `DeviceBuffer`. The launches are the 15 ops of C.2. The cut is one + trait of those ops, `ChainOps`, with an associated buffer type from the device layer, implemented by the CUDA kernels + (the existing wrappers, moved) and by the Metal ones; the walk is generic over it. A backend refuses an (op, weight + type) pair it does not implement by a named error, as `GateUpEntry`/`DownEntry` do today + (`crates/gpu/src/arch/qwen3moe/dispatch.rs:1093-1230`). The 316 `enqueue_*` wrappers of other families stay as they are: + they move to the op set only when a second family reaches Metal. The alternative, a Metal copy of the chain + (`dispatch.rs` is 1513 lines), is the per-model copy `AGENTS.md` allows only for a large measured speed gain; no such + gain is known, so the copy is not proposed. Part F asks the maintainer to confirm the cut. + +The serve side already has a seam of this size: `trait Seat` (`crates/gpu-gates/src/bind.rs:1268-1515`, 12 required methods) +is what `SeatEngine` (:1519, the only `serve::Engine` over a real model, impl at :2983) drives. `bind.rs` names 21 lines of +`gpu`/`cuda`/`nvidia` items and its `Vocab` (:81, the `serve::Tokenizer`) is in the device-linked crate `gpu-gates`. For the +first Metal decode a CLI that feeds ids needs none of it; for `bloomery-serve` on a Mac they lift into a device-free crate (section B.2). + +**Fault, snapshots, checkpoints.** The fault word moves as is (atomic min on a shared `atomic_uint`, read back with the +token). `Seat::snapshot`/`resume` (`bind.rs:1445,1465`) and `crates/gpu/src/checkpoint*` use CUDA events and pinned staging; +the Metal engine's first version answers them with a named refusal and the serve gates that need them are deferred by name. + +### C.2 The Qwen3-MoE decode kernel set `[derived from code]` + +Derived from `crates/gpu-gates/src/bin/generate_qwen3moe.rs` (the decode CLI, header :1-60), which opens a model through +`Qwen3moeModel` and runs `dispatch::enqueue_chain` (`crates/gpu/src/arch/qwen3moe/dispatch.rs:96`, a `Program` walk at +`m` = 1). The case read: a `qwen3moe` file (no shared expert, head 128 NeoX attention), whole on the card, f16 K/V, the +common Q4_K/Q6_K mix. A launch per layer is `pass_launches` (`dispatch.rs:303-382`): 1 (`norm_quant`) + 1 (fused q/k/v) ++ 3 (K/V append, flash segment, flash merge) + 2 (activation quantizer, fused output projection with the residual) + +5 (router, gate·up with SwiGLU, quantizer, down, combine) = **12 launches a layer**, plus 1 per layer whose V is Q6_K, +plus the front (1) and the head (3): **12 L + 4 launches a token** for L layers. Each kernel below has a `_prefetch` twin +for `Gpu::row_prefetch()` (`crates/gpu/src/lib.rs:2673`), chosen at launch; the table lists the base. + +| # | Kernel | Where | What it computes | Quant types it takes | +|---|---|---|---|---| +| 1 | `embed_rows_q4k` | `crates/gpu/src/elem.rs:682` | token row → `x`; also writes position and live key count | Q4_K; siblings `_q6k` :898, `_q5k` :849, `_q3k` :946, `_q5_0` :760; Q8_0 through `q38` (`dispatch.rs:504-519`) | +| 2 | `norm_quant` | `crates/gpu/src/fused.rs:99` | RMSNorm of `x` and the q8_1 quantization of the result | f32 → q8_1 | +| 3 | `qwen3moe_qkv_q4k` | `crates/gpu/src/arch/qwen3moe/proj.rs:56` | fused Q4_K q/k/v gemv against the q8_1 activation | Q4_K; V in Q6_K: `q6k_gemv` (`lib.rs:1340`); other: `q4k_gemv` (`lib.rs:1086`) | +| 4 | `head_norm_neox_append` | `crates/gpu/src/rope_neox.rs:875` | per-head Q/K RMSNorm, NeoX rope, K/V row appended to the cache | f16 K/V; Q8 cache: `_q8` :1265 | +| 5 | `gqa_flash_seg_mma` (scalar twin `gqa_flash_seg`) | `crates/gpu/src/flash_gqa.rs:3230` (:2956) | split-K flash attention segment over the live keys | f16 K/V | +| 6 | `gqa_flash_merge` | `crates/gpu/src/flash_gqa.rs:3477` | merges the segments' partial softmax | f32 | +| 7 | `q3k_quantize_q8_1` | `crates/gpu/src/lib.rs:912` | q8_1 quantization of the attention rows (the same kernel quantizes the SwiGLU rows at 11, and the head's input) | f32 → q8_1 | +| 8 | `qwen3moe_o_resid_q4k` | `crates/gpu/src/arch/qwen3moe/proj.rs:355` | Q4_K output projection plus the residual add | Q4_K | +| 9 | `qwen3moe_router_norm` | `crates/gpu/src/arch/qwen3moe/router.rs:1420` | FFN RMSNorm and q8_1 (every block), the router logits, then the ticketed routing (`publish_route`); one launch at `m` = 1 | f32 router weights | +| 10 | `qwen3moe_gate_up_swiglu_q4k` | `crates/gpu/src/arch/qwen3moe/experts.rs:136` | gate and up of the chosen experts with SwiGLU | Q4_K (other types: `elem::swiglu` :1137 and the `_sel` family: Q5_K, Q8_0, IQ2_XS/IQ3_XXS/IQ3_S/IQ4_XS, MXFP4, NVFP4, `dispatch.rs:1093-1230`) | +| 11 | (7 again) | | quantizes the SwiGLU rows | | +| 12 | `q4k_gemv_sel` / `q6k_gemv_sel` | `crates/gpu/src/q4k_sel.rs:344`, `q6k_sel.rs:166` | down projection of the chosen experts | Q4_K / Q6_K (also Q5_K, Q5_0, Q5_1, Q8_0, IQ*, MXFP4, NVFP4) | +| 13 | `qwen3moe_combine` | `crates/gpu/src/arch/qwen3moe/experts.rs:280` | weighted sum of the experts, plus the residual | f32 | +| 14 | `rms_norm` | `crates/gpu/src/elem.rs:1004` | the head's final RMSNorm | f32 | +| 15 | `qwen3moe_head_q6k_argmax` | `crates/gpu/src/arch/qwen3moe/head_argmax.rs:80` | Q6_K output projection with the argmax folded in | Q6_K (another head type: the shared head's gemv and `argmax_fault` :1410) | + +**Count: 15 distinct kernels** for a Q4_K/Q6_K file (16 with a Q6_K V), 7 of them (`qkv`, `o_resid`, `router_norm`, `gate_up_swiglu`, both `_sel`, the head) with a `_prefetch` twin, `[derived from +code]`. The family also holds the prompt-side set (`flash_gqa_prefill`, the grouped GEMM of `wide.rs`, `gemm/`) that the +decode does not launch; a file of other quant types adds the `_sel` variants named in the rows. The entry bodies of these 15 +total 1436 lines `[derived]` (summed `#[kernel]` extents of the rows above, scalar flash counted once); the device cores they +call (`cores.rs` 1455 lines, `quant_sel.rs` 1009, `route_body.rs` 806, `route_core.rs` 106) are shared with other kernels, +so the port is about 1.4k lines of entries plus the part of 3.4k lines of cores these reach, `[derived]` order of magnitude. +Qwen3.6 (`qwen35moe`, the fixture family, `crates/model/src/arch/registry.rs:124`) runs the same `dispatch::layer` and adds the +delta rule's three kernels, the gated router and the gated quantizer (`body.rs:68-76`, "The kernels only Qwen3.6's layers launch"). The `qwen3moe` registry row has no fixture (`registry.rs:100-104`). + +### C.3 Kernel authoring for Metal + +Six options, each judged on: exists and maintained, one source for CUDA and Metal, cost per kernel, risk to the ggml +bit band (`AGENTS.md`: kernels match ggml within each crate's RESULTS band), and the CUDA-graph equivalent. The Metal +equivalent of a CUDA graph for all of them is an indirect command buffer: https://developer.apple.com/tutorials/data/documentation/metal/mtlindirectcomputecommand.json , +"A compute command in an indirect command buffer."; https://developer.apple.com/tutorials/data/documentation/metal/indirect-command-encoding.json , "Store draw +commands in Metal buffers and run them at a later time on the GPU, either once or repeatedly." The contract matches the +graph's: fixed buffers in place, fixed grids, the per-replay values in a device record (`StepParams`, +`crates/gpu/src/arch/qwen3moe/body.rs:1000-1003`; the graph's rule at `crates/gpu/src/graph.rs:47-50`). Whether Metal 4 keeps the ICB +as the replay path for compute is in the pages above; the cost per replayed dispatch is a Mac measurement (C.5). + +| Option | Exists, maintained | One source for CUDA and Metal | Cost per kernel | Risk to the bit band | Tensor ops / graph | +|---|---|---|---|---|---| +| **Hand-written MSL** | MSL 4.1 spec dated 2026-06-04 (the PDF above) | no: a second source per kernel; the CUDA kernels stay | the 15 decode kernels are 1.4k entry lines in Rust (C.2); MSL is the same order. A kernel's loop order is copied, so its bits are controlled | lowest: `#pragma METAL fp math_mode(safe)` and `fp contract(off)` compile (measured). The spec says "The default is fast." and "By default, the compiler allows floating-point contractions. For example, a*b+c may be converted to a single fused-multiply-add." (MSL 4.1, the compiler options, pages 15-16 of the PDF) | full: `simdgroup_matrix`, MPP `matmul2d`, ICB, `MTLSharedEvent` | +| **Rust to SPIR-V (rust-gpu) to MSL (SPIRV-Cross)** | https://github.com/Rust-GPU/rust-gpu pushed 2026-10-16: "Currently only [SPIR-V] support is planned" (the README's own link markup); https://github.com/KhronosGroup/SPIRV-Cross pushed 2026-10-14: "Convert SPIR-V to readable, usable and efficient Metal Shading Language (MSL)" | no: kernels are rewritten against `spirv_std`, not `cuda_device` | the rewrite of a hand-written MSL, plus a two-step compile | high: SPIRV-Cross output is compiled by Apple's compiler at its default math mode; no way to name `simdgroup_matrix` or MPP from SPIR-V | none for tensor ops; ICB by the host | +| **CubeCL** | https://github.com/tracel-ai/cubecl pushed 2026-10-20: "A single `#[cube]` Rust function is compiled on demand and runs on CUDA, HIP, Metal, Vulkan, WebGPU, or CPU"; https://github.com/tracel-ai/cubecl/tree/main/crates/cubecl-metal README: "Metal 3+ required (atomic float and simdgroup matrix support)" | yes, but not with cuda-oxide: its own `#[cube]` DSL and codegen replace the 376 kernels' language | a rewrite per kernel into the DSL | high: the codegen's op order and contraction are CubeCL's; Metal tensor ops: not found in the crate list or README `[unverified]` | its own runtime, not ICB `[unverified]` | +| **Slang** | https://github.com/shader-slang/slang pushed 2026-10-21: Metal "experimental\*"; "Slang currently supports generating vertex, fragment, compute, task and mesh shaders for Metal." | yes, in a new language (CUDA output is C++) | a rewrite per kernel in Slang | high, as SPIRV-Cross | none known `[unverified]` | +| **MLX kernels as a library** | https://github.com/ml-explore/mlx pushed 2026-10-20 (MIT): "MLX is an array framework for machine learning on Apple silicon"; kernels in `mlx/backend/metal/kernels/` include `*_nax` files | no | none to write; but `quantized.h` is group-affine quantization (`group_size`, 0 hits for `q4_K`/`gguf`/`block_q`): GGUF K-quants are not served, so a model would be requantized | n/a: a different numeric (scales and biases per group) | uses TensorOps for the Neural Accelerators (https://machinelearning.apple.com/research/exploring-llms-mlx-m5 : "MLX leverages the Tensor Operations (TensorOps) and Metal Performance Primitives framework introduced with Metal 4 to support the Neural Accelerators’ features.") | +| **ggml-metal kernels as a library** | https://github.com/ggml-org/llama.cpp pushed 2026-10-21 (MIT): "Apple silicon is a first-class citizen - optimized via ARM NEON, Accelerate and Metal frameworks"; local checkout `36a73916e`: 978 `host_name` entries in 14841 lines of `ggml/src/ggml-metal/kernels/*.metal,.h` | no | none to write for 30+ quant types: `kernel_mul_mv_id_q4_K_f32` and the rest (`kernels/mul_mv.metal:3362` and below); the host side (`ggml-metal-ops.cpp`) is tied to the ggml graph and would be rewritten | different numerics: the K-quant mat-vec reads f32 activations (`kernels/mul_mv.metal:1546`, `float * y`), not q8_1; so its bits are not the card engine's, and their band against ik is not measured `[unverified]` | tensor API behind `GGML_METAL_HAS_TENSOR` (`kernels/mul_mm.metal:13`); ggml's own comment: "disable the tensor API by default for old chips because with the current implementation it is not useful - M2 Ultra: ~5% slower - M4, M4 Max: no significant difference" (`ggml-metal-device.m:1151-1153`) | + +**Recommendation: hand-written MSL, one kernel at a time, structured after ggml-metal's where the numerics allow.** +Reasons, each against the goal (speed on the M5 generation): + +- The two things that decide speed on M5, the Neural Accelerators and the dispatch chain, are reached only by MSL written + against MPP `matmul2d` and by an ICB the engine encodes itself. A cross-language route (SPIR-V, Slang) has no door to + either; CubeCL has none found. +- The numeric contract is ours to hold. A kernel that must match a CUDA kernel bit for bit (the integer-exact ones, C.4) can + only be written with the loop order and the fused-multiply-add placement under our control, which the compile flags + above give and a generator does not. +- ggml-metal is the structure to learn from and the oracle for what is fast on Apple silicon; it is MIT, so a kernel adapted + from it carries its notice (`THIRD_PARTY_NOTICES.md`). It is not the runtime: its activations are f32, so using it leaves + the engine's q8_1 path and makes Metal's results a different family from the CUDA engine's. +- Cost: the second source per kernel. It is bounded: 15 kernels for the first model (C.2), each gated against a reference + (Part D). The decision between "q8_1 activations as on CUDA" and "f32 or f16 activations as ggml-metal" is the first open + question (Part F), since it sets what the first kernels are. + +### C.4 Designing for the M5 Ultra + +**What M5 adds, and how Metal exposes it.** A Neural Accelerator in each GPU core, programmed through Metal 4 tensor APIs +(section 2). On the shader side that is `` and `MetalPerformancePrimitives`, in the SDK on this Mac: +`MetalPerformancePrimitives.framework/Headers/MPPTensorOpsMatMul2d.h` (macOS 27.0 SDK, copied to +`target/mac-portmap/MPPTensorOpsMatMul2d.h.txt`). Its header lists the data type combinations of `matmul2d` (left, right, +destination), among them `half half half`, `half int8_t half`, `int8_t int8_t int32_t`, `uint8_t uint8_t int32_t`, `bfloat bfloat float`, +`int8_t int4b_format int32_t`, `int8_t int2b_format int32_t`, `half metal_fp4_e2m1_format half` and the fp8 formats, and a descriptor +flag `relaxed_precision` ("set it to true to allow implementation to sacrifice accurancy for performance", sic). The host +side is `MTLTensor` and the `MTL4*` encoders (https://developer.apple.com/tutorials/data/documentation/metal/mtltensor.json , "A resource representing a +multi-dimensional array that you can use with machine learning workloads."; the shader can also build a `tensor_inline` from a device +pointer, as the header's example and the probe below do). + +**Measured on the M1 Pro** (`target/mac-portmap/mpp-int8-probe.log`, `mpp-types-probe.log`): an `matmul2d` of `int8_t × int8_t → int32_t` +(64 × 32 tile, 4 simdgroups, `tensor_inline` views of device buffers) compiles and builds a pipeline on an `apple7` GPU, runs, +and equals the CPU's integer product on all 8192 outputs (M = 128, N = 64, K = 96); `half×half→float`, `bfloat×bfloat→float`, +`float×float→float`, `half×int8_t→float` and `int8_t×half→float` also build. So the tensor-op path can be written and checked +bit for bit on the M1 dev machine. Whether it is fast there is not what the check shows; ggml's comment above says the tensor API is not +faster than `simdgroup_matrix` on M2 Ultra and M4. Whether the M5's Neural Accelerators run `int8` at a higher rate than `half`, and the +rate itself, is not stated on any page read: `[unverified]`. Tuning the matrix kernels therefore needs an M5-class machine (Part F). + +Apple's own description of where the matrix hardware pays (https://machinelearning.apple.com/research/exploring-llms-mlx-m5): "In LLM +inference, generating the first token is compute-bound, and takes full advantage of the Neural Accelerators." and "Generating +subsequent tokens is bounded by memory bandwidth, rather than by compute ability." That is this engine's split, so: + +| Choice | What it is | Speed: cost or gain | +|---|---|---| +| **Decode streams weights at bandwidth** | one-row (GEMV) kernels, no matrix hardware; the Q4_K block is read once, decoded in registers | the floor is bytes over bandwidth. For Qwen3.6-35B-A3B `Q4_K_M` the weight bytes a token are 1.8 GB (`dgx-spark.md:212`, from 3 B active parameters at 4.85 bits `[derived]`), so the ceiling is 1.8 GB / 1.2 TB/s = **1.5 ms a token on an M5 Ultra** `[derived]`, against 9.0 ms on this M1 Pro (200 GB/s), 3.3 ms on an M4 Max (546), 2.9 ms on an M5 Max (614). Datasheet figures; the achievable fraction is a measurement (C.5) | +| **Int8 dot without `dp4a`** | the q8_1 path's integer dot on the ALUs: 4 `int` multiply-adds a `dp4a`; or fp32 FMAs of the same products | exact either way. fp32 holds integers below 2^24 exactly; a 32-value q8 partial sum is at most 32 × 127 × 127 = 516128 `[derived]`, so an fp32 FMA chain per 32-value block gives the same integer as `dp4a`, and the scales come after. At 1.2 TB/s a Q4_K weight stream (256 values in 144 B, 1.78 values per byte `[derived]`) needs 2.1 × 10^12 multiply-adds a second `[derived]`: peak GPU throughput is not stated on the pages read (`[unverified]`), so whether decode stays memory-bound on M5 Ultra is the second open measurement | +| **Prefill and multi-row verify on tensor ops** | the prompt GEMM, the grouped expert GEMM and the draft verify rows (`m` up to `GEMV_COLS`, `crates/gpu/src/arch/qwen3moe/wide.rs`) as `matmul2d` over weight tiles decoded to `half` (or `int8` with the scales after) in threadgroup memory | the CUDA engine already does this with `mma_m16n8k16` over dequantized tiles (`gemm/kernels.rs`, `grouped.rs`); the Metal kernel is the same shape. Apple: "over 4x the peak GPU compute for AI" for M5 over M4. Prompt rows are compute-bound once the weights are read once per ubatch: a ubatch of P rows costs 2 × 3 B × P FLOP against one pass over the 21.2 GB file `[derived: dgx-spark.md:212]`, so the break-even P is `F_peak × 17.7 ms / 6 GFLOP` with the file read at 1.2 TB/s: for every 10 TFLOPS of achieved matrix rate it is about 30 rows `[derived]`. `F_peak` is `[unverified]` | +| **Chain structure: fewer, fatter dispatches** | the decode is 12 L + 4 launches (C.2). At 1.5 ms of weight time, a gap of g µs between dispatches costs (12 L + 4) g: for L = 48 `[unverified input: the layer count was not read from a file in this round]` that is 580 g µs, so g must stay under 2.6 µs `[derived]` for the chain not to dominate | the term the CUDA engine measured as `c_node` (`AGENTS.md`, "launch count only": Δt = ΔN × c_node). On Metal it is not known. Levers, in the order the CUDA engine used: fuse where it is free (the engine already fuses qkv, o+residual, router+norm, gate·up+SwiGLU), encode the whole step into an ICB once, and issue independent dispatches under `MTLDispatchTypeConcurrent` with barriers only between dependent ones. A fold costs the repeated work times the blocks (`AGENTS.md` "fold a launch into a neighbour") | +| **Weights streamed in place** | wrap the mmap'd GGUF in `MTLBuffer`s without a copy (`newBufferWithBytesNoCopy`; https://developer.apple.com/tutorials/data/documentation/metal/mtldevice/makebuffer(bytesnocopy:length:options:deallocator:).json , "Creates a buffer that wraps an existing contiguous memory allocation."), as ggml-metal does (`ggml-metal-device.m:2143,2212`) | removes the load copy (the file's bytes are the card's bytes). Costs: page alignment (16384 B here), `maxBufferLength` 20100448256 B on this machine, so a 21 GB file is several buffers (a tensor never spans two; an expert stack per layer is one) | +| **Residency** | keep the weights in an `MTLResidencySet` (https://developer.apple.com/tutorials/data/documentation/metal/mtlresidencyset.json , "A collection of resource allocations that can move in and out of resident memory.") | avoids a first-touch fault inside a replay; the pool can still page under pressure, so the wired limit matters (C.6) | +| **Numeric contract** | integer-exact kernels (the q8_1 dots, the quantizers, the combine's fixed order) bit-equal to the CUDA engine; transcendental ones (`exp` in SwiGLU and softmax, `tanh`) within a band | the integer sums are exact on any hardware; the float closes are IEEE fp32 with the same operation order, so they are equal when the contraction is pinned (`fp contract` and `math_mode(safe)`). `exp` differs in ulps between libdevice and Metal: a band, held by the same gates as the CUDA ones (`AGENTS.md`: accuracy measures are diagnostics, the pins are the bug catchers) | + +**Unified memory beyond the plan (section C.6).** The host tier on a unified machine shares DRAM bandwidth with the GPU: on +a Mac an expert computed by the CPU does not add a second memory, it competes for the one. So the CPU tier has no capacity +use on a Mac whose memory holds the model (an M5 Ultra's 512 GB, or this M1 Pro's 26.8 GB working set for the 21.2 GB +Qwen3.6 file), and only the NVMe tier past that. This is why M1 leaves the critical path. + +### C.5 Apple Silicon facts, the terms, and the resource timeline + +The chip facts the design needs (bandwidth by class, the largest memory, the M5 Ultra's GPU) are in section 2, the matrix hardware in C.4, +the measured device limits in section 2 and C.6. This section holds the terms. + +Per decode token, whole model on the GPU, weights resident (`[derived]`; `n/a` is a resource the machine has none of): + +| Resource | Demand a token | Note | +|---|---|---| +| GPU memory traffic | 1.8 GB (Qwen3.6-35B-A3B `Q4_K_M`) plus K/V reads at the depth | 1.5 ms at the M5 Ultra datasheet 1.2 TB/s; the K/V term is `2 × layers × kv_heads × head × depth × bytes` and not computed here (hparams not read) | +| GPU launch chain | 12 L + 4 dispatches | g µs each: unknown on Metal | +| GPU ALU | 2.1 × 10^12 multiply-adds a second to keep pace with the stream at 1.2 TB/s | peak not stated: `[unverified]` | +| Host CPU | write the step record, wait on the replay | no per-launch work if the step is one ICB | +| PCIe | n/a | one memory | +| NVMe | 0 after load | | +| Host DRAM | the same DRAM as the GPU's | a CPU tier steals from the GPU's bytes | + +The wall is not a sum: the host's two items run under the GPU's. It is the GPU's bytes plus its chain. A lever on the host CPU +predicts 0; a lever on the chain predicts up to the gap budget above; a lever on the bytes (a smaller weight format, drafted +rows that share one pass over the weights) predicts in proportion. The step is latency, not bandwidth, once the chain term passes the +byte term, and then the busiest unit is not the bound (`AGENTS.md`, the q3gemma case). The two terms a derivation cannot fix are +**c_d**, the cost of one replayed dispatch in an ICB on M-series GPUs, and **f_bw**, the fraction of the datasheet bandwidth a +streaming-read kernel reaches. Both need a Mac runner (Part D) and the first Metal round: an empty-kernel ICB of 580 dispatches +replayed N times, and a read-only kernel over a 2 GiB buffer. The expected values cannot be written from anything measured here; +that is the reason the first round's step 0 is these two measurements, taken before any kernel is ported. + +### C.6 Unified memory + +What of the Spark's placement applies (`crates/placement/src/placement.rs:280-289` `UnifiedPool`; `:3956-3960` `pool_card_bytes` set when +`machine.unified()` is some; `Violation::PoolOver` :1580, :4182; `tierengine.md`, section (h), line 264): + +| Spark (GB10) | On a Mac | +|---|---| +| one pool, `PoolOver` holds Σ card + host ≤ available | the same rule. The census fills `UnifiedPool` from `MTLDevice.hasUnifiedMemory` (https://developer.apple.com/tutorials/data/documentation/metal/mtldevice/hasunifiedmemory.json : "A Boolean value that indicates whether the GPU shares all of its memory with the CPU.") | +| the pool's room is `MemAvailable` | the GPU's share is capped lower: `recommendedMaxWorkingSetSize` ("An approximation of how much memory, in bytes, this GPU device can allocate without affecting its runtime performance.", https://developer.apple.com/tutorials/data/documentation/metal/mtldevice/recommendedmaxworkingsetsize.json) was 26800603136 B of 34359738368 on this machine (0.78 [derived]). `iogpu.wired_limit_mb` read 0 here (the default); a third party's reading of the default is two thirds of RAM up to 36 GB and three quarters above (`mistral.rs` `mistralrs-core/src/utils/memory_usage.rs:398-406`, `[unverified]` for Apple's behaviour). The pool is `min(working set, MemAvailable)` | +| `cuMemHostAlloc(PORTABLE\|DEVICEMAP\|WRITECOMBINED)` pinned arena | `MTLResourceStorageModeShared`: https://developer.apple.com/tutorials/data/documentation/metal/mtlstoragemode/shared.json , "The CPU and GPU share access to the resource, allocated in system memory." and "This is the default storage mode for … instances on integrated GPUs and both … and … instances on Apple silicon GPUs." (the type names are links in the page data) No flags, no registration; write-combined does not exist | +| `cudaHostRegister` of an anonymous arena | not needed: wrap page-aligned memory with `newBufferWithBytesNoCopy` | +| the card leg reads the arena in place | the same: a Metal kernel reads the shared buffer. The CPU and GPU must not touch a buffer at once (Apple: "you’re responsible for synchronizing access to the texture from the other participant"): the `HostFlags` role is `MTLSharedEvent` waits and signals (https://developer.apple.com/tutorials/data/documentation/metal/mtlsharedevent.json : "A shared event can synchronize memory operations across multiple Metal devices and the CPU.") | +| `O_DIRECT` reads land in the arena | `F_NOCACHE` and `pread` into the buffer's pages | +| `ROW_PAGE` 4096 | 16384: the page is a parameter | + +## Part B. The CPU path (cut to what M0 and M2 need) + +**What a CPU path is for on a Mac.** Three things. The development loop (the host chain builds and the pure crates' tests +run on the Mac today: 1477 passed). The oracle: `qdot`'s scalar mirrors are the definition the card kernels and the +reference sets are held to, and they run on a Mac as they are. And a host tier for experts that do not fit the GPU's +working set; on a unified machine that tier competes for the GPU's DRAM (C.4), so it matters only past the memory, for +the NVMe-backed tier. It is not on the path to a fast decode, so a NEON engine is not proposed ahead of the Metal one. + +### B.1 `qdot` off x86, in short + +- Every AVX2 item is under `cfg(target_arch = "x86_64")` (`crates/qdot/src/lib.rs:62`, 56 `#[target_feature]` fns in `lib.rs`, 5 in `generic.rs`). + Off x86, `supports` is "the type has a kernel" (`lib.rs:224-228`), `has_features` is false (`:277-280`), `quantize_col` and `dot_row` run the scalar + mirrors (`:431-434`, `:595-602`, `:605-636`), and the tile, row-lane and pack kernels are `unreachable!("no qdot … off x86_64")` + (`:893-902`, `:1133-1136`, `:1269-1272`, `:1377-1380`), guarded by `has_tile`/`has_lanes`, which are false (`:747-775`, `:1159-1163`). So every type + takes the fused path through its scalar mirror: the AVX2 kernels' bits, not their speed. +- Two kinds of arm. The **hand kernels** are the 16 types `dot_row_scalar_ty` matches (`lib.rs:610-627`: Q4_K, Q5_K, Q6_K, Q5_0, Q5_1, Q8_0, + Q2_K, Q3_K, IQ2_XS, IQ2_XXS, IQ3_XXS, IQ3_S, MXFP4, IQ4_NL, IQ4_XS, Q2_0), each with an AVX2 body and an `emul`/scalar mirror, plus the tiles + (14 types, `tile_kind` :747), the Q3_K row-lane tile (:1048), the Q4_K/Q5_K/Q6_K lanes (:1161-1292) and the AVX-512 bodies for Q4_K (6 `target_feature` + fns, `q4k_avx512` :292). Their non-x86 arm is the mirror or a refusal, as above. The **runtime-detected helpers** take an AVX2 path when + `is_x86_feature_detected!` says so and fall through to a scalar loop otherwise, so off x86 they are the scalar loop with no `cfg` arm of + their own: `card_q8_1_fill` (:2596), `card_q3k_dot_row_cols` (:3068), `card_q4k_dot_row_cols` (:3268), `card_swiglu_clamp` (:3493), `card_expf` (:3521), + `q_nope2_cells` (:9245, with a `cfg(not(x86_64))` arm at :9273), `dot_f32` (:9389), `sum_sq_f64` (:9452), `swiglu` (:9499), `swiglu_clamp` (:9524). +- The generic kernels (`crates/qdot/src/generic.rs`: `dot_avx2`/`dot_scalar`, Q8_K and Q8_2_X4 forms) serve no type today: + `ENTRIES` is empty (`generic.rs:129`), the three decoders (`Q2k`, `Iq4Xs`, `Iq4Nl`, `:572`, `:606`, `:636`) are the gate's proof types, and + `crates/qdot/tests/generic.rs:25` is `#![cfg(target_arch = "x86_64")]`, so the gate does not compile on a Mac. +- **NEON form of the two generic kernels (prototype, measured).** The AVX2 steps map as: `maddubs(|w|, sign(q,w))` then + `madd(1, ·)` to `vdotq_s32` (signed × signed, the sign fold is not needed); `madd(scale, pair)` to `vmlaq_n_s32` + of the per-16-value dot lanes; `_mm256_fmadd_ps` to `vfmaq_f32`; `_mm256_cvtepi32_ps` to `vcvtq_f32_s32`; `hsum_float_8` to + `vaddq_f32(lo, hi)`, `vadd_f32(low, high)`, then lane 0 + lane 1; the Q8_2_X4 group fold to a `vpaddq_s32` chain. The integer + sums are exact, and `maddubs` saturates only above 32767, which `|w| ≤ 127` and `|q| ≤ 127` (2 × 127 × 127 = 32258) never reach, so + `vdotq_s32` equals `maddubs` + `madd` on the types' legal range. The float closes are the same fused multiply-adds in the same lane order, so the result + is bit-identical. A scratch prototype (`target/mac-portmap/neonproto/main.rs`, built in the layer-3 copy against the real + `generic::Entry::dot_scalar`) gave **0 mismatches on 2400 rows** (Q2_K with mins, IQ4_XS, IQ4_NL; random, spiked and + alternating-sign activations; `target/mac-portmap/neonproto-run.log`). FAIL-first: an unfused multiply then add in the closing step + gave 341 mismatches, a sequential lane sum for `hsum` gave 1292 (`neonproto-mut1-run.log`, `neonproto-mut2-run.log`). The + intrinsics are Arm's (https://arm-software.github.io/acle/neon_intrinsics/advsimd.html : "int32x4_t vdotq_s32 ( int32x4_t r, int8x16_t a, int8x16_t b) r -> Vd.4S a -> Vn.16B b -> Vm.16B SDOT Vd.4S,Vn.16B,Vm.16B") and compile with no feature gate on the + pinned nightly (`#[target_feature(enable = "dotprod")]`). Apple lists the capability as https://developer.apple.com/tutorials/data/documentation/kernel/1387446-sysctlbyname/determining_instruction_set_characteristics.json : + "hw.optional.arm.FEAT_DotProd": "Advanced SIMD Int8 dot product instructions."; `FEAT_I8MM` ("Advanced SIMD Int8 matrix multiplication + instructions.", the `vmmlaq_s32` family) is 0 on this M1 Pro (measured) and not tested. +- What is missing for a CPU engine at speed: the type decoders for the hot types (only 3 exist, `ENTRIES` empty), a NEON tile kernel (up to + `TILE_COLS` columns, so a prompt batch reads each weight block once), and the hand kernels' lanes for Q4_K/Q5_K/Q6_K. The Spark's decision + stands (`dgx-spark.md:1399-1401`): "NEON only when a host-tier gate's wall asks for it." No Metal step waits for it. + +### B.2 `serve::Engine` and what the Metal engine needs from it + +`trait Engine` is `crates/serve/src/engine.rs:207-538`: 6 required methods and 32 with a default. The only implementation over a +real model is `SeatEngine` (`crates/gpu-gates/src/bind.rs:1519`, impl :2983), which drives a `Seat` (:1268, 12 required methods); the mock +is `crates/serve/src/mock.rs`. `crates/model` has no implementation, and `crates/serve/src/bin/bloomery-serve.rs:1` serves the mock only. + +| Required method (`engine.rs`) | What a CPU or Metal engine needs | What exists today | Missing | +|---|---|---|---| +| `tokenizer` :209 | an `Arc` over the file's vocabulary | `tokenizer::Tokenizer::from_gguf` (`crates/tokenizer/src/lib.rs:112`, pure); the adapter `Vocab` is `crates/gpu-gates/src/bind.rs:81-173`, in the device-linked crate | lift `Vocab` and `VocabDecoder` (:174) into a device-free crate: about 140 lines `[derived: bind.rs:81-220]`, a move | +| `prefill` :219 | evaluate ids, no token | CPU: `deepseek2::forward::step` (`crates/model/src/arch/deepseek2/forward.rs:307`) over a `KvCache` (`kv.rs:76`); GPU: `Seat::prefill` (`bind.rs:1276`) | CPU: every architecture but deepseek2. Metal: the engine itself | +| `next` :234 | evaluate `last`, return the argmax, optionally the logits | CPU: `step` then `argmax` (`forward.rs:267`); GPU: `Seat::step` :1296, `logits_into` :1298 | as above | +| `reset` :383 | drop the cache | `Seat::reset` :1421; CPU: a new `KvCache` | none | +| `ctx_max` :409 | cache positions | `Seat::ctx_max` :1272 | none | +| `describe` :411 | device and position for a crash report | `SeatEngine`'s impl | a Metal device string | + +The defaults carry a one-slot, no-draft, no-cache-reuse server (`slots` 1, `advance` = one `next`), so a Metal engine can answer +`bloomery-serve` with the six methods and `Seat`'s required twelve, of which `snapshot`/`resume`/`rollback`/`note`/`splits`/`keep` +can start as named refusals. Which architectures run end to end on the CPU path today: **only `deepseek2`** (`crates/model/src/bin/bloomery-decode.rs:33` +`RUNS = ["deepseek2"]`; `forward.rs` has `forward`, `step`, `new_cache`). The other architectures in `crates/model/src/arch/` hold +descriptions (`spec`, `place`, `roles`, `names`) and the host tier's expert legs (`host.rs` of `deepseek41`, `qwen35moe`, `glm5next`, `mimo2`); +`qwen3moe` has 869 lines and no forward. There is no host GQA attention, rope for GQA or QK-norm in `crates/model` (the only rope is the +MLA one, `deepseek2/attn.rs:326`), so a CPU engine for Qwen3-MoE is new code, not a wiring. + +**The threads pool on macOS.** The pool pins workers across L3 groups read from sysfs (`crates/threads/src/lib.rs:12-24`, `topology.rs:78-140`). +On a Mac: no sysfs, so one group of one cpu and a default of one thread (`topology.rs:99-112` `[derived]`); no affinity (`KERN_NOT_SUPPORTED`, +measured); two core classes (`hw.nperflevels` 2: 8 performance and 2 efficiency cores here, measured) and chunks split evenly +(`lib.rs:735-742`), so a worker on an efficiency core stretches every dispatch. The Mac form: the thread count is +`hw.perflevel0.logicalcpu`, workers take `QOS_CLASS_USER_INTERACTIVE`, topology is one group of the performance cores (their shared L2, +12 MiB here, stands where the L3 stands) `[derived: a design, not measured]`. + +## Part D. Gates on a Mac (M3) + +**What can run at each phase.** + +| Phase | Runs on a Mac | Evidence class | +|---|---|---| +| now | `just mac-check`, `mac-lint`, `mac-fmt-check`, `mac-static`, `mac-test` (`tools/mac-check.sh test`: 14 pure crates, 1477 passed, 0 failed, 210 ignored, 52 test binaries, measured, `target/mac-portmap/mac-check-test.log`) | development-loop evidence, never landing evidence (`AGENTS.md`) | +| M0 | the host library tests of the 7 crates that fail today (`threads`, `qdot`, `engram`, `model`, `model-hf`, `gguf`, `engram-lab`), natively; `qdot`'s non-hardware tests (the `hw_` ones are `#[ignore]`d: 62 of 94 in `crates/qdot/tests/qdot.rs`) | the pure-crates rule's native test, extended to those crates | +| M1 | `qdot` bit tests of a NEON kernel against its scalar mirror (the prototype's clause, above); the CPU engine's gates | needs `tests/generic.rs` un-gated from x86 (:25) or a new file | +| M2 | one Metal family gate: the first model's layer taps and token ids against the box's CUDA reference (a reference set dumped once on the box and read by `crates/refset`, which already runs on the Mac); each kernel against its scalar or CPU oracle (`qdot`'s q8_1 and Q4_K mirrors for the integer-exact ones) | a gate that needs Apple Silicon and a Metal device | +| M3 | fixture-tier Metal gates in a batch; a timing runner | a Mac runner (below) | + +**What assumes Linux, and where.** + +- `tools/gate.sh:74,97` runs `timeout --kill-after=10`; `tools/gpu-gate.sh` takes `flock` locks on `/root/...`, reads `/proc//stat` + and `nvidia-smi` (its header :16-17 says the self-test needs them and the Mac has neither; the self-test already carries Python stand-ins, :381-400). + `tools/box.sh` is the box. None of these run a gate on a Mac; a Mac gate runner is new: `tools/mac-gate.sh`, the lock a `lockf`-style + stand-in, the bound `gtimeout` or the Python one. +- The timing runners (`tools/ref/*`, `lease_take`) are the box's: a machine-wide lock, GPU-idle wait and witness blocks. A Mac equivalent + needs its own witness (chip, memory, power mode, thermal state, other GPU users): which Apple tool gives a usable thermal and power + reading without privilege was not tested `[unverified]`. Until it exists, `AGENTS.md` "Never time by hand" forbids a Mac number, so + the two terms of C.5 wait for it. +- `tools/recipes.py`: `pure-crates` (:5876) and its `_NOT_ON_MAC` list (:5512-5525) — its `#[target_feature(` pattern must take + a feature name (`neon`, `dotprod`, `i8mm`) as fine on aarch64, and the name list needs `O_DIRECT`; `HOST_DEVICE_LINKED` (:5905) holds one + crate (`bloomery-app`) and concerns the Linux host nightly (`tools/nightly/run.sh`), not the Mac; the `gate-*` recipes all start `./tools/box.sh` + (`justfile:1609`, `:1655`), so a Mac recipe needs its own `mac-` prefix, or a `BLOOMERY_LOCAL` mode as `dgx-spark.md` section 4 has for local gates. +- Crate classes: a new crate (`bloomery-device`, the platform crate, a Metal crate) needs rows in `tools/gate-class.tsv`, `tools/families.tsv`, + `tools/owners.tsv` and `just check-recipes` passing `[unverified: not run]`. + +**The `AGENTS.md` "Never" bullet that would change** (proposed text; `AGENTS.md` is not edited here). Today: + +> **Never run a gate where it cannot run.** Gates need x86_64 Linux and an NVIDIA card: `qdot` uses `std::arch::x86_64`, `threads` Linux affinity calls. On an arm64 Mac run only the static tier: … + +Proposed, from M0: + +> **Never run a gate where it cannot run.** The CUDA gates need Linux and an NVIDIA card. The host gates (`just mac-test`, the crates `tools/recipes.py pure-crates` selects) run on Linux and on aarch64 macOS; every Linux-only call sits behind `bloomery-os`, and a Mac result is development-loop evidence. From M2, the Metal gates (`gate-metal-*`) need Apple Silicon and run through `tools/mac-gate.sh`; they never run on Linux. On an arm64 Mac without the Metal tier run only the static tier: … + +And the first bullet ("Never build a device crate with plain cargo") gains one sentence at M2: "A Metal crate links no cuda-core and builds with plain cargo; `tools/recipes.py` treats it as a pure crate only when it builds without a Metal device." + +## Part E. The phased plan + +Ordered for the shortest path to a fast Metal decode on one model. M1 leaves the critical path: the Metal engine runs the model whole on +the GPU out of unified memory, where a CPU tier adds no capacity (C.4), so nothing in M2 waits on NEON or a CPU `serve::Engine`; the +`Vocab` lift (B.2) is the only part of M1 the serve step needs. M0 stays: the Metal crate reuses `bloomery-model`'s placement and +hparams, `gguf`, `refset` and `threads`, which do not build on a Mac today. + +| Step | Work items | Files touched; size `[derived]` | What proves it (change class, `AGENTS.md` "Derive first") | Depends on | +|---|---|---|---|---| +| **M0a** compile fence | the 45 errors at 38 lines: a leaf crate `bloomery-os` (name open, Part F) with `affinity`, `rusage`, `file_hint` (drop, populate, sync), `direct_open`, `rename_noreplace`, `mincore`, `free_bytes`; callers switched | 10 files, ~45 call lines changed; the crate ~15 functions × ~20 Linux-body lines + ~25 macOS lines ≈ 650 lines `[derived: counts × typical body]` | **move**: `just canon-check BASE` (both sides' `--no-jit` scans of the scan set equal; no kernel moved); the owning gates of the touched crates through `--round-ledger` (`gate-threads`, `gate-model-hf`, the engram and model gates); on the Mac `tools/mac-check.sh test` extended to the 7 crates | none | +| **M0b** runtime on macOS | `ROW_PAGE` a parameter (`placement/workstation.rs:141`, `model/.../workstation.rs:24-28`, `placement.rs:3055-3062`); `host_available` from `hw.memsize` and `vm_statistics64` or the Metal working set; the pool's default threads and QoS; `ru_maxrss` units | 6 files, ~120 lines | move on x86 (`canon-check`); a unit test at page 16384 and a Mac-only clause by name | M0a | +| **M2.0** two measurements | `c_d` (an empty-kernel ICB of 12 L + 4 dispatches replayed) and `f_bw` (a read-only kernel over 2 GiB), Swift or Rust probe, taken by a Mac runner | 1 probe crate, ~200 lines; the runner part of M3 | each with its prediction card named in the card's format, the expected value written from the datasheet bytes and the gap budget (C.4) | M3's runner, minimal | +| **M2.1** device layer | `bloomery-device`: `Buffer`, `Pinned`, `Stream`, `Event`, `Graph`/ICB, `HostFlags`, fault word, census; CUDA behind it, then Metal; family imports renamed | new crate ~1.5k lines (`crates/gpu/src/graph.rs` 1595 lines is the model); ~280 import lines in 9 crates | **move** on x86: `just canon-check BASE` plus the structural lines (graph node count, eager = replay, e2e set identical) | M0 | +| **M2.2** the decode set | the 15 kernels of C.2 in MSL, one at a time, each with a gate against its oracle: integer-exact kernels bit-equal, transcendental ones in a band | ~15 `.metal` files, 1.4k entry lines plus cores `[derived]`; one gate bin | **add**: the family gate green with FAIL-first per clause; the CUDA side unchanged (`just ptx-compare` rc 0) | M2.1 | +| **M2.3** the chain | `ChainOps` over `dispatch.rs`; the Metal chain with the ICB; a `generate`-style CLI; layer taps against the box's CUDA reference | `dispatch.rs` split (1513 lines) into walk and ops; ~600 lines new | **move** for the CUDA side (`canon-check`, the structural lines), **add** for Metal (the e2e token set equal to the box's) | M2.2 | +| **M2.4** prefill | MPP `matmul2d` GEMM and grouped expert GEMM, prefill flash | the prompt-side set of C.2's last paragraph | the family gate; the timing on an M5-class machine | M2.3, an M5 machine | +| **M3** Mac tier | `tools/mac-gate.sh`, a Mac timing runner with witnesses, `AGENTS.md` text (Part D), a Mac table in the README with the chip named | 3 scripts, 1 contract edit | the runner's self-test (`tools/gpu-gate.sh`'s model); a number is admissible only from the runner | M2.0 for the minimal part | +| **M1** (off the critical path) | `Vocab` lift (~140 lines `[derived]`, with the serve step); NEON generic kernels + decoders for hot types; a CPU engine for one architecture | ~400 lines for NEON and the tests; the engine is new code (B.2) | **add** and **integer-path reorder** classes: bit-equal to the scalar mirror (prototype above) | M0 | + +**The first implementation round (M0a), ready to use.** + +- *Boundary.* One owner for the new crate, `Cargo.toml` `members` (:3), the callers below; no other agent edits them. +- *Files.* New: `crates/os/Cargo.toml`, `crates/os/src/{lib,affinity,file,mem}.rs`. Changed: `crates/threads/src/lib.rs` (:750, :809, :839), + `crates/threads/src/helper.rs` (:216, :232), `crates/threads/tests/helper.rs`, `crates/gguf/src/lib.rs` (:1318), `crates/engram/src/{lib,direct}.rs` + (:542, :701, :749; :35), `crates/engram/tests/engram.rs` (:871), `crates/model/src/{fileio,r8file}.rs` (:19, :33; :510, :525), + `crates/model/src/placement/host_lock.rs` (:818, :901, :1207), `Cargo.toml`, `Cargo.lock`. +- *Rules.* Each Linux body moves unchanged into `crates/os` under `cfg(target_os = "linux")`; the macOS bodies are the forms of A.2; a function with no + macOS form returns a named `Unsupported` error, never a defined value (`AGENTS.md`: no silent failure). No behaviour change on Linux. +- *Proof.* Before: record the base scans, `just narrow BASE` then `just canon-check BASE` (rc 0 = every entry the same program), on the box. + After: `tools/mac-check.sh test` with the 7 crates added to the native set (expected: `threads`, `qdot`, `engram-lab`, `model`, `model-hf` build and their non-ignored tests run; the clauses skipped by name are + listed), `just mac-static BASE`, then the owning gates through `tools/gate-batch.sh --round-ledger`. +- *Completion commands and expected output.* + `cargo check -p bloomery-threads -p bloomery-qdot -p bloomery-engram -p bloomery-engram-lab -p bloomery-model -p bloomery-model-hf --all-targets` on the Mac: rc 0; + `cargo test -p bloomery-gguf --no-run`: rc 0; `just canon-check `: rc 0; `python3 tools/recipes.py pure-crates`: the 7 crates move from `rejected` to `pure` only after + the `_NOT_ON_MAC` edit; `python3 tools/recipes.py --self-test`: green. +- *Report.* The per-crate rc table again, the list of clauses skipped on macOS by name, the `canon-check` line. + +## Part F. Open questions for the maintainer + +1. **Numeric contract of the first Metal kernels.** q8_1 activations as on CUDA (integer-exact, bit-equal where the order is copied, two extra dispatches a layer for the quantizers) or f32/f16 activations as ggml-metal (fewer dispatches, no activation rounding, but a different family of bits and a band against ik that is not measured)? The kernel list depends on it. +2. **The first Metal family.** `qwen3moe` (15 kernels, no fixture row: `registry.rs:103`), or `qwen35moe` (fixture exists, `registry.rs:124`, plus the delta rule and gated kernels), or both with `qwen3moe` first and the fixture family as the gate input? +3. **The seam's shape.** The `ChainOps` cut over `dispatch.rs` (C.1), or a Metal copy of the chain until a second family arrives? +4. **Where the Linux-only calls live.** A new leaf crate `bloomery-os`, or a module in `bloomery-gguf` (the lowest crate they all reach)? (`threads` depends on `levers` only; `gguf` is below `engram` and `model` but not below `threads`.) +5. **An M5-class machine.** Tuning the tensor-op kernels and the two C.5 terms need one (the M1 Pro cannot show the Neural Accelerators' rate; ggml's own comment says the tensor API does not pay before M5). Who owns it, and how is it reached, as the box is? +6. **A Mac runner and witness.** `tools/mac-gate.sh` and a timing runner with a witness block: what reading of power and thermal state is accepted as the witness (`AGENTS.md` "Never time by hand" has no Mac form)? +7. **Minimum macOS.** MLX says the Neural Accelerator path needs "macOS 26.2 or later" (the research page above); Metal 4 tensor ops and `MTLResidencySet` set a floor of their own. Pin one? +8. **Metal bindings.** `objc2-metal` (https://github.com/madsmtm/objc2 , pushed 2026-10-17; its `Cargo.toml` lists `MTL4ComputeCommandEncoder` and `MTLTensor` features, the published crate's coverage `[unverified]`) or a small Objective-C++ shim as `ggml-metal` has? (https://github.com/gfx-rs/metal-rs says "Use of this crate is deprecated".) +9. **Is M1 dropped from the critical path** as proposed, with the `Vocab` lift the only part the serve step needs, and NEON only if a host-tier gate's wall asks (`dgx-spark.md:1399-1401`)? +10. **Public numbers.** One table a chip (`AGENTS.md`: "a table never mixes cards"), and the reference engines in the same session. Which Mac engines are the line to beat is a separate round's survey; the baseline rule for the README (llama.cpp, mistral.rs) stands. +11. **Adapting ggml-metal source.** MIT; a kernel taken from it carries the notice in `THIRD_PARTY_NOTICES.md`. Acceptable for the K-quant mat-vec structure? +12. **Release.** An aarch64-apple-darwin archive needs signing and notarization; who holds the keys (the secrets go to the vault, not a file)? + +## Appendix. What was run in this round + +All under the round's `target/mac-portmap/` (gitignored, not published): + +| What | File | +|---|---| +| native census, layer 1 | `check/*.log`, `check/*.rc` (21 crates) | +| layers 2 and 3 (scratch copies, stubbed libc) | `check-layer2/`, `check-layer3/` | +| the error sites, deduplicated | `sites.tsv` | +| `/proc` and `/sys` lines | `proc-sys-sites.txt` | +| kernel index (`#[kernel]`, 376 at line start; 385 grep lines count 9 comment mentions) | `kernels.tsv` | +| native tests of the pure crates | `mac-check-test.log` | +| NEON prototype, its run, two mutants | `neonproto/main.rs`, `neonproto-run.log`, `neonproto-mut1-run.log`, `neonproto-mut2-run.log` | +| MSL compile probe on the M1 Pro | `msl-probe.log` | +| MPP `matmul2d` probes | `mpp-int8-probe.log`, `mpp-types-probe.log`, `MPPTensorOpsMatMul2d.h.txt` | +| machine and device probes (sysctl, `MTLDevice`, thread affinity, QoS) | `machine-probe.log`, `probes/{dev.swift,aff.c}` | +| MSL and MPP probe sources | `probes/{msl,mpp,mpp2}.swift` | +| pages fetched for the sourced lines | `web/` | +| `man` lines used | `man-facts.txt` | diff --git a/justfile b/justfile index e03ab2df4..a2ac429ff 100644 --- a/justfile +++ b/justfile @@ -1656,6 +1656,11 @@ gate-arch-registry: gate-threads: ./tools/box.sh 'bash tools/gate.sh --release -p bloomery-threads --lib --test pool --test helper -- --include-ignored --nocapture' +# The OS-call owner (crates/os, host only): each call's Linux body runs here; its macOS bodies and their named +# refusals run on a Mac through `just mac-test`. +gate-os: + ./tools/box.sh 'bash tools/gate.sh -p bloomery-os --lib -- --nocapture' + # 레버 레지스트리 게이트(crates/levers, 호스트만): 모든 행이 형식에 맞는가, 종류가 받지 않는 값과 은퇴한 이름을 # 이름을 대고 거부하는가, 제자리 행과 tools/levers-direct.txt가 서로 맞는가. 레지스트리를 마크다운 표로 찍는다. [group('canary')] diff --git a/labs/hip-native/Cargo.lock b/labs/hip-native/Cargo.lock index c1e75e001..86f145205 100644 --- a/labs/hip-native/Cargo.lock +++ b/labs/hip-native/Cargo.lock @@ -45,6 +45,7 @@ dependencies = [ "bloomery-gguf", "bloomery-levers", "bloomery-models", + "bloomery-os", "bloomery-placement", "bloomery-qdot", "bloomery-runtime", @@ -62,6 +63,13 @@ dependencies = [ "thiserror", ] +[[package]] +name = "bloomery-os" +version = "0.1.0" +dependencies = [ + "libc", +] + [[package]] name = "bloomery-placement" version = "0.1.0" @@ -92,6 +100,7 @@ name = "bloomery-threads" version = "0.1.0" dependencies = [ "bloomery-levers", + "bloomery-os", "libc", ] diff --git a/tools/gate-class.tsv b/tools/gate-class.tsv index d56fa0323..9fdd954c3 100644 --- a/tools/gate-class.tsv +++ b/tools/gate-class.tsv @@ -190,6 +190,7 @@ gate-nemotron-h-moe-host family nemotron_h_moe host model_host_ops gate-nemotron-h-moe-meta family nemotron_h_moe meta arch_registry gate-nemotron-meta family nemotron_h meta arch_registry gate-ops common host model_host_ops +gate-os common host os_calls gate-placement common place placement_sizing gate-prepare tooling corpus - gate-prepare-corpus tooling corpus - diff --git a/tools/owners.tsv b/tools/owners.tsv index 3a9a6a629..4dceee986 100644 --- a/tools/owners.tsv +++ b/tools/owners.tsv @@ -45,6 +45,7 @@ tokenizer_jinja host crates/tokenizer/src/lib.rs crates/jinja/src/lib.rs levers lever crates/levers/src/registry.rs crates/levers/src/lib.rs refset_readers refset crates/refset/src/lib.rs crates/refset/src/family.rs crates/refset/src/fixture.rs crates/refset/src/columns.rs threads_pool host crates/threads/src/lib.rs +os_calls host crates/os/src/lib.rs engram_verbs host crates/engram/src/lib.rs crates/verbs/src/lib.rs arch_registry refset crates/model/src/arch/registry.rs crates/model/src/arch/mod.rs crates/model/src/arch/classify.rs crates/model/src/arch/names.rs crates/model/src/arch/stems.rs crates/model/src/arch/coverage.rs model_host_ops host crates/model/src/ops.rs crates/model/src/ffn.rs crates/model/src/moe.rs crates/model/src/head.rs crates/model/src/embed.rs crates/model/src/quant.rs diff --git a/tools/recipes.py b/tools/recipes.py index eaa8b3d7d..dcfb9509d 100755 --- a/tools/recipes.py +++ b/tools/recipes.py @@ -112,12 +112,14 @@ its dependencies as a dependent sees them. What counts is what the crates use: `std::arch::x86_64` / `core::arch::x86_64`, `#[target_feature(…)]`, `is_x86_feature_detected!`, `std::os::linux`, the Linux-only libc items the tree calls (cpu_set_t and CPU_*, sched_*affinity, - posix_fadvise, sync_file_range, RUSAGE_THREAD, renameat2, prctl, MADV_POPULATE_*, the huge-page - flags), memmap2's Linux-only advice (PopulateRead/Write, HugePage, NoHugePage), and a "/proc/" or + posix_fadvise, sync_file_range, RUSAGE_THREAD, renameat2 and RENAME_NOREPLACE, prctl, + MADV_POPULATE_*, O_DIRECT, the huge-page flags), memmap2's Linux-only advice + (PopulateRead/Write, HugePage, NoHugePage), and a "/proc/" or "/sys/" path literal (the file does not exist on the Mac, so the code that reads it fails there). Comments are not code. An item or statement under a `#[cfg(…)]` that is false on aarch64-apple-darwin (target_arch, target_os, target_family, target_vendor, target_env, - target_pointer_width, target_endian, unix, windows) is left out, up to its end — the `;` or `,` at + target_pointer_width, target_endian, unix, windows) is left out, up to its end — the item's own + end first (a signature's comma, in `Result`, does not end it), then the `;` or `,` at its own depth, its closing brace, or the close of the block around it, whichever comes first. A file whose `#![cfg(…)]` is false is left out whole. `cfg(test)` is unknown in the crate's own files (both sides are built) and false in a dependency's (its test modules are not built). Any other @@ -5882,11 +5884,14 @@ def oxide(features: str, bins_: str, extra: str = "", sub: str = "build") -> str ( re.compile( r"\blibc::(?:cpu_set_t|CPU_[A-Z_]+|sched_[gs]etaffinity|sched_getcpu|posix_fadvise|POSIX_FADV_[A-Z_]+|sync_file_range" - r"|SYNC_FILE_RANGE_[A-Z_]+|RUSAGE_THREAD|renameat2|RENAME_[A-Z_]+|prctl|PR_[A-Z_]+|MADV_POPULATE_[A-Z_]+" + r"|SYNC_FILE_RANGE_[A-Z_]+|RUSAGE_THREAD|renameat2|RENAME_(?:NOREPLACE|EXCHANGE|WHITEOUT)|prctl|PR_[A-Z_]+" + r"|MADV_POPULATE_[A-Z_]+|O_DIRECT" r"|MADV_(?:NO)?HUGEPAGE|MAP_POPULATE|MAP_HUGETLB)\b" ), "a Linux-only libc item", ), + # `RENAME_EXCL` and `RENAME_SWAP` are not in the list above: macOS has them + # (`renamex_np`), so only Linux's own rename flags are a Linux-only item. (re.compile(r"\bAdvice::(?:PopulateRead|PopulateWrite|HugePage|NoHugePage)\b"), "memmap2's Linux-only advice"), (re.compile(r"\"/(?:proc|sys)/"), "a Linux /proc or /sys path"), ] @@ -6066,8 +6071,10 @@ def _item_end(shape: str, start: int) -> int: @functools.cache def mac_view(text: str, test: bool | None) -> str: """`text` as aarch64-apple-darwin compiles it: comments blanked, and every item or statement under a - `#[cfg(…)]` false there (a file under a false `#![cfg(…)]`) blanked; newlines kept, so line numbers hold.""" - return _cfg_blanked(text, test, _MAC_CFG_KEYS, _MAC_CFG_NAMES)[0] + `#[cfg(…)]` false there (a file under a false `#![cfg(…)]`) blanked; newlines kept, so line numbers hold. + The item's own end is read (`items`), so a fenced item whose signature holds a comma — a `Result` + return type — is blanked whole and not only up to that comma.""" + return _cfg_blanked(text, test, _MAC_CFG_KEYS, _MAC_CFG_NAMES, items=True)[0] @functools.cache @@ -7170,7 +7177,23 @@ def _item(self, rel: str, span: tuple[int, int], head: str, body, attrs: list[tu inside = text[len("cfg_attr") :].strip() parts = _split_top(inside[1:-1]) if inside.startswith("(") and inside.endswith(")") else [] named = [re.sub(r"\s+", "", pm.group(1)) if (pm := _ATTR_PATH.match(p)) else p for p in parts[1:]] - if not parts or any(n in ("cfg", "ignore", "path", "test") or n.endswith("::test") for n in named): + if parts and named == ["ignore"] and re.match(r'^ignore\s*=\s*".*"$', parts[1]): + # A conditional ignore the box can decide: the test is + # skipped on the platform the predicate names and runs on + # the one it names not. The box is the platform this scan + # reads, so the ignore applies there when the predicate + # holds there; a predicate the box cannot decide (a + # feature) is refused with the rest. + pred = self._cfg_of(rel, aline, f"cfg({parts[0]})", errs) + if pred is None: + return + v = eval_cfg(pred[3], True, _BOX_CFG_KEYS, _BOX_CFG_NAMES) + if v is True: + ignored = True + elif v is None: + errs.append(f"{rel}:{aline}: #[{text}] — a conditional ignore whose predicate this checker cannot decide for the box") + return + elif not parts or any(n in ("cfg", "ignore", "path", "test") or n.endswith("::test") for n in named): errs.append(f"{rel}:{aline}: #[{text}] — a conditional cfg, ignore, path or test attribute this checker does not evaluate") return elif name == "test": @@ -12156,6 +12179,12 @@ def _st_pure_crates(fails: list[str], side: Side) -> None: "procp": 'pub fn f() -> String { std::fs::read_to_string("/proc/self/io").unwrap() }\n', "lastfield": 'pub struct S {\n #[cfg(target_os = "linux")]\n a: u32\n}\npub fn b() -> i32 { libc::RUSAGE_THREAD }\n', "tmod": "pub fn f() {}\n#[cfg(test)]\nmod tests {\n fn t() { let _ = libc::RUSAGE_THREAD; }\n}\n", + # The one owner of the OS calls: its Linux-only item sits behind the + # platform's own cfg, so the crate and every crate that reaches it + # stay pure. The same item outside that fence is a Linux-only item + # the Mac cannot build, and the rule says so. + "osfence": '#[cfg(target_os = "linux")]\npub fn f() -> i32 { unsafe { libc::O_DIRECT } }\n', + "osoutside": "pub fn f() -> i32 { unsafe { libc::O_DIRECT } }\n", "user": "pub fn f() {}\n", } deps = { @@ -12215,9 +12244,10 @@ def _st_pure_crates(fails: list[str], side: Side) -> None: ("lastfield", "lastfield/src/lib.rs:5: a Linux-only libc item"), ("procp", "procp/src/lib.rs:1: a Linux /proc or /sys path"), ("tmod", "tmod/src/lib.rs:4: a Linux-only libc item"), + ("osoutside", "osoutside/src/lib.rs:1: a Linux-only libc item"), ]: expect(any(want in w for w in verdict.get(n, [])), f"pure: {n} is not rejected with '{want}': {verdict.get(n)}") - for n in ("guarded", "inner", "user"): + for n in ("guarded", "inner", "user", "osfence"): expect(verdict.get(n) == [], f"pure: {n} is rejected: {verdict.get(n)}") # host nightly: the same lock reading over the synthetic workspace — every crate with no device root, whatever its # x86_64 or Linux-only code, and an admitted device-linked one only by name; the tables' stale entries are problems @@ -12262,10 +12292,13 @@ def _st_host_crates(fails: list[str], side: Side) -> None: real = dict(pure_crates(tree, load_lock(tree.root))) pure_now = sorted(n for n, w in real.items() if not w) expect( - pure_now == ["bloomery-decision", "bloomery-embed", "bloomery-hf", "bloomery-jinja", "bloomery-levers", "bloomery-models", "bloomery-placement", "bloomery-refset", "bloomery-runtime", "bloomery-safetensors", "bloomery-sampler", "bloomery-serve", "bloomery-tokenizer", "bloomery-verbs", "bloomery-vision"], + pure_now == ["bloomery-decision", "bloomery-embed", "bloomery-gguf", "bloomery-hf", "bloomery-jinja", "bloomery-levers", "bloomery-models", "bloomery-os", "bloomery-placement", "bloomery-refset", "bloomery-runtime", "bloomery-safetensors", "bloomery-sampler", "bloomery-serve", "bloomery-tokenizer", "bloomery-verbs", "bloomery-vision"], f"pure: the real tree's pure crates are {pure_now} — a new member of the set is proven by `just mac-test` before this list takes it", ) - expect(any("crates/gguf/src/lib.rs:" in w and "RUSAGE_THREAD" in w for w in real.get("bloomery-gguf", [])), f"pure: gguf's reason {real.get('bloomery-gguf')}") + # the host chain's Linux-only items are the one owner's now, behind its own cfg, so a crate that holds + # none of its own is pure; a crate that holds one is still named + expect(not real.get("bloomery-gguf"), f"pure: gguf is rejected: {real.get('bloomery-gguf')}") + expect(any("a Linux-only libc item" in w for ws in real.values() for w in ws), "pure: the rule names no Linux-only libc item in the tree") expect(any("bloomery-gpu -> cuda-" in w for w in real.get("bloomery-gpu", [])), f"pure: bloomery-gpu's reason {real.get('bloomery-gpu')}") expect(any("crates/threads/src/" in w for w in real.get("bloomery-qdot", [])), f"pure: qdot's reason {real.get('bloomery-qdot')}") # host nightly on the real tree: the crates the set runs are the device-free ones and the admitted app; every gpu crate is left out @@ -12275,12 +12308,14 @@ def _st_host_crates(fails: list[str], side: Side) -> None: # PIN(2026-10-20): bloomery-model-hf joined with the tests it took out of bloomery-model (the fixture fetch's, the GLM checkpoint # map's), which the nightly ran as part of that crate. # PIN(2026-10-21): safetensors keeps the mapped reader's fixture tests previously run by bloomery-hf. + # PIN(2026-10-21): bloomery-os joined with the OS calls the host chain now makes through it, which the nightly runs as part of + # that crate. expect( host_now == [ "bloomery-app", "bloomery-decision", "bloomery-embed", "bloomery-engram", "bloomery-engram-lab", "bloomery-gguf", "bloomery-hf", "bloomery-jinja", - "bloomery-levers", "bloomery-model", "bloomery-model-hf", "bloomery-models", "bloomery-placement", "bloomery-qdot", "bloomery-refset", "bloomery-runtime", - "bloomery-safetensors", "bloomery-sampler", "bloomery-serve", "bloomery-threads", "bloomery-tokenizer", "bloomery-verbs", + "bloomery-levers", "bloomery-model", "bloomery-model-hf", "bloomery-models", "bloomery-os", "bloomery-placement", "bloomery-qdot", "bloomery-refset", + "bloomery-runtime", "bloomery-safetensors", "bloomery-sampler", "bloomery-serve", "bloomery-threads", "bloomery-tokenizer", "bloomery-verbs", "bloomery-vision", ], f"host nightly: the real tree's set is {host_now} — a new member of the set is proven by a nightly run before this list takes it", diff --git a/tools/unsafe-ratchet.txt b/tools/unsafe-ratchet.txt index d320dfffd..4826889b7 100644 --- a/tools/unsafe-ratchet.txt +++ b/tools/unsafe-ratchet.txt @@ -24,7 +24,7 @@ embed 0 0 0 # line. engram 10 0 0 engram-lab 0 0 0 -gguf 9 2 0 +gguf 7 2 0 # PIN(2026-10-01): gpu 1354 -> 1356, a new kernel entry, not old code relaxed: kda_delta_lanes (GLM's two-row # verify, one state lane per row) adds its launcher's and its kernel's decay-row blocks, each with its SAFETY line. # PIN(2026-10-01): gpu 1356 -> 1363 blocks, 143 -> 144 fns, a new kernel entry, not old code relaxed: @@ -506,11 +506,14 @@ levers 0 0 0 # compute call (ops.rs dispatch_ccd, over each run of a block) and the warm walk's prefetch (ops.rs warm_touch), each # with its SAFETY line, and the cut test's RowSet double (ccd_tests::Pairs, whose `unsafe fn compute` the trait # makes unsafe). -model 115 31 13 +model 109 31 13 # model-hf: the four blocks that left model (119 -> 115) with the fixture bin and its test, the page-fault reads (`getrusage`), # each with its SAFETY line; no new unsafe. model-hf 4 0 0 models 0 0 0 +# PIN(2026-10-21): os 24 blocks, 5 fns, not old code relaxed: the host chain's OS calls moved here from threads, +# gguf, engram and model (their pins come down), each beside its macOS body, each block with its SAFETY line. +os 24 5 0 placement 0 0 0 # PIN(2026-10-02): qdot 69 -> 77 blocks, 36 -> 41 fns, new kernels, not old code relaxed: the Q4_K/Q5_K row # lanes (pack_lanes_avx2, dot_lanes_avx2 and their helpers transpose8_epi32, transpose8_ps, lane_codes: five @@ -583,7 +586,7 @@ sampler 4 4 1 # the zeroed sigaction and its registration, the handler's stderr write with _exit on a second # interrupt and its one-byte write on the first, each with its SAFETY line. serve 6 0 0 -threads 9 0 1 +threads 3 0 1 tokenizer 0 0 0 # PIN(2026-10-14): verbs 0 -> 57 blocks, 22 fns, 2 impls, new crate, moved from gpu-gates/src/bin/shared/verbs.rs (45 # blocks there) with the tier link's RDMA end as its second user: the calls into libibverbs loaded at run time