* server: add an ssh transport to the tools runtime
--tools-runtime ssh:<target> runs the built-in tools on a remote host,
where target is whatever ssh already resolves, a user@host or a config
alias, so no credentials live in llama.cpp.
Only build_argv and upload differ from the docker transport: the remote
shell re-parses the command line, so the argv travels through
shell_quote_join, and files go over scp with the same quoting on the
remote path. Authentication is key-based and the host key must already
be trusted, since the tools run without a console and any prompt would
hang them.
The target is validated before use. The spec can reach us from the
x-tool-runtime header, and a leading dash would turn it into an ssh
option, which is enough to run a command back on the host.
Nothing is created and nothing is reclaimed, so an ssh spec goes
straight to the tool call instead of through the container runtime.
Note that this is remoting rather than isolation: the tools can do
whatever the target account can do, and the isolation is whatever runs
them on the far side.
* server: support podman in the tools runtime
docker and podman expose the same run, exec, cp and inspect verbs with the
same argument order, so a single implementation drives both and the engine
is carried by the spec prefix: podman:<image> and podman-container:<id> sit
next to the docker forms.
tools_io_docker becomes tools_io_container and the runtime spawner becomes
server_tools_container_runtime, both holding the client binary chosen at
parse time. A single parse_container_runtime() resolves every spec, so
adding another engine is one string in the table.
make_tools_io() now rejects the spawning forms. The spec also reaches it
from the x-tool-runtime header, which is client controlled, and only the
runtime that owns a container is allowed to create one: a tool call can
attach to a running container, nothing more.
* ./build/bin/llama-gen-docs
* server: simplify the tools runtime and drop the file copy step
A server_tools_runtime base with one virtual spec() replaces the
container runtime and the bare spec string that ssh needed next to it,
so server_tools is back to a single pointer and neither setup nor the
handler tests which of the two is set.
write_file used to spill its content into a temporary file on the host
and copy it in, because run_subprocess had no way to feed a child. It
now takes an optional stdin payload and creates the parent directory
and the file in a single round trip through a shell in the isolate.
That removes the upload virtual and both implementations: no more
container cp or scp, no second binary on the host, no sftp subsystem on
the target, no predictable temporary in a shared tmp, and none of the
content reaching an argv the remote shell re-parses. It also fixes
write_file over ssh, which never worked: scp speaks sftp and takes the
remote path literally, so quoting it kept the quotes in the file name.
Writing the payload before reading the output relies on the child
draining stdin as it goes, which holds for cat, its only user today.
* ./build/bin/llama-gen-docs
* server: harden the tools runtime against argv injection and a stdin stall
Validate the container id from x-tool-runtime and --tools-runtime the
same way the ssh target already is, so an id shaped like an option
(docker-container:--privileged) is rejected before it reaches the
engine's exec command line instead of running against a hardened
container. Feed the child's stdin after the watchdog is armed, so a
transport that stalls mid-write is terminated at the deadline rather
than blocking the request forever.
Cover both guards and fix the unknown-scheme test, which used ssh: as
its example and now names a real runtime.
* tests: exercise the tools runtime tests on podman as well as docker
Follow-up #26507. The container runtime drives docker and podman
through one implementation, so parametrize the availability helper,
the container fixture and the attach test on the engine, and cover
both engine prefixes in the container id injection test. Each engine
skips on its own when it is not installed.
The spawn cleanup test stays docker only: it recovers the spawned id
from the container hostname, which docker sets to the short id and
podman rootless does not guarantee. Podman keeps its coverage through
the attach path.
* server: release the container handle before respawning
Follow-up #26507. create() writes over the handle it is given, so a
respawn after the container died on its own leaked the pipes and the
process handle of the previous one.
* server: trim the tools runtime comments
* server: read tool output as raw bytes and harden the runtime on Windows
The stdout pipe is read with read() instead of fgets(), so a chunk
can hold any byte, including NUL, and still streams as soon as data
is available. Past the size cap the pipe keeps draining so the child
never blocks on a full pipe. Both pipe fds are forced to binary mode
on Windows, where the CRT defaults them to text mode and translates
line endings in both directions. Stdin is now always closed after
the feed: the child reads a deterministic EOF, and the Windows
docker and ssh clients stop outliving their command on a stdin pipe
that never closes.
The attach form of --tools-runtime has no lifecycle to own, so it
becomes a static target validated once at startup. This removes the
subprocess that ran on every tool call and
serialized calls behind a mutex; a stopped container now surfaces
the engine's own error at exec time.
The cidfile path is passed as UTF-8, matching the encoding the
subprocess layer expects for the CreateProcessW command line, so
the spawn form works from a non-ASCII Windows profile.
The SIGPIPE note in server.cpp now names the tools runtime children
as well as the MCP ones.
* clean up comments
* less pollute global scope
* nits
* tests: name the container image after both engines
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* Get started with Onyx
* Add architecture
* Skip keys handled in super()
* Loading tensors
* Shorten
* Graph
* Apply suggestion from @pcuenca
* Remove norm now embedding in transformers weights
* Add eot
* Explicit output_multiplier
* Handle post_norm_eps
* No super call; unhardcode eot.
The pattern `self._set_vocab_gpt2()` seems preferred throughout the
codebase, and it allows `set_vocab()` to be called from a different part
of the Python class hierarchy: the drafter model converter that we may
need eventually.
* Register for drafting
* DFlash: inherit rope type from the linked target.
Another option would be to store it in the gguf file itself.
* mmproj conversion
Note: some fields to be renamed after the implementation works. We are
keeping compatibility with the reference Meta gguf for testing purposes.
* "clip" header declarations
* Load mmproj
* Pre-processing
* Graph
* Go back to using delimiters.
Otherwise our generations are worse.
Transformers does not use them. We need to trace inputs to verify
whether they are equivalent.
* downsample_factor -> merge_size
* Add vision graph
lol, forgot from a previous commit
* Additional renames, align with llama.cpp / transformers
* Prefer _size instead of independent _h and _w
* Fix token layout
Co-authored-by: Young Han <younghan@fb.com>
* onyx: bring the chat parser onto the onyx branch
common/chat.cpp on this branch has no Onyx handling, so a converted model
serves malformed chat: the assistant preamble leaks into content
("to=self<|message|>...") and tool calls fail with
HTTP 500 "The model produced output that does not match the expected
peg-native format"
common_chat_params_init_onyx exists on onyx-fair-patch, added there by
8bb73dd3d. It was never on this branch, so this is not a regression --
the two lines developed independently.
The code here is taken verbatim from that commit. It is the clean side of
`git merge origin/onyx-fair-patch`: chat.cpp is one of the files that
merges without conflict. The full merge is not viable -- it produces 13
conflicts, including add/add on conversion/onyx.py and src/models/onyx.cpp
where the q_norm-folding and metadata-scale approaches contradict each
other, and #4/#7 are stacked on this branch's side of that.
Verified on this branch: builds with 0 errors, converts an Onyx checkpoint,
and serving it gives "4" for "What is 2+2?" plus a correct
get_weather {"city":"Paris"} tool call, where the unported branch gives the
two failures above.
No converter or runtime changes are included, so this should not interact
with the q_norm work.
Co-authored-by: Beto de Paola <betodepaola@meta.com>
* Less params, bilinear pos-emb interpolation as a graph op instead of CPU
* Map to symbolic V_MMPROJ instead of strings
* Make a couple params explicit
* Patchify via build_inp()
* No param for rope_theta
* Small cleanup
* Restore blank line
* Unpermute, to adapt to the latest transformers checkpoint
* Apply norm after token embeddings
This follows the latest transformers approach.
* Remove duplicated function
* build_vit
* onyx: use the model rope theta on sliding-window layers
* DFlash: conversion from transformers drafter
* Revert rope_type derivation from target
NOTE: this breaks compatibility with Meta's distributed DFlash GGUFs, as
the Q/K are stored in "NEOX" (rotated half) format, like in
transformers.
* Apply suggestion from @pcuenca
* Set model type
* Remove comment that will become obsolete
* Hardcode post_norm_rms_eps instead of new param
* Derive SWA+RoPE pattern from gguf array or scalar
* Fix model type <-> number of layers
* Reorder
* Rename
* Fix typo
* DFlash: seed the draft KV cache from multimodal embedding batches
`common_speculative_impl_draft_dflash::process()` returned early on any batch carrying embeddings, so an image prefill never had its target-layer features fused through the DFlash encoder and injected into the draft's KV cache. That left a hole spanning the image's positions, and the next injection at a post-image position failed to initialize its batch:
```
decoding image batch 1/1, n_tokens_batch = 256
decode: failed to initialize batch
llama_decode: failed to decode, ret = -1
process: llama_decode(ctx_dft) failed rc=-1 (n_tokens=17, offset=0)
srv decode: failed to process speculative batch
```
Every image request with `--spec-type draft-dflash` failed with HTTP 500. Text-only was unaffected, since those batches carry token ids and were let through.
Restore the earlier condition, which admits a batch that is either tokens or embeddings and skips only the degenerate neither/both cases. The rest of `process()` is already layout-agnostic -- it gathers features via `llama_get_embeddings_layer_inp()` and indexes `batch_in.pos[]` / `batch_in.seq_id[]`, none of which assume token ids -- so this is the whole fix.
Validated against `muse-glimmer-30B-bf16.gguf` + `mmproj-muse-glimmer-30B-bf16.gguf` + a DFlash draft head, on an image describe-the-shapes request:
- before: HTTP 500, `failed to process speculative batch`
- after: HTTP 200, draft acceptance 0.34012 (167 accepted / 491 generated), mean len 3.04
Output equivalence holds, which is the property that matters: at temperature 0 the drafted response is byte-identical to the same request served with no draft attached (1213/1213 chars), so the draft is drafting correctly through the image context rather than merely not crashing.
* Conversion: prefer rewrite to mapping
* Revert "Conversion: prefer rewrite to mapping"
This reverts commit a92d0ac584.
* fix lint
* sliding_window metadata is not optional
* disable state save/load
* Apply suggestion from @pcuenca
---------
Co-authored-by: Young Han <younghan@fb.com>
Co-authored-by: Beto de Paola <betodepaola@meta.com>
Co-authored-by: Daniel Han <michaelhan2050@gmail.com>
Co-authored-by: ruanrms <ruanslv@gmail.com>
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Upstream merged the Windows argument quoting fix, the NetBSD build
fix and the chdir fallback for glibc older than 2.29, so pin the
vendored copy to a commit that carries all three and remove the
patch files along with the apply step in the sync script.
The new pin also brings the exec error report on glibc older than
2.24 and the ENOSYS mapping to a dedicated error code. Both are
additive and no caller inspects those values.
* Restore quantization of mmprojs
This was lost in the refactor undertaken in #22004.
* add noreturn
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
* ci: Add support for CUDA 13.4 ARM64 builds for Windows
Added an architecture-specific CUDA 13.4 Windows build entry targeting ARM64.
Added a CMake configuration to enable ARM64 CUDA cross-compilation from an x64 Windows environment using the x64-hosted CUDA and MSVC toolchain while linking against the ARM64 CUDA import libraries to produce ggml-cuda.dll.
Validated the self-hosted Windows x64 workflow, including toolkit acquisition, CMake configuration, ARM64 CUDA cross-compilation, and packaging. Runtime validation was performed separately on a native ARM64 RTX Spark system using TinyLlama 1.1B Q4_K_M to verify the generated binaries.
The ARM64 CUDA job builds only the ggml-cuda.dll backend (LLAMA_BUILD_SERVER=OFF). The release consists of two packages: the main ARM64 release package, which combines the existing ARM64 CPU outputs with ggml-cuda.dll, and a separate runtime package containing the required CUDA runtime libraries (cudart64_13.dll, cublas64_13.dll, and cublasLt64_13.dll).
The CUDA 13.4 setup uses NVIDIA Developer Preview component archives instead of the GA component downloads used by the existing CUDA setups and will require updates once CUDA 13.4 reaches GA.
* ci: cleans up to align with x64 CUDA setup
- Moves CUDA-specific CMake options into matrix defines.
- Keeps the CUB 3DOT2 option only for CUDA 12.4.
- Removes runtime argument construction and the unnecessary server option.
- Aligns ARM64 CUDA runtime packaging with the existing robocopy approach.
- Generalizes the ARM64 release label from CUDA 13.4 to CUDA 13.
* ci: Set CUDA job name as version-architecture pair
* mark as preview
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* granite-switch: add llama.cpp backend (POC, CPU)
New "granite-switch" architecture: a dense, all-attention Granite-4.1
model with N embedded LoRA adapters selected per-token by control tokens.
- gguf-py schema (arch, KV keys, stacked LoRA tensor names) + writer helpers
- conversion/granite.py: GraniteSwitchModel converter (stacks N adapters +
zero base slot into per-projection A/B tensors; emits switch metadata)
- C++ arch registration (llama-arch.{h,cpp}, llama-model.{h,cpp})
- src/models/granite_switch.cpp: load + per-token switched-LoRA graph via
ggml_mul_mat_id over stacked tensors; sticky per-token index + control-token
substitution in llm_graph_input_switch::set_input
- llm_graph_input_switch in src/models/models.h
Runs end-to-end on CPU: convert 3b checkpoint (842 tensors, stacked dim 13)
and generate on both base and control-token paths. Sticky switch state is
single-sequence (POC); full multi-sequence machinery is a follow-up.
* granite-switch: add Mac (Metal) build + mid-sequence switch demo script
Self-contained script to build llama.cpp on Apple Silicon (Metal),
convert the composed 3b checkpoint, and run the crisp mid-sequence
adapter-switch demos verified on Vela:
- answerability: <|answerability|> mid-seq -> "unanswerable"
- query_rewrite: <|query_rewrite|> mid-seq -> {"rewritten_question": ...}
Each demo runs the same prompt twice, differing only by a control token
placed before the assistant turn, so the per-token switch is visible.
* granite-switch mac demo: add -no-cnv so each run is one-shot
The composed model ships a chat template, so llama-completion auto-enables
interactive conversation mode and halts at a `>` prompt after generating,
stalling the script. -no-cnv disables conversation mode: generate once from
the raw prompt and exit (also prints special tokens, making the switch visible).
* granite-switch: replace global sticky index with in-graph router attention
The POC computed the per-token adapter index on the CPU and carried it
across ubatches in ONE global `mutable int32_t poc_sticky_index`, reset
only when a ubatch contained sequence position 0. That global had two
problems:
1. Concurrency: with multiple sequences in a batch it was last-writer-
wins — one sequence's adapter leaked into the others.
2. Multi-turn: an interactive `ollama run` chat continues one KV cache,
so turn 2 never saw position 0 and the index never reset — the
adapter stayed stuck on across turns.
Port the vLLM/HF backend mechanism faithfully: a single-head causal
"router" attention recovers the adapter index in-graph. Per token, only
dim 0 carries signal — Q[0]=1, K[0]=+gain for a control token / -gain
otherwise, V[0]=adapter slot / 0 — and the causal softmax over the single
visible control token recovers that adapter's slot (readback =
clamp(round(V[0]), 0, n_adapters)). gain=15 matches config.py and is
F16-safe (no F32 cache).
The router's K/V live in the model KV cache at an extra layer
R == hparams.router_layer (== n_layer). We bump n_layer_all to n_real+1
so the cache allocator gives the router its own per-sequence slot, and
set n_layer_nextn=1 so n_layer() stays n_real — the decoder loop and
tensor loading are untouched and never reference layer R. The router K is
exempted from the k-shift RoPE loop (its dim-0 value is a literal
magnitude, not a rotation).
Because the selection now lives in the per-sequence KV cache, CONCURRENT
requests are isolated for free (problem 1 fixed; verified by
scratch/concurrent_switch_test.cpp). set_input becomes stateless pure
per-token maps; the global is gone.
Single-switch contract / known limitation, identical to vLLM & HF: the
gain is flat (no recency), so within one sequence there is no mechanism to
revert to base mid-sequence — once an adapter fires it stays on until that
sequence ends (problem 2 is therefore NOT fixed by a faithful copy; vLLM/HF
avoid it only because each served request is a fresh sequence). A client
continuing one KV cache across turns must start a fresh sequence per turn,
or opt into a recency-biased router (a deliberate divergence, not done
here). Documented in granite_switch.cpp and asserted by
scratch/multiturn_leak_test.cpp.
Verified (CPU): both demos unchanged (answerability -> "unanswerable",
query_rewrite -> rewritten query); concurrent two-sequence isolation
passes; multi-turn carry-over matches the vLLM/HF contract.
* granite-switch: drop scratch tests and mac demo for upstream PR
Remove the local-only development artifacts that should not ship in the
upstream PR:
- granite-switch-mac-demo.sh (local Metal build + demo driver)
- scratch/concurrent_switch_test.cpp
- scratch/multiturn_leak_test.cpp
Also drop the now-dangling reference to the scratch tests from the
granite_switch.cpp header comment. Leaves only the core architecture
support (conversion, gguf constants, llama-arch/model/kv-cache, and the
granite_switch graph).
* granite-switch: trim comments to match native llama.cpp style
* granite-switch: trim conversion comments to match native style
* granite-switch: drop unused adapter_ranks metadata
* granite-switch: rename arch to graniteswitch and drop obid alias
* granite-switch: fix non-ASCII comments and document router gain assumption
* granite-switch: drop section comments from constants.py to match native style
* granite-switch: add functional tensor block comments matching Granite4 Vision style
* granite-switch: clarify n_expert_used comment
State the actual constraint: mul_mat_id needs n_expert_used == 1, and
since the GGUF carries expert_count = 0 the generic loader's
n_expert == 0 => n_expert_used == 0 assertion has already passed by the
time load_arch_hparams runs, so it is forced to 1 here.
* granite-switch: note n_layer_nextn reuse has no MTP
The router carving reuses n_layer_nextn, normally the MTP/next-token
count. Clarify in the comment that it is borrowed here purely as the
trailing-layers lever and that there is no MTP head, to spare readers
the double-take.
* granite-switch: rename source file and apply review nits
* granite-switch: don't force LoRA tensors to F16, follow --outtype instead
* granite-switch: drop redundant _permute_qk wrapper, call LlamaModel.permute directly
* granite-switch: read router gain from GGUF (control_token_gain) instead of hardcoding 15.0
* granite-switch: derive n_slots()
* granite-switch: move llm_graph_input_switch into granite-switch.cpp
* granite-switch: cut AI-style narration comments
* granite-switch: collapse multi-line comments
* granite-switch: rename control_token_* maps to adapter_token_*
* granite-switch: cut noise comments
* granite-switch: rename embedded LoRA tensors to <base>.lora_a/lora_b
* granite-switch: GGML_ASSERT token input to avoid UB on embeddings
* granite-switch: TODO for raw embedding input support
* granite-switch: collapse LoRA tensor constants to .lora_a/.lora_b suffix
* granite-switch: drop n_expert_used hack, guard mul_mat_id buft probe
* granite-switch: stop forcing dense expert counts, read from config
* granite-switch: renamed control_token_gain metadata key to router_gain
* granite-switch: trim header comments to match native style
* granite-switch: collapse LoRA tensors to base name + suffix
* granite-switch: inline suffix checks in tensor op resolution
* granite-switch: drop switch-lora struct comment
* granite-switch: guard router layer index and inline n_slots
* granite-switch: group adapter metadata under {arch}.adapters.* namespace
* granite-switch: add hparams.has_rope(il) for KV-shift rope skipping
* granite-switch: skip arch in test-llama-archs (adapter fixture missing, TODO)
* granite-switch: Keys.Adapters namespace + simplify n_slots
* granite-switch: validate substitute token ids against n_vocab
* granite-switch: bound adapter count and lora rank from GGUF
* granite-switch: reject MTP context type when router_layer is set
* granite-switch: throw on bad adapter metadata instead of GGML_ASSERT
* granite-switch: use ASCII +/- in router K signal comment
* granite-switch: document n_layer_nextn repurpose and its leak points
* granite-switch: gate lora_a/lora_b op mapping on router_layer
* granite-switch: label all three preview model sizes
docker info only proves the daemon answers, so the Windows CI passes
the check and then dies trying to run a linux image. The hosted
Windows runners cannot run one: GitHub states the VMs are not enabled
for nested virtualization and will not be, since they already sit one
level deep and the hypervisor does not support more levels
(https://github.com/orgs/community/discussions/25491). Probing the
image itself skips those tests there, and pulls it before the server
waits for the container id.
The saver called add_kv with LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH twice, the
second time passing n_ff_chexp. gguf_set_val_u32 removes-then-appends, so the second
call clobbers the first: the saved shared_feed_forward_length ends up as n_ff_chexp
(0 for every arch except GroveMoE), and expert_chunk_feed_forward_length is never
written at all.
So a save->load roundtrip of any MoE model with a shared expert loses n_ff_shexp. On
reload the arch falls back to n_ff for the shexp tensor shape, that no longer matches
the saved tensor, and the model FAILS to load. Hits qwen2moe, qwen3-next, granite-moe,
hunyuan-moe, ernie4.5, bailingmoe2, nemotron-h, and the other shared-expert MoEs.
Fix: the second call writes LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH.
test-llama-archs: set expert_shared_feed_forward_length to a value distinct from n_ff
in the MoE setup so the roundtrip exercises it. Without the fix the reload fails on a
shexp tensor-shape mismatch; with it, every arch roundtrips clean.
* Update build-sanitize.yml
* make it run on pr
* fix thread
* Update build-sanitize.yml
* Update build-sanitize.yml
* just run thread on github machine
The picker mounts whenever a cwd-aware builtin tool is enabled, so
it can open while file_glob_search is not served or was disabled by
the user. Every typed query then fired a search that could only
fail with a raw error.
Gate the debounced search on the tool state, the same way the
mention picker does, and show a message in place of the results
list that explains why search is unavailable. Manual entry with
Enter still commits a directory. The Browse button and the search
scope footer are hidden as well: Browse resolves the picked folder
name through file_glob_search, and the client-side toggle would not
stop that call.
* server: report the isolate working directory from get_info
Without an explicit cwd, get_info fell back to the server process
working directory even when a tools runtime was configured. That named a
host path no tool would ever run in, since an isolate starts in a
directory of its own.
It now asks the isolate for its working directory in that case, and
keeps the process one only when the tools run on the host.
* remove redundant comment
---------
Co-authored-by: Xuan-Son Nguyen <thichthat@gmail.com>
The working directory chip showed up as soon as the server exposed any
builtin tool, so a server started with just get_datetime, or a user who
turned every filesystem tool off in the settings, still got a control
that nothing would read.
Tools now declare whether they resolve their paths and run against the
working directory, next to the write permission they already publish in
the /tools listing. The WebUI shows the chip and enables the /cwd
command only when at least one such tool is both served and left
enabled.
get_output runs the waveform work the pipeline defers to it, from a
single trailing window to a full pass depending on the model. Measuring
it keeps the reported total and the audio to process ratio honest.
* feat: Add contenteditable tokenizer for badge/code-chip chat input
* feat: Add source-space undo/redo history for the rich input
* feat: Split text glued to a closing code fence onto its own line
* feat: Add ChatFormContenteditable rich input renderer
* feat : wire the contenteditable into ChatForm with auto-switch gating
* base : slash-command/misc foundation - model icon and focus-selector constants
* feat : slash-command picker and command parsing helpers
* refactor : wire command and @-mention pickers into the chat form
* ui : improve model selector keyboard navigation and load/dismiss
* feat: Unify markdown/raw-text rendering under one setting with migration
* fix: Misc fixes - tool-call subtitle, assistant wrap, progress guards
* feat: Clamp and style numeric settings inputs from registry bounds
ggml_metal_op_norm sized the threadgroup with
`nth = std::min(nth, args.ne00_t)`, which can leave nth not a multiple of
the simdgroup size. The kernels finish their row reduction with a
cross-simdgroup step where each lane of the last simdgroup reads one
per-simdgroup partial sum out of shmem_f32:
if (tiisg == 0) { shmem_f32[sgitg] = sumf; }
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
When the last simdgroup is partial it has fewer lanes than the
threadgroup has simdgroups, so the tail of the partial sums is never
read and the row sum is too small. For ne00_t = 33 nth becomes 33: two
simdgroups, but only one lane in the second, so one of the two partial
sums is dropped. The mean and variance are then wrong for the whole row.
Round ne00_t up to a whole number of simdgroups instead. Rounding up
rather than dropping the clamp keeps the threadgroup as small as
possible: deleting the line would raise nth to the next power of two
(ne00_t = 544 -> 1024 instead of 544), which costs idle lanes on 26 row
lengths below 8192 that were already correct, including 1536 and 3584.
GGML_OP_NORM is affected as well as GGML_OP_RMS_NORM - both dispatch
through ggml_metal_op_norm.
No mainstream LLM hidden size hits this: ne00_t is ne00/4 on the
vectorized path, so 4096, 8192, 2048 and friends all give a multiple of
32. It is reachable from other norm shapes, e.g. 320-channel norms.
Add NORM and RMS_NORM cases for ne0 = 33, 132 and 260 across the
existing eps values. 33 exercises the scalar path and 132/260 the
vectorized one, since only those divide by 4.
Before, on M3 Pro:
test-backend-ops test -b MTL0 -o NORM 25/50
test-backend-ops test -b MTL0 -o RMS_NORM 26/51
After:
test-backend-ops test -b MTL0 -o NORM 50/50
test-backend-ops test -b MTL0 -o RMS_NORM 51/51
test-backend-ops test -b MTL0 13943/13943
* base : @-mention picker foundation - glob search, picker nav, highlight
* feat : @-mention file/folder picker and mention badges in message bubbles
* fix: Imports
* feat : wire the @-mention picker into the chat form
* fix: Bound the glob-search result cache key and prune stale entries
* webui: load the model selected via ?model=
Opening the WebUI with ?model= selects the model but doesn't load it. The load only starts when you send your first message, so you wait for it then.
This loads it as soon as the page opens, while you're still typing your prompt. It's what the model dropdown already does, and it isn't awaited, so the UI still works while the model loads.
This is the path the Llama macOS app uses to open the WebUI, so it's a common way in.
* webui: gate the load behind ?load=true
Loading on landing is opt-in, so a plain ?model= link behaves as before and doesn't allocate memory on its own.
* webui: name the chat URL params
Collects the query params the chat routes read into a URL_PARAMS constant, instead of repeating the literals across three files. NEW_CHAT_PARAM folds into it.
The reference implementation has two mutually exclusive prompt layouts.
In non streaming mode the prefill carries the whole utterance text plus
tts_eos summed with codec_pad, and the trailing text hidden collapses to
a single tts_pad row. In streaming mode the prefill carries only the
first text token and the trailing rows stream the rest of the text
followed by tts_eos.
The pipeline built the non streaming prefill but the streaming overlay,
so the talker saw the utterance a second time during generation and read
it twice before emitting codec_eos.
The overlay is now the single tts_pad row that matches the prefill.
* ggml : add fallback definitions for missing aarch64 HWCAP bits
* ggml : require HWCAP_ASIMDHP for the aarch64 fp16 cpu variants
Also rename has_fp16_va to has_fp16, the field gates the whole FEAT_FP16
extension, scalar and vector half-precision arithmetic together.
* ui: read model modalities from the router model list
The router advertises input modalities for every model, loaded or not.
Reading them at list build time lets the UI accept image and audio
uploads for a model selected through ?model=, which has no /props yet.
* enum
The NVFP4 quantization format stores a scaling factor for every group of
16 weights, packed into a single UE4M3 byte.
The SYCL GPU code was converting these scale values using the E4M3 path,
but that's *signed*, and these are unsigned values.
* tests: add SWIGLU perf cases
perf mode had no GLU coverage. Adds SWIGLU at 17408 columns, 512 and
2048 tokens, f16 and f32, with the operands both fused and split.
* sycl: consolidate fused-GLU kernels
They differed only in which op_* they called, so take the op as an argument and share a common launcher.
Their block sizes were all 256, so launch geometry is unchanged;
SYCL_GELU_BLOCK_SIZE and SYCL_SILU_BLOCK_SIZE lose their last users so are dropped.
* sycl: contiguous fast path for the fused GLU ops
o0 == n and o1 == n collapse the de-interleave index math to the
identity, so dispatch a flat kernel in that case. It fires for
ggml_glu_split with packed operands; a fused [gate|up] tensor keeps the
strided path. test-backend-ops perf -o SWIGLU on an Arc Pro B70: split
+14% f16 and +4% f32, fused unchanged.
