Server Internals

Inference server internals

orangu-server (src/bin/orangu-server/) is a third binary in the same Cargo package as orangu and orangu-coordinator. Besides serving a GGUF model, it’s also the machine’s GGUF inventory tool (system/suggest/ list/show/download/delete/refresh) — stateless between runs for those seven: every invocation re-detects hardware and re-scans the models directory from scratch, so there is no cache, config-reload, or background process to reason about for them. system/suggest/show/delete stay entirely offline; download and refresh always talk to the Hub, and list does too, before printing its table, to check each Hugging Face-backed model for a newer commit (see latest_commits below) — swallowing the lookup silently rather than failing when the Hub can’t be reached. It does real tensor computation itself for serving — GGUF loading, dequantization, the transformer forward pass, sampling, and request scheduling are implemented in Rust with no dependency on any C or C++ inference library.

Module layout

The GGUF-inventory subcommands lean on library modules shared with the rest of the workspace rather than binary-local ones: orangu::gguf (the GGUF binary-format reader), orangu::model_spec (directory scan, shard grouping, and the Hugging Face repo-id/quant-tag reconstruction behind list’s MODEL column), orangu::model_download (download’s fetch logic), orangu::model_registry (the locked, versioned ~/.orangu/models inventory behind list’s LAST_USED column), orangu::os (OS detection) and orangu::hardware (CPU/GPU detection). Living in src/ alongside orangu’s and orangu-coordinator’s own shared code, rather than nested under src/bin/orangu-server/, is what let orangu-server absorb these subcommands from the now-removed orangu-gguf binary without duplicating any of this logic — orangu-server’s main.rs was already calling straight into orangu::model_spec::resolve_or_fetch_model for its own positional model argument, so list/show/download calling the same modules directly was additive, not a rewrite.

GGUF parsing (orangu::gguf)

GgufFile::read implements the header, metadata key-value, and tensor-info sections of the GGUF specification directly against a BufReader, without ever reading the tensor-data section itself — a Reader<R> wrapper tracks bytes_read as it goes, so GgufFile::data_offset (where tensor data would begin, aligned up to general.alignment, default 32) is computed for free without seeking into it. This is what keeps list/show fast against multi-gigabyte model files: parsing a file’s full metadata and tensor-info table costs only a few KB of reads regardless of the file’s total size. engine::loader (above) is a separate, mmap-based reader over the same format, built for loading tensor data rather than just metadata.

Only little-endian GGUF is read — the spec itself notes there is currently no reliable way to detect a big-endian file, and none exist in practice. GGUFv1 (32-bit tensor/metadata counts, long deprecated upstream) is rejected with a clear error rather than silently misread.

Two circuit breakers (MAX_STRING_BYTES = 100 MiB, MAX_ARRAY_ELEMENTS = 200M) guard string and array length prefixes: a corrupt or truncated download could otherwise claim an enormous length and force a huge allocation attempt before a single byte of it is verified to exist in the file.

GgufValue::display(preview_limit) renders a value for show; arrays longer than preview_limit print a truncated preview (... (N more)) rather than every element, since metadata arrays like tokenizer.ggml.tokens routinely hold well over 100,000 entries — --full passes usize::MAX to disable this.

ggml_type_name maps the ggml_type enum (ids 0–41, per ggml.h) to its canonical name; ids the format has since retired (e.g. Q4_0_4_4, whose numeric slot is never reused) print as reserved(N), and anything beyond the table (a type added after this was written) as unknown(N).

Quantization: element counts, not tensor counts (type_element_totals)

GgufFile::type_element_totals sums each tensor’s element count (dims.iter().product()) by ggml_type, rather than counting tensors. A model has far more small F32 bias/norm tensors than large weight matrices, but those matrices hold nearly all the parameters — a per-tensor-count majority would misreport a heavily quantized model as F32. This is a coarser signal than the true filename-derived quant tag (next section): it can’t distinguish Q4_K_S from Q4_K_M, since both use the Q4_K ggml type for most tensors, differing only in which few tensors (e.g. the output projection) get upgraded to a higher-precision type.

Shard grouping and the Hugging Face repo id (orangu::model_spec)

scan_models_dir walks the configured directory with walkdir::WalkDir::new(dir).follow_links(true). This is not optional: Hugging Face’s own hub cache — the layout an -hf/--hf-repo download produces — names every file under snapshots/<rev>/ as a symlink into blobs/. Without follow_links, entry.file_type().is_file() reports the symlink itself (never true), and every such model is silently skipped rather than listed.

Two further filters run in scan_models_dir itself, before any shard grouping, so only unique models are ever counted or listed:

group_models collapses a multi-part model’s shard files (name-00001-of-00004.gguf, …) into one ModelGroup, keyed by (parent directory, shard-suffix-stripped file stem) — so two files that merely share a name in different directories (e.g. two Hugging Face snapshot revisions of the same release) stay separate rows, while genuine shards of one model merge, with size_bytes summed and type_totals combined across every shard before picking one dominant type (a single shard’s own tensors are only part of the whole model).

shard_group_label and hf_tag_from_label deliberately mirror the reference implementation’s resolver in common/download.cpp byte-for-byte, rather than reinventing the convention:

group_models then drops that :TAG back off the MODEL label whenever QUANT already shows the same string — repeating it only widens the table’s widest column. Two quantizations of one repo consequently share a MODEL cell, so MODEL is no longer a unique key: ModelGroup::matches_label accepts the printed label and the reconstructed <repo>:<quant> form (which keeps every spelling ever printed resolving locally, instead of falling through to resolve_or_fetch_model’s download path), and a bare request takes the first matching row. The final sort_by is stable, so rows sharing a label keep their (parent directory, file stem) order and both NR and first-match resolution stay put between runs. delete prints the group’s quantization in its confirmation line for the same reason — the label alone can’t say which of two rows is about to go.

hf_repo_id_from_path recovers <user>/<model> by walking a file’s ancestor directories for one matching models--<user>--<model> (checking every ancestor, not just the immediate parent, since real files sit under snapshots/<rev>/, sometimes with a further per-quant subfolder). This directory-naming convention — folder_name = "models--" + repo_id.replace("/", "--") — is Hugging Face’s own, and is where a -hf download is documented to land (“models downloaded with -hf are now stored in the standard Hugging Face cache directory”). A file outside that layout has no repo_id to recover, so group_models falls back to the bare shard-stripped label.

resolve_show_target resolves whatever show was given, checking the fast, scan-free path first: resolve_model_path (a direct/relative/ absolute path, or a bare name under models) is tried before falling back to a full scan_models_dir + group_models for an NR or MODEL lookup — so the common case of show /path/to/file.gguf never pays the cost of scanning the whole directory. ModelGroup::representative_path (the first shard by sorted path order, which is also the one carrying full GGUF metadata under the standard shard-naming convention) is what show actually opens for a multi-shard model. resolve_or_fetch_model builds on top of resolve_show_target for the serving path’s own positional model argument: try resolving locally first, and only reach for orangu::model_download::download_model when nothing local matched — the same fallback main.rs’s prepare and select_model_interactively share.

Deleting a model (orangu::model_spec)

resolve_delete_target resolves delete’s argument to a full ModelGroup, not just resolve_show_target’s single representative path — delete_model needs every shard to remove a multi-shard model atomically, so this always scans and groups first rather than reusing resolve_show_target’s scan-free fast path for a plain file argument (that fast path only ever returns one file, with no way to tell whether it belongs to a larger group). Resolution order otherwise matches resolve_show_target: a direct/bare path first — returning that file’s whole group when group_models placed it in one, or a synthetic one-ModelGroup-of-one-path when it didn’t (an mmproj sidecar, which group_models deliberately excludes from every real group but delete should still be able to name directly) — then an NR, then a MODEL label.

delete_model removes every path in the resolved group, and, for each one that turns out to be a Hugging Face hub-cache symlink (models--<user>--<model>/snapshots/<rev>/<file>, resolved with std::fs::canonicalize before the symlink itself is unlinked), also removes its target blob under that same repo’s blobs/ — but only when blob_still_referenced finds no other symlink left under <repo>/snapshots/ still pointing at it. This matters for the same reason scan_models_dir’s own duplicate-file collapsing does: a repo’s ref can move without a file’s content changing, so the cache reuses (symlinks to) an already-downloaded blob from a second snapshot revision rather than re-fetching it — and scan_models_dir only ever lists the first, sorted-earliest occurrence of that shared content, so the second snapshot’s symlink is never part of any group delete was asked to remove. Scoping the reference check to just <repo>/snapshots/ (not the whole models directory) is both cheap and correct: blobs are already nested per-repo (models--<user>--<model>/blobs/), so cross-repo sharing can’t happen by construction — no walk of the full, potentially huge models directory is ever needed just to delete one model.

remove_empty_ancestors walks up from a path’s parent directory, removing it (and its own parent, and so on) as long as it’s empty, stopping the moment one isn’t or at models_dir itself (which is never removed, whatever’s left inside it). delete_model calls it twice per shard — once from the removed symlink’s own snapshots/<rev>/ chain, and, when a blob was also reclaimed, once more from that blob’s sibling blobs/ chain, since the two aren’t nested inside each other and either could be the one left holding the repo directory open. Together, deleting a repo’s last shard collapses the now-empty snapshots/<rev>/ and blobs/, and, once both are gone, models--<user>--<model>/ itself, rather than leaving a hollowed-out shell of empty directories behind.

main.rs’s Command::Delete arm always confirms before calling delete_model (confirm, a plain stdin Yes/No reader defaulting to No on an empty entry or closed stdin — the same fail-safe default a destructive filesystem action should have) unless --yes was passed, and resolves an omitted argument through select_model_for_deletion: the same format_list table list prints, followed by an NR prompt — the delete-time counterpart of main.rs’s own select_model_interactively (used to pick a model to serve), returning a full ModelGroup rather than just a path/label pair since that’s what delete_model needs.

Refreshing a model (refresh.rs)

refresh is delete followed by download of the same spec — the command that acts on a (Refresh) marker. It is its own module rather than another main.rs arm because three of its decisions are specific to it:

Delete first, then download. The reverse order would be safer against an interrupted transfer, but a refresh exists precisely because the repo’s files changed, so the new revision is a full second copy on disk rather than the blob-sharing symlink a re-download of the same commit produces. Downloading first would mean a 17 GiB model needing 34 GiB free to refresh. The cost — an interrupted download leaving the model missing rather than stale — is recovered by re-running refresh or download, which resumes from the .part file download_attempt left behind.

An ambiguous MODEL name is an error. model_spec::resolve_refresh_target mirrors resolve_delete_target (path, then NR, then MODEL label) with exactly one difference: where delete takes the first of several rows sharing a MODEL cell and names the quantization it picked in its confirmation line, refresh refuses and lists the quantizations on disk. It has to — refresh deletes what it then downloads, so a first-match would refresh the wrong quantization and leave the one the user meant untouched, which no confirmation line can undo. The error only suggests the <repo>:<quant> spelling when naming a quantization would actually disambiguate (two rows of the same repo and the same quant, two snapshots deep, can only be told apart by NR). A path that resolves to no group at all is a companion mmproj sidecar — delete synthesizes a one-file group for it, but download only ever fetches one alongside its base model, so refresh points at that model instead.

The download spec comes off the filename, not QUANT. ModelGroup::download_spec rebuilds <user>/<model>[:tag] from the representative path’s own hf_tag_from_label tag rather than from quantization: the QUANT column falls back to the dominant ggml type for a file whose name carries no tag, and that type names no file in the repo, so select_files_to_download would reject it as an unknown quant instead of re-fetching the model that is actually on disk. A group with no hf_repo at all has no spec, and run bails on it before deleting anything — a hand-copied .gguf has no repo to come back from.

--all uses the same proof once. refresh_all scans and groups the models, calls check_for_updates once for the distinct repositories, and keeps only groups for which plan returns Refresh. Current models, hand-copied files, and groups whose repositories could not be reached are never passed to deletion. Each selected group then goes through the same refresh_group delete-and-download path as a single explicit target; the first failure stops the run so a lost connection does not delete further models.

With no argument, select_model_to_refresh prints the same format_groups table list does, with Dimming::UpToDate and the same check_for_updates lookup, so the un-greyed rows are exactly the (Refresh) ones — then prompts for an NR, the counterpart of select_model_for_deletion. An unreachable Hub greys nothing (no row is known to be behind), which reads correctly: with no update information, no row is a better pick than another.

Downloading from Hugging Face (orangu::model_download)

download_model implements orangu-server download <user>/<model>[:quant] by directly mirroring the reference implementation’s common/download.cpp and common/hf-cache.cpp — read from that source rather than reimplemented from a guess at the Hugging Face API, since the whole point is producing a cache other GGUF tools recognize as already downloaded.

Resolving the commit. resolve_commit calls GET /api/models/<repo>/refs, which returns {"branches": [{"name", "targetCommit"}, ...]}; the branch named main wins, falling back to the first one listed. A repo that doesn’t exist can return 401 rather than 404 when unauthenticated (Hugging Face doesn’t distinguish “doesn’t exist” from “exists but is private” for a caller without access) — resolve_commit reports this as “repository not found … if it’s private or gated, set HF_TOKEN” when no token was supplied, or “authentication failed … check HF_TOKEN” when one was (a 401 with a token in hand means the token itself was rejected, not that the repo is missing).

Listing files. list_repo_files calls GET /api/models/<repo>/tree/<commit>?recursive=true, returning every file with its path, and either a top-level oid (the git blob sha1, for small files) or an lfs.oid (the LFS object’s sha256, for anything large enough to be stored as LFS — every real GGUF file). RepoFile::oid takes whichever is present; it doubles as the blob’s filename in the cache, so two snapshots referencing byte-identical content share one on-disk copy exactly like the real Hugging Face cache does.

Choosing what to download. select_files_to_download mirrors find_best_model + get_split_files:

Choosing a multimodal projector, if any. After the primary model file is picked, find_best_mmproj (calling the generic find_best_sibling with keyword = "mmproj") directly mirrors the reference find_best_sibling/ find_best_mmproj: among every .gguf path containing mmproj, it prefers the one sharing the deepest directory prefix with the primary file’s own path (rejecting any candidate whose directory list isn’t a prefix of the model’s), then — among ties at that depth — the one whose quantization bit depth (extract_quant_bits, reading the first run of digits in the filename’s trailing tag, e.g. Q4_K_M -> 4, BF16/F16 -> 16, F32 -> 32) is numerically closest to the primary file’s own. This is the same file orangu-server’s own -hf auto-fetches the first time a vision-capable model is launched with an image-related flag (verified against a real repo, unsloth/Qwen3.6-35B-A3B-GGUF, which offers three top-level mmproj variants — BF16/F16/F32 — alongside a Q4_K_M primary; both this code and a live orangu-server -hf ...:Q4_K_M --image-min-tokens 1024 run independently picked mmproj-BF16.gguf), so fetching it up front here means LLAMA_CACHE=<models> already has it ready offline. If found, it’s appended to the file list download_model fetches, alongside whatever shards the primary model itself has.

Planning the model before fetching it (RemoteModel). resolve_commit and list_repo_files answer everything engine::plan needs except the tensor tables, and those are not behind the download either — a GGUF file puts its header at the front, so the tables are the first few hundred kilobytes of each shard. resolve_remote_model performs exactly the two Hub calls above and stops there, returning a RemoteModel that names the commit and the selected shards with nothing fetched; RemoteModel::headers then streams each shard from /<repo>/resolve/<commit>/<path> into GgufFile::read_from and yields the parsed header.

Two properties make this cost the header rather than the file. The GGUF parser is strictly sequential and stops at the end of the tensor-info table — pinned by read_from_stops_at_the_end_of_the_tensor_table, which asserts the reader’s final position rather than merely that the parse succeeded, because a parser that read to EOF would return the same GgufFile and only the unread bytes reveal the difference. And dropping a reqwest response cancels the rest of its transfer, so returning from RemoteModel::header closes the connection. Planning a 1.3 TiB repo therefore transfers a few megabytes.

headers is lazy, so a consumer that fails on shard 1 never pays for the remaining ten, and it yields Result rather than swallowing failures: a header that will not parse is precisely the thing worth knowing before a multi-hour download. RemoteModel deliberately excludes the mmproj sidecar that download_model fetches — it is a separate CLIP-architecture model, and its tensors would be counted as this model’s weights by anything reading the tables.

main.rs’s plan_before_download is the only caller. It hands RemoteModel::headers to engine::plan::analyze_shards — the same classifier plan uses on local files, which cannot tell where the tables came from — prints the report, and consults Plan::dense_fits_in to decide whether to confirm. It confirms on the dense part alone, never on the experts: experts stream, so a model that overflows on them is slow rather than broken, and prompting there would cry wolf on the workload the expert path exists for. Every failure in this whole path reports one line and returns “go ahead”, since download_model is about to attempt the same repo and is the better place for the real error to surface. It sits outside download_model_reporting on purpose: that function also serves the web console’s model manager and model-spec resolution ahead of serving, neither of which has a terminal to confirm on.

Fetching bytes, concurrently. download_model first walks selected sequentially just to decide what needs fetching at all — a blob already present on disk with a matching size is skipped entirely rather than re-verified byte-for-byte (cheap and good enough; matches the practicality bar the rest of this tool holds to elsewhere, e.g. the element-count quantization guess), printed immediately with an [index/total] suffix. Everything left becomes a DownloadTask (label, URL, blob path, size, and that same (index, total) position), and download_all hands the whole batch to rayon’s par_iter().try_for_each — bounded by rayon’s global thread pool rather than one OS thread per file, so a model with dozens of shards doesn’t open dozens of simultaneous connections. This means a sharded model’s shards, and a bundled mmproj sidecar, download at the same time instead of one at a time; download_model only does the symlink-placement pass (link_or_copy, below) after every download has finished.

Each parallel task’s own download_with_resume streams its response body to a <blob>.part file, resuming from wherever that file left off via an HTTP Range request if one already exists from an interrupted attempt (falling back to a full restart if the server doesn’t honor it, signaled by a 200 instead of the expected 206). Progress is a plain percentage against the tree API’s own reported file size — not the response’s Content-Length, which would only cover the remaining bytes on a resumed request. Since several tasks report progress at once, each writes into its own line of a ProgressBoard shared behind a single Mutex (one mutex around the whole board, not one per line, so a “set this line, then redraw every line” update is atomic and two threads’ redraws can’t interleave); ProgressBoard::update redraws in place with \x1b[{n}A (cursor up n lines) followed by \x1b[2K (clear line) per row, so every in-flight file’s percentage stays visible at once until all are done, at which point its line switches from Downloading to a final Downloaded <label>: 100% [index/total] — kept at 100% rather than dropped, so every line stays in the same <verb> <label>: <percent>% [index/total] shape whether still in flight or finished. If a task fails, the others still run to completion rather than being cancelled (each writes its own .part file, so a later retry only re-fetches whatever actually failed); download_all surfaces the first error once every task has finished.

Placing the file. link_or_copy computes the same relative symlink target the real Hugging Face cache uses (../ once per path component between snapshots/<commit>/ and the file, plus two more to reach the repo root, then into blobs/<oid>) rather than an absolute path, so the whole models directory stays portable if moved. Falls back to a plain copy if symlinks aren’t available at all (e.g. Windows without developer mode enabled) — mirroring hf_cache::finalize_file’s own degraded-mode fallback.

Not implemented, out of scope for a first version: --mtp companion downloads (also a find_best_sibling call upstream, with keyword = "mtp-"), preset.ini-based repos (a repo-root manifest naming one specific file to fetch regardless of tag matching), and Docker registry sources.

Checking for updates (list’s (Refresh) marker)

list doesn’t just read local disk state — it also asks the Hub whether a newer commit exists for every model it found under a Hugging Face hub-cache directory. Two pieces make this work:

main.rs‘s Command::List arm wires the two together through check_for_updates (which does the dedupe-by-repo and is shared with refresh, so both commands agree on what’s stale): group_models runs first, its groups’ distinct hf_repo ids feed latest_commits, which returns a repo -> commit map — not a “these repos are stale” set — and format_groups (the renderer format_list itself now delegates to) asks ModelGroup::is_behind to compare each row’s own local_commit against that map when deciding whether to append (Refresh) after SIZE. Comparing per row rather than per repo matters: a repo can have two ModelGroup rows cached at different commits (e.g. :Q4_K_M downloaded weeks ago, :Q8_0 downloaded today), and only the one actually behind should be marked — a HashSet of “stale repos” would incorrectly mark both just because they share a repo id. The marker sits deliberately after SIZE rather than folded into MODEL, so the shell completion scripts (above), which only ever read list’s first two whitespace-separated columns, stay unaffected by a row growing a trailing marker.

The SUPPORTED column

list prints a SUPPORTED column reading Yes (<arch>) or No (<arch>) per row — so a user sees which models this build can actually load before selecting one, rather than only discovering it can’t once it’s loaded. model_spec::format_groups renders the column (and format_list’s signature carries the support/dim parameters through), but the lib deliberately doesn’t decide what is supported: that judgement lives in orangu-server, in engine::loader::model_load_support. So main.rs’s model_support opens each group’s representative file (header only — no tensor data, the same cheap read show does), calls model_load_support, and stores the result as one model_spec::ModelSupport { architecture, supported, unsupported_quant } per group before handing the slice to format_groups. Every shard of a group is inspected, not just its representative file: a split model’s later shards carry their own tensor directory and can use a quantization shard 1 never does. An empty slice omits the column entirely, which is what format_list (lib-side tests) and any caller without the loader pass.

model_load_support is deliberately allowed to be stricter than resolve_arch_family (whose family tables are the single source of truth for the architecture string): a model whose architecture is recognised can still carry tensors this build cannot read, so a bare resolve_arch_family “yes” would promise a load that then fails partway through. It therefore also checks every tensor’s ggml_type against quant::supports_type and reports the first unreadable one, which ModelSupport::cell renders as No (llama, TQ1_0) — distinct from No (glm4moe), because only the former is fixed by fetching a different quantization of the same model. Note this is not the same question as the arch module’s own tensor expectations: gemma MoE checkpoints (gemma-4-26B-A4B, blk.{i}.ffn_gate_inp.weight present) load via arch::gemma’s routed-expert path and report Yes (gemma4).

SUPPORTED answers “can this build read the file”, which is not quite the same question as “will it run on the backend you selected”. Every GPU backend covers fewer ggml_types than engine::quant does, so a row can read Yes and still be refused at startup by engine::backend::unsupported_tensor_types — see the CUDA/OpenCL/ROCm section for the coverage each backend has. The column deliberately does not fold that in: it is rendered before a backend is chosen, and the same file that one backend refuses runs on cpu.

A No row is greyed (dim ANSI SGR), not hidden: a user can still pick it and will hit the same clear “not yet supported” error prepare gives for any other unsupported model — the greying just deprioritizes it visually.

Which rows are greyed is format_groups’ dim parameter, a model_spec::Dimming rather than the boolean it started as, because refresh wants the same table deprioritizing a different set of rows: Dimming::Unsupported greys what this build can’t load (list, show, delete, the serve-time picker), Dimming::UpToDate greys what isn’t behind its repo (refresh), and Dimming::Off emits no escapes at all. Every call site passes its mode through main.rs’s dimming helper, which returns Dimming::Off unless std::io::stdout().is_terminal(). Piped or redirected output — including what the shell completion scripts parse with awk '{print $1; print $2}' — therefore stays escape-free, so an ANSI prefix can never corrupt the NR/MODEL columns those scripts read. One renderer serves every table this binary prints, so they all carry the column consistently.

OS detection (orangu::os)

orangu::os::detect gathers the OS section orangu::hardware:: format_report prints first, and orangu::os::format_section formats it — the report’s other two sections stay in orangu::hardware, so neither module has to know how the other’s fields are gathered. OsInfo is a flat struct of Option fields precisely because platform coverage is uneven: format_section skips any field that came back None, which is what lets one formatter serve three platforms without a cfg in it.

Nothing here runs a subprocess. Every field comes from a Rust API:

is_redundant_long_version decides whether the Full name line is worth printing. sysinfo’s long version is Linux (Fedora Linux 44) on Linux — Name and Version rearranged, plus the kernel name and punctuation — but MacOS 15.1 Sequoia and Windows 11 Pro on the other two, where the codename and edition appear nowhere else. Comparing word by word (ignoring punctuation, and treating linux as already-known) drops the first and keeps the other two, without a cfg deciding it per platform.

CPU/GPU detection (orangu::hardware)

CPU statistics (brand, vendor, architecture, physical/logical core counts, peak frequency, total/available RAM) come from sysinfo, used with its system and component features only (no disk/network/user) to keep the dependency footprint small — the same dependency orangu::os uses above.

Power and thermals (detect_power). component is the second feature, and it is there for temperatures: sysinfo::Components reads sensors on every platform this targets, which is worth a dependency rather than writing three hwmon/SMC/WMI readers by hand. Sensors reporting None, or exactly 0.0, are dropped — the latter is a channel that is not wired up rather than a component at freezing, and an integrated GPU’s memory channel does report it. What survives is sorted warmest first, because the only sensor anybody acts on is the one nearest its limit.

The power source is not from sysinfo, which has no battery or AC-line API at any version, so each platform is read natively:

classify_power_source is split out from the Linux walk and compiled into the tests on every platform, because it holds the one real trap: a laptop plugged in with a full battery reports its battery status as Not charging, not Charging. Deciding from the battery alone therefore means reading a double negative that is easy to invert, and inverting it would tell every desk-bound laptop it was running down a battery. Asking the AC line first removes the question — an adapter reporting itself online is authoritative whatever the battery says. A machine with neither an online adapter nor a draining battery is a desktop, a server, or a container, and all three answer Mains: every caller is really asking “is my power about to run out”.

power_advisories is the only advisory source left, and everything it returns is a condition: it holds on every platform, and neither of its two findings has a command as an answer — one is answered by a cable and the other by airflow. Machine settings used to sit beside them in a performance_advisories function that no longer exists, because a setting has a value on every start and an advisory can only speak when that value is wrong. The CPU governor became cpu_governor, a banner row; AMD’s power_dpm_force_performance_level became a documented check, since a machine with a discrete card and an integrated one printed a line per card, on every start, to say the same thing about each. The thermal one fires only against a threshold the platform itself declared and only within THERMAL_ADVISORY_FRACTION of it, which means a hot sensor that declares no threshold is reported in the POWER section and never warned about. That is deliberate: silicon runs hot under load, Tctl on AMD parts is a control offset rather than a junction temperature, and a fixed limit invented here would fire on machines that are working perfectly.

GPU detection has no single cross-platform API, so detect_gpus layers several best-effort, independent sources and concatenates whatever each finds — a card no source recognizes simply doesn’t appear, rather than the whole command failing:

  1. NVIDIA (detect_nvidia_gpus, Linux and Windows): shells out to nvidia-smi --query-gpu=... --format=csv,noheader,nounits, the one interface guaranteed to exist wherever an NVIDIA driver is installed. A missing binary or non-zero exit returns an empty list, not an error — “no NVIDIA GPU” is the expected common case. memory_kind is always MemoryKind::Dedicated — no consumer NVIDIA GPU is anything else.
  2. AMD/Intel/other, Linux only (detect_linux_sysfs_gpus): enumerates /sys/class/drm/card*/device, the kernel interface every Linux GPU driver exposes. NVIDIA vendor ids (0x10de) are skipped here — already reported by nvidia-smi above, and mem_info_vram_total is an amdgpu-specific sysfs attribute this path can’t get for NVIDIA anyway. VRAM total/used come from mem_info_vram_total/mem_info_vram_used when present (AMD only; Intel iGPUs report no separate VRAM, being shared system memory). The device’s marketing name is looked up in the system’s pci.ids database (load_pci_ids, checking /usr/share/hwdata/pci.ids first — the hwdata package’s path on Fedora/RHEL — then the pciutils paths used elsewhere), the same file lspci itself reads; if it isn’t installed, the raw vendor:device PCI ids are shown instead of a name, rather than failing.
  3. macOS (detect_macos_gpus): system_profiler SPDisplaysDataType -json, parsed with serde_json (already a workspace dependency).
  4. Windows (detect_windows_gpus): PowerShell’s Win32_VideoController WMI class via Get-CimInstance | ConvertTo-Json. A single result comes back as a bare JSON object rather than a one-element array, which the parser normalizes explicitly. AdapterRAM is a well-known 32-bit field that misreports (often as 0 or wrapped) for cards with more than ~4 GiB of VRAM; it’s still the best zero-dependency source available on Windows, so a 0 reading is treated as “unknown” rather than shown literally.

Dedicated vs. shared memory (MemoryKind)

Every GpuInfo carries a memory_kind: MemoryKind (Dedicated / Shared / Unknown), derived by a different signal per platform — there is no single cross-platform API for this either:

MemoryKind::Unknown is only ever constructed on macOS/Windows, whose detection functions are cfg’d out on other build targets — hence the variant carries a blanket #[allow(dead_code)] rather than one scoped per target.

Shared memory’s total is system RAM, not the raw query result

detect_gpus(total_memory_bytes) takes the system’s total RAM — CpuInfo::total_memory_bytes, computed once by the caller so this doesn’t pay for a second sysinfo query — and, after concatenating every platform’s GPUs, runs apply_shared_memory_total over the result: any GpuInfo with memory_kind == MemoryKind::Shared has its vram_total_bytes overwritten with total_memory_bytes, unconditionally.

This matters because a shared GPU’s own reported figure (where one exists at all) drastically understates what it can actually use: amdgpu reports an APU’s tiny BIOS-reserved carve-out via mem_info_vram_total (as little as a few hundred MiB — 512 MiB on the Renoir APU this was verified against), and Intel/Windows sources often report nothing at all. System RAM is the real ceiling on how much such a GPU can draw on, so it’s the only figure worth showing as its total; vram_used_bytes is left untouched (whatever the platform reported, or None), since “how much of the shared pool is currently claimed as graphics memory” is a real and distinct figure from the override, unlike the total.

Hardware-based model-size suggestion (suggest.rs)

main.rs’s Command::Suggest arm calls the same orangu::os::detect and orangu::hardware::detect_cpu/detect_gpus trio Command::System does, then passes the result to suggest::format_suggestion, which appends two size-suggestion tables after orangu::hardware::format_report’s own OS/CPU/GPU listing (via the shared push_suggestion_block helper). There is no separate detection path — suggest is purely a second interpretation of the same inventory system already knows how to gather (and the same report printed at the top of every attached orangu-server startup — see the Inference server chapter’s Quick start section).

The memory-estimation formula. estimate_total_vram_bytes mirrors Sam McLeod’s GGUF VRAM Estimator’s own calculateMemoryBreakdown function (read directly from its published vram-calculator.min.js, not guessed) and the general shape of erans/selfhostllm’s calculator:

Since suggest runs before any model is chosen, there’s no real GGUF file to read hidden_size/layers from. estimate_hidden_dims instead estimates both from the parameter count alone. The standard transformer parameter-count approximation (params ≈ 12 × layers × hidden_size²) is one equation with two unknowns, so the split is underdetermined; it’s resolved by putting everything into the hidden size (hidden_size = sqrt(params / 12)), which makes layers work out to exactly 1 by construction. The KV-cache estimate built on it therefore scales as context × √params — which tracks modern GQA-era models well (their per-layer KV width shrinks as depth grows, so total KV grows sublinearly in parameters), and matches the fallback smcleod’s own calculator uses when it has no real GGUF metadata to read either.

DEFAULT_BITS_PER_WEIGHT (4.83, Q4_K_M) and KV_CACHE_BITS (8, Q8_0) match this project’s own established defaults (orangu::model_download’s DEFAULT_TAG_PREFERENCE, and the same Q8_0 KV-cache quantization engine::kv_cache itself stores) rather than assuming full FP16 throughout.

A table, not a single guess. Actual context usage varies far too much to guess well from hardware alone, and bits-per-weight depends on which quantization tag you end up downloading — so instead of picking one of each, push_suggestion_block prints a row per context length in CONTEXT_LADDER (1K up to a generous long-context ceiling, 1048576) and a column per quantization in QUANT_LADDER (Q2_K at 3.00 bits/weight, Q4_K_M at DEFAULT_BITS_PER_WEIGHT, and Q8_0 at 8.5 — all three bits-per-weight figures read from smcleod’s own table, the same source as the formula itself). Each cell is independently computed by suggest_param_count, so the suggested size correctly shrinks along a row as quantization gets heavier, and down a column as context grows.

Picking a size. suggest_param_count walks PARAM_LADDER_BILLIONS — a curated list of common open-weight parameter counts, largest first — and returns the first whose estimate_total_vram_bytes result (at that cell’s context length and bits-per-weight) fits within the budget, or None if even the smallest rung (1B) doesn’t (rendered as -).

Where the estimate ends. format_suggestion closes with NEXT_STEP, a fixed paragraph naming download and plan as the two commands that do not have to estimate. It is there because every figure above it is derived from a parameter count and the approximation described above — at this point no model has been chosen, so there is no file to read, and engine::plan (which reads real tensor tables) has nothing to read from. suggest therefore cannot be made exact; what it can do is hand the user the command that is. download plans the repo’s real tables before fetching it and plan does the same for a local model, so the size class suggest produces gets checked against reality before any bandwidth is spent on it.

Two budgets, (up to) two tables. format_suggestion computes two separate budgets and prints a labeled push_suggestion_block for each, "Suggested model size (Dedicated)" and "Suggested model size (Total)". Both read each eligible GPU’s own vram_total_bytes — deliberately not reduced by vram_used_bytes, since suggest estimates the hardware’s own capability (this file’s module doc — “likely to run comfortably on this machine”, picked before any model is chosen), not how much happens to be free at the exact moment it runs; whatever else is transiently using VRAM (a compositor, a browser, an already-running orangu-server) shouldn’t shrink a hardware-based estimate:

Unknown-kind GPUs: a Windows-specific path. On Linux/macOS, is_dedicated_for_budget/is_total_budget_eligible only ever see Dedicated/Shared GPUs — MemoryKind is already reliably known there (see above), so both functions have a plain, cfg-free body for those targets. Windows is different: windows_memory_kind classifies any AMD adapter Unknown, discrete Radeon and integrated APU alike, since that distinction only exists in DXGI’s DXGI_ADAPTER_DESC — unreachable from the WMI query detect_windows_gpus uses. Rather than counting every Unknown GPU (overcounts an APU’s tiny carve-out as if it were a hard VRAM ceiling) or none (undercounts a real discrete Radeon card), the #[cfg(target_os = "windows")] variants of both functions trust an Unknown GPU’s own vram_total_bytes only above WINDOWS_UNKNOWN_DEDICATED_THRESHOLD_BYTES (1 GiB — comfortably above a typical integrated carve-out, comfortably below any real discrete card). Below the threshold it’s treated like a Shared GPU: excluded from both budgets, since its real ceiling is system RAM, which total_budget_bytes already considers on its own.

Shell completions (shell.rs)

Mirrors orangu’s own -s/--shell-completions (src/bin/orangu/ shell.rs): hand-written bash/zsh/fish scripts embedded as &str constants, selected by inspecting $SHELL, rather than clap-generated completions. The selection — and the error naming the supported shells when $SHELL is none of them — is orangu::shell_completions in the library, shared by all five binaries (orangu-coordinator, orangu-bench and orangu-gguf each carry a shell.rs of the same shape); each binary’s COMPLETION_SCRIPTS hands its three scripts to it. The same module’s unoffered is what each binary’s every_flag_is_offered_by_every_completion_script test runs, so a flag added to clap and forgotten in the scripts fails the build rather than Tab. The positional model argument, and show’s, delete’s and refresh’s own arguments, complete the same way orangu’s own scripts complete session UUIDs — the shell function shells back out to orangu-server list itself (2>/dev/null, so a missing config yields no candidates rather than an error) and reads its first two columns with awk. This keeps the completion logic entirely in the shell script, depending on nothing but orangu-server itself being on $PATH — no dynamic-completion protocol or extra binary flag is needed. The bash and fish scripts also list the subcommand names as literal completion candidates alongside the dynamic model list at the first argument position; the zsh script achieves the same with _alternative combining a _values list (subcommand names) and a compadd-based function (model candidates) for that position.

An earlier version of this explored clap_complete’s unstable-dynamic feature for this instead; it was backed out in favor of the approach above once orangu’s own precedent was found, since introducing a genuinely unstable (semver-exempt) dependency wasn’t warranted when a small, self-contained shell script does the same job with zero new dependencies.

prune’s own argument completes differently from model/show/delete/refresh above: directly against ~/.orangu/server/sessions/* (each entry a UUID directory) plus the literal all, with no process invocation at all — this time genuinely the same trick orangu’s own -r/--resume completion uses (_orangu_sessions/__orangu_sessions in src/bin/orangu/shell.rs), not just the same general shape. Shelling out to orangu-server prune itself the way model completion shells out to list isn’t an option here: prune with no argument prints its table and then reads a selection from stdin, so piping its output into a completion function would risk the completion hanging on that prompt — list never reads stdin, which is exactly why it’s safe to use as a completion source and prune isn’t.

The PowerShell script registers a native argument completer (Register-ArgumentCompleter -Native) whose script block does what the three Unix scripts do — the model list from orangu-server list, the session UUIDs from the sessions directory, fixed value lists for --host, --sort and --device-split, and nothing at all for a path-taking flag, since a native completer that returns nothing hands the word to PowerShell’s own file completion. It is kept to ASCII on purpose: Windows PowerShell decodes a native command’s output in the console’s code page, and orangu-server -s | Out-String | Invoke-Expression would otherwise garble every dash in a tooltip.

GGUF loading and dequantization

engine::loader memory-maps the file and reads hyperparameters using the same <arch>.* key names every GGUF loader reads (confirmed directly against the reference llama-arch.cpp’s LLM_KV_* table). Weight tensors are not eagerly dequantized into RAM — each row is read straight from the mmap and dequantized on demand, so even a large model’s memory footprint stays close to its file size.

engine::quant’s dequantization struct layouts and algorithms are taken directly from ggml’s own ggml-common.h/ggml-quants.c (dequantize_row_*), not reimplemented from a description, so the CPU path is bit-for-bit compatible with what other GGUF loaders read. Supported types: the floats (F32, F16, BF16), the legacy quants (Q4_0, Q4_1, Q5_0, Q5_1, Q8_0), the whole K-quant family (Q2_K through Q6_K), and the IQ* codebook quants (IQ1_S, IQ1_M, IQ1_XS, IQ1_XXS, IQ1_XXXS, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S, IQ4_NL, IQ4_XS) — any other ggml_type fails to load with a clear “not yet supported” error rather than misreading it.

IQ4_NL is worth calling out because it turns up in files whose name promises a pure K-quant. It is the one IQ* type that blocks at 32 elements rather than 256, and a K-quant needs 256 to divide the row, so upstream’s llama_tensor_get_type substitutes it per tensor wherever a row is too narrow — every 896-wide row of a Qwen2.5-0.5B, for instance. Any code that reads “IQ*” as “QK_K-blocked” gets its stride wrong for this one type alone; quant::block_layout, vecdot::supports, and every backend’s block-size table each place it with the 32-element quants deliberately.

The six repacked layouts — Q4_0_4_4/_4_8/_8_8 (ggml ids 31-33) and IQ4_NL_4_4/_4_8/_8_8 (36-38) — are handled differently from every other type, in engine::loader rather than engine::quant’s per-row dequantizers. They are ARM-SIMD pre-repacked Q4_0/IQ4_NL: ggml retired the ids and never shipped a to_float for any of them (only gemv/gemm kernels), which is why upstream refuses such a file. The packing is lossless though — repack_{q4_0,iq4_nl}_to_*_bl plus the make_block_* functions interleave 4 or 8 rows into shared records — so quant::deinterleave_repack inverts it and LoadedModel::open rewrites every such tensor to its plain base type before anything reads one. Past the loader these ids do not exist, so the CPU fused kernel and all four GPU backends serve them with no new code.

IQ1_XS/IQ1_XXS/IQ1_XXXS (ids 64-66) are the other three that sit outside ggml’s own enum, and the reason gguf::GGML_TYPE_NAMES now runs past its end. They are IQ1_S with a narrower codebook index — 10, 9 and 8 bits instead of 11, into 1024-, 512- and 256-point subsets of the very same 2048-point lattice — reaching 1.4375, 1.3125 and 1.1875 bits per weight, and a “dynamic” 1-bit release of a trillion-parameter mixture-of-experts model stores its expert stacks as the last of them. Because only the index changes, dequantize_iq1_xs/_xxs/_xxxs are the IQ1_S dequantizer with a different index reassembly, ending in the same push_iq1_grid call, and iq1_narrow_sub_scale is shared by the two whose scale and sign live in a nibble array rather than in the index byte itself. Nothing downstream needed a change: the expert stacks these appear in are read on the CPU whatever the backend, so SUPPORTED_TYPES for the GPU backends stays as it was.

The ids matter as much as the layouts. 42-63 are left free for ggml to grow into, which is what makes 64+ safe for a quantizer to claim: a build that does not know these types rejects the file instead of reading it at the wrong stride. GGML_TYPE_NAMES therefore spells the gap out as None, printed as reserved(N), so list distinguishes “an id ggml has not used yet” from unknown(N), “a type newer than this build”.

testdata/ggml-dequant-reference.bin covers 20 types rather than 16 for this: the three new ones plus IQ1_S, whose presence is what says the harness covers this family rather than agreeing vacuously. Regenerating it leaves the 16 that were already there byte-identical, but needs a ggml build that implements the narrow types — that fixture is local rather than checked in, so on a checkout without it the cross-check skips and dequantizing_every_type_is_unchanged’s checksums are what hold the three dequantizers still.

quant::repack_layout is the whole specification, as (base type, rows, run, xor). Two entries in it do not follow from the type names and would corrupt every weight if guessed:

IQ4_NL_4_8 is the one layout with no executable upstream definition: make_block_iq4_nlx4’s 8-byte branch is commented out and marked “this branch seems wrong”. Its entry is the straightforward generalization — the x4 index math with 8-byte runs — and no released file is known to carry the id.

Two properties follow from the row interleaving. The conversion is eager, not lazy: a row’s blocks are strided across its 4- or 8-row group, so there is no row-shaped slice of the mapped file to defer — TensorLocation holds an owned buffer for these tensors instead of an Mmap slice (both behind the TensorBytes trait object, so readers are unchanged). And quant::dequantize refuses these ids outright rather than decoding a “row”, since bytes for one row are not a thing that exists in this layout.

Correctness is pinned by testdata/ggml-dequant-reference.bin: random blocks of every quantized type paired with the f32s ggml’s own ggml_get_type_traits(t)->to_float produced from them, compared bit-for-bit. Regenerate it with testdata/ggml-dequant-reference.c (see quant.rs’s read_ggml_reference for the command) when adding a type — appending to that file’s type list leaves every existing entry byte-identical.

The fused int8 CPU path (engine::vecdot)

The obvious way to multiply a quantized weight matrix on the CPU is to dequantize each row to f32 and take an f32 dot product. engine::vecdot does neither: it quantizes the activations to int8 once (32 elements per shared f32 scale), then dots them against the weight bytes while those are still quantized, using integer SIMD. That removes the dequantize, removes the per-row allocation, and replaces scalar f32 multiplies with 16- or 32-wide int8 ones (NEON vmull_s8/sdot on aarch64; AVX-512 VNNI, AVX2 or SSE4.1 on x86-64, all chosen by runtime feature detection).

Prefill on i8mm cores takes a different shape. A dot kernel — one row against one token, reduced to a scalar every 32 elements — is the right unit for decode and the wrong one for a prompt or a picture, where the same weights meet hundreds of tokens. On an ARMv8.6 core with i8mm, vecdot::dot_k_rows runs the K-quant GEMM as smmla tiles instead: four weight rows against eight tokens, each instruction a 2×8·8×2 int8 product accumulated into a 2×2 of i32 lanes, the per-sub-block scale folded in with one vector mla, and nothing reduced horizontally until the tile is stored. Its activations (ActQ8Mm) are the q8_K layout re-ordered so each smmla operand is one aligned load, and its rows are re-paired the same way once per task (PairedRows). It is the pair kernel’s arithmetic to the bit — the integer sums are identical and the f32 steps are the same operations in the same order — which the test i8mm_quad_kernel_is_bit_identical_to_the_pair_kernel asserts as equality rather than tolerance. Measured on a Cortex-A720: 2.3× per core over the sdot kernel at a Qwen-Image linear’s shape (3072×3072, 256 tokens). ORANGU_EXPERT_I8MM=0 turns it off for an A/B on one binary; ORANGU_EXPERT_K_ROWS sets the rows per rayon task (default 16, the best of 4–64 on that core: more rows reuse each activation tile from L1 across more smmla tiles, until the unpacked rows crowd it out).

A picture’s attention scores run on smmla too. Attention is the part of a denoising step that grows with the square of the picture: profiled with orangu-bench --image … --flamegraph, a Qwen-Image 2.1 step spends 14% of its time there at 512 × 512 and 39% at 1024 × 1024, almost all of it in the f32 tile (gemm_f32_rows) — while the linears around it run on the int8 kernel. So a step’s q · k is now int8 (image::transformer::step_attention, vecdot::i8_scores_4rows): each head’s keys lose their mean over the sequence (every score of a query moves by the same amount, which the softmax does not see) and each query and key row is quantized with its own scale; the softmax and the value product stay f32. A 1024 × 1024 Qwen-Image 2.1 step went 122 → 104 s (attention 39% → 29%), a 512 × 512 one 21.6 → 20.7 s; Qwen-Image 2512, whose joint attention takes the same path, 139.7 → 119.8 s at 1024 × 1024 and 24.7 → 23.7 s at 512 × 512. The picture is the same picture: the same seed drew the same composition and lettering, differing in fine detail (30.6 dB PSNR against the f32 one; 29.1 dB for 2512) — the size of difference diffusers itself documents between two equally valid samples. ORANGU_IMAGE_ATTENTION=f32 keeps the exact path, for an A/B on one binary. The value product itself runs on bfmmla where the CPU has bf16 (probabilities and values rounded to bf16, summed in f32, sixteen multiply-adds an instruction; 1024² step 72.4 → 63.0 s), else on rten-gemm’s f32 kernel (a packed 4 × 16 tile, ~150 G MAC/s at its shape against ~100 for gemm_f32_rows); ORANGU_IMAGE_PV=rten|orangu picks the f32 paths. The prompt prefix, the vision tower and the VAE keep f32 attention; they run once a picture.

Every one of these kernels exists on every architecture. The smmla tile product has a portable twin (vecdot::dot_k_rows_portable, the same arithmetic in i32 and f32 scalars, bit-identical to the pair kernel and to smmla, which the same test asserts on every machine), and it is what dot_k_rows_tiles is off aarch64 — a complete definition, not a stub, though the callers there prefer the pair kernels’ AVX2/VNNI forms through have_i8mm. The float prefill GEMM (gemm_f32_rows, dot_f32_slices, widen_float_row), its two-dimensional split in CpuBackend::matmul_float_into, and the blocked joint attention of the picture transformer (joint_attention_blocked) are architecture- independent, with a NEON tile on aarch64 and a portable eight-lane tile elsewhere (f32_tile_portable, plain arrays LLVM vectorizes under the x86_64 build’s AVX2 baseline); the tests that check them against the per-token dot and the per-query attention run everywhere.

The serial half of a prefill matmul. After the kernel above, a per-thread profile of the 12-thread microbenchmark showed every worker busy about half the wall time: the phases that run on one thread while eleven sleep were 36% of it — finish_into’s transpose of the [out_dim][n_tokens] accumulator into the token-major output a quarter on its own, the activation quantizer a tenth. On the i8mm path both are gone: ActQ8Mm::quantize is one rayon task per tile, and matmul_k_gemm_into gives each task a private L1-resident block for its 16 rows and has it write, per token, one contiguous 64-byte run straight into the output (SharedOut, a raw pointer whose column ranges no two tasks share) — the transpose done in parallel and for free. Measured: a 256-token 3072×3072 matmul 205 → 346 G MAC/s, a Qwen-Image step 18.3 → 14.4 s, 7B prefill 35 → 43 tok/s, decode unchanged.

Big and little cores. hardware::big_cores reads Linux’s cpu_capacity per core and the [cpu] line reports the mix. Pinning the pool to the big cluster (ORANGU_EXPERT_BIG_CORES=1) was measured both ways on the development board — it matched the microbenchmark, lost prefill and pictures in the server (the little cores fill the slack of a forward pass’s many short parallel regions) and won decode by 18% — so it is opt-in, and a decode-specific pool is the open task (doc/PERF-IMAGE.md).

The float weight types get a tiled kernel too. F32/F16/BF16 weights have no int8 form, and their per-token dot (dot_row_f32) was load-bound. vecdot::gemm_f32_rows is a 4-row × 6-token NEON tile — the largest that does not spill — and matmul_float_into splits the work over token blocks × row groups so a VAE convolution (96 rows against 65,536 pixels) and a bottom-of-the-decoder one (384 rows against 1,024) both get a few tasks per worker; F32 weights are read in place. It halved a Qwen-Image VAE decode and raised gemma-4-E2B’s prefill by a quarter, whose per_layer_model_proj is the one float matmul in that model.

What the machine can use, not what it advertises. have_dotprod, have_i8mm and have_vnni each require the CPU to report the feature and the instruction to reproduce a scalar reference on known inputs before the kernel is dispatched — the assembly was written on one machine and runs on many. The startup [cpu] line names the level that passed (vecdot::int8_kernel_label: i8mm (smmla), dotprod (sdot), NEON; AVX-512 VNNI, AVX2, SSE4.1, scalar), so a benchmark’s header says which kernel produced its number.

Five types have a fused kernel — Q8_0, Q5_0, IQ4_NL, Q4_K, Q6_K. Anything else returns false from vecdot::supports and the caller keeps the ordinary dequantize path, so the module is strictly additive.

supports takes the row length as well as the type, and that second check is load-bearing rather than defensive. A GGUF row is only guaranteed to be a whole number of blocks, and the block sizes differ: Q8_0, Q5_0 and IQ4_NL need 32 | in_dim, while Q4_K and Q6_K need 256 | in_dim. Real small models mix both within one file — Qwen2.5-0.5B is 896 wide and SmolLM2-360M 960, each a whole number of 32-element blocks but neither a multiple of 256, so their attention weights take the 32-block kernels while only the 256-divisible ffn_down reaches a K-quant one.

Every supported type reduces to the same shape, which is what lets one dot loop serve all five:

weight[i] = scale[i / GROUP] * q[i]  -  min[i / GROUP]

min is zero for the symmetric types (Q8_0, Q5_0, Q6_K and IQ4_NL all fold their bias — or, for IQ4_NL, their codebook level — straight into the signed int8 weight, so no correction term survives); only Q4_K is genuinely asymmetric. GROUP is 16 because Q6_K carries one scale per 16 weights; every other type repeats its scale across the two halves of its 32-element block, which the dot loop exploits.

There are two entry points, because decode and prefill want different things. dot_row (GEMV) walks a row’s blocks once for a single token. unpack_row + dot_unpacked_multi (GEMM) unpack a row into plain int8 plus per-group scale metadata once per matmul rather than once per (row, token), which is what stops prefill from re-unpacking the same row for every token in the batch.

Quantizing activations to int8 is lossy, exactly as it is anywhere else — that is the accepted tradeoff of this kernel family, not an oversight. The tests check every kernel against quant::dequantize plus an exact f32 dot, to within the error int8 activation quantization can introduce, and separately require the two entry points to agree with each other far more tightly than either agrees with the reference (they quantize identically, so only summation order differs).

Model forward passes

One ModelForward implementor per architecture family (engine::arch:: mod), so adding a family is additive rather than a rewrite:

Image generation (engine::image)

A qwen_image GGUF (Qwen-Image; unsloth/Qwen-Image-2512-GGUF) is not a ModelForward and does not go through engine::arch at all. It holds a diffusion transformer — diffusers’ QwenImageTransformer2DModel, read from the file the ComfyUI-GGUF tooling writes, whose tensor names are diffusers’ own and whose metadata is three keys and no hyperparameters, so the loader (read_image_model_config) sizes it from tensor shapes: img_in.weight gives the stream width, attn.norm_q.weight the head width, and the transformer_blocks.N prefixes the depth.

Its role is config::Role::Image, and unlike the other five that is a fact about the file rather than a choice: Role::required_by(architecture) names it for qwen_image and nothing else, Role::fixed_by_model marks it as never prompted for and never narrowed by an x-orangu-role header, and main::fix_role_for_model reconciles it at every start with whatever was asked — the CLI flag, or under --daemon the config’s role key — refusing a conflict in either direction. The startup picker (select_model_interactively) draws the table in one of three model_spec::Dimming modes — ImageModels under --image, TextModels under any other flag, Unsupported with none — from ModelSupport::image, and refuses a greyed row; --init reads the same column to answer its own role prompt and to ask the image_* keys plus vae_precision and image_lora_merge (init::ImageKeys, written only where an answer differs from what the server would take: image_lora [auto] names the adapter find_qwen_image_lightning finds, and the step and guidance prompts after it offer lightning_steps of that adapter and 1, the server’s own rule). The Model: prompt itself pre-selects the config’s model = only if ModelSupport::selectable says the flag serves it, else the first row that is.

Which backend runs it is measured, not assumed. Under backend = auto, main::prepare opens the transformer before the encoder is uploaded and calls engine::image::calibrate: the first block’s img_mlp input projection (3072×12288, the widest linear and over half of every pass) through the selected device and through CpuBackend, 256 tokens of synthetic activations, each run twice and the second timed — the first pays pipeline compilation and the weight upload, a per-picture cost, not the per-step one being compared. If the CPU is faster the whole pipeline (encoder included) is rebuilt on CpuBackend. On the development board the Mali-G720 measured 352 ms against the CPU’s 39 ms, and a real 256-pixel step 111.7 s against 14.4 s: the GPU would have made the picture eight times slower, and auto used to pick it. An explicit backend is never overridden.

Its attention (transformer::joint_attention) is blocked: 24 queries per (head, task), the scores and the value product each a vecdot:: gemm_f32_rows tile — the values gathered transposed so the second product is a GEMM over the keys — and the softmax per row in between. The per-query form (joint_attention_per_query) is the reference it is tested against and the path on non-ARM targets. Measured at 512 pixels, attention went from 36% of a step to 17% and the step from 35.9 to 27.7 s; the 24-query block keeps a block’s score rows in L2 through the softmax.

The timestep-only vectors — every block’s modulation and the output norm’s — are cached on the transformer keyed by sigma, so guidance’s second pass at the same sigma skips sixty weight-bound one-token linears. The MLP’s GELU and the attention softmax’s exponentials go through tensor::gelu_inplace and tensor::exp_inplace, the engine’s vector polynomial exp; the scalar tensor::softmax_inplace the sampler and the language models use is untouched.

engine::image::lora reads a low-rank adapter — <linear>.lora_down/ lora_up/alpha in safetensors, keyed by the linear’s diffusers name and shape-checked against it at load — and by default merges it (LoraPair::merge_into): every adapted row dequantized (engine::quant), scale · B·A added, and re-encoded by the encoder that wrote the file — orangu::quantize, the GGUF encoders moved out of orangu-gguf into the library for exactly this — into a QuantMatrix::from_encoded_rows the linear then runs at full speed; 165 s on the development board for 720 linears — once, because lora::MergedCache writes the merged tensors to <models>/orangu-merged/<hash>.bin (a magic, a JSON header of identities and tensor ranges, the rows) and the next start maps them (QuantMatrix::from_shared_bytes) in seconds, keyed by both files’ path, size and mtime. image_lora_merge = no applies it instead, y += scale · (x · Aᵀ) · Bᵀ after each linear: on the CPU a token-blocked down product through vecdot::gemm_f32_rows and the backend’s two-dimensional float matmul for the up product, the two-matmul form kept for other backends and as the test reference, Stages::lora accounting its time apart. What it is for is Qwen-Image-Lightning — see the server manual’s Eight steps instead of fifty — and that is the default: config::ImageLora::Auto (the key absent) has prepare serve model_spec::find_qwen_image_lightning’s file (a Lightning…<N>steps name under models, the most steps first), which model_download::download_image_companions fetches with the model as a fourth companion; ImageLora::None is the base model. The picture defaults the pipeline is given are ServerConfiguration:: image_defaults_under(adapter): image_steps and image_cfg_scale as written, else a Lightning adapter’s step count (model_spec:: lightning_steps) and 1.

The VAE’s convolutions (engine::image::vae::conv_matrix) are encoded as Q6_K at load under the default VaePrecision::Int8, each im2col row zero-padded to a multiple of 256 so the K-quant kernel takes it, and matmul_k_gemm_into splits such a matmul over token tiles as well as row groups (vecdot::dot_k_rows_tiles) — 96 rows against 65,536 pixels was six tasks on twelve cores before. ActQ8Mm::quantize is NEON (vcvtaq, the same ties-away rounding as f32::round, so the bit-identity test still holds); it was a tenth of a decode’s samples as a scalar loop. And the im2col band no longer exists on this path: ActQ8Mm::quantize_with takes a closure that fills one token’s row into a scratch buffer, vae:: gather_window is that closure for a pixel’s 3×3 window, and CpuBackend::matmul_k_mm_into runs the kernel on the result — a full-resolution band was 268 MiB of f32 written and read back per convolution. Decode at 512 px went 10.6 → 2.7 s over the two steps.

engine::image::Pipeline is diffusers’ QwenImagePipeline, step for step, over four parts that each map onto a piece of the reference:

Latents travel between the parts in diffusers’ packed layout: [h/16 * w/16, 64] tokens, each a 2x2 patch of the 16-channel latent with features ordered (channel, dy, dx), normalised by the VAE’s latents_mean / latents_std. The pipeline runs one picture at a time (its Mutex): a step is a full pass over 20 billion parameters per image token, and two interleaved would be slower than two in turn. It runs on a blocking thread and reports ImageEvents — progress per step, then the picture — over a channel whose reader holds a CancelOnDrop, so a client that disconnects stops the work at the next step.

image::codec is the edge: PNG and JPEG through the image crate the tree already carried (via printpdf), GIF and lossless WebP through the same crate’s two further codecs, SVG drawn with resvg and then read as the raster it became; a picture starts from the attachment’s own format, so a JPEG comes back as JPEG. http::images is /v1/images/generations, and http::openai::image_chat turns a chat turn into the same request; the web console’s send_image_message does the same for its own SSE shape and keeps every picture as a file under the session’s files/. All three build the request against Pipeline::defaults() — a RwLock<ImageDefaults> behind a clone, so one request is built from one consistent set — which http::images::apply_settings replaces for POST /props (and the console’s POST /api/props), checked as config::parse_image_defaults checks the same keys; Pipeline::configured is what reset restores. http::images::props_json is the image object both /props and /api/props answer with, the rate model included so a client can cost settings before sending them.

What is verified, and what is not. The unit tests pin each piece to the reference’s arithmetic — the rotary layout, the modulation split order, the LayerNorm, the timestep embedding, the schedule’s values at 1024x1024, the convolution tap order and padding, the latent packing — the way every other architecture’s tests do, and the real-checkpoint test (ORANGU_TEST_QWEN_IMAGE_*) generates a small picture end to end. What does not yet exist is a captured-tensor fixture from diffusers itself (the convention in doc/DEVELOPERS.md): the machine this was written on cannot run PyTorch, so the transformer’s output has been checked structurally and visually rather than to the fifth significant figure. It is the first thing to add.

Constrained decoding (engine::constraint, engine::sampling::Constraint)

JsonPrefix is a byte-level recogniser for prefixes of valid JSON. At every point it answers “could this still become valid JSON”, which is what a mask needs, and separately whether what it has is already a complete document (is_complete), which is what a stop condition needs. Both are required and they are different questions: {"a": satisfies the first and not the second.

It is a hand-written state machine rather than a parser generator because the state has to be cheap to clone — testing a candidate token is “clone, feed its bytes, see if it survived”, and that happens per candidate per step. The state is a small enum, a stack of {/[ frames, and two flags.

Two details are easy to get wrong and are pinned by tests, having both been wrong first:

Constraint joins the grammar to a vocabulary. It holds the tokenizer’s token_bytes table — built once, on the first constrained request, and shared by Arc, so a deployment that never constrains anything never pays for it. The table comes from the same append_token_bytes that Tokenizer::decode uses, because a constraint reasoning about text the model is not actually emitting would be worse than none.

Where the mask is applied, and what it costs. On the greedy path the plain argmax is tested first and returned if it is legal, which it almost always is — one grammar probe, nothing sorted. Only when the model’s preferred token is illegal is the vocabulary ordered and walked. On the sampled path the mask is applied after top_k and before the softmax, so probabilities are renormalized over the allowed set; if every one of the top k is rejected the field is refilled from the whole vocabulary rather than failing.

Three things a constrained request gives up, all because they choose tokens without consulting the mask: the device-side argmax fast path, the batch coordinator’s greedy sampling, and prompt-lookup speculation. Each is gated on Sampler::is_constrained. A speculative draft accepted on “this is what greedy would have produced” would sail straight past the grammar.

Stopping. End-of-sequence ids are masked until the document is complete, and once it is complete they are the only ids left. The second half is not symmetry for its own sake: measured without it, a model that had just written {} went on emitting blank lines to max_tokens, because whitespace after a finished document is legal JSON and far more probable to that model than end-of-sequence.

Whitespace before the document is deliberately still legal. Banning whitespace-only tokens was tried, and it made the output worse rather than better: with no room to hesitate the model opens { on the first step and, having committed with nothing planned, immediately closes it — {} where allowing it to pause produced {"Name": "John Doe", "Age": 32}. The cost is that a model which refuses to emit JSON at all stalls on whitespace until max_tokens. A constraint can make invalid output unreachable; it cannot make a model cooperate, and pretending otherwise would trade a visible failure for a worthless document.

Request scheduling and continuous batching

engine::scheduler’s SlotPool bounds how many requests generate concurrently (slots in the config) and tracks each one’s progress for /slots. Each slot’s prefill+decode loop (engine::generate::run) runs on its own blocking-pool thread against its own KV cache — real concurrency, bounded fairly by slot count, but not a single fused multi-sequence GEMM by default.

Cross-sequence decode batching used to be an opt-in alternative for that last part (ORANGU_BATCH_DECODE), collecting concurrently-decoding requests inside a short window and fusing every sequence’s QKV/wo/FFN/PLE/lm_head matmuls into one backend call each. It has been removed, because it was measured and it lost at every size it was tried at.

window mean batch aggregate tok/s
4 ms 2.21 65.02
25 ms 13.86 45.47
100 ms 22.55 44.67
400 ms 26.98 42.75

At 32 concurrent streams and a mean batch of 27 the fused path was 34% slower than not batching. Unbatched decode reads ~78 GiB/s of weights here, at the card’s ceiling; batched at mean 27 it read ~1.9 GiB/s and still delivered fewer tokens. Fusing does remove the weight-bandwidth bottleneck — what remains costs more than it saves: the per-sequence attention, RoPE and KV write the fused path left per-sequence, plus the window every sequence paid on every token.

The first row is why an earlier comparison was wrong, and the reason worth keeping: it concluded “batching is 3% slower” at the default window, when the mean batch size was 2 — it had measured a rendezvous rather than a batch. That is why the mean batch size, not the window setting, is the number a batching claim has to carry.

A rendezvous can only build a batch by making sequences wait, and the waiting lost at every size tried, so no window setting wins. Closing this gap needs a scheduler that batches whatever is ready on each step without waiting for stragglers — not a larger constant, and not the code that was here.

Serving over TLS (tls.rs)

Built in rather than delegated to a reverse proxy, because “one static binary, nothing to install” is the property this project trades other things for, and “put nginx in front of it” spends exactly that property in the deployments the binary exists for — air-gapped, sovereign, one machine, no package manager. Terminating in front stays valid and is what most fleets will do; it just is not the only way to reach the server safely over a network.

It cost glue and no cryptography: rustls and tokio-rustls were already in the tree for reqwest‘s HTTPS, so what was added is an acceptor and a PEM reader — the latter through rustls_pki_types’ own PemObject, already present as a rustls dependency, rather than rustls-pemfile, which is the historical spelling and is now flagged unmaintained.

TlsListener implements axum::serve::Listener — accept a connection, report the local address — so it drops into the existing axum::serve call and keeps one serving path, one shutdown select!, and the ConnectInfo<SocketAddr> the loopback-only routes rely on. A hand-rolled accept loop feeding hyper would have duplicated all of that to add one wrapper.

accept cannot fail, and that shapes the error handling. A refused handshake is not a server error: on any network-reachable port it happens constantly — scanners, plain HTTP sent to an HTTPS port, clients with no shared cipher — so the loop drops those and continues. Giving up would turn background noise into an outage; logging each one would flood the log with it.

ConnectInfo needed an idiom. axum implements Connected<IncomingStream<'_, L>> for SocketAddr only for its own TcpListener, plus a blanket impl for anything wrapped in TapIo. Writing the impl here is not allowed — both types are foreign and a local type appearing only as a type parameter does not satisfy the orphan rule — so TlsListener::with_connect_info wraps in a TapIo that does nothing and earns the blanket impl. Worth the indirection because the alternative was dropping ConnectInfo to make it compile, which would have quietly widened what /model-cache/drop accepts.

Both keys or neither: tls_cert alone is a configuration error, because the alternative is a server that starts in the clear while the config looks like it does not. Certificate loading happens before anything is served, so a bad file is a startup failure naming it.

A note at startup when the server is exposed without either gate. Binding off-loopback is deliberate; serving an inference engine there unauthenticated and in the clear usually is not. The default bind is loopback precisely so that cannot happen by accident, so the only way to reach that line is to have widened host — at which point the omission is worth naming rather than leaving to be discovered from outside.

Authentication (http::require_api_key)

A route_layer over the whole API router. With no [orangu-server].api_key (or ORANGU_API_KEY) configured it passes everything through, which is the behaviour before it existed and the right default for the loopback address the server also defaults to.

Authorization: Bearer <key> rather than a scheme of this project’s own, because the orangu client already sent one — orangu::llm’s bearer_auth, from the client’s own api_key config — and every OpenAI-shaped client sends one. The server was the only half of this that was missing, so the work was checking a header that was already arriving, not designing a protocol.

The comparison is constant-time. A == on secrets returns at the first differing byte, so its duration reports how many leading bytes were right — enough, over many attempts, to recover a key one byte at a time. The cost of avoiding that is a few nanoseconds on a path about to run a forward pass.

OPEN_PATHS is /health and nothing else. /health names no model, reports no load, and returns the same bytes to everyone, so requiring a secret for it buys nothing and costs what every deployment needs: a probe that works from a load balancer with no credentials.

/v1/models is deliberately closed, and the coordinator was changed rather than the exemption widened. Its ensure_reachable and startup health check both probed /v1/models and required is_success() — so the first server given an api_key would have answered 401, been read as “stopped answering”, and been restarted on every request, in a loop. Both now accept any HTTP response as proof of life, which is what they were actually asking: a 401 proves a process is there just as well as a 200. Conflating reachability with authorization is the bug; exempting the path would only have hidden it.

Proxied requests need nothing: proxy forwards every header verbatim, so a client’s Authorization reaches the backend unchanged.

Admission: a bounded, first-come-first-served queue

SlotPool::try_acquire takes a ticket before it waits, and only the holder of the head ticket may claim a slot. Without that the wait is a scramble: a release wakes one sleeper, which then rescans, and a request arriving in that gap can call try_take_any first and take the slot it was woken for. Under steady load that is unbounded waiting for whoever is unlucky, with nothing in the system that would ever report it.

A global semaphore would be the obvious fix and does not work here, fair though tokio’s is: acquire_slot bypasses admission by design — a pinned request is waiting for one specific slot’s warm cache, not competing for one — so a global permit count drifts out of step with the slots actually free. That is the same failure this pool’s own comment records having removed a global semaphore for once already.

The ticket’s Drop is the load-bearing part. A waiter can vanish at any moment: the client hangs up, the request times out, the runtime cancels the task. A ticket that outlived its holder would leave the head pointing at a number nobody will ever claim, stalling every request behind it for the life of the process. Releasing on drop makes cancellation safe by construction rather than by remembering, and an_abandoned_waiter_does_not_stall_the_queue fails — by hanging — without it.

[orangu-server].queue_limit bounds the depth. Past it try_acquire returns None, which becomes a StreamEvent::Overloaded and then 503 with Retry-After. It is its own event rather than an Error because the two mean different things to a caller: an error is a request that cannot be served, and this is one that could be served later. /metrics exports orangu_server_queue_depth and orangu_server_queue_limit, because a bound nobody can see is half a feature.

Client disconnects already free their slot, and did before any of this. Both the streaming and non-streaming paths consume the same channel, so a dropped receiver makes the next tx.send fail and generation stops at the next token. Measured: a non-streaming client killed three seconds into a 2000-token request left the slot busy: false one second later, 125 tokens in, and the next request was served immediately. No cancellation plumbing was needed.

What this does not do is stop one long generation from occupying a slot while short requests wait. That needs preemption, which is a scheduler question rather than an admission one — see the batching section above for why the scheduler is the real gap.

The host-resident bound on prefix reuse

LayerCache::len counts every position the cache logically holds. The host buffers do not always hold all of them: the fused GPU decode path writes a token’s key and value straight into the device mirror and calls advance_gpu_only, which moves len without pushing anything host-side. After N decode steps len is N rows ahead of k/v.

Everything that reads the host side — cross-request prefix reuse, per-slot retained caches, slot save, CPU attention — therefore has to bound itself by KvCache::host_committed_len, not by committed_len. Bounding by the wrong one is not a subtle inaccuracy: it indexes off the end of the buffer.

That was a live crash, on the commonest usage there is — the same slot, a growing conversation, a GPU backend. The second turn panicked with range end index 31488 out of range for slice of length 30592: 246 rows claimed against 239 held, the difference being exactly the tokens the first turn generated. advance_gpu_only’s own doc comment had named the hazard in advance (“If prompt-prefix reuse (slot save/restore) is ever built, this becomes unsafe”), and before the host buffers were sized to their contents it read zeros instead of panicking, which is worse.

Two details of the bound are load-bearing. It is derived from the buffer length rather than tracked in a second field, so it cannot drift from what it describes. And it takes the shortest layer that holds anything, skipping permanently-empty ones: engine::arch::gemma’s cross-layer KV donor is empty by design, and a plain minimum would read it as “no prefix is reusable” and switch reuse off for the whole architecture.

What this costs is reuse of the generated tail, which the next turn re-prefills. Measured on a four-turn conversation: 695 of 923 tokens reused, the shortfall being the eight tokens the previous turn generated.

engine::slot_store also stopped locking its cells with lock().unwrap(). One panic under that lock poisoned it, and every later request on every slot then panicked on the unwrap rather than being served — a process that stayed alive and answered 500 to everything. It is a cache: a poisoned entry is at worst stale, the caller already handles absence, so it recovers the guard and reads as a miss.

Latency histograms and readiness (engine::metrics)

/metrics was three gauges and a limit. Gauges answer “what is happening right now”, which is the wrong question for latency — the useful questions are all about the tail, and a mean is dominated by whichever requests happened to be long. Four Prometheus histograms answer them instead: queue wait, time to first token, inter-token, and request duration.

Bucket counts are per-bucket and made cumulative only at render. A cumulative write would mean touching every bucket at or above the observed value on every observation, and the inter-token histogram is observed once per generated token. The format requires cumulative buckets, so the conversion has to happen somewhere; doing it on the read side costs a scrape and not a decode step.

Sums are microseconds in a u64, not a float. AtomicU64 is lock-free everywhere this runs, while a float sum accumulated by compare-and-swap would be slower and non-deterministic in its last digits across runs — which makes two measurements of the same workload disagree for no reason.

Two bucket sets, because one cannot span both. A request is milliseconds to a minute; a decode step is milliseconds to a second. Putting inter-token gaps on the request bounds lands almost every observation in one bucket, which is a histogram that reports nothing.

Time to first token counts from arrival, so it carries the queue wait and the prefill together — what an interactive caller actually waits through. It is the first token produced, not the first the client sees: a chat format’s structural prefix is filtered out of the stream (MessageHeader), and charging the template’s shape to the server’s latency would make two models of the same speed report different numbers.

Outcomes are recorded by a Drop guard, not at each exit. run has five places that report an error and return, one that returns on a client disconnect, and a catch_unwind above it that turns a panic into a reply. A counter with a path that forgets to increment it is worse than no counter, because the total silently stops matching the request count. OutcomeGuard defaults to error and is told otherwise on the paths that know better, so a failure path added later is counted correctly without anyone noticing it needed to be. A refusal is counted through its own entry point rather than as a zero-duration request: folding a handful of microseconds into the latency histogram would drag every quantile down exactly when the server is under the load those quantiles are being watched for.

Verified by arithmetic rather than by looking plausible. One 24-token request produced one time-to-first-token observation and twenty-three inter-token observations, and the two sums added to the request duration exactly (1.724549 + 1.157920 = 2.882525 s). Parts that add up to the whole is the check worth having, because every individual number here looks reasonable whether or not it is right.

/ready is not /health, and the split is the point. /health asks “is this process alive” — a supervisor’s question, and one that must stay 200 while the server is merely busy, because restarting a loaded server is the worst possible response to load. /ready asks “would a request sent now be served” — a load balancer’s question, where busy is exactly the case worth reporting. It is 503 when the admission queue is already at queue_limit (the request would be refused anyway, so saying so saves a round trip) or when the GPU device has been lost. The decision is a free function taking (device_lost, queued, limit) rather than reading AppState, so the rule can be tested at all — the alternative is standing up a loaded model to assert five comparisons, which is why rules like this usually go untested.

/ready joins /health in OPEN_PATHS, which is a deliberate widening: it does disclose load, where /health discloses nothing. That is the one fact the probe exists to report, it is bounded (depths and slot counts, no model name and no request content), and anyone able to send a request learns the same thing from the 503 they would get instead.

What this immediately found, and the care the finding needed. outcome="cancelled" sat at zero under a test that should have produced one, which looked like the disconnect abort not working at all. It was not that. The abort rode on sending a token — tx.send(…).is_err() — so it only ran for a token that produced visible text, and the model under test rendered one character for forty tokens. Almost every iteration skipped the check entirely.

That is a real hole rather than a quirk of one model: a chat format’s structural markers are suppressed, and under --review a whole reasoning body is, so a request could generate to max_tokens — holding the slot every other request is queued for — for a client that hung up at the first token.

The check is now tx.is_closed(), asked every token whether or not there was anything to send. Measured on the model that leaked: a 2000-token request whose client sent a TCP reset at four seconds used to run to completion and finish as length; it now stops within about four seconds and finishes as cancelled. On a model with ordinary visible output the behaviour was already correct, and still is — which is why the first reading of the zero counter, that disconnect abort did not work at all, was itself wrong.

Draft-model speculative decoding

[orangu-server].draft_model puts a second, smaller model in front of the served one: it proposes draft_tokens continuations, the target verifies all of them in a single multi-position forward, and the longest prefix the target would itself have produced is kept. Everything after that prefix is thrown away, so the emitted text is byte-for-byte what greedy decoding alone would have emitted — a drafter can only change how fast an answer arrives.

The verification machinery already existed for prompt-lookup drafting, and the two now differ only in where the candidate tokens come from. That is the whole abstraction: one enum with a draft and a commit, and the source with no model behind it simply has nothing to commit. Their cost profiles are opposite, which is why both are kept — prompt lookup produces nothing unless the context repeats itself but its misses are free, while a draft model always produces a draft and always pays a forward pass per token for it.

Both models need forward_all_logits, and the reason is not the obvious one. The target needs it because verification is a multi-position forward. The draft needs it because it is the only entry point that keeps the KV rows on the host: a single-token forward takes the fused GPU decode path, which writes key and value straight into the device mirror and leaves the host rows unpopulated. KvCache::advance_gpu_only’s own doc comment had predicted the consequence in as many words — “a resumed cache could need this position’s real data for a later multi-token prefill’s CPU attention path, which would silently read zeros instead” — and a draft cache is exactly that: rolled back to the accepted prefix and re-read on every step. Running the draft through forward produced both failures the warning describes, one loud (a read one row past the end of the host buffer) and one silent (garbage where zeros were read). Both models are probed with a one-token forward at startup rather than checked against a list of architectures, because the method is defaulted and there is nothing to ask a type about.

The draft’s cache follows the target’s exactly. After verification the target truncates to start_pos + 1 + accepted, and the draft truncates to the same count; when everything was accepted the draft is left one token short and the next step’s catch-up fills it in. Getting this wrong has no visible symptom — the target re-derives every token either way — beyond an acceptance rate quietly falling to zero, which is why commit is asserted directly rather than through output.

Two guards worth naming. The draft is capped by the room left in the KV cache, because a verification forward appends the whole draft at once and the GPU mirror is sized to the request’s capacity — a bound only reachable in the last few tokens of a full context, which is why it survived unnoticed on the prompt-lookup path. And the pair’s vocabularies are compared by token string, not by size: two tokenizers of the same size whose contents diverge is exactly what a same-family model at a different scale can be, and speculation compares token ids.

Measured, and it lost. On a 4 GiB card serving a target that overflows it (gemma-4-12B-it:Q4_K_M, 1.43 tok/s unassisted), a gemma-4-E4B draft achieved a better acceptance rate than prompt lookup — 2.15 accepted per verification against 1.67 — and still ran at 1.02 tok/s where prompt lookup ran at 3.01. Drafting deeper made it worse in exactly the way that identifies the cost: eight tokens per step raised acceptance to 2.56 and dropped throughput to 0.67, so the cost is per drafted token, not per step. Each draft forward is a pass through a second set of weights competing for device memory the target already overflows. The condition for a draft model to pay is therefore not about prediction quality at all — it is that the draft must be small enough not to disturb the target’s residency.

Not built: the DFlash sidecar format’s own draft-and-verify loop. A DFlash draft is not a standalone model — it reads the target’s hidden states at named layers (dflash.target_layers) and fuses them through its own fc — so running one needs ModelForward to expose intermediate hidden states, which would be an interface change landing on every architecture at once for one sidecar format. engine::arch::dflash still resolves such a file to its paired target and serves that. Sampled (non-greedy) acceptance is also not built: it needs the draft’s own probability distribution and a rejection-and-resample rule, which is a different algorithm from “did the draft match the argmax”, not a relaxation of it.

Durable slot persistence (engine::slot_store)

orangu-server implements the POST /slots/{id_slot}?action=save|restore endpoints — its equivalent of the --slot-save-path prompt-cache save/restore, and the receiving end of the orangu client’s per-session slot persistence (orangu::llm::SlotRegistry, driven from tab park/activate).

The important structural difference from that design is that an orangu-server slot is a concurrency permit (SlotPool), not a long-lived owner of one KV cache. A completed request’s cache otherwise survives only inside the in-RAM engine::prefix_cache pool (opt-in, bounded, cross-slot). engine::slot_store adds the durable layer: each slot retains the (tokens, KvCache) of the last request that ran on it, and

Prefix reuse from the cross-request pool moves the source’s buffers rather than copying them. PrefixCache::take_best_match removes the entry it returns, so the request owns it outright and nothing else can still be reading it; KvCache::adopt_prefix trims the source to the reused length and takes its layer buffers, keeping this request’s own capacity, kv_dim and stride. On a 2001-token conversational prefix that is 0.26 ms against 18.1 ms for the copy it replaced, and it removes the transient doubling of the prefix’s resident footprint at the moment the source was about to be dropped. The host buffers grow with the committed length instead of being pre-filled to the context ceiling, which is what makes the move possible.

engine::slot_store still copies, and the asymmetry is deliberate: it retains its snapshot for the same slot’s next request, so it does not own the source and cannot move out of it. Both paths go on sharing CachedPrefill::reusable_prefix_len, so the matching rules stay single-sourced even though the transfer differs.

A committed length is a token count, and on a block-compressed slot that is not its row count — one row stands for stride tokens. KvCache::committed_len converts through the stride, and reusable_prefix_len calls it rather than recomputing the maximum inline, which is what it used to do. Both readings gave the same answer on every architecture built so far, because deepseek4 is the only one that strides and it always carries ordinary per-token slots that are longer — a coincidence of the current models rather than a rule, and exactly the sort that a change to how rows are allocated would remove silently.

This is what survives the cases the RAM pool cannot: eviction under cache pressure, a server restart, and — most relevant behind orangu-coordinator — a model swap that tears the server down. The client saves a tab’s slot before the coordinator activates the next tab’s model, so the save reaches the still-active server and lands on disk; the later restore repopulates a freshly (re)started server.

<fingerprint> is a SHA-256 of the architecture, the model label, and the KV structure tag (layer count, per-layer kv_dim and stride, recurrent specs). The stride was missing until it was audited for: two caches differing only in how many token positions one row stands for laid their rows out differently and shared a signature, which the model label hashed alongside happened to mask — a signature that is right only because something else is also checked is not one worth relying on. A snapshot saved for one model therefore resolves to a different directory than any other model’s, and every file also carries the fingerprint internally — a mismatched or corrupt file is treated as “nothing to restore” (n_restored: 0, a normal prefill next request), never a hard error, so a stale sidecar never trips the client’s fallback notice. Client-supplied filenames are validated to a single safe path component before touching the filesystem.

The feature is on by default; set ORANGU_NO_SLOT_SAVE to disable it (it also stays off when $HOME can’t be resolved). While off, the endpoints report “not supported” exactly as a server started without --slot-save-path does — which the orangu client already degrades against, falling back to a full reprefill. The opt-out exists for the same reason ORANGU_PREFIX_CACHE is itself opt-in — a bug in prefix reuse would produce a silently wrong generation, not merely a slow one — but persistence is only ever exercised when a client explicitly saves or restores a slot, so it stays dormant unless used.

Requests routed through orangu-coordinator carry the session’s model in the slots request body (the orangu client adds it), so the coordinator proxies each save/restore to that model’s backing server rather than its default profile; a direct orangu-server or plain OpenAI-compatible server ignores the extra field.

Saved files accumulate one directory per distinct model. orangu-server prune sweeps slot files untouched for over 30 days on every run (slot_store::sweep_stale_slot_files), alongside its empty-session sweep — see Session management.

GPU backend architecture

engine::backend::Backend (backend/mod.rs) is the trait every backend implements — matmul/matmul_batch plus a downcast hook (as_wgpu) the model forward pass uses to reach VulkanBackend’s much larger fused surface when it’s the active backend. Six implementors exist: CpuBackend (scalar with runtime AVX2 dispatch via engine::tensor::dot, parallelized across output rows with rayon; always available, and the fallback when no GPU backend is found), VulkanBackend, MetalBackend, CudaBackend, OpenClBackend, and RocmBackend.

The hook is named for wgpu rather than for Vulkan because two backends answer Some to it: VulkanBackend and MetalBackend, which is that same engine on another wgpu API (see The Metal backend below). Every fused path reached through it is therefore live on Apple GPUs too.

main.rs’s select_backend implements the backend = auto cascade: Vulkan, then CUDA, then OpenCL, then ROCm (if built with the rocm feature), falling back to CpuBackend if none of them initialize. On Apple targets the cascade starts with Metal instead — not a preference but a cost: macOS ships no Vulkan driver, so leading with Vulkan there is four retry rounds of guaranteed failure before reaching the API the machine actually has, and MetalBackend gives up nothing, being the same kernels. On Windows, DX12 sits behind Vulkan and ahead of CUDA/OpenCL — also the same wgpu engine and the same WGSL (via naga’s HLSL output), so it reaches every fused path the matmul-only backends do not, but behind Vulkan because that is the API this engine was tuned on. An explicit backend = <name> instead brings up that one backend and fails to start if it can’t, rather than falling back — useful when GPU inference was asked for specifically and a silent CPU fallback would be the wrong failure mode.

Device selection

backend picks the API; engine::backend::device picks the device within it. Every GPU backend exposes the same two entry points — devices() -> Vec<DeviceCandidate> (enumeration only, no device creation, so it is safe and cheap on a machine with no driver) and try_init_index(index) / try_init_selected(&[index]) — and select_backend drives both through the one shared policy, resolving --device first, then ORANGU_DEVICE, then [orangu-server].device (requested_device).

Selection returns a set, not a device. device::select_all answers with every candidate the request admits, best first: under auto that is the whole ranked hardware list, and under an index or a name it is exactly one — which is what makes a named device exclusive rather than merely preferred. The head runs the model; the tail is carried into VulkanBackend::device_selection purely so the startup inventory and /props can report it. The distinction is load-bearing before any placement pass exists: “orangu chose this card out of three” and “orangu was told to use this card and nothing else” are different runs, and only the second one stays correct when a second card appears in the machine.

DeviceRole is what the inventory reports per device — InUse, Idle (selected but not running the model), or Excluded with a reason. An idle device is named as idle rather than as “available”: a second card sitting unused beside a slow first one is a question worth answering on the same screen as the number that prompted it.

The enumeration itself goes through the same short retry as bring-up (devices_with_retry). A driver whose previous context has just been torn down can briefly report no adapters, and request_adapter used to hide that inside the retried call by both choosing and creating in one step — splitting the two apart would otherwise have turned a restart race into “no device was found” on a machine that has one.

The policy ranks by class, then by size: discrete > unclassified > virtual > integrated, with software rasterizers never selected automatically (orangu’s own CpuBackend is faster than llvmpipe, and a software adapter reporting itself as a GPU run is worse than useless). Class beats size deliberately — an iGPU reports the machine’s whole system RAM as its memory and would otherwise win on a laptop. Unknown size ranks last within a class but never demotes a device out of it: “unknown” is not “zero”, which is llama.cpp’s own rule for its --fit accounting.

What this replaced was request_adapter(PowerPreference::HighPerformance) — a hint, answered by the loader, and routinely answered with the integrated GPU on a machine that also has a card. llama-server has the same trap and the same fix (--device Vulkan1).

Three properties are worth keeping if this code is touched:

VRAM comes from vulkan_replay::adapter_device_local_bytes, which reaches through Adapter::as_hal to vkGetPhysicalDeviceMemoryProperties — wgpu has no memory query on any backend. It reports heap size, not VK_EXT_memory_budget’s free figure: ranking wants a property of the card, and a card that happens to be driving a compositor must not be demoted below an iGPU for it. On non-Vulkan APIs it answers None, which the policy already handles.

Each backend’s try_init() (no index) still exists but is now #[cfg(test)]: tests want a device and don’t care which, while the server proper always goes through enumerate → select → report → try_init_selected.

The KV mirror grows with the sequence

LayerCache::sync_gpu allocates the GPU-side mirror for the rows in use, doubling from a 256-row floor and capped at the layer’s capacity. It used to allocate the whole capacity — which is prompt + max_tokens — on first use, so a request that asked for a large budget reserved it in VRAM whether or not it generated anything. On a 3.98 GiB card, the same two-token prompt and one-word answer took 2191 MiB at max_tokens = 64 and 3727 MiB at 32768; it is now flat at ~2185 MiB across both.

Host buffers never had this problem and were not changed: a large zeroed Vec is mmaped, and the kernel commits pages only as they are written — two gigabytes of vec![0.0f32; ..] adds no resident memory at all. Device memory is not overcommitted, which is why the same over-reservation that costs nothing on the host costs the whole card on the GPU.

Growth copies the rows already on the device across with copy_buffer_to_buffer rather than re-uploading them, so a doubling costs device-local bandwidth instead of putting the whole cache back over the bus; synced_len carries over unchanged. Cached attention bind groups name the old buffer, so they are dropped on growth and rebuilt through the caller’s existing cache-miss path.

The mirror is sized to len + 1, and that is load-bearing. The fused decode path binds these buffers and then writes the current token’s key and value at row len, before the host-side push that commits it. Sizing to exactly len puts that write one row past the end of the k region — and because k and v are two sub-ranges of a single buffer, it lands on row 0 of v rather than outside the allocation, so no validation fires and the corruption is silent. Reserving the full capacity hid this indefinitely; it appeared one growth step after demand-sizing landed, as an overrun on the grow-copy.

Device footprint

engine::footprint::DeviceFootprint measures the loaded model against the chosen device, at startup, and prints it under the inventory.

It reports, it does not refuse. Weights reach the device lazily and are never evicted, the KV cache is sized per request rather than at the context limit, and the arenas grow to the widest prefill — so “does it fit” is not decidable at startup, while headroom and what that headroom buys are. A model whose weights exceed VRAM gets a warning naming the shortfall and keeps running: the driver pages, which is slow rather than broken, and refusing would convert working configurations into failures. Resist adding a verdict here.

Splitting a model across devices

engine::placement decides which device holds which layer; engine::backend::multi::MultiDeviceBackend makes it happen. The whole design rests on one observation: a QuantMatrix can carry its own device.

LoadedModel::matrix is the single place every architecture obtains a weight, and it knows the tensor’s name — so it stamps each matrix with the device its blk.<n>. layer was placed on (LoadedModel::layer_device, set by main between loading the weights and building the model, because building the model is what calls matrix). MultiDeviceBackend::matmul then reads w.device() and forwards. Not one line of any forward pass changes, and there are eleven of them.

The cross-device transfer falls out of the same shape rather than being written: Backend::matmul already takes host &[f32] and returns host Vec<f32>, so a layer on device 0 ends with its output in host memory and the next layer’s first matmul uploads it to device 1. There is no peer-to-peer path and no residual to shuttle by hand.

Fitting a context (--context <TOKENS>, ORANGU_CONTEXT, [orangu-server].context). Unset, every layer goes to the card when the weights fit and the context is whatever the card has left after them — the pool’s share of the headroom, sized from the driver’s budget — and a longer prompt is refused. Set, the KV cache that many tokens need (in the device’s own storage width, n_layer · n_head_kv · head_dim · 2 per token) plus the prefill’s scratch reserve is taken out of the head card’s budget before the layers are placed, so a model whose weights fit the card with no room for that context has its last layers moved to the host until there is (device_split = cpu, reported at startup with the plan). It is a choice rather than a default because it costs decode: a layer on the host is read at the host’s rate every token, and the refusal is what keeps a model that fits the card at its full rate. Measured on a 4 GiB card with a 3.2 GiB Q8_0 model asked for 8192 tokens: 12 of 28 layers to the host, 9344 tokens of context where there were 1728, decode 43 → 9 tok/s at depth 0. A model already split for its size gets the same reserve on top of its placement. The card’s other tenants count: what the compositor holds is planned around, not over.

Windowed layers keep a window. A gemma-4 file’s sliding-window layers — five of every six — never attend further back than their window (512 or 1024 tokens), so their device copy is a ring of the window plus one prefill chunk rather than one row per position, whatever the context: kv_elements_per_token counts only the full-attention layers, and the rings are a fixed charge per slot taken out of the pages’ budget first. The host copy keeps every position, so rollback, prefix reuse and slot save are unchanged, and on a paged model the pool holds those layers’ rows for sharing while allocating them no device pages. A layer that shares another’s cache holds nothing at all. Together on a 4 GiB card: E4B’s ceiling 2944 → 15424 tokens, and --context 8192 needs no layer moved. ORANGU_SWA_FULL=1 keeps a full-context copy on every layer, for comparison.

Two hooks, and the line between them is layer scope.

Backend::as_wgpu_on(device) is for work scoped to one layer: fused attention, the fused post-attention/FFN chain, the device-side KV mirror. Each takes host input, returns host output, and touches only that layer’s weights and cache, so it runs happily on whichever card the layer is on. device is always read off a weight the call is about (QuantMatrix::device) rather than tracked separately — one map, living on the weights, so it cannot disagree with where matmul sends the same layer’s operands.

Backend::as_wgpu() still answers None on a split, and now means something narrower: work that spans layers. The whole-step decode submission (~37 submissions down to 1), GPU sampling, the logits readback. Those assume one device holds the whole chain. multi::tests::a_split_model_never_exposes_a_wgpu_backend and per_layer_work_asks_the_layer_s_own_device hold both halves.

What that None costs the reporting, and what is done about it. Three things hang off as_wgpu() that have nothing to do with running the model: the kernel/tuning report, the device footprint, and /gpu-timings. A split therefore used to lose all three at once, silently — /props carried a placement plan in the gpu slot with no tuning fields, the footprint was never measured, and the timings endpoint answered enabled: false in the same words it uses for “you did not switch timestamps on”.

Two of the three are now answered per device instead of not at all:

orangu-bench renders both, and the last of the three — the kernel report — is printed as explicitly absent rather than omitted.

The KV mirror is safe by construction: a layer’s device never changes, sync_gpu is only ever called from inside a VulkanBackend (so the mirror lands on that backend’s own device), and LayerCache::copy_prefix_from drops the mirror, so no buffer survives into a cache reused elsewhere.

Measured on the dev machine, release build, a 0.5B model split 3:1 over two GPUs: 11.9 tok/s with per-layer fusion off, 14.9 with it on, against 27.8 unsplit. Per-layer fusion is worth about a quarter. ORANGU_NO_SPLIT_FUSION=1 is what makes that A/B possible from one binary — measuring it by building a second binary is how a stale copy ends up being the thing timed.

The whole-layer decode chain, per device run. A split used to lose record_fused_layer for the whole model, not just at boundaries: the recorder is reached through as_wgpu(), which answers None. Every layer therefore fell to the step-by-step path, which round-trips through host memory between individual ops.

LlamaModel::record_split_decode restores it. record_decode_run takes a layer range, a host input vector, and whether to append the tail; a single-device model is one run over every layer with the tail — the same code that ran before it took a range — and a split model is one run per device, with the hidden state crossing to host in between (VulkanBackend::submit_and_read_at). The vocab projection runs where output_weight is, which is device 0, so a model whose last layers are elsewhere pays one more hand-off.

Two measurements decided that design, and the first killed the obvious alternative:

Result, one batch, 0.5B model, 3:1 split:

decode
unsplit 41.8 tok/s
split, all split GPU work on 21.2 tok/s
split, all off (ORANGU_NO_SPLIT_FUSION=1) 12.1 tok/s

+75%. And the diagnostic that exposed the bug is monotonic again: one iGPU layer now costs 26.4 tok/s against 41.4 unsplit, twelve cost 19.1 — consistent with a slower second device and a per-boundary hand-off, which is what is left and is inherent.

llama, phi and mistral all have it, each verified live on a split with byte-identical output to the same run with the paths disabled:

family split, on split, off
llama/qwen2 (0.5B Q4_K_M) 21.2 tok/s 12.1
phi (Phi-4-mini Q4_K_M) 14.0 tok/s 9.5
mistral (Ministral-3B IQ3_XXS) 5.0 tok/s 2.9

Consistently +47% to +75%. gemma is verified for correctness on a split (gemma-4-12B across both cards) rather than for speed — 36 of its 48 layers land on a 4 GiB card, so that run is dominated by driver paging and says nothing about the chain. The one piece shared across the three is arch::decode_device_runs — the grouping of consecutive layers by device, and the rule that any layer without a GPU behind it declines the whole chain (which is how a CPU overflow tier opts out). It is shared precisely because its failure mode is silent: a mis-grouped run would record a layer’s fused chain against the wrong card’s weights.

gemma has it too, with one exclusion. Its recorder takes an encoder from the caller — the cross-sequence batched path shares one across sequences — so the range and with_tail slot in without disturbing that, and record_split_decode brings its own encoder per run.

The exclusion is per-layer embeddings. record_one_sequence_decode projects the token embedding once into a [n_layer, per_layer] ple_buf that every layer reads a slice of, and that buffer belongs to one device. Worse, a later run’s x is a mid-model hidden state, not the token embedding, so recomputing it per run needs the original vector threaded through as well. A model with PLE therefore declines the chain and takes the step-by-step path; that is gemma-3n (gemma-4-E2B and relatives). Dense gemma-4 has per_layer == 0 and is unaffected. has_ple also requires layers.start == 0, so the guard is in the code as well as in the caller.

(Absolute numbers move between measurement batches with GPU clock state. Only compare within one batch.)

Points worth preserving if this is extended:

The host tier’s decode. Three things bound a split model’s decode once most of its layers are on the host, and each was measured on a 48-layer model split 13:35 over a 4 GiB card and the host (5.7 tok/s against 4.3 before, level with the reference engine):

The CPU as a device

SplitMode::Cpu appends CpuBackend to the device set as the last entry and lets placement place layers on it. Nothing else was needed: MultiDeviceBackend holds Arc<dyn Backend>, and the CPU is one.

Three things make it a fill (placement::fill_in_order) rather than a share, and they are the parts worth keeping:

WEIGHTS_SHARE_OF_DEVICE (0.8) is the only invented constant in the device work. It cannot be computed: the KV geometry needs a built model, and the model cannot be built until placement is decided, because building it is what stamps each tensor’s device. Explicit ratios are the escape hatch, and the footprint report says afterwards what the choice left.

project_expert reads the quantized bytes directly, through the same engine::vecdot integer-dot kernels CpuBackend uses for every dense matmul: dot_row/dot_k_row when one token is routed to an expert, and unpack_row once plus dot_unpacked_multi when several are. Only the types vecdot has no unpacking for fall back to dequantizing a row to f32 and dotting it there.

That was worth finding. A routed expert’s rows are read once and thrown away, so the f32 row was pure overhead — materialized, dotted, discarded — and on a nemotron_h_moe decode profile quant::dequantize_into under this one function was 60% of all CPU time in the process. It is the same shape for every mixture-of-experts family here, since they all share this path. Note what it also does to the arithmetic: an expert now quantizes its activations exactly as the dense path always has, so an expert’s numerics are consistent with the rest of its layer rather than more precise than it.

configure_cpu_threads sizes rayon’s global pool once, before anything parallel runs — CpuBackend’s matmul, project_expert, and the per-expert fan-out all share it, so the knob belongs there rather than on any one of them. Unset leaves rayon’s own default, so a config that says nothing keeps the behaviour it had. 0 is rejected: rayon reads num_threads(0) as “the default”, which would make a typo silently mean the opposite of what it looks like.

NUMA and P-core/E-core awareness stay out of scope. That is a decision rather than an oversight: rayon’s pool is unaware of core topology, and pinning workers to a socket or to P-cores only pays once there is a measurement saying the default placement is the bottleneck.

Device expert tiers — the seam, and why it stops there

Routed experts are host-resident (is_cpu_only_tensor), and a hot subset lives in owned RAM under engine::expert_store’s budget, with an LRU/LFRU policy and a learned-heat sidecar that survives a restart. A device tier would be the same idea in VRAM. Two pieces of it exist:

The dispatch now exists, behind ORANGU_GPU_EXPERTS=1 and off by default. arch::gpu_project_expert views one expert as a QuantMatrix (ExpertQuantMatrix::expert_matrix) and hands it to Backend::matmul — no new kernel, which is what the seam was for: every GPU backend already has one for every quantization an expert is stored in. It declines to the host path when the backend has no GPU or no kernel for that type (the IQ* types large MoE models often ship in are exactly the gap).

It is a measurement knob before it is a feature, and it does no residency management at all: every expert it touches lands in the weight arena, which never evicts. Point it at a device that cannot hold them and the number it produces is driver paging.

What it exists to answer is R2’s own premise — whether a GPU expert matmul beats engine::vecdot’s tuned AVX2/rayon path at all, given that this dispatch is a blocking submit-and-readback per (expert, projection, layer).

Measured, and batching is what decides it. Qwen3.6-35B-A3B with its dense part on the iGPU, page cache warm, same prompt throughout:

routed experts on decode
CPU (engine::vecdot, AVX2 + rayon) 2.12–2.14 tok/s
GPU, one dispatch per expert 1.39–1.42 tok/s
GPU, batched across experts 2.67–3.71, settling ~3.3 tok/s

One dispatch per expert loses to the host by 1.5×; batching them — arch::evaluate_routed_experts_batched, one matmul_batch per (layer, projection) rather than one blocking submit-and-readback per expert — turns that into a 1.55× win. matmul_batch requires a uniform token count, so ops bucket by n_tokens; at decode that is one bucket.

Three traps in measuring this, all hit while doing so:

The tier is bounded. main::plan_expert_tier chooses the resident set up front from expert_tier::plan — half the device’s free memory after the dense weights — and stamps it per expert (ExpertQuantMatrix::is_device_resident). Non-resident experts stay on the host path through engine::expert_store as always, so the batch never pulls an expert into an arena that cannot evict it. On the model above that is 15978 of 30720 experts in a fixed 9.40 GiB, measuring 2.68–2.74 tok/s: some of the unbounded win given back, because 48% of routed experts fall back, in exchange for a tier that does not grow until the device is full.

Five MoE architectures take the batched path — qwen35moe, qwen3next, glm, deepseek4, kimi3 — each supplying its own activation closure. gemma’s MoE does not: it projects a fused gate_up tensor by row range rather than separate gate/up tensors, so it needs a variant of the helper.

The resident set is filled from the routing profile (ORANGU_EXPERT_USAGE, expert_store::learned_heat) when one exists and by size otherwise; the startup line says which, because that choice is most of what the tier is worth — colibri measured the same tier 3–5× apart depending on it. Every number above is the by size floor.

The batched path bypasses engine::expert_store’s residency tier, which is correct for weights in VRAM and wrong for the host path, so non-resident experts keep project_expert.

What is still not built is the tier, and the projection is why. On this project’s dev machine a 20.6 GiB MoE model on a 4 GiB card leaves room for 5% of the experts — a tier that cannot pay for itself whatever the kernels look like. Three things would have to be true before it should be:

  1. Enough coverage to matter, measured on the target machine, and ORANGU_GPU_EXPERTS=1 beating the host path on it. If the naive dispatch loses badly there is still hope in step 2; if it loses by an order of magnitude, there is not.
  2. Batched dispatch. project_expert is called per expert per projection per layer, and the knob above issues one blocking submit-and-readback for each. The shape that works is one matmul_batch per (layer, projection) over every routed expert — MatmulOp already carries a per-op x, so the operands fit, but evaluate_routed_experts would have to be restructured into gather-then-batch without disturbing its bit-identical accumulation order, across six architectures.
  3. A bounded residency. VulkanBackend::weight_buffer’s arena never evicts, so experts reaching it on demand grow without limit. Placement has to choose the resident set up front, from the profile — which is exactly what expert_tier::plan returns, and which ORANGU_GPU_EXPERTS deliberately skips.

The honest prior is not favourable: colibri’s own finding is that a GPU expert tier “earns its VRAM only when the CPU is the weak link”, and orangu’s host expert path is tuned AVX2 over rayon.

The Vulkan backend

VulkanBackend (engine::backend::vulkan, via wgpu’s Vulkan backend — ash dlopens the system Vulkan loader at runtime, so no Vulkan SDK is needed to build, only a driver to run against a GPU) is the mature, hardware-verified backend. Each supported ggml_type gets two WGSL compute pipelines sharing the same per-type dequantization math (dequant_element in vulkan_shaders.rs, a line-for-line port of engine::quant’s dequant algorithm restated in WGSL), dispatched differently by n_tokens:

A weight tensor is uploaded once (still quantized) and cached on the GPU for the model’s lifetime. For Gemma-family models, VulkanBackend:: fused_attention chains QKV projection, Q/K-norm, RoPE, the KV-cache write, and the attention kernel itself into one GPU submission; fused_post_attention similarly chains the residual add, RMSNorm, and GEGLU; record_fused_layer/fused_layer fold a whole layer (attention + FFN) into one command encoder; and GemmaModel::forward chains every layer plus output_norm/lm_head into one shared encoder per decode step. Together these collapse the number of GPU submissions per decode token from one per matmul/op down to a small constant (as low as one for a fully-fused Gemma decode step), removing the per-submission submit/poll/ readback latency that otherwise dominates a many-layer forward pass. With round trips largely eliminated, the remaining cost is per-kernel compute and weight-memory bandwidth, which the alternative decode kernels below target.

Vulkan backend environment variables

The Vulkan backend reads these environment variables at startup to select between alternative compute kernels. Each is read once when the backend initializes; changing one takes effect on the next server start. All are correctness-verified against CpuBackend.

Boolean flags read 0, false, no, off and the empty string as OFF, and anything else as on — engine::env::flag_on, which every one of them goes through. That matters because these knobs exist to be swept: they were previously read for presence, so FLAG=0 switched the feature on and a sweep of 0,1 ran it on both arms, reporting the difference between a thing and itself. The handful of variables that carry a value rather than a boolean — ORANGU_NORM_WG, ORANGU_COOP_GEOM, ORANGU_DUMP_SHADERS, ORANGU_EXPERT_USAGE, ORANGU_PREFIX_CACHE_DIR — are still presence-checked or parsed, and are noted as such where they appear.

Variable Default Effect
ORANGU_PACKED_GEMM off (1 turns it on) Runs the float prefill matmuls (F32, F16, BF16 weights) through a packed, cache-blocked GEMM from the rten-gemm crate instead of this engine’s own tiled kernel. The switch is read where the kernel is chosen, not inside it, so the shape rules can be tested independently of the default. Off because it is not faster on the hardware it was measured on. With normal float values the two kernels are level — between 0.92x and 1.28x across every shape a forward pass runs, and the spread between repeats of one shape is wider than the difference between the kernels at that shape; end to end on a mixture prefill it is worth about 1% at a 624-token prompt and nothing at 2,236, where the closing control of an interleaved A/B came back at the treatment’s rate. The packed kernel’s advantage depends on the host’s instruction set, and the machine this was measured on has AVX2 but neither AVX-512 nor VNNI, so the crate selects its weakest x86 kernel — the comparison is worth repeating where it can select its best, which is what this knob is for. It never applies to quantized weights: this engine multiplies the packed bytes directly, and dequantizing them so a float GEMM can read them costs more than the whole product. It never applies at one token either, where a mat-vec has no reuse to block for.
ORANGU_PREFILL_EVEN on (0 turns it off) Whether a prompt’s chunks are evened out. The sizer answers how wide the next chunk may be and never what that leaves for the ones after it, so under a 2048 ceiling a 2239-token prompt came out [2048, 191] — two passes, the second a runt paying a whole pass’s fixed cost for a tenth of a pass’s work. Evened, the same prompt is [1120, 1119]: the same number of passes, so the same weights streamed, and the widest chunk halved. Peak prefill scratch is set by the widest chunk, and on a card with little free VRAM that scratch is what pages. Never widens a chunk past what the sizer allowed and never adds a pass, so neither the driver’s submission limit nor the streaming cost can be made worse. A diagnostic control arm rather than a tuning knob; 0 takes the sizer’s width as it comes.
ORANGU_PREFILL_BATCH unset (chosen per prefill; see below) (integer, not a presence flag) How many prompt tokens go into one forward pass. Setting it fixes the width in every regime and overrides the choice described here; 0 means the whole prompt in one pass. Left unset, the width depends on the backend and on how much of the model is in RAM. On a backend with a driver submission timeout the value is a ceiling the adaptive sizer works below (see ORANGU_PREFILL_CHUNK_MS), and it is 512: measured on a 4 GiB card holding 2.5 GiB of weights with a 17.5k-token prompt, no chunking lost the device after 21 s, 2048 after 3m54s, and 512 completed in 4m13s with peak VRAM identical (3.67 GiB) in all three. On a backend with no such timeout the width is flat — no probe, no shrink — and is 512 when the model is resident and the whole prompt when it is not. Resident, narrow is faster and smaller: at 8,001 tokens, three interleaved pairs, 512 prefilled 10.6% faster than one pass and peaked at 1,988 MB against 2,967. Streamed from disk the ordering reverses, because every extra pass re-reads what is not cached: at ~1,030 tokens one pass ran at 45.7 tok/s reading 1.29 GiB, against 35.4 tok/s and 4.97 GiB at width 256. A mixture-of-experts file whose routed experts run on the card gets 2048 instead of 512, for the same reason one level down the memory hierarchy: each pass streams a layer’s whole expert stack across the bus, so a prompt taken in five passes moves five stacks. Measured on a 30B-A3B file at a 2,236-token prompt — 512: 86.5 tok/s moving 0.34 MiB of expert weight per token; 1024: 94.7 and 0.23; 2048: 96.3 and 0.17. A file with a sliding window is excluded, because its ring mirrors are sized as the window plus one chunk: widening the chunk doubles them, and on a 26B-A4B file that cost 18.7% at 626 tokens and 32.5% at 2,238 even though the expert traffic halved.
ORANGU_PREFILL_CHUNK_FLOOR unset (integer; 0 is no floor) The narrowest chunk the adaptive sizer may choose, overriding the floor the model itself asks for. Two kinds of model ask for one. A model with layers on the host is floored at ORANGU_HOST_STREAM_TOKENS, below which a host layer’s projections run on the CPU rather than streaming to the card — a 16-token probe measured there says a 512-token chunk would take minutes, and every chunk after it stays narrow. A mixture whose routed experts run on the card is floored at ORANGU_EXPERT_GEMM_TOKENS for the same reason: below it the chunk takes the host’s per-expert path, which is slower, which sizes the next one narrower still. A model with both reasons gets the larger of them. What the floor does in practice is make the opening probe a real pass: measured on a 30B-A3B file, +6.7% at a 2,236-token prompt (two expert-stack streams rather than three) and −2.7% at 624 (where the 16-token probe on the host path was the cheaper opening); on a 26B-A4B file, level at both. It exists as a knob so both settings can be compared on one binary — a floor is chosen by what the model is, so otherwise the two arms would be two builds.
ORANGU_PREFILL_CHUNK_MS 3000 (integer) Wall-clock target for one prefill submission, on a backend that has a submission timeout only. A chunk is timed and the next is sized from what that chunk revealed about this machine’s cost, because cost per token climbs with context: a fixed 512-token chunk measured 2.3 s at position 512 and 10.1 s at position 6 656, past the ~10 s amdgpu allows before it resets the device (see Losing the GPU device). It is ignored where there is no device to lose, and deliberately: the quotient it adapts on is a per-token rate only for a resident model. When weights stream from disk each pass costs about the same whatever it contains, so a narrow chunk reports a huge apparent rate, the next chunk shrinks, and the fixed cost is paid again over fewer tokens. Measured, that does not settle — it lands on the 16-token floor on the first chunk and stays: a 1,016-token prompt became 63.8 passes and 78.5 GiB of reads on a 1.23 GiB model, prefilling at 0.8 tok/s, against one pass, 1.29 GiB and 45.7 tok/s once the adaptation is dropped.
ORANGU_PREFILL_CHUNK_SPREAD 4 (integer; 0 turns it off) How a chunk’s wall time is counted against ORANGU_PREFILL_CHUNK_MS. The budget guards against one submission outliving the driver’s timeout, and a chunk is many submissions — a handful on the device-resident stream, one per matmul on a host-orchestrated model. A chunk of n submissions counts as elapsed × spread / n (its mean submission time with a margin for the longest), never more than elapsed; a chunk of spread submissions or fewer counts in full, which leaves the stream’s chunks exactly as they were. Before this, a model split onto an integrated GPU — where a 128-token chunk takes ~34 s across ~200 submissions — was cut to eleven-token chunks, each paying every layer’s fixed cost, and prefilled a 572-token prompt at 1.5 tok/s where one chunk ran at 5.7. 0 counts the whole wall time, the rule this replaces, the control arm.
ORANGU_CHUNK_POLICY unset (chosen from the backend) adaptive or flat, forcing the prefill splitting strategy that the backend would otherwise pick — adaptive is the timed sizer above, flat keeps every chunk the configured width. Exists so both strategies can be measured on one machine, since otherwise the choice is a property of the hardware and an A/B would need two of them. An unrecognised value, empty included, is ignored rather than treated as one of the two: a typo in a harness must not silently select an arm.
ORANGU_NO_CHUNK_COST_FIT unset (the fit is on) Size each prefill chunk from the previous one’s rate alone, as this did before the cost of a submission was modelled as fixed + per_token · n. The control arm, kept so the change can be measured against what it replaced in one orangu-bench --sweep rather than across two sessions. Why the fit exists: a submission pays for itself before it processes a token, so reading elapsed / n as a per-token rate charges all of that to per-token work — worst on the opening 16-token probe, which is almost entirely fixed cost. Its apparent rate then comes out several times the real one and every following chunk is sized short. Two observations at different widths separate the terms, and the estimate is kept across requests, so only the first prompt a server sees pays a ramp. Measured on gemma-4-E2B-it:Q4_K_M, a 1,120-token prompt went from 5 submissions (16, 122, 351, 473, 158) to 3 (512, 512, 96). Both numbers are measured here rather than carried in as constants, which is what makes the sizing portable to a machine with a different round-trip cost.
ORANGU_PREFILL_TRACE unset (1 turns it on) Every 256 readbacks of one size, prints the wall clock of a device submission’s phases, in microseconds each: alloc (the staging buffer), finish (the encoder’s own — the commands recorded since it was opened, translated for the driver; ~7 µs a dispatch on the CIX P1), submit (the queue’s), wait (the device, plus launch and readback latency) and copy (the mapped bytes out). The small size is a decode step’s; a device idles through finish and submit unless an earlier submission is still running (see ORANGU_HYBRID_FLUSH_LAYERS).
ORANGU_PREFILL_CHUNKS unset (off) Print one line per prefill naming the chunk widths the sizer chose, e.g. 1120 tokens from 0 in 3 submissions: [512, 512, 96]. Answers “what did the sizer actually do” without ORANGU_PREFILL_TRACE, which writes inside the submission loop and therefore changes the cost it is measuring — under it the sizer reads inflated chunk costs and picks narrower chunks, so it cannot be used to check the sizer’s own choices.
ORANGU_EXPERT_TIER_FLOOR 25 (percent) The share of recorded routing a device expert tier must cover before it is built. Below it the tier is declined and the startup line says so, naming the coverage, the floor and the VRAM not spent — that VRAM goes to the KV cache and the transient arenas instead, which on a small card are what a tier competes with. Coverage is only known once a routing profile exists (ORANGU_EXPERT_USAGE); with no profile there is no evidence either way and the tier is built, since the profile can only be recorded by running with it on. The default rejects the obviously futile case rather than marking a measured optimum: the only coverage measured end-to-end here was 44.4%, and it was 9.4% slower on decode, so clearing the floor is not a promise that the tier pays. 0 never declines.
ORANGU_NO_MLP_UNROLL unset (block-unroll on) Set to disable the block-unroll reduce kernel for K-quant (Q4_K/Q5_K/Q6_K) decode and fall back to the scalar per-element reduce kernel. The block-unroll iterates whole super-blocks, loading each block header once and issuing several weight/activation loads before the dependent dot; it is the default decode path.
ORANGU_NO_FLOAT_WIDE unset (wide float on) Set to serve F32 and F16 weights at decode with the generic per-element reduce instead of the vec4, loads-in-flight kernel — the generic one reads each weight byte by byte through the type’s dequant_element and each activation alone, one round trip per element. Matters for a model that carries F32 tensors on the decode path (the per-layer-embedding gate and projection of the gemma families): both dropped to about half their time. Same arithmetic, different summation order. The control arm of its measurement.
ORANGU_Q6K_V4 unset (word reads on) Set to have the Q6_K light decode kernel read its 210-byte blocks through 16-byte loads with a per-load alignment test, as it originally did, instead of as words with each thread’s misalignment resolved once (a thread’s blocks are all congruent modulo four). Same values in the same order, so output is byte-identical; about a third faster per byte. The control arm of its measurement.
ORANGU_NO_DEVICE_TOPK unset (device top-k on) Set to have a sampled decode step (temperature above zero, no grammar, top_k in 1..=64) read the whole vocabulary’s logits back and scan them on the host, as it originally did, instead of having the device apply the repeat penalty and return the sampler’s top_k candidates. The host still applies temperature, top_p, min_p and draws — from the same candidates in the same order, so the distribution is unchanged. The control arm of its measurement (+15% at temperature 0.8 on the model this was measured on; the client sends no temperature, so this is the path a review takes).
ORANGU_COPY_INPUTS unset (direct on) Set to have every fused decode layer copy its per-call inputs — the previous layer’s output, attention’s output, the per-layer-embedding slice — into its own buffers before reading them, as it originally did, instead of binding them where they already are. Each copy is a command outside any compute pass, so each also split the layer’s pass. The control arm of its measurement (+5% on the model this was measured on).
ORANGU_NO_BLOCK_HOISTED unset (block-hoisted on) Set to disable the block-hoisted decode kernel and fall back to the scalar per-element reduce for every type that has one. The block-unroll above is fast and covers only Q4_K/Q5_K/Q6_K (it hardcodes their 256-element super-block as a fixed 4x64 geometry); the element-wise reduce covers everything and re-decodes each block’s header once per element, which for a 32-element block is 32 times over. The block-hoisted kernel is the third algorithm: several adjacent lanes share one block and take contiguous byte slices of it, so the header is amortized while the coalesced access shape is kept. It covers every remaining quantized type — Q4_0, Q4_1, Q5_0, Q5_1, Q8_0, Q2_K, Q3_K, IQ1_S, IQ1_M, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S, IQ4_NL, IQ4_XS — and is the control for any A/B of it. Term-for-term identical arithmetic to the element-wise path but a different summation order, so output is not bit-identical and a greedy argmax can flip on a near-tie. Ask a running server which kernel a type resolves to: /props -> gpu.kernels.decode.
ORANGU_BLOCK_DOT_WORDS on (0 turns it off) The block-hoisted kernel’s per-type block_dot for Q2_K, Q3_K, IQ4_XS, IQ3_S and IQ2_S as word reads: a lane’s sixteen elements taken as whole words (one qs word and one hmask word for the two K-quants, the four 2-bit fields of each byte being four sub-blocks’ worth), four activations per load, in straight-line code — where the byte-reading forms decode one byte at a time in a rolled loop, sixteen word loads with a shift and mask each and sixteen scalar activation loads, in order. Measured on the E2B UD-Q2_K_XL file the decode step’s device time went from 36 to 15.5 ms (24.8 → 50.5 tok/s); output differs from the byte-reading forms only in summation order. 0 is the control arm, and takes the wide pipelines below with it.
ORANGU_BLOCKS_PER_TRIP 1 (1, 2 or 4) Blocks a lane of the block-hoisted decode kernel takes per trip round its block loop, their loads issued together. Measured over the kernel’s own hot weights in isolation, two halves the FFN gate’s time; inside a decode step the same op loses 28% at two and 130% at four, so the default is one and the knob is for measurement.
ORANGU_KQ_LIGHT unset (off) The word-reading types’ decode mat-vec on the Q4_K light kernel’s skeleton (32 lanes, sixteen per super-block, two in flight) instead of the block-hoisted one. Measured level with it inside a decode step — −4% on the down projection, +6% on the gate, the token unchanged — which is what it was built to find out.
ORANGU_Q2K_STUB unset (loads or alu) A wrong Q2_K decode routine for splitting the kernel’s time: loads does every load with the arithmetic reduced to a sum of the words, alu the full arithmetic on words made up from the lane with no weight loads at all. Timing only; the start-up check retires a wrong kernel, so pair it with ORANGU_VK_NO_SELFCHECK=1. Measured on the E2B UD-Q2_K_XL FFN gate: 66 µs real, 49.5 loads only, 65.6 arithmetic only — the kernel is bound by its arithmetic, not the weight stream.
ORANGU_DECODE_MMVQ on where the device accelerates the packed 8-bit dot (0 turns it off) The decode mat-vec of Q2_K, Q3_K, IQ4_XS, IQ3_S, IQ2_S and the legacy Q4_0, Q4_1, Q5_0, Q5_1, Q8_0 on the integer dot: the weight bytes as stored against an 8-bit activation through dot4I8Packed. In the fused decode chain the activation is quantized by the kernel that produces it — the post-attention norm pair writes the FFN input’s 8-bit form beside the float one for the gate and up projections, the fused activation·multiply writes the product’s for the down projection, and the attention norm writes its output’s for the Q/K/V projections — so no quantize dispatch and no copy is added and every fusion stays. That is the difference from ORANGU_Q4K_MMVQ‘s form, a quantize dispatch per projection and the copy-in input path, which lost 3 ms a token to its dispatches and the fusions it gives up. Measured on the E2B UD-Q2_K_XL file: gate and up 65 → 43 µs a layer, down 100 → 63, wq 32 → 28 (the norms’ epilogues cost 0.5 and 1.7 µs), the token’s device time 14.8 → 11.9 ms; on the Q4_1 file 18.7 → 13.7 ms (gate/up 112 → 59), on Q8_0 19.6 → 18.4. The output projection and the per-layer-embedding ones stay on the float kernels. Output differs from the float kernels by the 8-bit rounding of the activation, as the prefill’s integer-dot GEMM does; 0 is the control arm and the float path’s numbers. The CPU-orchestrated matmul batch quantizes per op itself.
ORANGU_Q4K_I8 on where ORANGU_DECODE_MMVQ is (0 turns it off) Q4_K’s decode mat-vec on the integer dot (vulkan_shaders::shader_source_q4k_i8), for projections of 512 rows or more: the fused chain’s 8-bit activation (the same the types above read) against the weight nibbles through dot4I8Packed, each super-block’s 6-bit scales and minimums applied once to the integer sums — on the skeleton of the ternary integer-dot kernel, 64 lanes, eight to a super-block, eight super-blocks in flight, the activation words loaded once and reused by every row of the workgroup. Measured on the CIX P1’s Mali-G720 with gemma-4-E2B at Q4_K_M: ffn_gate and ffn_up 368 → 265 µs a layer, ffn_down 321 → 284, attn_q 100 → 88; narrower projections (the 256-row K and V) measured slower on it and keep the float kernel. Checked at start-up against the float kernel at the 8-bit rounding bound (decode_kernel_agrees), as the other integer-dot kernels are. 0 is the control arm.
ORANGU_Q4K_I8_ROWS 4 (1..=16) Rows per run of ORANGU_Q4K_I8’s kernel: the activation words of a block are loaded once and reused by every row.
ORANGU_Q4K_I8_GROUPS 4 (1..=16) Runs of ORANGU_Q4K_I8_ROWS rows one workgroup takes in sequence. Swept on the Mali-G720 for ffn_gate: 4 × 4 265 µs, 4 × 8 274, 8 × 4 292, 8 × 2 308, 2 × 8 303, 8 × 1 399, 16 × 1 453, 1 × 16 453.
ORANGU_BLOCK_HOISTED_WIDE on (0 turns it off) For the word-reading types, a second block-hoisted pipeline at four rows per workgroup, taken for a projection of 2048 rows or more: one row per workgroup leaves each workgroup too short-lived for its fixed cost on a wide projection, while on the 256-wide per-layer-embedding projections four rows leave too few workgroups (0.41 → 1.58 ms per token), hence the split by width rather than ORANGU_REDUCE_N_ROWS for everything. Measured −12% on the E2B Q2_K FFN gate, neutral elsewhere.
ORANGU_TERNARY_IDOT on where the device executes dot4I8Packed natively (0 turns it off, 1 insists) The integer-dot decode kernels for Prism’s ternary types (PQ2_0, PTQ1_0), taken ahead of the block-hoisted words kernel for a projection of 2048 rows or more. The activations are quantized to int8 per 128-block inside the kernel (a subgroupShuffleXor block maximum, pack4x8snorm), the trits packed four to a word and dotted with dot4I8Packed, the -1 of a 0/1/2 trit one Σq subtraction per block. On the Mali-G720 PTQ1_0 7.5 → 16 GB/s at the FFN gate shape, 24 at the down shape; the 27B’s Vulkan decode 0.98 → 1.54 tok/s. Off, the float words kernel is the control arm; the cross-checks compare each against its own activation rounding.
ORANGU_TERNARY_I8 on (0 turns it off) With ORANGU_TERNARY_IDOT, the integer-dot ternary kernels read the activation already quantized to int8 by a dispatch per token over the norm’s or the previous kernel’s output (shader_source_quantize_q8’s per-32 layout, the one the block_hoisted_i8 kernels read), rather than quantizing it inside the kernel per row group. On the Mali-G720 the PQ2_0 FFN projections are 1.3× faster for it (gate 901 → 698 µs, down 708 → 535; the 27B’s decode 3.28 → 3.64 tok/s); PTQ1_0, bound by its trit remainder chain, gains nothing. Off, the kernels take the f32 input — the control arm.
ORANGU_TERNARY_REPACK on (0 turns it off) On a Vulkan device, a PTQ1_0 file’s weights are repacked as PQ2_0 when the model loads (QuantMatrix::repacked_two_bit): the same trits and scales, two bits an element instead of base-3 packing — lossless, 21% more bytes on the device (the 27B: 5.5 → 6.7 GiB), a few seconds at load. The device’s PTQ1_0 kernel is bound by the trit unpacking where its PQ2_0 kernel is a shift and a mask, so a PTQ1_0 file then decodes at the PQ2_0 rate. The banner’s Kernels line and the weights … on device figure describe the served form (pq2_0, the larger size). Off, the file’s format is served as is.
ORANGU_TERNARY_IDOT_ROWS 8 (1..=16) Rows per workgroup of the integer-dot ternary kernels: the activations are quantized once per block per workgroup and reused for every row. Measured 4 → 8 as 1387 → 1255 µs on the PTQ1_0 gate projection, 16 no better.
ORANGU_FUSED_ATTN_LAYER on (0 turns it off) The Qwen hybrid trunk’s full-attention sub-layer as one submission at decode (VulkanBackend::fused_attention_layer: the projections, per-head norms, RoPE, cache write, attention, sigmoid gate, output fold and wo). Off, the sub-layer runs as a batched projection submission, the host’s norms/RoPE/attention, and a wo submission. Measured level at depth 0 and +3% at depth 512 on the 27B ternary model; the control arm.
ORANGU_HYBRID_FLUSH_LAYERS 8 (0 for one submission) How many layers of the Qwen hybrid trunk’s device token go into one submission: the device runs the layers already submitted while the host records and finishes the next, where one submission per token left it idle through the encoder’s finish (~7 µs a dispatch, 14 ms for the 27B’s token on the CIX P1). On that board PQ2_0 decode 3.64 → 3.86 tok/s at 8; 16 and 4 are within 1% of it.
ORANGU_HYBRID_DEVICE_TOKEN on (0 turns it off) The Qwen hybrid trunk’s decode token recorded on the device (Trunk::forward_token_on_device): the residual stream, every residual add fused with the norm that follows it, each sub-layer’s input fold fused with its quantize, the recurrent and attention sub-layers, the dense FFN (its activation, fold and quantize one dispatch), the final norm and the output head — in submissions of ORANGU_HYBRID_FLUSH_LAYERS layers, one readback of the logits. Off, each sub-layer and FFN is its own submission with the norms and adds on the host. Bit-identical output; the control arm.
ORANGU_HYBRID_PREFILL_CPU unset (1 turns it on) [orangu-server].prefill_backend = cpu from the environment (it overrides device and auto): the Gemma and Qwen hybrid trunks run a multi-token pass on the CPU backend and single-token decode on the selected device. On the CIX P1 with the 27B ternary model: prefill 1.6 → 10.3 tok/s at 64 tokens (11.3 at 256), decode unchanged. The attention caches and the recurrent state move between the two through the accessors that already keep one copy current.
ORANGU_DECODE_PROBE on (0 turns it off) Under backend = auto on one GPU, time a few one-token decode steps of the model on the device and on the CPU backend (the same mapped weights) at load, and decode on the cores only when they are at least 1.1× faster (engine::decode_backend). 0 keeps the device without measuring, as before.
ORANGU_DECODE_PROBE_MIN_GAIN 1.1 (a positive number) How much faster the cores must decode for ORANGU_DECODE_PROBE to move decode to them. A measurement knob: a value below 1 makes a slower CPU win, which is how the switch itself is exercised on a machine whose GPU is the faster (on the CIX P1, 0.1 served the model from the CPU at backend = cpu’s 8.3 tok/s).
ORANGU_HEAD_SPLIT on (0 turns it off) When decode moved to the cores and left a GPU idle, measure sharing the output head’s rows with a fresh backend on that GPU, and keep the split only if whole decode steps get ≥ 1.05× faster (engine::head_split). 0 keeps the head on the cores and skips the ~1 s measurement.
ORANGU_HEAD_SPLIT_SHARE unset (a number in 0..1) The device’s share of the head’s rows for ORANGU_HEAD_SPLIT, kept whatever the measurement says. A measurement knob.
ORANGU_SDOT_K_ROWS on (0 turns it off) On aarch64 with sdot, decode’s one-token Q4_K/Q6_K row dot takes the super-block-wide kernels (vecdot::dot_k_row_q4_k_sdot/_q6_k_sdot; bit-identical, ~1.5× per core). 0 keeps the per-32 form, for measuring (doc/PERF-ALL.md, task 15).
ORANGU_CPU_POOLS on (0 turns it off) With the default worker pool, the pools a request runs inside (engine::cpu_pools): every core, and on a machine with big and little cores the big ones, pinned — each CPU phase timed in both after load (prompts on the model’s widest FFN GEMM at 256 tokens, decode on its one-token steps) and run in the faster. Running inside a pool matters by itself: a request on a tokio thread hands each parallel region to the global pool from outside, and on the CIX P1 decode on the same eight cores went 12.6 → 17.5 tok/s from that alone; on the big cores rather than all twelve, 8.2 → 17.7. 0 runs requests on the global pool, as before. Not built when threads or ORANGU_EXPERT_BIG_CORES names the pool.
ORANGU_ATTENTION_MIXED on where the CPU has i8mm and bf16 (0 turns it off) A prompt’s attention on the CPU backend (a pass of at least ORANGU_ATTENTION_TILED_MIN_TOKENS tokens) in blocks of 32 rows — query tokens times the heads of one K/V group — with the scores in int8 on smmla (keys less their per-channel mean, quantized per row: the SageAttention recipe the picture transformers use) and the value product in bf16 on bfmmla, each block against only the keys its windows cover. Measured on the CIX P1 at gemma-4-E2B’s shapes over 2048 tokens: 5.4× the one-query loop on the full-attention layers, 3.9× on the sliding ones; prompts 127 → 154 tok/s at 2048 tokens and 115 → 147 at 4096. Within int8/bf16 rounding of the f32 loop, not bit-identical to it. 0 falls back to ORANGU_ATTENTION_TILED’s f32 tiles.
ORANGU_ATTENTION_TILED on (0 turns it off) Where ORANGU_ATTENTION_MIXED cannot run (no i8mm/bf16, or turned off): a prompt’s CPU attention in f32 tiles — a block of queries and their K/V group against a block of keys, online softmax. Within float reassociation of the one-query loop; about 1.06× it on the CIX P1, where that loop was already compute-bound. 0 leaves the one-query loop for every pass.
ORANGU_ATTENTION_TILED_MIN_TOKENS 32 (integer) The narrowest pass the two blocked attention kernels above take; narrower passes, and decode, keep the one-query loop, which spreads over the pool by token.
ORANGU_PROMPT_WEIGHTS unset ([orangu-server].prompt_weights, auto) auto, copy or file over the configuration for one run: whether the weights a prompt multiplies on the CPU run from copies for the cores’ matrix instructions — per-row int8 for the K-quants (vecdot::RowI8, the smmla tile), bf16 for the unquantized matrices (vecdot::Bf16Weights, the bfmmla tile) — or from the file’s weights. auto measures each kind’s widest matrix both ways at load and needs a 10% gain and the memory; see prompt_weights in the server configuration.
ORANGU_PREFIX_WARMUP unset ([orangu-server].prefix_warmup, on) 0/off turns off, anything else on, over the configuration for one run: the record of reused prompt prefixes and their background prefill after a restart. The [adapt] prefix warm-up: line says what was warmed and how long it took.
ORANGU_TERNARY_IDOT_GROUPS 4 (1..=16) Runs of ORANGU_TERNARY_IDOT_ROWS rows one workgroup of the integer-dot ternary kernels takes in sequence, so the lane’s role constants, the launch and the drain are paid once per run of four. Measured on the Mali-G720: the PTQ1_0 FFN gate projection 1070 → 901 µs, the down projection 774 → 708; eight hurt the down projection (too few workgroups).
ORANGU_NO_DUAL_NIBBLE unset (dual on for Q4_K and Q6_K decode) Set to disable the dual decode kernels for both Q4_K and Q6_K and fall back to the two-wave block-unroll. Each two-wave kernel splits a 64-thread workgroup into two halves that re-read shared weight bytes — Q4_K streams every qs byte twice (once per nibble half), Q6_K re-reads every qh byte (once per w_lo half). The dual kernels use a 32-thread (single-subgroup) workgroup that loads each such byte once, cutting decode GPU-execution time (~22% for Q4_K, a further ~4–8% for Q6_K) with identical greedy output. They reorder the per-lane float adds, so they cross-check against the CPU backend within a tolerance rather than bit-for-bit. No effect on Q4_K when ORANGU_PACKED_DOT=1, or on either when ORANGU_NO_MLP_UNROLL=1.
ORANGU_NO_Q6K_DUAL unset (Q6_K dual on) Set to disable only the Q6_K dual kernel (reverting Q6_K tensors — e.g. ffn_down — to the two-wave block-unroll) while leaving the Q4_K dual kernel on. For A/B isolation of the Q6_K kernel; ORANGU_NO_DUAL_NIBBLE=1 disables both.
ORANGU_Q4K_CONTIG unset (off) Set to select an alternative Q4_K decode kernel with a contiguous thread→element mapping: each lane loads a u32 of four consecutive qs bytes (all used) plus two vec4<f32> activations, ~3× fewer VMEM load instructions than the default dual kernel. Correctness-verified (byte-identical greedy output) but measured no faster on this hardware — the Q4_K matmul is memory-latency-bound rather than load-issue-bound — so it is off by default. Kept for other GPUs where issue rate may bind.
ORANGU_REDUCE_N_ROWS 2 (integer, not a presence flag) Output rows one decode matmul-vec workgroup computes (reusing each activation element across all of them). Lower values launch more, smaller workgroups — more independent wavefronts in flight per compute unit to hide VRAM latency — at the cost of re-reading each activation in more workgroups. Clamped 1..=16. Was 4 historically; re-swept to 2 after the dual-nibble kernel and chunked submission made GPU occupancy (not CPU submission cost) the decode critical path. Applies to every K-quant reduce/block-unroll kernel and its dispatch-count math together.
ORANGU_NORM_WG 128 (must be 64, 128, or 256) Workgroup size (thread count) of the default tree-reduce RMSNorm and RMSNorm+residual-add kernels. These run one dispatch_workgroups(1,1,1) workgroup over the whole n_embd row, so they are occupancy-starved; more threads shorten each thread’s grid-stride loop and light up more of one work-group processor’s SIMDs. Raised from 64 to 128 after measurement (halving to 32 had previously doubled the time, so these are compute/load-bound, not launch-bound): decode GPU-execution time dropped ~11% (35.4 → 31.5 ms/token) on real E2B/RX 5500M with byte-identical output; 256 was no better than 128 (deeper reduction tree, more barriers). Only affects the default (non-subgroup) norm path, and only rows the wide kernel below does not take.
ORANGU_NORM_WIDE on (0 turns it off) The whole-row RMSNorm family (norm, norm+residual, norm+residual×scale) on its wide kernel: vec4 loads, every thread’s loads issued straight-line and waited for once, held in registers for the rescale, a two-round shared-memory reduce. Taken for any row width that is a multiple of four; other widths keep the grid-stride kernel ORANGU_NORM_WG sizes. The grid-stride kernel spent its time in a serial chain of one-at-a-time loads, and measured 13 µs a dispatch where the wide one measures 3.5 — a decode layer runs five of these. 0 is the control arm of a sweep.
ORANGU_SPLIT_PASSES unset (off) Records a fused decode layer the old way, one compute pass per recording function (about ten per layer) instead of one pass per layer. A pass boundary costs the device several microseconds on top of the dispatch it precedes, and the merge is worth about a quarter of decode throughput on the model this was measured on; this flag exists so that number can be re-measured. Inside a pass the driver still orders dependent dispatches, so the two are the same computation.
ORANGU_LAYER_PASSES unset (off) Records each fused decode layer in its own compute pass instead of carrying one pass across a whole submission chunk of layers. The first dispatch of every layer used to measure about 6 µs more than the same kernel elsewhere — the cost of the pass boundary it followed — and nothing between two layers of a chunk needs the encoder. 1 is the control arm of a sweep.
ORANGU_HEAD_WIDE on (0 turns it off) The per-head kernels after the Q/K/V projections on their wide forms: the Q norm+RoPE with the head loaded straight-line into vec4 registers and a two-round shared reduce, and one dispatch for the whole key/value epilogue — K’s norm and RoPE, V’s norm, and both rows’ write into the KV mirror — where four dispatches ran before. Taken for heads up to 512 wide and a multiple of four, on an f16 or f32 KV cache; a q8_0 cache, a layer that borrows its V from K, and projection biases keep the separate dispatches. Same arithmetic in the same order for the rotation; output is byte-identical on the served model. 0 is the control arm of a sweep.
ORANGU_NORM_PAIRS on (0 turns it off) The attention post-norm with its residual add and the FFN norm that reads its result as one dispatch, wherever the wide norms apply: the first result stays in registers for the second reduce and is stored once, for the FFN’s own residual. 0 runs the two whole-row dispatches, the control arm of a sweep.
ORANGU_ADD_NORM_PAIR on (0 turns it off) The norm pair on a layer without an attention post-norm (the llama family, phi3, mistral3): the residual add and the FFN norm as one dispatch, with the same 8-bit epilogue for the gate and up projections. Measured level with the two dispatches on SmolLM2-360M and Llama-3.2-1B (a merged dispatch is a few microseconds; the token is milliseconds) — kept for the count, and 0 is the control arm.
ORANGU_FFN_GATE_UP on (0 turns it off) A layer’s gate and up projections as one decode dispatch over [n_embd, 2 * n_ff] where the checkpoint stores the two tensors adjacently with one type (ffn_gate then ffn_up, which a name-ordered file does; phi3’s single ffn_up tensor always) — a view over the bytes both already occupy, placed on the device once with the halves bound as sub-ranges of it, so nothing is uploaded twice. One dispatch’s fixed cost instead of two, and a longer dispatch on a small model whose [960, 2560] matrices could not fill the card. Output is unchanged: the rows are the same rows in the same order. 0 runs the two dispatches, the control arm.
ORANGU_BUSY_POLL 1 (spin; 0 block, auto predicted) How the decode thread passes the wait for a token’s readback. 0 blocks on the fence for the whole wait: the thread is parked while the card runs, its core drops into a deep idle state (hundreds of microseconds to leave on a laptop part), and the fence’s completion reaches it through an interrupt and a scheduler wake-up — so the next token’s recording starts late, and the card idles between tokens long enough for its clock governor to settle below its top state. 1, the default, spins on a non-blocking poll for the whole wait: back within microseconds of the card finishing, the core hot, at the price of one core busy for the whole wait. auto blocks for most of the wait as the last waits of the same kind measured it and spins through the rest (ORANGU_BUSY_POLL_BLOCK, the blocked share in percent, default 75) — measured at about half the spin’s gain for a quarter of its core time, since the core kept awake through the wait is most of the effect. Measured on real RX 5500M at depth 0, the spin against the block: SmolLM2-360M 131 → 171 tok/s, Llama-3.2-1B 125 → 132, 3B 58 → 60, E2B 72 → 74; level on the 26B-A4B and Nemotron, whose experts run on the host. 0 is the choice for a machine that serves other work beside the model. ORANGU_WAIT_TRACE=1 prints every 32nd wait: its length, the prediction, the blocked part and the polls spun.
ORANGU_ATTN_REDUCE_NARROW unset (off) Runs the split-k attention merge on its original 64-thread kernel, a rolled loop of one dependent load per partial per element, instead of the 256-thread one that issues its loads four at a time and stages the softmax weights in shared memory once. The merge is the same arithmetic in the same order; the flag is the control arm of its measurement (+3% at short context, +10% at depth 1024 on the model this was measured on).
ORANGU_DELTA_HEADS_PER_TASK derived (heads / 4, at least 1) Heads handed to one task by the delta rule’s fan-out. A head’s work is one head_dim × head_dim update, and at one head per task the hand-off is a real fraction of it — profiling a prompt on a delta-net mixture puts about a ninth of all CPU in the kernel scheduler and in crossbeam’s steal and epoch paths. Grouping amortizes that over more arithmetic without changing any of it: the tasks stay disjoint and each does as many heads as this says. The default is a floor on the task count, not a match to the worker count — four tasks across sixteen workers measured faster than sixteen, so the pool is not what limits this. A layer with four heads or fewer is never grouped, since grouping it would leave cores idle for no dispatch saved. Worth +1.4% on the one model here that exercises it, four of four alternating pairs. Set a number to override.
ORANGU_ATTN_TOKEN_TILE 8 Prompt tokens sharing one pass over the key buffer in the host’s batched causal attention. The key traffic divides by this directly, but the tile holds tile × heads × key_dim of queries and as much again of accumulator, and both must stay resident across the sweep over keys or the saving is handed back. Measured on a delta-net mixture: 4 is level with 8, 16 is 1.8% behind, and 32 loses 29% as the resident set stops fitting a core’s L2. Eight sits one doubling below that cliff, and 4 costing nothing says the kernel is not traffic-bound at this width — so this is a knob for a different architecture’s geometry, not a tuning dial for this one.
ORANGU_ATTN_SPLIT_K 8 for the classic kernel, 32 for the cooperative one, 64 for the cooperative one at a 16-lane subgroup (a power of two, 1..=64) Split-k factor for decode attention: how many workgroups each query head’s KV-position range is split across (n_head × k_num phase-1 workgroups, merged by a phase-2 pass). Each head attends a KV range that grows with context, so more splits expose more parallelism the longer a generation runs, at the cost of more phase-2 merge overhead when the range is short. Raised from 4 to 8 after re-sweeping in the full decode chain at real context lengths (the earlier sweep was on the isolated dispatch at short context): on real E2B/RX 5500M, k_num=8 cut per-token GPU time at ~245 tokens of context from 35.6 to 32.4 ms while staying neutral at short context, and had the best end-to-end throughput of {4,8,16}; 16 wins only once context is very long. Byte-identical output across values. Workloads dominated by very long contexts may prefer 16. Pure runtime uniform — no shader rebuild. 64 is a hard ceiling: the wide split merge holds a split’s running max and sum in 64-entry workgroup arrays, and at 128 every head past the first read another’s partials (NaN).
ORANGU_ATTN_COOP16 on where the device’s subgroups are narrower than 32 (0 turns it off) Decode attention on the cooperative kernel (shader_source_attention_split_coop) at the subgroup’s own width: a workgroup of one subgroup, each lane owning every 16th dim of the head so K and V are read coalesced and the score is one subgroupAdd — where the 32-lane cooperative kernels are not built (the Mali-G720’s subgroups are 16 wide) and decode fell back to the classic split kernel’s serial per-thread dot. Measured on the CIX P1 with gemma-4-E2B: decode attention 4.4 → 1.1–1.8 ms a token at depth 0 and 27.8 → 7.3 at depth 2048. Prefill keeps its own kernels. 0 is the control arm.
ORANGU_PACKED_DOT unset (off) Dequantizes Q4_K weight elements in pairs and accumulates the dot product as vec2<f16> instead of two scalar f32 multiplies. Requires an adapter with WGSL f16 support. When set together with the block-unroll, selects the combined unroll+packed Q4_K decode kernel.
ORANGU_WIDE_LOAD unset (off) Binds the weight buffer as array<vec4<u32>> (16-byte reads) instead of array<u32> (byte-wise reads), consolidating each Q4_K/Q5_K block header into one 16-byte read. Covers all supported quant types.
ORANGU_NO_KV_F16 unset (f16 on when the adapter supports it) 1 disables storing the per-request KV-cache GPU mirror as f16 and fall back to f32. f16 (the default on an adapter with WGSL f16 support) halves KV-read memory traffic per attention dispatch, with a per-write cast, and matches the ecosystem’s default KV cache type.
ORANGU_KV_Q8_0 unset (off) 1 stores the per-request KV-cache GPU mirror as q8_0 (8-bit block-quantized) instead of f16, dequantized inline in the attention shader. Halves KV-read bytes again vs f16, directly cutting attention’s cost at long context — measured −32% attention GPU time at ~295 tokens on real E2B/RX 5500M (5.87 → 3.99 ms), the saving growing with context, at a slight cost at short context (per-write quantize overhead). Takes precedence over f16. Lossy (unlike f16), so off by default; the recommended lever for long-context / long-generation workloads, where per-token decode slows as the KV cache grows. Superseded by [orangu-server].kv_cache, which names every value rather than one; kept because existing repro lines and sweeps use it.
ORANGU_KV_CACHE unset (the config file’s kv_cache) f16, q8_0 or f32 — the same three values [orangu-server].kv_cache takes, overriding it for one run. The variable a sweep should use: it is the only one that can name all three arms, so a sweep does not have to express “not the other one”.
ORANGU_NO_TILED_PREFILL unset (tiled prefill on) Set to disable the 16×64-output-tile GEMM for prefill (n_tokens >= 64) and fall back to the plain cooperative kernel (one workgroup per output row, looping over the whole prompt internally) — measured on real hardware to drive real requests into GPU-driver hangs at ordinary prompt lengths (~170-450 tokens) and, even where both complete, ~10x slower. Not recommended; kept for A/B comparison.
ORANGU_COOP_MIN_TOKENS 64 (integer, >= 1) The token count at or above which a matmul takes the weight-amortizing tiled-GEMM path instead of the reduce family. The reduce kernels dispatch n_row_groups × n_tokens workgroups, so each token re-streams the whole weight matrix (their cost grows with n_tokens); the tiled kernel loads a 16×in_dim weight tile into shared memory once and reuses it across a 64-wide token tile (cost ~flat in n_tokens). Below the threshold the per-token re-stream is cheaper because it keeps far more wavefronts in flight to hide this GPU’s weight-load latency. Lowering it to route the “missing middle” (K ≈ 2…63, e.g. speculative or small-chunk prefill) onto the tiled path was measured slower on RX 5500M (this matmul is latency-bound, not bandwidth-bound — see SERVER_ROADMAP.md Step 13), so the default keeps that split at 64; exposed as a knob for GPUs where the crossover sits lower and as the A/B harness for a future small-K kernel. 1 forces every matmul tiled; a value past the longest batch keeps everything on the reduce path.
ORANGU_NO_GPU_SAMPLE unset (GPU sampling on) Set to disable running greedy (temperature-0) argmax sampling with repeat penalty on the GPU in the same submission as the forward pass (reading back one token id instead of the full [n_vocab] logits vector) and fall back to a CPU-side readback + sample.
ORANGU_DECODE_CHUNKS unset (integer; set, 3 equal chunks) How many queue.submit() calls one decode step’s layer loop is split across. 1 records the whole token and submits once (the historical behaviour); > 1 submits the first chunks - 1 groups of layers as soon as they are recorded, so the GPU starts executing them while the CPU is still recording and validating the later ones — overlapping the CPU-side submission cost with GPU execution instead of serialising it. Clamped to 1..=n_layers. On real E2B/RX 5500M this raised decode throughput from 14.4 tok/s (1) to 18.8 tok/s (7, the default), with byte-identical output; 35 (one submit per layer) reaches 19.4 tok/s but adds per-submission overhead for a marginal gain. Unset, the split is not equal: the plan is the probed geometric split ORANGU_DECODE_FIRST_CHUNK describes; setting this pins equal chunks.
ORANGU_DECODE_FIRST_CHUNK probed (integer; 0 = equal chunks) Layers in a decode step’s first submission; each following chunk is ORANGU_DECODE_CHUNK_RATIO times the last. Only the recording ahead of the first submission is host time the device cannot hide, so the default plan carries two layers first and grows each chunk threefold so the submissions stay few. Two rather than one: a single-layer first burst after a pause, when the card has parked its clocks, was measured holding a whole request at a third of its rate; two layers never did. The split applies to every family’s fused decode chain (the llama family, phi3, mistral3, muse-glimmer and gemma). Unset, the plan is probed: the server times its decode steps on three plans — the default, the same with each chunk twice the last, and equal chunks of four layers — over three rounds of a block of steps per plan (after 32 settling steps), and keeps the plan that won most rounds by 5%; the first answer is asked again after 1024 steps, later ones every 4096. The plans differ per model with no rule read off the file: a 360M model ran +11% on the twofold plan and +20% in fours, a 2B +13% twofold, a 1B −7% twofold. Setting this, ORANGU_DECODE_CHUNKS or ORANGU_DECODE_CHUNK_RATIO pins the plan and stops the probe. ORANGU_CHUNK_TRACE=1 prints each round’s medians and the vote.
ORANGU_DECODE_CHUNK_RATIO probed (integer ≥ 1) The growth of the decode step’s submissions after the first: with a first chunk of two layers, 3 submits a 32-layer token as 2, 6, 18 and 6 layers; 2 as 2, 4, 8, 16 and 2. A larger ratio submits fewer times, a smaller one keeps the device closer behind the host’s recording. Setting it pins the plan (see ORANGU_DECODE_FIRST_CHUNK).
ORANGU_PREFILL_FUSED_ATTN on (0 turns it off) A prefill layer’s whole pre-attention half as one recording — the Q/K/V projections, the per-head Q/K norms, RoPE, V’s weightless norm, the KV-cache write for the whole batch, and attention itself — with nothing between the projections and attention reaching the CPU, and attention’s output left on the device for the post-attention chain. 0 takes the step-by-step path (projections read back to the CPU, norms and RoPE on the CPU, a per-token cache push, then a separate attention dispatch) and with it the whole host-orchestrated prefill: every layer’s result comes back to the CPU. It was off for a while — measured as a loss against a prefill that computed only part of its output, then unable to run under the paged KV cache at all (a layer reading a donor’s cache asked for pages past the sequence’s reservation) — and is on now that both are fixed: 163/176/182 → 251/303/354 tok/s at 128/512/2048 prompt tokens with the device-resident stream below, output byte-identical.
ORANGU_MAX_TOKENS_PER_SUBMISSION 512 on the Vulkan backend, 256 elsewhere (integer) The token stripe of one prefill submission: a chunk wider than this runs its chains in stripes of it. Raised to 512 for wgpu once the integer-dot GEMMs’ 128-token tiles existed (the same kernel at 256 tokens runs at three quarters of its 512-token rate) and every element-wise dispatch in the chains had become grid-stride, so nothing bounds the stripe but memory; before that a stripe was capped by the device’s workgroups-per-dimension limit on the flat activation dispatch, which on the layers with a 12288-wide FFN cut a 512 stripe into 320 + 192.
ORANGU_LLAMA_RESIDENT_PREFILL on (0 turns it off) On the llama family (Llama, Qwen2, Granite and the others that architecture serves), on phi3 and on mistral3, whether a prefill chunk keeps its residual stream on the device: every layer’s attention chain and post-attention chain recorded back to back, a few layers per submission, the K/V rows read back once per chunk and, for a chunk before the last, only after the next chunk’s chains are submitted. The step-by-step path brought attention’s output and the residual home at every layer, two readbacks and a host turn each with the device idle through them; measured on a 1B file this is 1068 → 1938 tok/s at 2238 tokens and 736 → 1263 at 156. Only whole-device runs take it (every layer on the first card, no NPU tier, no activation dump); 0 is the control arm. ORANGU_RESIDENT_PREFILL_CHECK=1 runs every chunk on both paths and logs how far apart their residuals are (the root-mean-square difference over the chunk and the worst row): the two differ by the placement of their 8-bit activation roundings, ~1.5% RMS on the files measured, and a wrong parameter shows as whole rows off.
ORANGU_PREFILL_LAYERS_PER_SUBMIT 4 (integer) While the residual stream stays on the device, how many layers’ chains are recorded into one submission. Each submission costs the host about a millisecond of driver validation and leaves the device idle for about as long, and a layer makes three chains; but a single submission for a whole chunk makes the device wait while the host records everything first, so 35 measured slower than 4 (311 vs 334 tok/s at 2048 tokens). 0 submits every chain on its own, the control arm; 1 is one per layer.
ORANGU_PREFILL_KV_WAIT off (1 turns it on) While the residual stream stays on the device, a KV-owning layer’s attention chain no longer waits for its key/value rows to come back to the host before the next layer is recorded: the positions are counted into the cache at once (the device pages already hold the values, which is all the next stripe and the next layer read) and the rows land in one mappable stage per chunk, filled after the chunk with a single wait. The host copy is what prefix reuse and slot save read, and it is complete by the time either can run. The per-layer wait had been the one place the device-resident prefill still drained the device: on a 42-layer model with 24 KV-owning layers it was 335 ms of a 158-token chunk. 1 restores the wait per layer, the control arm.
ORANGU_PREFILL_STREAM_DENSE on (0 turns it off) Whether a dense gemma layer without a per-layer-embedding stage keeps the residual stream on the device: its post-attention chain writes the layer’s result — output scale included, folded into the FFN post-norm’s kernel — straight into the other stream buffer, as the per-layer-embedding chain does for a layer that has one. Before, the stream only held a layer with that stage; every other layer (the embedding model’s, the 31B’s) landed its rows on the host and took its chains one waited-for call at a time. Measured 2× on embeddinggemma at 64–128 tokens; a bidirectional batch above the stripe width still takes the host path (the fused attention chain declines it), so 256 tokens and up are unchanged. 0 restores the per-layer-embedding requirement, the control arm.
ORANGU_HOST_STREAM_TOKENS 32 (integer; 0 turns it off) On a model split with the CPU as a device, the narrowest batch at which a host layer’s projection is sent to the first card rather than run on the CPU. A host layer’s weights are already in RAM, and at one token the card would spend longer fetching them than multiplying — but a prefill batch multiplies the same weights by every token, so above this width the layer’s Q/K/V, gate/up and down projections upload their weights into the card’s bounded streaming region (the same region routed experts stream through, ORANGU_EXPERT_STREAM_BYTES), run the integer-dot GEMM there, and read the rows back to the host exactly as the CPU’s product would arrive. Nothing else about the layer moves: its attention and its KV cache stay on the host, and the card’s residency plan is untouched since the region recycles. A batch wider than one submission’s stripe (ORANGU_MAX_TOKENS_PER_SUBMISSION) streams stripe by stripe with the weights uploaded once. Decode never takes it, whatever the width — a decode step’s numerics must not depend on how many sequences share it. Output differs from the CPU kernel’s by the activation rounding, and is closer to the float product: on real weights and activations of the 12B model’s first host layer, the card’s error against the dequantized product is 0.4–0.6 % where the CPU kernel’s is 0.9–1.5 %. Measured on gemma-4-12B split 18:30 over a 4 GiB card and the host, a 512-token chunk’s host layer went from 1045 ms (gate/up 586, down 258, Q/K/V 182, attention 19) to 187 ms (83, 56, 29, 19) against 115 for a resident layer’s whole fused chain; the prompt went from 12.5 to 49 tok/s at 574 tokens, the 31B from 3.7 to 16. 0 keeps every host layer on the host, the control arm.
ORANGU_EXPERT_GEMM_TOKENS 256 (integer; 0 turns it off) The narrowest prefill batch at which a mixture-of-experts layer’s routed experts run on the card, as one submission per layer: the layer’s whole expert stack (every expert’s gate/up rows as one tensor, ~270 MiB on a 128-expert model) crosses the bus into the streaming region, every expert is multiplied by the rows routed to it in a single indexed integer-dot GEMM dispatch per projection (a table names, per token tile, which expert’s rows and which activation rows; the activations are quantized once for the layer and never gathered per expert), the activation and multiply run on the card, the product is quantized there, the down stack is copied over the gate/up one inside the same submission and multiplied the same way, and the rows are weighted by their routing weights and summed per token on the card — the host reads back the layer’s [n_tokens, n_embd] result and nothing else. Per-expert output scales (a QAT file’s) are applied by the GEMM as it stores. Uploads go through a persistently mapped staging pair filled by every core and copied by the DMA engine (2.3 → 4.8 GB/s against write_buffer‘s one-thread copy), and the next layer’s stacks are staged while this layer’s submission runs. Below the width the host’s per-expert GEMM is faster: the stack crosses the bus whatever the batch, the host’s cost is per token; measured on a 26B-A4B file the two are level near 220 tokens (host 0.55 ms per token per layer; card ~120 ms per layer fixed plus ~10 ms per 128 tokens), the card 9% behind at 158 tokens and 29% ahead at 574 (49 → 63 tok/s, the reference engine at 66). The region is reserved at model load so it lands in the card’s own memory: created later, once the weights had filled a 4 GiB card, it landed in host memory and every expert dispatch read its weights across the bus at a fixed 40 ms (6 ms resident). ORANGU_EXPERT_GEMM_TOKS (16/32/64/128, default 32) is the indexed kernel’s token tile — one expert’s routed tokens are a few dozen, and the dense kernels’ 128-token tile would be three quarters padding. Output differs from the host path’s by the activation roundings, both 8-bit; a wrong tile or table is off by whole rows. 0 keeps the experts on the host at every width, the control arm. Applies to the SwiGLU/GEGLU expert families — gemma, qwen35moe/qwen3next, and every architecture on the shared SwiGLU mixture path (qwen3moe, glm-dsa, glm5next, bailingmoe3) — whose whole routed feed-forward runs as one submission; a gate-less family (nemotron_h_moe) has the same one-submission form with its squared ReLU fused on the card, but keeps its experts on the host at every width unless ORANGU_GPU_EXPERTS is set: it picks six of 128 experts per token, so a wide chunk streams 650–830 MiB of stack per layer and the device form measured level with the host GEMM, which leaves the card’s memory to the KV pool.
ORANGU_EXPERT_PACK on (0 turns it off) Whether a layer’s routed experts reach the card packed: only the experts the batch routed to, copied from the staged stack into the streaming region in stripes the region’s size, each stripe multiplied before the next is copied over it, all inside the one submission — rather than the whole stack uploaded and the region sized to hold it. A 512-token chunk routes to 70–95 of a 128-expert layer, so it moves 60–75% of the bytes, and the region is ORANGU_EXPERT_REGION_MIB instead of the largest stack; the staging pair becomes a trio so the next layer’s stack can be filled while the stripes’ copies are still in flight. Measured on a 26B-A4B file: 100 → 112 tok/s at 512 tokens, the region 333 → 128 MiB, and the card’s KV pool 1408 → 1600 tokens. Needs the mapped staging (a discrete card whose host-visible heap is known and is a window smaller than the card’s memory — on unified memory, an integrated GPU, there is no window for a half to be sized past, so the pair is not used and uploads take the belt); otherwise, or with 0, the stack goes whole.
ORANGU_EXPERT_REGION_MIB 128 (integer, MiB) The streaming region’s size when the experts are packed (ORANGU_EXPERT_PACK): one stripe of routed experts, and what the card holds for the region between prompts. Larger means fewer stripes per layer (a stripe is a copy and a dispatch per projection) and less room for the KV pool; smaller than one expert’s rows falls back to the whole stack.
ORANGU_HOST_PRELOAD on (0 turns it off) On a model whose host-resident weights (a mixture-of-experts model’s routed experts) fit in memory, reads them into the page cache at start — through the file in 8 MiB reads, in file order, twice — before the first request would otherwise bring them in by demand faults (measured at 30 MB/s on a drive that reads sequentially at 400: eleven minutes of a first prompt on a 20 GiB file). Twice because a page read once sits on the kernel’s inactive list and is the first evicted under the server’s own allocations when other files hold the cache; the second read keeps it. Logged with the bytes and the time (57 s cold for 24 GiB, 17 s mostly cached). A model larger than memory is the streaming regime and takes ORANGU_EXPERT_WILLNEED instead; the per-selection read-ahead also applies to the first pass of a resident model. 0 leaves the first request to do it.
ORANGU_EXPERT_STREAM_EARLY unset (integer, MiB) Creates the expert streaming region at start-up at this size, ahead of the weights — what ORANGU_EXPERT_GEMM_TOKENS does by itself for a model whose stacks it can size at load; for a sweep of the region’s placement on a card the model fills.
ORANGU_MOE_HEATER unset (integer, workgroups) A probe, not a setting: a spin kernel of that many workgroups after each routed feed-forward, to tell a dispatch slowed by an idle card’s clock from one slowed by its buffers’ placement (the latter was the case: the heater changed nothing, the region’s placement changed everything).
ORANGU_MOE_HEAD on (0 turns it off) On a mixture-of-experts model whose layers run on the card, the part of a layer between attention and the routed experts — the output projection, the post-attention norm and residual, the router’s norm and logits, and the shared expert’s whole MLP — runs as two submissions per layer instead of five separate calls each with its own upload and readback: the first returns the residual and the router logits the host needs to pick the experts, the second (the shared MLP) is left running while the host multiplies the routed experts, and its rows are collected afterwards. Attention’s own result stays on the card for the first submission to read. A decode step was five submissions and four host round trips per layer; this makes it two and one, and a decode token one deep in its context takes the same position-parallel attention kernel the whole-step decode chain uses rather than the prefill kernel’s one workgroup per query (0.9 → 0.3 ms a layer at a thousand positions on the page pool). 0 runs the five calls, the control arm.
ORANGU_MOE_DECODE_BATCH off (1 turns it on) At decode, runs a mixture-of-experts layer’s routed experts on the host as two parallel regions over every chosen expert’s rows (gate and up, then down) instead of one per expert. Off because it measured slower on the models tried; kept as the arm of a sweep.
ORANGU_EXPERT_REGION_RELEASE off (1 turns it on) Gives the expert streaming region back to the card after each prompt’s prefill and takes it again at the next prompt that needs it. Off now that the KV pool is sized from the driver’s budget with the region already left out: holding the region costs the pool nothing, while releasing it made its return depend on what the card had left — less than the region plus its headroom once the pool held its share, so from the second prompt on the routed experts ran on the host. The return is judged against the driver’s budget rather than the card’s raw used bytes, which count other tenants the budget excludes. Taking it again is refused, and that prompt’s routed experts run on the host, when the card has less left than the region plus ORANGU_EXPERT_REGION_HEADROOM_MIB (256): a region created into a full card lands in host memory and reads across the bus at every dispatch, and what it displaces reads that way for the rest of the process — a card filled to its last page stops moving anything back. Measured on a 26B-A4B file on a 4 GiB card: with the region held across prompts, a deep prompt after a shallow one decoded at 7.3 tok/s for the rest of the process against 10.7 with it released. 0 holds the region for the process, the control arm.
ORANGU_CLOCK_HOLD on (0 turns it off) On a model whose routed experts live in host memory, holds the card’s clock up through a decode step. Such a step alternates between the card and the host — the attention or recurrent projections and the shared expert on the card, the routed experts on the host — and each host turn leaves the card idle for a few milliseconds, after which its governor parks the clock and every device call runs at a fraction of its rate (measured: the card at 500 MHz through a whole decode, more than half the token in device calls that would take a tenth of it warm). The holder is a one-workgroup spin on a second queue of the same card — a compute-only family where the card has one, its own hardware ring, so it never sits in the working queue’s order — armed for the step and released after it. Measured on a 30B-A3B hybrid: 8.9 → 11.0 tok/s, the card at its full clock through the step and drawing ~17 W more. A dense model, whose step never leaves the card, arms nothing — nor does a split dense model, whose token is mostly the host’s row dots: measured level to 12% behind there, the holder’s power coming out of the host’s share of a laptop’s budget. 0 is the control arm.
ORANGU_CONTEXT unset (tokens) The context one request must fit on the device — the environment form of --context / [orangu-server].context; see “Fitting a context” above. Layers move to the host until the card has room for that many tokens of KV cache beside the weights it keeps.
ORANGU_KV_POOL_BYTES an eighth of the machine’s RAM, at least 2 GiB (bytes) Host memory the paged KV pool may take — its rows live in system memory in full precision, the card holds a copy in its own width. A fraction of total memory rather than of what is free, so the pool is the same size on every run of a machine and larger on a larger one; the flat two gigabytes it replaced held a 26B file to 4736 tokens on a machine whose card could serve far more. The pool is also bounded by the card (ORANGU_KV_DEVICE_BUDGET below), and the startup line names the resulting capacity.
ORANGU_MMQ_TOKS 0 (the rule; 16, 32, 64 or 128 fixes it) The token tile of the dense integer-dot GEMM: the longest of 128, 64, 32 and 16 that pads the batch by at most a quarter. The full tile made a 156-token prompt 256 rows of work in every projection; the rule takes 64 there and 128 where the batch fills it. Measured on a 3B file: 504 → 647 tok/s at 156 tokens, 761 → 831 at 572, level at 2236; a fixed 32 is 28% slower than the rule at 572.
ORANGU_MMQ_ROWS 0 (the rule; 64 or 128 fixes it) The row tile of the integer-dot prefill GEMM, the twin of ORANGU_MMQ_TOKS: the two together pin the whole tile, so a shape’s geometry can be swept in situ rather than only in the kernel probe. A shape whose output width is not a multiple of the pinned tile falls through to the narrow kernel as an unpinned call would. Measured on gemma-4-E2B-it Q4_1, the 128-row tile is 1.6× slower than the 64-row one on the down projection at 512 tokens, which is the rule’s own answer.
ORANGU_MMQ_TRACE unset (off) Prints one line per distinct shape the integer-dot prefill GEMM serves — weight type, [in_dim × out_dim], batch width, the kernel chosen and the dispatch grid — the first time each is dispatched. Answers “which kernel is this op actually running”, which a device timestamp on its own cannot.
ORANGU_KV_POOL_RESERVE_MIB 448 (MiB) Device memory kept out of the KV pool’s pages, and out of a requested context’s placement, for what a prefill allocates as it runs — a chunk’s pooled regions, the K/V readback stage, the attention scratch. A sweep knob: the prefill’s transient set is what it is, and the answer to a card at its last page is to shrink that set (every layer of a chunk recycles its scratch), not to grow the reserve.
ORANGU_KV_DEVICE_BUDGET on (0 turns it off) Sizes the KV pool’s device pages from the memory budget the driver reports for this process (what it may still allocate without moving anything to host memory): the budget less the planned weights, the streaming region and a reserve for the prefill’s transient scratch, never less than the quarter of the headroom below. On a 4 GiB card with a 1.9 GiB model that is 9344 tokens of context where the quarter gave 4736; with a 2.3 GiB model 5760 where it gave 3264; a model that fills the card keeps the quarter. 0 is the quarter alone, the control arm.
ORANGU_SWA_FULL off (1 turns it on) Keeps a full-context device copy on the sliding-window layers instead of a ring of their window — see “Windowed layers keep a window” above. Costs what the ring saves (most of a gemma-4 file’s KV bytes) and buys nothing but a comparison.
ORANGU_KV_DEVICE_SHARE 25 (percent, 1–95) The floor of the pool’s share of the device headroom, and the whole of it where the driver reports no budget. Past what the budget rule allows it is harmful, not merely wasteful: a pool that takes the card past the driver’s budget has part of the process moved into host memory and read across the bus at a fraction of the rate.
ORANGU_TAIL_MARGIN 10 (percent; 0–50) How much faster the host arm must be before the vocabulary projection leaves the card — the margin in the probe ORANGU_TAIL_HOST describes. It exists so a near-tie does not churn between arms every time the probe re-runs, and it is per-probe: the chunk-plan probe beside it uses a different value. Worth lowering only where the two arms are genuinely close and the wider margin is seen to decline a real win; measured on a delta-net mixture, lowering it to 3 changed nothing because the host arm lost every round outright.
ORANGU_TAIL_HOST probed (0 device, 1 host) Whether the vocabulary projection runs on the host after an n_embd-sized readback or on the card before an n_vocab-sized one. Unset, the server times the two: after 32 settling steps, three rounds each of a block of consecutive decode steps on the device arm then a block on the host arm; the host wins a round only when its median step is 10% faster, and takes the tail only when it wins most rounds. The first answer is asked again after 1024 steps, later ones every 4096. The blocks matter: arms alternating step by step let a host step park the card’s clock and the device step after it was timed at the parked clock, so the host won by its own presence. The settling and the vote matter too: one pair of blocks timed on a process’s first tokens saw the device arm at twice its settled step and lost to the host one start in six, and that process then ran 40% behind for its whole window. ORANGU_TAIL_TRACE=1 prints each round and the vote, and GET /props reports tail as device, host, probing or host (forced) — the probe takes about seventy steps to vote, so a short benchmark can spend its whole window alternating, and orangu-bench compares that field before and after each point to say when a rate is a blend of the two arms rather than either of them.
ORANGU_VULKAN_MEMORY_HINT blocks of at most 16 MiB (usage for 8–64, performance for 128–256, or an integer for that many MiB) The block size the Vulkan backend’s allocator packs small buffers into. A block stays with the card while one live allocation sits in it, and the chains’ per-call uploads and scratch churn: with 64 MiB blocks a 4446-token prompt on a 1.9 GiB model held 1280 MiB of them with 383 live. Every large buffer (a weight chunk, the streaming region) is its own allocation either way; the difference is what the activation arenas and the KV pool’s per-layer buffers leave reserved and empty — ~400 MiB of a 4 GiB card at the larger size, which on a model that fills the card was the difference between its streaming region fitting and the driver moving the KV pages out.
ORANGU_VRAM_REPORT unset (1 turns it on) Prints the backend’s live device allocations grouped by label, largest first, and its allocator’s memory blocks by size class with what is live in them, after the KV pool is placed and with every ORANGU_GPU_TIMESTAMPS=ops report — what is competing for the card when a kernel that read fast an hour ago reads slowly now. The weight arena also reports each chunk it opens, with the plan it was cut to.
ORANGU_KV_TRACE unset (1 turns it on) Prints a sequence’s block table each time it grows — which of the pool’s pages it holds, in order — and every per-request contiguous KV mirror the device builds, with its size and the path that asked for it. A request served from the pool builds none past the one-row scratch; a mirror of hundreds of MiB beside a pool is a request the pool could not table.
ORANGU_MUSE_FUSED on (0 turns it off) Whether a muse-glimmer model runs its decode step as one submission per device through the fused layer chain (the output norm and the logits projection behind it), and its prefill through the two fused chains. The chain takes what this architecture adds to the llama block: no rotation on a full-attention layer, a sigmoid gate on attention’s output projected from the same normed input (attn.gate, attn.gate_mul in the per-dispatch stamps), and the post-norms’ own epsilon. Measured on a model split between a card and an integrated GPU: the server’s CPU per generated token 265–309 → 21–26 ms, decode +36%; prefill of 328 tokens 1.4 → 5.0 tok/s. 0 restores the step-by-step path — one Backend call per matmul, ~200 submissions a token — the control arm.
ORANGU_PREFILL_MMQ on where the device accelerates the packed 8-bit dot (0 turns it off, 1 forces it) The prefill’s Q4_K projections on the integer-dot GEMM: the activation rows are quantized to 8 bits per 32-element block once per stage, the 4-bit weights are read as bytes, and each dot instruction does four multiply-adds — against one per element on the float kernel, which dequantizes every weight first. Taken for any Q4_K (and, see the next row, Q6_K) projection of 32 or more tokens whose width is a multiple of 256 and height of 64. Not bit-identical: the 8-bit activation changes the low bits, the way it does on every engine with this kind of kernel — a 1,563-token prompt’s greedy answer differed in wording at the first sentence while a 2,846-token one did not. 0 keeps the float kernels everywhere, the control arm and the way to get the float path’s exact numbers back. Every architecture’s prefill takes it, not only the fused chains: the generic matmul batch every CPU-orchestrated model goes through picks the integer-dot kernel for each projection it applies to and the float kernel for the rest, in the same submission. Unset, the default is per device, from the driver’s own word on whether the packed dot is a native instruction: where it is not (an integrated GPU holding a split model’s overflow layers, say) the emulated dot loses to the float kernel — measured 16–18% slower at a 128-token prefill on a 30B dense model — so that device keeps the float GEMM while the card beside it uses the integer dot. /props reports which under gpu.flags.prefill_mmq; 1 forces the kernel on regardless, for the control arm of exactly that measurement.
ORANGU_MMQ_STUB unset (off) A probe, not a setting: builds the integer-dot GEMM with its dots replaced by one xor, so the kernel’s time splits into the arithmetic and everything around it. The output is meaningless under it.
ORANGU_PREFILL_MMQ_Q6K on (0 turns it off) The Q6_K projections (ffn_down on most layers of a Q4_K_M file) on the integer-dot GEMM too. Its blocks are 210 bytes, so a row’s odd blocks start two bytes into a word and the staging reads them a word at a time and unpacks to bytes; 21% faster than the float kernel on the device clock, about 3% on a whole prefill. 0 keeps that projection on the float kernel while the Q4_K ones stay on the integer dot.
ORANGU_PREFILL_MMQ_WIDE on (0 turns it off) The integer-dot GEMMs’ wide tiles: a 128 × 128 (or 64 × 128) output tile per workgroup, the reduction staged 32 elements at a time rather than a whole 256-element super-block, and 128 (or 64) accumulators per thread — each staged byte serves 32 tokens × 4 rows where the narrow 64 × 32 tile served 2 × 8. Chosen per projection by shape (ORANGU_MMQ_WIDE_MIN_WG, default 32: the half-height tile needs that many workgroups, the full one twice as many; a projection too narrow for either keeps the narrow tile). Together with the unchecked indexing below, the Q4_K FFN GEMM at 512 tokens went from 3.4 to 5.1 TFLOP/s-equivalent on rotating weights and the Q6_K one from 1.9 to 4.0. 0 is the control arm.
ORANGU_CHECKED_MMQ unset (off) Builds the integer-dot GEMM kernels with the shader compiler’s injected array-bounds clamps. They are built without them by default, whatever ORANGU_TRUSTED_SHADERS says: every index those kernels form is bounded by construction (the grid from validated dimensions, the activation buffer padded to whole tiles, the one guarded store guarded), and the clamps — a compare and a clamp per shared-memory read on a kernel already at the register limit — cost 46% of the wide tile’s time. 1 puts them back for a bisect.
ORANGU_MMQ_WIDE_MIN_WG 32 (integer) The workgroup floor behind ORANGU_PREFILL_MMQ_WIDE’s tile choice.
(the other weight types) — The wide tiles exist for Q5_K, Q8_0, Q4_0, Q4_1, Q5_0, Q5_1, Q2_K, Q3_K, IQ4_XS, IQ3_S and IQ2_S too (vulkan_shaders::shader_source_mmq_wide_bytes: one generator, a staging routine per type, each row’s sub-block staged as 32 bytes and two scale terms so one dot loop serves them all; the types whose scales cover sixteen values — Q2_K, Q3_K, IQ2_S — stage two pairs and take the dot per half; the lattice types look their bytes up in the codebook buffer at staging), built on first use per type rather than at start-up. They have no narrow tile, so a projection too small for the half-height tile (the 256-wide K/V and per-layer-embedding projections) stays on the float kernel. Measured on rotating weights at 512 tokens: 3.4–5.8 TFLOP/s-equivalent against the float kernel’s 0.7–1.5, which is what made every non-Q4_K file’s prefill a third of the Q4_K_M one’s.
ORANGU_PREFILL_PLE_INPUTS on (0 turns it off) gemma4’s per-layer-embedding inputs for a prefill chunk computed on the device — the projection of the scaled embeddings as one GEMM, then one dispatch per (token, layer) row for the scale, RMS-norm, gathered-embedding add and scale, written layer-major so every layer’s stage reads its operand as a slice — recorded ahead of the chunk’s first chain. 0 computes them on the host as before: a device GEMM read back (18 MiB at 512 tokens), normed and added on the CPU, and uploaded again a layer at a time, with the device idle throughout.
— — Reading the ops table. A stamp’s time is the interval since the previous stamp, so the header separates the two things that interval can be: span X = device Y + host Z. host.turn is the gap between one device chain’s last dispatch and the next chain’s first — the previous chain’s readbacks and, on a mixture whose routed experts run on the cores, the whole host expert turn. On such a model it is the largest line in the table and the device’s share is about a quarter of the span, so a figure taken from span alone will overstate the card’s work several-fold.
ORANGU_GPU_OPS_RAW unset (off) With ORANGU_GPU_TIMESTAMPS=ops, also prints every stamp in order with its interval — the per-layer structure the totals fold away, for the question “which dispatch of the seventy is the slow one”.
ORANGU_DECODE_BATCH unset (off) Set to have concurrent slots’ decode steps run as one forward pass — the weights read once for all of them — instead of one after another. Each decoding slot joins a rendezvous; a step runs when every joined slot has offered one, and a slot that finishes leaves so no step waits for it. Correct (concurrent answers are the solo answers), and off because it is not yet faster: the batched step runs on the prefill chains, whose per-call resource building outweighs the saved weight reads at two to four rows (two streams 48 tok/s aggregate against 54 for one).
ORANGU_PREFILL_ATTN unset (off) Set to run the standalone prefill attention dispatch where the fused path above declines, instead of the CPU attention loop. Unlike the fused path this pays a Q upload and an attention-output readback per layer, which the CPU loop does not, so it wins only at long prompts. A measurement aid rather than a recommended setting.
ORANGU_PREFILL_TILED_ATTN on where subgroups are supported (0 turns it off) The prefill’s attention as a tiled kernel: one workgroup per (tile of queries, head), the tile’s queries staged once in workgroup memory, the window walked in blocks of positions, every thread owning a few rows × a few columns × a slice of head_dim so each K or V element read feeds one multiply-add per row and the only cross-lane work per position is a three-step shuffle of the scores. The online softmax runs per thread over its own columns and the lanes sharing a row merge their partials once at the end. The per-query kernels it replaces (ORANGU_NO_PREFILL_GQA, ORANGU_GQA_HEADS below) walk the window one position at a time with a reduction and an exp per head per position, and every query re-reads K and V; measured on the E2B prefill at 2048 tokens the attention dispatch went from 91 to 74 ms per 512-token chunk (−18%). Output is the same up to summation order. 0 is the control arm, and the fallback for an adapter without subgroups or with subgroups narrower than 32 lanes (the shuffles fold 32 lanes unconditionally; the Mali-G720’s subgroups are 16, and there the kernel is not built — [vulkan] … subgroups are 16 lanes wide; the cooperative attention kernels need 32 and are not built at startup — the same rule that keeps the cooperative decode kernels off such a device). ORANGU_FA_ROWS (1–4, default 2 at head_dim 256, more for narrower heads), ORANGU_FA_COLS (2/4/8, default 8), ORANGU_FA_LANES (32/64 lanes per row group, default 32) and ORANGU_FA_DLANES (8/16 lanes across head_dim, default 8) pin the tile for measurement: on this card every four-row tile lost to the two-row one, whichever way the lanes were split, because the registers those rows hold cost more occupancy than the halved loads buy.
ORANGU_NO_PREFILL_GQA unset (GQA sharing on) Set to disable sharing each KV head’s reads across the query heads that use it in the prefill attention kernel, falling back to one workgroup per (head, query). Only affects models whose n_head exceeds n_head_kv.
ORANGU_GQA_HEADS unset (chosen by register budget) Pins how many query heads of one KV group a prefill attention workgroup owns. Must divide the group size. Sharing a KV read across more heads and keeping enough waves resident to hide that read pull in opposite directions; the default picks from inside the measured band. A tuning knob for other GPUs.
ORANGU_GPU_TRACE unset (off) Logs the number of GPU submissions per decode step to stdout — a diagnostic for round-trip counting, no effect on the computation.
ORANGU_HOST_MATMUL_KIB 6144 Weight bytes, in KiB, below which a decode matmul (n_tokens == 1) runs on the host instead of the device; 0 switches the rule off and is the control arm a sweep needs. A device matmul costs a large fixed amount before it computes anything — building a command buffer, handing it to the driver, waiting for the device to take it — and that cost is several times the host’s, so the two backends’ cost lines cross: above the crossing the device’s bandwidth wins, below it the device’s fixed cost does. The threshold is compared against the whole batch’s weights, not one op’s, because the fixed cost is per submission and a batch pays it once. It is machine-dependent — the card, the driver and the clock all move it — so measure it where the model will run: cargo test --release --bin orangu-server decode_matvec -- --ignored --nocapture prints the crossing directly. The default is this project’s own card, where the probe put the crossing at 5.3–7.7 MiB and an end-to-end sweep agreed (4 MiB −6%, 6 MiB and 8 MiB within 1% of each other, 20 MiB −5%). Measured worth: +4.7% on Ornith-1.5-35B-A3B, +10.1% on Llama-3.2-3B, +0.3% (neutral) on gemma-4-12B, whose tensors are all above the threshold. Unlike most tuning here this changes the numbers a model produces, because a host matmul sums in a different order than a device one; it is a routing choice between two backends the engine already trusts, the same choice matmul_host_fallback makes for a type the device has no kernel for.
ORANGU_NPU_WEIGHTS_GB 1.5 How many GiB of requantized weights the Rockchip RKNPU backend may hold on the device; 0 keeps every weight on the CPU and is the control arm. The single biggest lever on what backend = npu is worth, because everything that does not fit runs on the CPU exactly as before. The ceiling is the device’s, not the machine’s: an RK3588 stops handing out memory near 2 GiB across all contexts — measured, a loop allocating 4 MiB at a time is refused at 2048 MiB, and 254 contexts holding a weight plus scratch each are refused at 2034 MiB. The default leaves headroom under that, because the allocation that fails when it runs out is as likely to be scratch for an already-resident weight as a new weight. Measured on an RK3588 prefilling a fresh 1976-token prompt through gemma-4-E2B-it at Q4_K_M, against backend = cpu’s 16.4 tok/s: 0 → 16.3 tok/s (0.99×, the wrapper costs nothing), 1.0 → 19.4 (1.18×), 1.5 → 21.2 (1.29×), 1.75 → 22.2 (1.35×). Still climbing at the top of that range, so the limit is the device and not diminishing returns; the quarter-gibibyte between 1.5 and 1.75 is what a model with a vocabulary-sized output projection needs for scratch at a wide rung. Running out is handled — the weights already placed stay, the rest run on the CPU — and says so once.
ORANGU_NPU_MODE int8 Which arithmetic the RKNPU is asked for: int8 (int8 × int8 → int32, with a symmetric scale per output channel on the weight and one per token on the activations, all recombined on the host) or fp16 (float16 × float16 → float32). int8 is the default on throughput and density: measured at K = N = 2048, 926 GFLOP/s at 64 tokens and 1190 at 512, against fp16’s 265 and 266 — 4.5× — at half the bytes per weight, which doubles how much of a model fits under ORANGU_NPU_WEIGHTS_GB. fp16 needs no requantization of a GGUF weight at all and is the fallback if int8 ever costs visible quality. The platform’s mixed mode (fp16 × int8) and every int4 variant are rejected by this runtime as unsupported matmul dtype ... in this platform, so there is no third option on RK3588. An unrecognized value warns and uses int8 rather than silently selecting something.
ORANGU_NPU_MIN_TOKENS 16 The fewest tokens in a matmul worth sending to the RKNPU; anything narrower goes to the CPU backend, and 0 sends everything. This is what makes backend = npu not a regression. The device’s fixed cost is about 0.4 ms per call, so at one token it manages 20 GFLOP/s against 1190 at 512, and the CPU — which reads the Q4_K weight straight out of the mapped file rather than a requantized copy — wins a narrow call outright. Overriding Backend::matmul_decode is the semantically right statement and was not sufficient: a decode step reaches matmul too, and with only that override in place backend = npu decoded at 2.15 tok/s against cpu’s 5.27. This width rule closed it to 5.12. Same shape of argument as ORANGU_HOST_MATMUL_KIB above, applied to a device whose crossing is in tokens rather than bytes.
ORANGU_NPU_FFN measured (0 never, 1 always) Whether a prompt’s feed-forward blocks go to the NPU’s precompiled graphs. Unset, it is measured once the CPU prompt path is final (prompts on the CPU, prompt_weights prepared): one block is timed on the NPU at the width its graphs were compiled for and through the CPU’s own prompt path at 384 tokens — the width it takes a prompt in — and the NPU is used only when it is at least 1.25× faster a token, since it also holds every chunk of the pass to its width and leaves the cores idle while it runs. The [adapt] npu: line says what was measured. The first form of the check compared both sides at the NPU’s 128 tokens and flipped between starts on the CIX P1 (11.7 ms against 16.4 in one); the start that took the NPU prefilled gemma-4-E2B at 74 tok/s where the cores alone do 150. When prompts run on a device the NPU is kept. 0 turns the NPU’s feed-forward blocks off, 1 uses them without measuring.
ORANGU_RKNN_LIB unset (/usr/lib/librknnrt.so, then /usr/lib64, /usr/local/lib, /usr/lib/aarch64-linux-gnu, then the bare soname) An explicit path to Rockchip’s RKNN runtime, for an SDK installed somewhere the search does not guess. It replaces the search rather than being tried first, the same rule ORANGU_NPU_LIB has for the NOE stack and for the same reason: a probe pointed at one library must not quietly succeed against another that happens to be installed. Honoured identically by orangu-server system’s report and by the backend that opens the library, because both go through one resolver — a machine where those two named different files would be a machine whose report is a lie.
ORANGU_EXPERT_DOT_TILED on (0 turns it off) Whether a routed expert’s prefill projection — many tokens against one expert’s rows, the call that was 87% of a mixture-of-experts model’s prefill CPU — takes the token-tiled integer-dot kernel. The kernel loads each 32-byte weight block once and dots it against four tokens with the products kept in registers, one horizontal reduction per output; the kernel it replaces reduced every block for every token and wrote the partial sums back to memory. Same integer products, so the sums match; only the float order of the block scales differs (within 1e-4 of the per-block reduction). Measured on the 26B-A4B mixture: prefill 37.5 → 43.8 tok/s at 128 tokens, 37.4 → 47.1 at 512; decode untouched — one token takes the single-token kernel, not this one. 0 restores the per-block reduction, the control arm. Takes effect only where the host has the vector extensions the kernel is built for; elsewhere the per-block kernel runs regardless.
ORANGU_SSM_FANOUT on (0 turns it off) Whether a selective state-space block (the nemotron_h_moe recurrence) updates its heads in parallel. Each head owns a [head_dim, state] matrix and its own slice of the output and reads only its group’s B/C, so the update is a fan-out with nothing to synchronize; before, the sixty-four heads ran one after another on the calling thread while the rest of the pool idled between two expert stages. Bit-identical — the per-head arithmetic and its summation order are unchanged, only which thread runs which head. Measured on a 30B-A3B hybrid: recurrent.delta 16.4 → 7.7 ms per token, out of ~145. 0 runs the heads serially, the control arm.
ORANGU_Q5_UNPACK_VECTOR on (0 turns it off) Whether the host Q5_0, Q5_1 and Q4_0 kernels unpack a 32-element block with the AVX2 form — one byte shuffle and one compare give every lane its fifth bit — instead of the scalar loop that shifted qh once per lane and could not be vectorized by the compiler. Exact: the same bytes come out, checked lane by lane against the scalar loop. Reaches the routed experts’ decode dot, the many-token prefill dot, and any of those types on a host-side projection; on a file whose experts are Q5_0 the host Q5_0 rate was 20 GB/s against Q8_0’s 36, and the unpack was the difference. 0 keeps the scalar loop, the control arm. Only on x86-64 with AVX2; elsewhere the scalar (or NEON) form runs regardless.
ORANGU_EXPERT_WILLNEED on (0 turns it off) Whether a mixture-of-experts layer asks the kernel to start reading every expert it has just selected — across the gate, up and down tensors — before the first of them is touched. Only acts on a model whose routed experts are larger than the machine’s memory (the startup line [experts] … streamed from the disk says when); on a model the page cache holds, the hint would find every page in place and cost a page walk per expert per layer, so it is skipped. Without it, each expert is faulted in page by page from under the dot that needs it: a thread per expert, each fault a small read the drive answers one at a time, and the kernel’s read-ahead never sees a pattern it recognizes. With it the whole selection is one batch of large sequential reads the device can queue, and the down projection’s bytes arrive while the up projection is computing. 0 leaves the reads to the page faults, the control arm. Ignored by a read route that bypasses the page cache (ORANGU_EXPERT_READ=direct).
ORANGU_MATMUL_MIN_ROWS 64 The fewest output rows a thread-pool task takes when a host matmul fans out over its rows (with_min_len, a floor rather than a chunk size — rayon still splits down to it and no further). One row per task is a few hundred nanoseconds of arithmetic inside a job that has to be created, possibly stolen and joined; on Ornith-1.5-35B-A3B 40% of all CPU samples were thread-pool machinery and 24% were sched_yield alone, after the host-routing rule multiplied a decode step’s parallel regions. Worth +2 to +4%, and it takes the pool’s share from 40% to 34%. 1 disables the floor and is the control arm. A flat number rather than one scaled to the pool, and that was measured, not assumed: scaling it as out_dim / threads is 6.7% slower than no floor at all, because a routed expert’s projection already sits inside an eight-wide fan-out over experts and this call cannot see it. Machine-dependent, like every other threshold here. Bit-identical at any value — each output row is computed and written independently, so this changes only which worker does what.
ORANGU_DECODE_STAGES unset (off) 1 measures the one-token passes a generated token is made of; prefill (or pp) measures the wide passes a prompt’s chunks are instead, and the two are exactly complementary — whichever width is asked for, the other is subtracted, because a multi-second prompt averaged into a 90 ms token is a breakdown of neither. A prefill window’s header carries the chunk width the sizer chose, since a per-chunk cost cannot be read without it. Accumulates the elapsed time of each forward-pass stage, drained by GET /decode-stages and printed by orangu-bench as its stages lines. Answers what a CPU profile cannot: a stage running alone on one core costs the whole of its wall time while the same cycles spread over sixteen cores cost a sixteenth, and a stage parked in a GPU driver waiting for a submission barely samples at all. forward, attn, ffn.router and ffn.routed are timed in code every architecture goes through, so they are reported for any model; the finer stages are timed per-architecture and an architecture without them reports the remainder under other; a token recorded as one device submission (the Qwen hybrid trunk on Vulkan) reports its host-side recording as device.record and the submissions, the wait and the readback as device.submit. Off by default because two clock reads and two atomic adds per stage entry are not free at the rate a decode step enters them, and a measurement that perturbs what it measures whenever it is linked in is worse than none. A diagnostic; no effect on the computation.
ORANGU_DUMP_SHADERS unset (a directory path) Set to a directory to write the generated WGSL of the decode-path kernels (Q4_K/Q6_K matmul-vec, RMSNorm, split-attention), the tiled prefill GEMM and the integer-dot prefill GEMM (per weight type and row tile) into it as .wgsl files at startup, then continue normally. A profiling aid: pair it with the driver’s own RADV_DEBUG=shaders,shaderstats (which dumps the compiled ACO ISA and register/occupancy stats for this GPU), or hand the WGSL to an offline analyzer such as Radeon GPU Analyzer on an RDNA3+ machine. No effect on the computation.
ORANGU_SPECULATIVE unset (off) Enables prompt-lookup speculative decoding: each step drafts the next few tokens by matching recent output against an earlier point in the context and verifies the whole draft in one forward, so the weights stream once for several tokens. Greedy-only (the output is identical to non-speculative greedy decoding); ignored for non-greedy sampling and for multi-slot batched decode. Currently slower on this GPU — see SERVER_ROADMAP.md Step 12 — because the multi-token verify runs the CPU-orchestrated forward; kept for hardware/paths where a resident multi-position forward makes it a win. Off by default.
ORANGU_SPEC_NGRAM 2 With ORANGU_SPECULATIVE, how many trailing tokens must match an earlier point in the context to trigger a draft. Lower drafts more often (more speculative work, more misses); higher drafts only on a longer exact echo.
ORANGU_SPEC_DRAFT 4 How many tokens to draft (and verify in one forward) once a match is found. The ceiling on tokens a single accepted step can produce. Also overrides [orangu-server].draft_tokens for a run using a draft model, so a sweep of drafting depth reads the same whichever drafter is in play.
ORANGU_NO_MTP unset (head on) Turns off the multi-token-prediction head the server attaches when one ships beside the model it is serving (MTP/mtp-*.gguf; see engine::arch::qwen4exp_mtp). The control arm for measuring what a head is worth, and the way out if one ever costs more than it saves. The head only ever changes how fast an answer arrives, so this changes throughput and nothing else.
ORANGU_CPU_TIMESTAMPS unset (off) Logs, per decode step, how its host time splits: since-return (from the previous step’s token to this step’s start — the generate loop’s own share), record (building the step’s commands), submit+wait (the submissions and the wait for the device), and each intermediate submission’s own call time. For a prefill chunk it logs the host’s phases instead — stage (the token and per-layer embeddings gathered), ple (the per-layer inputs recorded), record (the layers’ chains), submit, fill-previous (the wait for the previous chunk’s K/V rows and their fill into the host cache — with the device pipelined this is most of the chunk’s wall time, which is the point), land (the last chunk’s residual read back) and kv. Read with ORANGU_GPU_TIMESTAMPS=ops’s idle before to see how much of the host’s turn the device waits for.
ORANGU_GPU_TIMESTAMPS unset (off; 1 or ops) Logs a per-decode-step GPU timing breakdown to stderr — the per-layer-embedding (PLE) projection, the sum/average/slowest across all model layers, and the output-norm-plus-lm_head tail, in milliseconds. Also logs a [gpu-op-breakdown] line splitting each token into qkv-side (Q/K/V matmuls + norm/RoPE + KV write), attention (split-k), and ffn-side (wo/gate/up/down matmuls + their norms + GELU/mul + PLE + copies) — so matmul vs attention vs overhead is measured, not estimated. Requires an adapter with TIMESTAMP_QUERY and TIMESTAMP_QUERY_INSIDE_ENCODERS; a diagnostic, no effect on the computation. Unavailable on a split model: a query set belongs to one device and a split resolves none, so GET /gpu-timings answers {"enabled": false, "timings": null, "unavailable": "split"} — the reason is named rather than left as an empty result, since a client that reports nothing when it receives nothing makes a split run look like one whose GPU stages cost nothing. --flamegraph still profiles the CPU side. ops additionally stamps every dispatch of the step from inside its compute pass and prints a [gpu-ops] block per token: one line per kind of dispatch (ffn.down, attn.norm, ple.gate, …) with its total, count and average — the cost as the step pays it, the barrier and any pass boundary in front of the dispatch included, which is what separates a slow kernel from an expensive schedule around a fast one. The header line also carries idle before, the device’s gap between the previous step’s last dispatch and this one’s first — the host’s turn as the device sees it — and a memory: line with the allocator’s block and allocation counts. ops replaces the per-layer brackets of 1 rather than adding to them. Needs TIMESTAMP_QUERY_INSIDE_PASSES as well; says so and prints nothing otherwise. ops is not a cheap observer: it resolves a timestamp query set in a submission of its own every step, which on a small model costs about a third of the token (measured: 184.7 tok/s with no instrument, 183.2 under ORANGU_DECODE_STAGES=1, 128.7 under this). A rate measured under it is therefore not the machine’s rate, and a device total from its table cannot be subtracted from one measured without it — its numbers are comparable with each other and nothing else. GET /props reports gpu_op_timer, and orangu-bench prints a warning line beside the rate whenever it is on.

Shader compilation is cached to disk across restarts (~/.orangu/server/<adapter-key>/cache.bin, keyed by a vendor/device- derived string so a cache built for one GPU is never handed to another) — a startup-time optimization only, with no effect on decode/prefill throughput once running.

The Metal backend

engine::backend::metal::MetalBackend is the section above, on Apple hardware. It is not a reimplementation of anything: nothing in VulkanBackend is Vulkan-specific. Every compute pipeline, every kernel in vulkan_shaders, the weight/op/uniform arenas, the fused decode and prefill submissions, split-k attention and GPU sampling are written against portable wgpu and WGSL, and only ever ran on Vulkan because VulkanBackend::try_init asked wgpu for a Vulkan adapter and nothing else.

So try_init now delegates to try_init_backends(wgpu::Backends), and MetalBackend calls that with METAL and wraps the result. naga translates the same WGSL to MSL rather than SPIR-V. MetalBackend:: as_wgpu returns the inner engine, so every fused submission engine::arch reaches for through that hook runs here too — the point of the newtype is to keep the two distinct at the type level while sharing all of the code.

Bring-up is entirely feature-negotiated in try_init_backends, so the Metal module has no adapter logic of its own. On Apple silicon SHADER_F16 (so the f16 KV cache is on), SUBGROUP (so the cooperative-reduction attention kernel is on) and TIMESTAMP_QUERY (so ORANGU_GPU_TIMESTAMPS works) are all present. PIPELINE_CACHE is not: wgpu::util::pipeline_cache_key returns None for anything but Vulkan, so no on-disk pipeline cache is built and cold start recompiles the kernels, which Metal’s own shader cache largely absorbs.

One kernel needs a different form on Metal, and it is a correctness difference rather than a tuning one. The tiled prefill GEMM is the only kernel that fills shared memory by having many threads each write a single dynamically-indexed component of a shared vec4 (store_w/store_x into tile_w/tile_x; four different threads write the four components of each vector). On Vulkan/RADV that lowers to a 4-byte store, which is what lets the read side take four values per load. On Metal it does not: a probe kernel on CI’s Apple Paravirtual device read back [1, 0, 0, 0, 5, 0, 0, 0, …] — exactly one surviving component per vector, the signature of a read-modify-write of the whole 16 bytes with the four writing threads clobbering one another. Every tiled-path cross-check disagreed with the CPU backend there, including plain f32 with no dequantization in play, while every scalar-shared-memory kernel passed.

vulkan_shaders::coop_vec4_tiles therefore enables vec4 tiles only where they are known-correct, and everything else gets a scalar tile: four loads where the vec4 form takes one, same layout, same arithmetic, same results. An unrecognized backend defaults to the safe form on purpose. Whole-vec4 stores would also have worked and are the faster-looking fix, but the tiled kernel cannot use them without changing its tile layout — store_w’s four consecutive slots are four consecutive rows at one k, while the fill deliberately gives one thread RUN consecutive k of one row so a quantized block’s scale and min hoist once per run. tests::coop_tiles_are_vec4_only_where_component_stores_work asserts the invariant, and its whole-vec4 twin is the control that pins the fault to the component store rather than to shared memory or the barrier.

Two further things genuinely are Vulkan-only, and both check the API rather than assuming:

Correctness rests on the same tests as the Vulkan path, not on a separate suite: vulkan::shared_test_backend asks for whichever wgpu API the platform has, so the per-ggml_type cross-checks against CpuBackend — ordinary tests, not #[ignore]d ones — run against a real Metal device on CI’s macOS runner. An #[ignore]d real-model test (gemma4_predicts_paris_after_capital_of_france_metal) adds the whole-model claim on top, and CI fails rather than passes if it skips for want of an adapter.

CUDA, OpenCL, and ROCm backends

engine::backend::cuda::CudaBackend, engine::backend::opencl:: OpenClBackend, and engine::backend::rocm::RocmBackend each implement the same Backend trait, at a deliberately smaller scope than Vulkan: one dequantizing matmul kernel per ggml_type, a direct port of vulkan_shaders’s MAIN_REDUCE_SUFFIX reduction strategy restated per kernel language (CUDA-C, OpenCL-C, HIP-C), cross-checked against CpuBackend the same way VulkanBackend’s own tests are. Deliberately not ported: VulkanBackend’s cooperative/tiled dispatch, GPU-resident attention/RoPE/norm fusion, fused whole-layer submissions, GPU-side argmax sampling, and the disk pipeline cache — none of the three has been run against real hardware during development (no NVIDIA GPU, no ROCm install, no OpenCL ICD on the project’s dev machine), so correctness rests on the kernel math matching engine::quant’s already-verified dequant code line-for-line, plus the same CPU cross-check test pattern vulkan.rs uses (which, like those tests, skips gracefully rather than fails when no matching device is found).

Their ggml_type coverage is a subset of what engine::quant reads on the CPU, and each backend’s SUPPORTED_TYPES is the authority: the float types, the legacy quants, Q2_K through Q6_K, and IQ4_NL. The remaining IQ* types are absent because each indexes a lattice codebook that would need its own uploaded buffer — VulkanBackend has one (vulkan_shaders’s IQ_GRID_PRELUDE, bound at @binding(4)) and these three do not. IQ4_NL is the exception that fits: a 16-entry level table, small enough to inline into the kernel source, and the reason a Q2_K download of a model whose rows aren’t 256-divisible (see the Scope section of the server chapter) is runnable here at all.

VulkanBackend’s own coverage is a subset too — it has no shader for the three narrow IQ1_* types (ids 64-66). Every GPU backend therefore overrides Backend::supports_type, and engine::backend::unsupported_tensor_types walks every shard’s tensor directory once at startup so a gap is reported as an error naming each missing type. Before that check existed, the gap surfaced as a panic from inside matmul partway through the first request.

That list used to include IQ1_S, IQ1_M and IQ2_XXS, and dropping them from it is what makes a “dynamic” 2-bit release runnable on a GPU at all: unsloth/Qwen3.8-27B-GGUF:IQ2_XXS is 96 IQ1_M tensors and 48 IQ2_XXS ones, so the startup check refused the device for the whole model rather than for those tensors. iq_grid_words now uploads IQ2XXS_GRID and IQ1S_GRID alongside the four it already carried, taking the codebook buffer from ~15 KiB to ~33 KiB — IQ1S_GRID is 2048 eight-byte lattice points, 16 KiB by itself. The iq1* pair needed one thing the others did not: their codebook values are signed and they carry no sign field at all, the per-group freedom being whether a ±0.125 delta is added or subtracted, so IQ_GRID_PRELUDE gained iq_grid8_signed beside iq_grid8. Reading an iq1* grid byte as unsigned is well formed and produces plausible output, which is why matmul_matches_cpu_backend_for_iq1_m exists rather than being folded into a generic sweep.

cudarc and the resolved opencl3 version both dlopen their vendor library (libcuda.so/libnvrtc.so, libOpenCL.so) at runtime and return a real error if it can’t be found, so cuda/opencl are always compiled in — nothing extra is needed to build orangu-server. cubecl-hip-sys (ROCm’s underlying bindings) is different: it directly links -lamdhip64 -lhiprtc at build time whenever its build script finds a ROCm install, which would break a plain build on a machine without ROCm — so rocm sits behind its own Cargo feature, off by default (see BUILDING.md).

cudarc has one notable wrinkle: unlike every other fallible step here, it panic!s (rather than returning a Result) the first time a driver/NVRTC call is made and no libcuda.so is found. CudaBackend::try_init runs try_init_inner under std::panic::catch_unwind (with the panic hook silenced for the call) specifically so a non-NVIDIA machine gets the same graceful None/CPU-fallback outcome every other missing-backend path already has, not a crashed server.

The startup banner’s Mode row

Mode names which of --all/--code/--review/--explorer/--embedding/ --image this process came up as (image whenever the model is a qwen_image file, flag or no flag) — the other half of “what will this server do with my request”, beside Model. It decides the sampling defaults, whether reasoning is suppressed, and whether the generation endpoints answer at all, and without it on the banner the only way to tell a --review server from an --all one was to send a request and read the answer. /props reports the same value as role.

A coordinator starts a real --review process for the review role rather than labelling a request on another one, so Mode review on the banner is what an /auto_review run is served by. all, code and explorer do share a process — nothing about them differs that a request cannot carry — and a request served under one of those on another’s process says so on its own completion line. A request served under another role says so on its own completion line, which is how a run can be checked rather than assumed:

orangu-server: [slot 0 as review] prompt 27 tokens in 1.50s (17.94 tok/s), generated 60 tokens in 2.41s

Refusing a context the device cannot hold

The cheapest way not to lose a device is not to hand it work that loses it. A request’s KV cache is the one part of its footprint that scales with what the caller sent, and past a point the device stops answering a submission inside the driver’s reset timeout — not because an allocation failed, but because it is paging while it works. What follows is a reset, an exit 75, and every other request on the process dying with the one that caused it.

So the size is refused where an answer is still free, at the HTTP boundary, before a token is processed:

$ curl -s localhost:8100/v1/chat/completions -d @big.json
prompt (41833 tokens) plus max_tokens (512) needs 42345 tokens of KV cache,
more than the 16384 this server has room for on its device — send a shorter
prompt, lower max_tokens, or serve this model on a device with more memory

400, not 500: retrying changes nothing and only a shorter prompt can work, which is what the status code tells a client to do. /auto_review does exactly that — a file whose prompt is refused is re-sent without its whole-file section and with its diff bounded, and is reviewed from that.

The number is derived per machine, never written down for one. With the paged KV pool on (the default) it is the pool’s own token capacity, which KvPool::pages_within computes from this device’s headroom after the weights and this model’s KV geometry — so it rises with the card. Without a pool it is the headroom figure the startup report already prints. With no device at all there is no ceiling and nothing is refused. The line naming it is printed at startup, next to the pool’s own:

orangu-server: [kv] paged cache — 256 pages of 64 tokens (16384 tokens total …)
orangu-server: [kv] one request may hold up to 16384 tokens of context; a
longer prompt is refused rather than risking the device

A request the pool cannot promise room for is not refused on that account alone — it falls back to a per-request contiguous cache, which is what every request used before the pool existed, and keeps prefix reuse through slot_store. It is refused only when it is also over the ceiling above.

Losing the GPU device (device_lost.rs)

Every wgpu readback in VulkanBackend funnels its failure paths through one place, crate::device_lost::fail, which records the loss, writes the real detail to the server’s log, arms a 75/EX_TEMPFAIL exit two seconds out, and panics to unwind the request that was in flight. Three things it replaced are worth naming, because each was a separate way the old code made a driver reset worse than it had to be:

One funnel is not enough on its own, though, because wgpu does not always hand a lost device back as an Err at all: Device::poll routes it through handle_error_fatal, which panics from inside wgpu (Error in Device::poll: Validation Error / Caused by: Parent device is lost), so the Result the engine checks never arrives. Every wgpu call made after the device dies ends that way. panic_capture’s hook therefore reads every panic’s message and calls device_lost::note_panic, which marks the loss and arms the same exit without panicking again — process-wide, so a wgpu panic on a thread nothing catches (a rayon worker) still takes the process down cleanly instead of leaving it up with a dead GPU. The message match is on the condition (device is lost, DeviceLost), not on any one call’s name.

engine::generate’s catch_unwind recognizes the loss (device_lost:: is_lost) and swaps the panic’s captured detail — meaningless to a caller, since the backtrace describes the driver rather than their request — for device_lost::CLIENT_MESSAGE (panic_report). Requests that arrive in the window before the process exits get that same sentence without being started at all.

The exit code is the contract with orangu-coordinator, which names the same number (process::SERVER_EXIT_DEVICE_LOST) so it can report the restart as the recovery it is instead of an unexplained crash. Nothing about the mechanism is Vulkan-specific: MetalBackend is the same engine, so a lost Metal device takes the identical path.

Correctness testing

VulkanBackend’s dequant math (each quant type, bit-for-bit against the CPU backend, across both dispatch paths), fused post-attention chain (including a dedicated test that calls it twice for one layer with different inputs each time, to catch cache-reuse bugs specifically), and fused attention (including GQA head-grouping, sliding-window attention, proportional RoPE, and Gemma4’s cross-layer KV-donor case — two different layers sharing one KV cache) are covered by cross-check tests in engine::backend::vulkan::tests, run on real hardware whenever it’s present and skipped otherwise. Those same tests are the Metal backend’s tests: the device they run against comes from shared_test_backend, which asks for whichever wgpu API the platform has, so they are a Vulkan cross-check on Linux and a Metal one on macOS. The CUDA/OpenCL/ROCm backends follow the same skip-if-no-device pattern.

A second set of tests runs a full forward pass against a real downloaded model and is marked #[ignore] so the normal suite doesn’t require one. These read the model path from an environment variable, and each panics with a clear message if its variable is unset when the test is run (cargo test -- --ignored):

Variable Used by Points to
ORANGU_TEST_MODEL Gemma/qwen35moe/qwen35 real-model forward-pass tests A local .gguf chat model file
ORANGU_TEST_EMBEDDING_MODEL embedding-model tests A local .gguf embedding model file
ORANGU_TEST_QWEN3VL_MODEL qwen3vl tokenizer/embedding tests A local qwen3vl .gguf file
ORANGU_TEST_LLAMA_MODEL llama-architecture forward-pass test A local Llama-3.x Instruct .gguf file
ORANGU_TEST_QWEN3MOE_MODEL qwen3moe routed-feed-forward forward-pass test A local Qwen3-MoE .gguf file (e.g. Qwen3-Coder-30B-A3B-Instruct)
ORANGU_TEST_MISTRAL_MODEL mistral3 forward-pass test A local Ministral-3 .gguf file
ORANGU_TEST_PHI_MODEL phi3 real-model forward-pass test A local Phi-3/Phi-4-mini .gguf file
ORANGU_TEST_QWEN4EXP_MODEL qwen4exp forward-pass test A local Qwen3.8-Flash-Next .gguf (the first shard of a split model)
ORANGU_TEST_QWEN4EXP_MTP qwen4exp draft-head acceptance test One of that repo’s MTP/mtp-*.gguf heads; measures how often the head guesses what the model goes on to say
ORANGU_TEST_QWEN_IMAGE_MODEL qwen_image end-to-end picture test A local Qwen-Image .gguf; its qwen2vl encoder and the VAE are found under ORANGU_TEST_QWEN_IMAGE_MODELS_DIR (default: the model’s directory’s ancestors)

HTTP layer and web UI

http::mod assembles the router and shared AppState (model, scheduler handle, config, workspace root, start time); http::openai and http::native hold the OpenAI-compatible and native handlers respectively; http::files holds the file-lifecycle API (see the next section); /v1/shutdown lives in http::mod itself since it’s neither. Ctrl+C, SIGINT, and POST /v1/shutdown all converge on the same shutdown path via tokio::select!, mirroring orangu-coordinator’s own pattern. What that path does first is tell everything long-running to stop where it stands, because the runtime’s drop joins the blocking threads: engine::image:: request_shutdown sets a flag transformer::forward reads before every block and the VAE before every layer (ForwardInput::cancel, decode_unless/encode_unless, both also carrying a request’s own cancel flag); npu_tool::stop_preparation sets its own flag, which prepare reads between stages and block_error between sampled blocks, kills the compile child, and sets orangu::npu_ffn::request_stop, which NpuFfn::open reads between cached graphs read off disk and the device thread between binds (releasing what it bound). A shutdown measured on the development board went from 107 s with a 1024-pixel picture in flight, and 26 s with NPU graphs binding, to under three seconds.

web::mod serves a small server-rendered chat UI (vanilla HTML/CSS/JS, no build step) on its own [web].port, sharing the same in-process Engine as the API so a chat turn never makes an HTTP hop. web::render renders markdown to HTML (including syntax-highlighted code blocks) with the same markdown/syntect crates orangu’s terminal UI uses. web::mermaid draws ```mermaid blocks (below). web::sessions persists each chat as ~/.orangu/server/sessions/<uuid>/chat.json. web::models is the model manager (below); web::mcp the MCP inventory (GET, and PUT /api/mcps, which config::rewrite_mcp_section writes one section at a time — a header through the line before the next — into the file WebState::config_file names, McpConfiguration::validate having refused what the loader would); /api/props the picture settings. The three are the panes of the console’s one Settings dialog (index.html’s #settings-overlay, app.js’s openSettings/showSettingsPane): a brand-brown column of white buttons picks the pane, a brand-brown footer holds Save and Cancel (saveSettings posts the Image form when it differs from the server’s defaults, then closes; Cancel just closes), the model manager’s polling runs only while its pane is shown, and the Image pane is greyed when /api/props answers image: null.

Downloading one code block

render::render_code_block wraps every fenced block in <div class="code-block"> with a .code-footer under it carrying the file name and a .code-dl button. The footer shares .gen-time’s dim/hover/ icon-size rules rather than restating them: a download in this console is a small dimmed icon at the lower right of the thing it saves, and reusing the rules is what keeps that true as they change. It is deliberately not a bar welded onto the <pre> — the highlighter’s theme is a fixed dark one that paints its own background inline while an unhighlighted block has none, so a strip with a background of its own would match one and clash with the other. Two more details are worth keeping:

The text is not sent twice. The obvious shape — a data: URI on an anchor, which is what a diagram’s download uses — is wrong here. A diagram is already a data: URI in the <img>, so its download reuses one that had to exist; a code block’s source is plain text in the <pre>, and a data: URI would be a base64 second copy of it. The transcript is re-rendered and re-sent on every streamed token, so that copy would ride along in every frame of the reply. app.js reads textContent back out of the <pre> on click instead, and the click is delegated from the transcript rather than bound per button — assistantEl.innerHTML = payload.html on each token would throw away a listener attached to the button itself.

The name is derived, and refused when unsure. declared_file_name reads the fence (mdast splits the info string into lang — the first word — and meta, so ```rust src/main.rs, ```rust:src/main.rs and ```rust title="src/main.rs" all land somewhere readable); file_name_from_first_line reads a first-line comment; otherwise Renderer::code_blocks_seen numbers it and default_extension extends it from the resolved SyntaxReference::file_extensions, which is why rust becomes .rs without a hand-written language table.

Every candidate goes through clean_file_name, which keeps only the last path component and rejects anything that isn’t plainly a file name. The strict version is the point: the first-line rule only fires on a comment holding a single token, so # Install the dependencies and #include <stdio.h> fall through to the generated name, and an extension has to start with a letter so # roughly 3.14 doesn’t save as 3.14. A confidently wrong name on a saved file is worse than a numbered one. The name is also model output, so it is escaped into the attribute like everything else in a reply.

A block routed to web::mermaid or web::plantuml never reaches render_code_block, so a diagram doesn’t consume a snippet number — it isn’t a code window, and it has a download of its own.

The licence header (orangu::license)

The module is orangu::license, in the library rather than in this binary, because the web console is not the only thing that generates code: orangu’s create_file tool writes a licence header onto a new file with the same call (see the Tools chapter). One module, so the two cannot disagree about the year, the licence, or where the header goes.

The licence is the project’s, not this program’s. Project::detect reads it from the workspace root — the license field of Cargo.toml, pyproject.toml or package.json, or failing that the LICENSE/COPYING file — together with the copyright holder, and Licence is the small set of licences whose per-file header is well established enough to write: the popular set from https://opensource.org/license, with the GNU family split the way SPDX splits it. A LICENSE file that is none of them is used verbatim — the project has said what its licence is — and a project that says nothing gets [DEFAULT_LICENCE]. Only a project with no copyright holder to name produces no header, because inventing an attribution is worse than omitting one. Choice is /license’s answer laid over all of it, held by the client’s ToolExecutor and read on every create_file.

This used to be a single hard-coded MIT constant, which meant every file generated into a GPL project — orangu’s own repository among them — arrived carrying an MIT header attributed to “orangu”, and git added. files::create already refused to put a header on an existing file for exactly that reason; a file being created is no more this tool’s to relicense.

One trap is worth naming, because it is the reason the copyright holder is not simply read out of LICENSE: a GPL or Apache project ships the licence’s own boilerplate, whose copyright line names the Free Software Foundation or the Apache Software Foundation. orangu’s own LICENSE opens Copyright (C) 2007 Free Software Foundation, Inc.. For those licences the manifest is the only source consulted; MIT and BSD put the project’s real holder in that line, so for those it is used when the manifest has nothing.

The licence texts are raw string constants in that module. Not data files, and not include_str! of any: there is nothing to ship beside the binary and nothing an install can lose. PLACEHOLDERS is a table of the process-wide tokens in them that get filled in at request time — <YEAR> today, anything else by adding a line there and a token to the file. Request time, not build time, is the point: a server left running past midnight on 31 December would otherwise keep stamping last year onto everything it saves. <HOLDER> is not in that table: it belongs to the project being written into, not to this process.

comment_style maps the saved file’s extension — not the fence’s language tag — onto a CommentStyle, since the extension is what the file will be read as. Line markers get one per line, with a bare marker on a blank line rather than a marker plus a trailing space; Block styles (<!-- -->, /* */, (* *)) open and close around the text. A test asserts the licence text itself contains none of those closing delimiters — one */ inside it would end the comment early and spill the rest of the licence into the file as code.

Two refusals are deliberate. An unknown extension gets no header rather than a guessed marker: # at the top of a JSON file, or % at the top of Objective-C, doesn’t produce a differently-licensed file, it produces a broken one, and .m (Objective-C or MATLAB) and .s (assembler-dependent) are ambiguous enough to be left out on purpose. And app.js puts the header below a shebang or an XML declaration, which have to stay on line one or the file stops being executable/well-formed.

render_code_block applies the licence to the source it hands the highlighter, so the header is in the block the reader sees — highlighted, selectable, and already in the textContent the download button reads back out. There is no endpoint and nothing for app.js to fetch or reassemble: what gets saved is what was on screen.

Order matters in that function. The file name is derived before the licence goes on, because file_name_from_first_line reads line one and line one is about to become // MIT License.

Only the rendering is licensed. The message stored in the session, replayed as context on the next turn, and written by Save as Markdown is the raw text the model produced — render_markdown_to_html is a pure render step, and that separation is what keeps the licence out of the model’s own context.

license::apply is the whole job — header, plus the shebang/XML-declaration placement — and is what both surfaces call: web::render on a block it is about to display, files::create on a file it is about to write. Each adds the one decision that is its own. For files::create that is only a file that did not exist before the call, since rewriting or editing an existing file must not stamp a licence onto somebody’s project. For web::render there is no such qualifier — a code block in a reply is generated code by construction.

render::LANGUAGE_EXTENSIONS and license::comment_style are two halves of one path — the first turns a fence tag into a file name, the second turns that name into a comment. A tag missing from the first (syntect ships Sublime’s syntaxes, which cover neither TypeScript nor Kotlin nor PowerShell, so typescript would otherwise save as .typescript) silently costs the snippet its licence header too. web::render’s every_name_render_generates_for_a_common_language_gets_a_header walks both halves for 64 fence tags and is what caught that.

Mermaid diagrams (web::mermaid)

Diagrams are rendered by merman, a headless Rust implementation of Mermaid that parses, lays out, and emits SVG without Node, Puppeteer, or a JavaScript runtime — the same requirement that made KaTeX a vendored asset rather than a CDN link. It is the opposite arrangement to the math path: $...$ ships raw TeX to the browser because no server-side TeX engine exists in Rust, whereas a diagram is finished server-side and the client only ever receives a picture.

Three decisions carry the design, each forced by something measured rather than assumed:

The SVG is embedded as an <img> data URI, never inlined. merman strips scripts and event-handler attributes from diagram labels, but it does not escape a literal </svg> inside one: a label of A["</svg>…"] emits that tag raw. Inlined, the HTML parser reads it as the real end tag, closes the diagram early, and lets the rest of the label escape into the transcript. An <img> makes the SVG a separate document — scripts inert, ids unable to collide with the page or with a second diagram on it — and keeps the same escape-everything stance web::render takes for every other node kind. Base64 rather than percent-encoding, because an SVG is full of #, <, " and &, and one missed escape silently truncates the image.

Labels are SVG <text>, not <foreignObject>. Mermaid’s usual HTML labels get no layout engine inside an <img> document and would render as nothing at all; HostThemeOutput::resvg_safe_editor is what converts them, and is therefore not optional. A test asserts no foreignObject survives, since dropping that setting fails silently as blank labels.

Each diagram is rendered twice, once per theme. An <img> cannot inherit the page’s CSS variables, so the palette is baked in and app.css shows whichever of the two finished pictures matches the current theme. The theme roles map onto the console’s own custom properties, so a diagram is styled like the transcript around it rather than arriving in Mermaid’s stock lavender.

Sizing, alignment, and getting the original out

A message is max-width: 50% of the transcript and diagrams are large — the ER diagram used to shake this out measures 2734×3571 — so the picture is scaled to fit with max-width: 100%. Rendering at natural size instead was tried and is wrong at this width: it puts most of the diagram behind a scrollbar inside a half-width bubble. The viewBox dimensions still go on the <img> as width/height, not to force full size but so the browser reserves the correct aspect ratio before the image decodes rather than reflowing the transcript when it lands.

Scaling means the displayed picture is well below full resolution, so each diagram carries a download control — a plain anchor onto the same data: URI the <img> already holds, so saving needs no JavaScript and no round trip and the file is exactly what is displayed. One per theme, toggled by the same rules that pick the image, so the saved SVG matches the screen instead of always being the light variant. A test decodes the href and asserts it is both the image’s own URI and a complete SVG document, since a broken link here fails silently.

Alignment: the <img> is display: block. As an inline element it sat in a line box, subject to inline alignment and carrying the baseline’s descender gap beneath it, which pulled diagrams off the left edge; the theme rules therefore switch between block and none, never inline. .mermaid-diagram’s margin: 0.6em 0 also clears the UA stylesheet’s margin-inline: 40px on <figure>, which would indent every diagram.

Two properties of the streaming path shape the rest. The transcript is re-rendered from scratch on every token, and a diagram costs roughly 2 ms to lay out — so completed diagrams are cached by source hash (failures too, or a mislabelled block would be retried once per token), and render::unterminated_fence_start withholds the one block the document ends inside of.

That second guard matters more than it looks. A half-written diagram usually still parses — flowchart TD plus one edge is valid Mermaid — so without it the reader would watch a diagram redraw, reflow and jump on every token until the fence closed, each throwaway state costing a full layout the cache can never hit. The guard therefore keys on the fence, not on whether the source happens to parse.

A source that doesn’t parse falls back to the ordinary highlighted code block. This is a common path, not an edge case — models emit near-miss Mermaid regularly — and it is why merman was chosen over mermaid-rs-renderer, the other pure-Rust candidate: the latter answers malformed input with a 16×16 blank SVG rather than an error, in strict mode as well as lenient, leaving no way to tell a diagram from a failure and putting an empty frame where the source should be.

Detecting a diagram without a tag

A ```mermaid tag is not always there — models emit diagrams into bare fences, and an attached .mmd file has no fence at all — so mermaid::looks_like_diagram decides from the content.

It cannot simply ask merman. Mermaid’s parsers are extremely permissive, and merman inherits that faithfully: handed the sentence graph is a data structure of nodes and edges, it detects a flowchart and renders one, with is as a node. classDiagram is what you want and a log line reading info: build succeeded behave the same way. Measured over a corpus of realistic non-diagram blocks, merman’s own detector produced false positives on three of seventeen — and each would have turned someone’s prose or logs into a nonsense picture.

The gate is therefore a table of the diagram headers with the tokens each may be followed by (flowchart takes a direction, pie takes title or showData, most stand alone). The first meaningful line — after front matter and %% comments, both legal above a header — must be a bare header and nothing else. That admits all 24 header forms tested, including pie title A Very Long Descriptive Title, and rejects all seventeen non-diagrams including the three merman renders. Successful rendering is still required on top.

The gate applies only where there is no explicit tag. A block tagged bash or json is left alone even when its contents would parse: the tag is the author saying what they wrote, and overriding it is exactly how a shell transcript ends up drawn as a flowchart.

Diagrams in attachments

mermaid::find_in_text runs the same detection over an attachment’s extracted text, handling both a file that is a diagram (no fence — the case the header gate exists for) and a document that contains them (found by parsing as markdown, so fence lengths and info strings follow the same CommonMark rules as the transcript). Capped at MAX_PER_ATTACHMENT, with the cap reported to the reader rather than silently truncating.

The results ride on AttachmentView, which get_session builds on load and send_message emits as an attachments SSE event before the first token — so a diagram is on screen while the reply is still generating, and a reload is not what makes it appear. Cost is bounded by construction: attachment text doesn’t change while a reply streams, so this runs once per send and once per load, never per token.

The view also carries the extracted text, and the browser turns a chip into a collapsed disclosure holding it plus the diagrams. text: None — a binary or otherwise unreadable format — is what tells the client to render a bare chip with no expand control, so one is never offered with nothing behind it. The text is what the model received verbatim, already bounded by attachments::MAX_TEXT_CHARS which marks its own truncation inline, so displaying it adds no undisclosed cap.

This closes a real gap rather than adding a flourish. An attachment is otherwise invisible to its sender — the text goes to the model and the UI shows only a chip with the file’s name and size — so a diagram someone attached was the one part of their own message they could not see.

Fencing an attachment into the prompt

compose_content inlines an attachment’s text into the prompt as a fenced block, and the fence has to outgrow the body: attachments::fence_width returns one more backtick than the longest run inside it.

Three backticks is only safe for a body containing no fences of its own, and the documents most worth attaching do contain them. CommonMark ends a fenced block at the first fence at least as long as the opening one, so a 3-backtick wrapper around a Markdown file holding a ```mermaid block was closed by that file’s own closing fence — the rest of the document escaped the block, and the wrapper’s real closing fence went on to open a new, unterminated one. Measured on a real 200-line file, the model received four fence transitions where there should have been two, with the document split into fragments. Asked to render the diagram in it, the model described it instead. With the fence widened the same file parses back out of the composed prompt as exactly one code block with its mermaid fence intact.

Whether the model then re-emits the diagram is still the model’s call — nothing here can force that — but it is now working from an intact document rather than a scrambled one.

Putting the attachment’s diagram in the answer

Measured across four sessions against the same file, the answer to “Please, render this” contained a Mermaid fence zero times. The replies open with “Here is the rendered content” and then describe the diagram in prose — a correct explanation, and no picture. Only an explicit follow-up (“You have a Mermaid diagram”) produced a fence. Fixing the prompt fencing above did not change this; it is how models answer that request.

appendAttachedDiagramsToAnswer in app.js therefore appends the turn’s attachment diagrams below the answer. Two rules keep it honest:

Nothing is written into the message. The persisted content stays exactly what the model generated, which is also what Save-as-Markdown exports and what the next turn’s context replays — this is a presentation-layer addition, not a rewrite of model output. It runs on the done event (every token reassigns innerHTML, which would wipe an earlier append) and again on session load, so a reloaded answer carries the same picture a live one did.

PlantUML diagrams (web::plantuml)

web::plantuml is a clean-room, offline implementation of the commonly generated PlantUML UML families. It parses guarded @startuml documents, lays them out directly, emits theme-specific SVG, and rasterizes the same SVG to PNG with resvg. It never starts Java or Graphviz, downloads a jar, expands an include, or contacts a PlantUML server.

The parser deliberately accepts a smaller language than PlantUML itself: sequence, class/object/interface, component/deployment/use-case/state, and modern activity diagrams. Unsupported structural input returns None, so web::render leaves the original fenced block visible instead of presenting an incomplete diagram as if it were authoritative. Presentation-only syntax such as common skinparam blocks may be accepted when ignoring it cannot change topology. Source size, item count, output dimensions, and raster pixel count are capped before expensive work.

Each successful render produces four assets: light and dark SVG plus light and dark PNG. The streamed HTML contains short /api/diagrams/<sha256>/<asset> URLs rather than repeating base64 PNG data on every token. A 256-entry, 128-MiB LRU holds owned diagrams and cached failures; eviction drops the Arc and its SVG and PNG buffers instead of leaking them for the rest of the server process. Attachment discovery shares the same renderer, recognises plantuml, puml, and pu fences case-insensitively, preserves mixed Mermaid/PlantUML source order, and applies the common per-attachment cap.

Bundling a model into the binary (bundle.rs)

orangu-server bundle writes a new executable: this binary’s program image, byte for byte, then the model’s .gguf bytes, then a JSON manifest and a fixed 32-byte footer.

[ program image                       ]  base_len bytes, byte-identical
[ padding to a 4 KiB boundary         ]
[ shard 1 .gguf                       ]
[ padding, shard 2 .gguf, ...         ]  only a split model has these
[ manifest (JSON)                     ]
[ manifest_offset: u64                ]  ─┐
[ manifest_len:    u64                ]   ├ the footer
[ MAGIC:           16 bytes           ]  ─┘

The obvious alternative, include_bytes!, would put a multi-gigabyte array through rustc on every build, tie one binary to one model at compile time, and make a bundle something only whoever can build the project could produce. Appending instead makes the bundle a file operation on a finished binary, so anyone with one can make a bundle in seconds — and, because the manifest records base_len, a bundle can be bundled again, replacing its model rather than stacking a second one behind the first.

The footer being fixed-size and last is what makes “is this a bundle?” a seek and a 32-byte read at startup regardless of payload size — cheap enough that bundle::embedded() runs unconditionally on every start. A file whose footer matches but whose manifest doesn’t parse is reported on stderr and then treated as unbundled: a corrupt bundle should not take away the one thing that might still work, orangu-server <model> against a real file.

Alignment is why each shard starts on a 4 KiB boundary. The mapping is of the executable, not of a .gguf, and a model should not read differently for having been carried in one — a page-aligned start gives every tensor the same alignment relative to a page that it has in a file of its own.

Reading it back needed two small generalizations rather than a second load path:

At startup main::prepare picks a ModelSource — File or Embedded — and everything past it is written against the result. A bundled binary with no config file uses config::bundled_configuration: 127.0.0.1:8100, 127.0.0.1:8200, the Hugging Face hub cache as models, and the bundle’s own role. Loopback rather than all because a bundle is a binary somebody downloaded and ran, not a deployment somebody configured; --host all is how that gets opted out of for one run.

bundle records --host/--port/--web in the manifest as config::BundledListen, which bundled_configuration then layers over those defaults. Every field is Option and #[serde(default, skip_serializing_if)], so a bundle written before they existed parses unchanged and keeps exactly the behaviour it had — the manifest is a format other builds read, and adding a key to it must never be a reason an older bundle stops starting. The value is a default, not a lock: a run-time flag and a config file both still win over it. bundle validates the host itself (all/* or a literal IpAddr) rather than leaving it to the target machine’s bind, since the machine that would report the failure is not the machine that could fix it.

--host is also why ServerConfiguration carries web_host_explicit. The console follows the API’s address unless something says otherwise, so --host has to move it too — but a config that deliberately separated them (an API on the network, the console on loopback) must keep them separated, or exposing the API would be a way to expose the console by accident. The two addresses being equal cannot answer that question, so whether the key was written is recorded rather than inferred.

The bundle’s default output name, orangu-server-bundle-<arch>, comes from bundle::detect_target, which reads the architecture out of the binary being bundled — ELF e_machine, Mach-O cputype (a fat binary with more than one slice is universal), PE Machine, which also decides the .exe suffix. Reading the header rather than using std::env::consts::ARCH is what makes --binary honest: cross-bundling an aarch64 build on an x86_64 host has to produce a file named for the machine that can run it. An unrecognized format falls back to the host’s own architecture rather than failing — the name is a label, and nothing resolves against it.

Two interactions are worth naming. A handover (reexec.rs, below) that fails and falls back must not name the embedded model by its label — that is a Hugging Face repo id, and the fallback would go to the network for a model already inside the file it is falling back into — so bundle::EMBEDDED_SPEC ("bundled") is a reserved spec meaning “the model in this binary”, and that is what travels in FALLBACK_MODEL_VAR. And the model manager’s listing has no row for the embedded model, since it isn’t a file in the models directory: nothing is marked loaded and no Delete button exists for it, so CurrentView.bundled tells the panel to say bundled rather than leave an unexplained gap.

On macOS the copied program image is re-signed ad-hoc (codesign --force --sign -), and the ordering is the point: codesign writes the new signature at the end of the image it is pointed at, so it runs before a single payload byte follows, leaving the model outside the signed range where the kernel never looks. Signing can change the image’s length, so the length the image actually ended up with — not the base_len asked for — is what the manifest records and what a re-bundle truncates back to. ELF and PE images need none of this. Failure is noisy but not fatal: the file is written and correct, codesign is the one step that depends on the developer tools being installed, and the manifest is read back off disk afterwards, which is also what would catch a codesign that rewrote more than it was asked to.

Loading a different model (reexec.rs)

POST /api/models/select does not swap the model inside the process. It execves this binary again with the new model in argv, which means the model is loaded by main::prepare — the code that already runs at startup — and there is never a second load path to keep in step with it. Three properties turn that from a restart into a handover:

The listening sockets survive. Rust opens every socket SOCK_CLOEXEC, so Handover::exec clears FD_CLOEXEC on both listeners before the exec and names their descriptors to the new image in ORANGU_INHERIT_FDS (api:<fd>[,web:<fd>]). prepare calls reexec::adopt_or_bind instead of TcpListener::bind: given a descriptor it verifies still open (F_GETFD, so a number recycled after a failed handover can’t be adopted by mistake) it takes it, otherwise it binds. The port is therefore never released — a client connecting mid-load is queued in the listen backlog rather than refused. A 400-request probe across a live handover saw every request answered but the single one already in flight.

The process identity survives. execve keeps the pid, so a supervisor goes on watching the same process, and a --daemon server inherits its own already-detached session. Handover::argv therefore deliberately omits --daemon: passing it again would fork a second time and orphan the pid being watched.

argv is rebuilt from what this process resolved, not from what it was given: the workspace as an absolute path (a --daemon process has since moved to /) and the role as an explicit flag (it may have been answered at an interactive prompt). The flag is the next model’s role, not this one’s (reexec::role_for_model): a qwen_image file is always --image (Role::required_by, Role::fixed_by_model), which no language model can take — so a handover onto one says --image whatever this process is, and one off it back onto a language model says --all, nobody being at a prompt to choose. --config is passed only if it was passed to this process, so a server that found its config by the default search makes the new image repeat that search rather than pinning a path it never chose; --host/--port/--web follow the same rule (reexec::Listen). The address is normally moot on a handover — both listeners are inherited, so nothing is bound — but it matters in the one case the adoption check exists for, a descriptor that didn’t survive, where the new image binds instead and must bind where this server has been answering.

A failed load falls back. FALLBACK_MODEL_VAR carries the previous model spec and FALLBACK_ROLE_VAR the role it was working in (the failed image’s own flag was chosen for the model that failed, and may be exactly the one the fallback refuses); if the new image’s prepare fails, main execs once more with them and without the variables, which is what bounds the retry to one. This matters because the pre-check cannot be exhaustive: reexec::precheck reads the header and applies the same judgement as the SUPPORTED column (architecture resolvable, every tensor type decodable), but a GPU backend with no kernel for one of those types, or a model too large for the machine, can only be found by loading it. prepare binds its listeners after loading the model precisely so that case leaves the inherited descriptors untouched for the fallback to hand on again.

Both environment variables are read once, at the very top of main, by reexec::take_inherited, which also removes them — the same only-thread-that-exists-yet window that makes main’s own set_var("RUST_BACKTRACE", ...) sound. Nothing afterwards reads the environment for them, so a stale value can’t reach a child process or a second handover.

select answers 202 and arms the handover on a 300 ms timer, because execve leaves no “after” to answer from. The timer is best-effort UX, not correctness: a client whose connection is reset instead of receiving the 202 is looking at the same event, and its next poll lands on the new image either way. WebState::arm_handover allows one per process — there is only one process to replace.

[web].reexec (default true) and reexec::supported() (cfg!(unix)) gate the whole thing. When either is false serve builds no Handover and GET /api/models reports can_load: false.

Model manager (web::models)

Served on the web port: orangu-server list as the view, plus show, download and delete as the things that can be done from it. Each endpoint calls the same shared code the matching subcommand does — orangu::model_spec for the scan, grouping and delete, crate::format_show for the metadata dump, orangu::model_download for the fetch — rather than a second implementation that could drift from it.

That extends to the table itself. ModelView is one row of list, column for column, and carries the strings the CLI would print, not the raw numbers: quant already fell back to -, size has been through format_bytes, and supported is ModelSupport::cell verbatim (Yes (llama), No (llama, TQ1_0)) — which is why that method is pub. The client only decides layout. A row this build can’t load is greyed, an unreadable file’s error: replaces its last three cells, and a repo behind its Hub revision is marked, all exactly as format_groups does the same three things.

Two things shape the module beyond that:

The download runs detached. It takes minutes to hours, so POST starts a Job on a blocking thread and returns 202 immediately; the panel polls GET /api/models for its progress. ModelJobs holds one job slot — a second POST while one runs is refused with the name of the one holding it, since two fetches into one models directory would compete for the same disk and the same free-space check. A finished job doesn’t hold the slot but stays readable, so a completed download’s result survives a page refresh; DELETE /api/models/job clears it. There is no cancel: the worker is detached precisely so a closed browser tab doesn’t abandon a download part-way.

Progress comes from orangu::model_download::DownloadProgress, a sink the existing ProgressBoard publishes into. Passing one also turns all of the board’s printing off, interactive and logged alike: a running server has no terminal to draw an in-place-updating block on, and the per-file log lines would go somewhere nobody is watching.

The listing is cached, not re-scanned per request. GET /api/models costs a scan_models_dir plus a model_support pass, which between them open every GGUF header under the directory and every shard of every group — seconds on a directory holding a few dozen models. The panel polls once a second for download progress, and that progress is in memory, so ModelCatalog serves a cached scan and only rebuilds on ?rescan=true (the panel opening, its Rescan button) or after invalidate() (a delete, a finished download). Which row is loaded is not cached — that is about this process, not about the directory, so it is decided per request against WebState::model_path.

[web].delete (default true) gates removal the same way, reported as can_delete.

It gates models only. History’s own DELETE /api/sessions/{id} and DELETE /api/sessions (its per-row cross and Clear all footer) are unconditional, and GET /api/sessions carries no capability flag for the page to check. The two are not the same kind of thing: a model is a file on disk that a download, or a human with scp, put there, and a deployment can reasonably want that directory read-only while still allowing a model switch. A chat session is this console’s own scratch data, written by the page that is now asking to delete it — a console unable to clear its own transcripts is not a posture worth a config key.

Clear all does not spare an active session the way sweep_empty_sessions does. The console’s own current chat is one of the rows being cleared, and leaving behind exactly the one the user is looking at is not what the button says; the browser starts a fresh session immediately afterwards, so nothing goes on writing into a directory that just went away. Both delete paths also replace the on-screen transcript when what went was the session it belongs to — otherwise the next message would POST against an id that no longer resolves. And both stop generation first when a reply is still streaming into the session being removed: save_session on the stream’s done event recreates the directory, so the chat would otherwise reappear seconds after being deleted.

Both switches remove their button rather than disabling it. A disabled control with a tooltip is the right shape for something conditional — a handover already in flight, a model this build can’t load — where the same button works a moment later or on the row below. A capability the config has switched off is not a condition of any row; it is what this server does, and a column of permanently dead buttons explains less than their absence. When can_load is false the loaded row’s check mark goes too: which model is serving is already on the row, as the loaded badge beside its name.

Three things are refused rather than attempted:

These endpoints are neither authenticated nor loopback-restricted, matching the rest of the web port and the file-lifecycle API on the API port: the whole server assumes a trusted network.

File-lifecycle API (http::files)

The eight endpoints, their request and response fields, the Git integration and the error codes are documented in the HTTP endpoints chapter, under File-lifecycle API — one place for a contract three surfaces share (these endpoints, orangu’s file tools, and its typed commands of the same names).

What is worth saying here is where the code lives. The operations are orangu::files; the axum router that exposes them is orangu::files_http::router, mounted by orangu-server beside its inference endpoints and generic over a WorkspaceState that names the workspace root. Both are library code rather than server-binary code precisely because orangu’s own tools and typed commands call the same functions, so a tool call and an API request cannot drift apart in fields, defaults or errors. orangu-coordinator has no implementation of its own and proxies them.

Session activity tracking and prune (web::sessions, prune.rs)

save_session (called by both create_session and append_turn, so both creating a session and appending a turn to one trigger it) writes a second file alongside chat.json: session.json, recording this process’s own pid and — critically — its sysinfo::Process::start_time(). Recording pid alone would be enough as long as the writing process stays alive, but not once it exits: the OS is free to hand that same pid number to an unrelated later process, and without a way to tell the two apart, is_active would read the old session as still active forever. start_time is what closes that gap — a different process at the same pid almost never has the same start time down to the second, so a mismatch (or the pid not running at all) both read as “not active,” never as an error. mark_active’s own write is best-effort: a failure doesn’t fail the session save itself, since chat.json — already written by the time mark_active runs — is the data that actually matters; a session that never got a marker (or whose marker write failed) just reads as not active, the same as one from a build predating this.

is_active is read from an entirely separate process: orangu-server prune (prune.rs), a plain CLI invocation with no connection to whatever server process actually owns a session. That separation is the whole point — it’s what makes “keep track of which sessions are active” correct even for a session created long after some other still-running server’s own startup: is_active re-queries the live process table every time prune runs, rather than consulting anything cached or computed once earlier, so the answer is always current relative to this invocation, not relative to whenever the server happened to start.

prune itself needs no config file and loads no model — a pure filesystem operation against a fixed path, the same shape as system/suggest. Every invocation first calls sweep_empty_sessions (deletes every non-active session whose chat.json is empty, missing, or fails to parse — the last two read as “empty” too, so an interrupted-write leftover doesn’t linger forever uncleaned), then lists what’s left via list_sessions_for_prune (unlike list_sessions, the web UI’s History source, this includes zero-message sessions too — only ones is_active protected from the sweep, which prune needs to show, not hide) and hands off to one of three flows: no argument (prints the table, prompts for an NR or all), all (deletes every remaining non-active session, partition-ing active from inactive first), or a specific NR/id (resolved against the same listing). main.rs’s confirm — the same Yes/No stdin reader delete uses — is reused here rather than duplicated (pub(crate) in main.rs); prune’s own relative-time formatter (format_relative, “2h ago”) is hand-rolled rather than pulling in a date/time dependency, the same reasoning web::current_year already used for the copyright year.