Add the Models Discover explorer behind a new sidebar action and full-screen
dialog: a searchable HuggingFace GGUF list (org avatars, capability icons,
catalog size ranges) and a detail pane with header, metadata chips, the
sanitized model-card readme and the download area - per-quant action chips
grouped by bit depth, sidecar queuing and a scrollable serve-command preview
with dynamic quant selects. Searchable dropdowns gain sticky headers/footers.
Assisted-by: pi:GLM-5.3-Flash
Extract ModelCapabilityIcons (canonical Tools/Reasoning/Vision/Video/Audio
order) out of ModelId and reuse it there, add the shared DialogConfirmDownload
for destructive download actions, the discover org avatar with dark-mode
inversion and the thin download progress bar, and rework ModelId badges to take
thinking/tool-use support directly.
Assisted-by: pi:GLM-5.3-Flash
Add an allowHtml prop to MarkdownContent: raw HTML found in the markdown is
rendered after DOMPurify sanitization instead of being escaped as literal
text. Default stays escaped.
Assisted-by: pi:GLM-5.3-Flash
Track HuggingFace downloads end to end: the server download/cancel endpoints,
a status manager fed by the /models/sse download progress events, and a
models-discover store holding the catalog and detail state for the discover
view. Downloaded and in-flight entries are excluded from the loadable model
list.
Assisted-by: pi:GLM-5.3-Flash
Replace the raw runtime-memory estimate with the app's compatibility check:
the smallest real Mac memory tier that fits a model file, budgeted as
RAM x 0.75 minus fixed overhead with headroom on the file size. The constants
move to lib; the unused runtime-memory estimate is dropped. browser-info's
OS detection is exported for reuse.
Assisted-by: pi:GLM-5.3-Flash
Add HuggingFaceService and its constants/enums/types: GGUF repo search, file
tree and model detail fetching, quant/sidecar filename analysis, shard-set
collapsing and the llama.app catalog feed, plus an orgOf() helper on the model
name utils.
Assisted-by: pi:GLM-5.3-Flash
Extend the shared model id parser with sidecar tokens (draft variants and
auxiliary imatrix/mmproj files), weight-file and custom-quant regexes, and add
the tools capability to ModelCapabilities; the selector option row picks it up
from the model's declared capabilities.
Assisted-by: pi:GLM-5.3-Flash
The download monitor thread acquires the mutex on its way out, so joining
it while holding the lock in server_models::remove deadlocks once the
status has flipped to DOWNLOADED. Join outside the lock, same pattern as
load_models().
Assisted-by: pi:zai-org/GLM-5.3
A Q4_0-mtp style tag now resolves the sidecar file when no model file matches it, so a solo draft or mmproj download actually pulls the file. Cached sidecar files list as their own entries so the state survives a restart, and removing such a tag deletes only the sidecar.
Assisted-by: pi:zai-org/GLM-5.3
The server caches the most recent compute graph per device so that
GRAPH_RECOMPUTE can re-execute it without resending tensor data. The
cached graph nodes hold direct pointers to backend buffers that were
live at graph_compute() time. If any of those buffers is later
released via FREE_BUFFER, the next GRAPH_RECOMPUTE re-executes the
cached graph through the dangling pointers (use-after-free).
The bug is reachable by an unauthenticated remote client. The
dangling pointers point into chunks an attacker can reshape via
subsequent ALLOC_BUFFER/SET_TENSOR commands, and the resulting
read/write through the cached graph is sufficient to leak libc
addresses and hijack the buffer iface vtable used by BUFFER_CLEAR,
yielding remote code execution.
Discard all cached graphs in free_buffer(). The existing null-check
in graph_recompute() then rejects the request and the client falls
back to GRAPH_COMPUTE on the next call.
No protocol or API change.
It's found the MoE ncols_opt tile heuristic needs to be broadened
to include the RDNA3.5 architecture.
The code change is implemented in ggml/src/ggml-cuda/mmq.cu
and just change the GGML_CUDA_CC_IS_RDNA3_0 to
GGML_CUDA_CC_IS_RDNA3 in the condition.
The dense dispatch logic remains unchanged.
The Test machine configuration we used is
AMD Radeon 8060S, gfx1151 (RDNA3.5), 20 CU, wave32
+ AMD Ryzen AI MAX+ 388, 8C/16T, 23.79 GB RAM
we complete the Correctness verification and performance evaluation as follows:
test-backend-ops test -b ROCm0 -o MUL_MAT -p type_a=<q4_K|q5_K|q4_0|q5_0>
test-backend-ops test -b ROCm0 -o MUL_MAT_ID -p type_a=<q4_K|q5_K|q4_0|q5_0>
all pass: MUL_MAT 64/64, 29/29, 48/48, 14/14;
MUL_MAT_ID 84/84, 3/3, 74/74, 3/3
Performance result on target machine:
LFM2.5-8B-A1B-UD-Q4_K_M (Q4_K MoE) +16.198% [+12.704, +19.799] 8/8
Qwen1.5-MoE-A2.7B-Q2_K (Q2_K MoE) +6.189% [ +5.245, +7.141] 8/8
pooled (16 pairs) +11.081% [ +7.972, +14.279] 16/16
Token generation (tg128) is unchanged on the Q4_K MoE model and +2.188%
[+0.905, +3.488] on the Q2_K one.
Use BN/2 as the default for BNover2 and as the disabled fallback for BNover4, and remove the enable gate from the MUL_MAT_ID BN/2 branch. The BN/4 branch remains gated by enable_smaller_matrices, while the p.N path is unchanged.
* test-backend-ops: reproduce MUL_MAT_ID NaN for activations beyond f16
The Metal mul_mm_id path narrows src1 to `half` for the simdgroup MMA
(`S1 = half` in every instantiation; ggml-metal.metal:10582 and :10595,
mirrored at :10643/:10654 in the tensor-ops path). f16 saturates at
65504, so a model whose activations exceed that produces inf, and
`simdgroup_multiply_accumulate` then turns the whole 8x8 accumulator
tile into NaN. The mul_mv_id path used below `ne21_mm_id_min` (32)
carries the same values in f32 and is correct, as is every CPU path.
This was untestable before: `init_mul_mat_id_tensors` initializes
uniform [-1, 1], so no existing case can drive an operand out of f16
range. `test_mul_mat_id` gains an `amax` parameter (default 1.0f,
preserving the historical init exactly) that scales only the f32
activations, leaving the quantized weights in their normal range.
Six cases: n=16 sits below the mul_mv_id -> mul_mm_id switch and is the
control that must stay green; n=32 and n=64 are above it and fail on
Metal today. Two shapes, because this is not model- or size-specific —
q4_K at 128 experts / 4 active / 4096x2048 mirrors a real model, and
q8_0 at 8 experts / 2 active / 512x256 shows the same failure at
minimal size.
Observed on Apple M2 Max, macOS, llama.cpp b10156:
MUL_MAT_ID(type_a=q8_0,...,n=32,k=256,amax=100000.000000):
[MUL_MAT_ID] NaN at index 0 (MTL0=nan CPU=583442.375000) FAIL
The real model behind this is Mistral Small 4 (arch mistral4, 128
experts / 4 active), one of whose layers reaches ~1e5 activations: on
Metal every prefill of >=32 tokens returns an entirely NaN vocabulary,
while <32 tokens is correct.
Note kernel_mul_mm (dense) has the identical conversion at :10273 and
:10286 and is expected to fail the same way; it is not covered here.
Found and written by Claude Opus 5 (via Claude Code).
* metal: fix NaN in mul_mm_id when activations exceed f16 range
kernel_mul_mm_id narrows src1 to `half` for the simdgroup MMA operands
(`S1 = half` in every instantiation). f16 saturates at 65504, so a model
whose activations exceed that produces inf on load, and
simdgroup_multiply_accumulate then propagates NaN across the whole 8x8
accumulator tile. The result is an entirely NaN output — not a precision
loss, a total loss. The mul_mv_id path taken below ne21_mm_id_min (32)
keeps the same values in f32 and is correct, as is every CPU path, so
the same model produces correct logits for short inputs and NaN for
long ones.
Fix: rescale src1 by a power of two so it fits, and undo the scale on
the f32 accumulator at the store. A two-stage reduction computes
max(|src1|) and writes the pair (1/scale, scale) into scratch chained
off the destination buffer, in the same style as the existing tpe/ids
id-mapping scratch. The matmul multiplies on load and on store.
This is exact, not approximate, for two reasons: the dot product is
linear, so one tensor-wide factor commutes through the accumulation;
and the factor is a power of two, so both multiplications are exact in
binary floating point. When max(|src1|) already fits — every model that
works today — the factor is exactly 1.0 and the output is bit-identical
to before. Accumulation was already f32 and is unchanged; only the
operand narrowing was ever the problem.
The reduction is two-stage (256 threadgroups into partials, then one
threadgroup folding them) specifically so it stays bandwidth-bound. A
single-threadgroup version was measured first and cost up to +451%
median on prefill — the scan serialized against an otherwise idle GPU.
It is also dispatched only on the mm path, so decode never pays for it.
Measured on Apple M2 Max, `test-backend-ops perf -o MUL_MAT_ID -b MTL0`,
99 cases, versus the same build without this change:
n=1/4/8 (mul_mv_id, decode) : -0.8% / -0.8% / -0.4% median (noise)
n=32 (mul_mm_id, prefill) : +1.73% median
n=64 : +1.30% median
n=128 : +1.80% median
n=256 : +3.98% median
n=512 : +3.74% median, +7.20% worst
overall : +1.14% median
Correctness, same machine:
- the six new test-backend-ops cases go from 4 FAIL / 2 OK to all OK,
with the n=16 controls (mul_mv_id path) unchanged;
- `test-backend-ops -b MTL0` full run: 0 failures, no regression;
- Mistral-Small-4-119B (arch mistral4, 128 experts / 4 active) now
generates correctly at the default n_ubatch of 512, in both
UD-IQ3_S and UD-Q4_K_XL quantizations. Before this, every prefill of
>= 32 tokens returned an all-NaN vocabulary and only n_ubatch <= 31
(forcing the mul_mv_id path) worked.
Likely fixes#25722 (mistral4 empty output on Metal above ~300 tokens,
FA on and off, generation degenerating to a single control token — the
signature of argmax over an all-NaN distribution). #20668 may be the
same defect attributed to a bad GGUF.
Note kernel_mul_mm (dense) has the identical narrowing at the
corresponding load sites and is expected to fail the same way; it is
left alone here to keep this change reviewable. Also possible, and left
for later: scaling per output column rather than per tensor, which
would preserve more precision when a single token is the hot one.
Found, diagnosed and fixed by Claude Opus 5 (via Claude Code).
* metal : make requested edits
- remove verbose comments
- explain rationale as requested
Generative AI disclosure: Claude made the edits as requested.
* metal : stack mul_mm_id map0 with amax_part
Implement @ggerganov suggestion to stack amax_part + map0. Mean 2.6% faster (worst -0.7%, best -4.1%). Win grows with batch size. Benchmarked on a hot M2 Max after reboot.
Generative AI disclosure:
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* cont : fix var scope
* cont : comment out tests temporarily
Comment out tess to not break CI temporarily
Assisted-by: Claude Fable 5.1
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* add self-hosted vulkan and webgpu to hf-jobs
* try t4-medium
* cont : adjust cpu backend threads
* try t4-small again
* restore cm jobs
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* rpc : hash-cache only weights
ggml_backend_rpc_buffer_set_tensor and ggml_backend_rpc_set_tensor_async
hashed every transfer above HASH_THRESHOLD and let `rpc-server -c` serve it
from its file cache. The cache is meant for weights, but the activations
ggml_backend_sched copies between backends took the same path: with a
two-node split of Qwen3.8-Flash-Next every prefill ubatch above 10 MB was
hashed, written to the worker's cache directory (1.4 TB after a day) and
later served from there. Use the hash path only for tensors in buffers
marked GGML_BACKEND_BUFFER_USAGE_WEIGHTS.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
* rpc : save a cache entry only for the tensor that missed the hash check
With the client hashing weights only, the server still wrote every
SET_TENSOR above HASH_THRESHOLD to the cache directory, so the compute
data the scheduler sends kept filling the disk. Remember the hash of the
last SET_TENSOR_HASH that missed and save only the SET_TENSOR that
follows it with that hash - the weight the client is re-sending.
* rpc : signal the cache decision in the SET_TENSOR payload
Replace the server-side `pending_cache` state with a `cache_flag` byte
in the SET_TENSOR message: the client sets it when SET_TENSOR_HASH
reported a miss, the server saves a cache entry only when it is set.
Bump RPC_PROTO_MAJOR_VERSION since the wire format changes.
---------
Co-authored-by: Patrick Hoffmann <patrickhoffmann@MacBook-Pro-14-HOP.local>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
* cuda: support row-contiguous SUM_ROWS
* organize the code and add GGML_OP_MEAN to support row-contiguous tensors using the same shared kernel, and add a test to MEAN permute/slice
* Keep original comments and add if/else branch
* exclude GPU/NPU failing POOL_2D case
* Fix pool case
* ggml-openvino: fix stateful decode for Gemma-4 per-layer-type head sizes
* ggml-openvino: fix MSVC narrowing error in permute
* ggml-openvino: classify sliding-window layers structurally on interleaved-SWA models
* ggml-openvino: add GGML_OPENVINO_REQUANT_KQUANT to select a 4-bit requant target
* ggml-openvino: add GGML_OPENVINO_SPILL_DIR to spill weight buffers to disk
* Stateful Performance: Added pass::KVStateSeqAxis to change KV layout
* ggml-openvino: fix stateful decode past the sliding-window size
Assisted-by: Claude Sonnet
* ggml-openvino: refuse stateful decode that cannot resume from the KV state
The stateful path seeds its KV state from ggml's cache when the decode position
is ahead of what the state holds. That only works when ggml's cache is a plain
prefix, where cell i holds position i. A sliding-window layer keeps just the last
n_swa positions and drops the rest, so past the window cell i no longer holds
position i and the seeded state is wrong.
Slicing the state to the decode position also had no bounds check, so a position
past the end surfaced as a bare ov::Exception from the ROI constructor
(llama_decode ret = -3, with no reason given at default verbosity).
Refuse both cases with a clear message instead, and refuse on the compile path
too, where a new model starts with an empty state and so can only serve a
sequence from its beginning. Reproducible with llama-bench -d, which restores a
saved sequence state rather than recomputing the depth prefill.
Assisted-by: Claude Opus 5
* ggml-openvino: use the per-layer KV head count for the stateful KV state
The stateful path reinterprets ggml's KV buffer [1, 1, seq, n_heads_kv * head_size]
as [1, seq, n_heads_kv, head_size]. The head size is already taken from the
tensor's own combined dim, because gemma-4 varies it per layer type, but the head
count still came from a model-level scalar that compute_llm_params() overwrites
per attention node, so it ended up holding whatever the last layer said.
gemma-4 varies the head count per layer too: 12B has 8 x 256 sliding layers and
1 x 512 full layers, 31B has 16 x 256 and 4 x 512. So 40 of 12B's 48 layers were
split as 1 x 2048 instead of 8 x 256, and attention read the state with the wrong
head split - both models decoded garbage on CPU and GPU. E2B is unaffected, its
head count is 1 everywhere.
Record the count per layer instead and look it up by the cache_k_l<N> leaf name.
Key it by layer, not by layer type: the sliding/full classification comes from
cache extents, which tie at a small -c, while the head count does not.
The stateful state trim now derives its sequence axis per state for the same
reason, since pass::KVStateSeqAxis matches per state on the head count.
Assisted-by: Claude Opus 5
* ggml-openvino: apply the KV state relayout to any KV head count
pass::KVStateSeqAxis was limited to states with a single KV head, where moving
the sequence axis from dim 1 to dim 2 is a pure metadata change. The limit was
also based on a measurement showing no gain for a multi-head model, but that was
taken at depth 0, which is the one depth where this change does nothing.
With several heads the pass does more than move metadata: it drops the reader
side transpose of the whole accumulated state, which the graph otherwise redoes
every token at a cost that grows with the context length, and replaces it with a
transpose of the single new row. Measured on GPU, tg128, alternating arms:
gemma-4-12B 6.27 -> 9.11 t/s at depth 8192 (stateless is 7.69, so stateful now
wins at depth instead of losing), Llama-3.2-1B 47.8 -> 59.6 t/s. Both are within
noise at depth 0, which is why the earlier check saw nothing.
The state refill needs the rows copied rather than reinterpreted now: ggml stores
[seq][n_heads_kv * head_size], and a relayout state with several heads is a
different element order. Without that, a refill would seed wrong data - it is
reachable today through llama-bench -d.
Assisted-by: Claude Opus 5
* ggml-openvino : support ggml_rope_set_offset and simplify op support gating
* add more cpy cases
* reject BF16 cpy on NPU
* Remove mul_mat_id fallback, gate large mul_mat_id only for mxfp4
* ggml-openvino: fuse the MoE expert block into MOECompressed on GPU
* ggml-openvino: skip GPU MUL_MAT_ID for unbound expert tensors
* ggml-openvino: requantize grouped 8-bit MoE experts on GPU
* Enable special strided CPY for conv state writeback
* openvino: support cacheless encoder models on NPU
Packed QKV views used by mmBERT were rejected by the ROPE support check. This split Q/K RoPE onto CPU, prevented cacheless attention detection, and sent fragmented encoder graphs through the decoder-oriented NPUW path.
Accept packed QKV RoPE views, detect cacheless attention from its mask, and run these models as a single full-sequence prefill without NPUW or a decode graph. Also provide static mask, output index, and mean-pooling shapes and inputs.
* openvino: optimize norm and RoPE translation
Replace the decomposed mean/variance normalization graph with an opset6 MVN operation. This preserves the GGML epsilon placement while allowing OpenVINO plugins to compile normalization as one operation with fewer intermediate tensors.
Cache RoPE sine and cosine outputs in the graph-wide tensor map. Build the cache key from all RoPE parameters and the optional frequency-factor input so compatible Q/K and layer nodes share one subgraph without mixing different RoPE configurations.
Expose NodeContext::put_shared() to publish translator-created outputs for graph-level reuse.
* ggml-openvino : simplify op translators and enable IMROPE/NEOX RoPE fusion
* remove unnecessary include and clean up PAD
* fix mulmat bug
* use ov::as_type_ptr instead of std::dynamic_pointer_cast
* ggml-openvino: fix mixed-dtype ADD/SWIGLU_CLAMP, gate unsupported ROPE/SOFTPLUS cases
- translate_add: upcast mismatched operand types (e.g. f16/f32 in fused
ADD_ADD) to f32, add, then cast once to the output type. opset1::Add
requires matching input types and downcasting first lost precision.
- translate_glu_swiglu_clamp: same fix, f16 Swish/Clamp rounding was
drifting past the test tolerance.
- supports_op: reject ROPE with ne[3] > 1 (multi-sequence) since the
cos/sin tables only cover one sequence, and SOFTPLUS on GPU since the
OpenVINO GPU kernel overflows to inf for large inputs (CPU is fine).
- ci/run.sh: serialize test-backend-ops on OpenVINO GPU; running two
workers concurrently crashes the GPU plugin (CL_OUT_OF_RESOURCES).
* openvino: share compiled models with per-context inference state; fix thread-safety
* ggml-openvino: gate MoE expert-sum ReduceSum shortcut past 8 experts
The ReduceSum shortcut for the MoE expert-plane-sum ADD chain drifts past
the 1e-7 test tolerance for >8 experts (f32 accumulation order vs CPU
reference), intermittently, like the existing Q4_K/Q5_K NMSE case.
Expose is_moe_expert_sum_add() so supports_op can gate on expert count
and fall back to CPU for just that reduction op.
* ggml-openvino: gate degenerate m=1,n=1 MUL_MAT on GPU
CI hit ERR=1.8e-3 (> 5e-4 tolerance) for a scalar-output f32 dot product
(m=1,n=1,k=2048); didn't reproduce locally in 8 tries, so likely an
internal fp16 accumulation path the GPU plugin picks for this tiny
shape. m=1 output dim doesn't occur in real model weights, so gate it.
* ggml-openvino: make SoftPlus decomposition opt-in native
Assisted-by: Codex
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
Co-authored-by: ravi9 <ravi.panchumarthy@intel.com>
* metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3)
MiniCPM3 sets attention.key_length to 96 and does not set
attention.value_length, which defaults to n_embd / n_head = 64. Metal had no
(96, 64) instantiation, so -fa auto aborted on the missing
kernel_flash_attn_ext_vec_f16_dk96_dv64.
Instantiate the tile kernel at (96, 64) for every K/V type that already has
(96, 96), and the vec kernel for the NE=4 configurations. Of the NE values the
vec dispatch considers, only NE=4 works here, because NL = 32/NE has to divide
both DK/4 = 24 and DV/4 = 16.
* tests : avoid redundant FA vec slice coverage
This commit updates cmake to use PROJECT_SOURCE_DIR instead of CMAKE_SOURCE_DIR for paths in function calls.
The motivation for this is that when using add_subdirectory,
CMAKE_SOURCE_DIR is fixed to the top-level projects source directory,
that is the caller of add_subdirectory and not the llama.cpp root
which means that common/common.h header will not be resolved.
Refs: https://github.com/ggml-org/llama.cpp/pull/28091#issuecomment-5636106377
When /tools returns 403 (server started without tools), the web UI
refetched the tool list before every chat message, since the guard
treated an empty tool list as "not yet fetched". Each retry returned
403 and could trip fail2ban.
Skip the refetch once the store flags the endpoint as disabled, and
detect that state via the response status code instead of string-
matching the error message. The tools panel keeps probing on open so
the UI recovers once the server is restarted with tools enabled.
Fixes#28299
* release : add ubuntu-cuda build job (12.8/13.3, x64+arm64)
* Add GCC 14 for CUDA arm64 builds in CI
* Eplicit bash
* Install git for CCCL fetch
* Install git before we clone/checkout
* Match CI names for WIndows
* Whitelist llama.cpp repo to git
* Use $GITHUB_WORKSPACE
* Also ship dependent libs on Ubuntu
Need NCCL additionally as it's pre-built available on Linux
* Avoid duplicate files in packaged cudart
* Copy NCCL license
* Install CURL to fetch NCCL license
* Update .github/workflows/release.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
* Remove NCCL until licensing has been confirmed
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
This commit removes the precompiled headers that I added in Commit
3bcfeb700 ("cmake : add PCH and unity build to improve build times
(#28091)").
The motivation for this is that this looked good when developing this
but has caused multiple issues that I had taken into consideration and
we have decided to remove it and only keep the unity builds from the
above commit.
Refs: https://github.com/ggml-org/llama.cpp/pull/28882#issuecomment-5662272126
Aleksander Grygier