cherenkov/qwen4_exp/
gpu.rs

1//! Metal decode path for qwen4-exp: row-batched steps (prefill chunks and
2//! speculative verification of MTP drafts) over an explicit wired expert
3//! pool.
4//!
5//! Dense weights are one zero-copy buffer over `dense.bin`. Routed experts
6//! live in a fixed pool of page-aligned slots (Metal wires whole buffers per
7//! command buffer), LRU managed here
8//! and filled by uncached parallel reads. One command buffer per step: after
9//! each layer's router the GPU signals a shared event and waits for the CPU
10//! to publish that layer's expert slots; a one-layer lookahead router lets
11//! most fetches run in the background.
12//!
13//! A step runs `nb` consecutive positions at once (activations are
14//! row-major `[nb][...]`). The gated DeltaNet scan snapshots its state after
15//! each row so a partially accepted verify batch can roll back.
16
17use super::cpu::CpuModel;
18use super::packed::Packed;
19use crate::kernels::{BATCH_MSL, FORWARD_MSL};
20use crate::metal::MetalContext;
21use crate::options::{Options, PoolBudget};
22use crate::units::BYTES_PER_GB;
23use anyhow::{Context, Result, ensure};
24use objc2::rc::Retained;
25use objc2::runtime::ProtocolObject;
26use objc2_metal::{
27    MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue, MTLComputeCommandEncoder,
28    MTLComputePipelineState, MTLSharedEvent, MTLSize,
29};
30use std::ffi::c_void;
31use std::ptr::NonNull;
32
33mod activity;
34mod attention;
35mod budget;
36mod decode;
37mod deltanet;
38mod dispatch;
39mod experts;
40mod hyperconnection;
41mod load;
42mod memory;
43mod mtp;
44mod params;
45mod phases;
46mod ple;
47mod sampling;
48mod state;
49
50pub use activity::layers::{PhaseStats, PredictionStats, QuantStats};
51pub use activity::reads::ReadStats;
52pub use activity::{ExpertActivity, ExpertCounters, LayerStats};
53pub use memory::MemoryStats;
54pub use phases::GpuTiming;
55
56use params::*;
57
58pub mod prefill;
59mod residency;
60mod streaming;
61
62use streaming::PendingRead;
63
64const ARGMAX_TGS: usize = 1024;
65const ATTN_TB: usize = 512;
66const ATTN_MAX_WG: usize = 256;
67/// Rows per step (prefill chunk, or 1 + drafts when verifying).
68pub const MAX_NB: usize = 4;
69/// Rollback planes: verify batches may have up to this many draft rows.
70pub const MAX_SNAP: usize = 3;
71/// Slots per layer row of the expert slot table (last entry = union size).
72const SLOT_STRIDE: usize = 64;
73/// ids buffer layout: trunk inputs, trunk argmax, MTP inputs, MTP argmax.
74const IDS_IN: usize = 0;
75const IDS_OUT: usize = 8;
76const IDS_MTP_IN: usize = 16;
77const IDS_MTP_OUT: usize = 24;
78
79type Buf = Retained<ProtocolObject<dyn MTLBuffer>>;
80type Pso = Retained<ProtocolObject<dyn MTLComputePipelineState>>;
81type Enc = ProtocolObject<dyn MTLComputeCommandEncoder>;
82
83/// Tokens, expert IDs, routing weights, and misses for one step.
84/// Keep the tuple layout for compatibility with existing research dumps.
85pub type RouteStep = (Vec<u32>, Vec<Vec<u32>>, Vec<Vec<f32>>, Vec<Vec<u32>>);
86
87/// Affine-Q4 projection: dimensions and byte offsets in its bound weight buffer.
88#[derive(Clone, Copy)]
89struct Q {
90    w: usize,
91    s: usize,
92    b: usize,
93    out: u32,
94    inp: u32,
95}
96
97impl Q {
98    /// Resolve a record-relative projection against a bound buffer.
99    fn at_offset(self, bytes: usize) -> Self {
100        Self {
101            w: self.w + bytes,
102            s: self.s + bytes,
103            b: self.b + bytes,
104            ..self
105        }
106    }
107}
108
109/// Byte offset of a bf16 tensor in dense.bin.
110#[derive(Clone, Copy)]
111struct T(usize);
112
113struct Hc {
114    norm: T,
115    down: Q,
116    up: Q,
117    inject: Option<Q>,
118}
119
120struct Attn {
121    q: Q,
122    k: Q,
123    v: Q,
124    o: Q,
125    qn: T,
126    kn: T,
127    kc: Buf,
128    vc: Buf,
129    /// QSA indexer: the index q/k projection, its norms, the raw index
130    /// key cache `[max_t][ihd]` and the block keys `[max_t/ratio][ihd]`.
131    iqk: Q,
132    iqn: T,
133    ikn: T,
134    ikc: Buf,
135    blk: Buf,
136}
137
138struct Delta {
139    qkv: Q,
140    z: Q,
141    a: Q,
142    b: Q,
143    conv: T,
144    a_log: T,
145    dt_bias: T,
146    norm: T,
147    o: Q,
148    state: Buf,
149    hist: Buf,
150    /// Snapshot planes of `state` / `hist` after each verify row.
151    mid: Buf,
152    mid_hist: Buf,
153}
154
155enum Mix {
156    Attn(Attn),
157    Delta(Delta),
158}
159
160struct Moe {
161    router: T,
162    shared_gate: T,
163    sg: Q,
164    su: Q,
165    sd: Q,
166    record_layer: usize,
167}
168
169struct Ple {
170    key: Q,
171    value: Q,
172    norm_key: T,
173    norm_query: T,
174    norm_conv: T,
175    conv: T,
176    kernel: u32,
177    dilation: u32,
178    span: u32,
179    hist: Buf,
180    /// `[nb][ple_embed_dim]` n-gram rows gathered on the CPU.
181    e: Buf,
182    multipliers: Vec<i64>,
183    head_offsets: Vec<u64>,
184    head_sizes: Vec<u64>,
185}
186
187struct GLayer {
188    attn_hc: Hc,
189    mlp_hc: Hc,
190    mix: Mix,
191    moe: Moe,
192    ple: Option<Ple>,
193}
194
195/// The MTP draft head (see cpu.rs `MtpWeights`).
196struct Mtp {
197    enorm: T,
198    hnorm: T,
199    fc_e: Q,
200    fc_h: Q,
201    layer: GLayer,
202    mixer: Hc,
203}
204
205struct Pipes {
206    // Common kernel library (kernels/common/).
207    qmv_h: Pso,
208    qmv_hn: Pso,
209    prep_h: Pso,
210    embed_rows: Pso,
211    qk_norm_rope_b: Pso,
212    conv_b: Pso,
213    delta_norms: Pso,
214    delta_gates: Pso,
215    delta_scan2: Pso,
216    kv_append_q8: Pso,
217    attn_part2_q8: Pso,
218    attn_combine: Pso,
219    argmax_partial: Pso,
220    argmax_final: Pso,
221    add: Pso,
222    copy_f32: Pso,
223    // qwen4-exp kernel library (kernels/qwen4_exp/).
224    zero: Pso,
225    replicate_b: Pso,
226    group_norm_b: Pso,
227    norm_prep_b: Pso,
228    qmv_silu_b: [Pso; MAX_NB],
229    hc_mix_b: Pso,
230    inject_b: Pso,
231    bf16_matvec_b: Pso,
232    topk_softmax_b: Pso,
233    moe_gate_up_b: [Pso; MAX_NB],
234    moe_act_b: Pso,
235    moe_down_b: [Pso; MAX_NB],
236    moe_combine_b: Pso,
237    ple_gate_b: Pso,
238    ple_conv_b: Pso,
239    mtp_fold: Pso,
240    gate_norm_sigmoid_b: Pso,
241    index_append: Pso,
242    index_blocks: Pso,
243    index_q: Pso,
244    index_score: Pso,
245    index_select: Pso,
246    attn_sel: Pso,
247    // prefill engine
248    qmm_n8: Pso,
249    qmm_n16: Pso,
250    qmm_w: Pso,
251    /// [precision 2/3][token tile 8/16/32].
252    expert_qmm: [[Pso; 3]; 2],
253    silu_mul: Pso,
254    attn_q_stage: Pso,
255    attn_kv_stage: Pso,
256    gemm_hh: Pso,
257    attn_o_scatter: Pso,
258    softmax_sel: Pso,
259    silu_rows: Pso,
260    gather_rows: Pso,
261    scatter_add_rows: Pso,
262    shared_add_rows: Pso,
263    qmv_small_b: Pso,
264}
265
266/// One lookahead prediction and what became of it (CHERENKOV_DUMP_LA).
267#[derive(Clone, Copy, serde::Serialize)]
268pub struct LaEntry {
269    pub step: u64,
270    /// Layer the prediction was for.
271    pub layer: usize,
272    pub expert: u32,
273    /// Largest lookahead router weight over the rows, and best rank.
274    pub weight: f32,
275    pub rank: u32,
276    /// Already in the pool when predicted (no fetch needed).
277    pub resident: bool,
278    /// Used by that layer in this step.
279    pub hit: bool,
280}
281
282/// Half even/odd streams plus per-32 group sums of a row batch, the input
283/// format of the multi-row Q4 matvecs.
284struct HalfSet {
285    xe: Buf,
286    xo: Buf,
287    xsum: Buf,
288}
289
290/// Buffers of one gated-residual read (the main set, or the lookahead's).
291struct HcBufs {
292    normed: Buf,
293    d: Buf,
294    u: Buf,
295    mixed: Buf,
296    inj: Buf,
297    h1: HalfSet,
298    h2: HalfSet,
299}
300
301struct Scratch {
302    ids: Buf,
303    e: Buf,
304    hyper: Buf,
305    hc: HcBufs,
306    la: HcBufs,
307    mix_out: Buf,
308    qg: Buf,
309    k: Buf,
310    v: Buf,
311    attn_out: Buf,
312    attn_parts: Buf,
313    qkv: Buf,
314    z: Buf,
315    a: Buf,
316    b: Buf,
317    kqn: Buf,
318    gbuf: Buf,
319    delta_y: Buf,
320    router: Buf,
321    topk_idx: Buf,
322    topk_w: Buf,
323    la_router: Buf,
324    la_idx: Buf,
325    la_w: Buf,
326    gate_e: Buf,
327    hx: HalfSet,
328    y_e: Buf,
329    moe_out: Buf,
330    ple_key: Buf,
331    ple_keyn: Buf,
332    ple_value: Buf,
333    ple_query: Buf,
334    ple_gated: Buf,
335    ple_gvn: Buf,
336    ple_out: Buf,
337    logits: Buf,
338    mtp_logits: Buf,
339    /// Logits of chained (single-row) MTP passes, kept apart so the first
340    /// pass's rows stay readable for checks.
341    mtp_logits2: Buf,
342    partials: Buf,
343    mtp_hyper: Buf,
344    fe: Buf,
345    fh: Buf,
346    /// Indexer: projection `[nb][qk_dim]`, roped queries `[nb][inh*ihd]`,
347    /// block scores `[nb][max_blocks]`, visible tokens `[nb][vis_stride]`.
348    iqk: Buf,
349    iq: Buf,
350    bscore: Buf,
351    vis: Buf,
352    nvis: Buf,
353    vmask: Buf,
354}
355
356fn set_bytes<T>(enc: &Enc, index: usize, value: &T) {
357    unsafe {
358        enc.setBytes_length_atIndex(
359            NonNull::from(value).cast::<c_void>(),
360            std::mem::size_of::<T>(),
361            index,
362        )
363    };
364}
365
366fn kv_q8_side(max_t: usize, kv_row: usize) -> (usize, usize) {
367    let align = |v: usize| v.div_ceil(16384) * 16384;
368    let qs = align(max_t * kv_row);
369    let sc = align(max_t * kv_row / 32 * 2);
370
371    (qs, qs + sc)
372}
373
374/// Debug: CHERENKOV_LAYERS=N runs only the first N decoder blocks (for
375/// bisecting one execution path against another).
376fn layer_cap() -> usize {
377    static CAP: std::sync::OnceLock<usize> = std::sync::OnceLock::new();
378
379    *CAP.get_or_init(|| {
380        std::env::var("CHERENKOV_LAYERS")
381            .ok()
382            .and_then(|v| v.parse().ok())
383            .unwrap_or(usize::MAX)
384    })
385}
386
387/// Unique entries in first-seen order.
388fn union_of(ids: &[u32]) -> Vec<u32> {
389    let mut u: Vec<u32> = Vec::with_capacity(ids.len());
390
391    for &e in ids {
392        if !u.contains(&e) {
393            u.push(e);
394        }
395    }
396
397    u
398}
399
400pub struct Gpu<'a> {
401    pub p: &'a Packed,
402    ctx: MetalContext,
403    pipes: Pipes,
404    dense: Buf,
405    /// LRU expert pool: wired copy slots by default, or mapped residency
406    /// set entries with the developer override. Cached reads serve mapped
407    /// entries; uncached reads fill copy slots and the prefill ring.
408    res: residency::Pool,
409    activity: ExpertActivity,
410    read_tracker: activity::reads::ReadTracker,
411    activity_started: std::time::Instant,
412    phase_timer: Option<phases::PhaseTimer>,
413    pool_file: std::fs::File,
414    pool_file_nocache: std::fs::File,
415    /// The n-gram table, for touching rows' pages ahead of the gather.
416    ngram_file: std::fs::File,
417    ngram_prefetch: std::cell::RefCell<Option<std::thread::JoinHandle<()>>>,
418    step_no: u64,
419    /// Background read for the lookahead prediction of the next layer,
420    /// with the records to add to the set once it lands.
421    pending: Option<PendingRead>,
422    /// Decode synchronous misses and prefill ring misses read the low-bit record
423    /// (--miss-experts 3|2) when the store exists; lookahead reads
424    /// stay 4-bit unless `all_low_bits`.
425    low_bit_store: Option<super::lowbit::Layout>,
426    /// --experts 3|2: every expert record is low-bit, so the pool
427    /// holds more of them. Prefill uses the same records and retains its
428    /// most-used experts for decode.
429    pub all_low_bits: bool,
430    /// GPU <-> CPU handshake: the GPU signals after each layer's router and
431    /// waits for the CPU to publish that layer's pool slots in `slot_tab`.
432    event: Retained<ProtocolObject<dyn objc2_metal::MTLSharedEvent>>,
433    event_base: u64,
434    /// CPU -> GPU signals of the prefill expert stream (a second event,
435    /// so each event has one writer and its values stay monotonic).
436    event_cpu: Retained<ProtocolObject<dyn objc2_metal::MTLSharedEvent>>,
437    /// GPU -> CPU: a block's resident expert part is done (value seq+2).
438    event_res: Retained<ProtocolObject<dyn objc2_metal::MTLSharedEvent>>,
439    event_cpu_base: u64,
440    /// [layer][SLOT_STRIDE] GPU addresses (u64) of the union of routed
441    /// experts' records (entry 62: resident count, 63: union size) and
442    /// [layer][MAX_NB][SLOT_STRIDE] per-row routing weights, written by
443    /// the CPU. Row `n_layers` is the MTP's.
444    slot_tab: Buf,
445    wmap: Buf,
446    layers: Vec<GLayer>,
447    mtp: Option<Mtp>,
448    embed: Q,
449    final_mixer: Hc,
450    lm_head: Q,
451    scratch: Scratch,
452    pub max_t: usize,
453    /// Committed positions.
454    pub pos: usize,
455    /// Committed tokens, plus the rows of the step in flight.
456    tokens: Vec<u32>,
457    /// The step in flight: first position and row count.
458    batch_pos: usize,
459    batch_nb: usize,
460    batch_snap: bool,
461    /// Positions held by the MTP head's KV cache.
462    mtp_len: usize,
463    /// Router choices of the last step, per layer (union over rows).
464    pub last_experts: Vec<Vec<u32>>,
465    /// Router choices of every step, `[step][layer][union]` (cache studies).
466    pub expert_history: Vec<Vec<Vec<u32>>>,
467    /// Per step: the rows' tokens and, per serviced layer, every row's
468    /// top-k expert ids (nb * k, row-major), for offline predictor studies.
469    pub route_history: Vec<RouteStep>,
470    /// CHERENKOV_DUMP_STATES: per step, row 0's router input at every
471    /// serviced block (hidden floats) alongside its top-k, for training an
472    /// expert predictor offline.
473    dump_states: bool,
474    pub state_history: Vec<Vec<(Vec<f32>, Vec<u32>)>>,
475    last_states: Vec<(Vec<f32>, Vec<u32>)>,
476    last_routes: Vec<Vec<u32>>,
477    last_route_w: Vec<Vec<f32>>,
478    /// Expert ids of each serviced layer's misses (records read in-step).
479    last_miss: Vec<Vec<u32>>,
480    /// Seconds spent gathering n-gram rows from the mapped table (CPU,
481    /// page faults included), cumulative; `ngram_ms` is the per-step view.
482    pub ngram_gather_s: std::cell::Cell<f64>,
483    pub ngram_ms: Vec<f64>,
484    pub step_ms: Vec<f64>,
485    /// Command-buffer spans per step, including event waits.
486    pub gpu_ms: Vec<f64>,
487    /// Time spent waiting for synchronous expert fetches, per step.
488    pub io_ms: Vec<f64>,
489    /// Records fetched from disk at exact routing time, per step.
490    pub misses: Vec<usize>,
491    pub miss_bytes: Vec<usize>,
492    step_misses: usize,
493    step_miss_bytes: usize,
494    /// Seconds in residency-set bookkeeping (acquire, add, commit) and
495    /// blocked on miss reads, per step.
496    step_set_s: f64,
497    step_read_s: f64,
498    /// Argmax per row of the MTP head folded into the last trunk step
499    /// (`step_rows` with fold), consumed by `mtp_draft`.
500    folded_mtp: Option<Vec<u32>>,
501    /// Records needed this step whose pages were still in the page
502    /// cache (no read).
503    step_warm: usize,
504    pub warm: Vec<usize>,
505    pub set_ms: Vec<f64>,
506    pub read_ms: Vec<f64>,
507    /// One-layer lookahead routing: fraction of a layer's exact experts
508    /// predicted by the approximate router run one layer earlier, and the
509    /// records the prediction fetched.
510    pub lookahead_hit: Vec<f64>,
511    pub lookahead_issued: Vec<usize>,
512    /// Every lookahead prediction with its outcome (diagnostics), and the
513    /// ones awaiting the next layer's exact routing.
514    pub la_log: Vec<LaEntry>,
515    la_pending: Vec<LaEntry>,
516    log_la: bool,
517    /// Deadline policy (--cut-weak w, default 0 = off): once the
518    /// GPU has finished a block's resident experts, a missing expert whose
519    /// weight is below w in every row and whose read has not landed is cut
520    /// from the late part instead of waited for; stronger ones are waited
521    /// for. Cut reads finish in the background and are joined at step end.
522    cut_w: f32,
523    step_cut: usize,
524    pub cut: Vec<usize>,
525    inflight: Vec<(residency::Landed, usize)>,
526    lookahead: bool,
527    /// CHERENKOV_SPIN=0: block on the shared event instead of spinning.
528    spin_wait: bool,
529    /// Diagnostics: CHERENKOV_FAKE=experts skips disk fetches (garbage
530    /// numerics, true GPU timing); CHERENKOV_SKIP=stage,... skips stages
531    /// (mixer, experts, shared, lmhead) to attribute GPU time.
532    fake_experts: bool,
533    skip: Vec<String>,
534    /// Kernel dispatches per step and GPU idle time waiting on the CPU
535    /// within a step.
536    pub dispatches: Vec<usize>,
537    /// Legacy name: CPU service wall time, which overlaps resident GPU work.
538    pub gpu_idle_ms: Vec<f64>,
539    /// Rows per trunk step and time of MTP draft passes, per step.
540    pub rows: Vec<usize>,
541    pub mtp_ms: Vec<f64>,
542    dispatch_count: std::cell::Cell<usize>,
543    /// Prefill engine scratch, allocated for the first long prompt and
544    /// released after it, and the ring the expert stream cycles through
545    /// (the pool itself keeps each layer's most-used experts).
546    pf: Option<prefill::PrefillScratch>,
547    ring: Buf,
548    /// Per prefill chunk: tokens, seconds, expert records streamed,
549    /// seconds waiting on fetches, and GPU seconds in (DeltaNet blocks,
550    /// attention blocks, expert streams, MTP).
551    prefill_reserved_bytes: usize,
552    pub prefill_stats: Vec<prefill::ChunkStats>,
553}
554
555/// Compact sequence checkpoint; no weights, expert slots, or device event counters.
556#[cfg_attr(test, derive(PartialEq))]
557pub(crate) struct PrefixState {
558    pos: usize,
559    mtp_len: usize,
560    tokens: Vec<u32>,
561    data: Vec<Vec<u8>>,
562}
563
564impl<'a> Gpu<'a> {
565    pub fn has_mtp(&self) -> bool {
566        self.mtp.is_some()
567    }
568
569    fn skips(&self, stage: &str) -> bool {
570        self.skip.iter().any(|s| s == stage)
571    }
572
573    fn record_id(&self, record_layer: usize, expert: u32) -> usize {
574        record_layer * self.p.manifest.experts.experts + expert as usize
575    }
576
577    /// Records resident in the expert pool (the residency set).
578    pub fn pool_resident(&self) -> usize {
579        self.res.resident()
580    }
581
582    pub fn pool_slots(&self) -> usize {
583        self.res.budget()
584    }
585
586    /// Dependent-FMA chain time in ms (higher = deeper clock throttle).
587    pub fn throttle_ms(&self) -> Result<f64> {
588        self.ctx.throttle_probe()
589    }
590
591    /// Trunk logits of row `r` of the last step.
592    pub fn logits_row(&self, r: usize) -> &[f32] {
593        let v = self.p.cfg.vocab_size;
594        let ptr = self.scratch.logits.contents().cast::<f32>();
595
596        unsafe { std::slice::from_raw_parts(ptr.as_ptr().add(r * v), v) }
597    }
598
599    pub fn logits(&self) -> &[f32] {
600        self.logits_row(0)
601    }
602
603    /// Batched prefill writes one final row; row stepping retains every row.
604    pub(crate) fn last_logits_row(&self) -> usize {
605        self.batch_nb.saturating_sub(1)
606    }
607
608    /// MTP logits of row `r` of the last draft pass.
609    pub fn mtp_logits_row(&self, r: usize) -> &[f32] {
610        let v = self.p.cfg.vocab_size;
611        let ptr = self.scratch.mtp_logits.contents().cast::<f32>();
612
613        unsafe { std::slice::from_raw_parts(ptr.as_ptr().add(r * v), v) }
614    }
615
616    pub fn allocated_gb(&self) -> f64 {
617        self.allocated_bytes() as f64 / BYTES_PER_GB as f64
618    }
619
620    /// What Metal will keep resident at once on this machine.
621    pub fn pool_bytes(&self) -> usize {
622        self.res.bytes()
623    }
624
625    /// Positional KV-cache bytes for one attention layer at `n` positions:
626    /// the q8 key/values (kc, vc) with their q4 scales, the QSA index key
627    /// cache and the compressed block keys. Mirrors `prefix_regions`.
628    fn kv_layer_bytes(&self, n: usize) -> usize {
629        let c = &self.p.cfg;
630        let kv_row = c.num_key_value_heads * c.head_dim;
631        let ihd = c.indexer_head_dim;
632        let ratio = c.indexer_compress_ratio.max(1);
633
634        2 * (n * kv_row + n * kv_row / 32 * 2) + n * ihd * 4 + n.div_ceil(ratio) * ihd * 4
635    }
636
637    /// Trunk attention (KV) layers; DeltaNet layers hold no positional KV.
638    fn trunk_attn_layers(&self) -> usize {
639        self.layers
640            .iter()
641            .filter(|l| matches!(l.mix, Mix::Attn(_)))
642            .count()
643    }
644
645    /// Whether the MTP head contributes an attention KV cache (0 or 1).
646    fn mtp_attn_layers(&self) -> usize {
647        usize::from(
648            self.mtp
649                .as_ref()
650                .is_some_and(|m| matches!(m.layer.mix, Mix::Attn(_))),
651        )
652    }
653
654    /// Bytes of the positional KV caches now in use by the sequence on the
655    /// GPU (trunk layers at `pos`, the MTP head at `mtp_len`).
656    pub fn kv_cache_bytes(&self) -> usize {
657        self.trunk_attn_layers() * self.kv_layer_bytes(self.pos)
658            + self.mtp_attn_layers() * self.kv_layer_bytes(self.mtp_len)
659    }
660
661    /// Total capacity of the positional KV caches at `max_t` positions.
662    pub fn kv_cache_capacity(&self) -> usize {
663        (self.trunk_attn_layers() + self.mtp_attn_layers()) * self.kv_layer_bytes(self.max_t)
664    }
665
666    /// Fraction (0..=1) of the positional KV caches in use by the sequence.
667    pub fn kv_cache_fullness(&self) -> f64 {
668        let cap = self.kv_cache_capacity();
669
670        self.kv_cache_bytes() as f64 / cap.max(1) as f64
671    }
672
673    /// Fraction (0..=1) of the context window in use by the sequence.
674    pub fn context_fullness(&self) -> f64 {
675        self.pos as f64 / self.max_t.max(1) as f64
676    }
677
678    pub fn working_set_limit_gb(&self) -> f64 {
679        use objc2_metal::MTLDevice as _;
680
681        self.ctx.device.recommendedMaxWorkingSetSize() as f64 / BYTES_PER_GB as f64
682    }
683}
684
685#[cfg(test)]
686#[path = "../../tests/unit/qwen4_exp/gpu/mod.rs"]
687mod tests;
cherenkov/qwen4_exp/gpu.rs

cherenkov/qwen4_exp/
gpu.rs
