From 1886a4a2bd2ac215ddc88f1396560bf58cdc9b10 Mon Sep 17 00:00:00 2001 From: midagedev Date: Tue, 3 Nov 2026 17:46:16 +0900 Subject: [PATCH 1/2] metal: the encode trace numbers a load's buffers as one thread running its jobs in order would, so two runs write one trace The weights load on several threads (`in_order`), and the trace named a buffer by the order it was made, so a weight buffer's `#n` moved from run to run: a compare of a commit against itself renumbered half the buffers, and `--compare-trace` could not return identical for any Metal change. `in_order` now runs each job under a trace scope (`trace::Load::job`). A buffer made inside job j is placed (j, k), and when the load ends the jobs below the earliest failure are numbered from the counter at the load's start, in job order. Buffers made after it continue from there. A buffer with no deterministic place is counted on a `refused REASON: N` line, which the compare refuses (rc 2): - one made on another thread outside a load; - one made by no job inside a load; - a nested load; - a stale job. The trace format is v2. A v1/v2 pair compares through the renumbering path with a note and never returns 0. The engine is unchanged; without the `encode-trace` feature the scope is a no-op. Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01CjHkkntc2L1xRN21xB39mj --- crates/metal/src/arch/qwen3moe/arena.rs | 143 ++++++- crates/metal/src/bin/metal_scan.rs | 173 ++++++-- crates/metal/src/trace.rs | 503 +++++++++++++++++++++++- tools/metal-compare.sh | 5 +- tools/spawn-allow.txt | 1 + 5 files changed, 764 insertions(+), 61 deletions(-) diff --git a/crates/metal/src/arch/qwen3moe/arena.rs b/crates/metal/src/arch/qwen3moe/arena.rs index e94cd80eb..23cdf9797 100644 --- a/crates/metal/src/arch/qwen3moe/arena.rs +++ b/crates/metal/src/arch/qwen3moe/arena.rs @@ -632,11 +632,34 @@ struct Loads { layers: Vec, } +/// The encode trace's view of a load: its buffers are numbered as the jobs +/// run in index order would make them. With no trace built, a load is its jobs. +#[cfg(feature = "encode-trace")] +use crate::trace::Load; + +#[cfg(not(feature = "encode-trace"))] +struct Load; + +#[cfg(not(feature = "encode-trace"))] +impl Load { + fn begin() -> Self { + Self + } + + fn job(&self, _index: usize, run: impl FnOnce() -> R) -> R { + run() + } + + fn end(self, _counted: usize) {} +} + /// Jobs `0..jobs` run by up to `workers` threads, each taking the next job in /// order; their results in job order, or the error of the earliest job that /// failed. A worker takes no job past the earliest failure, and every job /// before it was taken, so the error is the one a run in order on one thread -/// returns. +/// returns. The buffers the jobs make are named by job and place in the job, +/// so the encode trace numbers them the same on every thread count: the jobs +/// up to the earliest failure, in order. fn in_order( jobs: usize, workers: usize, @@ -644,6 +667,7 @@ fn in_order( ) -> Result, E> { let next = AtomicUsize::new(0); let failed = AtomicUsize::new(usize::MAX); + let load = Load::begin(); let work = || { let mut done = Vec::new(); loop { @@ -651,7 +675,7 @@ fn in_order( if j >= jobs || j > failed.load(Ordering::Relaxed) { return done; } - let r = job(j); + let r = load.job(j, || job(j)); if r.is_err() { failed.fetch_min(j, Ordering::Relaxed); } @@ -671,6 +695,13 @@ fn in_order( slots[j] = Some(r); } }); + // Jobs past the earliest failure may have run, which a run in order would + // not have: their buffers are not counted. + let counted = slots + .iter() + .position(|r| matches!(r, Some(Err(_)))) + .map_or(jobs, |failed| failed + 1); + load.end(counted); // Every job up to the earliest failure ran; the first `None` is past it. let mut out = Vec::with_capacity(jobs); for r in slots.into_iter().map_while(|r| r) { @@ -1463,6 +1494,9 @@ mod in_order_tests { /// came first in time returned (job 9 here, with two workers or more). #[test] fn jobs_return_in_order_and_the_earliest_failure_wins() { + // A load numbers buffers in the process's trace, one at a time. + #[cfg(feature = "encode-trace")] + let _session = crate::trace::testing::session(); for workers in [1, 2, 3, 8] { let got = in_order(20, workers, |j| Ok::<_, usize>(j * 3)).unwrap(); assert_eq!(got, (0..20).map(|j| j * 3).collect::>(), "{workers}"); @@ -1481,6 +1515,111 @@ mod in_order_tests { } assert_eq!(in_order(0, 4, Ok::<_, ()>).unwrap(), Vec::::new()); } + + /// The buffers the jobs make through the trace's hook, as the weights + /// load makes them: job `j` makes `MADE[j]` of them, buffer `k` at + /// `address(j, k)`. + #[cfg(feature = "encode-trace")] + const MADE: [u64; 5] = [2, 1, 0, 3, 2]; + + #[cfg(feature = "encode-trace")] + fn address(j: usize, k: u64) -> usize { + 0x10_0000 + j * 0x1000 + usize::try_from(k).expect("a small count") * 0x10 + } + + /// The names the trace gives after a load of [`MADE`]'s jobs on `workers` + /// threads, one buffer made before it and one after; job `j` fails when + /// `fails` says so. A script fixes the order the buffers are made in + /// across the threads. + #[cfg(feature = "encode-trace")] + fn loaded( + workers: usize, + script: Option<&crate::trace::testing::Script>, + fails: Option, + ) -> std::collections::BTreeMap { + use crate::trace::buffer_made; + let _session = crate::trace::testing::session(); + buffer_made(0x100, 8); + let result = in_order(MADE.len(), workers, |j| { + for k in 0..MADE[j] { + let make = || buffer_made(address(j, k), 16 * (j + 1)); + match script { + Some(script) => script.step(j, k, make), + None => make(), + } + } + if fails == Some(j) { Err(j) } else { Ok(j) } + }); + assert_eq!(result.err(), fails); + buffer_made(0x300, 32); + crate::trace::testing::names() + } + + /// The buffers a load's jobs make are numbered as one thread running the + /// jobs in index order would number them, on one worker and on four, + /// whichever order the four make them in. Mutant: one counter for every + /// buffer (the four workers' order shows). + #[cfg(feature = "encode-trace")] + #[test] + fn a_loads_buffers_are_numbered_alike_on_one_worker_and_on_four() { + use crate::trace::testing::Script; + + let one = loaded(1, None, None); + let mut expected = std::collections::BTreeMap::from([(0x100, (0, 8)), (0x300, (9, 32))]); + let mut n = 1; + for (j, &made) in MADE.iter().enumerate() { + for k in 0..made { + expected.insert(address(j, k), (n, 16 * (j + 1))); + n += 1; + } + } + assert_eq!(one, expected); + // The last jobs make their buffers first, the first jobs' last. + let script = Script::new(vec![ + (4, 0), + (3, 0), + (0, 0), + (4, 1), + (1, 0), + (3, 1), + (0, 1), + (3, 2), + ]); + assert_eq!(loaded(4, Some(&script), None), one); + } + + /// A load that fails numbers the jobs up to the earliest failure and no + /// other, as a run in order would have: a later job that ran beside it + /// left no number behind. Mutant: every job that ran is counted. + #[cfg(feature = "encode-trace")] + #[test] + fn a_failed_load_numbers_the_jobs_up_to_its_earliest_failure() { + use crate::trace::testing::Script; + + // Job 3 makes its buffers, then job 1 makes its and fails, then job 0 + // makes its own. Job 4 runs or not, depending on when a thread frees. + let script = Script::new(vec![ + (3, 0), + (3, 1), + (3, 2), + (1, 0), + (0, 0), + (0, 1), + (4, 0), + (4, 1), + ]); + let names = loaded(4, Some(&script), Some(1)); + assert_eq!( + names, + std::collections::BTreeMap::from([ + (0x100, (0, 8)), + (address(0, 0), (1, 16)), + (address(0, 1), (2, 16)), + (address(1, 0), (3, 32)), + (0x300, (4, 32)), + ]) + ); + } } #[cfg(test)] diff --git a/crates/metal/src/bin/metal_scan.rs b/crates/metal/src/bin/metal_scan.rs index 393bca5f2..8c3ad4ef3 100644 --- a/crates/metal/src/bin/metal_scan.rs +++ b/crates/metal/src/bin/metal_scan.rs @@ -91,18 +91,24 @@ //! `--compare-trace` compares two encode traces //! (`generate_metal --encode-trace`) and prints the first differing event: //! exit 0 identical, 6 identical up to buffer numbering, 3 they differ, 2 a -//! file that is no trace, two format versions or a line it cannot read. A -//! binding names its buffer by allocation order, so a change that allocates -//! one more buffer early shifts every later id and moves every later line. -//! When the lines differ, each side's buffer ids are renumbered by first use -//! (the first buffer a binding names is buffer 0, the next new one buffer 1) -//! and the traces compared again with every other field as text: equal -//! renumbered traces bind their slots in the same pattern under a bijection of -//! the two sides' buffers (exit 6, the pairs printed), anything else is a real -//! difference (exit 3, the first differing event printed with both sides' raw -//! lines). Exit 6 is a diagnostic, not a proof: a trace names a buffer by its -//! allocation order only, so two buffers of one length whose roles are -//! swapped at every use renumber the same. +//! file that is no trace of a format version this reads, or holds a line it +//! cannot read (a `refused` line, which a trace carries when a buffer had no +//! place in the numbering, is one). A binding names its buffer by allocation +//! order, so a change that allocates one more buffer early shifts every later +//! id and moves every later line. When the lines differ, each side's buffer +//! ids are renumbered by first use (the first buffer a binding names is +//! buffer 0, the next new one buffer 1) and the traces compared again with +//! every other field as text: equal renumbered traces bind their slots in the +//! same pattern under a bijection of the two sides' buffers (exit 6, the +//! pairs printed), anything else is a real difference (exit 3, the first +//! differing event printed with both sides' raw lines). Exit 6 is a +//! diagnostic, not a proof: a trace names a buffer by its allocation order +//! only, so two buffers of one length whose roles are swapped at every use +//! renumber the same. Two traces of the same format version with equal lines +//! are identical. A v1 trace numbered the buffers of the weight load in the +//! order its threads made them, so a v1 and a v2 trace are never identical: +//! they are compared renumbered, with a note, and equal lines are exit 6 at +//! best. use std::process::ExitCode; @@ -1188,7 +1194,11 @@ mod compare { /// The trace format's first line (`bloomery_metal::trace::VERSION`, which /// exists only in an encode-trace build). - const TRACE_VERSION: &str = "# bloomery encode trace v1"; + const TRACE_VERSION: &str = "# bloomery encode trace v2"; + + /// The first line of the older format, read for the compare of a trace + /// made before the move: its buffer numbers followed thread scheduling. + const TRACE_VERSION_V1: &str = "# bloomery encode trace v1"; /// The bytes an archive may change size by between two sessions of the /// device compiler (one 16-byte alignment unit, twice) and still be the @@ -1485,17 +1495,23 @@ mod compare { (lines, rc) } - /// The events of a trace: its lines but the `#` ones. - fn events(path: &str) -> Result, String> { + /// The events of a trace and its format's first line: its lines but the + /// `#` ones. + fn events(path: &str) -> Result<(&'static str, Vec), String> { let text = std::fs::read_to_string(path).map_err(|e| format!("{path}: {e}"))?; let mut lines = text.lines(); - if lines.next() != Some(TRACE_VERSION) { - return Err(format!("{path}: not a {TRACE_VERSION} trace")); - } - Ok(lines - .filter(|l| !l.starts_with('#')) - .map(str::to_string) - .collect()) + let first = lines.next(); + let version = [TRACE_VERSION, TRACE_VERSION_V1] + .into_iter() + .find(|v| first == Some(v)) + .ok_or_else(|| format!("{path}: not a {TRACE_VERSION} or {TRACE_VERSION_V1} trace"))?; + Ok(( + version, + lines + .filter(|l| !l.starts_with('#')) + .map(str::to_string) + .collect(), + )) } /// A buffer as a binding names it: `#n` for the nth buffer the engine @@ -1573,6 +1589,11 @@ mod compare { lines.push(line.clone()); continue; } + Some("refused") => { + return Err(format!( + "event {at}: the trace found a buffer no place in the numbering, so its ids are no names: {line}" + )); + } Some("dispatch") => {} kind => { return Err(format!( @@ -1605,6 +1626,11 @@ mod compare { Ok(Renumbered { lines, first_use }) } + /// A format version's name in a note: the last word of its first line. + fn version_tag(version: &str) -> &str { + version.rsplit(' ').next().unwrap_or(version) + } + /// The first position where two event lists differ, a list ending early /// differing where it ends. fn first_difference(a: &[String], b: &[String]) -> Option { @@ -1618,12 +1644,13 @@ mod compare { /// The most buffer pairs the renumbered verdict lists. const PAIRS_SHOWN: usize = 8; - /// Two traces: rc 0 the events are equal, 6 equal once each side's - /// buffer ids are renumbered by first use, 3 they differ (the first - /// differing event, in that comparison, printed with the events before it), - /// 2 a file that is no trace or holds a line this cannot read. + /// Two traces: rc 0 the events are equal and the traces number buffers + /// alike (one format version), 6 equal once each side's buffer ids are + /// renumbered by first use, 3 they differ (the first differing event, in + /// that comparison, printed with the events before it), 2 a file that is + /// no trace or holds a line this cannot read. pub(super) fn traces(base: &str, head: &str) -> u8 { - let (a, b) = match (events(base), events(head)) { + let ((va, a), (vb, b)) = match (events(base), events(head)) { (Ok(a), Ok(b)) => (a, b), (a, b) => { for e in [a.err(), b.err()].into_iter().flatten() { @@ -1643,6 +1670,16 @@ mod compare { dispatches(&b), buffers(&b) ); + // Equal lines of two formats are not one program: the weights' ids + // follow thread scheduling in one and job order in the other. + let alike = va == vb; + if !alike { + println!( + "note: base is a {} trace and head a {} trace: they number buffers differently, so the events are compared renumbered and are never identical", + version_tag(va), + version_tag(vb) + ); + } // Every line is read before any verdict, so two equal traces with a // line this cannot read are refused, not identical. let (ra, rb) = match (renumber(&a), renumber(&b)) { @@ -1656,10 +1693,11 @@ mod compare { return 2; } }; - let Some(raw_at) = first_difference(&a, &b) else { + let raw_at = first_difference(&a, &b); + if raw_at.is_none() && alike { println!("identical"); return 0; - }; + } let Some(at) = first_difference(&ra.lines, &rb.lines) else { let moved: Vec = ra .first_use @@ -1673,16 +1711,20 @@ mod compare { } else { "" }; + let pairs = if moved.is_empty() { + "the lines are equal as written".to_string() + } else { + format!("{}{more}", moved[..moved.len().min(PAIRS_SHOWN)].join(", ")) + }; println!( - "identical up to buffer numbering: {} of {} buffers renumbered ({}{more})", + "identical up to buffer numbering: {} of {} buffers renumbered ({pairs})", moved.len(), ra.first_use.len(), - moved[..moved.len().min(PAIRS_SHOWN)].join(", ") ); return 6; }; println!("first difference at event {at}:"); - if raw_at != at { + if let Some(raw_at) = raw_at.filter(|&raw_at| raw_at != at) { println!( " (the raw lines first differ at event {raw_at}; buffer ids are compared renumbered by first use)" ); @@ -2272,16 +2314,26 @@ mod compare { events } - /// The exit code of comparing two traces made of `a` and `b`. - fn trace_rc(a: &[String], b: &[String]) -> u8 { - let text = |events: &[String]| format!("{TRACE_VERSION}\n{}\n", events.join("\n")); - let (x, y) = (write("a", &text(a)), write("b", &text(b))); + /// The exit code of comparing two traces made of `a` and `b`, of the + /// format versions `versions` (base, head). + fn trace_rc_as(versions: [&str; 2], a: &[String], b: &[String]) -> u8 { + let text = + |version: &str, events: &[String]| format!("{version}\n{}\n", events.join("\n")); + let (x, y) = ( + write("a", &text(versions[0], a)), + write("b", &text(versions[1], b)), + ); let code = traces(&x, &y); let _ = std::fs::remove_file(&x); let _ = std::fs::remove_file(&y); code } + /// The exit code of comparing two traces of the current format. + fn trace_rc(a: &[String], b: &[String]) -> u8 { + trace_rc_as([TRACE_VERSION; 2], a, b) + } + #[test] fn equal_traces_are_identical() { let t = program(&[ @@ -2291,13 +2343,60 @@ mod compare { assert_eq!(trace_rc(&t, &t), 0); } + /// Equal lines under two format versions are not one program: v1 numbered + /// the weights' buffers in the order threads made them. They are compared + /// renumbered (6 when that is equal, 3 when not), never 0, whichever side + /// is the older. Mutant: the version ignored (0). + #[test] + fn equal_lines_of_two_format_versions_are_never_identical() { + let t = program(&["b0=#1+0/64 b1=#2+0/64", "b0=#2+0/64 b1=#1+0/64"]); + for versions in [ + [TRACE_VERSION_V1, TRACE_VERSION], + [TRACE_VERSION, TRACE_VERSION_V1], + ] { + assert_eq!(trace_rc_as(versions, &t, &t), 6, "{versions:?}"); + // Two traces of the old format are compared as they were. + let shifted = program(&["b0=#2+0/64 b1=#3+0/64", "b0=#3+0/64 b1=#2+0/64"]); + assert_eq!(trace_rc_as(versions, &t, &shifted), 6, "{versions:?}"); + let other = program(&["b0=#1+0/64 b1=#2+0/64", "b0=#1+0/64 b1=#2+0/64"]); + assert_eq!(trace_rc_as(versions, &t, &other), 3, "{versions:?}"); + } + assert_eq!(trace_rc_as([TRACE_VERSION_V1; 2], &t, &t), 0); + // A format this does not know is refused, not read as the nearest one. + let later = "# bloomery encode trace v3"; + assert_eq!(trace_rc_as([later, TRACE_VERSION], &t, &t), 2); + assert_eq!(trace_rc_as([TRACE_VERSION, later], &t, &t), 2); + } + + /// A trace that says a buffer had no place in the numbering is refused + /// by name: its ids cannot be read, equal or not. + #[test] + fn a_trace_that_refused_a_buffer_is_refused() { + let mut t = program(&["b0=#1+0/64"]); + t.insert(0, "refused a buffer made by no job: 1".to_string()); + assert_eq!(trace_rc(&t, &t), 2); + assert!( + renumber(&t).err().is_some_and( + |e| e.contains("no place in the numbering") && e.contains("refused") + ) + ); + } + + /// The format version the engine writes is the one this reads. + #[cfg(all(target_os = "macos", feature = "encode-trace"))] + #[test] + fn the_engines_trace_version_is_the_one_read() { + assert_eq!(bloomery_metal::trace::VERSION, TRACE_VERSION); + } + #[test] fn a_trace_with_its_ids_shifted_is_the_same_program() { let base = program(&["b0=#1+0/64 b1=#2+0/64", "b0=#2+0/64 b1=#3+0/64"]); // Every id one higher. let head = program(&["b0=#2+0/64 b1=#3+0/64", "b0=#3+0/64 b1=#4+0/64"]); assert_eq!(trace_rc(&base, &head), 6); - // The shift begins after the first dispatch, as one more buffer made early does. + // The shift begins after the first dispatch, as one more buffer made early does + // (a trace numbers buffers in the order one thread would make them). let head = program(&["b0=#1+0/64 b1=#3+0/64", "b0=#3+0/64 b1=#4+0/64"]); assert_eq!(trace_rc(&base, &head), 6); let (a, b) = (renumber(&base).unwrap(), renumber(&head).unwrap()); diff --git a/crates/metal/src/trace.rs b/crates/metal/src/trace.rs index adba1d33d..1a5d40dfc 100644 --- a/crates/metal/src/trace.rs +++ b/crates/metal/src/trace.rs @@ -8,8 +8,9 @@ //! ([`start`], [`finish`]). //! //! A buffer is named by its allocation order (`#n`, the nth buffer -//! [`crate::MetalBuffer`] made in the process), an offset and its length, never -//! by an address; a command buffer by the order it was made. A dispatch line +//! [`crate::MetalBuffer`] made in the process, counted as one thread would +//! make them), an offset and its length, never by an address; a command buffer +//! by the order it was made. A dispatch line //! states the whole binding table the dispatch runs under (Metal keeps a slot's //! binding until it is set again), so the order of the `set` calls between two //! dispatches is not part of the trace. The host calls no @@ -24,11 +25,25 @@ //! dispatch N.M.K LABEL max=… width=… tg_static=… grid=GXxGY tpg=T b0=… b1=… //! barrier N.M //! host settle V | host sync +//! refused REASON: COUNT //! ``` +//! +//! The order threads make buffers in is not the order a name may rest on, so +//! two things number them. A load whose jobs run on several threads ([`Load`]) +//! names a buffer by its job and its place in that job, and numbers them when +//! it ends: jobs in index order, as if one thread had run them in that order. +//! Every other buffer is numbered in the order the one thread that numbers +//! them (the first to make a buffer outside a load) makes it. A buffer that +//! has no such place, made by no job inside a load or on another thread +//! outside one, gets no number: it is bound as `#?n`, and the trace carries a +//! `refused` line that `metal_scan --compare-trace` does not read as an event +//! and so refuses the trace by name. use core::fmt::Write as _; +use std::cell::Cell; use std::collections::BTreeMap; use std::sync::{Mutex, MutexGuard, PoisonError}; +use std::thread::ThreadId; use objc2::runtime::ProtocolObject; use objc2_metal::{MTLBuffer, MTLCommandBuffer, MTLComputePipelineState, MTLDispatchType}; @@ -36,17 +51,53 @@ use objc2_metal::{MTLBuffer, MTLCommandBuffer, MTLComputePipelineState, MTLDispa use crate::encode::MAX_BINDINGS; use crate::library::Pipeline; -/// The format's version, the first line of a trace. -pub const VERSION: &str = "# bloomery encode trace v1"; +/// The format's version, the first line of a trace. A version moves when a +/// buffer's name means something else (v2: a load's buffers are numbered in +/// job order, not in the order threads made them). +pub const VERSION: &str = "# bloomery encode trace v2"; + +/// Why a buffer got no number. +const BY_NO_JOB: &str = "a buffer made inside a load by no job"; +const OFF_THREAD: &str = + "a buffer made outside a load on a thread other than the one that numbers them"; +const NESTED: &str = "a load began while another was open"; +const LOST_JOB: &str = "a buffer made by a job of a load that is not the open one"; + +/// How a buffer is named. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +enum Name { + /// The nth buffer, in order. + Number(u64), + /// Buffer `k` of job `job` of the open load, numbered when the load ends. + Placed { job: usize, k: u64 }, +} + +/// A load whose jobs are making buffers. +struct OpenLoad { + id: u64, + /// The number the load's first buffer takes: the counter when it began. + base: u64, + /// The buffers each job has made so far. + made: BTreeMap, +} struct State { /// A trace is being written. active: bool, lines: Vec, - /// Buffer address to (allocation order, length): every buffer made since - /// the process began, the last one at an address winning. - buffers: BTreeMap, + /// Buffer address to (name, length): every buffer made since the process + /// began, the last one at an address winning. + buffers: BTreeMap, + /// Buffers numbered so far; a load's are counted when it ends. made: u64, + /// The thread that numbers the buffers made outside a load. + sequence: Option, + /// The load now open, if one is. + load: Option, + /// Loads begun, the next one's id. + loads: u64, + /// Buffers that got no number, by reason. + refused: BTreeMap<&'static str, u64>, /// Buffers bound but never made by [`crate::MetalBuffer`], by first sight. strangers: BTreeMap, /// Command buffer address to its order. @@ -56,22 +107,86 @@ struct State { encoders: BTreeMap, } -static STATE: Mutex = Mutex::new(State { - active: false, - lines: Vec::new(), - buffers: BTreeMap::new(), - made: 0, - strangers: BTreeMap::new(), - command_buffers: BTreeMap::new(), - command_buffers_made: 0, - encoders: BTreeMap::new(), -}); +static STATE: Mutex = Mutex::new(State::new()); fn state() -> MutexGuard<'static, State> { STATE.lock().unwrap_or_else(PoisonError::into_inner) } impl State { + const fn new() -> Self { + Self { + active: false, + lines: Vec::new(), + buffers: BTreeMap::new(), + made: 0, + sequence: None, + load: None, + loads: 0, + refused: BTreeMap::new(), + strangers: BTreeMap::new(), + command_buffers: BTreeMap::new(), + command_buffers_made: 0, + encoders: BTreeMap::new(), + } + } + + fn refuse(&mut self, reason: &'static str) { + *self.refused.entry(reason).or_insert(0) += 1; + } + + /// The name of a buffer just made on `thread` inside job scope `scope` + /// (a load's id and a job's index), or why it has none. + fn name( + &mut self, + scope: Option<(u64, usize)>, + thread: ThreadId, + ) -> Result { + match (scope, self.load.as_mut()) { + (Some((id, job)), Some(load)) if load.id == id => { + let made = load.made.entry(job).or_insert(0); + let k = *made; + *made += 1; + Ok(Name::Placed { job, k }) + } + (Some(_), _) => Err(LOST_JOB), + (None, Some(_)) => Err(BY_NO_JOB), + (None, None) => { + if *self.sequence.get_or_insert(thread) != thread { + return Err(OFF_THREAD); + } + let n = self.made; + self.made += 1; + Ok(Name::Number(n)) + } + } + } + + /// Load `id` ends: the buffers of its jobs below `counted` are numbered + /// from the load's base, jobs in index order and each job's in the order + /// it made them; the buffers of the jobs from `counted` on are dropped + /// (their jobs would not have run in order). Nothing when `id` is not the + /// open load. + fn close(&mut self, id: u64, counted: usize) { + let Some(load) = self.load.take_if(|l| l.id == id) else { + return; + }; + let mut first = BTreeMap::new(); + let mut next = load.base; + for (&job, &n) in load.made.range(..counted) { + first.insert(job, next); + next += n; + } + self.buffers.retain(|_, (name, _)| match *name { + Name::Number(_) => true, + Name::Placed { job, k } => first.get(&job).is_some_and(|&at| { + *name = Name::Number(at + k); + true + }), + }); + self.made = next; + } + fn emit(&mut self, line: impl FnOnce() -> String) { if self.active { let line = line(); @@ -107,6 +222,9 @@ pub fn finish() -> String { state.active = false; let mut text = String::from(VERSION); text.push('\n'); + for (reason, n) in &state.refused { + let _ = writeln!(text, "refused {reason}: {n}"); + } for line in state.lines.drain(..) { text.push_str(&line); text.push('\n'); @@ -114,12 +232,94 @@ pub fn finish() -> String { text } +thread_local! { + /// The load and the job of it that the buffers this thread makes belong to. + static SCOPE: Cell> = const { Cell::new(None) }; +} + /// A buffer of `length` bytes at `address` was made. pub(crate) fn buffer_made(address: usize, length: usize) { + let scope = SCOPE.get(); + let thread = std::thread::current().id(); let mut state = state(); - let n = state.made; - state.made += 1; - state.buffers.insert(address, (n, length)); + match state.name(scope, thread) { + Ok(name) => { + state.buffers.insert(address, (name, length)); + } + Err(reason) => { + state.refuse(reason); + // An earlier buffer at this address is gone: bound, this one is a + // stranger. + state.buffers.remove(&address); + } + } +} + +/// A load whose jobs make buffers on several threads. The threads' order is +/// not the buffers' order: each job runs under [`Load::job`], and when the +/// load ends its buffers are numbered as one thread running the jobs in index +/// order would have made them. Loads do not overlap: one begun inside another +/// is refused by name and numbers nothing. +pub(crate) struct Load { + id: u64, +} + +impl Load { + /// A load begins; its buffers are numbered from the counter as it stands. + pub(crate) fn begin() -> Self { + let mut state = state(); + let id = state.loads; + state.loads += 1; + if state.load.is_some() { + state.refuse(NESTED); + } else { + let base = state.made; + state.load = Some(OpenLoad { + id, + base, + made: BTreeMap::new(), + }); + } + Self { id } + } + + /// Job `index` runs `run`, on this thread, start to end. + pub(crate) fn job(&self, index: usize, run: impl FnOnce() -> R) -> R { + let _scope = Scope::enter(self.id, index); + run() + } + + /// The load ends with its jobs below `counted` having run: the buffers of + /// those are numbered, the others dropped. + pub(crate) fn end(self, counted: usize) { + state().close(self.id, counted); + } +} + +/// A load dropped without [`Load::end`] (it unwound) counts every job it ran. +impl Drop for Load { + fn drop(&mut self) { + state().close(self.id, usize::MAX); + } +} + +/// The job scope of the running thread, restored when it ends. +struct Scope { + previous: Option<(u64, usize)>, +} + +impl Scope { + fn enter(load: u64, job: usize) -> Self { + Self { + previous: SCOPE.replace(Some((load, job))), + } + } +} + +impl Drop for Scope { + fn drop(&mut self) { + SCOPE.set(self.previous); + } } /// A command buffer was made. @@ -214,7 +414,9 @@ impl EncoderTrace { let address = core::ptr::from_ref(buffer).addr(); let mut state = state(); let name = match state.buffers.get(&address) { - Some(&(n, length)) => format!("#{n}+{offset}/{length}"), + Some(&(Name::Number(n), length)) => format!("#{n}+{offset}/{length}"), + // Bound while its load runs: no number yet, and no id a compare reads. + Some(&(Name::Placed { job, k }, length)) => format!("#!j{job}.{k}+{offset}/{length}"), None => { let next = state.strangers.len() as u64; let n = *state.strangers.entry(address).or_insert(next); @@ -257,3 +459,262 @@ impl Drop for EncoderTrace { state().emit(|| format!("enc {} end", self.name)); } } + +/// What the numbering's tests share: the process-wide state is reset and held +/// by one test at a time. +#[cfg(test)] +pub(crate) mod testing { + use std::collections::BTreeMap; + use std::sync::{Condvar, Mutex, MutexGuard, PoisonError}; + use std::time::Duration; + + use super::{Name, State, state}; + + /// The state, reset, held until the guard drops. + pub(crate) fn session() -> MutexGuard<'static, ()> { + static LOCK: Mutex<()> = Mutex::new(()); + let guard = LOCK.lock().unwrap_or_else(PoisonError::into_inner); + *state() = State::new(); + guard + } + + /// Address to (number, length) of every buffer named; a buffer a load has + /// not numbered yet is a failure. + pub(crate) fn names() -> BTreeMap { + state() + .buffers + .iter() + .map(|(&address, &(name, length))| match name { + Name::Number(n) => (address, (n, length)), + Name::Placed { job, k } => { + panic!("buffer {k} of job {job} at {address:#x} was never numbered") + } + }) + .collect() + } + + /// The buffers that got no number, by reason. + pub(crate) fn refused() -> BTreeMap<&'static str, u64> { + state().refused.clone() + } + + /// An order the buffers of a load are made in, across threads: each + /// `(job, k)` event waits for its turn, so an interleaving is the script's + /// and no timing's. + pub(crate) struct Script { + order: Vec<(usize, u64)>, + at: Mutex, + turn: Condvar, + } + + impl Script { + pub(crate) fn new(order: Vec<(usize, u64)>) -> Self { + Self { + order, + at: Mutex::new(0), + turn: Condvar::new(), + } + } + + /// Runs `event` as the script's next step, which must be `(job, k)`. + pub(crate) fn step(&self, job: usize, k: u64, event: impl FnOnce()) { + let at = self.at.lock().unwrap_or_else(PoisonError::into_inner); + let (mut at, waited) = self + .turn + .wait_timeout_while(at, Duration::from_secs(60), |at| { + self.order.get(*at) != Some(&(job, k)) + }) + .unwrap_or_else(PoisonError::into_inner); + assert!( + !waited.timed_out(), + "the script never reached ({job}, {k}): step {at} is {:?}", + self.order.get(*at) + ); + event(); + *at += 1; + self.turn.notify_all(); + } + } +} + +#[cfg(test)] +mod tests { + use std::collections::BTreeMap; + + use super::testing::{Script, names, refused, session}; + use super::{BY_NO_JOB, LOST_JOB, Load, NESTED, OFF_THREAD, buffer_made, finish}; + + /// The buffers each job of the tests makes; job 2 makes none. + const MADE: [u64; 5] = [2, 1, 0, 3, 2]; + + /// Where buffer `k` of job `j` lives, and its length. + fn address(j: usize, k: u64) -> usize { + 0x10_0000 + j * 0x1000 + usize::try_from(k).unwrap() * 0x10 + } + fn length(j: usize, k: u64) -> usize { + 64 * (j + 1) + usize::try_from(k).unwrap() + } + + /// What one thread running the jobs in index order, after `before` buffers, + /// names them: job-major. + fn in_index_order(before: u64) -> BTreeMap { + let mut names = BTreeMap::new(); + let mut n = before; + for (j, &made) in MADE.iter().enumerate() { + for k in 0..made { + names.insert(address(j, k), (n, length(j, k))); + n += 1; + } + } + names + } + + /// The names after a load of [`MADE`]'s jobs, one thread each, made in the + /// order of `order`, with two buffers made before it and one after. + fn load_in(order: &[(usize, u64)]) -> BTreeMap { + let _session = session(); + buffer_made(0x100, 8); + buffer_made(0x200, 16); + let load = Load::begin(); + let script = Script::new(order.to_vec()); + std::thread::scope(|s| { + for (j, &made) in MADE.iter().enumerate() { + let (load, script) = (&load, &script); + s.spawn(move || { + load.job(j, || { + for k in 0..made { + script.step(j, k, || buffer_made(address(j, k), length(j, k))); + } + }); + }); + } + }); + load.end(MADE.len()); + buffer_made(0x300, 32); + assert!(refused().is_empty(), "{:?}", refused()); + names() + } + + /// A buffer's name does not depend on which thread got to it first: the + /// same names for jobs made in order, in reverse, and interleaved, and + /// they are the names one thread running the jobs in order gives. Mutant: + /// one counter for every buffer (the interleavings differ). + #[test] + fn a_loads_buffers_are_numbered_in_job_order_whatever_the_interleaving() { + let mut expected = in_index_order(2); + expected.extend([(0x100, (0, 8)), (0x200, (1, 16)), (0x300, (10, 32))]); + let in_order = [ + (0, 0), + (0, 1), + (1, 0), + (3, 0), + (3, 1), + (3, 2), + (4, 0), + (4, 1), + ]; + let reversed = [ + (4, 0), + (4, 1), + (3, 0), + (3, 1), + (3, 2), + (1, 0), + (0, 0), + (0, 1), + ]; + let mixed = [ + (4, 0), + (3, 0), + (0, 0), + (4, 1), + (1, 0), + (3, 1), + (0, 1), + (3, 2), + ]; + for (name, order) in [ + ("in order", in_order), + ("reversed", reversed), + ("mixed", mixed), + ] { + assert_eq!(load_in(&order), expected, "{name}"); + } + } + + /// Two buffers at one address: the one made later, in time, holds the + /// name, whichever job it belongs to. Mutant: the first buffer made at an + /// address keeps it. + #[test] + fn the_buffer_made_last_at_an_address_holds_the_name() { + let _session = session(); + let load = Load::begin(); + // Job 5 made a buffer there and freed it; job 3's then took the address. + load.job(5, || buffer_made(0x1000, 64)); + load.job(3, || buffer_made(0x1000, 128)); + load.end(6); + assert_eq!(names(), BTreeMap::from([(0x1000, (0, 128))])); + } + + /// The jobs a failed load did not need to run are not counted, and nor + /// are the buffers a load that unwound made after the point it stopped. + #[test] + fn a_load_counts_the_jobs_it_is_told_to() { + let _session = session(); + let load = Load::begin(); + for j in 0..4 { + load.job(j, || buffer_made(0x1000 + j * 0x10, 8)); + } + load.end(2); + assert_eq!( + names(), + BTreeMap::from([(0x1000, (0, 8)), (0x1010, (1, 8))]) + ); + // The next buffer follows the two counted. + buffer_made(0x2000, 8); + assert_eq!(names()[&0x2000], (2, 8)); + // A load dropped before it ends counts every job it ran. + let load = Load::begin(); + load.job(0, || buffer_made(0x3000, 8)); + drop(load); + buffer_made(0x4000, 8); + assert_eq!((names()[&0x3000].0, names()[&0x4000].0), (3, 4)); + } + + /// A buffer with no place in the order is not numbered and the trace says + /// so: made on another thread outside a load, inside a load by no job, by + /// a job of a load that is not the open one, or by a load begun inside + /// another. The names around it do not move. Mutant: it takes a number. + #[test] + fn a_buffer_with_no_place_is_refused_by_name() { + let _session = session(); + buffer_made(0x100, 8); + std::thread::scope(|s| { + s.spawn(|| buffer_made(0x200, 8)); + }); + buffer_made(0x300, 8); + assert_eq!(names(), BTreeMap::from([(0x100, (0, 8)), (0x300, (1, 8))])); + assert_eq!(refused(), BTreeMap::from([(OFF_THREAD, 1)])); + + let load = Load::begin(); + buffer_made(0x400, 8); + let nested = Load::begin(); + nested.job(0, || buffer_made(0x500, 8)); + load.job(0, || buffer_made(0x600, 8)); + nested.end(1); + load.end(1); + assert_eq!( + names(), + BTreeMap::from([(0x100, (0, 8)), (0x300, (1, 8)), (0x600, (2, 8))]) + ); + assert_eq!( + refused(), + BTreeMap::from([(OFF_THREAD, 1), (BY_NO_JOB, 1), (NESTED, 1), (LOST_JOB, 1)]) + ); + let text = finish(); + assert!(text.starts_with(super::VERSION), "{text}"); + for reason in [OFF_THREAD, BY_NO_JOB, NESTED, LOST_JOB] { + assert!(text.contains(&format!("\nrefused {reason}: 1\n")), "{text}"); + } + } +} diff --git a/tools/metal-compare.sh b/tools/metal-compare.sh index 1dfe6c0d5..852e81074 100755 --- a/tools/metal-compare.sh +++ b/tools/metal-compare.sh @@ -47,7 +47,10 @@ # 6 identical up to buffer numbering (a binding names its buffer by allocation order, so a change that allocates one more buffer # shifts every later id; the compare renumbers each side by first use and finds the same binding pattern: the pairs are printed), 3 # differ, 2 refused. A trace rc 6 counts as a difference for the run's code: the trace holds no identity of a buffer but its order, -# so a swap of two buffers' roles at every use also renumbers the same; the summary line keeps the 6 for the reader. +# so a swap of two buffers' roles at every use also renumbers the same; the summary line keeps the 6 for the reader. A trace numbers +# its buffers in the order one thread would make them, whatever threads loaded the weights, so two runs of one binary write one +# trace; a base whose trace is the older format (v1, the weights numbered in thread order) is never rc 0 against a head (rc 6 at +# best, with a note), and a trace with a `refused` line (a buffer that had no place in the numbering) is rc 2. # Runs no gate and takes no lock: it opens this Mac's GPU, which nothing it does changes. set -uo pipefail ROOT=$(cd "$(dirname "$0")/.." && pwd) diff --git a/tools/spawn-allow.txt b/tools/spawn-allow.txt index 301f0ab79..b747330a1 100644 --- a/tools/spawn-allow.txt +++ b/tools/spawn-allow.txt @@ -65,3 +65,4 @@ crates/device/src/pages.rs 2 a test's writer and reader threads, spawned in one crates/metal/src/bin/gate_metal_device/mac.rs 1 the stream clause's two waiter threads, scoped in the gate main, never pinned (macOS has no affinity call) crates/model-hf/src/remote.rs 1 a set's header reads, up to eight at once in a thread scope (network waits, no compute), from the start by name's pick and the fixture bin's main thread, never pinned crates/metal/src/arch/qwen3moe/session.rs 1 the open's pipeline build, scoped beside the weight load on the open's own thread, never pinned (macOS has no affinity call) +crates/metal/src/trace.rs 2 a test's scripted load workers and a test's off-thread buffer, scoped by the test harness, never pinned (macOS has no affinity call) From 52943decd538f00d84097e31173c5ebf093c04bd Mon Sep 17 00:00:00 2001 From: midagedev Date: Tue, 3 Nov 2026 18:28:19 +0900 Subject: [PATCH 2/2] metal: lab-metal runs the trace's numbering tests, and the load fixture has one owner The trace's numbering tests and metal_scan's version test sit behind the encode-trace feature, which no recipe built, so tools/recipes.py orphan-tests found seven tests no gate runs. lab-metal now runs the lib and metal_scan tests with encode-trace. Every default-feature test also runs there, and metal_scan's features become the ones metal-compare builds it with. The review's must-fix (R14): MADE, address and the mixed order now live once, in trace::testing. trace's tests and arena's in_order tests both import them. Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01CjHkkntc2L1xRN21xB39mj --- crates/metal/src/arch/qwen3moe/arena.rs | 19 ++--------- crates/metal/src/trace.rs | 43 ++++++++++++++----------- justfile | 5 +-- 3 files changed, 29 insertions(+), 38 deletions(-) diff --git a/crates/metal/src/arch/qwen3moe/arena.rs b/crates/metal/src/arch/qwen3moe/arena.rs index 23cdf9797..70a011c0b 100644 --- a/crates/metal/src/arch/qwen3moe/arena.rs +++ b/crates/metal/src/arch/qwen3moe/arena.rs @@ -1520,12 +1520,7 @@ mod in_order_tests { /// load makes them: job `j` makes `MADE[j]` of them, buffer `k` at /// `address(j, k)`. #[cfg(feature = "encode-trace")] - const MADE: [u64; 5] = [2, 1, 0, 3, 2]; - - #[cfg(feature = "encode-trace")] - fn address(j: usize, k: u64) -> usize { - 0x10_0000 + j * 0x1000 + usize::try_from(k).expect("a small count") * 0x10 - } + use crate::trace::testing::{MADE, MIXED, address}; /// The names the trace gives after a load of [`MADE`]'s jobs on `workers` /// threads, one buffer made before it and one after; job `j` fails when @@ -1574,17 +1569,7 @@ mod in_order_tests { } } assert_eq!(one, expected); - // The last jobs make their buffers first, the first jobs' last. - let script = Script::new(vec![ - (4, 0), - (3, 0), - (0, 0), - (4, 1), - (1, 0), - (3, 1), - (0, 1), - (3, 2), - ]); + let script = Script::new(MIXED.to_vec()); assert_eq!(loaded(4, Some(&script), None), one); } diff --git a/crates/metal/src/trace.rs b/crates/metal/src/trace.rs index 1a5d40dfc..1eddc33f7 100644 --- a/crates/metal/src/trace.rs +++ b/crates/metal/src/trace.rs @@ -468,6 +468,27 @@ pub(crate) mod testing { use std::sync::{Condvar, Mutex, MutexGuard, PoisonError}; use std::time::Duration; + /// The buffers each job of the numbering tests makes; job 2 makes none. + pub(crate) const MADE: [u64; 5] = [2, 1, 0, 3, 2]; + + /// An order across threads that [`MADE`]'s buffers are made in: the last + /// jobs make theirs first, the first jobs' last. + pub(crate) const MIXED: [(usize, u64); 8] = [ + (4, 0), + (3, 0), + (0, 0), + (4, 1), + (1, 0), + (3, 1), + (0, 1), + (3, 2), + ]; + + /// Where buffer `k` of job `j` lives. + pub(crate) fn address(j: usize, k: u64) -> usize { + 0x10_0000 + j * 0x1000 + usize::try_from(k).expect("a small count") * 0x10 + } + use super::{Name, State, state}; /// The state, reset, held until the guard drops. @@ -541,16 +562,10 @@ pub(crate) mod testing { mod tests { use std::collections::BTreeMap; - use super::testing::{Script, names, refused, session}; + use super::testing::{MADE, MIXED, Script, address, names, refused, session}; use super::{BY_NO_JOB, LOST_JOB, Load, NESTED, OFF_THREAD, buffer_made, finish}; - /// The buffers each job of the tests makes; job 2 makes none. - const MADE: [u64; 5] = [2, 1, 0, 3, 2]; - - /// Where buffer `k` of job `j` lives, and its length. - fn address(j: usize, k: u64) -> usize { - 0x10_0000 + j * 0x1000 + usize::try_from(k).unwrap() * 0x10 - } + /// The length of buffer `k` of job `j`. fn length(j: usize, k: u64) -> usize { 64 * (j + 1) + usize::try_from(k).unwrap() } @@ -623,20 +638,10 @@ mod tests { (0, 0), (0, 1), ]; - let mixed = [ - (4, 0), - (3, 0), - (0, 0), - (4, 1), - (1, 0), - (3, 1), - (0, 1), - (3, 2), - ]; for (name, order) in [ ("in order", in_order), ("reversed", reversed), - ("mixed", mixed), + ("mixed", MIXED), ] { assert_eq!(load_in(&order), expected, "{name}"); } diff --git a/justfile b/justfile index 1ae5d5227..634b0dc27 100644 --- a/justfile +++ b/justfile @@ -90,14 +90,15 @@ mac-test: # The Metal backend's gate, on an Apple Silicon Mac: the device contract's clauses on this Mac's Metal device, run by # tools/mac-gate.sh (the Metal lock, the bound, the refusals). crates/metal links no cuda-oxide, so plain cargo builds it. -# The crate's lib tests (the qwen35moe session's dispatch counts) and its bins' unit tests first, then the device +# The crate's lib tests (the qwen35moe session's dispatch counts) and its bins' unit tests first, the lib's and +# metal_scan's with encode-trace so the trace's numbering tests run beside every default one, then the device # gate, then the qwen35moe kernels' gate (the qwen35moe fixture under $BLOOMERY_FIXTURE_ROOT, or --fixture), then that # gate's mutants, each of which its clause must catch. A lab-* recipe, not a gate-* one: the box has no Metal device, so no landing batch # selects it. The gate and its mutants step each run every registered instance in one process under one bound; # PIN(2026-11-02): on an M1 Pro (16-core GPU) at fourteen instances the gate took 1185 s and the mutants step (780 mutants) # 5461 s, over the 900 s default, so they run at 1800 and 8400 (each about 1.5 times its wall). lab-metal: - cargo test -p bloomery-metal --lib && cargo test -p bloomery-metal --features pipeline-scan --bin metal_scan && cargo test -p bloomery-metal --bin metal_attention_time --bin metal_gemm_time --bin metal_matvec_time --bin generate_metal --bin gate_metal_qwen35moe_kernels && cargo build --release -p bloomery-metal --bin gate_metal_device --bin gate_metal_qwen35moe_kernels && ./tools/mac-gate.sh gate_metal_device && BLOOMERY_GATE_BOUND=1800 ./tools/mac-gate.sh gate_metal_qwen35moe_kernels && BLOOMERY_GATE_BOUND=8400 ./tools/mac-gate.sh gate_metal_qwen35moe_kernels --mutants-all + cargo test -p bloomery-metal --lib --features encode-trace && cargo test -p bloomery-metal --features pipeline-scan,encode-trace --bin metal_scan && cargo test -p bloomery-metal --bin metal_attention_time --bin metal_gemm_time --bin metal_matvec_time --bin generate_metal --bin gate_metal_qwen35moe_kernels && cargo build --release -p bloomery-metal --bin gate_metal_device --bin gate_metal_qwen35moe_kernels && ./tools/mac-gate.sh gate_metal_device && BLOOMERY_GATE_BOUND=1800 ./tools/mac-gate.sh gate_metal_qwen35moe_kernels && BLOOMERY_GATE_BOUND=8400 ./tools/mac-gate.sh gate_metal_qwen35moe_kernels --mutants-all # The Metal session's e2e gate on an Apple Silicon Mac, run by tools/mac-gate.sh over eight rows (`--row`): the qwen35moe # MLX split, then the qwen3moe MLX split (attention layers only, no shared expert, no recurrent store; the clauses that