From e7e42cd70ae91ce912433de5bd3176360bd0483e Mon Sep 17 00:00:00 2001 From: midagedev Date: Sun, 25 Oct 2026 06:55:04 +0900 Subject: [PATCH 1/3] metal session: the open makes the rows arena and builds the prefill pipelines; no host zero of new buffers Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01Q5zTV27LbaVZFTjnkdmTjJ --- crates/metal/src/arch/qwen3moe/rows.rs | 53 ++++--- crates/metal/src/arch/qwen3moe/session.rs | 175 ++++++++++++++-------- 2 files changed, 145 insertions(+), 83 deletions(-) diff --git a/crates/metal/src/arch/qwen3moe/rows.rs b/crates/metal/src/arch/qwen3moe/rows.rs index acdc90612..d85a66f88 100644 --- a/crates/metal/src/arch/qwen3moe/rows.rs +++ b/crates/metal/src/arch/qwen3moe/rows.rs @@ -1,5 +1,5 @@ -//! The rows arena: the buffers a pass of up to `cap` rows touches, opened on -//! the first pass and kept. +//! The rows arena: the buffers a pass of up to `cap` rows touches, opened +//! with the session and kept. //! //! A pass walks the chain at `m` rows with the residual stream, the normed //! input, the P1 projection, the delta numerator or the attention output, @@ -13,8 +13,15 @@ //! the down's output at their widths: they are most of the arena. //! //! The copies a pass leaves after rows inside it ([`SnapBufs`]) are opened -//! on the first pass that takes one, apart from the arena: a pass that takes -//! none never pays for them. +//! apart from the arena, by the open of a session that marks or else by the +//! first pass that takes one: a session that takes none never pays for them. +//! +//! No buffer is written by the host at its open. A new buffer's pages are +//! zeros (Metal clears a buffer it makes), and no kernel of a pass reads a +//! cell no earlier op of the same pass wrote: the GEMMs clamp their row +//! loads to the pass's rows, the routed GEMMs theirs to the tile's places, +//! the route table counts in threadgroup memory it clears, and the pick +//! reads only the router's columns. use bloomery_device::{Buffer, Elems}; @@ -42,8 +49,8 @@ fn refuse(detail: impl Into) -> SessionError { } } -/// A buffer of `n` elements, zeroed. -fn zeroed( +/// A buffer of `n` elements, as Metal makes it. +fn fresh( device: Alloc<'_>, n: usize, what: &str, @@ -51,9 +58,7 @@ fn zeroed( if n == 0 { return Err(refuse(format!("{what} of no elements"))); } - let mut buffer = device.buffer(Elems(n))?; - buffer.zero(); - Ok(buffer) + Ok(device.buffer(Elems(n))?) } /// `a * b`, refused by name when it wraps. @@ -187,28 +192,28 @@ impl RowsArena { // the GEMM's output rows are aligned to. let proj_stride = arena.proj.len().get().next_multiple_of(4); let words = route_table_words(slots, arena.n_experts); - let ids = zeroed(device, cap, "the ids")?; - let x = zeroed( + let ids = fresh(device, cap, "the ids")?; + let x = fresh( device, product(cap, hidden, "the residual rows")?, "the residual", )?; - let xn = zeroed( + let xn = fresh( device, product(cap, hidden, "the normed rows")?, "the normed input", )?; - let proj = zeroed( + let proj = fresh( device, product(cap, proj_stride, "the P1 rows")?, "the P1 output", )?; - let mix = zeroed( + let mix = fresh( device, product(cap, MIX_STRIDE, "the mixer rows")?, "the mixer rows", )?; - let delta_partials = zeroed( + let delta_partials = fresh( device, product(cap, DELTA_PARTIALS_STRIDE, "the delta partials")?, "the delta partials", @@ -225,27 +230,27 @@ impl RowsArena { .ok_or_else(|| refuse(format!("a delta layer of {n_k} key heads"))) }) .try_fold(0, |w, e| e.map(|e| w.max(e)))?; - let delta_prep = zeroed( + let delta_prep = fresh( device, product(cap, delta_prep_stride.max(1), "the delta's prepared rows")?, "the delta's prepared rows", )?; - let logits = zeroed( + let logits = fresh( device, product(cap, LOGIT_STRIDE, "the logit rows")?, "the logits", )?; - let pick_ids = zeroed(device, slots, "the picks")?; - let pick_weights = zeroed(device, slots, "the pick weights")?; - let sorted = zeroed(device, slots, "the sorted slots")?; - let inv = zeroed(device, slots, "the slots' places")?; - let tiles = zeroed(device, words, "the route table")?; - let h = zeroed( + let pick_ids = fresh(device, slots, "the picks")?; + let pick_weights = fresh(device, slots, "the pick weights")?; + let sorted = fresh(device, slots, "the sorted slots")?; + let inv = fresh(device, slots, "the slots' places")?; + let tiles = fresh(device, words, "the route table")?; + let h = fresh( device, product(slots, EXPERT_K, "the pair rows")?, "the pair outputs", )?; - let y = zeroed( + let y = fresh( device, product(slots, hidden, "the down rows")?, "the down outputs", diff --git a/crates/metal/src/arch/qwen3moe/session.rs b/crates/metal/src/arch/qwen3moe/session.rs index 2ec2f417c..9bf59e24c 100644 --- a/crates/metal/src/arch/qwen3moe/session.rs +++ b/crates/metal/src/arch/qwen3moe/session.rs @@ -40,12 +40,13 @@ //! //! [`Session::pass`] runs more than a step's window of ids as one pass: every //! layer once over all the rows, its GEMMs and row kernels over the rows -//! arena, opened on the first pass and kept. The pass is the token at its last -//! position as far as the record goes: its faults land in, and its head writes -//! its pick to, the step record slot of its last row, and the delta state it -//! leaves is the lane a step at the position after it reads. A pass is a -//! command buffer an op, the last one waited: there is no host wait inside the -//! pass, and none past the last op's. A prompt call cuts its ids into passes +//! arena, which the open makes beside the token's. The pass is the token at +//! its last position as far as the record goes: its faults land in, and its +//! head writes its pick to, the step record slot of its last row, and the +//! delta state it leaves is the lane a step at the position after it reads. A +//! pass is a command buffer a layer, each opened at the token boundary and +//! only the last one waited: there is no host wait inside the pass, and none +//! past the last layer's. A prompt call cuts its ids into passes //! at its checkpoint marks ([`Session::prefill`]). A pass may also leave //! copies of the delta layers' lanes and rings after rows inside it: each is //! the state the pass held after that row, and the pass's own bits do not @@ -53,9 +54,10 @@ //! //! Every buffer the kernels bind is registered into the stream's residency //! set when the session allocates it, and the open waits for the set to be -//! requested resident ([`crate::residency`]): the first token finds the -//! weights wired, and a token after a pass or an idle gap finds them still -//! wired. [`Session::residency`] reports the set. +//! requested resident ([`crate::residency`]): the first token and the first +//! pass find the weights and the rows arena wired, and a token after a pass +//! or an idle gap finds them still wired. [`Session::residency`] reports the +//! set, and [`Session::open_split`] how long the open's parts took. //! //! [`Session::step_traced`] is the instrument the gates read nodes through: //! the same ops of a host-fed step, each on a command buffer of its own, @@ -72,6 +74,7 @@ use core::sync::atomic::Ordering; use core::time::Duration; use std::path::Path; use std::sync::{Mutex, PoisonError}; +use std::time::Instant; use bloomery_device::fault::Fault; use bloomery_device::{Buffer, DeviceError, Elem, Elems, Stream}; @@ -240,8 +243,8 @@ pub enum StepMode { Ahead { next: Option }, } -/// What the session's steps did with the token ahead, the host's waits and -/// the command buffers committed. +/// What the session's steps did with the token ahead, the host's waits, the +/// command buffers committed and the passes run. #[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] pub struct StepStats { /// The host's waits on the session's stream: every settle and sync. @@ -258,6 +261,8 @@ pub struct StepStats { /// Steps at the cache's last row that would commit a token ahead, and /// commit none. pub last_row: u64, + /// Passes that ran to their settle. + pub passes: u64, } /// The token committed ahead of its call. @@ -580,12 +585,42 @@ pub struct Session { /// Each delta layer's state lane and conv ring, at the positions a /// prompt call takes. checkpoints: Checkpoints, - /// The device the rows arena is opened on, on the first pass. + /// The device the mark copies are opened on, by the first pass that + /// takes one when the open did not make them. device: Handle, - /// The buffers of a pass, once a pass has run. - rows: Option, + /// The buffers of a pass. + rows: RowsArena, /// How a prompt call feeds its ids. prefill: PrefillMode, + /// How long the open's parts took. + open_split: OpenSplit, +} + +/// When a session opens the copies a pass leaves after rows inside it +/// ([`StateOut`]). +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum MarkCopies { + /// With the session, beside the rows arena: a server's session, whose + /// chat calls mark their messages. + AtOpen, + /// On the first pass that takes one: a session that never marks never + /// holds them. + OnFirstMark, +} + +/// How long each part of a session's open took, on the host's clock. +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +pub struct OpenSplit { + /// The plan read from the file and the weights copied into the arena. + pub load: Duration, + /// The rows arena's buffers allocated, and the mark copies' when the + /// open makes them. + pub rows: Duration, + /// The residency set's commit: the driver wiring every buffer the open + /// made. + pub residency: Duration, + /// The kernels' pipelines built, the prefill attention's among them. + pub pipelines: Duration, } /// What a pass returns of its head. @@ -631,8 +666,10 @@ pub struct Passed { impl Session { /// The session of the file at `path` on `device`, its pipelines from /// `library`, with a cache of `ctx` rows: the plan read from the file's - /// header, every weight copied into the arena, every store at rest. - /// Checkpoints are taken every [`CHECKPOINT_EVERY`] positions. + /// header, every weight copied into the arena, every store at rest, and + /// the rows arena of [`Session::pass_rows`] rows, all resident before + /// the open returns. Checkpoints are taken every [`CHECKPOINT_EVERY`] + /// positions. pub fn open( device: &MetalDevice, library: &Library, @@ -650,6 +687,28 @@ impl Session { ctx: usize, every: u32, ) -> Result { + Self::open_as(device, library, path, ctx, every, MarkCopies::OnFirstMark) + } + + /// [`Session::open_with`] with the mark copies opened as `copies` says. + /// Every buffer is made and the residency set committed here, once: a + /// refusal is a named error and leaves the set as it was, and no pass + /// allocates or builds a pipeline but the first marked pass of a session + /// opened [`MarkCopies::OnFirstMark`]. + pub fn open_as( + device: &MetalDevice, + library: &Library, + path: &Path, + ctx: usize, + every: u32, + copies: MarkCopies, + ) -> Result { + let started = Instant::now(); + let kernels = Kernels::from_library(library)?; + // Built here so a refused build fails the open, never a pass. + kernels.prepare_prefill()?; + let pipelines = started.elapsed(); + let started = Instant::now(); let split = Split::open(path) .map_err(|e| refuse(format!("{} did not open: {e}", path.display())))?; let read = model::arch::qwen35moe::spec::read(&split) @@ -666,23 +725,26 @@ impl Session { &SITES, Read35::Whole, )?; - let kernels = Kernels::from_library(library)?; let mut stream = device.new_stream()?; stream.enable_residency(device.metal())?; let residency = stream .residency() .ok_or_else(|| refuse("the stream made no residency set"))?; let staged = residency.stage(); - let arena = Arena::open( - Alloc::new(device.metal(), &staged), - &split, - &spec, - &plans, - ctx, - )?; - // The open waits for the driver to wire the arena, so the first token - // finds it resident. + let alloc = Alloc::new(device.metal(), &staged); + let arena = Arena::open(alloc, &split, &spec, &plans, ctx)?; + let load = started.elapsed(); + let started = Instant::now(); + let mut rows = RowsArena::open(alloc, &arena, pass_rows_max().min(arena.ctx))?; + if copies == MarkCopies::AtOpen { + rows.snaps = Some(SnapBufs::open(alloc, &arena)?); + } + let rows_t = started.elapsed(); + // The open waits for the driver to wire the arenas, so the first + // token and the first pass find them resident. + let started = Instant::now(); staged.commit(); + let residency_t = started.elapsed(); let ctx_positions = u32::try_from(arena.ctx).map_err(|_| { refuse(format!( "a cache of {} rows: more than a u32 position names", @@ -711,11 +773,23 @@ impl Session { clock: None, checkpoints, device: device.handle().clone(), - rows: None, + rows, prefill: PrefillMode::Batch, + open_split: OpenSplit { + load, + rows: rows_t, + residency: residency_t, + pipelines, + }, }) } + /// How long the open's parts took. + #[must_use] + pub fn open_split(&self) -> OpenSplit { + self.open_split + } + /// The token as the count models it. #[must_use] pub fn token_layers(&self) -> &[Layer] { @@ -895,7 +969,7 @@ impl Session { /// and no more than the cache's rows. #[must_use] pub fn pass_rows(&self) -> usize { - pass_rows_max().min(self.arena.ctx) + self.rows.cap } /// What the residency set of the session's kernel buffers holds and how @@ -908,10 +982,10 @@ impl Session { .stats() } - /// The bytes the rows arena holds, once a pass has opened it. + /// The bytes the rows arena holds. #[must_use] - pub fn rows_arena_bytes(&self) -> Option { - self.rows.as_ref().map(|r| r.bytes) + pub fn rows_arena_bytes(&self) -> u64 { + self.rows.bytes } /// `ids` as one pass at the next positions: every layer once over all the @@ -1006,38 +1080,20 @@ impl Session { if let Some(a) = self.ahead.take() { self.drop_ahead(a)?; } - if self.rows.is_none() { + let rows = &mut self.rows; + if !after.is_empty() && rows.snaps.is_none() { let residency = self .stream .residency() .ok_or_else(|| refuse("the stream has no residency set"))?; let staged = residency.stage(); - let rows = RowsArena::open(Alloc::new(&self.device, &staged), &self.arena, cap)?; + let copies = SnapBufs::open(Alloc::new(&self.device, &staged), &self.arena)?; staged.commit(); - self.rows = Some(rows); + rows.snaps = Some(copies); } - if let Some(rows) = self.rows.as_mut() { - if !after.is_empty() && rows.snaps.is_none() { - let residency = self - .stream - .residency() - .ok_or_else(|| refuse("the stream has no residency set"))?; - let staged = residency.stage(); - let copies = SnapBufs::open(Alloc::new(&self.device, &staged), &self.arena)?; - staged.commit(); - rows.snaps = Some(copies); - } - if let Some(copies) = rows.snaps.as_mut() { - copies.holds = [None; SNAPS]; - } + if let Some(copies) = rows.snaps.as_mut() { + copies.holds = [None; SNAPS]; } - // Built before the first op commits, so a refused build leaves the - // state, rings and cache as they were. - self.kernels.prepare_prefill()?; - let rows = self - .rows - .as_mut() - .ok_or_else(|| refuse("the rows arena did not open"))?; rows.ids.write_at(Elems(0), ids)?; let slot = slot_of(from + rows_n - 1); let head = out == PassOut::Last; @@ -1096,6 +1152,7 @@ impl Session { ) } }; + self.stats.passes += 1; if let Some(fault) = self.arena.faults[slot].read_settled(&settled)? { self.poisoned = Some(fault); return Err(SessionError::Fault(fault)); @@ -1105,7 +1162,7 @@ impl Session { } else { ids[m - 1] }; - if let Some(copies) = self.rows.as_mut().and_then(|r| r.snaps.as_mut()) { + if let Some(copies) = self.rows.snaps.as_mut() { for (held, &r) in copies.holds.iter_mut().zip(after) { *held = Some(from + r + 1); } @@ -1385,7 +1442,7 @@ impl Session { /// The last pass's copies hold no position any more. fn drop_copies(&mut self) { - if let Some(copies) = self.rows.as_mut().and_then(|r| r.snaps.as_mut()) { + if let Some(copies) = self.rows.snaps.as_mut() { copies.holds = [None; SNAPS]; } } @@ -1450,8 +1507,8 @@ impl Session { pub fn checkpoint_copy(&mut self, at: u32) -> Result<(), SessionError> { let held = self .rows + .snaps .as_ref() - .and_then(|r| r.snaps.as_ref()) .and_then(|c| c.holds.iter().position(|&h| h == Some(at))); match held { Some(k) if at <= self.pos => self.take_point(at, Some(k)), @@ -1480,8 +1537,8 @@ impl Session { Some(k) => Some(( k, self.rows + .snaps .as_ref() - .and_then(|r| r.snaps.as_ref()) .ok_or_else(|| refuse("a pass's copy with no copies open"))?, )), }; From 5183926e13cab9d383a76333465f3fe7032ba6ea Mon Sep 17 00:00:00 2001 From: midagedev Date: Sun, 25 Oct 2026 06:55:04 +0900 Subject: [PATCH 2/3] metal bins: the server opens its mark copies with the session; generate_metal prints the open's split Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01Q5zTV27LbaVZFTjnkdmTjJ --- crates/metal/src/bin/bloomery_serve_metal.rs | 19 ++++++++++---- crates/metal/src/bin/generate_metal.rs | 27 +++++++++++++++----- 2 files changed, 34 insertions(+), 12 deletions(-) diff --git a/crates/metal/src/bin/bloomery_serve_metal.rs b/crates/metal/src/bin/bloomery_serve_metal.rs index 7b47f1fae..2a39a46d0 100644 --- a/crates/metal/src/bin/bloomery_serve_metal.rs +++ b/crates/metal/src/bin/bloomery_serve_metal.rs @@ -21,8 +21,10 @@ //! the server marks inside it (`Target::prompt_marked`: a copy the pass //! leaves at that row, the call uncut) and at its end. Any other prefix runs //! from position 0. A prompt span of more than a step's window runs in passes -//! of up to `Session::pass_rows` ids, a shorter one a host-fed step an id. No -//! prompt cache, no slot files, no image input. +//! of up to `Session::pass_rows` ids, a shorter one a host-fed step an id. The +//! session opens the copies a marked pass leaves with its rows arena +//! (`MarkCopies::AtOpen`), so the first request pays nothing the later ones +//! do not. No prompt cache, no slot files, no image input. //! //! A bad flag or value is a named error and exit 64; a refusal of the file, //! the device or the session is a named error and exit 1; an engine failure @@ -42,7 +44,7 @@ mod cli { use std::time::Instant; use bloomery_metal::MetalDevice; - use bloomery_metal::arch::qwen3moe::session::Session; + use bloomery_metal::arch::qwen3moe::session::{CHECKPOINT_EVERY, MarkCopies, Session}; use bloomery_metal::server::TargetEngine; use serve::flag::{self, ApiKeys}; use serve::sampling::sampler_factory; @@ -163,8 +165,15 @@ bloomery-serve-metal -m PATH [--host H] [--port P] [--ctx C] [--chat-template-fi let vocab = Arc::new(Vocab::new(tok)?.with_chat_template(&template)?); let device = MetalDevice::open().map_err(|e| e.to_string())?; let t0 = Instant::now(); - let session = - Session::open(&device, device.library(), path, ctx).map_err(|e| e.to_string())?; + let session = Session::open_as( + &device, + device.library(), + path, + ctx, + CHECKPOINT_EVERY, + MarkCopies::AtOpen, + ) + .map_err(|e| e.to_string())?; eprintln!( "bloomery-serve-metal: opened {} in {:.1} s, ctx {ctx}, checkpoints every {} ({} bytes a point, capacity {})", path.display(), diff --git a/crates/metal/src/bin/generate_metal.rs b/crates/metal/src/bin/generate_metal.rs index 1e0c2462e..5a88bb879 100644 --- a/crates/metal/src/bin/generate_metal.rs +++ b/crates/metal/src/bin/generate_metal.rs @@ -22,8 +22,12 @@ //! `--tokens` spelling, so a run's output feeds back as an input), for //! gates. The summary lines go to stderr, `key=value` one a line: //! -//! model= open_s= device_allocated_bytes= prompt_ids= +//! model= open_s= open_load_s= open_rows_ms= +//! open_residency_ms= open_pipelines_ms= device_allocated_bytes= rows_arena_bytes= +//! prompt_ids= //! prompt_s= prompt_tok_s= decode_tokens= //! decode_passes= decode_s= decode_tok_s= //! stop= decode_first_ms= @@ -42,8 +46,8 @@ //! //! `--prefill steps` feeds the prompt a decode step an id, the same-binary arm //! of the default (`batch`: the prompt in passes of up to -//! [`Session::pass_rows`] ids, a short one by steps); the summary prints the -//! rows arena's bytes (`rows_arena_bytes`) once a pass has opened it. +//! [`Session::pass_rows`] ids, a short one by steps). The session opens its +//! rows arena either way. //! //! `--profile [N]` (default N 32; `--profile-layers` implies it) steps N //! more tokens after the summary, each group of each token on a command @@ -786,12 +790,21 @@ generate_metal (--prompt | --tokens a,b,c) [-n N] [--ctx C] [--print-ids] eprintln!("# development-loop numbers: no lease, no witness"); eprintln!("model={}", a.model.display()); eprintln!("open_s={open_s:.3}"); + let split = session.open_split(); + eprintln!("open_load_s={:.3}", split.load.as_secs_f64()); + eprintln!("open_rows_ms={:.2}", split.rows.as_secs_f64() * 1e3); + eprintln!( + "open_residency_ms={:.2}", + split.residency.as_secs_f64() * 1e3 + ); + eprintln!( + "open_pipelines_ms={:.2}", + split.pipelines.as_secs_f64() * 1e3 + ); eprintln!("device_allocated_bytes={allocated}"); let resident = session.residency(); eprintln!("resident_bytes={}", resident.bytes); - if let Some(bytes) = session.rows_arena_bytes() { - eprintln!("rows_arena_bytes={bytes}"); - } + eprintln!("rows_arena_bytes={}", session.rows_arena_bytes()); eprintln!("prompt_ids={}", prompt.len()); eprintln!("prompt_s={prompt_s:.3}"); eprintln!("prompt_tok_s={:.1}", prompt.len() as f64 / prompt_s); From 1ea0d76ee6b6bc6492b28004312ee4f7358c97dd Mon Sep 17 00:00:00 2001 From: midagedev Date: Sun, 25 Oct 2026 06:55:04 +0900 Subject: [PATCH 3/3] gate_metal_qwen35moe_e2e: r_resident holds that the first pass allocates nothing; passes counted by StepStats Co-Authored-By: Claude Opus 5.5 Claude-Session: https://claude.ai/code/session_01Q5zTV27LbaVZFTjnkdmTjJ --- .../src/bin/gate_metal_qwen35moe_e2e/chain.rs | 1 + .../src/bin/gate_metal_qwen35moe_e2e/pass.rs | 17 +++++++------- .../bin/gate_metal_qwen35moe_e2e/resident.rs | 22 ++++++++++--------- 3 files changed, 22 insertions(+), 18 deletions(-) diff --git a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/chain.rs b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/chain.rs index 8e7e58258..2f80ada3d 100644 --- a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/chain.rs +++ b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/chain.rs @@ -79,6 +79,7 @@ pub fn moved(s: &Session, before: StepStats) -> StepStats { kept: now.kept - before.kept, discarded: now.discarded - before.discarded, last_row: now.last_row - before.last_row, + passes: now.passes - before.passes, } } diff --git a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/pass.rs b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/pass.rs index d9d29588b..de8c3cb5a 100644 --- a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/pass.rs +++ b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/pass.rs @@ -5,7 +5,7 @@ //! - `q_short` (p′): a prompt of at most a step's window of ids is the step //! path: the batch session and the steps session leave the same pick, logits //! row bits, delta lanes, conv rings and key and value rows, and the batch -//! session's rows arena was never opened. +//! session ran no pass. //! - `q_state` (u): the prompt as one pass against the same ids stepped one at //! a time: layer 0's recurrent state and conv ring within [`U_L0_REL`] of //! the stepped run's; every delta layer's state distance printed. @@ -361,13 +361,14 @@ pub fn q_short(device: &MetalDevice, library: &Library, path: &Path) -> Clause { } notes.push(n.to_string()); } - if let Some(bytes) = batch.rows_arena_bytes() { + let passes = batch.stats().passes; + if passes != 0 { return Err(format!( - "a prompt of at most a step's window opened the rows arena ({bytes} bytes)" + "prompts of at most a step's window ran {passes} passes; they are steps" )); } Ok(Verdict::Noted(format!( - "prompts of {} ids leave the steps run's pick, logits, states, rings and cache rows bit for bit; no rows arena opened", + "prompts of {} ids leave the steps run's pick, logits, states, rings and cache rows bit for bit; no pass ran", notes.join(", ") ))) } @@ -378,7 +379,7 @@ pub fn q_state(device: &MetalDevice, library: &Library, path: &Path) -> Clause { let ctx = ids.len() + 1; let mut pass = open(device, library, path, ctx, 2048, PrefillMode::Batch)?; let a = run_prompt(&mut pass, &ids)?; - if pass.rows_arena_bytes().is_none() { + if pass.stats().passes == 0 { return Err("a prompt of 1,024 ids did not take the pass".to_string()); } drop(pass); @@ -718,8 +719,8 @@ pub fn q_layers(device: &MetalDevice, library: &Library, path: &Path) -> Clause 2048, PrefillMode::Batch, )?; - // The rows arena opens at a first pass of other ids, outside the counted - // one: a row an op reads before the op ahead of it wrote holds those. + // A first pass of other ids, outside the counted one, leaves its rows in + // the arena: a row an op reads before the op ahead of it wrote holds those. s.pass(other, PassOut::None, StateOut::default()) .map_err(says)?; s.reset().map_err(says)?; @@ -890,7 +891,7 @@ pub fn q_mark_cut(device: &MetalDevice, library: &Library, path: &Path) -> Claus let mut s = open(device, library, path, ctx, MARK_EVERY, PrefillMode::Batch)?; s.reset().map_err(says)?; Target::prompt_marked(&mut s, &ids, &MARK_CUTS, Want::Argmax).map_err(says)?; - if s.rows_arena_bytes().is_none() { + if s.stats().passes == 0 { return Err("a prompt of 1,024 ids did not take the pass".to_string()); } let layers = s.layer_count(); diff --git a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/resident.rs b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/resident.rs index b68f9c5ab..9b69ac87b 100644 --- a/crates/metal/src/bin/gate_metal_qwen35moe_e2e/resident.rs +++ b/crates/metal/src/bin/gate_metal_qwen35moe_e2e/resident.rs @@ -11,7 +11,8 @@ //! across `Session::open` by the set's bytes, to within one page a buffer //! (the two counts round a buffer differently), so a kernel buffer made //! outside the set shows as bytes the set lacks; -//! - the rows arena of the first pass joins the set the same way; +//! - the open makes the rows arena too, so the first pass allocates nothing +//! on the device and adds nothing to the set; //! - a submit after a gap of [`KEEP`] renews the request once, and a submit //! straight after another does not. @@ -78,14 +79,13 @@ pub fn r_resident(device: &MetalDevice, library: &Library, path: &Path) -> Claus let before_pass = held(device)?; s.pass(&ids, PassOut::Last, StateOut::default()) .map_err(says)?; - let rows = held(device)? - before_pass; let after_pass = s.residency(); - let joined = after_pass.buffers - at_open.buffers; - if !agree(rows, after_pass.bytes - at_open.bytes, joined) || rows == 0 { + let grew = held(device)?.abs_diff(before_pass); + if grew != 0 || after_pass.buffers != at_open.buffers || after_pass.bytes != at_open.bytes { return Err(format!( - "the first pass allocated {rows} bytes on the device; the residency set grew by {} in \ - {joined} buffers", - after_pass.bytes - at_open.bytes + "the first pass moved the device's bytes by {grew} and the residency set from {} \ + buffers of {} bytes to {} of {}; the open makes every buffer a pass binds", + at_open.buffers, at_open.bytes, after_pass.buffers, after_pass.bytes )); } @@ -112,9 +112,11 @@ pub fn r_resident(device: &MetalDevice, library: &Library, path: &Path) -> Claus )); } Ok(Verdict::Noted(format!( - "the open allocated {opened} bytes; the residency set holds {} buffers of them to within a \ - page each; the first pass's rows arena added {rows} in {joined}; a step after {:?} idle renewed the request once, two \ - steps straight after the pass renewed none", + "the open allocated {opened} bytes, the rows arena's {} among them; the residency set \ + holds {} buffers of them to within a page each; the first pass allocated nothing and \ + added nothing to the set; a step after {:?} idle renewed the request once, two steps \ + straight after the pass renewed none", + s.rows_arena_bytes(), at_open.buffers, KEEP + GAP_MARGIN )))