Struct Gpu 

pub struct Gpu<'a> {

    pub p: &'a Packed,
    ctx: MetalContext,
    pipes: Pipes,
    dense: Retained<ProtocolObject<dyn MTLBuffer>>,
    res: Pool,
    activity: ExpertActivity,
    read_tracker: ReadTracker,
    activity_started: Instant,
    phase_timer: Option<PhaseTimer>,
    pool_file: File,
    pool_file_nocache: File,
    ngram_file: File,
    ngram_prefetch: RefCell<Option<JoinHandle<()>>>,
    step_no: u64,
    pending: Option<PendingRead>,
    low_bit_store: Option<Layout>,
    pub all_low_bits: bool,
    event: Retained<ProtocolObject<dyn MTLSharedEvent>>,
    event_base: u64,
    event_cpu: Retained<ProtocolObject<dyn MTLSharedEvent>>,
    event_res: Retained<ProtocolObject<dyn MTLSharedEvent>>,
    event_cpu_base: u64,
    slot_tab: Retained<ProtocolObject<dyn MTLBuffer>>,
    wmap: Retained<ProtocolObject<dyn MTLBuffer>>,
    layers: Vec<GLayer>,
    mtp: Option<Mtp>,
    embed: Q,
    final_mixer: Hc,
    lm_head: Q,
    scratch: Scratch,
    pub max_t: usize,
    pub pos: usize,
    tokens: Vec<u32>,
    batch_pos: usize,
    batch_nb: usize,
    batch_snap: bool,
    mtp_len: usize,
    pub last_experts: Vec<Vec<u32>>,
    pub expert_history: Vec<Vec<Vec<u32>>>,
    pub route_history: Vec<RouteStep>,
    dump_states: bool,
    pub state_history: Vec<Vec<(Vec<f32>, Vec<u32>)>>,
    last_states: Vec<(Vec<f32>, Vec<u32>)>,
    last_routes: Vec<Vec<u32>>,
    last_route_w: Vec<Vec<f32>>,
    last_miss: Vec<Vec<u32>>,
    pub ngram_gather_s: Cell<f64>,
    pub ngram_ms: Vec<f64>,
    pub step_ms: Vec<f64>,
    pub gpu_ms: Vec<f64>,
    pub io_ms: Vec<f64>,
    pub misses: Vec<usize>,
    pub miss_bytes: Vec<usize>,
    step_misses: usize,
    step_miss_bytes: usize,
    step_set_s: f64,
    step_read_s: f64,
    folded_mtp: Option<Vec<u32>>,
    step_warm: usize,
    pub warm: Vec<usize>,
    pub set_ms: Vec<f64>,
    pub read_ms: Vec<f64>,
    pub lookahead_hit: Vec<f64>,
    pub lookahead_issued: Vec<usize>,
    pub la_log: Vec<LaEntry>,
    la_pending: Vec<LaEntry>,
    log_la: bool,
    cut_w: f32,
    step_cut: usize,
    pub cut: Vec<usize>,
    inflight: Vec<(Landed, usize)>,
    lookahead: bool,
    spin_wait: bool,
    fake_experts: bool,
    skip: Vec<String>,
    pub dispatches: Vec<usize>,
    pub gpu_idle_ms: Vec<f64>,
    pub rows: Vec<usize>,
    pub mtp_ms: Vec<f64>,
    dispatch_count: Cell<usize>,
    pf: Option<PrefillScratch>,
    ring: Retained<ProtocolObject<dyn MTLBuffer>>,
    prefill_reserved_bytes: usize,
    pub prefill_stats: Vec<ChunkStats>,
}

Fields

p: &'a Packed

ctx: MetalContext

pipes: Pipes

dense: Retained<ProtocolObject<dyn MTLBuffer>>

res: Pool

LRU expert pool: wired copy slots by default, or mapped residency set entries with the developer override. Cached reads serve mapped entries; uncached reads fill copy slots and the prefill ring.

activity: ExpertActivity

read_tracker: ReadTracker

activity_started: Instant

phase_timer: Option<PhaseTimer>

pool_file: File

pool_file_nocache: File

ngram_file: File

The n-gram table, for touching rows’ pages ahead of the gather.

ngram_prefetch: RefCell<Option<JoinHandle<()>>>

step_no: u64

pending: Option<PendingRead>

Background read for the lookahead prediction of the next layer, with the records to add to the set once it lands.

low_bit_store: Option<Layout>

Decode synchronous misses and prefill ring misses read the low-bit record (–miss-experts 3|2) when the store exists; lookahead reads stay 4-bit unless all_low_bits.

all_low_bits: bool

–experts 3|2: every expert record is low-bit, so the pool holds more of them. Prefill uses the same records and retains its most-used experts for decode.

event: Retained<ProtocolObject<dyn MTLSharedEvent>>

GPU <-> CPU handshake: the GPU signals after each layer’s router and waits for the CPU to publish that layer’s pool slots in slot_tab.

event_base: u64

event_cpu: Retained<ProtocolObject<dyn MTLSharedEvent>>

CPU -> GPU signals of the prefill expert stream (a second event, so each event has one writer and its values stay monotonic).

event_res: Retained<ProtocolObject<dyn MTLSharedEvent>>

GPU -> CPU: a block’s resident expert part is done (value seq+2).

event_cpu_base: u64

slot_tab: Retained<ProtocolObject<dyn MTLBuffer>>

layer GPU addresses (u64) of the union of routed experts’ records (entry 62: resident count, 63: union size) and layerSLOT_STRIDE per-row routing weights, written by the CPU. Row n_layers is the MTP’s.

wmap: Retained<ProtocolObject<dyn MTLBuffer>>

layers: Vec<GLayer>

mtp: Option<Mtp>

embed: Q

final_mixer: Hc

lm_head: Q

scratch: Scratch

max_t: usize

pos: usize

Committed positions.

tokens: Vec<u32>

Committed tokens, plus the rows of the step in flight.

batch_pos: usize

The step in flight: first position and row count.

batch_nb: usize

batch_snap: bool

mtp_len: usize

Positions held by the MTP head’s KV cache.

last_experts: Vec<Vec<u32>>

Router choices of the last step, per layer (union over rows).

expert_history: Vec<Vec<Vec<u32>>>

Router choices of every step, [step][layer][union] (cache studies).

route_history: Vec<RouteStep>

Per step: the rows’ tokens and, per serviced layer, every row’s top-k expert ids (nb * k, row-major), for offline predictor studies.

dump_states: bool

CHERENKOV_DUMP_STATES: per step, row 0’s router input at every serviced block (hidden floats) alongside its top-k, for training an expert predictor offline.

state_history: Vec<Vec<(Vec<f32>, Vec<u32>)>>

last_states: Vec<(Vec<f32>, Vec<u32>)>

last_routes: Vec<Vec<u32>>

last_route_w: Vec<Vec<f32>>

last_miss: Vec<Vec<u32>>

Expert ids of each serviced layer’s misses (records read in-step).

ngram_gather_s: Cell<f64>

Seconds spent gathering n-gram rows from the mapped table (CPU, page faults included), cumulative; ngram_ms is the per-step view.

ngram_ms: Vec<f64>

step_ms: Vec<f64>

gpu_ms: Vec<f64>

Command-buffer spans per step, including event waits.

io_ms: Vec<f64>

Time spent waiting for synchronous expert fetches, per step.

misses: Vec<usize>

Records fetched from disk at exact routing time, per step.

miss_bytes: Vec<usize>

step_misses: usize

step_miss_bytes: usize

step_set_s: f64

Seconds in residency-set bookkeeping (acquire, add, commit) and blocked on miss reads, per step.

step_read_s: f64

folded_mtp: Option<Vec<u32>>

Argmax per row of the MTP head folded into the last trunk step (step_rows with fold), consumed by mtp_draft.

step_warm: usize

Records needed this step whose pages were still in the page cache (no read).

warm: Vec<usize>

set_ms: Vec<f64>

read_ms: Vec<f64>

lookahead_hit: Vec<f64>

One-layer lookahead routing: fraction of a layer’s exact experts predicted by the approximate router run one layer earlier, and the records the prediction fetched.

lookahead_issued: Vec<usize>

la_log: Vec<LaEntry>

Every lookahead prediction with its outcome (diagnostics), and the ones awaiting the next layer’s exact routing.

la_pending: Vec<LaEntry>

log_la: bool

cut_w: f32

Deadline policy (–cut-weak w, default 0 = off): once the GPU has finished a block’s resident experts, a missing expert whose weight is below w in every row and whose read has not landed is cut from the late part instead of waited for; stronger ones are waited for. Cut reads finish in the background and are joined at step end.

step_cut: usize

cut: Vec<usize>

inflight: Vec<(Landed, usize)>

lookahead: bool

spin_wait: bool

CHERENKOV_SPIN=0: block on the shared event instead of spinning.

fake_experts: bool

Diagnostics: CHERENKOV_FAKE=experts skips disk fetches (garbage numerics, true GPU timing); CHERENKOV_SKIP=stage,… skips stages (mixer, experts, shared, lmhead) to attribute GPU time.

skip: Vec<String>

dispatches: Vec<usize>

Kernel dispatches per step and GPU idle time waiting on the CPU within a step.

gpu_idle_ms: Vec<f64>

Legacy name: CPU service wall time, which overlaps resident GPU work.

rows: Vec<usize>

Rows per trunk step and time of MTP draft passes, per step.

mtp_ms: Vec<f64>

dispatch_count: Cell<usize>

pf: Option<PrefillScratch>

Prefill engine scratch, allocated for the first long prompt and released after it, and the ring the expert stream cycles through (the pool itself keeps each layer’s most-used experts).

ring: Retained<ProtocolObject<dyn MTLBuffer>>

prefill_reserved_bytes: usize

Per prefill chunk: tokens, seconds, expert records streamed, seconds waiting on fetches, and GPU seconds in (DeltaNet blocks, attention blocks, expert streams, MTP).

prefill_stats: Vec<ChunkStats>

Implementations

impl Gpu<'_>

pub(super) fn record_cut_eligible( &mut self, layer: usize, records: &[usize], missing: &[bool], weights: &[f32], )

pub fn expert_activity(&self) -> &ExpertActivity

pub(super) fn record_routed_experts( &mut self, layer: usize, ids: &[u32], experts: &[u32], )

pub(super) fn record_read_plan( &mut self, plan: &mut ReadPlan, need: &[usize], source: ReadSource, )

pub fn copy_expert_activity(&self, snapshot: &mut ExpertActivity)

pub(super) fn record_quant_selection( &mut self, layer: usize, ids: &[u32], experts: &[u32], )

pub(super) fn record_prediction( &mut self, layer: usize, predicted: &[u32], selected: &[u32], )

impl Gpu<'_>

pub(super) fn attention_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, a: &Attn, base_pos: usize, nb: usize, )

Full attention over nb rows at positions base_pos..: dense while a row’s context fits the indexer budget, otherwise over the blocks the QSA indexer selects for that row. Input prepped in scratch.hc.h1; output in scratch.mix_out.

impl Gpu<'_>

pub fn step_rows( &mut self, tokens: &[u32], snap: bool, fold_mtp: bool, ) -> Result<Vec<u32>>

Run tokens at positions pos.. as one batch; returns the greedy argmax after each row. Nothing is committed until commit. snap keeps rollback snapshots after each row but the last. With fold_mtp, the MTP head’s first pass runs in the same command buffer over all nb rows, pairing row b with the trunk’s own prediction for it (the token that follows any accepted row), so mtp_draft needs no round trip for the first draft.

fn encode_ple_before_block( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, layer: &GLayer, nb: usize, pos: usize, pending: &mut Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, ) -> Result<()>

PLE needs the pending expert output injected before its convolution.

fn lookahead_layer( &self, index: usize, count: usize, fold_mtp: bool, ) -> Option<&GLayer>

One-block lookahead; deeper routing raised misses about 40%. The last trunk layer predicts the folded MTP block when enabled.

pub fn step(&mut self, token: u32) -> Result<u32>

One token, committed: the greedy next token.

impl Gpu<'_>

pub(super) fn deltanet_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, d: &Delta, nb: usize, snap_after: usize, )

Gated DeltaNet over nb rows; snapshots the state after rows 0..snap_after-1 for rollback. Input prepped in scratch.hc.h1.

impl Gpu<'_>

pub(super) fn bind( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, index: usize, buf: &Retained<ProtocolObject<dyn MTLBuffer>>, offset: usize, )

pub(super) fn dispatch( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, pso: &Retained<ProtocolObject<dyn MTLComputePipelineState>>, setup: impl FnOnce(&ProtocolObject<dyn MTLComputeCommandEncoder>), grid: usize, tg: usize, threadgroups: bool, )

pub(super) fn prep_h( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, x: &Retained<ProtocolObject<dyn MTLBuffer>>, x_off: usize, in_dim: u32, nb: usize, set: &HalfSet, )

Half even/odd streams + group sums of nb rows of x (f32, in_dim wide, starting at byte offset x_off) into set.

pub(super) fn qmv_h( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, q: &Q, y: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, set: &HalfSet, )

y[b] = W x[b] over nb prepped rows (chunks of up to 8 rows).

pub(super) fn zero( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, x: &Retained<ProtocolObject<dyn MTLBuffer>>, n: u32, )

pub(super) fn add( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, x: &Retained<ProtocolObject<dyn MTLBuffer>>, r: &Retained<ProtocolObject<dyn MTLBuffer>>, n: usize, )

pub(super) fn read_u32( &self, buf: &Retained<ProtocolObject<dyn MTLBuffer>>, n: usize, ) -> Vec<u32>

pub(super) fn read_f32( &self, buf: &Retained<ProtocolObject<dyn MTLBuffer>>, n: usize, ) -> Vec<f32>

impl Gpu<'_>

pub(super) fn router_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, moe: &Moe, nb: usize, x: &Retained<ProtocolObject<dyn MTLBuffer>>, logits: &Retained<ProtocolObject<dyn MTLBuffer>>, idx: &Retained<ProtocolObject<dyn MTLBuffer>>, w: &Retained<ProtocolObject<dyn MTLBuffer>>, k: usize, )

Router over nb rows of x: softmax over experts, top-k.

pub(super) fn experts_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, moe: &Moe, slot_row: usize, nb: usize, part: u32, )

Routed experts of nb rows (the union published in slot table row slot_row at execution time) plus the shared expert, summed per row into scratch.moe_out. Part 0 covers the records resident when the table was published (plus the shared expert); part 1 the ones fetched meanwhile. Input: scratch.hc.mixed (prepped in h1).

impl Gpu<'_>

pub(super) fn group_params(&self, eps: bool, shift: f32) -> GroupParams

pub(super) fn group_norm_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, x: &Retained<ProtocolObject<dyn MTLBuffer>>, x_off: usize, w: T, y: &Retained<ProtocolObject<dyn MTLBuffer>>, n: u32, groups: u32, shift: f32, nb: usize, )

Grouped RMSNorm of nb rows (each groups slices of n), weight shifted by shift (1.0 for raw HF norm weights).

pub(super) fn hc_read_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, hc: &Hc, nb: usize, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, hyper_off: usize, pending: Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, bufs: &HcBufs, inject: bool, )

Gated residual read of nb rows of hyper for one block: the block input lands in bufs.mixed and, prepped, in bufs.h1; injection logits in bufs.inj. pending is a block output not yet injected into the stream; the fused norm applies it first (in place).

pub(super) fn inject_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, out: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, )

impl<'a> Gpu<'a>

pub fn load(p: &'a Packed, max_t: usize, options: &Options) -> Result<Self>

pub(crate) fn load_bounded( p: &'a Packed, max_t: usize, options: &Options, reserve: usize, memory_bytes: Option<usize>, prefill_rows: usize, ) -> Result<Self>

impl Gpu<'_>

pub fn allocated_bytes(&self) -> u64

pub fn memory_stats(&self) -> MemoryStats

fn kv_index_capacity_bytes(&self) -> usize

Reserved KV and attention-index bytes, including the enabled MTP layer.

impl Gpu<'_>

pub fn mtp_draft(&mut self, next: &[u32], chain: usize) -> Result<Vec<u32>>

MTP draft head after a commit of next.len() rows: row r pairs the trunk’s wide residual at batch_pos + r with next[r], the token that follows it. Fills the head’s KV cache for those positions and returns chain greedy drafts for the positions after the last committed token (each chained draft re-enters the head on its own residual).

pub(super) fn mtp_pass( &mut self, tokens: &[u32], rows: usize, base_pos: usize, from_trunk: bool, src_row: usize, ) -> Result<Vec<u32>>

One MTP pass over rows rows: hidden input from the trunk’s wide residual rows 0.. (from_trunk) or from the head’s own residual row src_row; tokens in tokens; positions base_pos… Returns the argmax per row (logits stay in scratch.mtp_logits).

pub(super) fn encode_mtp_prelude( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, mtp: &Mtp, rows: usize, ids_in: usize, src: &Retained<ProtocolObject<dyn MTLBuffer>>, off: usize, )

The MTP head’s input stream for rows rows: tokens from scratch.ids[ids_in..], residual rows from src (byte offset off), folded into scratch.mtp_hyper.

pub fn mtp_chain(&mut self, draft: u32) -> Result<u32>

A second draft chained on the MTP residual left by prefill_chunk.

impl Gpu<'_>

pub(super) fn phase_encoder( &self, cb: &ProtocolObject<dyn MTLCommandBuffer>, start: Option<usize>, end: usize, ) -> Result<Retained<ProtocolObject<dyn MTLComputeCommandEncoder>>>

pub(super) fn phase_clock(&self) -> Option<ClockSample>

pub(super) fn phase_cpu_mark(&self, slot: usize, phase: CpuPhase)

pub(super) fn collect_trunk_phases( &mut self, layers: usize, mtp: Option<usize>, clock: Option<ClockSample>, )

pub(super) fn collect_phases( &mut self, slots: &[(usize, usize)], before: Option<ClockSample>, )

impl Gpu<'_>

pub(super) fn ngram_ids_at(&self, pl: &Ple, t: usize) -> Vec<u64>

The 16 n-gram row ids of the token at index t of self.tokens.

pub(super) fn ngram_prefetch_start(&self, t0: usize, n: usize)

Start pulling the n-gram rows of tokens t0..t0+n into the page cache on a thread pool, so the block’s gather (which faults the mapping one row at a time) finds them warm. The ids only depend on the tokens, which are known when a step or chunk starts.

pub(super) fn ngram_prefetch_join(&self)

pub(super) fn ple_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, pl: &Ple, nb: usize, base_pos: usize, ) -> Result<()>

PLE block over nb rows at positions base_pos..: n-gram rows are gathered on the CPU from self.tokens (which must hold the batch).

impl Gpu<'_>

pub(super) fn head_b( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, mixer: &Hc, nb: usize, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, hyper_off: usize, pending: Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, logits: &Retained<ProtocolObject<dyn MTLBuffer>>, ids_out: usize, )

Final mixer read of nb rows of hyper (injecting pending first), LM head into logits [nb][vocab], greedy argmax of each row into ids[ids_out + b].

impl Gpu<'_>

pub fn commit(&mut self, n: usize) -> Result<()>

Keep the first n rows of the step in flight: roll the recurrent state back to after row n-1 if rows were rejected, advance the position. KV rows and the PLE ring are positional and get overwritten by the next step.

pub fn reset(&mut self)

Forget the sequence: recurrent states, conv histories and the PLE ring are zeroed; KV rows are overwritten positionally.

pub(super) fn prefix_regions( &self, pos: usize, mtp_len: usize, ) -> Vec<(&Retained<ProtocolObject<dyn MTLBuffer>>, usize, usize)>

pub(crate) fn prefix_state_bytes(&self) -> usize

pub(crate) fn state_bytes_at(&self, pos: usize, mtp_len: usize) -> usize

pub(crate) fn save_prefix(&self) -> PrefixState

pub(crate) fn save_into(&self, checkpoint: &mut Option<PrefixState>)

Reuse suspended-request buffers; only initialized prefixes are copied.

pub(crate) fn restore_prefix(&mut self, state: &PrefixState) -> Result<()>

pub(crate) fn reset_request(&mut self)

Start another request without retaining unbounded profiling history. Weights and expert residency remain loaded; sequence state is independent.

pub(crate) fn clear_profile(&mut self)

Server scheduling publishes counters each turn instead of retaining a log.

impl Gpu<'_>

pub(super) fn pf_alloc( &self, rows: usize, all_logits: bool, ) -> Result<PrefillScratch>

pub fn prefill_rows_fit(&self, all_logits: bool) -> Result<usize>

Maximum chunk that fits the configured budget, or the device recommendation.

impl Gpu<'_>

pub(super) fn pf_attention( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, a: &Attn, base: usize, t: usize, pf: &PrefillScratch, )

Full attention over t rows at positions base..: GEMM form in sub-chunks of QS queries, with the QSA mask for rows past the budget. Input pf.mixed, output pf.mix_out.

impl Gpu<'_>

pub(super) fn pf_deltanet( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, d: &Delta, t: usize, pf: &PrefillScratch, )

Gated DeltaNet over t rows (sequential scan). pf.mixed -> pf.mix_out.

impl Gpu<'_>

fn prepare_expert_jobs( &mut self, record_layer: usize, t: usize, pf: &PrefillScratch, ) -> Result<Vec<ExpertJob>>

Group prompt rows by expert and reserve the resident/ring destinations. This preserves expert order, stable usage ranking and CSR row order.

pub(super) fn pf_experts( &mut self, moe: MoeRef, t: usize, pf: &PrefillScratch, ) -> Result<(usize, usize, f64, f64)>

The MoE of one block over t rows: shared expert over all rows, then the routed experts one at a time (their tokens gathered, computed, scatter-added). The layer’s most-used experts (up to the pool’s per-layer share) go into the residency set and stay for decode; the rest stream through the ring. Returns (records fetched, bytes fetched, seconds waiting on ring slots, GPU seconds).

fn read_expert_batch( &mut self, record_layer: usize, batch: &[ExpertJob], ring_base: usize, stride: usize, miss_kind: u8, ) -> Result<(usize, usize)>

Read one batch into its reserved destinations before publishing its event.

fn ring_read( &mut self, layer: usize, job: &ExpertJob, destination: usize, ) -> RecordRead

impl Gpu<'_>

pub(super) fn pf_inject( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, out: &Retained<ProtocolObject<dyn MTLBuffer>>, inj: &Retained<ProtocolObject<dyn MTLBuffer>>, t: usize, )

pub(super) fn pf_hc_read( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, hc: &Hc, t: usize, pf: &PrefillScratch, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, pending: Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, inject: bool, )

Gated residual read over t rows: pf.mixed (block input), pf.inj.

impl Gpu<'_>

pub(super) fn pf_ple( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, pl: &Ple, base: usize, t: usize, pf: &PrefillScratch, ) -> Result<()>

PLE block over t rows. Nothing else is live between blocks, so its buffers alias the stream-wide scratch: e -> mix_out, key -> normed (and the output, once the key is normed), keyn -> u, value -> mixed, query -> qkv, gated -> fh, gvn -> mtp_hyper.

impl Gpu<'_>

pub(super) fn qmm_from( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, wb: &Retained<ProtocolObject<dyn MTLBuffer>>, q: &Q, x: &Retained<ProtocolObject<dyn MTLBuffer>>, y: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, )

y[b] = W x[b] for nb rows with the simdgroup-matrix GEMM: the x buffer must hold rows padded to the token tile (32).

pub(super) fn expert_qmm_from( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, wb: &Retained<ProtocolObject<dyn MTLBuffer>>, q: &Q, x: &Retained<ProtocolObject<dyn MTLBuffer>>, y: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, bits: u32, )

Low-bit expert records share the tiled dispatch and output layout, but use their own unpacking kernels. Dense projections stay Q4.

fn qmm_tiled( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, wb: &Retained<ProtocolObject<dyn MTLBuffer>>, q: &Q, x: &Retained<ProtocolObject<dyn MTLBuffer>>, y: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, pipes: [&Retained<ProtocolObject<dyn MTLComputePipelineState>>; 3], )

pub(super) fn qmm( &self, enc: &ProtocolObject<dyn MTLComputeCommandEncoder>, q: &Q, x: &Retained<ProtocolObject<dyn MTLBuffer>>, y: &Retained<ProtocolObject<dyn MTLBuffer>>, nb: usize, )

impl<'a> Gpu<'a>

fn pf_layer( &mut self, li: Option<usize>, base: usize, t: usize, pf: &PrefillScratch, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, pending: Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, ) -> Result<(usize, usize, f64, f64, f64)>

One decoder block over t rows of hyper: everything up to the router in one command buffer, then the expert stream. Returns (records fetched, bytes fetched, ring wait seconds, block GPU seconds, expert stream GPU seconds).

pub fn prefill_chunk( &mut self, tokens: &[u32], next_after: Option<u32>, all_logits: bool, ) -> Result<(u32, u32)>

Run one prompt chunk at positions pos.. through the trunk and the MTP head, committing everything. next_after is the token that follows the chunk when known (the next chunk’s first token); the trunk’s own prediction pairs with the last row otherwise. With all_logits every row’s logits are kept for checks. Returns the greedy token after the chunk and the MTP’s first draft.

pub fn prefill_release(&mut self)

Release the prefill scratch (a few hundred MB to GB).

pub fn debug_engine_blocks(&self, r: usize) -> Vec<u32>

Debug: blocks selected for row r of the last prefill chunk (its last attention layer), from the block bitmask.

pub fn debug_row_vis(&self, r: usize) -> Vec<u32>

Debug: visible tokens of row r of the last row-batched step (its last attention layer).

pub fn pf_logits_row(&self, r: usize) -> &[f32]

Row r of the last prefill chunk’s logits (all-logits mode).

impl Gpu<'_>

pub(super) fn join_inflight(&mut self) -> Result<()>

Wait for reads the deadline policy stopped waiting on, then make their records usable (nothing reuses their slots before this).

pub(super) fn join_pending(&mut self) -> Result<()>

Join the background read of predicted records and add them to the set (they are in memory once the read is done).

fn prefetch_async(&mut self, record_layer: usize, ids: &[u32]) -> Result<usize>

Read predicted records for the next layer on a background thread while the GPU runs the current layer’s experts; they join the set at the next service. Returns records read.

fn wait_for_router(&self, ready: u64, slot_row: usize) -> Result<()>

Wait until the GPU releases the routing scratch for CPU reads.

fn log_lookahead( &mut self, record_layer: usize, slot_row: usize, experts: &[u32], indices: &[u32], weights: &[f32], )

pub(super) fn encode_block( &self, cb: &ProtocolObject<dyn MTLCommandBuffer>, enc: &mut Retained<ProtocolObject<dyn MTLComputeCommandEncoder>>, layer: &GLayer, slot_row: usize, base_pos: usize, nb: usize, snap_after: usize, hyper: &Retained<ProtocolObject<dyn MTLBuffer>>, pending: Option<&Retained<ProtocolObject<dyn MTLBuffer>>>, lookahead: Option<&GLayer>, seq: u64, ) -> Result<()>

One decoder block (attention or DeltaNet, then MoE) over nb rows of hyper, with the event handshake around the expert dispatch. The block’s MoE output is left in scratch.moe_out for the next fused norm to inject.

fn miss_record_bytes(&mut self, need: &[usize]) -> usize

Fetched records use the miss precision; resident records keep their kind.

fn wait_for_deadline( &mut self, slot_row: usize, resident_done: u64, flags: &[Landed], need: &[usize], rids: &[usize], order: &[usize], n_res: usize, wmax: &[f32], ) -> Result<()>

Wait for required reads and truncate late weak records at the deadline. The caller has published the resident table and has not released misses.

fn finish_miss_reads( &mut self, need: &[usize], flags: &[Landed], deadline: bool, ) -> Result<()>

Cut reads stay in flight until step end; publish only completed records.

pub(super) fn service_block( &mut self, record_layer: usize, slot_row: usize, nb: usize, seq: u64, predicted: &mut VecDeque<(usize, Vec<u32>)>, lookahead: Option<usize>, io_s: &mut f64, turn_s: &mut f64, ) -> Result<(usize, usize, usize)>

CPU side of one block’s handshake: wait for the router, publish the union of the rows’ experts with the resident ones first. Read misses while the GPU computes the resident part, then issue lookahead and release the fetched part. Returns (lookahead hits, total experts, records read for lookahead). See BlockSignals for the event contract.

impl<'a> Gpu<'a>

pub fn has_mtp(&self) -> bool

fn skips(&self, stage: &str) -> bool

fn record_id(&self, record_layer: usize, expert: u32) -> usize

pub fn pool_resident(&self) -> usize

Records resident in the expert pool (the residency set).

pub fn pool_slots(&self) -> usize

pub fn throttle_ms(&self) -> Result<f64>

Dependent-FMA chain time in ms (higher = deeper clock throttle).

pub fn logits_row(&self, r: usize) -> &[f32]

Trunk logits of row r of the last step.

pub fn logits(&self) -> &[f32]

pub(crate) fn last_logits_row(&self) -> usize

Batched prefill writes one final row; row stepping retains every row.

pub fn mtp_logits_row(&self, r: usize) -> &[f32]

MTP logits of row r of the last draft pass.

pub fn allocated_gb(&self) -> f64

pub fn pool_bytes(&self) -> usize

What Metal will keep resident at once on this machine.

fn kv_layer_bytes(&self, n: usize) -> usize

Positional KV-cache bytes for one attention layer at n positions: the q8 key/values (kc, vc) with their q4 scales, the QSA index key cache and the compressed block keys. Mirrors prefix_regions.

fn trunk_attn_layers(&self) -> usize

Trunk attention (KV) layers; DeltaNet layers hold no positional KV.

fn mtp_attn_layers(&self) -> usize

Whether the MTP head contributes an attention KV cache (0 or 1).

pub fn kv_cache_bytes(&self) -> usize

Bytes of the positional KV caches now in use by the sequence on the GPU (trunk layers at pos, the MTP head at mtp_len).

pub fn kv_cache_capacity(&self) -> usize

Total capacity of the positional KV caches at max_t positions.

pub fn kv_cache_fullness(&self) -> f64

Fraction (0..=1) of the positional KV caches in use by the sequence.

pub fn context_fullness(&self) -> f64

Fraction (0..=1) of the context window in use by the sequence.

pub fn working_set_limit_gb(&self) -> f64