* vulkan: optimize IQ4_XS matmul kernels
Assisted-by: OpenAI Codex
* vulkan: address IQ4_XS review nits
- drop the dead LOAD_VEC_A != 8 branch in the IQ4_XS shmem load; iq4_xs is
in lut_load_vec_a()'s "8" list, so that path is never generated
- disable MMVQ for IQ4_XS on Intel (27.3% tg regression on A770)
- remove a stray empty line in types.glsl
Assisted-By: Claude Opus 5 <[email protected]>
Since #28732 our internal symbols are exported. A duplicate copy dlopened and
dlclosed by ggml_backend_load_all() then interposes them, so its destructors
destroy the live vk_instance and later device queries hit the GGML_ASSERT on
vk_instance.device_indices. Hidden visibility exports only GGML_BACKEND_API,
as before #28732.
Fixes#29138
Assisted-by: henk:claude-fable-5
Avoid returning references through lambdas that hold a local cast pointer, which triggers -Werror=dangling-reference in some CI compilers. Reuse the precomputed select_expr pointer directly.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* server: route every model load through the queue
A model loaded by the fast path has no queue entry, so tick() evicts
it at its LOADED transition before its own request is proxied. Every
load now joins the queue, whose entry protects the model until its
waiters leave.
* server: do not admit requests into a stopping model
A request for a model that is being stopped still sees it LOADED and
is proxied into the dying child. Such a request now joins the queue
and is served by the next instance. The stopping mark is cleared
under the same lock that sets UNLOADED, so no request can see a
model that is neither stopping nor unloaded while its child is gone.
* vulkan : Intel FA kernel optimization for split k path
* vulkan : Host code update for Intel split k FA kernel path selection, fix A770 Linux op test failures
* vulkan : use symmetric coopMatMulAdd() in flash_attn_decode_phase_1 shader to resolve test op failre on A770 Linux with 26.2.3 mesa driver
* vulkan : fix editorconfig issue in flash_attn_decode_phase_2.comp
---------
Co-authored-by: Liu, Russell <[email protected]>
* metal : gate mul_mm_id src1 rescale behind ggml_prec
Assisted-by: Claude Fable 5.1
* ggml-webgpu: reject MUL_MAT_ID when src1 precision is F32
* cuda/vulkan: reject MUL_MAT_ID in supports_op when src1 prec is F32
fix `supports_op` to return false for failing backends when the specified src1 precision is f32
Assisted-by: Claude Fable 5.1
---------
Co-authored-by: yomaytk <[email protected]>
* model : add DFlash layer-input taps for HunyuanVL
DFlash speculative decoding needs the target graph to expose the residual
stream entering each layer (res->t_layer_inp[il]) - the draft model reads
those tensors to build its cross-context. Qwen3 and the other DFlash-capable
targets register them, but the Hunyuan graphs do not, so serving a DFlash
draft against a HunyuanOCR target aborts during the first graph build:
GGML_ASSERT(t_layer_inp[il] != nullptr && "layer input tensor is null")
Register the tensor at the top of the layer loop, mirroring qwen3. The
layer input is the residual stream entering layer il, i.e. the output of
layer il-1, which is what the draft's target_layers metadata refers to
(the converter writes target_layer_ids+1). hunyuan-dense.cpp reuses this
graph, so it is covered as well; hunyuan-moe has a separate graph and is
untouched.
The vector is only read when a speculative implementation enables those
layer ids, so there is no behaviour change without a draft model.
Tested with tencent/HunyuanOCR 1.5 and its DFlash draft: image requests now
run, draft acceptance is ~0.5 and the OCR output is byte-identical to the
non-speculative run.
Co-authored-by: wendadawen <[email protected]>
* convert : fix DFlash draft conversion against HunYuan targets
Converting a DFlash draft with a HunYuan target failed in two ways.
1. DFlashModel.set_vocab() reuses the target class' vocab handling by
calling it unbound with the draft instance, but HunYuanModel.set_vocab()
called self._fix_special_tokens(), a method that only exists on
HunYuanModel, so the conversion always aborted with
AttributeError: 'DFlashModel' object has no attribute '_fix_special_tokens'
Make the vocab helpers module-level functions taking the model
explicitly, so they do not depend on the instance being a HunYuanModel.
They have no other callers, so the two id lookups are folded into
_fix_special_tokens().
2. The delegated call runs with self.dir_model pointed at the target but
keeps the draft's self.hparams, so config lookups inside the target's
vocab code (the pad_token_id < 0 guard, eod_token_id) read the draft's
config instead of the target's. That aborts on targets with
pad_token_id = -1 (e.g. the HunyuanOCR v1.0 checkpoint) and otherwise
writes special token ids that disagree with the target.
Add _vocab_hparams(): it returns the target's config (with text_config
merged to the root, as TextModel does) when the model is a draft
converted with --target-model-dir, and the model's own hparams
otherwise, so a normal conversion is unaffected.
Tested: converting tencent/HunyuanOCR/dflash succeeds with both the 1.5 and
the v1.0 target; converting the base model without --target-model-dir
produces a byte-identical GGUF to before.
Co-authored-by: wendadawen <[email protected]>
* convert : fix DFlash draft vocab against HunYuan targets
Switch hparams to the target config for the duration of the borrowed
set_vocab(), matching the existing dir_model swap, instead of teaching
HunYuanModel::set_vocab about draft models.
* convert : fix HunYuan special token ids for DFlash drafts
* convert : use load_hparams for HunYuan special token ids
* convert: add MiMo-V2.6 support
Hoist the K3 mxfp4 conversion repack into base.py so it can be reused
Remove decoder from mmproj convert
* Update conversion/mimo.py
* fix: use autoparser
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
Co-authored-by: Piotr Wilkin <[email protected]>
The snapdragon CI builds packages only to feed the QDC device tests, so
Hexagon NPU binaries never reached the releases page. Build both targets
in release.yml and attach them as release assets.
This commit tweaks the Toaster element to include a close button.
These toasts often cover other UI elements like the model selector, and
this change avoids having to wait for them to disappear on their own
(e.g. after a load failure).
* hex-gdn: start putting together HMX support for GDN
* hex-gdn: working hmx but not-pipelined and slow for now
* hex-gdn: re-write vtcm layout handling and prep for pipelining
* hex-gdn: starting to pipeline hmx and dmas
* hex-gdn: add hvx threading for most pipeline stages
* hex-gdb: add detailed trace events
* hex-gdn: vectorize expfs and use aligned hvx reads/writes
* hex-gnd: vectorize the rest of expf
* hex-gdn: optimize tail processing (pad partial chunks)
* hex-gdb: avoid float up/down casts in hot loops
* hex-fa: remove float up/down casts from inner loops
* hex-gdn: do exp() in f16 to improve HVX utilization
* hex-gdn: optimize tiler
* hex-hmx: bump hmx-queue to 128 and dispatch all GDN gemms at once
* hex-gdn: further pipeline improvements
* hex-gdn: optimize gdn prep stage
* hex-gdn: yet more tweaks to optimize GND_SOLVE task and pipeline
* hex-gdn: improve accuracy and optmize gdn-prep further
* hex-gdn: fix rebase conflict
* hex-bufs: revert max_bufsize enforcement, it is enough to just enforce max_vmem
* hex-scripts: improved inspect script to avoid false alarms in reg spill detector
* hex-fa: improve inline softmax with in-reg VKQ32 accum
* hex-fa: minor improvement for dma pipeline in hvx kernel
* hex-fa: reduce ddr reads by 20-30% during token gen
* hex-gdn: proper alignment for hvx vtcm spads
* llama-context : report graph inputs and input tensors during sched reserve
- fix the tg (token generation) graph bs label to use n_seqs instead of a hardcoded 1
- report the number of graph inputs from llm_graph_result::inputs for both the pp and tg graphs
- report the number of input tensors (nodes and their src tensors flagged with GGML_TENSOR_FLAG_INPUT)
- log a warning when an input tensor has an op other than GGML_OP_NONE
- log a trace line for each input tensor and the nodes (name and op) that use it
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* cont : count input tensors before reserving the sched
* wip
* llama-graph : name the unnamed graph input tensors
- name the kv-cache idxs input tensors (attn_inp_k_idxs, attn_inp_v_idxs)
- name the recurrent state copy idxs input tensor (rs_s_copy)
- report the input tensor shape in the sched_reserve trace
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : rename "graph inputs" to "graph input objects"
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* llama-context : report the sched reserve graph stats on a single line
- print nodes, splits, input objects and input tensors in one line
- when the pp and tg graphs differ, print each value as 'pp / tg'
and annotate the line with the batch sizes used for each graph
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : pad logs
The 5-argument load_ldmatrix added in 1884824fd only defines tile<16,8>, so the Volta tile<8,4> does not match. See https://github.com/ggml-org/llama.cpp/issues/29222 for details. Building on 1884824fd, generalize the tile shape of the 5-argument load_ldmatrix from <16,8> to <I,J>, so the non-swizzle branch forwards to the 3-argument loader for any shape. Local compilation and testing passed.
Assisted-by: DeepSeek V4.1 Flash (OpenCode)
convert_unary handles the contiguous case through the general strided kernel,
one element per thread: each lane reads 4 bytes and writes 2. Converting the
activations for a bf16 matrix multiplication that way moves 126 MB in 1021 us
on gfx1151, about 65% of what the memory system can do.
Give the contiguous path its own kernel that takes four elements per thread
through a vector type, so a warp loads 512 bytes at a time instead of 128. It
is used only when the element count is a multiple of four and both pointers
carry the alignment the vector type needs, and falls back to the strided
kernel otherwise.
Model level, Qwen3.8-Next-Flash IQ3_XXS on gfx1151, llama-bench -ub 2048 -r 6,
mean of the last 3 reps, ABBA counterbalanced:
pp2048 688.0 680.0 -> 694.3 691.1 +1.26%
tg128 24.8 24.8 -> 24.8 24.8 +0.14%
Every conversion in a prefill takes the new kernel (kernel trace: 1146
convert_unary_cont_vec4, no convert_unary). Output is bit identical; MUL_MAT,
MUL_MAT_ID, CPY, CONT, GET_ROWS and SET_ROWS pass.
Assisted-by: Claude Opus 5
* tests/test-backend-ops : allow regex entries in the -o filter
so far -o only accepted a comma separated list of exact op names or
full test case strings. entries that are not plain op names are now
treated as regexes matched against the op name (e.g. "MUL_MAT.*"),
while plain names keep their exact-matching behavior so that
"-o ADD" does not match ADD_EX etc.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : don't print the FA vec slice log when not needed
* tests/test-backend-ops : reformat the help text
use the same style as the other tools, with separate sections for
modes, options, and examples
Assisted-by: pi:llama.cpp/Qwen3.8-27B
Allow configuring --temp, --top-p, --min-p, --repeat-penalty,
--presence-penalty and --frequency-penalty via LLAMA_ARG_* so
llama-server can be fully controlled from an EnvironmentFile
(e.g. systemd on Debian).
Use `llama-gen-docs` to regenerate the readme files.
In router mode, authentication belongs to the router. unset_reserved_args()
already unset LLAMA_API_KEY, but did not unset LLAMA_ARG_API_KEY_FILE.
When --api-key-file was passed, children re-validated against file keys only,
causing clients using --api-key to 401 on chat completions (#28820).
In addition, router internal calls without auth headers (such as
POST /v1/streams/lookup and DELETE /v1/stream) were silently rejected with 401.
Unset LLAMA_ARG_API_KEY_FILE in unset_reserved_args() so no API keys reach
child instances. This keeps keys out of child argv, ensures all keys the router
accepts work end-to-end, and prevents router internal stream calls from 401ing.
Fixes#28820
* Fixed json enum handling
Added common_json_value handling for enum values.
Added tests/test-json.cpp to cover testing of some aspects of common_json.
* Removed tests as requested.
* Applied recommended style and simplification
Simplified by delegating enum constructor to the constructor of the underlying type
Matched style of surrounding templating code
* Implemented ARM NEON DP q1 4x4 repack
* Hoisted out scaling by b_d in gemm
* Added 4x8 NEON I8MM repack kernels
* Cleanup for q1 arm repack
* Added missing aliases for arch fallback
* Corrected unused var statements
* Extended table guard condition to account for i8mm w/o dp build
Co-authored-by: Copilot Autofix powered by AI <[email protected]>
* Moved new declarations and references to groups' top
* Moved declarations for uniformity
---------
Co-authored-by: Copilot Autofix powered by AI <[email protected]>
* hex-dma64: enable support extended buffer mappings and 64bit dma
hex-dma64: expand binary ops to support more DMA scenarios
hex-dma64: add binary-ops.h
hex-dma64: add --hex-dma64 to run.py and fix minor issues
hex-dma64: update SSM_CONV to use dma with proper support for 64bit
hex-ops: remove obsolete gate for % 128 in binary ops
hex-l2: dont check weight tensors against dirty ranges
hex-dma64: most binary ops now support dma
hex-dma: use dma_addr_t instead of plain uint64_t to avoid overhead on older targets
hex-dma: update all dma users to use dma_data (instead of pointers)
hex-dma64: simplify lazy buffer mapping and clonning
hex-fusion: factor out try_fuse_common that checks for dma64 buffers
hex-bufs: minor cleanup for mmaping logic
hex-bufs: simplify buffer clonning
hex-ssm-conv: tighten gating checks and check vtcm size in kparams
hex-binary: fix incorred mod/wrap in scalar ops
hex-binary: make sure to call precompute kparams in support checks
hex-dma64: update addr handling in mm,concat,binary
hex-dma64: fixing up leftover of dma_addr_t conversion
hex-binary: redo the kernel selection again and fix regressions in MOEs
hex-binary: specialize per-type/per-op
hex-binary: vtcm-layout and per-src dma-queue
hex-dma64: update dma_push to transparently handle 64bit/extended
* hex-cpy: fix improper rebase with the fixes for cont. tensors
* hex-dma-cpy: update CPY to use safe dma rows/size limits
* hex-mmap: bump number of mmaps to 64 to allow avoid eviction in larger models
* hex-dma: add support for the secondary ring as a fallback for too-large transactions
* hex-rope: fix freq_factors access with 64bit dma
* hex-dma: audit all ops for proper use/gards for 64bit addresses
* hex-dma64: uninline glu-compute funcs to avoid register pressure due to 64bit addr math
* hex-dma64: refactor binary ops to separate dma loops
* hex-devel: add inspect script to help with dbg and analysis
* hex-dma: refactor dma-pipelines in unary-ops
* hex-dma: rewrite softmax to use dma
* hex-dma: rewrite GDN dma loops and improve HVX register usage
* hex-gdn: fuse GDN+CPY
* hex-mm: factor out HVX solver
* hex-mm: remove hvx-flat kernels, the chunked version now handles vtcm limits much better
* hex-buffs: reject huge buffer allocations that we cannot memory map
* hex-inspect: add logic to look for float promo calls
* hex-mm: reduce HVX register spills in HVX prompt kernels
* hex-bufs: do not double count buffers from tensors in the same op
* hex-roll: fix merge conflict
* hex-dma: reroute all matmul ddr kernels to new chunked dma/vtcm kernels
* hex-dev: update developer docs to include inspection for register spils and float promos
* hex-ops: forgot to add new headers
* hex-softmax: fix gpt-oss dims
* hex-dma64: cleanup dma_addr_t casts
* hex-dma64: add support for dma/vtcm for flash-atten with sinks
* hex-mm-add: fix MUL_MAT+ADD fusion with bias.weights in extended bufs
* hex-add-id: add support for dma for src1 (exp. table)
* hex-dma: imrpove v73 fallback paths
* hex-bufs: do not drop extended mappings during va defrag
* hex-scripts: fix flake8 warnings
* hex-docs: fix editor-config warnings
* hex-inspect: fix warnings from ty
the dsv4_hc_pre kernels hardcoded hc = 4 via a constexpr used with
simd_shuffle, so the op was rejected by supports_op for any other hc
and fell back to CPU. Kimi-K3 uses dsv4_hc_pre with hc equal to the
number of banked checkpoints in the cross-layer residual stack, which
grows with the layer index.
pass n_hc as a function constant (FC_DSV4_HC) with per-n_hc pipeline
variants, and loop over it in both pre kernels with direct loads
add test-backend-ops cases for hc = 1, 2, 3, 5, 8 and 65, gated and
not gated
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* ui : let the chat column shrink below its content width
The chat column is a flex item, so its automatic minimum size kept it as wide as the widest row inside it. Message rows cap at max-w-3xl plus padding, so a narrower window pushed a page-level horizontal scrollbar.
Set min-w-0 on the column so the inner scroll containers take over.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : wrap markdown tables in a scroll container
Markdown tables render as a bare <table>, which keeps its content-driven minimum width and can stretch the chat column past the window. The table-wrapper CSS already existed, but nothing produced the wrapper.
Add a rehype plugin that wraps each table in div.table-wrapper, following the existing enhance-* plugins.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : scroll long inline content inside markdown blocks
Long unbreakable content (inline code, paths, hashes) widened the message row and spilled over the neighbour elements. Give each markdown block a horizontal scroll container, and the content root one as well, since the trailing block renders with display: contents and has no box of its own.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : use exact transition properties for markdown images
transition: all repainted every property and 300ms felt sluggish. Name transform and box-shadow at 200ms ease-out, and gate the hover scale behind (hover: hover) and (pointer: fine) so touch taps do not trigger it.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : fit wide image attachments to the message width
Attachment thumbnails used a fixed height with w-auto, so a wide image kept its aspect-driven width and, being flex-shrink-0 in a right-aligned bubble, overflowed to the left of the message row.
Cap the thumbnail with max-height and max-width instead of a fixed height so it scales down proportionally, and let it shrink outside the single-row carousel.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : keep long tool call titles inside the message row
A tool title could not shrink below its content, so a long path escaped the message row. Let the title span shrink and scroll, and for the file tools put the value on its own line only when it does not fit, with the value as the only scroll container.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui : render get info as a collapsible block with a table
get_info rendered its own always-open row with the values trailing the label. Use the shared ToolCallBlock chrome so it collapses like the other tools, and list os and cwd as table rows with the key as a row header.
The error and pending states now show inside the body, including the plain-string errors the server tools path produces.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* test : pin the server mode in the add menu a11y story
The story asserts the add menu's first enabled item is the reasoning submenu, which is mounted only outside router mode. The vitest dev server proxies /props to whichever server is running, so the assertion depended on the machine's server mode and failed whenever a router was up.
Pin the mode in the story, including props.role so a re-detection cannot flip it back.
Assisted-by: pi:deepseek-ai/DeepSeek-V4.1-Flash
* ui: wrap long markdown tokens instead of scrolling every block
Making each markdown block and the content root a horizontal scroll
container turns any hover transform into a scrollbar: the blockquote
translate and the image zoom overflow their block and flash a scrollbar
under it. Each block also becomes a block formatting context, so the
paragraph margins stop collapsing across blocks and the spacing doubles.
Drop both overflow-x rules and let long unbreakable tokens wrap with
overflow-wrap: break-word on the content root. break-word leaves the
min-content width untouched, so wide tables and code blocks keep
scrolling inside their own containers.
---------
Co-authored-by: Pascal <[email protected]>
* metal: add F16 input to the FWHT
The Metal FWHT kernel accepts F32 input only. This change makes the source
type a template parameter, so the kernel reads an F16 source directly instead
of requiring a converted copy. The F32 instantiations are unchanged.
The pipeline name now carries the source type, and supports_op accepts an F16
src1 for the Hadamard hint at the four sizes the kernels cover. Every other
F16 src1 path still goes through ggml_metal_supports_mul_mat_op.
These are the test cases mentioned in #27779.
test-backend-ops on M5 Pro: MUL_MAT_HADAMARD 16/16, MUL_MAT 1265/1265.
* metal: ask the same FWHT question in supports_op and the dispatch
supports_op admitted an F16 src1 on the type, the hint and the width alone, but the
dispatch also requires src1 and dst to be contiguous and the same shape. A Hadamard
hinted MUL_MAT that passed the first and failed the second reached the generic path,
which has no F32 src0 by F16 src1 kernel, and aborted on a nil pipeline:
kernel not found in any metal library: base = 'kernel_mul_mv_f32_f16_4'
ggml_metal_encoder_set_pipeline: nil Metal pipeline
ggml_metal_use_fwht now holds the whole condition and both callers use it, so they
cannot drift apart again. The added test case has src1 and dst of different shapes,
which aborted before this change and is declined by the Metal backend after it.
* metal: branchless butterfly select in the FWHT simdgroup kernel
Review suggestion. Replaces the ternary in the shuffle stages with
val2 - val + 2*((lane & i) == 0)*val, which is the same value without the
select.
Measured on M5 Pro, interleaved A/B, five rounds, first discarded, on a
Hadamard matmul with block 512 and 65536 rows so the kernel rather than the
launch dominates: 1324.6 us before, 1285.0 us after, a 3.0% gain, and faster
in every round. At the shapes already in the perf suite the op runs 1.6 to
3.9 us against a 1.6 us launch floor, so the difference is not visible there.
FOR_UNROLL on the same loops was also measured and made no difference, the
delta changing sign between rounds, so it is not included.
* metal: move the FWHT dispatch predicates to ggml-metal-common
Review feedback. ggml_metal_use_fwht and ggml_metal_fwht_supported_size were
static inline in ggml-metal-device.h. They now follow the
ggml_metal_op_mul_mat_use_mm pattern: declared in ggml-metal-common.h and
implemented in ggml-metal-common.cpp, which is already the home for helpers
shared between supports_op and the op dispatch. The predicate is named
ggml_metal_op_mul_mat_use_fwht to sit alongside the _use_mm pair it parallels.
This also fixes the macos-latest-arm64 build. The header needed ggml-impl.h
for ggml_get_op_params_i32, but ggml-metal-device.h is reached from
tools/tuning through ggml-metal-tuning.h, and that target does not have
ggml/src on its include path. ggml-metal-common.cpp already includes
ggml-impl.h, so the accessor is used normally there and the header goes back
to needing nothing extra.
* metal: keep the FWHT size check internal and group the dispatch helpers
Applies the patch from the review. ggml_metal_fwht_supported_size becomes
static in ggml-metal-common.cpp since nothing outside it needs the size list,
which also drops stdint.h from the header again, and
ggml_metal_op_mul_mat_use_fwht joins the existing _use_mm declarations under
their shared comment instead of carrying its own block.
* tests: drop the mismatched-shape Hadamard case
I added a case with m != k to cover an abort, but the hint is a promise that
src0 is a Hadamard matrix, so src0 is square and dst has the same shape as
src1. Every other case in the suite holds to that. The case was not a valid
op, and on CPU it compared the FWHT against a real matmul of a non-square
src0, which cannot agree.
The supports_op and dispatch conditions still come from one predicate, which
is what keeps them from disagreeing on contiguity.
* chat: add dedicated Ling 3.0 (Bailing V3) parser
Ling 3.0 Flash templates pre-open the think block in the generation
prompt, so the model never emits an opening <think>, and a tool call can
arrive before any </think>. The generated autoparser terminated reasoning
only at the close tag, which classified such tool calls entirely as
reasoning_content: clients received content="" with no tool_calls and
agent loops died as reasoning-only turns.
Adds a specialized parser that terminates reasoning at the think close
tag or at a <tool_call> start, mirroring the hand-written Qwen3-Coder and
Kimi K3 parsers and the reference vLLM/SGLang Ling3 parser (which treats
<tool_call> as an implicit reasoning terminator). Detection is gated on
the <role>...</role> section markers, unique to this family among the
tagged-argument templates.
Adds the Ling 3.0 Flash chat template and tests covering the
unclosed-think tool call (full parse and streaming), healthy closed-think
paths, trailing prose, parallel calls, marker-like strings in argument
values, string-union and non-string argument types, and
reasoning_format=none.
Assisted-by: Kimi Code
* tests : move Ling 3.0 test
---------
Co-authored-by: aetherbird <[email protected]>
Co-authored-by: Alde Rojas <[email protected]>
* server-models : show source per model in log
- Show [source] tag (preset/models_dir/cache) per model instead of cryptic * marker
- Show HF hub cache path in the 'Loaded cached model presets' log
- Add hf_cache::get_cache_dir() public accessor
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : pad log
* metal : add top-k MoE fusion
Adds a Metal fusion for SOFT_MAX + ARGSORT + GET_ROWS with optional
routing-weight normalization and scale, matching the top-k MoE fusion
available in the CUDA and Vulkan backends. The fused kernel writes the
selected expert ids and routing weights directly, eliding the separate
softmax, argsort, get-rows, sum-rows, clamp, div and scale kernels.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add MoE weighted reduction fusion
Fuses MUL(experts, weights) plus the expert VIEW/ADD chain into one kernel
that computes the weighted sum directly. The graph_optimize hook keeps the
expert and weight buffers alive until the fused output so the allocator cannot
reuse them while the kernel is still reading them.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* tests : expose MoE weighted reduction in fusion baseline
Use 2 experts per token in the generated MoE test models so the Metal
MoE weighted reduction fusion (MUL + ADD) is exercised by test-fusion.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : fuse RMS_NORM + SCALE
Adds NORM/RMS_NORM + SCALE fusion to the Metal backend by reusing the
norm+mul kernel with a scalar scale flag. Adds test coverage for both
NORM+SCALE and RMS_NORM+SCALE and regenerates the fusion baseline.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constant for RMS_NORM + SCALE
Replaces the runtime use_scale karg with a Metal function constant. The
norm+mul kernel is compiled with FC_norm_use_scale=false for MUL fusion and
FC_norm_use_scale=true for SCALE fusion, so the fused kernel has no runtime
branch.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constant for top-k MoE with_norm
Replaces the runtime with_norm karg with a Metal function constant. The
top-k MoE kernel is compiled separately for the normalized and non-normalized
routing variants, removing the runtime branch.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename moe_weighted_reduction suffix to moe_reduce
Shortens the MoE weighted-reduction fusion identifiers, kernel, pipeline,
matcher, args struct, and test op name from moe_weighted_reduction to
moe_reduce.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion
Adds dense mat-vec activation fusion for sigmoid/silu and bias+softplus.
The mat-vec kernels apply the activation/bias epilogue via function
constants, avoiding the separate unary/add passes.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : revert MUL_MAT + UNARY and MUL_MAT + ADD + UNARY fusion
The mat-vec activation fusion regressed decode throughput on Qwen3.6-35B-A3B
by ~8% (tg32 81.5 vs 88.5 t/s). The regression is caused by loss of
concurrency: the standalone unary kernels previously overlapped with other
mat-vec work, while fusing the activation into the mat-vec kernel serializes
it on the critical path.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add SSM_CONV + UNARY (silu) fusion
The SSM_CONV kernels apply silu directly via a function constant, eliding
the separate unary pass. Regenerates the fusion baseline.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : address fusion review comments
- Fix declaration/table alignment
- Rename top-k MoE kargs fields to val_clamp / val_scale
- Move moe-reduce alloc-deps handling into a general fusion helper
- Remove the public moe-reduce matcher API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : fix unused parameter in top-k MoE fusion check
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : guard SSM_CONV fusion lookup behind use_fusion
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : track all fused outputs in graph reorder
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : keep top-k MoE logits alive until fused output
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : refactor alloc deps to pattern-driven approach
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : check fused kernel destination in concurrency tracking
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* meta : forward graph_optimize to underlying backends
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use vector for fusion table
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* meta : keep graph_optimize unimplemented
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* parallel : fix non-deterministic prompt selection
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* parallel : support dummy models and add global logits run hash
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : sync cross-device copies with destination completion event
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : avoid const_cast in fusion alloc deps
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : skip fusions with aliased sources
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : hide fusion pattern definition
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use vector fusion op sequences
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : drop redundant struct keywords
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : add alloc deps comment separator
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : generalize fusion output memory ranges
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename fusion out_offsets to outs
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : avoid dst vector in memory range check
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : optimize fusion matching and multi-output handling
- use pointer arithmetic for fusion info count lookup
- avoid heap allocations in top-k MoE and MoE reduce pattern matchers
- use fusion outs for multi-output subgraph checks
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* Revert "parallel : support dummy models and add global logits run hash"
This reverts commit 57c7caf941c1b43c270fd5009c9f175063522e96.
* fusion : update MTL.csv
* metal : unroll constant loops in top-k MoE kernel
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constants for top-k MoE n_expert and top_k
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : rename fusion kargs to scale and clamp
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use function constants for moe_reduce and ssm_conv
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* fusion : update MTL.csv
Add support for the new DSV4 HC op variants used by qwen4exp:
- hc_pre with per-element sigmoid gate (gated variant)
- hc_post with identity mixing (comb == nullptr)
Assisted-by: pi:llama.cpp/Qwen3.8-27B
argsort_f32_i32_cuda_cub called the one-shot DeviceRadixSort::SortPairs
API with d_keys_in == d_keys_out (temp_keys, temp_keys). CUB's internal
double-buffer ping-pong requires distinct key buffers: with aliased
buffers the sort partially overwrites its own input mid-pass and emits a
corrupted permutation, surfacing as intermittent garbage indices (e.g.
backend top_k over a 248k-column vocab on Maxwell/CUDA 12.5/CCCL 2.x,
which then triggered out-of-bounds gathers in downstream get_rows).
Use a distinct keys-out buffer for all six call sites (plain and
segmented, ascending and descending, size-query and execute).
---------
Co-authored-by: Claude Opus 4.6 <[email protected]>
Co-authored-by: Oliver Simons <[email protected]>
Allow HMX flash-attention to run with head_dim not a multiple of 64
(e.g. SigLIP head_dim=72), by operating on DK/DV rounded up to 64 with
zero-filled tail lanes.
* ggml-cpu: add F16 input to the FWHT
The CPU FWHT accepts F32 input only. This change makes the source type a
template parameter. The CPU path now accepts F16 input and F32 input.
The CPU MUL_MAT reference now converts an F16 src1 to F32. It does this when
the caller sets the Hadamard hint.
No backend has an F16 FWHT kernel yet. The test cases come with the backend
changes that add one.
* ggml-cpu: assert the F16 FWHT input path, and use the bulk converter
Address review feedback.
The F16 branch writes plain floats into wdata, which is only correct when
vec_dot_type is F32. That invariant held because supports_op only accepts an
F16 src1 for the Hadamard hint with F32 src0 and dst, but nothing enforced it.
Assert it next to the existing src1 type check so widening supports_op cannot
silently break the write.
Replace the hand-rolled conversion loop with ggml_cpu_fp16_to_fp32.
* llama: read the SWA pattern as a period or a per-layer array
Add llama_model_base::load_swa_pattern(), which reads
sliding_window_pattern either as one flag per layer or as a period
expanded by set_swa_pattern(), and use it in every loader that reads
the key as a period.
These loaders silently ignored an array and applied their default
period, although the converters of olmo2, gemma3n and exaone4 write
arrays. The published GGUFs match the defaults, so their outputs do
not change. The loaders that already accepted both forms lose their
duplicated scalar-then-array block, and use their declared default
period when the key is absent.
* model-saver: write the SWA pattern and the MLA SWA geometry
Write sliding_window_pattern as one flag per layer, nextn layers
included, for every model using SWA. The array is never collapsed to
a scalar, since the loaders read a scalar as a period.
Also write the MLA key/value lengths and KV LoRA rank of the SWA
layers, required by dots3note.
This enables the saver for plamo3, gemma3, cohere2, cohere2moe,
olmo2, exaone-moe, afmoe, mimo2, spark2_5, muse-glimmer, mellum,
laguna, granite_swa, dots3note and maple, all passing the bit-exact
roundtrip of test-llama-archs.
* vulkan: add IQ3_S MMQ matmul kernels
* Make block_a_to_shmem do 2-byte loads (110 bytes is divisible by 2)
* Align the check, IQ3_S is also using K tile size
* vulkan: raise the hoisted row-id limit for mul_mat_id to 512 experts
The expert-count shader (count_experts.comp) sizes its shared arrays
with BLOCK_SIZE, which is 256. Because of that, row-id hoisting is
switched off for any model with more than 256 experts, and every
mul_mat_id workgroup has to rescan the whole ids tensor on its own.
Qwen3.8-Flash-Next has 512 experts and was quietly running on that
slow path.
This change sizes the arrays with a separate MAX_EXPERTS constant (512),
clears them in a loop instead of one entry per thread, and raises the
matching limit on the host side.
On Strix Halo at batch 2048 the expert matmuls drop from 12.5 to 9.5 ms
(iq3_s) and from 14.0 to 7.5 ms (iq4_nl) per op, and prompt processing
gets about 19 % faster at 8k tokens. test-backend-ops MUL_MAT_ID passes
(891/891) with new 512-expert test cases.
Assisted-by: Claude Fable 5.1
* vulkan: raise the hoisted row-id limit for mul_mat_id to 1024 experts
Follow-up to review feedback: 1024 matches LLAMA_MAX_EXPERTS instead of
stopping at 512. The three shared arrays in count_experts.comp grow to
3 * 1024 * 4 = 12 KiB, which fits the 16 KiB that Vulkan guarantees for
maxComputeSharedMemorySize.
Adds mul_mat_id test cases at 1024 experts alongside the existing 512
ones. test-backend-ops MUL_MAT_ID passes on Vulkan (RADV, Strix Halo,
Radeon 8060S): 889/889.
* Update to openvino-2026.4
* Update OV docs
* ggml-openvino : fix clangd and MSVC warnings
* fix int to ptr cast, more internal linkage enforcement, and avoiding duplicate switch case
---------
Co-authored-by: Mostafa Faheem <[email protected]>
* ui: fix accidentally removed reasoning menu in single model mode on desktop
* ui: formatting task run to fix storybook test
* ui: mount the add menu reasoning submenu outside router mode only
The models selector already owns the reasoning submenu in router mode,
so the add menu only mounts it in single model mode. The first enabled
item of the add menu is now the reasoning submenu, the accessibility
story expects it.
---------
Co-authored-by: Ben Babik <[email protected]>
Co-authored-by: Pascal <[email protected]>
* ci : add API/ABI check to make-release workflow [no ci]
This commit adds an API/ABI compatibility check to the make-release
workflow.
The motivation for this to allow us to detect any potential breaking
changes in API/ABI compatibility between releases and fail the the
release if there are any.
The workflow can be triggered manually as before and this check can be
skipped if needed as it does take some time which might be useful when
doing a dry-run and not specifically interested in the API/ABI check.
By default this will check the current release against the latest
release, but this can also be configured in the workflow, or in the
script run on the command line, to check a different tag.
* add check for minor version bumps [no ci]
This commit also changes the build type to be RelWithDebInfo so that the
reported information is more useful.
* gguf : align the data section relative to the GGUF start, not the file
gguf_init_from_file_ptr reads a GGUF from the current file position, but padded
the data section from file offset 0, so a GGUF embedded at an offset that is not
a multiple of the alignment loaded without error and returned wrong tensor data.
Also adds llama_adapter_lora_init_from_file_ptr, and disables mmap with a warning
when an embedded data section is not aligned, instead of asserting in ggml.
Assisted-by: Claude Opus 5
* llama : load lora from path through the FILE* variant
The test now checks that mmap is disabled only for an unaligned offset.
Assisted-by: Claude Fable 5.1
* Update ggml/src/gguf.cpp
Co-authored-by: Johannes Gäßler <[email protected]>
* Update include/llama.h
Co-authored-by: Johannes Gäßler <[email protected]>
* llama : error on unaligned mmap of an embedded GGUF, drop test-load-file-ptr
---------
Co-authored-by: Johannes Gäßler <[email protected]>
Both im2col.comp and im2col_3d.comp declare D_ptr without an explicit
buffer_reference_align, so glslang emits writes through it as Aligned
16. The shaders advance the pointer by D_SIZE, a per-variant define
set to 4 for float and 2 for float16_t, so most write addresses are
not 16-byte aligned. This triggers
VUID-RuntimeSpirv-PhysicalStorageBuffer64-06315 under GPU-AV.
Declaring buffer_reference_align = D_SIZE matches the alignment to the
actual write stride and takes validation hits from 20 to 0 for both
IM2COL and IM2COL_3D.
Fixes#28960
* model: calculate split states for attn_qkv from n_head * n_embd_head_k
required for gemma4 with --fuse-qkv, where n_embd is 5376 but Q is 8192.
* model: handle fused full attention layers for qwen35/qwen35moe
* model: add TODO: [TAG_SPLIT_QGATE_QWEN]
llama probes weight placement with a rope where all params are 0, so rejecting
n_dims == 0 or freq_base == 0 puts rope_freqs on the CPU. That splits the decode
graph at every full-attention layer (gemma-4-E2B: 5 splits instead of 2).
Assisted-by: Claude Opus 5
* model : add support for HrmTextForCausalLM (DFM Mimir 1B)
HRM-Text runs two transformer stacks (low, high) in an alternating cycle over the same token stream. The low-cycle state z_l starts from a learned [n_embd] tensor and is broadcast over positions.
- conversion: new writer for the fused gqkv projection (order gate,q,k,v) remapped to llama.cpp q/k/v plus a separate sigmoid gate tensor
- loader: block_count = lps * h_cycles * (l_cycles + 1) cache slots aliasing 2*lps physical blocks via struct copies
- graph: looped build with sigmoid-gated attention, SwiGLU FFN and parameterless RMS norms; learned embedding_scale applied in build_inp_embd
- saver: pointer-deduplicated layer loop (looped archs alias tensors)
- tests: hrm_text fixture (lps 1, h 2, l 3) in test-llama-archs
Limitations:
causal attention only - the upstream prefix-LM mode is not implemented (the prefix_lm GGUF key round-trips unused).
The KV cache holds one entry per pass: 128 layers for Mimir 1B, i.e. 4x a same-width 32-layer model - about 3072 MiB at ctx 4096 in F16 (halves with q8_0 KV + FA).
Every token runs all 128 block passes, so decode cost is roughly 4x a dense model of equal width (2.65 t/s BF16, 8-thread desktop CPU).
Verified against the HF reference: identical argmax at 334/334 positions across 20 prompts (BF16 GGUF vs FP32 golden).
q8_0 requant: 95.8% top-1, all remaining misses inside the HF top-5 (accumulated error over 128 sequential blocks).
AI usage disclosure: YES
Used GLM-5.3 for the majority of code AI-generated under my direction, all gates verified locally.
All in all I could say that I have written less than 20% of the code and most of the heavy lifting has been done by the model. As such, this should be considered experimental.
* Update conversion/hrm_text.py
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* convert : add gguf_writer methods for hrm_text metadata
replace raw add_uint32/add_bool calls with dedicated GGUFWriter methods, following the add_embedding_scale pattern
Assisted-by: GLM-5.3
* convert : map regular hrm_text tensors via tensor_mapping
delegate unfused checkpoints to the base tensor mapping; training-style attn. names are renamed to self_attn. so the patterns match
Assisted-by: GLM-5.3
* model : format hrm-text build_* calls as in other models
one argument group per line, matching sibling model files
Assisted-by: GLM-5.3
* llama : move hrm z_l_init table entries out of the nemotron group
place the name and tensor-info entries with the other global input tensors
Assisted-by: GLM-5.3
* convert : slim down hrm_text comments
Assisted-by: GLM-5.3
* convert : build hrm_text block tensor names from the {bid} template
The tensor map holds concrete per-block names, so format the template
with the computed layer index before handing it to super().
* llama : name hrm metadata keys in their own hrm. namespace
The four keys are arch-independent, unlike the arch-substituted
Keys.LLM entries, so group them under Keys.HRM (like Keys.Split) and
rename the llm_kv entries to LLM_KV_HRM_*. Only our own GGUFs carry
the old hrm_text.* keys; they are regenerated.
* Update src/llama-model-saver.cpp
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* llama : keep hrm metadata keys arch-substituted
Per review: the GGUF keys stay "{arch}.h_cycles" style, so the Python
members drop the LLM_KV_HRM_ prefix and keep arch templates; C++ keeps
the LLM_KV_HRM_* enums. GGUF output is unchanged - existing files and
HF uploads stay valid.
* Update gguf-py/gguf/constants.py
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* Update src/llama-arch.cpp
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* convert : rename hrm writer methods to add_hrm_*
Generic names like add_h_cycles/add_prefix_lm are too broad on the
shared GGUFWriter; prefix them with hrm_ like the metadata keys.
* model : fix meta-split lookup for archs with aliased cache slots
Cache tensors of archs that alias physical blocks across looped slots
(hrm_text, nanbeige with num_loops > 1) can reference block indices
without weight tensor names. Take the output projection from the layer
array instead of asserting; all other lookups are unchanged.
* model : replicate hrm_text tensors on meta devices instead of splitting
The aliased cache slots rotate split states differently from their
physical weights, so the meta-split execution invariants (set_rows
requires the cache state to match the token indices) cannot hold for
any device count. Replicate all hrm_text tensors on every meta device
instead; single-device and non-meta paths are unchanged.
Assisted-by: Claude Sonnet
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
The `sizeof(int16_t)` branch in `permute_transpose_impl` calls
`rvv_transposed_s32_mn_to_nm` instead of `rvv_transposed_s16_mn_to_nm`.
This is a copy-paste bug from the `sizeof(int32_t)` branch above it.
The s32 function uses 32-bit segment load/stores (`vssseg8e32.v`) on 16-bit
data, reading 2x bytes per element and producing completely wrong
transposition results -- 14 out of 16 positions are corrupted for a 4x4
int16 matrix.
The correct function `rvv_transposed_s16_mn_to_nm` already exists (line 390)
and is used elsewhere in flash attention (line 1488).
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 <[email protected]>
* 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 <[email protected]>
Co-authored-by: Georgi Gerganov <[email protected]>
* 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 <[email protected]>
* 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 <[email protected]>
* 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 <[email protected]>
Co-authored-by: Claude Opus 5 <[email protected]>
* 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 <[email protected]>
Co-authored-by: Mustafa Cavus <[email protected]>
Co-authored-by: zhaixuejun1993 <[email protected]>
Co-authored-by: ravi9 <[email protected]>
* 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 <[email protected]>
* Remove NCCL until licensing has been confirmed
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
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
* tests : add README for updating the per-backend fusion baselines
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* ci : trigger fusion on changes to test-llama-archs.cpp and src/models
the dummy models and their architectures drive the fusion baselines, so a
change to either can alter the per-fusion counters and should re-run the
fusion job.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : merge the fusion build commands in the README
assisted-by: pi:llama.cpp/Qwen3.8-27B
* pi : require explicit permission before posting PR/issue comments
assisted-by: pi:llama.cpp/Qwen3.8-27B
* gguf-py: add Maple tensor constants
Add MODEL_ARCH.MAPLE, its "maple" name, and the tensor list for the
Maple 20B-A1B ternary MoE architecture: token embeddings, output,
attention with Q/K RMS norms, and per-expert FFN tensors.
* convert: add Maple HF->GGUF converter
Register MapleForCausalLM in the HF architecture map and add the
converter for the Maple 20B-A1B ternary MoE model: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, partial rotary factor 0.5, and
per-expert weight stacking into merged 3D tensors.
* llama: add Maple architecture (20B-A1B ternary MoE)
Add the Maple 20B-A1B ternary MoE architecture: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, and ternary TQ1_0/TQ2_0
quantization support.
- register LLM_ARCH_MAPLE between MAMBA2 and JAMBA
- implement llama_model_maple: Q/K RMS norms after projection (GEMMA4
style), rope applied only on SWA layers (nope_on_global_attention),
ISWA KV cache, and MoE FFN with swiglu gate clamp at +7 (DEEPSEEK4
style)
- mark MAPLE as unsupported by the model saver (roundtrip skipped)
* tests: mark Maple as MoE-mandatory
Maple is always-MoE: the model throws when n_expert == 0, so the
test harness must only run the MoE config for LLM_ARCH_MAPLE.
* maple: apply review feedback (n_ff_exp_arr, get_arr, rope params)
- load_arch_hparams: use n_ff_exp_arr + n_ff_exp() accessor (upstream
changed these from a scalar member during the rebase)
- sliding_window_pattern: get_arr, the pattern is mandatory for this arch
- partial_rotary_factor: read only from rope_parameters (base.py mirrors
the top-level key automatically)
- document why TOKEN_EMBD/OUTPUT are forced to F16 (they are the two
dense tensors in Maple, and the reference GGUFs ship them as F16)
- add @ModelBase.example("deepgrove/maple-preview")
* tests: add Maple to the SWA pattern array list
get_arr for maple.attention.sliding_window_pattern requires an array, but
the harness only emitted a per-layer array for the arches in its list, so
test-llama-archs -a maple failed to load the model.
Assisted-by: DeepSeek Harness
* maple: move swiglu_clamp_exp to the converter
The loader prefilled 7.0 and read the key optionally. The converter now
writes it and the loader reads it as required, because llama-graph.cpp
skips the clamp when the limit is 0 and an optional read would silently
run unclamped. The test harness provides the key for the same reason.
Also drops tensor_force_quant: base.py already forces FFN_GATE_INP to F32
and TOKEN_EMBD/OUTPUT to F16 for ternary file types.
Assisted-by: DeepSeek Harness
* convert: fix the LazyBase func signature in the Maple converter
ty flagged the stack() closure: it takes no argument, while LazyBase is
annotated with func: Callable[[Any], Any]. Pass the tensor list through
args instead of closing over it, the same way kimi_k3 does, so the
callable shape matches.
Assisted-by: DeepSeek Harness
* sycl: GPU-resident TOP_K for large k, parallelised over the device
The SYCL backend refused GGML_OP_TOP_K above k = 32 and let it fall back to
the CPU, a backend round-trip per call. The limit was not conservatism: the
scan-merge kernels keep (split_block + 1) * k candidate (value, index) pairs
in SLM, so at k = 128 a work-group already needs 132 KB and cannot launch.
qwen4exp's sparse-attention indexer asks for k = 2048 in 12 layers on every
token, so this fired at every context length.
Add a radix select for large k. The k-th largest is found by four
most-significant-first passes over an order-preserving unsigned key: histogram
the digit over the candidate set, walk the buckets from the top, and recurse
into the one where the running count reaches what is still needed. SLM holds
the histogram rather than candidates, so the footprint is independent of k.
A final pass emits every column beating the pivot plus exactly as many
pivot-equal columns as are still missing, so duplicate keys still yield
exactly k distinct indices. Output order is not required and is not paid for:
ggml-cpu/ops.cpp swaps its first two outputs to say so.
The key folds -0.0 onto +0.0 so its equivalence classes match the reference
comparator, under which the two tie. NaN has no defined order in the reference
(its comparator is not a strict weak order there); here +NaN keys above +inf
and -NaN below -inf, which at least makes the result deterministic.
One work-group per row leaves the device idle whenever a graph has fewer rows
than it has cores, which at batch size 1 means one work-group full stop:
qwen4exp tops-k a tensor of shape [n_kv, n_tokens/n_stream, n_stream], so
token generation gives nrows == 1, and the backend sampler reshapes logits to
a single row as well. Measured, ne=[200000,1] and ne=[200000,16] cost 358.0 us
and 363.4 us -- sixteen rows for 1.5% more wall-clock.
So also spread a row over several groups when there are too few rows to cover
the device. Per-pass state moves to global memory and each digit pass becomes
its own launch, since a work-group barrier can no longer span the row. Groups
accumulate in SLM and contribute 256 global atomics each, keeping global
traffic per-group rather than per-element, and the last group of a row -- the
one whose fetch_add returns G-1 -- performs that pass's scan, holding the
launch count at one per digit plus one emit. The group count comes from the
device and is floor-divided by nrows, so a row count that already covers the
device is left whole and pays nothing. Below 64K columns the single-group
kernel finishes inside the cost of the extra launches and stays in charge.
Reading the row's prefix/mask/need through a device-scope atomic_ref costs
more than the sweep it guards: those loads are uncached, so passes 2-4 ran at
49 us against 12 us for pass 1. One lane reads them into SLM and the group
takes them from there -- 208 us -> 44.6 us at ne=[131072,1], k=2048.
The block size now takes the device's max_work_group_size instead of a cap of
512. The cap was never a floor, so a device reporting 512 is unaffected; one
allowing 1024 was being given half its width.
Finally, put the scan-merge gate where the two paths actually cross. That
kernel's cost climbs with k while the radix select's does not; measured over
widths from 2 to 200K columns and row counts from 1 to 8192, radix is ahead
everywhere from k = 8 up and behind at k <= 2, where scan-merge's smaller
fixed cost wins. The short-row corner (ncols=2, nrows=65536, as in bailingmoe2
group selection) is exactly where radix loses at low k, and the gate keeps it
on scan-merge.
Op-level against the CPU-fallback path this replaces, and against the
single-group radix select for the split: 4.98x at ne=[131072,1] k=2048,
6.65x at ne=[151936,1] k=40, 13.35x at k=20, 118x at ne=[65000,16] k=32.
No measured shape regressed. End to end on 3x Arc Pro B60 with
Qwen3.8-Flash-Next UD-IQ4_XS, llama-bench tg64, the parallelisation is worth
5.91 -> 6.05 t/s at d=131072 and a wash at shallower depths. Perplexity over
wikitext-2 is unchanged within noise at both 512 and 81920 context.
test-backend-ops: 525/525 TOP_K (previously every k > 32 case was refused),
880/880 MUL_MAT_ID. Perf coverage added for k > 32 at large widths and for the
short-row corner, neither of which was exercised before.
* move topk-select to topk-radix.{cpp|hpp}
---------
Co-authored-by: cwriter <cwriter@localhost>
Disable the ggml-cpu precompiled header and remove the
std::hardware_destructive_interference_size branch from CACHE_LINE_SIZE.
The PCH force-includes ggml-impl.h before ops.h, which pulls in <new>
via <array>/<vector> and defines __cpp_lib_hardware_interference_size.
This makes the C++ kernels use CACHE_LINE_SIZE = 256 (hardware
destructive interference size) while the C work-buffer sizing code in
ggml-cpu.c always uses the fallback 64. The mismatch undersizes the
rope work buffer by (CACHE_LINE_SIZE/4 - 16) * n_threads * 4 bytes,
causing a heap-buffer-overflow that corrupts the heap and later crashes
in ggml_compute_forward_rope_flt.
Disabling the ggml-cpu PCH restores the natural include order so
ops.h is processed before <new>, keeping CACHE_LINE_SIZE consistent.
Removing the std::hardware_destructive_interference_size branch makes
the value deterministic and include-order independent.
ref: https://github.com/ggml-org/llama.cpp/issues/28858
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
The workflow's push/pull_request path filters did not include the
ci/run.sh script that all of its jobs execute, so changes to it never
re-triggered the self-hosted CI.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
This commit moves the llama_n_rs_seq function call to before the
llama_decode call and returns directly if the check is true, removing
the setting of res and the goto statement.
The motivation for this change is to avoid the llama_decode call if it
is not needed.
* ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA)
* apply logic to NVIDIA as well
---------
Co-authored-by: Johannes Gäßler <[email protected]>
1) Combine two consecutive lookups (find + insert) into a single insert-attempt/lookup routine so that we don't per
form two O(log(n)) lookup operations in a row anymore -- we only need to do it once and then see if the insert succeeded.
2) Instead of copying every potential stack (expensive) and then moving it (cheap) to new_stacks when it's a final output state, we switch the order so that we move every potential stack (cheap), and then only copy it (expensive) to new stacks when it's a final output state. There are a LOT of intermediate states that get generated, and unless they become final output states, then all of these expensive intermediate copies are wasted.
Before: lookup -> lookup/insert + copy -> optional move to output
New: lookup/insert + move -> optional copy to output
The NextN/MTP tail loop derives the expert FFN size as n_ff/n_expert_used
when expert_feed_forward_length gives nothing for the layer. Both values come
from per-layer arrays that legitimately hold 0 on layers that are not MoE, so
a checkpoint whose predict layers hold 0 in both divides by zero and dies with
SIGFPE at load time, with no error message. Report the malformed metadata
instead.
Corrects a typo in `tests/test-quant-type-selection` for the
Nvidia Nemotron 3 Nano 30B A3B model, which was referred to as
*nvidia-nemotron-nano-3-30b-a3b*.
The error made the test skip that test case, rather than failing
the test.
[no release]
* fix for unsupport zes API
* optimize the code
* adjust the log level
* rm unused head files
* Update docs/backend/SYCL.md
Co-authored-by: Titaniumtown <[email protected]>
* fix the error to detect level zero SDK/dev package, stop build after detect the error
* update the message
* fix the build error when missed to install level zero dev package
* rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions
---------
Co-authored-by: Neo Zhang Jianyu <[email protected]>
Co-authored-by: Titaniumtown <[email protected]>
Co-authored-by: Neo Zhang <NA>
Move the EditorConfig Checker and Code Style Checker workflows from the
`[self-hosted, fast]` runners to `ubuntu-slim`, which is an established
runner label in the repo.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
- Clamp the -j parallelism to min(nproc, 2) so a single-core runner
uses -j 1 and multi-core runners use at most -j 2, instead of
unconditionally using $(nproc).
- Add a 3600s timeout to both test-backend-ops runs (the high-perf CPU
path and the default path) so a hung test cannot stall CI indefinitely.
- Note a TODO to reduce the timeout to 1800s in the future.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
There is a driver bug where two queues on the same VkDevice simultaneously
submitting can break some internal synchronization. Until it's fixed, add a
mutex around queuesubmit.
Clang stores the modification time of the precompiled header sources
inside the header and refuses the header when they differ. A cached
header restored from another checkout carries the timestamps of that
checkout, so the build fails. The option covers the compilers ccache
treats as MSVC while they are clang underneath, clang-cl and the Intel
LLVM drivers.
The child writes its state commands on stdout while the logger writes
on stderr, and both share a single pipe. The logger emits the trailing
color reset after the newline of a debug, warn or error entry, so that
escape sequence has no newline of its own and the router reads it glued
in front of the next command. The line prefix check then fails and the
command is forwarded as a log line instead of being handled, which
leaves a finished download stuck in the downloading state.
Writing the command with a leading newline closes the pending line so
it always starts at a line boundary.
Walk the binding offset back until the distance to the tensor is a
whole number of blocks, so block quantized views get a valid element
offset in the shader.
* hex-row-split: add support for multi-device row spliting
Co-authored-by: Max Krasnyansky <[email protected]>
* hex-mdev: add work splitting to fused kernels
* hex-mdev: use mdev_ prefix for all multi-device state
* hex-mdev: make device configuration more expressive to support device groups
* hex-mdev: fix mdev session init
* hex-mdev: fused nx (2x,3x) matmuls must update row counts for each w/o
* hex-mdev: fix MUL_MAT work partitioning bugs introduced by mdev
* hex-cont: fix crashes with new tests due to wrong striding
* hex-mdev: move fences after l2flushes
* hex-cont: fix work splitting for mnpu -- align chunks to cachelines
* hex-mdev: fix CPY tests with multi-dev
* hex-mmid: fix work partitioning with mnpu
* hex-mm: fix test failures with mdev
* hex-binary: fix work partitioning for mdev
* hex-argsort: fix mdev partitioning
* hex-mdev: fix work partitioning and general updates for all simple ops
* hex-fa: fix mdev work splitting issues
* hex-mdev: fixing more failing ops test
* hex-mdev: update the rest of the ops
* hex-mdev: refactor all mdev splitting logic to be contained within if (mdev_count > 1) {...}
* hex-mdev: fix macros
* hex-mdev: simplify session flush logic
* hex-sync: fix recursion in session flush
* hex-mdev: factor out fence buffer and allocator
* hex-fence: make fence allocation more robust with reserved slots for mdev
* hex-mdev: keep all mdev state in htp_mdev_group
* hex-mdev: further cleanup mdev group handling at the host
* hex-mdev: update group idx in the opbatch before serializing
* hex-batch: remove separate op_pending and use batch_req/rsp_seq
* hex-async: workaround another missing tensor_init in ggml-meta
* hex-fence: cleanup and robustify fences and error handling in multi-device scenarios
* hex-ar: improve ALLREDUCE error handling
* hex-async: robust error handling for op_cpy_fence
* hex-async: use seq0 from allreduce context to allocate fence_seq
* hex-mdev: fix remaining issues with fence and barrier clearing in CPY_FENCE
* hex-misc: realign macros and fix misplaces trace events
* hex-misc: align macros
* hex-mdev: fix unclone buffer re-entrancy
* hex-glu: fix mdev partitioning logic
* hex-mdev: make buffer uncloning/cleanup work with tensor-split scenarios
* hex-mdev: tighten up the can_split check in act-ops
* hex-mdev: factor out common bits of the partitioning logic
* hex-mm: minor realignment of the macros
* hex-bufs: fix incorrectly placed assert for MAX_BUFS
* hex-pad: tighten up gating checks for PAD
* hex-kparams: make sure all kernels properly use kparams->n_threads
* hex-docs: update user and developer docs with new features and detailed guide for ops development
* hex-scripts: update run script to properly parse dev groups
* hex-misc: formatting
* hex-sess: minor cleanup for session init
* hex-ar: fix vtcm size calc in allreduce kparams
* hex-scripts: fix flake8 warnings
* hex-rope: update ROPE to support mdev work split
* hex-ops: remove redunant checks and minor reformat
* hex-dev-guide: update dev-guide to avoid redundant null checks
* hex-async: improve event_wait, event_sync and fence implementations
* hex-async: remove synchronous flush from event_sync
* hex-async: symplify fence recovery protocol and make sync more robust
* hex-async: futher simplify error recovery for fences
* hex-err: return status instead of just -1
* hex-async: print all seq nums in hex
* hex-async: make sure fences flush dirty ranges
* hex-async: add dirty ranges merging to reduce fence flushes
* hex-async: properly sync before freeing the event
* hex-async: make sure fence owner session is not overriden
* hex-async: more fence write order more robust
* hex-async: make sure not to fuse ALLREDUCE+ADD if their dsts overlap
* hex-fusion: cleanup redundant checks
---------
Co-authored-by: Alexander Lu <[email protected]>
* ggml-webgpu: Update to a recent version of Dawn
* No module scanning
* Accept review suggestion to update comment
Co-authored-by: Masashi Yoshimura <[email protected]>
---------
Co-authored-by: Masashi Yoshimura <[email protected]>
* server: refactor subproc handling
* fix Windows build
* download: keep concurrent downloads of one blob apart
Every process writes the same path + .downloadInProgress, so a second
download of the same blob finds that file, takes it for its own partial
transfer and asks for the bytes after it, which produces a corrupt
result. The in-progress file now carries the pid of the process writing
it.
std::rename also replaces an existing destination on POSIX but fails on
Windows, so a download whose blob appeared in the meantime is dropped
after every retry and an etag rewrite silently keeps the old value.
std::filesystem::rename has the POSIX behaviour everywhere, and the
error now carries the reason reported by the system.
* Revert "download: keep concurrent downloads of one blob apart"
This reverts commit 917b83f149.
* tests: serialize the router tests that download the same model
Parallel workers share one cache, so the two tests fetch the same blob
into the same in-progress file and race to rename it. They now take a
file lock around the download, like the session fixture does for the
preset models.
* Revert "tests: serialize the router tests that download the same model"
This reverts commit c368a4a98c.
---------
Co-authored-by: Pascal <[email protected]>
* ci : run test-backend-ops as a dedicated gg test
Run test-backend-ops as a separate gg test in ci/run.sh so it is executed outside ctest. With GG_BUILD_HIGH_PERF it keeps the existing CPU-only invocation (-b CPU); otherwise it runs all available backends without a backend filter.
Remove the dedicated backend-ops workflow and keep test-backend-ops as a built target that is not registered with ctest to avoid duplicate runs.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : run test-backend-ops earlier and enable high-perf on kleidiai
Move the test-backend-ops gg test before test-llama-archs.
Enable GG_BUILD_HIGH_PERF and LLAMA_ARG_THREADS on the Graviton4 KleidiAI job and use the standard self-hosted results/mnt paths.
Add TODO markers for decoupling tests from libllama.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : run test-backend-ops in parallel
Pass -j $(nproc) to test-backend-ops in both high-perf and all-backend modes.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : disable parallel tests for ROCm
* cont : disable parallel tests with MoltenVK
* Test for nrc=2 as well | i8mm kernels
* Trigger only on supported HW
* Remove trailing whitespace
* Address review comment
* test: properly prepare nrc=2 inputs with independent data per row
* tests : make nrc=2 dot product inputs distinct
Assisted-by: Kiro
* tests : use non-trivial strides in nrc=2 dot product test
* tests : fail nrc=2 dot product test on non-finite errors
The expected success table holds when the four requests enter the shared
pool together. On a loaded runner they are admitted tens of milliseconds
apart, the slot lifetimes overlap differently and the pool overflows
while a short request is still resident. The decode failure aborts every
slot, so a request the table marks as successful comes back with the
context error instead of its generation.
Such a request now passes on that error too, while any other status, a
different error or a truncated generation still fails the test.
kernel_mul_mm_id splits its NR1 = 32 token tile into two 16-row halves and skips
the upper half when the expert did not fill it, on both the tensor and simdgroup
paths. The tB extents are corrected to (NK, NR1H) for the [NR1][NK] row-major tile.
The B tile is staged unconditionally, as on master: rows past nr1 restage a clamped
duplicate of a valid row, lie in the output-row dimension so they never contribute
to a valid row, and are dropped by the final store loop.
test-backend-ops: re-draw the expert ids between perf iterations of test_mul_mat_id
so MoE perf numbers are not warm-cache, and add token-tile boundary coverage using
n_used == n_mats, which routes every token to every expert so each expert receives
exactly n rows; n = 32, 33, 47, 48, 49 reach mul_mm_id and leave a last tile of 32,
1, 15, 16 and 17 rows.
* scripts : add initial profiling script (wip)
* src : add precompile headers (PCH) for models.h
* common : add common.h as PCH
* ggml : add PCH for ggml-impl.h
* mtmd : use PCH for models.h
* scripts : add script to build with Server/Tools/Tests
* server : add PCH for common.h
* docs: add profiling progress notes (wip)
* ggml : add exclude for GCC + SVE on ARM
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33393906061/job/99493756214?pr=28091
* ggml : attempt to fix use of std::hardware_destructive_inference_size
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/33396221677/job/99501265689?pr=28091
* squash! ggml : attempt to fix use of std::hardware_destructive_inference_size
Add a version check for GCC 12 to conditionally apply the `-Winterference-size`
pragma.
* editorconfig : exclude profiling reports dir
This directory will not be included in the merge later and this commit
can be ignore at that point. Just fixing to keep CI happy.
* ggml : skip PCH for gcc on non-x86 architectures
* tests : add PCH for peg-parser/tests.h
There are 7 peg-parser tests that can share one PCH instead of then each
parsing the full tests.h.
* common : add PCH for chat.h
* docs : update linux build profiling full results
Just updating after a number of PCH additions. These are not exact
figures and will vary a bit from run to run, but they give a general idea
of the performance impact of PCH.
* cmake : introduce unity build for models
This commit introduces a unity build for the models to improve
compilation time.
The improvements were roughly the following:
```console
+------------------------+-----+------------+------------+------------+
| Build | TUs | Frontend | Backend | Total |
+------------------------+-----+------------+------------+------------+
| Full, master | 396 | 811.0 s | 692.2 s | 1,503.2 s |
| Full, with PCH | 405 | 380.0 s | 664.7 s | 1,044.7 s |
| Full, with PCH + UB | 264 | 357.7 s | 635.7 s | 993.4 s |
+------------------------+-----+------------+------------+------------+
TU = Translation Unit.
Full = includes Server, Tools, and Tests.
PCH = precompiled headers.
UB = unity build for models.
```
* docs : update linux profiling table with unitiy build results
* docs : update mac profiling results to include unity build [no ci]
* docs: remove profiling reports
* scripts : merge build profile scripts into one script
I was lazy before and just copied the first script to enable Tests,
Server, and Tools. This now merges them into a single script.
* Revert "editorconfig : exclude profiling reports dir" [no ci]
This reverts commit 2922a12118.
* src : rename ggml_view_2d_slice to gemma3n_view_2d_slice
This is to be consistent with the rename in gemma4.cpp which was
required to avoid a name clash.
* cmake : add build profile script for windows [no ci]
This commit adds a port of the scripts/build-profile.sh script to
windows powershell.
This was developed on Windows on ARM but should work on X64 as well but
needs to be tested there as well.
* metal : rework fusion patterns into a single table
All fusable op patterns for the Metal backend are now declared once in a
fusion table (ggml-metal-fuse.cpp) and consumed by both the graph optimizer
(ggml_metal_fuse_max, packing) and the op encoders (ggml_metal_fuse_next,
compute). The two phases share the same pattern table plus ggml_can_fuse_subgraph_ext
for the structural checks, and differ only in the mode used for the pattern
check (STRUCTURAL at optimize time, since tensors are not allocated yet, and
FULL at compute time, including Metal buffer placement). This also protects the
snake activation (MUL + SIN + SQR + MUL + ADD) from being reordered during graph
optimization, which was previously unprotected.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : fix absolute output indices in fusion patterns
ggml_can_fuse_subgraph_ext expects the outputs array to contain absolute graph
node indices (it indexes cgraph->nodes[outputs[i]]), but the fusion table query
was passing a relative index (n_ops - 1). As a result the last node of every
pattern was not recognized as an output and was subjected to the elidable
use-count check, which failed for essentially all fusions. This silently
disabled the norm/MUL fusion and caused a ~5% token-generation regression.
Pass the absolute graph index of the last node instead.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : fuse gated_delta_net with cache cpy
Add GGML_METAL_FUSE_GDN_CACHE to the fusion table: when the gated_delta_net
kernel is followed by a cpy that scatters its recurrent state snapshots into
the KV cache, the kernel writes the snapshots straight into the cache buffer
and the trailing cpy is elided.
The gdn output has other consumers (the attn scores view), so unlike the
elision-chain patterns this is not a simple chain: a 'raw' flag on the fusion
pattern skips the generic chain/shape and ggml_can_fuse_subgraph_ext checks,
making the pattern-specific check callback the sole validator. Packing
(ggml_metal_fuse_max) now matches on the same view-transparent node sequence
that the compute phase uses, so the gdn + cache cpy group is packed along with
any intermediate views and stays adjacent through the reorder.
The fused cpy is a view consumer of the gdn (it writes the cache directly),
so its mem-range is skipped in the encoder; the skip is restricted to CPY
nodes consuming the previous fused node through a view so other fusions are
unaffected.
Add test_gated_delta_net_cache_fusion and register 5 cases.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : drop is_view_consumer mem-range skip
The is_view_consumer skip was carried over from the upstream gated_delta_net
cache-fusion draft, but it is not needed: keeping the elided cpy's mem-range in
the concurrency tracker only ever adds a (conservative) memory barrier at the
fusion point. It can never remove a barrier, so it cannot introduce a race. The
worst case is one spurious barrier per gdn+cache-cpy fusion, which is within
run-to-run noise on Qwen3.5-0.8B Q8_0.
Dropping the check keeps the mem-range loop uniform for all fused groups.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : rename gated_delta_net fused state output args
Rename the fused cache-write kernel argument to match the rest of the kargs:
state_out_stride -> nb_out (and widen it to uint64_t), and the local buffer id
bid_state_out -> bid_out.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : rename raw fusion flag to unsafe
raw did not convey that the flag opts a fusion pattern out of the generic
elision-chain safety net (ggml_can_fuse_subgraph_ext + chain/shape checks).
rename it to 'unsafe' to make explicit that the pattern's check callback is the
sole validator and must re-establish the safety guarantees itself.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : tidy fusion pattern checks and table
- const-correct ggml_metal_fuse_outputs buffer
- annotate unused check-callback parameters
- drop a redundant size_t cast
- align the ops/table initializers and add blank-line separation
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* metal : add generic fusion stats via ad-hoc proc-address API
Add a device-owned fusion context that lets a test tool count how many
times each fusion pattern fires and toggle fusion. It is exposed through
the ad-hoc ggml_backend_reg_get_proc_address mechanism with generic names
so the testing tool is backend-agnostic:
- ggml_backend_fusion_stats_init: start collecting fusion stats; when a
context is created afterwards it registers the labels/counters and
encodes single-threaded (n_cb == 0) so the counters are race-free
- ggml_backend_fusion_stats_reset / _get_stats / _set_enabled
The context lives on the metal device (not on the last backend context),
so counters accumulate across contexts and reads are always consistent.
The enable/disable toggle is initialized from GGML_METAL_FUSION_DISABLE
and can be overridden by the test through set_enabled. Labels are
synthesized from the fuse table via ggml_metal_fuse_label (e.g.
"GATED_DELTA_NET+CPY").
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : add fusion count regression test with per-backend baseline
test-fusion runs every dummy model generated by test-llama-archs on a
single backend (single-threaded encoding, n_cb == 0) with fusion enabled
and disabled, and for each mode (prefill / decode) reports the per-fusion
counters and the NMSE between the fused and unfused logits, plus the NMSE
against a CPU reference.
A fusion pattern that silently stops matching (or fires when it should
not) is caught as a regression by comparing the counters against a
committed per-backend TSV baseline:
- --record writes the golden baseline, --check (default) validates it
- the unfused run doubles as a control: its counters must be all-zero
- NMSE is skipped when it is NaN or the arch is already broken on the
device (e.g. plamo2 on Metal), so the count check is the hard gate
- baseline counts depend only on graph structure, not weights (verified
stable across weight seeds)
- the fusion stats API is resolved through the ad-hoc get_proc_address
mechanism with generic names; a backend that does not export it makes
the test fail with an error
The committed MTL0.tsv baseline covers 110 dummy archs (298 rows).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : rename fusion api helpers to match stats_init signature
Align the test with the ad-hoc fusion stats API: fusion_stats_init no
longer takes an enable bool (stats are turned on by calling it), so the
proc-address wrappers and typedefs are renamed to the api_* convention.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : rename backend to device in fusion test CLI
The fusion test operates on a compute device (e.g. MTL0), not a backend,
so rename the --backend argument to --device and the backend_name
variable to device_name. Keep "backend" where it refers to the ggml
backend interface (the ad-hoc proc-address mechanism).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : add --model and --help to fusion test
--model FILE runs the fusion regression test over a single model file
instead of enumerating a --models DIR. --models and --model are mutually
exclusive. Also add a --help/-h option that prints the usage.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : use backend base name for fusion baseline output
The fusion test is invoked with a specific device name (e.g. MTL0), but
its output - the recorded baseline and the header it writes - should be
named after the backend base name (e.g. MTL, via ggml_backend_reg_name),
since the counters depend on the backend, not on the specific device
index. Rename the committed baseline to MTL.tsv.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : run fusion test from ci instead of ctest
The fusion test needs Metal and generates a lot of dummy models, so it
does not belong in the generic ctest suite. Move it to ci/run.sh as
gg_run_test_fusion, gated on GG_BUILD_METAL like
gg_run_test_llama_archs_tensor_split: it generates the dummy models with
test-llama-archs -o and then validates the fusion counts against the
committed baseline. test-fusion.cpp is still built (llama_build) but no
longer registered as a ctest.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : align fusion baseline TSV columns
Pad the TSV fields to fixed widths so the columns line up regardless of
the variable arch and fusion-label lengths, and trim each field on parse
so the padded file is still accepted. Regenerate the committed MTL.tsv
baseline in the padded format (data unchanged, verified identical modulo
padding).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : widen label column and align fusion TSV header
Give the label column more room (28 chars) and fix the column header
widths so they match the data rows (moe/mode/label), keeping the header
aligned with the values. Regenerate the MTL.tsv baseline in the new
format (data unchanged).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : switch fusion baseline from TSV to CSV
Use comma-separated values like the rest of the project, keeping the
padded, aligned columns. Split on ',' and trim on parse. Rename the
committed baseline to MTL.csv (data unchanged, verified identical modulo
padding/separator). Update the ci/run.sh check path accordingly.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* cont : rebase + update MTL stats
* tests : avoid graph reallocations for some archs
* metal : tidy fusion debugging context and op init
- simplify the shared fusion debugging context comments
- shorten the ggml_metal_fusion struct comment
- align the ggml_metal_fuse struct fields and comments
- move the fusion parameter of ggml_metal_op_init right after dev
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : dedup fusion baseline into any mode
prefill and decode always produce the same per-graph fusion count, so
store a single row per label with mode = "any" and the per-graph count
instead of two rows. this halves the baseline size and keeps the check
stable.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* ci : move fusion model generation to a separate step
the dummy models generated by test-llama-archs are reused by other tests,
so generate them once in their own step instead of inside test_fusion.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : bump nmse thold
* models : fix plamo2 graph
* tests : remove "skip" logic from test-fusion
* tests : set qwen3tts dummy vocab to codec head size
the dummy qwen3tts model used a vocab of 4096 while the codec head is
3072, so the graph padded the output with -inf which made the NMSE in
test-fusion produce NaN. use the exact codec head size instead so the
padding is not generated at all.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* tests : regen fusion baseline
reflect the plamo2 graph fix, which changed its fusion pattern split
(RMS_NORM+MUL 11->10, RMS_NORM+MUL+ADD 3->4; same total).
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* ci : skip dummy model generation on OpenVINO
test-llama-archs does not build on the OpenVINO platform, so do not try
to generate the dummy models there.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* cont : minor
* tests : enable test-llama-archs on windows
* cont : disable on windows + workaround
* metal : naming nits
* test-fusion : add instructions to update baseline
* context : fix Kimi-K3 graph reserve
* fusion : update MTL
* cont : fix naming
* metal : rework fusion info storage
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : align fusion info API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* metal : use opaque fusion handle in ad-hoc API
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* ci : move fusion test to dedicated workflow
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
* cont : run only on ggml changes
* cont : simplify
* fusion : remove multi-output stuff for now
* ci : fix typo
* metal : fix idle threads in the remaining iq mul_mv kernels for ne00 < 1024
Generalize the row split from #28086 to the six other kernels that use the
same lane-to-block mapping: iq1_s, iq1_m, iq2_xxs, iq2_xs, iq2_s and iq3_s.
Each of them assigns one 32-element chunk per thread, so when a row has
fewer than 32 chunks the rest of the simdgroup is idle. When nb32 < 32 and
nb32 divides 32, 32/nb32 threads now share each chunk and each takes a
slice of the rows, reusing the FC_mul_mv_split function constant and the
dispatch wrapper introduced for iq3_xxs.
The plain path is untouched: wide matrices keep one thread per chunk and
N_R0_<TYPE> = 4. Only the split path uses N_R0_<TYPE>_SPLIT = 8. The
K-quants have the same idle-thread issue but a different lane mapping, so
they are left for a separate change.
* metal : offset the src0 row pointer once in the iq mul_mv kernels
q2, dh, sc, qh and signs are all derived from xr, so the row slice
offset only has to be applied to xr.
* metal : fold iq mul_mv row split into offset0
Compute row0 and row1 before initializing the source pointers and apply
the row slice directly to offset0.
This keeps x and its derived pointers on the existing path while applying
the split row offset once.
* server: fix speculation after an image
Pass the actual position to the drafter after an image, instead of the
token count. Affects every drafter, not just DFlash.
* rename draft n_past to pos0
n_past is used to denote number of tokens and this parameter is meant to be a position
* HIP: enable mma FA for head size 256 on RDNA4, tune configs
Assisted-by: Claude
Assisted-by: Codex
* HIP: prefer whole-tile FA grids over stream-k on AMD WMMA
Assisted-by: Claude
Assisted-by: Codex
* revise stream_k logic
* revise kernel selection logic
---------
Co-authored-by: Johannes Gäßler <[email protected]>
argsort had a data race in the inner loop, which VVL caught. But I don't think
this was causing failures in practice.
argsort_large has OOB accesses which might explain the failures in CI, but I
couldn't reproduce it locally and I don't think it's a convincing explanation
of the failures.
* vulkan: optimize m=1 mul_mat by swapping A/B
* vulkan: Improve small M perf
Allow split_k with small M.
Make small vs med tile selection (for coopmat2) depend on M, not just N.
* speculative: fix failed to decode mtmd chunk with DFlash
When using DFlash w/ vision models, the drafter memory fails to
allocate new tokens because images report a fixed offset. Stop copying
them to allow the drafter to continue.
* address PR feedback
limit M-RoPE skip to images only, allow audio to pass through. Clean up
comments to align to the updated implementation
This commit updates the version parsing in make-release-checks.sh to use
sed instead of grep. The motivation for this is that currently when
running this script on macos it errors:
```console
$ ./scripts/make-release-checks.sh --dry-run
grep: invalid option -- P
usage: grep [-abcdDEFGHhIiJLlMmnOopqRSsUVvwXxZz] [-A num] [-B num] [-C[num]]
[-e pattern] [-f file] [--binary-files=value] [--color=when]
[--context[=num]] [--directories=action] [--label] [--line-buffered]
[--null] [pattern] [file ...]
```
With the changes in this commit it is possible to run this without
failure.
Some of these if statements were copypastaed in a former refactor and
never cleaned up to remove the cases that could never happen anymore. The
only thing that's shared between these relatives anymore is
llama_model_bert::graph::graph, so the rest of the code doesn't need the
conditionals.
The Imagination proprietary Vulkan compiler returns VK_ERROR_UNKNOWN from
vkCreateComputePipelines for every dequant mul_mat_vec shader built with the
subgroup-only reduction that requires a subgroup size >= 16. That covers the
k-quants, the i-quants, TQ2_0, MXFP4 and NVFP4. ggml rethrows, so the first
generated token of any such model kills the process.
Reproduced on a Pixel 11 Pro (PowerVR C-Series CXTP-48-1536 MC1, driver
1.662.3024, subgroup size 128, min 32, max 128). The failure is independent of
subgroup size: 32, 64 and 128 all fail, as does dropping the full-subgroups
flag and the required-subgroup-size pNext. The legacy quants, which use the
plain subgroup reduction, compile and run fine.
The shared-memory reduction variant compiles and matches the CPU reference for
q2_K, q3_K, q4_K, q5_K and q6_K. The hybrid variant also compiles but costs
27% of token throughput (3.78 vs 5.20 t/s on Qwen3.5-2B-Q4_K_M).
* vulkan: use spec constant for mul mat type_a
vulkan: use map for mul_mm shapes
cleanup
fix indentation
fix cm2 and shmem init
fix cm2 spec constants
fix cm2 bindings
consolidate shmem tables and reduce size by type spec constant
fix compiler warning
fix missing Q2_0 type
fix unused warning when integer dot glslc support is missing
use minimal shmem size 8 instead of 1 to workaround cm2 compiler bug
fix missing Q2_0 type in cm2 matmul
fix types
* remove LUT quants from unified shader
* clean up
* restore coopmat2 q4_k/q5_k optimization
* split out q4_k/q5_k cm2 shader to fix Ampere regression
* revert iq shmem table renames
* simplify cm2 code with single uint8_t buffer
* fix fp4 extension use switch being overwritten by generic shader
* clean up
* adapt TQ1_0 changes
* adapt #27471 f16 Intel tuning changes
* divide workload to 2D
This is to workaround FILL exceeding maxComputeWorkGroupCount for Intel GPUs on Qwen 3.8 flash next
* minor change
* Fixed comment
Templates that default an optional variable to none and then test its
membership in a map hit an error, while the same expression is a normal
lookup returning false in Jinja. The undefined counterpart of this case
was already handled just above.
* vulkan: add dedicated iq4_xs mat-vec shader
Dedicated mul_mat_vec_iq4_xs for the dmmv path, replacing the generic fallback. ~+6-17% token generation on RDNA4 depending on model.
Assisted-by: Pi agent with Qwen3.8 27B
* vulkan iq4_xs: remove dead n_it unroll branch
Remove the n_it <= 8 experimental branch that attempted to fully unroll
the block loop. Since n_it is a runtime value, [[unroll]] is ignored by
the compiler, making both branches equivalent. Kept the simple loop
matching mul_mat_vec_iq3_s.comp.
The spacing eviction in create_checkpoint() keeps the oldest checkpoint and
erases every later one within checkpoint_min_step of it. For prompts shorter
than checkpoint_min_step this drops the checkpoint at n_tokens - 4 that the
next request resumes from, so hybrid/recurrent models re-prefill from the
previous checkpoint instead. Apply the spacing rule only once the list is at
n_ctx_checkpoints, and replace an existing checkpoint at the same n_tokens
instead of appending a duplicate.
* metal : fix half-idle simdgroup in kernel_mul_mv_iq3_xxs_f32 for ne00 < 1024
* metal : keep N_R0_IQ3_XXS = 4, dispatch a separate 8-row split kernel for ne00/32 < 32
The plain kernel is unchanged from master (4 rows per simdgroup, one thread per
chunk). The row-split mapping now lives in a separate kernel_mul_mv_iq3_xxs_f32_split
instantiation with N_R0_IQ3_XXS_SPLIT = 8, and the host selects it only when
ne00/32 < 32 and divides 32, so wide matrices keep the master kernel bit for bit.
* metal : select the iq3_xxs row split with a function constant instead of a separate kernel
* ci : add PYTEST_WORKERS=1 to fix server-self-hosted job
This commit adds the `PYTEST_WORKERS=1` environment variable to the
hf-jobs-t4-small:cuda13 runner steps.
This is an attempt to address CI failure of this job that I might have
introduced in Commit 42f0225fea
("server : use pytest-xdist for server tests (#28298)").
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/34126971262/job/101757819134
* apply same changes to server-metal steps
* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL
* vulkan : fuse UNARY(SIGMOID|SILU|SOFTPLUS) + MUL
- implement fusion in unary.comp behind UNARY_MUL_FUSION ifdef,
specialized pipelines per op instead of runtime branching
- fuse adjacent nodes only, ordering handled by graph_optimize
- drop runtime consumer scan and pending_unary_mul deferral
* vulkan : fuse UNARY(GELU|SIGMOID|SILU|SOFTPLUS) + MUL
1. GELU: gelu_mul_f32/f16 pipelines registered, CREATE_UNARY_MUL(gelu), GELU in dispatch + fuse gate + perf fusion name
2. Renamed/moved: gate is now ggml_vk_can_fuse_unary_mul(cgraph, unary_idx, mul_idx), placed with the other can-fuse helpers
3. norepeat both variants: each op gets plain (spec {0}) + _norepeat (spec {1}) pipelines from the same SPIR-V, selected via ggml_are_same_shape(src0, src1); the shape gate now allows broadcast (other dims equal-or-1)
4. graph_optimize: lambda deleted; standard "// UNARY + MUL: pull the consuming MUL forward" block added alongside the SSM_CONV/ROPE/MUL_MAT reorderings, with the same "other src must be weights or already processed" readiness check
* vulkan : align unary_mul fusion with binary kernel layout, relax gelu test tolerance
- schedule the fused kernel like mul.comp (256 threads x 2 unrolled
iterations), recovering a 10-18% prompt-processing regression
- allow 5e-7 f32 error for gelu_mul: the shader evaluates gelu with an
exp-based tanh identity while the CPU reference uses tanhf (~1 ulp)
* vulkan : use ggml_can_repeat in UNARY+MUL fusion shape check
The fused kernel indexes src1 via per-dim fastmod (generic_binary_head.glsl),
which is exact whenever the other operand tiles into the unary result -- not
just when its dims are equal or 1. Replace the hand-rolled loop with
ggml_can_repeat(other, unary) so the check matches the kernel's actual
capability and reuses the standard helper. Argument order matters: reversed,
it would wrongly admit graphs where the unary result is mul->src[1] and the
other operand is larger, producing truncated output.
Also add a rep_ne0 layout to the fused unary+mul backend tests covering a
non-1 repeat factor along dim 0.
* vulkan : fuse UNARY+MUL pairs separated by zero-compute nodes
gemma4's per-layer embedding gating builds gelu -> view_2d_slice -> mul,
where the intervening view is a zero-compute node aliasing an input that
was computed much earlier. Strict adjacency requirements meant neither
CUDA nor the vulkan unary+mul fusion handled this pattern.
Extend ggml_vk_graph_optimize to detect a UNARY whose consuming MUL is
separated only by unscheduled zero-compute nodes (GGML_OP_NONE, VIEW,
RESHAPE, TRANSPOSE, PERMUTE) and schedule those nodes ahead of the pair,
making it adjacent so the existing fusion applies. The reorder is guarded
by ggml_vk_can_fuse_unary_mul, a source-availability check for every
interleaved node, and the protected fusion patterns (topk_moe*, snake);
if fusion is later rejected the reordered graph still executes correctly,
just unfused.
Add a view_mid layout to the fused unary+mul backend tests replicating
the gemma4 pattern.
* vulkan : support OP-on-B in UNARY+MUL fusion
Some models apply the unary activation to the smaller MUL operand, e.g.
qwen3next/qwen35moe shared-expert gating builds ffn_shexp * sigmoid(gate)
with a [1,n_tokens] gate tensor. This shape was correctly rejected before:
the fused kernel derives its iteration extent from the unary tensor and
would leave most of the destination unwritten, and the generic same-shape
requirement in ggml_can_fuse blocked the pair outright.
Add UNARY_MUL_B_FUSION shader variants computing dst = src0 * OP(src1):
the OP operand rides the existing per-dim fastmod indexing, while the
iteration extent now comes from mul. Route {UNARY, MUL} pairs through a
local can-fuse variant that drops the generic same-shape rule and instead
requires the unary result to tile into mul->src[0] (ggml_can_repeat);
pairs with the unary as src0 keep the previous direction check, and
equal-shape pairs keep using the original pipelines.
Add a "gate" layout to the fused unary+mul backend tests covering the
shared-expert gate shape for gelu/sigmoid/silu/softplus in f32 and f16.
* vulkan : fold unary+mul view-hoisting into graph_optimize dep checks
Replace the dedicated UNARY + EMPTY* + MUL scanning block with two small
extensions to the existing scheduling logic:
- a consuming MUL may now join its in-set UNARY across a gap of unused
zero-compute nodes (NONE/VIEW/RESHAPE/TRANSPOSE/PERMUTE), instead of
requiring strict adjacency
- while doing so, such zero-compute blockers are ignored for this pair
Fusion validity is still decided later by ggml_vk_can_fuse at dispatch
time, so a rejected pair simply executes adjacent-but-unfused. Note the
relaxation must stay scoped to this pattern: exempting zero-compute
blockers globally reproduces silent output corruption on gemma3n.
* vulkan : select unary_mul OP-on-B via specialization constant
Replace the UNARY_MUL_B_FUSION compile-time shader variants with an
op_on_b specialization constant on the existing unary_mul SPIR-V,
mirroring how the norepeat flag is handled. The four {op}_mul_b_{f32,f16}
shader artifacts are gone - the OP-on-B pipelines reuse the base SPIR-V
with two-entry {norepeat, op_on_b} spec lists - and the duplicated store
expression is collapsed into a single runtime branch that the driver
prunes per specialization.
The constant is declared only under UNARY_MUL_FUSION so every other
binary pipeline keeps its single-entry specialization list.
* vulkan : replace unary_mul pipeline switches with a lookup table
Collapse the four nested selection switches in ggml_vk_unary_mul into a
single indexed lookup against a pipeline_unary_mul[4][2][2][2] table
([unary op][f16][norepeat][op_on_b]), whose trailing dims mirror the
{norepeat, op_on_b} spec constant list. The op axis uses a small shared
index helper that also replaces the switch in ggml_vk_can_fuse_unary_mul,
making it the only place that maps ops to the table.
Pipeline names are unchanged. Adding another supported op now requires
one macro invocation line and one helper case instead of edits in four
separate switches.
* vulkan : use ggml_can_fuse_subgraph for unary_mul pairs
Replace the hand-rolled pair validation in ggml_vk_can_fuse_unary_mul_pair
(bounds, op match, compute flags, single-use elision) with the shared
ggml_can_fuse_subgraph helper; backend-specific shape/type rules remain in
ggml_vk_can_fuse_unary_mul. Unlike ggml_can_fuse, the subgraph helper has
no same-shape requirement, so it covers both operand slots including
OP-on-B gates, and additionally rejects intermediates flagged as graph
outputs and validates view-source confinement.
The outputs parameter takes absolute node indices into the cgraph.
* Fix Whitespace
* vulkan : drop redundant unary_mul gap check in graph_optimize
The zero-compute nodes separating a UNARY from its consuming MUL are
already scheduled ahead of the pair by pass 2 of an earlier
optimization window, so the scoped gap tolerance added for this pattern
is unreachable in practice - disabling it leaves gemma-3n dispatch
counts unchanged (841 GELU_MUL per pass). Remove the flag, the empty
blocker exemption, and the now-unused gap helper, restoring the strict
adjacency requirement of the UNARY -> MUL pull-forward.
Keep the relaxation scoped out entirely: generalizing "zero-compute
nodes never block" beyond this pattern previously reproduced silent
output corruption on gemma3n.
* vulkan: fix whitespace (tab in indent)
* vulkan: fix whitespace (extra blank line)
* vulkan : move op_on_b spec constant to unary.comp
op_on_b is only used by the fused unary*mul path. Keep
generic_binary_head.glsl generic by defining it in unary.comp
instead. Same constant_id=1 and guard, no functional change.
* vulkan : make RMS_NORM/UNARY fusion gap-tolerant for views
Strict j==c+1 blocked RMS_NORM->MUL and UNARY->MUL when a
VIEW sits between (e.g. rms_norm -> view -> mul). Allow
c==back() with an empty-or-scheduled gap, matching the
review suggestion to check src linkage instead of adjacency.
Scoped to the two blessed pairs; safe because gaps can only
contain zero-compute nodes.
* vulkan : trim comments in UNARY+MUL fusion
Assisted-by: Muse Spark
On 32-bit platforms, Vulkan non-dispatchable handles such as VkBuffer are
represented as uint64_t, and Vulkan-Hpp disables implicit conversions for
type safety. This exposes two issues in ggml-vulkan:
1. vk::Buffer is streamed directly into std::ostream in debug/memory logs.
2. vk::Buffer is cast to VkBuffer before being passed to Vulkan-Hpp
CommandBuffer::copyBuffer APIs.
Fix these by add the operator<< for vk::Buffer, and
by passing vk::Buffer directly to Vulkan-Hpp copyBuffer calls.
* chat : split specialized parsers into common/parsers
Move the 14 dedicated template parsers out of chat.cpp into one file each under
common/parsers, mirroring the src/models split. chat.cpp keeps the template
detection in common_chat_try_specialized_template() and drops from 3915 to 1513
lines.
common/parsers/parsers.h holds the shared helpers and one declaration per
parser. foreach_function/foreach_parameter become inline there since nothing in
chat.cpp uses them any more; common_chat_template_direct_apply_impl and
common_chat_template_generation_prompt_impl lose static and carry their default
arguments in the header. Parser-specific helpers move with their parser:
is_lfm2_template, deepseek_v4_sort_tool_results and the gemma4 turn builder.
No functional change.
Assisted-by: Claude Opus 5
* chat : enumerate parser sources instead of globbing
file(GLOB) does not re-run CMake when a source file is added or removed, so an
incremental build silently keeps building the old set. List the parsers in
common/parsers/sources.cmake and include it from common/CMakeLists.txt.
Assisted-by: Claude Opus 5
* split helpers, add newlines
* tests: bind the L2_NORM batch count to a local
GCC cannot prove the loop fills norms up to the index read after it
while the bound is a class member, so it reports a maybe uninitialized
use. Reading the count once into a local restores the tracking.
* tests: initialize the L2_NORM batch array
The read after the fill loop is only provably defined once the array
carries an initializer, which GCC 12 requires on the aarch64 Release
build where warnings are fatal.
* server: fix LRU hang on multiple requests same model
* server: keep a queued model out of the victim pool until its waiters leave
A waiter that gave up while its model was still loading left the
model idle with no request behind it, and nothing recounted the free
slots, so a second request queued behind it stayed queued forever.
tick() was only driven by requests: join, claim and the end of a
proxied request.
Keep the queue entry alive after a successful claim so the model
coming up is never picked as a victim before its waiters use it, and
recount the slots on every status change and whenever a waiter
abandons the queue. The model is then evicted as soon as it comes up
with nobody left to serve.
---------
Co-authored-by: Pascal <[email protected]>
* vulkan: add DeepSeek-V4 hyper-connection fused ops (DSV4_HC_COMB/PRE/POST)
CUDA has these ops from the DeepSeek-V4 merge and Metal gained them in
PR 26459. Vulkan was the last major backend running the unfused primitive
chain. On DeepSeek-V4-Flash the unfused Sinkhorn comb chain alone takes
about 32% of decode op time on gfx1151 (Strix Halo), spread over roughly
16k dispatches per token.
dsv4_hc_comb runs the full 20-iteration Sinkhorn in registers. A token's
4x4 comb matrix lives in 16 consecutive subgroup lanes, with idst in bits
0-1 and isrc in bits 2-3 to match the CPU reference layout, so
subgroupShuffleXor by 1|2 reduces rows and by 4|8 reduces columns. One
dispatch replaces about 137 strictly ordered node executions per site.
The shuffle masks never cross a 16-lane boundary, so a subgroup of size
64 packs 4 independent tokens.
dsv4_hc_pre and dsv4_hc_post handle the elementwise stream collapse and
fan-out, with per-token coefficients staged in shared memory.
GGML_VK_DISABLE_DSV4_HC disables all three ops. The _COMB, _PRE and
_POST variants gate each op independently so a single kernel can be
bisected against the unfused graph.
Adds eval cases at the production n_iter=20 across batch sizes that
cross subgroup and workgroup boundaries.
* vulkan: dsv4 hc review fixes
Drop the per-op env-var disables and device flags, the stride divisibility
check (ggml guarantees it) and the workgroup-count fallback in supports_op.
Trim the comb shader comments to the lane layout.
---------
Co-authored-by: Kevin Hopper <[email protected]>
* adjust ncols_picker for routed MoE in mul_mat_q_case function
* Adding CDNA, RDNA2 and RDNA4
* fix: update mmq_use_routed_moe_ncols_picker to include NVIDIA + Volta support
* feat: enhance mmq configuration for various architectures with moe_ncols_min_cc support
* refactor: replace moe_ncols_min_cc with use_typical_moe_ncols in mmq configuration files
* HIP: mmq: enable typical moe ncols on RDNA4
---------
Co-authored-by: Carl Philipp Klemm <[email protected]>
* Update Q4_K and Q5_K to use branchless computation, which stops the scale unpack being re-executed for every column in mmvq, improving perf at batch sizes > 1
* Gating the change off from DGX Spark due to no gain
* Adding prefetch gated to Spark, making branchless change in Q4_K and Q5_K general and modifying switch points based on latest perf data
* Guard the mmvq L2 prefetch against MUSA as well as HIP
* Define the mmvq L2 prefetch only under the Spark guard
* Update switch point for Q4_K to accommodate more models
* Remove stale comments
* Add block_size to ggml_cuda_type_traits and create a separate mmvq_should_prefetch function
* Rename block_size to bs for cleaner indentation
* Fix build error on non-Spark CUDA arch with appropriate conditional around new function added
---------
Co-authored-by: praneshgo <[email protected]>
* vulkan: fall back to CPU for GET_ROWS with misaligned offsets
The Vulkan GET_ROWS shader asserts when a tensor's backing-buffer offset
plus view_offs is misaligned w.r.t. minStorageBufferOffsetAlignment
(see init_pushconst_tensor_offsets). Previously this caused a hard crash
on models using ggml_view + ggml_get_rows (e.g. Qwen3-TTS, Qwen3-VL).
Return false from supports_op() in the misaligned case so the scheduler
falls back to CPU, matching the existing pattern for PAD_REFLECT_1D and
other unsupported op/shape combinations.
Repro: llama-tts -m Qwen3-TTS-*.gguf -mm mmproj-*.gguf -ngl 99
Crash: GGML_ASSERT(dst->op != GGML_OP_GET_ROWS || (a_offset == 0 && ...)) failed
* vulkan: trim comment for GET_ROWS misalign fallback
* vulkan: fix file corruption in gated_linear_attn struct
* vulkan: properly handle misaligned offsets in GET_ROWS quantized path
- get_rows_quant.comp was missing get_aoffset()/get_boffset()/get_doffset()
calls that are already present in get_rows.comp, causing GGML_ASSERT crashes
when GET_ROWS operates on views with non-zero view_offs, as produced by
KV cache slices in Qwen3-TTS and Qwen3-VL.
- Remove the defensive misalignment GGML_ASSERT in init_pushconst_tensor_offsets
for the binary push-constants specialization, since both get_rows.comp and
get_rows_quant.comp now correctly apply per-tensor base offsets.
- Remove the workaround CPU fallback in supports_op() for GET_ROWS, since the
Vulkan backend now handles misaligned offsets natively (no more bailout).
- Add backend test coverage with view_src0=true (ggml_view_4d into a padded
tensor) for F32, F16, Q4_0, Q4_K, Q8_0, and I32 types, exercising both the
non-quantized (get_rows.comp) and quantized (get_rows_quant.comp) paths
with non-zero view_offs that reproduce the original Qwen3-TTS crash.
* tests: trim redundant comments in test_get_rows vs0 region
* tests: trim redundant comments in test_get_rows vs0 region (follow-up)
* vulkan: bind tensor base for binary ops, pass full view_offs via push constants
For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, MUL, etc.),
bind the view_src base and pass the full view_offs divided by type_size via
push constant misalign_offsets. This avoids truncation when misalign_bytes is
not a multiple of quantized block size.
ggml_vk_tensor_subbuffer gains a use_view_offs parameter. When false, the
binding points to vk_tensor_offset (base) and size includes view_offs.
init_pushconst_tensor_offsets<binary> computes a/b/d_offset directly from
tensor->view_offs, which is always row-aligned and therefore exact.
Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).
All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).
* vulkan: bind aligned offset for binary ops, pass adjusted misalign via push constants
For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.
ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types).
ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (target - adjusted_misalign) instead of the view_src base, keeping the
offset small enough for 16-bit/8-bit push constant fields.
Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).
All 223 GET_ROWS tests pass on Vulkan (NVIDIA RTX 5060 Ti).
* vulkan: bind aligned offset for binary ops, fix UMA offset mismatch
For ops using vk_op_binary_push_constants (GET_ROWS, ADD, SUB, etc.), bind
the buffer to an aligned position near the view offset (not the tensor base)
and pass the adjusted misalignment via push constants.
Added ggml_vk_tensor_physical_offset to unify physical offset lookup across
UMA and non-UMA devices. On UMA, resolves via ggml_vk_host_get(tensor->data);
otherwise uses vk_tensor_offset(t) + t->view_offs. Both get_misalign_bytes and
the new ggml_vk_get_adjusted_misalign helper build on top of this function,
so buffer bindings and push constant offsets are always consistent regardless
of device memory model.
ggml_vk_get_adjusted_misalign finds the smallest misalign that is both a
multiple of minStorageBufferOffsetAlignment and type_size, ensuring
misalign/type_size is exact (no truncation for quantized block types) while
remaining small enough for 16-bit/8-bit push constant fields
(adjusted_misalign < lcm(align, type_size)).
ggml_vk_tensor_subbuffer gains use_view_offs parameter. When false, binds
to (physical_offset - adjusted_misalign) on both UMA and discrete GPUs,
fixing a bug where the UMA host_get path previously skipped the adjusted
misalign binding and returned the target offset directly.
Added non-zero view offset (offset_rows=3) backend tests for GET_ROWS across
all_types with be1={1,7}, v={false,true}, skipping gradient setup for view
tensors (GGML_OP_VIEW fails ggml_set_param).
All 223 GET_ROWS tests pass on Vulkan (NVIDIA GeForce RTX 5060 Ti).
* finish misalignment fix
* supports_op changes for openvino/webgpu
---------
Co-authored-by: AiChiTuDouPian <[email protected]>
This commit adds the printing of the ggml version and commit to the
test-cmake example.
The motivation is just to be able to quickly verify that the correct
version of ggml is being used.
Example output:
```console
test-cmake] llama.cpp version: 0.4.0-dev, build: 10837 (5202104b5)
[test-cmake] ggml version: 0.23.0, commit: 5202104b5
[test-cmake] Initializing backend...
...
```
Problem
- Loader prefers `<arch>.attention.recurrent_layers`, falls back to `full_attention_interval` if missing
- Converter only ever writes the interval. gguf-py has no constant/writer for the array
- Interval can only describe evenly spaced full-attention layers. Any non-uniform `layer_types` gets reconstructed wrong
- No error, no warning. Model loads, runs, wrong layers get wrong ops. Full-attn layers marked recurrent lose their KV cache
- Every published Qwen3.5 checkpoint is uniform so nobody's hit it yet
Repro
12 layers, periods 4/3/5:
layer: 0 1 2 3 4 5 6 7 8 9 10 11
actual: L L L F L L F L L L L F
loader: L L L F L L L F L L L F
^ ^
Layer 6 is full attn, loaded as recurrent. Layer 7 the reverse.
52-layer non-uniform stack: 15/52 mis-typed.
Fix
- `constants.py`: add `Keys.Attention.RECURRENT_LAYERS` (name already registered in llama-arch.cpp)
- `gguf_writer.py`: add `add_recurrent_layers()`, same shape as `add_rope_pattern()`
- `conversion/qwen.py`: emit array from `layer_types` in `Qwen3NextModel.set_gguf_parameters` (covers 3-Next, 3.5, 3.5-MoE)
Notes
- Array is padded with `false` for MTP blocks. `get_key_or_arr` checks length against `n_layer_all`, which includes MTP. Matches the fallback's `i < n_layer()` guard
- Interval is still written. Old builds only understand the interval
- `layer_types` length != `num_hidden_layers` now raises in converter instead of producing a GGUF that fails at load
Tested
- End-to-end on a 62-layer non-uniform Qwen3.8-27B (2 linear layers removed). Loader reads the array, 62 blocks, 0 mismatches. Without fix: interval fallback, mis-typed
- MTP padding NOT tested on a real MTP model. Reasoned from qwen35.cpp + get_key_or_arr. Would appreciate a check
Co-authored-by: Claude Opus 5 <[email protected]>
Support RMS_NORM + MUL + ADD (+ MUL) and RMS_NORM + VIEW + SET_ROWS.
Extend ROPE + VIEW + SET_ROWS to support IMROPE.
Worth around 4% in gemma4 on my system.
This commit contains a suggestion for handling container images which
are currently not semver tagged, they only have build numbers in there
tags.
The proposed solution here is to first add a check to make sure that
there are container images built for the build number of the release and
if not fail the build. The container images are build nightly but they
can be triggered manually as well.
If the the container images check passes then the make-release workflow
will re-tag the images with the semver.
* vulkan: add TQ1_0 support (mm, mat-vec, dequant, get_rows)
* vulkan: pack TQ1_0 powers of 3 into a 32-bit constant
Replaces the constant array with a packed 32-bit value (7 bits per entry,
max 81 < 128) extracted with shift/mask, as suggested in review — avoids a
constant array that may not be kept in registers.
test-backend-ops on gfx1151: tq1_0 MUL_MAT 11/11, MUL_MAT_ID 6/6,
GET_ROWS 4/4, unchanged.
* vulkan: address review - shared TQ1_0 decode helpers, fix standalone dequant shader
Review feedback from jeffbolznv, all points:
- Move the packed-pow3 decode into shared helpers in types.glsl
(tq1_0_byte_of / tq1_0_digit_of / tq1_0_trit) and use them from
dequant_funcs.glsl, mul_mm_funcs.glsl, dequant_funcs_cm2.glsl and
dequant_tq1_0.comp instead of repeating the logic. The cm2 path also
drops its constant array for the packed-constant extraction.
- Translate all remaining comments to English.
- dequant_tq1_0.comp: use dequant_head.glsl. The shader previously declared
its own single-field push constant while the pipeline is created with the
5-field layout, so p.ne read the wrong field - confirmed broken, as
suspected in review.
- Fix wg_denoms for the standalone dequant pipeline: one invocation decodes
4 elements with local_size 256, so a workgroup covers 256*4 elements, not
256*16. With the old value the dispatcher launched a quarter of the
required workgroups.
Verified by temporarily forcing the dequant + f16 matmul path for TQ1_0
(hack not committed): test-backend-ops MUL_MAT passes through the rewritten
standalone shader, and the standard MUL_MAT / MUL_MAT_ID / GET_ROWS
tq1_0 cases still pass on Vulkan (AMD gfx1151).
* vulkan: address review — English comments, shared tq1_0_trit, trim TQ1_0 test cases
- mul_mat_vec_tq1_0.comp: drop leftover non-English comment and the local
POW3_PACKED constant; all decode sites now call tq1_0_trit() from types.glsl
- types.glsl / dequant_funcs_cm2.glsl: ASCII-only, drop stale reviewer note
- test-backend-ops: remove the oversized MUL_MAT_ID case (432 MiB A tensor,
~172 GFLOP reference); move the two remaining ones next to the other
backend-specific mul_mat_id one-offs and document why they are needed
* metal: decline TQ1_0 for GET_ROWS and mat-mul in supports_op
The new TQ1_0 cases in test-backend-ops exposed that the Metal backend
claimed support for GET_ROWS/MUL_MAT/MUL_MAT_ID with TQ1_0 sources while
having no such kernels (ggml_metal_library_compile_pipeline aborted on the
missing kernel_get_rows_tq1_0). Decline the type so the ops fall back to
the CPU, matching the existing NVFP4 handling on the same lines.
Assisted-by: Claude Fable 5
* vulkan: trim the TQ1_0 comments
Addresses @0cc4m's review: keep only what the code does not already say.
Removed the block-format recaps (the layout is right there in the struct) and
the step-by-step decode walkthrough. Kept the two facts a reader cannot infer:
the 8-bit truncation is part of the format, not an optimisation, and the powers
of 3 are packed into one uint so they do not end up in a constant array that
may miss the registers.
No functional change.
* vulkan: address review — trim comments, fold Metal check, drop unused _v
Per @0cc4m's review:
- dequant_funcs.glsl, dequant_funcs_cm2.glsl: drop the "see types.glsl"
pointers — they apply to every quant and say nothing specific.
- dequant_tq1_0.comp: drop the wg_denoms note. It is a precondition, not
information.
- mul_mm_funcs.glsl: same pointer removed.
- types.glsl: the comment on tq1_0_trit is down to the one fact the code
cannot show — the 8-bit truncation is part of the format, matching the C
reference, not an optimisation.
- dequant_funcs_cm2.glsl: removed dequantFuncTQ1_0_v and its define. You were
right that it is optional: it wrapped four scalar decodes and vectorised
nothing, and mul_mm_cm2.comp already guards the path with
`#if defined(dequantFuncA_v)` (DATA_A_F32 omits it the same way).
- ggml-metal-device.m: folded TQ1_0 into the existing NVFP4 check instead of a
separate block, and dropped both comments.
- test-backend-ops.cpp: the two mul_mat_id cases stay — they cover the
block-stride loop and the per-expert base offset that k == 256 alone never
reaches — but the comment is now one line instead of five.
Kept: the one-line labels on the three block regions in mul_mat_vec_tq1_0.comp
and on tq1_0_byte_of(). Those state the 5-trits-per-byte packing, which the
loop bounds do not show. Happy to remove them too if you prefer.
Re-verified on AMD gfx1151 (Vulkan), test-backend-ops, 2/2 backends passed:
MUL_MAT 9 TQ1_0 cases, MUL_MAT_ID 5, GET_ROWS 4 — all OK, no failures.
The coopmat2 path is unchanged apart from the removed _v define.
* models: use flash-linear-attention's l2norm for gated delta net q/k
The GDN q/k normalization is defined by flash-linear-attention as
l2norm(x) = x * rsqrt(sum(x*x) + eps)
with eps inside the root. Every GDN call site in the tree uses ggml_l2_norm
instead, which is x / max(sqrt(sum(x*x)), eps), i.e.
torch.nn.functional.normalize - its CUDA kernel cites that page.
The clamp never engages at these magnitudes, so in practice llama.cpp
normalizes with no epsilon at all where the reference has one inside the
root.
transformers made the same substitution when it first added Qwen3-Next and
corrected it three days later in huggingface/transformers#40842, 'Fix the
misalignment between the l2norm in GDN of Qwen3-Next and the implementation
in the FLA library'. vLLM and SGLang vendor FLA rather than reimplementing
it, so neither ever had the clamp.
eps keeps coming from the checkpoint, exactly as every call site already
passed it. The references hardcode 1e-6 for this norm; that is a separate
question and the two agree on every GDN checkpoint in the wild.
ggml_l2_norm itself is correct and unchanged, as is rwkv7-base, its original
caller, which passes normalize's own default eps of 1e-12.
No new ggml op: rms_norm already carries eps inside the root, so
rms_norm(x, eps/n) * (1/sqrt(n)) is exactly x * rsqrt(sum(x*x) + eps).
* Update src/models/models.h
Co-authored-by: Georgi Gerganov <[email protected]>
---------
Co-authored-by: Georgi Gerganov <[email protected]>
* ui : update active conversation fields in place
updateCurrentNode, applyConversationUpdate, updateConversationTimestamp
and the pin toggle replaced the whole activeConversation object, so its
identity changed on every send, tool result and rename. ChatMessages
tracks that identity to refresh sibling info, so each replacement
triggered a full refetch of every message in the conversation. Write the
changed fields instead, mirroring updateMessageAtIndex.
Assisted-by: pi:zai-org/GLM-5.3
* ui : reuse the conversation load read for sibling info
Opening a conversation read every message from the database twice: once
in loadConversation for the active path, once in ChatMessages for the
sibling map. Hand the freshly read array over once so the chat screen
builds sibling info from it, and set the conversation and its messages
in one sync block so effects never see the new conversation paired with
the previous one's messages.
Assisted-by: pi:zai-org/GLM-5.3
* ui : memoize leaf walks in sibling map build
buildSiblingInfoMap resolves each sibling's leaf by walking the last-child
chain, once per sibling per message, so the walk repeats along the same
chains for every message in the conversation ( O(messages^2) on long
chats ). Memoize leaf resolution per build with path compression so each
edge is walked once.
Assisted-by: pi:zai-org/GLM-5.3
* ui : skip sibling refetch for in-place message edits
refreshAllMessages refetches every message of the conversation just to
rebuild sibling info, but preserve-responses and non-branching assistant
edits never create branches, so the sibling map stays valid. Refresh only
after actions that branch (editWithBranching kept) or delete.
Assisted-by: pi:zai-org/GLM-5.3
* ui : drop unused currentResponse reactive writes
Nothing reads chatStore.currentResponse, but setChatStreaming reassigned
it on every streamed chunk, so each token paid a reactive write and string
assignment for nothing. Remove the field and the clearUIState wrapper
that only reset it.
Assisted-by: pi:zai-org/GLM-5.3
* ui : reuse completed agentic turn sections during streaming
deriveAgenticSections runs in a $derived invalidated per streamed chunk,
but re-derived every turn of the session each time, so per-chunk cost grew
with session length. Cache completed turns keyed by their assistant message
plus reference checks on every field that feeds derivation; only the
streaming turn recomputes. Cache hits return the same section objects, so
tool block props stay stable and skip their per-chunk re-derive.
Assisted-by: pi:zai-org/GLM-5.3
* ui : share markdown block infrastructure
Every markdown block duplicated shared work: a full copy of the hljs
theme CSS per instance, and the remark/rehype plugin chain rebuilt on
every processMarkdown call ( once per block at mount, again per coalesced
chunk while streaming ). Use the single theme style element already
maintained by SyntaxHighlightedCode, and build pipelines once - shared
process-wide for attachment-less blocks, cached by attachments identity
otherwise.
Assisted-by: pi:zai-org/GLM-5.3
* ui : measure assistant layout only for the last message
Every assistant message ran getComputedStyle, getBoundingClientRect and
a ResizeObserver over the previous user bubble at mount, even off-screen
ones, forcing a layout pass per message while a long conversation
renders. The measured vars only feed the :last-child min-height rule, so
gate the effect on isLastAssistantMessage; one measurement and one
observer remain, and the effect re-runs when the last message changes.
Assisted-by: pi:zai-org/GLM-5.3
* ui : trim whole-blob scans in tool block headers
Tool block headers parsed their entire blobs at mount, even collapsed,
and most tool results and args are large plain text or embedded file
content: skip JSON.parse unless the blob starts with a JSON container,
prefilter search-result extraction with a Title:/URL: substring check,
and match the end-anchored exit-code marker against only the tail of exec
outputs.
Assisted-by: pi:zai-org/GLM-5.3
* ui : parse write_file and edit_file titles without the content blob
Both block headers parsed the full args JSON at mount, even collapsed, and
write_file and edit_file args embed the whole file content or edit
strings, so every block paid a full-blob JSON parse just to read the path.
Split the meta into a title tier that extracts the path with a targeted
key match (full parse only as fallback) and a body tier that keeps the
full parse; Svelte deriveds are lazy, and the body snippet renders only
while the block is expanded, so collapsed blocks no longer parse args.
Assisted-by: pi:zai-org/GLM-5.3
* ui : mount chat messages lazily near the viewport
Every message row mounted its full component tree on load, so the cycle
collector, GC and layout invalidation kept walking every live object and
DOM node even for rows the user never scrolls to - which dominated the
profile of long conversations. Wrap each row in a placeholder with an
IntersectionObserver ( two viewport heights of runway ) that swaps in the
real ChatMessage when the row approaches the viewport; the row shell
keeps the content-visibility sizing, and rows stay mounted once
realized. Rows targeted by the pending-edit flow mount eagerly.
Assisted-by: pi:zai-org/GLM-5.3
* ui : smooth the chat navigation animations
Slide the centered new-chat form to the bottom edge with a transform
instead of a bottom offset - layout-property transitions need the main
thread every frame and stutter while a long conversation loads, while
transform transitions run on the compositor. Fade the message list in
with a CSS animation keyed to the conversation id, disabled under
prefers-reduced-motion.
Assisted-by: pi:zai-org/GLM-5.3
* ui : follow the svelte runes guidance in chat message code
Two effects detected changes with manual previous-value refs and reset
flags. The permission request carries object identity, so its dismissal
is now a derived comparing the dismissed request; the continue request
is a bare boolean, so its dismissal only shrinks to a reset while no
request is pending. Also drop a dead if (browser) guard in the markdown
theme loader - effects never run on the server.
Assisted-by: pi:zai-org/GLM-5.3
* test : pin the chat perf invariants in the unit suite
Cover the fixes whose silent regression would be stale or wrong UI rather
than a crash: the turn-section cache must reuse unchanged turns yet
recompute on every field it compares; the sibling map must resolve the
same leaves after the leaf-walk memoization; the active conversation must
keep its identity through field updates; and the blob gates ( exec tail
window, plain-text result gate, search prefilter ) must keep accepting
what they gate. Only the risky invariants are pinned - no coverage for
coverage's sake.
Assisted-by: pi:zai-org/GLM-5.3
* refactor : address review remarks
Name the tool-arg string-field pattern, move the file tools' path field
aliases and the JSON container gates into lib/constants, and export the
write_file / edit_file meta types from $lib/types instead of the parser
modules.
Assisted-by: pi:zai-org/GLM-5.3
Remove the build-time C++ helper and external gzip dependency,
simplifying cross-compilation. Keep the generated C++ in templates for
readability and preserve fully embedded UI assets.
Signed-off-by: Adrien Gallouët <[email protected]>
* Reapply "sycl : add Kronecker product FWHT support for sizes 384, 640, 768, 12…" (#28184)
This reverts commit c845263f8b.
* tests : fix unused variable M in test-backend-ops
* tests: fix trailing space error and isolate kronecker tests for sycl backend only
define two new environment variables to better understand how much
memory is being allocated, and when. This has been invaluable in
inproving the --fit algorithm, and is likely to be useful when debugging
other memory-related issues.
`-lv 4` will be required to enable the following:
GGML_SYCL_MEMTRACE=1 will show per-site memory usage, updated whenever
it increases by more than 64MiB.
GGML_SYCL_MEMTRACE=2 will show every allocation and deallocation.
To change the default 64MiB threshold for reporting memory usage increases, use
GGML_SYCL_MEMTRACE_STEP.
A sample log line:
[SYCL-MEMTRACE] device memory query (dev): total 59493 MiB, free 4494, in use 54998; allocated 0 (buffers 0 + scratch 0), peak 0 MiB
* ui : fix MCP image attachments not displayed in tool block (#25789)
Fixes regression from #25450 where ChatMessageAgenticContent passed
message.extra instead of section.toolResultExtras to tool blocks,
leaving tool images invisible. Also fixes TOOL_RESULT_JSON_OPEN_REGEX
which misclassified "[Attachment saved: ...]" as JSON.
Fixes#25789
Assisted-by: Muse Spark
* Addressed PR comments: 1.- Removed ·?? mesage?extra· as it has no case left to cover 2.- Added ·[\· to cover the case of ·[[1, 2], [3, 4]]· case suggested in the PR comment 3.- Added unit test for covering up this regex case
* ui : fix MCP image attachments not displayed in tool block (ggml-org#25789) - Addressed lint error on regex (redundant \)
* opencl: add extended elementwise unary ops (sgn, step, elu, hardswish, hardsigmoid, floor, ceil, round, trunc)
Adds nine GGML_UNARY_OP_* elementwise ops that were falling back to CPU on the
OpenCL backend, following the same variant shape as the existing ABS op: f32,
f32_4 (vec4), f16, f16_4 (vec4), and stride-addressed f32_nc / f16_nc for
non-contiguous inputs. New kernels/unary_ext.cl (macro-generated), a shared
ggml_cl_unary_ext dispatch helper mirroring ggml_cl_abs, the supports_op cases,
and the compute-forward cases.
Values are computed in float (the f16 variants read/write half and convert), so
the conditional ops (step, elu) match the CPU reference; the vec4 forms use
select() for the branch.
Validated with test-backend-ops on Adreno 840 and 850 (E17): all nine ops pass
every case including the vec4 and non-contiguous variants (8/8 or 14/14).
* opencl: dispatch a contiguous f32 copy over the whole device
kernel_cpy_f32_f32 maps one workgroup to each (i01,i02,i03) row and strides the
row across that workgroup's lanes, and the host launches ne01*MIN(64,ne00) work
items. A tensor with few long rows therefore runs on a single workgroup. The
mamba2 and gated-delta-net recurrent state cache is one row of 524288 floats,
copied once per layer per graph, and lands on 64 work items.
When both sides are contiguous the copy is a linear move, so dispatch it over
the whole device: one work item per float4. Gated on ggml_is_contiguous for both
tensors and equal element counts, so copies already spread over many rows keep
the existing path. The kernel is created optionally, so a driver that rejects it
falls back rather than aborting.
vload4/vstore4 rather than a float4 cast: they require only the scalar type's
alignment, and these buffers carry an arbitrary 4-byte view offset.
CPY, DUP and CONT are 217/217 on Adreno 840 and 740 with the path enabled and
disabled. GGML_OPENCL_CPY_FLAT=0 forces the old kernel.
* opencl: support all easy-copy types in CONCAT
CONCAT was F32-only. Extend it to every "easy-copy" type -- any non-quantized
type with a block size of 1 and an element size of 1, 2, 4 or 8 bytes, i.e.
f16/bf16/i8/i16/i32/i64 as well as f32.
The kernels are keyed by element SIZE rather than by type, which is what CUDA
already does for the same op: one kernel per byte width (b1/b2/b4/b8) plus the
packed b4 fast path, instead of one per ggml type. supports_op gates on the
same property, so a new type of a supported width is picked up with no further
work.
Validated with test-backend-ops on Adreno 840 / A8X and X2-90 / X2E.
* model: add Tencent Hy 4 (hy_v4) preview architecture support
Adds support for the Tencent Hy 4 model (Hugging Face architecture
HYV4ForCausalLM, GGUF arch hy_v4):
Add HF -> GGUF conversion script (conversion/hy_v4.py) and wire it into the conversion registry
Register hy_v4 GGUF constants, arch enum, and writer support
Implement the hy-v4 model graph, hparams, vocab and context changes
Register the new arch in llama-arch and models registry
Extend arch tests to cover hy_v4
Assisted by Claude Opus 5
* Update convert_hf_to_gguf_update.py
Co-authored-by: fairydreaming <[email protected]>
* Update conversion/base.py
Co-authored-by: fairydreaming <[email protected]>
* convert : move hy_v4 entry to the same place as in convert_hf_to_gguf_update.py
* model : apply changes related to n_ff_exp becoming per-layer in Hy4-preview
* n_layer_all
---------
Co-authored-by: fairydreaming <[email protected]>
Co-authored-by: Stanisław Szymczyk <[email protected]>
Co-authored-by: Sigbjørn Skjæret <[email protected]>
This commit adds a cmake version configuration file to replace the
current compile definition solution for the version.
The motivation for this change is that I made a mistake and did not take
into consideration that the compile definition means that this will
become a compiler flag for all sources in the target. This means that
when a version update happens that will recompile all sources in the
target even if they have not changed.
Refs: https://github.com/ggml-org/llama.cpp/pull/28278
Use std::error_code overloads of fs::current_path() and
fs::directory_iterator in ggml_backend_load_best() so an
inaccessible search path (WebDAV mount, removed CWD) is
skipped instead of terminating the process with an uncaught
filesystem_error.
Signed-off-by: Adrien Gallouët <[email protected]>
Let llama_print_build_info write to a caller-provided FILE* instead of
hardcoding stderr. The parameter defaults to stderr so existing callers
keep their current behavior.
The version command in llama-app now passes stdout, so plain version
output goes to stdout where users expect it.
Signed-off-by: Adrien Gallouët <[email protected]>
* src : add n_expert_used_max function
With Commit c61b98b875 ("model: add
NVIDIA Nemotron-3-Puzzle-75B-A9B (NemotronHPuzzle) support (#25444)") it
is now possible for each layer to have a specific number of experts but
there are a few checks that need to be updated to handle this upon model
loading. For example:
```console
llama_model_load: error loading model: model has expert layers but no expert layers are used
```
And later:
```console
/llama.cpp/src/llama-model-loader.cpp:955: GGML_ASSERT(n_ids_used > 0) failed
```
This commit adds the n_expert_used_max function so that these checks
can use it.
Refs: https://github.com/ggml-org/llama.cpp/pull/25444#issuecomment-5524976031
* src : use hparams.n_expert_used_max in llama_model_base::load_hparams
* src : use 0 as initial value for n_expert_used_max
Fuse RMS_NORM+MUL+ADD and ADD+ADD under GGML_SYCL_ENABLE_FUSION.
ADD+ADD uses the same binbcast indexing and type matrix as standalone
add() (f32, f16, f16/f32, i32, i16, bf16, including broadcast and
non-contiguous). Unsupported combinations fall back to two add() launches.
* opencl: quant lm_head / decode GEMV and medium-batch GEMM optimizations
* opencl: guard q4_K/q6_K tiled_ns convert-kernel registration for non-Adreno build
* opencl: gate q4_K MUL_MAT+GLU fusion dispatch to Adreno
* opencl: require the noshuffle weight layout in the q4_K GLU fusion gate
* opencl: do not take the vectorized f16 mrow GEMV path on an unaligned row stride
* opencl: pass the new get_scale_min_k4 stride argument at the row-major call sites
* opencl: enable the q4_K split-K decode GEMV only where it is measured to win
* opencl: record the X1-85 split-K datapoint (neutral, exclusion confirmed)
* opencl: restrict the tiled lm_head/embed GEMV default to X2E/A8X
* opencl: fix q4_K variant kernels to read the transposed scales layout
* opencl: keep the flat-GEMV large-m escape opt-in
* opencl: guard the o4 GEMV store against the rounded-up dispatch tail
* opencl: restore the tiled q4_K/q6_K layout on tensor read-back
* opencl: split-K for the q8_0 decode GEMV at small M
* opencl: keep the q6_K noshuffle correctness escape ahead of the opt-in gate
* server : use pytest-xdist for server tests
This commit adds pytest-xdist to the server tests. This is pytest
plugin that distributes test execution across multiple CPU cores.
Assisted-by: pi:llama.cpp/qwen3.8-27B
Refs: https://github.com/ggml-org/llama.cpp/pull/26734#issuecomment-5220707042
* remove server_base_port and BASE_PORT
* use worksteal and pytest builting tmp_path
* metal : support n_kv_max sparse mask hint in flash attention vec kernel
- add kernel_flash_attn_ext_vec_idx: compacts finite mask entries into
a per-row index list (Hillis-Steele scan, one threadgroup per row)
- extend vec FA kernel with optional sparse index gathering (FC slot 5)
- add host-side gate: sparse path when n_kv_max > 0, mask present,
supported head sizes / KV types, n_kv_max <= 4096
- new buffer region extra_idx for the index list
- pipeline getter extended with has_sparse param
- add test cases: head sizes, quant types, nb>1, nr23 variants,
sinks, ALiBi, softcap, permute, v_view_of_k, no-mask fallback
Note: multi-row (nb*nr23[1] > 1) cases still failing - rid mapping
in the store phase needs revisiting for the sparse path.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : fix sparse flash attention row addressing
- kernel_flash_attn_ext_vec_idx: mask param is half* but nb31 is a byte
stride, so the per-row mask offset was scaled by 2x; cast to char*
before applying the byte strides
- kernel_flash_attn_ext_vec: sparse pidx param is char* so the per-row
element offset was under-scaled by sizeof(int); scale it by sizeof(int)
to get the correct byte offset
- fixes the multi-row (nb*nr23[1] > 1) sparse flash attention failures
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* cont : use sparse vec FA for prefill
* metal : single-pass flash attention sparse index compaction
The idx kernel previously read the mask row twice: once to count the finite
entries (for the prefix scan) and again to recover their positions. Since the
kernel is memory-bound, this doubled the mask traffic.
Keep the finite positions in a per-thread register array during the count
pass and write them out directly, avoiding the second mask read. A dense
mask with more than NLOCAL finite entries in a slice falls back to re-reading
the mask to write the remaining positions.
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* tests : add perf cases for sparse flash attention prefill
Measure the sparse vec FA kernel across KV sizes, n_kv_max hints and batch
sizes. Run with:
./build/bin/test-backend-ops -b MTL0 -o FLASH_ATTN_EXT -p "n_kv_max=[1-9]" perf
Assisted-by: pi:llama.cpp/DeepSeek-v4-0731
* qwen4 : enable sparse attention
* cont : adjust nsg
* cont : sync test-backend-ops
* cont : disable Qwen4 for now
* cont : clean-up + tests
* mtmd : mark context as const in more methods
Mark `mtmd_context` as `const` in:
- mtmd_bitmap_init_lazy
- mtmd_tokenize
- mtmd_tokenize_from_parts
- mtmd_helper_support_video
- mtmd_helper_bitmap_init_from_file
- mtmd_helper_bitmap_init_from_buf
- mtmd_helper_video_init
- mtmd_helper_video_init_from_buf
- mtmd_helper_model_can_chat
The tokenization functions in particular are useful to have marked
`const`, as that allows more easily telling the compiler that we can
safely tokenize from multiple threads (`mtmd_tokenize` is already
documented as thread-safe, this just reifies that in the signature).
* mtmd : mark tokenization input pointer as const
Mark the `bitmaps` and `parts` pointers in `mtmd_tokenize` and
`mtmd_tokenize_from_parts` as `const`. This allows more easily calling
these with immutable arrays / vectors.
* mtmd : mark llama_context as const in mtmd_helper_model_can_chat
* CUDA: Allow CUDA optimization per split for multi-GPU.
Previous guard caused multi-GPU to skip the graph optimization. The
graph is already split per device and the optimization doesnt run
over the whole model but once per split, and thus should be allowed.
However, the CUDA event ggml_cuda_concurrent_event belongs to
whichever GPU was "current" when created. If the pass ran while
GPU 0 was current, it would stick and during event creation for the
second GPU it would land on GPU 0.
The fix: set the device explicitly ggml_cuda_set_device(cuda_ctx->device);
Default behaviour remains unchanged, only active for GGML_CUDA_GRAPH_OPT=1.
Explicit device setting pattern re-used from ggml_backend_cuda_graph_compute.
* Update ggml/src/ggml-cuda/ggml-cuda.cu
Co-authored-by: Aman Gupta <[email protected]>
---------
Co-authored-by: tannerbruhn <[email protected]>
Co-authored-by: Aman Gupta <[email protected]>
Skip the nb[3] check when ne[3] == 1, the shader never reads it for a
single stream. Cache views carry the full-buffer stride there, so the old
check reduced to n_kv == kv_size and the path only engaged with the
cache full.
* convert : skip bias_vl tensor in DeepSeek-V4 DSpark conversion
The DFLASH arch does not include FFN_EXP_PROBS_B_VL, so the DSpark
conversion failed when it tried to write the mtmd-only hash routing
tensor ffn.gate.bias_vl. Drop it like the tid2eid tensor; the DFLASH
draft only consumes ffn.gate.bias via FFN_EXP_PROBS_B.
Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-0731
* cont : fix
Co-authored-by: Sigbjørn Skjæret <[email protected]>
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* sycl: Q4_K Weight unpack optimization and reuse between destination Columns
* sycl: Q4_K small N (N=2..4) + two output rows by subgroup reuse of activation between two rows.
* sycl: gate Q4_K two-row reuse for small N=2
* sycl: Fix on magic number now uses Q4_K_MMVQ_ROW_PAIR_MIN_NROWS=6272 for it, added tests for coverage around Q4_K_MMVQ_ROW_PAIR_MIN_NROWS with perf support to test Q4_K MUL_MAT, applied the same reuse pattern to the activation as the weights.
Assisted-by: GPT-5.6 Sol
---------
Co-authored-by: RaulAbejonDelgado <[email protected]>
* hparams: add per-layer n_ff_exp/n_expert_used arrays with scalar-or-array loading
G1/G2 infrastructure for variable-per-layer expert FFN size and top-k routing
(required for Puzzle-75B which has 5 distinct n_ff_exp values and 7 top-k values
across its 40 MoE layers).
Design: rename scalar members to _impl suffix (following existing convention),
add LLAMA_MAX_LAYERS arrays, add n_ff_exp(il)/n_expert_used(il) accessors with
scalar fallback. No new GGUF keys: reuses existing expert_feed_forward_length and
expert_used_count keys via get_key_or_arr (scalar -> broadcast, array -> per-layer).
- llama-hparams.h: n_ff_exp -> n_ff_exp_impl, n_expert_used -> n_expert_used_impl;
add n_ff_exp_arr / n_expert_used_arr arrays; add per-layer accessor declarations.
- llama-hparams.cpp: implement n_ff_exp(il) and n_expert_used(il); out-of-range
il returns impl safely (shared code, no abort).
- llama-model.cpp: central n_expert_used load changed to get_key_or_arr; derive
impl as max-of-array for validations and backward compat; zero both new arrays;
HunyuanVL override also zeroes n_expert_used_arr.
- llama-graph.cpp: aggregation loop in build_moe_ffn uses hparams.n_expert_used(il)
so per-layer top-k bounds the ggml_view loop correctly.
- All other files: mechanical rename hparams.n_{ff_exp,expert_used} -> *_impl.
Scalar arches are unaffected (broadcast fills all array slots with the single value).
(cherry picked from commit 269a81e03d66e1c353e1a203a0c03a03eb2b1a4e)
* nemotron-h: use per-layer n_ff_exp(il) and n_expert_used(il) at MoE call-sites
Load n_ff_exp via get_key_or_arr into hparams.n_ff_exp_arr in load_arch_hparams;
derive impl as max for existing uniform GGUFs.
In load_arch_tensors, compute n_ff_exp_i = hparams.n_ff_exp(i) with fallback to
n_ff(i)/n_expert_used(i) for GGUFs that omit expert_feed_forward_length.
In build_ffn_layer, pass hparams.n_expert_used(il) to build_moe_ffn so per-layer
top-k is used for expert routing selection.
All other nemotron-h behaviour (mamba2, attention, shared-exp, latent projection,
routed_scaling_factor, expert_weights_norm, sigmoid gating) is unchanged.
(cherry picked from commit b1878a101793cd4e59868ac72635a86ea694987c)
* arch/*.cpp + gguf-py: mechanical rename n_ff_exp->n_ff_exp_impl, n_expert_used->n_expert_used_impl
All non-nemotron arch files continue using the scalar impl member directly.
Behaviour is identical: the impl value is the broadcast value from the GGUF scalar.
gguf_writer: add_expert_feed_forward_length and add_expert_used_count now accept
int | Sequence[int], mirroring add_feed_forward_length, so converters can write
per-layer arrays with the same existing GGUF keys.
(cherry picked from commit 8f009f54bea5ef9a6a354123bd25e9d5ea2d5e03)
* convert: support NemotronHPuzzleForCausalLM (per-block MoE config)
Parse block_configs/mtp_block_configs into per-layer arrays (scalar-or-array
keys), append the MTP [attention, moe] sub-blocks as blk.88/blk.89 with
nextn tensors, accept the backbone.* prefix, and register the arch.
Also fix a pre-existing undeclared _experts attribute on NemotronHModel.
(cherry picked from commit d1a592f278336e78457454eb6c96bca917135f10)
* nemotron-h: distinguish Nemotron 3 Puzzle (75B.A9B) from Super (120B.A12B)
Both have 88 layers; the per-layer expert_used_count array (heterogeneous
for Puzzle, broadcast-uniform for Super) is the discriminator.
(cherry picked from commit f824e09dc812169589cf5662c92d149a4c18c30a)
* convert: accept the official Puzzle BF16 checkpoint's tensor naming
The officially distributed BF16 checkpoint (NVIDIA-Nemotron-Labs-3-Puzzle-
75B-A9B-BF16) names the trunk model.* (model.layers.*, model.embeddings,
model.norm_f) where the original release used the NemotronH-style
backbone.*, and spells the router bias e_score_correction_bias instead of
e_score_correction.bias. Normalize both at the top of
NemotronHPuzzleModel.modify_tensors so either checkpoint converts; every
tensor name in the official index (42683 keys, MTP head included) resolves
through the tensor map after normalization.
(cherry picked from commit 189b67fc2c9d50970416c94b3317a6e7baa49b03)
* laguna: use n_ff_exp_impl for the uniform-MoE FFN size
Laguna landed after this branch was cut and reads hparams.n_ff_exp as a
scalar. This series turns it into a per-layer array with an n_ff_exp(il)
accessor, so the three scalar reads no longer compile. Laguna is a
uniform MoE, so point them at the scalar fallback n_ff_exp_impl, same as
deepseek2/qwen3moe/gemma4 in this series. No behaviour change.
(cherry picked from commit dbedc9e19c50dca0acdfb402362e2707bee424ae)
* arch: extend the n_ff_exp/n_expert_used rename to archs added upstream
kimi-k3, dflash, bailingmoe3, deepseek4, granite-swa and the nemotron-h MTP
block still referenced the scalar fields by their old names. n_ff_exp and
n_expert_used are accessors now, so those reads no longer compile; point the
non-per-layer archs at the _impl scalars and use the indexed form where the
call site is per-layer.
* convert: keep Puzzle opted out of the NemotronH MTP export path
#26725 added MTP export to NemotronHModel, keyed on num_nextn_predict_layers.
Puzzle's config carries that key, but NemotronHPuzzleModel bypasses
NemotronHModel.__init__ (its per-block config needs a different setup), so
_mtp_bid was never assigned and modify_tensors raised AttributeError on any
mtp.* tensor. Puzzle's head is also laid out by mtp_block_configs, not the
mtp.layers.* form the base maps.
Set _mtp_bid to None, drop mtp.* in filter_tensors, and declare
supports_mtp_export = False so --mtp / --no-mtp fail at the CLI.
* llama: replace n_ff_exp/n_expert_used scalars with per-layer accessors
Follow-up to review feedback: the previous revision kept the scalar
hparams fields alongside the new per-layer arrays, which duplicated
state that get_key_or_arr already handles by broadcasting a scalar
value over every layer.
Drop both scalars and expose n_ff_exp(il) / n_expert_used(il) built
exactly like the existing n_head_kv(il) and n_ff(il) accessors: they
index the array and GGML_ABORT out of range, with il defaulting to 0
so genuinely uniform call sites stay a plain n_ff_exp().
Arch loaders now read both keys through get_key_or_arr over
n_layer_all, and the n_expert_used validation checks the maximum
across layers instead of a single field.
* llama: restore per-key required flags on the expert hparam reads
The scalar-to-array conversion passed required=false at every call site,
which silently made mandatory keys optional. Each read now carries the
same required flag it had before the conversion.
* server : accept data: URLs for input_video and input_audio
input_video and input_audio passed accept_base64_uri=false to
handle_media(), so data: URLs got treated as raw base64 strings and
failed later with a confusing media probe error (#27724).
pass true for these two content types the same way image_url already
does, and allow video/audio mime types in the data: url check instead
of image only. data URL validation now throws std::invalid_argument so
malformed input comes back as 400 instead of 500, matching the other
input validation in this file.
* server : simplify handle_media and drop unused accept_base64_uri flag
* server : update comment and add unit test for invalid data URI MIME
Extend the HTP backend's F16 unary op coverage to include ABS on top
of the existing NORM/RMS_NORM/L2_NORM/SCALE/CLAMP/SQR/SQRT set.
- Add hvx_abs_f16_{aa,au,ua,uu} + dispatcher in hvx-arith.h, mirroring
the sqr_f16 kernel structure and using the existing hvx_vec_abs_f16()
sign-bit-clear helper
- Add abs_f16() row-wise dispatch and DEFINE_UNARY_TASK_F16(unary_abs, ...)
in unary-ops.c, wired into execute_op_unary()'s op_type/task_func
switches
- Register HTP_OP_UNARY_ABS in htp_op_is_unary() (unary-ops.h) so that
ggml_hexagon_precompute_unary_params() fills kernel_params (n_threads,
VTCM layout) for ABS nodes -- required for the F16 path to function
- Narrow the F16 GGML_OP_UNARY gate in ggml_hexagon_supported_unary()
(ggml-hexagon.cpp) to allow GGML_UNARY_OP_ABS specifically, instead of
rejecting all GGML_OP_UNARY ops for F16
- Merge the separate execute_op_unary_f32()/execute_op_unary_f16()
functions into a single execute_op_unary(), branching on an is_f16
flag for the parts that actually differ by type (elem_size, the
early F16 op-support check, and which task_func table to use) while
keeping the F32-only tiled/RMS_NORM_MUL paths intact -- per review
feedback to avoid duplicating the shared VTCM/DMA plumbing
Verified on-device (QRD8850, Hexagon v81) via test-backend-ops -o ABS:
8/8 passing (F16 + F32, HTP0, no CPU fallback). Regression-checked
SQR/CLAMP/SQRT (F16+F32) and NORM/RMS_NORM/L2_NORM/SCALE (F32; their F16
paths have no CPU reference kernel in test-backend-ops and cannot be
correctness-tested there independent of this change).
* common, server : enable preserve_reasoning kwarg by default, log its effective state
If the preserve_reasoning chat template kwarg is not specified explicitly
via --reasoning-preserve / --no-reasoning-preserve, it is enabled by
default after argument processing. The server logs the effective state of
the kwarg, warns that it is enabled by default when the template supports
it, and only warns "has no effect" when it was enabled explicitly on a
template that does not support it. Setting the kwarg via
--chat-template-kwargs is deprecated.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : update comment
Co-authored-by: Xuan-Son Nguyen <[email protected]>
---------
Co-authored-by: Xuan-Son Nguyen <[email protected]>
* mtmd: load the qwen3-tts code predictor proj_in as optional
The talker and the code predictor share the hidden size on the 0.6B
checkpoints, so the reference builds no small_to_mtp_projection and
the conversion emits no tensor for it. The graph already falls back
to identity when the weight is missing, the loader now agrees.
* mtmd: keep the qwen3-tts code predictor ffn_down in F32
The code predictor carries a massive activation: its layer 2 FFN
intermediate peaks around 1.5e5, well past the 65504 ceiling of F16.
mul_mat casts its input to the weight type, so an F16 ffn_down turns
that peak into inf, the residual follows, and the next rms_norm yields
NaN. Reference forward in float32 gives 145109 against 145396 measured
in the graph.
* hex-mm: fuse QKV and FFN matmuls that land on HMX
* hex-mm: remove hardcoded ne[1] < 32K restriction
* hex-get-rows: explicitly reject repacked Q8_0 just in case somebody decided to add an override
* hex-mm: correct overhead sizing to make sure we dont exceed vtcm budget for large dims
* hex-mm: fuse MUL_MAT_ID into MUL_MAT_ID_NX (2x,3x,...) where possible
* hex-fusion: update opbatch and opqueue sizing to acount for new fusion and reduce overhead for trace buffer alloc
* hex-bufs: sort buffers while finalizing opbatch, helps avoid va space fragmentation
* hex-bufs: add simple va defrag to make sure we dont abort just because the va space is fragmented
* hex-mm: replaced more scalar divs with fastdiv and minor cleanup
* hex-mm: tighten up supported fusion checks to exactly match supported kernels
* vulkan: handle larger batch sizes (>4) efficiently for IQ3_S mat-vec when NUM_COLS > 4. 5x perf at n=8
Assisted-by: Claude Opus 5
* adds 2 cases per quant type at `k=16*256` to the `all_types` mat-vec sweep
---------
Co-authored-by: Marshall <[email protected]>
When building with gcc < 15, CMakeLists.txt unconditionally adds
ime2_kernels.cpp, which fails to compile. FindSMTIME.cmake only defines
RISCV64_SPACEMIT_IME2 when the IME2 instructions are detected, and gcc 14
only has IME1, so ime2_kernels.cpp hits its #error.
This PR fixes it by using IN_LIST to add each kernel source according to
the spec that was actually detected.
* opencl: clamp the q4_K decode GEMV's fetch row on a padded x-grid
* opencl: enforce the tiling contract of the image KQ/KQV GEMMs
* opencl: decide the image KQ/KQV split at the dispatch, not from strides
* cuda : fuse MoE weighted reduction (mul + view + add)
The MoE combine tail currently writes weighted expert outputs to
global memory before reducing them. That intermediate global-memory
traffic is the main cost. The production baseline generally runs two
physical fused kernels; this path runs one.
This change matches the full expert-weighting plus ordered-reduction
subgraph and replaces it with one weighted-reduction kernel.
Supported graphs:
- unscaled: experts * router_weights
- scaled: (experts * expert_scale) * router_weights
k = 2..15 is handled by one runtime-k kernel.
Matching is structural: op sequence, shapes, strides, expert views,
and the left-to-right ADD chain. The fused kernel keeps that same
reduction order. Results are not claimed bit-identical; CUDA FP32
contraction can change rounding slightly.
Allocator integration uses add_alloc_dep from the graph-optimizer
API so experts, router weights, and optional expert scales stay live
until the fused destination is written. Memory ranges are rechecked
before the fused kernel runs.
Unrecognized or unsafe graphs are left alone and keep the existing
per-op path. Set GGML_CUDA_MOE_WEIGHTED_REDUCTION=0 to disable the
fusion.
test-backend-ops covers scaled/unscaled, aligned/unaligned, and
representative values across k=2..15, plus a k=16 case that must
stay on the per-op path.
* Pruned the test matrix from 15 to 6
* Addressed the aman and olivers review comments
get_prev_tokens() rebuilt a (seq, pos) -> token hash map on every
ubatch by walking all used cells, while llama_kv_cells already keeps
an ordered index of the positions of each sequence in seq_pos, updated
on every cell mutation to serve seq_pos_min() and seq_pos_max().
The index now stores (pos, cell) pairs in a std::set instead of a
position -> count map, so a repeated position (cache reuse via rm + add,
vision inputs with shared positions) yields distinct entries and the
removal of a cell erases its own pair. The new seq_pos_tok_le() returns
the token of the cell at the largest position <= p in logarithmic time,
which is exactly what the old window lookup and its M-RoPE gap fallback
computed together.
get_prev_tokens() shrinks to a direct lookup per (token, offset) and
for_each_token_in() goes away with its only caller. The kv-cache keeps
no n-gram logic of its own.
Measured on Qwen3.8-Flash-Next UD-Q4_K_XL at 71k context, alternating
two binaries with the first run discarded: tg 69.3 -> 72.7 t/s (+4.9%),
pp unchanged at ~2720 t/s, greedy output identical, needle retrieved.
* metal : fix more leaks due to missing autoreleasepools
* metal : rename variable
* metal : fix another missing pool warning
Co-authored-by: YiChen Lv <[email protected]>
---------
Co-authored-by: YiChen Lv <[email protected]>
* qwen4exp: follow up fixes
* -kvu NaN collapse fix
Assisted-by: Claude
* indexer cache ext.x/ext.y restore fix
Assisted-by: Claude
* kv-cells: rename seq_set to seq_get_all
seq_get is already taken by the single-id getter, so the suggested name
cannot be overloaded on return type alone.
Assisted-by: Claude
* memory-hybrid-idx: implement set_input_qsa on the memory class
The context held the whole implementation, where the pattern elsewhere is a
thin context forwarding to the memory class, as llama_kv_cache_context does
for set_input_kq_mask. The body reads no context state, so it moves unchanged
and the context keeps a forwarder.
Also shortens the seq_get_all comment as suggested.
* tests: check that a sequence state survives a save/restore round-trip
Saves seq 0, erases it, restores the blob and saves again, requiring the two
blobs to match. Compares blobs rather than generated text, which cannot see a
field dropped on the way back in.
Note this passes on master for qwen4exp, so it does not demonstrate the
ext.x/ext.y drop this PR fixes; reaching that needs 2D mrope content.
* tests: give the synthetic qwen4exp a PLE so the state test bites
has_cell_ext() is n_pos_per_embd() > 1 || ple_n_heads > 0, and the indexer
cache sets rope_type = NONE, so without a PLE it serializes no cell ext at
all and the round-trip test cannot see a dropped ext.x/ext.y. With one,
removing the ext_set restore in state_read_meta fails the test: 198 of
335692 bytes differ, first at offset 282092.
Loading such a model needed two fixes:
- the row count of per_layer_token_embd came from require_weight(), which a
model synthesised from metadata alone has no file to answer. Derive it
from the head ranges and prefer the file's padded count where there is one.
- the PLE conv history is a row of the recurrent cache, so a PLE on a full
attention layer dereferenced a null p_l. Reject it at load time instead.
The meta mirror is skipped for qwen4exp. It returned NaN logits before this
fixture carried a PLE, which the nmse check passes since a NaN comparison is
false, and aborts with one. -sm tensor on real devices works.
Assisted-by: Claude
* llama: disable -sm tensor for qwen4exp
test-llama-archs skipped the tensor split for this arch from inside the
test, so the arch still advertised support it does not have. Declare it in
llm_arch_supports_sm_tensor instead and drop the test-side exception; the
existing llm_arch_supports_sm_tensor branch then does the skipping.
Assisted-by: Claude
* metal : request Metal 4.0 language version for the tensor API
* metal : load the tensor API kernels from a separate metallib
* tests : add external-metallib tensor API regression test
* metal : fix metallib build order for the tensor API kernels
MTP speculative decoding needs the target state to move back by the
number of rejected draft tokens. Without rollback support the context
is classified as SEQ_RM_TYPE_FULL and the server serializes the whole
recurrent state to host memory on every round, which costs more than
the drafting saves.
The recurrent cache already holds n_rs_seq + 1 snapshot planes and the
delta net writes its SSM state into them, but build_conv_state_at wrote
a single plane, so a rollback restored a convolution history that was
never captured. It now writes one snapshot per slot, each ending one
token earlier, for the delta net QKV convolution and for the PLE
convolution alike.
Measured on Qwen3.8-Flash-Next UD-Q4_K_XL with the standalone MTP
draft, n-max 3 and a single slot: decoding reaches 183 tok/s on code
and 144 tok/s on prose. The same branch before this change, where the
server falls back to checkpointing the state to host memory, reaches
123 and 83 tok/s, for 108 tok/s without a draft.
* qwen4exp: sum the indexer heads by slices
The head reduction went through a transpose and a sum_rows over ne[1],
which left sum_rows with ne0 = 4, one block per row for a four element
reduction, and the transpose copied the whole block by token surface
twice on the way in.
The heads are adjacent on ne[1], so each one is a strided view and the
sum is a short chain of adds.
RTX PRO 6000, Qwen3.8-Flash-Next UD-Q4_K_XL, fa on, 55k context, warm
runs on top of #28011:
prompt processing 2170 -> 2366 t/s
Generation is unaffected. The removed work scales with n_blocks by
n_tokens, so the gain grows with context and with ubatch size.
* qwen4exp: drop the redundant cont on the indexer query
rope returns a freshly allocated, contiguous tensor, so the reshape that
feeds the matmul does not need a copy. ggml_reshape_3d asserts
contiguity, so a layout that would need the cont cannot slip through
silently.
Greedy output is unchanged token for token.
Address review from @ggerganov
* Batched gemm for grid IQ quants
Style updates and a bit more performance
Clean up comments
Move code around
Vectorize IQ panel decode, lower threshold for speedup
IQ panel: single-source gather layout, gate bias, vectorize interleave
Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer
Move IQ panel code out of repack into iqp.cpp, clean up comments
Another comment sweep
* Add myself as iqp.* codeownder
* Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size
* Renaming and moving
* The other half of renaming and moving
* Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition
* Update ggml/src/ggml-cpu/iqp.h
Co-authored-by: Georgi Gerganov <[email protected]>
* Add iqp_rows work buffer
* Revert "Add iqp_rows work buffer"
This reverts commit 425542991e.
* Add NUMA fallback
* Add 10 row batch tests for IQP coverage on all grid IQ types
* Swap assert for return false in support check
* Move IQP mul_mat_id test
---------
Co-authored-by: Georgi Gerganov <[email protected]>
* kv cache : batch state restore scatter reads per contiguous run
When restoring state into non-contiguous destination cells (e.g. a
prompt-cache snapshot into a fragmented ring), state_read_data issued
one small copy per KV cell - ~1.4M copies of a few KiB each for a
40k+ token restore, taking 25-63 s on the CUDA backend.
The snapshot stores cell rows in cell order, so a maximal run of
consecutive destination indices maps to one contiguous block and can
be restored with a single copy. Precompute the runs once and use them
in all three scatter loops (K, V, transposed V). Byte-identical.
The on-device reader copies with a byte cursor when the read and
write chunking differs, so the batched reads are safe for it as well.
Batching makes equal tensor counts with a different split reachable
(save ranges [2,1] vs restore runs [1,2]); the next commit teaches the
reader's 1:1 path to fall back to the byte cursor in that case.
Verified in a production setup: 1,363,616 copies / 25-63 s -> 224
copies / 221-424 ms for the same restores (42,603 cells, 4 runs).
Assisted-by: Claude Code (unsloth/qwen3.8-27b)
* context : fall back to the byte cursor when read and write chunking differ
the on-device reader copies saved state back with a 1:1 copy by tensor
index whenever the write and read sides recorded the same number of
tensors, guarded by a per-tensor size assert.
equal tensor counts do not imply equal chunking: a state restore may
batch its reads per contiguous run of destination cells while the save
used per-range reads, so both sides can record two tensors that split
the same data differently, and the assert aborts in all builds.
compare the per-tensor sizes and only take the 1:1 path when the
chunking actually matches, otherwise fall through to the existing
byte-cursor copy. both sides enumerate the same logical data in the
same order, so the cursor copy is well-defined across tensor
boundaries.
Assisted-by: Claude Code (unsloth/qwen3.8-27b)
* tests : cover state restore scatter reads on host and on-device paths
decode the same prefix on two sequences, interleaving the seq 0 cells
between the seq 1 cells, so the seq 1 cells are isolated from each
other in the kv cache (three cells, two saved ranges). save the seq 1
state, free the interleaved seq 0 cells, and restore: the destination
is then non-contiguous (two runs), and the restore-side chunking has
the same tensor count as the save-side with a different split, so the
scatter path is batched per contiguous run and the on-device reader's
byte-cursor fallback is exercised.
the restored state is saved again on the host and compared byte for
byte with the first save: the blob is serialized in sequence cell
order, so the two saves are identical if and only if the scatter
restore wrote exactly the same KV content. this documents the
byte-identical guarantee of the run-batched scatter reads.
one test per io backend: the host (CPU) path and the on-device path.
Assisted-by: Claude Code (unsloth/qwen3.8-27b)
* ui: copy the displayed text of grouped agentic responses
Agentic sessions render as a single entry anchored on the first
assistant turn, whose content is typically just the first tool call,
so the copy button wrote an empty string to the clipboard. Derive the
text sections of the whole session and copy them joined, matching the
visible response. Plain messages keep the previous behavior.
* const
* dflash : fuse the encoder into the KV injection decode
The encoder is a single fc + norm, but running it as a separate
llama_encode forced a device-to-host round trip of its output before the
injection decode could re-upload it, plus a second graph build per
round. Fold the encoder into the decoder's embd branch and feed the
target features directly to one llama_decode.
Assisted-by: Claude Fable
* nit
* Apply batched suggestions from code review
Co-authored-by: Ruixiang Wang <[email protected]>
* Fix missing references from renaming
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
Co-authored-by: Ruixiang Wang <[email protected]>
* vulkan: RDNA3 static mat-vec rows above four columns
On RDNA3 above four columns a static 4 rows for all types benches faster than
the default.
* vulkan: RDNA3 static mat-vec-id rows
mul_mat_vec_id has no column dimension to switch on. On my Strix Halo machine,
a static 4 is faster here than the defaults across types and batch sizes.
* rpc: avoid serializing buffers from other servers
Only include remote buffer pointers when the buffer belongs to the RPC dispatcher receiving the graph. Add a two-server regression test for cross-server tensor serialization.
Assisted-by: Codex
* cont : add ref
---------
Co-authored-by: Georgi Gerganov <[email protected]>
for_each_token_in tested all LLAMA_MAX_SEQ sequences for every used cell,
while a cell almost always belongs to one. The scan now stops once the
cell's own sequences have been seen. Same visit order, same callback
arguments, so behaviour is unchanged.
get_prev_tokens is the only caller, so this affects the n-gram path.
RTX PRO 6000, Qwen3.8-Flash-Next UD-Q4_K_XL, fa on, warm runs:
55k context generation 56.3 -> 74.3 t/s
132k context generation 33.6 -> 50.9 t/s
Prompt processing is unchanged, the scan is amortised over the ubatch
there. The gain follows the number of used cells, so it grows with
context and is invisible on short prompts.
The optimized path grouped warp lanes by token and required
warp_size % n_expert_used == 0, with a single hardcoded exception
padding 6 up to 8. Every other count fell back to the generic path,
which walks the tokens one at a time with a warp reduction per token,
for each of the n_expert blocks.
The lane group only has to divide the warp, and the loop body already
guards the padded lanes with iex < n_expert_used, so the padding
generalizes to the next power of two. The 6 -> 8 case and every count
already dispatched keep the exact same padding as before.
n_expert_used = 10 now reaches the fast path. Measured on
Qwen3.8-Flash-Next (512 experts, 10 used) at 55k context on an
RTX PRO 6000, warm runs with the first one discarded:
prompt processing 2334 -> 2600 t/s
Token generation is unaffected, since a single token leaves nothing to
walk. Other expert counts reach the fast path by adding their case to
the dispatch.
- DFlash2 NVFP4 draft models produced almost no accepted speculative
tokens because the Q, K, V, and output projection scales were not
passed to the corresponding graph operations.
Rename the --tensor-read-lazy CLI argument to --lazy-mode, to match the
internal lazy_mode parameter, and add a -lzm shorthand. Sync the READMEs.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
some backends (Metal, SYCL, WebGPU) require additional memory for
fleeting data for certain ops, which is reflected in their
get_alloc_size implementations.
add ggml_backend_op_alloc_size_may_expand() to the backend utils,
listing these ops, and assert in ggml_backend_buft_get_alloc_size
that a backend expanding the alloc size of a compute op only does so
for ops listed in the helper.
use the helper in the RPC backend to decide whether to query the
remote server for the actual alloc size, instead of a hardcoded list.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* hex-devices: add support for lazy session allocation and cleanup dev interfaces
Co-authored-by: Marco Colombo <[email protected]>
* hex-devices: support for runtime discovery of available NPU cores
Co-authored-by: Alexander Lu <[email protected]>
Co-authored-by: Ehsan Bateni <[email protected]>
* hex-devices: reject non-existing devices early during init
---------
Co-authored-by: Marco Colombo <[email protected]>
Co-authored-by: Alexander Lu <[email protected]>
Co-authored-by: Ehsan Bateni <[email protected]>
* metal : fail closed on mul_mat shapes with missing F16 kernels
* metal : abort on nil pipeline in encoder_set_pipeline
* metal : address review comments
* metal : share mul_mat mm dispatch with supports_op
* opencl: default the Adreno xmem F16xF32 GEMM on for X2E
kernel_mul_mm_f16_f32_l4_lm is the slowest matmul this backend has on Adreno: on
the X2-90 it runs the gpt-oss-20b attention projections at roughly a quarter of
what the tuned dense q4_0 GEMM reaches on the same device. That matters for any
model whose non-expert weights stay f16 -- the stock gpt-oss-20b release is
exactly that, and its prefill spends 40.8% of GPU time in that one kernel. The
xmem route already existed but was left opt-in, so nobody hit it.
Worth about 25% prefill on gpt-oss-20b on an Adreno X2-90. Gated to X2E: the
Adreno 840 measures neutral. Decode is untouched -- the dispatch gate needs
N >= 16. It is worth nothing on the q8attn variant, whose attention weights
already take the dp4a dense GEMM.
The env var was presence-tested before, so =0 previously enabled it; it is now
atoi()'d. MUL_MAT 963 OK / 0 FAIL on both arms.
* opencl: bypass the tiled f32 GEMM on the Adreno A7X
The A7X (E031.41) compiler executes kernel_mul_mm_f32_f32_l4_lm at roughly a
tenth of what the same silicon reaches in its own f16 and q4_K kernels. It
allocates 488 B/WI of private memory against 304 for the same source on the
following generation, i.e. the older register allocator spills in the K-loop.
Models with per-layer F32 projection pairs kept F32 by quantization policy land
on this kernel twice per layer, and it dominates their prefill on that part.
Route batched f32xf32 (ne11 > 8) around the tiled path on the A7X and let it
fall through to the per-row f32 kernel, which that compiler handles fine; small
batches keep the tiled path. Weights stay GPU-resident, so decode placement is
untouched -- declining the op in supports_op instead was measured first and
rejected, because the per-layer CPU round-trips cost more decode than the
prefill it gained.
Worth about 9% prefill on gemma-3n-E4B on an Adreno 740, with MUL_MAT counts
identical on and off. No other generation is affected. Override with
GGML_OPENCL_A7X_F32_LM_BYPASS=0.
* opencl: enable xmem GEMM for adreno by default
---------
Co-authored-by: Li He <[email protected]>
improve the --fit algorithm to take into account the actual peak
required VRAM for a given context size on a SYCL backend.
This includes both properly accounting for how much VRAM is required
when the allocated context is fully used (which makes the reported
context drop below what it did before, but stop it OOMing) as well
as preventing some overly-conservative calculations which meant too much
VRAM was being reserved.
Tested on a Arc b70 with unsloth's qwen3.8 (Q4_K_XL), able to get 262144 context,
fully usable, with q8_0 KV and MTP and 4k ubatch size using --fit-target 1
The N padding is needed for mul_mat, but not mul_mat_id. For mul_mat_id,
we indirect the row index through a shared memory lookup table which avoids
any OOB row coordinate. But that callback doesn't bounds check K, so we
actually need K padding instead.
* port setup-build.ps1 to setup_sdk.py, to facilitate installation of Hexagon and OpenCL SDKs on Windows
* rename setup_sdk.py -> setup-sdk.py
* flake8 fix: print() -> logger.info()
---------
Co-authored-by: Kristopher Urquhart <[email protected]>
* vulkan: fix missing view-alias dependencies in ggml_vk_graph_optimize
is_src_of doesn't treat two views of one tensor as dependent, so the optimizer reorders nodes across aliased reads and writes.
Result: silently wrong tokens under greedy decoding, different output on every server start, and invalid speculative-decoding acceptance, with nothing logged.
Hits Qwen3.8's recurrent state (and any model with view-aliased state) on AMD and NVIDIA Vulkan. CUDA is clean.
Compare view_src bases on both sides.
Fixes#27805
* vulkan: don't treat view/no-op nodes as aliasing dependencies
Nodes whose op is NONE, RESHAPE, TRANSPOSE, VIEW or PERMUTE execute nothing, so aliasing through them is not a real dependency. The previous base comparison matched them anyway, which only costs the optimizer reordering freedom.
Co-authored-by: Jeff Bolz <[email protected]>
* vulkan: make the lambda parameter const and capture is_empty in is_src_of
Code will not compile without these changes.
is_src_of has an empty capture list, so is_empty was not visible inside it, and is_empty took a non-const pointer, while is_src_of receives const ones. Other call sites pass non-const pointers, which still convert as usual.
---------
Co-authored-by: Jeff Bolz <[email protected]>
* ggml : fix conv_transpose_2d for multiple batches
ggml_compute_forward_conv_transpose_2d_impl only computed the first
batch (ne[3] of the destination); every batch after the first was left
as zero. Both the src1 permutation and the main compute loop now iterate
over the batch dimension, and the work buffer size in ggml_graph_plan is
scaled by the src1 batch count so the extra permuted batches fit. A
multi-batch test case is added to test-backend-ops.
Fixesggml-org/ggml#1448
* metal : fix conv_transpose_2d for multiple batches
The kernel only computed batch 0 of the input (src1->ne[3]); every
output batch after the first was left as zero, so multi-batch
conv_transpose_2d results diverged from the CPU reference.
The grid now covers all batches (OW x OH x OC x N), the kernel decodes
the batch from the grid z coordinate and offsets both the input and
destination indices accordingly. nb3 is passed in the kernel args.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
---------
Co-authored-by: Georgi Gerganov <[email protected]>
* vulkan: add hoisting support for row IDs and expert count in shaders
* use hoisted row ids in coopmat2
* vulkan: address review feedback on count_experts
- use vk_op_count_experts_push_constants instead of a raw uint vector
- apply the fastdiv trick to the ne00 div/mod in count_experts
- compute the per-expert offsets with subgroupExclusiveAdd when the
device supports it, keeping the serial path as fallback
- document the data_d layout and the hoisted_row_id_words bound
- drop a leftover debug print in ggml_vk_matmul_id
* vulkan: use init_pushconst_fastdiv for count_experts push constants
* vulkan: refine comments for row ID hoisting and data layout in count_experts shader
* Whitespace
---------
Co-authored-by: Jeff Bolz <[email protected]>
This adds fa_vec_tuned_table records for Apple M4 to ggml-metal-tuning.cpp.
Includes F16, Q4_0, Q4_1, Q5_0, Q5_1, and Q8_0. (M4, 10 GPU Cores)
Co-authored-by: Strongtut <[email protected]>
* OpenVINO Backend: Fuse IM2COL + MatMul convolution into OpenVINO convolution
* ci:ggml-ov: Skip recurrent state rollback tests
* ci:ggml-ov: Skip recurrent state rollback tests
* Update OPENVINO.md
* ggml-openvino : add env-var gated op support debugging
* Fix ggml_rope_set_offset case
* OpenVINO backend: Support Whisper.cpp
* Fix code style
* openvino : enable qwen35 on NPU
Static shapes:
- get_graph_input_shape() left the s_copy / s_copy-leaf inputs dynamic
([1,1,1,-1]) even in static mode, which propagated a dynamic slot dim through
GET_ROWS into the conv/GDN state, the state reshapes and the GDN output.
- With -np 1 the s_copy defrag remainder gathers zero rows; short-circuit that
CPY to the untouched cache instead of emitting a degenerate Slice/Concat, and
skip binding its zero-byte ggml tensor as an output (the dynamic path already
did the latter, the static path wrote the full cache over a 0-byte buffer).
Token-count independence:
- In static mode the compiled model's token count is the prefill chunk size or
1, not the captured cgraph's. Offsets derived from the captured count were
therefore wrong. Anchor the GDN state slice at the end of the packed
[attn | state] output and drop the rs_src_begin runtime inputs, and make
VIEWs over the GDN output / conv_input pass through so the consumer does the
slicing.
- CONT could not identify its token axis when the graph was captured with a
single token (every trailing dim has the same stride and size 1) and baked
the captured shape into the prefill model.
Chunked prefill:
- The last chunk is padded with fabricated tokens. Attention masks them, but
the recurrent path folded them into cache_r/cache_s permanently. Add a
chunk_valid_len runtime input, use it to zero g and beta for padded steps
(making the recurrence an exact identity) and to end the conv snapshot window
at the last valid token, and disable the recurrent-cache reset after the
first chunk so earlier chunks are not wiped.
- get_is_prefill() and the chunk loop bound read inp_pos->ne[0] directly, but
IMROPE stacks 4 position planes, so every decode step was run through the
padded prefill model and the loop ran extra out-of-bounds chunks.
cache_rs_reset_idx/len now stay runtime Parameters in static mode, since
can_reuse_statically() does not invalidate the cached model on ComputeParams
changes. Add GGML_OPENVINO_FORCE_STATIC to exercise the static path on CPU.
* Update to OpenVINO 2026.3.1
* ggml-openvino: forward NPU compilation mode parameters
Add GGML_OPENVINO_NPU_COMPILE_CONFIG to the backend's cached environment so callers can configure the NPU compiler without using the generic property escape hatch.
When the value is non-empty, pass it to OpenVINO as NPU_COMPILATION_MODE_PARAMS. This enables settings such as optimization-level=3 for NPU compilation while preserving the existing behavior when the variable is unset and leaving CPU and GPU configuration unchanged.
Document the variable, its NPU-only scope, and the optimization-level=3 example in the OpenVINO backend runtime configuration table.
* ggml-openvino : support RELU, POOL_2D, QUICK_GEGLU, and ROLL ops
* reorder op table
* exclude GPU/NPU failing POOL_2D case
* move op type detection to compute_op_case
* Relax rope supported cases
* Fix pool case
* Update openvino doc, gpu driver in ov docker
* openvino: remove unused static remote context branch
* openvino: parallelize static model build
* Apply editorconfig
---------
Co-authored-by: Mostafa Faheem <[email protected]>
Co-authored-by: Ravi Panchumarthy <[email protected]>
Co-authored-by: zhaixuejun1993 <[email protected]>
Measured at a live KV length of 34816 (32768 depth plus one 2048 ubatch),
on Qwen3.8 27B Q4_K_S:
per tensor 4 * 34816 * 256 * 2 B = 71.3 MB
staged per call K and V, so 2x = 142.6 MB
traffic per call read once, write once = 285.2 MB
traffic per ubatch 285.2 MB * 16 calls = 4.56 GB
One ubatch is one ggml_cgraph submission (llama_context::process_ubatch ->
graph_compute), so that 4.56 GB is the cost of a single 2048-token prefill
chunk, and it scales with the live KV length: the first ubatch of the same run,
at seq = 2048, moves 0.27 GB.
Reproduce the two measured inputs with:
GGML_SCHED_DEBUG=2 llama-bench -m MODEL -p 8 -n 0 -r 1 -ngl 0 \
-fa on -ctk f16 -ctv f16 -v > nd.txt 2>&1
grep -E 'n_layer|n_head_kv|n_embd_head_k' nd.txt
awk '/node # 0 /{g++} g==1 && /\(FLASH_ATTN\)/{n++} END{print n+0}' nd.txt
* metal : add fa-vec tunings for M5
This is a followup contribution to efeda76b94 as requested in https://github.com/ggml-org/llama.cpp/discussions/27668 to add support for additional Apple GPUs. I generated this output using the provided instructions:
```sh
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
cmake -B build -DGGML_METAL=ON
cmake --build build --target ggml-metal-tuning -j
./build/bin/ggml-metal-tuning fa-vec --dtype f16,q8_0 > fa_vec_rows.txt 2> fa_vec_sweep.log
```
This ran on a machine with Apple M5.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : add fa-vec tunings for M5 Pro
This adds fa_vec_tuned_table records for Apple M5 Pro to ggml-metal-tuning.cpp.
Contributed by SerayaEryn in https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18157544 (F16, Q4_0, Q8_0; M5 Pro, 20 GPU cores).
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* metal : add fa-vec tunings for M3 Max
This adds fa_vec_tuned_table records for Apple M3 Max to ggml-metal-tuning.cpp.
Contributed by TeeAaTeeUu in https://github.com/ggml-org/llama.cpp/discussions/27668#discussioncomment-18175220 (F16, Q8_0; M3 Max, MacBook Pro 64GB, low power mode).
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* cont : whitespaces
This is a followup contribution to efeda76b94 as requested in https://github.com/ggml-org/llama.cpp/discussions/27668 to add support for additional Apple GPUs. I generated this output using the provided instructions:
```sh
git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
cmake -B build -DGGML_METAL=ON
cmake --build build --target ggml-metal-tuning -j
./build/bin/ggml-metal-tuning fa-vec --dtype f16,q8_0 > fa_vec_rows.txt 2> fa_vec_sweep.log
```
This ran on a MacBook Pro (14-inch, Nov 2024) with Apple M4 Pro. The `ggml-metal-tuning` command completed successfully in 1h 13m 1s with no other notable load on the system.
* tests : run test-save-load-state across all architectures
test-save-load-state previously only ran in ctest against a single
downloaded model (tinyllamas/stories15M), i.e. only the llama arch.
Add a --models DIR mode to test-save-load-state that runs the full
save/load suite over every *.gguf in a directory, reporting a
per-model PASS/FAIL and exiting non-zero if any model fails, and wire
a ctest to run it over all architectures using the existing
generate-models fixture (test-llama-archs). The single-model -m mode
is preserved (still used by ci/run.sh).
Also bump the dummy-model training context in test-llama-archs from
128 to 256 so that the per-sequence context (which is padded up to a
multiple of 256) no longer exceeds n_ctx_train and emits the
"possible training context overflow" warning.
The test is expected to fail until the affected arches are fixed:
deepseek4 (host seq-copy), gemma2/gpt-oss/lfm2 (device seq-copy),
minimax-01 (state load). It aborts at the first arch that crashes.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : match dummy DSA indexer to fused Lightning Indexer kernel
The dummy DSA indexer (deepseek32, glm-dsa, ...) used key_length=64 and head_count=1, so the fused Lightning Indexer op's q tensor was shaped [64, 1, ...]. The Metal fused kernel is fixed to DK=128, NH=64, so it rejected the op and the scheduler fell back to CPU, emitting a 'layer assigned to MTL but Lightning Indexer on CPU' warning. Bump key_length to 128 and the DSA head_count to 64 so the fused op runs on the GPU.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : add --help and document -o in test-llama-archs
Add a --help/-h flag to test-llama-archs and list the existing -o/--out option in the usage text, which was previously missing.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : use 64 indexer heads for deepseek4
deepseek4's indexer head count was set to n_head (8), which does not match the fused Lightning Indexer kernel's fixed NH=64, so the fused op fell back to the CPU backend and emitted a device-mismatch warning. Give it the same fixed 64 as the other indexer archs by dropping it from the n_head ternary (only minimax-m3 keeps n_head, since it does not use the fused Lightning Indexer op).
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* tests : fix dsv4 save-load n_stream mismatch
The dsv4 KV cache keeps per-sequence KV/state streams even in unified mode, so its n_stream equals n_seq_max. The test saved the state in the baseline with n_seq_max=1 but loaded it in the seq-copy tests with n_seq_max=2, so state_read threw an n_stream mismatch. Use n_seq_max=2 in the baseline and state-load tests so the save and load agree.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* context : relax on-device seq-copy chunk alignment
The on-device state seq copy (llama_state_seq_set_data with LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) copied the write-side cpy tensors to the read-side targets 1:1 by index, requiring the writer and reader to emit the same number of chunks in the same order with the same per-chunk sizes. state_write_data chunks per cell-range while state_read_data chunks contiguous-or-per-cell, so the counts diverged for non-contiguous sources (dsv4, SWA) and the copy aborted with "memory buffer mismatch".
All state writers and readers enumerate the same logical data in the same order, differing only in chunking. Copy the flat write-side data into the read-side targets with a byte cursor that walks both tensor lists across their boundaries, so the chunking no longer needs to match. Keep the total-size guard; drop the n_tensors equality check.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* model : fix dangling hparams ref in minimax-01 LA graph input
llm_graph_input_la stored const llama_hparams & hparams, bound to the llm_graph_params temporary in llama_context::process_ubatch. The input object outlives that temporary (it is kept in llm_graph_result::inputs for graph reuse), so set_input() read destroyed stack memory on every graph reuse - test-save-load-state crashed for minimax-01 when the stack region was overwritten (n_layer_all read as 0, abort in llama_hparams::n_head). Store a copy like every other graph input class.
Assisted-by: pi:llama.cpp/Qwen3.8-27B
* context : handle "worst case" graph and add TODO
Add HVX-accelerated implementations for GGML_OP_LOG and
GGML_UNARY_OP_ABS on the HTP backend.
- Register HTP_OP_UNARY_ABS and HTP_OP_UNARY_LOG in op_remap_to_htp()
- Add ABS and LOG to ggml_backend_hexagon_device_supports_op()
- Implement hvx_abs_f32_aa() in hvx-arith.h using hvx_vec_abs_f32()
- Implement hvx_log_f32_aa() in hvx-log.h using hvx_vec_log_f32()
- Add abs_f32() and log_f32() row-wise dispatch in unary-ops.c
- Define tiled and non-tiled task functions via DEFINE_UNARY_TASK and
DEFINE_UNARY_TILED_TASK macros
- Route HTP_OP_UNARY_ABS and HTP_OP_UNARY_LOG through execute_op()
in main.c
* Add ctx-per-slot argument for unifid KV cache
* Swap out ctx fractions for ctx pool slots
* Formatting cleanup
* Remove ctx-pool-slots, make ctx-per-slot an int
* refactor it
---------
Co-authored-by: Xuan Son Nguyen <[email protected]>
* gguf: add qwen4exp (Qwen3.8-Flash-Next) arch and converter
Adds the GGUF-side plumbing for HF model_type qwen4_exp:
- MODEL_ARCH.QWEN4EXP plus tensors for the low-rank hyper-connection
variant (hc_*_norm/down/up/inject) and the PLE n-gram hash embeddings.
The DeepSeek-V4 hc_*_fn/base/scale tensors are a different
parameterisation, so these are separate entries rather than reuse.
- Reuses the existing indexer, per_layer_token_embd, SSM and
compress_ratios keys unchanged.
- conversion/qwen4exp.py inherits the Qwen3.5 linear-attention V-head
reorder and interleaved mrope, concatenates the 128 PLE embedding
shards, and splits index_qk_proj into separate indexer q/k tensors.
The PLE hash multipliers reach ~2.4e13. prepare_tensors() casts every
non-float dtype to float32 before modify_tensors() runs, and GGUF array
writes infer INT32 from Python ints, so both paths are bypassed: the
constants are read from the pre-cast lazy tensors and written as
explicit UINT64 arrays.
Additive only; no existing arch changes behaviour.
* llama: load qwen4exp (Qwen3.8-Flash-Next) hparams and tensors
Adds LLM_ARCH_QWEN4EXP with its hparams and tensor loading. The graph
comes in the next commit; this makes the model load and report correct
metadata.
- hyper-connections set n_embd_out_impl = hc_count * n_embd, so the
residual stream is 4x wide and there is no output_norm: the final
mixer's hc_norm is the last norm in the model.
- registered as hybrid and given the same recurrent/attention memory
filters as Qwen3-Next and Qwen3.5.
- reuses the existing indexer, per_layer_token_embd, SSM and
compress_ratios keys as-is.
- the PLE table row count is read back from the file rather than
recomputing the vocab padding rule.
llama-model-loader gains UINT64 array support. That branch previously
threw, so no existing caller changes behaviour; it is needed because the
PLE hash multipliers do not fit in int32.
* qwen4exp: shorten comments
* llama: qwen4exp text graph with hyper-connections, GDN and MoE
Implements the decode graph for Qwen3.8-Flash-Next: the hyper-connection
residual stream, gated delta net layers, the MoE block with its gated shared
expert, and dense full attention. The QSA indexer and the PLE n-gram embedding
are not wired up yet and land in later commits.
Hyper-connections are implemented here rather than shared with deepseek4.cpp.
The two formulations agree on the [n_embd, hc, n_tokens] layout and little
else: DeepSeek-V4 mixes with a full-rank projection and Sinkhorn-normalises
it, whereas this model uses a low-rank down/silu/up sigmoid gate and collapses
by a plain mean. Only the ~10 line stream mean is genuinely common, so sharing
would mean touching DSV4's hot path and its three fused CUDA ops to reuse very
little. What is reused is the substantive part: the LLM_KV_HYPER_CONNECTION_*
keys, the n_embd_out_impl wide-residual support already in the loader, and the
layout convention.
Also allows a checkpoint to carry no PLE layers at all, which makes it
possible to bring the graph up and validate it in stages.
Validated against vLLM, the only working reference implementation. On a
scaled-down model with an init scale large enough to give non-uniform logits,
agreement with vLLM sits at the numerical noise floor: llama.cpp f32 against
its own bf16 gives 84.3% top-1 agreement over 255 positions, and this graph
against vLLM gives 85.1%. The comparison was calibrated by seeding three
deliberate bugs (silu instead of sigmoid on the delta net gate, dropping the
1/hc scale in the mix, dropping the 2x in the combine); each drops top-1 to
between 0% and 11%, an order of magnitude below the floor.
* llama: qwen4exp PLE n-gram hash embedding
Adds the per-layer embedding: a custom I32 graph input hashes each token with
its ngram_size-1 predecessors host-side and the result is a plain row gather
over the shared table, the same shape gemma3n's per-layer embedding uses. The
hash has to run on the host because the splitmix64-derived multipliers reach
2^45, so the products need 64-bit integers and an xor, neither of which ggml
has.
Predecessors that fall outside the ubatch come from a small per-sequence
history on the model, mirroring the per-request ngram_context the reference
carries. It is only trusted when contiguous with the incoming position, so a
fresh prompt or a rewound cache falls back to EOS padding rather than hashing
against stale tokens.
The depthwise conv is written out as a sum of shifted, per-channel-scaled
copies rather than through ggml_conv_1d_dw, which carries a correctness
warning upstream.
Verified two ways. The row indices match a transcription of the reference's
tensor formulation exactly, 1024 of 1024 rows, including sequences with EOS
tokens sprinkled through them to exercise the segment reset. Separately, with
PLE placed on layer 0 so its input is just the token embedding, ple_embd and
ple_gated_value match a PyTorch computation from the same checkpoint to every
printed digit.
End to end over 1023 scored positions the port sits the same distance from
vLLM with PLE as without it, 6.3 points of top-1 against 6.0, so PLE costs no
accuracy relative to the rest of the model. That common offset is vLLM's bf16
activations, which cannot be removed: its QSA kernel refuses float32.
Two bugs found along the way, both caught by the row-index check. The history
was read and updated in the same pass, so a token early in a ubatch could pick
up an earlier token of that same ubatch as prior context; it is now snapshotted
first. And an EOS token was cutting its own context, where the reference takes
the last EOS strictly before the position, so a boundary only hides tokens from
the positions after it.
Known gap: the conv carries no state across ubatches, so it is exact only for a
prefill that starts at position 0. Chunked prefill and decode need the conv
state wired into the recurrent memory, and the conv branch itself is still
numerically unverified because the fixture zeroes its weights.
* llama: carry the qwen4exp PLE conv state across ubatches
The PLE depthwise conv was zero-padding on the left, which is only right for a
prefill that starts at position 0. Decode and chunked prefill saw a truncated
history for the first (kernel-1)*ngram_size positions of every ubatch.
The PLE module sits on a layer that is also a delta-net layer, so both need a
conv history in the same recurrent row. Rather than plumb a per-layer state
size through build_rs and build_conv_state, the row is widened once and each
convolution addresses its own slice through a local helper. n_embd_r() gains
the extra span, which is zero for every other architecture because it is
derived from ple_n_heads.
Verified by feeding the same 1024 token sequence in chunks instead of one
shot: at 64 tokens per decode the logits are bit-identical to the single-shot
run, 1023 of 1023 top-1 and a maximum logprob deviation of exactly zero. At
one token per decode they differ slightly, but the no-PLE model differs more
under the same test (94.6% against 97.1%), so that is the usual gemv-versus-
gemm accumulation difference and not the state.
The conv branch is also no longer unverified. With non-zero conv weights the
port sits 6.3 points of top-1 below the numerical floor, the same distance as
with the weights zeroed and as the model with no PLE at all, so the branch
adds no error of its own.
test-llama-archs passes every existing architecture at 0.00e+00, including the
delta-net models that share this code path.
* llama: fix the qwen4exp PLE conv state and unblock test-llama-archs
build_rs writes into the state tensor in place, zeroing one row and copying the
carried-over states, so calling it twice for the same layer let the second call
clobber the first write-back. The PLE layer is also a delta-net layer, so that
is exactly what happened: both convolutions gathered the same row. They now
share a single gather per layer.
The earlier claim that the conv state was carried correctly was tested on a
fixture whose conv weights are zero, where the branch contributes nothing and
chunking matches trivially. Re-running with non-zero conv weights showed the
divergence, growing with the number of ubatch boundaries: 97.1% top-1 at one
boundary down to 90.2% at seven. With the shared gather it is bit-identical to
the single-shot run at every chunk size tried, 512, 128 and 64, with a maximum
logprob deviation of exactly zero over 1023 positions. The delta-net-only model
stays bit-identical too, so nothing regressed there.
Also derive the delta-net conv channel count the way load_arch_tensors sizes
wqkv instead of from ssm_d_inner. The two agree for this model, but n_embd_r()
only bounds the row and the convolution has to match the tensor feeding it.
test-llama-archs previously aborted on this architecture and took every later
architecture with it. qwen4exp is marked MoE-only, given the hyper-connection
keys and an ssm_d_inner consistent with its tensor derivation, and skipped for
now: the hyper-connection keys written by get_gguf_ctx are not reaching the
synthesised file, which needs a separate look. The suite completes again, 124
architectures at 0.00e+00.
* llama: optional indexer key cache in llama_memory_hybrid
Groundwork for qwen4exp's QSA sparse attention. Its indexer needs a per-token
key history for the full-attention layers, but a hybrid model cannot use
llama_kv_cache_dsa: that class derives from llama_memory_i rather than
llama_kv_cache, and llama_memory_hybrid constructs its attention cache
directly. No existing architecture pairs recurrent state with a sparse
indexer, so there was nothing to reuse wholesale.
llama_memory_hybrid therefore gains a third, optional cache, shaped the same
way llama_kv_cache_dsa shapes its lightning-indexer cache: a copy of hparams
with n_head_kv forced to 1 and n_embd_head_k_full set to indexer_head_size.
It is built only when a filter_idx callback is passed, which defaults to
nullptr, so every existing architecture gets exactly what it got before. The
per-sequence operations and the batch preparation forward to it under a null
check, matching how the DSA cache prepares its two caches over the same
ubatches.
test-llama-archs passes all 124 architectures at 0.00e+00, including the 12 in
the hybrid family that share this code. The qwen4exp fixtures are unchanged:
same logits against vLLM, and chunked evaluation still bit-identical to
single-shot.
* llama: QSA sparse attention for qwen4exp
The full-attention layers of this model do not attend to everything. An
indexer scores one mean-pooled key per block of compress_ratio tokens and
keeps a budget of the best blocks, plus the tail of tokens that do not yet
form a complete block. Below indexer_top_k + compress_ratio - 1 cached
tokens every block fits in the budget, so the result is exactly dense.
What is reused rather than rebuilt:
- the mask machinery. build_attn's DSA overload already turns a list of
token indices into a KQ mask via ggml_set_rows, so that block is lifted
out verbatim into build_attn_mask_top_k and shared with a new overload
on llm_graph_input_attn_kv. DSA's node sequence is unchanged; the new
overload exists because llama_kv_cache_dsa assumes MLA and cannot be
dropped into a hybrid model.
- the indexer key cache, which is the optional third cache added to
llama_memory_hybrid in the previous commit. It holds raw keys, because
pooling happens before the norm and the rotation.
The graph expands block scores rather than block indices: giving every
token of a block its block's score needs only a gather, where expanding
indices would need an integer multiply-add that ggml has no op for. Since
the budget is a whole number of blocks and a block's members tie exactly,
the cut still lands on a block boundary.
Everything that depends on cache layout is computed host-side in
set_input_qsa. Blocks are cuts of the position line rather than of the cell
array, so nothing assumes the cache is contiguous.
Measured on the tiny fixture against vLLM, comparing the selected token
indices directly rather than the logits:
below the budget selection identical, and 1024-token logits are
bit-identical to the pre-QSA dense path
above the budget mean jaccard 0.975
The direct index comparison is what made this correct. The reference
rectifies each head's dot product before summing over heads, which an
earlier reading of it had missed; on logits alone the resulting port looked
fine, because on a randomly initialised fixture the known-correct dense
path already disagrees with vLLM by more than the bug did. Comparing the
indices showed 0.794, and fixing the ReLU moved it to 0.975.
* llama: give the qwen4exp indexer cache the attention cache's slots
The indexer cache found its own slots, independently of the attention
cache. Both are the same size and see the same ubatches, so in a
straight-through prefill they agree, which is why every fixture and every
single-shot parity run passed. They drift once the context is being
rewritten between turns, and then the QSA top-k indices, which are applied
against the attention mask, point at the wrong cells.
The seven-turn chat test caught it on the third turn: llama-server aborted
on the assertion that the two caches report the same n_kv.
The cache is a side buffer addressed by the attention cache's cells, so it
now takes that cache's slot layout instead of computing one. Applying that
layout also marks its cells identically, so the two agree cell for cell by
construction rather than by coincidence, and the assertion can no longer
fire.
Inert where the caches already agreed: test-llama-archs green at 126 archs
and 0.00e+00, and the 4096-token tiny fixture is unchanged at max logit
delta 0.0.
* tests: record what the qwen4exp arch-test skip actually observes
The old note guessed that the hyper-connection keys never reach the file.
They do: dumping the gguf_context handed to llama_model_init_from_user
shows both among its 67 KVs, and the loader still reports one missing.
* tests: cover qwen4exp in test-llama-archs
The arch was skipped with a note guessing that the hyper-connection keys
never reached the synthesised file. They did. The suite builds a model, then
saves and reloads it, and llama_model_saver did not re-emit those keys, so
the failure was in the roundtrip leg rather than the first load. Three gaps,
all in shared code and all additive:
- add_kv_from_model wrote no hyper-connection, compress-ratio or PLE keys.
The PLE group only means anything whole, so it is written or omitted
together; the rest follow the file's existing style of writing every key
unconditionally, since an architecture that does not read one is
unaffected by a zero.
- the saver had no uint64 path at all, which the PLE hash constants need.
- add_tensors_from_model enumerates model-level tensors by hand and was
missing per_layer_tok_embd and the three final-mixer tensors.
Two smaller fixes on the qwen4exp side, both found by running the test:
- build_qsa_top_k divided by the compression ratio before asserting it was
non-zero, so a file without the key crashed instead of reporting.
- a layer with no compression ratio now falls back to dense attention,
which is what the model computes below the budget anyway. The test then
has to write a ratio to reach QSA at all, and an indexer key length no
narrower than n_rot, since the indexer ropes with the main attention's
rotary width.
Full suite: 126 archs, qwen4exp at 0.00e+00 with roundtrip OK. The tiny
fixture is unchanged, max logit delta 0.0 against the pre-QSA dense run.
* convert: stream the qwen4exp PLE table instead of concatenating it
The n-gram table arrives as 128 shards that were held in a dict and then
torch.cat-ed, so the peak was the shards plus the concatenation: around
300 GB of RSS on the real checkpoint, which rules out machines that could
otherwise convert this model.
Each shard is now written straight into a memory-mapped file at its final
row offset and dropped, so the resident set is one shard and the rest is
the page cache's problem. The temporary file sits beside the output and is
removed once the write finishes, including on failure.
Shards other than the last must be uniform for direct placement, which is
asserted rather than assumed, and a shard arriving before the stride is
known is held instead of misplaced.
Verified on the tiny fixture: the resulting GGUF is byte-identical to the
one the concatenating path produced (md5 2d274efac91ad1e9a6007efb0687e597).
* quantize: fall back to F16 for 32-block types with an odd ncols
tensor_type_fallback demotes a tensor whose ncols is not a multiple of the
target's block size, but its switch only enumerates the 256-block types. A
target that is already a 32-block type (iq4_nl, q4_0, q5_0, q8_0, ...) falls
into default: and throws, even though the function already knows how to answer
that case: the ncols check right below the switch resolves an unrepresentable
shape to F16.
Route those types into that check instead of throwing. Only paths that abort
today change, so no quantization that currently succeeds is affected.
Found on a 4-wide depthwise conv kernel. llama-quantize reported nothing but
"failed to quantize model from ...", with no tensor name and no exception text,
which made a quant recipe that had simply not pinned the tensor look like a
corrupt model. It now names the tensor and continues.
* quantize: let --tensor-type name per_layer_token_embd
per_layer_token_embd shares the TOKEN_EMBD category with token_embd.weight, so
--token-embedding-type is returned for it before any --tensor-type pattern is
consulted, and there is no way to give it a tier of its own.
That grouping is fine as a default and stays the default. It is a poor fit for
the size, though: on qwen4exp the table is 97.7 GiB of a 337.6 GiB BF16 file and
about 46% of a 4-bit one, roughly eighty times token_embd.weight, and it is
read by ggml_get_rows rather than a matmul so no imatrix ever covers it.
Allow an explicit --tensor-type pattern to name it, and only it. Nothing
changes unless such a pattern is passed, and token_embd.weight keeps the old
precedence in either case.
Measured on Qwen3.8-Flash-Next, Q4_K_M with an imatrix: the table lands at q8_0
(51.9 GiB, 113.5 GiB total) by following --token-embedding-type, and pinning it
q4_1 gives 30.5 GiB for 92.1 GiB total, 19% off the file.
* quantize: size the output buffer exactly instead of nelements * 4
The per-tensor output buffer was sized `nelements * 4`, described as an upper
bound. It is a very loose one: the output is at most 2 bytes per element
(f16/bf16) and usually well under 1.1 (q8_0 and below), so between 2x and 4x of
it is never touched. The exact size is already known here, since it is what the
quantization loop writes, what new_size sums to, and what the GGUF metadata is
asserted against a few lines later.
On a model whose largest tensor is a few GB none of this matters. On
Qwen3.8-Flash-Next it does: per_layer_token_embd is 51.2 G elements, so the
buffer was 205 GB where 54 GB is needed at q8_0 and 32 GB at q4_1.
Measured on that model, VmHWM of a live llama-quantize was 485 GB per process.
Three of them fit in 2 TB and five did not, which is what an OOM-killed quant
ladder looks like. This removes about 150 GB of that.
Byte-identical output, verified against the same binary built at the parent
commit: q4_K, q8_0, q5_K, q6_K and IQ4_XS, over BF16 and F32 sources, with and
without a PLE table present. Six cases, six matching md5s.
* qwen4exp: hash the image placeholder for multimodal batches
The PLE row indices are computed host-side from ubatch->token, and set_input
returned early when that was null. A multimodal ubatch is exactly that case:
the mtmd layer consumes the image placeholder ids and hands llama_decode
embeddings instead. The early return left the I32 index tensor uninitialised,
so ggml_get_rows indexed a 320 M row table with whatever the buffer happened to
contain, and aborted:
GGML_ASSERT(i01 >= 0 && i01 < ne01) failed
ggml_compute_forward_get_rows
mtmd_helper_decode_image_chunk -> llama_decode
Every image request crashed. Nothing caught it because the vision work had only
ever been verified by converting an mmproj, never by running one.
The reference computes the hash over input_ids, where those positions still
hold the image placeholder, so carry that id through as qwen4exp.ple.image_token_id
and hash it. The key is optional: a file converted before it existed falls back
to the PLE EOS token, which is defined and treats the image as a segment
boundary rather than crashing.
Verified end to end with llama-mtmd-cli, a Q4_K_M base and the F16 mmproj, on a
generated image with known content. The model names the red circle, the blue
square, the inverted green triangle and reads "UNSLOTH 42", each with the right
position.
* qwen4exp: support a non-unified KV cache in QSA
set_input_qsa asserted n_stream == 1, so llama-server could not serve this
model with more than one slot unless -kvu was passed. With a non-unified
cache each sequence owns its own cells, and a cell index means a different
token in each stream, so a single shared mapping is wrong.
- cell_blk, blk_cells and bias gain a stream dimension. At n_stream == 1
these collapse to the shapes they had, so the unified path is unchanged.
- Scoring is now batched over streams. ggml_mul_mat matches ne[2] on both
operands, so stream s's queries only ever meet stream s's blocks; without
this sequences would score against each other's context.
- set_input_qsa loops per stream and resolves cells through
v_cells[seq_to_stream[seq_id]], following set_input_kq_mask_impl, instead
of hardcoding v_cells[0].
- llama_kv_cache_context::get_n_stream() is added, mirroring the ns that
get_k and get_v already derive from the slot info.
build_attn_mask_top_k needed no change: it already expects
[n_top_k, n_batch, 1, n_stream], so the top-k result is reshaped to meet it.
set_input_qsa has exactly one caller, so the blast radius is qwen4exp only.
Validation, UD-Q4_K_XL on one B200:
- unified cache unchanged within noise: 1802.9/68.85 -> 1807.2/69.11 t/s at
batch 1, 2262.5/192.43 -> 2270.1/193.75 at batch 4.
- non-unified now runs at npl 1, 4, 16 where it previously aborted, and is
22% faster than the -kvu workaround at batch 16 (1205 vs 984 t/s total),
since per-stream cells avoid the cross-sequence masking a unified cache
pays for.
- no cross-stream contamination: four concurrent sequences each carrying a
distinct secret all recall their own and no other, on both cache modes.
- test-llama-archs green on qwen4exp, deepseek2, gemma3n, qwen3next, llama.
Note on testing: comparing concurrent output against solo output exactly is
not a valid check. It failed 0/4 with no bug present, and the unified-cache
control failed the same way, because batch composition changes the
floating-point reduction order and near-tied tokens flip. The contamination
test above is what the exit code gates on.
* llama: keep the qwen4exp top-k attention mask arch-local
The QSA graph needed a build_attn that attends only to the cells named by a
top_k tensor, and the first version got it by adding a llm_graph_input_attn_kv
overload to llm_graph_context and factoring the mask construction out of the
existing MLA sparse path into a shared build_attn_mask_top_k.
That put a new arch on the shared attention path and made the deepseek32 and
glm-dsa attention build depend on a helper introduced for qwen4exp. Build the
mask in src/models/qwen4exp.cpp instead and leave llama-graph.{h,cpp} exactly as
they were: the MLA path keeps its own copy of the same node sequence.
The nodes emitted are unchanged, so this is bit-identical.
* llama: hold the qwen4exp indexer cache in a new llama_memory_hybrid_idx
The indexer key cache was added by extending llama_memory_hybrid with an
optional third cache, and the host-side cell/block mapping that drives QSA was
added as set_input_qsa on llama_kv_cache. Both are shared classes that every
hybrid and every attention model goes through.
Move both into a new memory type, llama_memory_hybrid_idx, following
llama_kv_cache_msa: the indexer cache and the pos<->cell translation live with
the sparse-attention memory rather than in the classes that serve every other
architecture. llama-kv-cache.{h,cpp} and llama-memory-hybrid.{h,cpp} are
restored to their unmodified state.
init_batch is repeated from llama_memory_hybrid because the indexer cache has to
be handed the attention cache's slot infos, and those are not reachable through
the context the base returns. Allocating them separately lets the two caches
drift, which is what pointed QSA's top-k at the wrong cells before.
The context derives from llama_memory_hybrid_context so build_inp_mem_hybrid
keeps working unchanged, and get_n_stream is computed from the slot infos
exactly as llama_kv_cache_context did.
Behaviour is unchanged: logits over an 8192-token sequence are bit-identical to
the previous implementation, sparse and dense alike.
* llama: save and restore the qwen4exp indexer KV cache
llama_memory_hybrid_idx forwarded clear, seq_rm, seq_cp, seq_keep, seq_add and
seq_div to the indexer cache but not state_write / state_read, so a saved
session dropped the indexer keys and a restored one selected QSA top-k against
an empty cache. The effect is invisible until the context passes
indexer_top_k + compress_ratio - 1 cells, because QSA is exactly dense below
that and the indexer contents cannot change the result.
The indexer section is written last rather than next to the attention cache it
mirrors. As a suffix, a reader that does not expect it stops early and the
trailing bytes are caught by the size check in state_load_file; placed between
the attention and recurrent sections it would instead be parsed as recurrent
state, which can succeed and restore silent garbage. It follows the same
LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY gate as the attention cache, since a partial
checkpoint deliberately skips the token-level attention caches.
The indexer restores its own cells instead of taking the attention cache's
restored slots. The two caches share size, padding and every sequence
operation, and init_batch hands the indexer the attention cache's slot infos,
so both state_read_meta calls run find_slot over identical occupancy and land
on identical cells.
The overrides live on llama_memory_hybrid_idx, the only memory type that owns
an indexer cache, so llama_memory_hybrid and every architecture that uses it
write and read exactly the bytes they did before.
The session and sequence state versions are bumped because the qwen4exp state
layout changed. The session path already rejects a short read via its size
check, but llama_state_seq_load_file accepts one silently, so only the version
check stops a pre-fix blob from being half-restored by a fixed build.
(cherry picked from commit 2721542354f8e158c3217625f4e2e7b83e51e3fe)
* llama: make the qwen4exp PLE n-gram history per context and serialise it
The PLE hash of a token mixes in the ple_ngram_size - 1 tokens before it, which
a decode ubatch does not carry, so they were remembered in a map on
llama_model_qwen4exp. That is the wrong owner twice over.
A llama_model is shared by every context that loads it, and the map was keyed
only by llama_seq_id, so two contexts running the same sequence id - two server
instances on one model, or a draft/target pair - overwrote each other's window.
The next_pos guard turned that into EOS padding instead of a crash, so it
degraded quality silently.
The map was also in no state blob: grep found ple_hist in neither
llama-kv-cache.cpp nor llama-memory-*.cpp nor llama-context.cpp. A restored
context therefore failed the next_pos check on its first ubatch and hashed the
first tokens after the restore against EOS padding. This is why a session blob
round-tripped byte for byte while the restored context computed different
logits: the state was never in the bytes.
It moves to llama_memory_hybrid_idx, which is per context, is the memory type
qwen4exp always builds, and already does the per-sequence bookkeeping this
needs. Every sequence operation now carries the window with it:
seq_rm a rewind (p1 < 0) truncates the window to the surviving prefix and
moves next_pos to p0, so a rollback keeps exact context; a hole
punched in the middle leaves the window non-contiguous, so it is
dropped
seq_cp the destination inherits the source's window, truncated to the
copied position range - a copied sequence continues with the same
n-grams the source would have used
seq_keep every other sequence's window is dropped, like its cells
seq_add a shift that moves the whole window keeps it and moves next_pos with
it, which is the context-shift case; one that cuts through it drops
it
seq_div positions stop being consecutive, so an overlapping window is
dropped
clear everything is dropped
Dropping means next_pos = -1, which set_input turns into full EOS padding: the
same thing a fresh sequence gets, and the same thing this code did before it
followed the sequence operations at all, so no case is worse than before.
The state payload is a self-delimiting list, u32 count then per entry
{ i32 seq_id, i32 next_pos, u32 n_toks, i32 toks[n_toks] }, so a whole-context
save and a single-sequence save share one format and a single-sequence restore
can retarget the window at its destination seq_id. It is written after the
indexer section, last, for the same reason that one is: as a pure suffix an
older reader stops early instead of parsing these bytes as something else.
Unlike the indexer section it is not under LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY.
The window is recurrent state - it is the input the PLE convolution's own
recurrent state is derived from - and the recurrent cache beside it is written
for partial checkpoints too. Gating it would leave the server's speculative
decoding checkpoints restoring the conv state without the window that produced
it.
No further version bump: LLAMA_SESSION_VERSION 10 and LLAMA_STATE_SEQ_VERSION 3
were introduced for the indexer section in the same unreleased series, and both
changes are qwen4exp-only additions to the same blob layout.
Also fixes the padding of a short window. set_input pads a window shorter than
ngram_size - 1 up to that length, but prev() indexes the snapshot with the most
recent token last, and resize() pads at the back, so the filler EOS landed where
the immediately preceding token belongs. It now pads at the front. A window is
short at a sequence start after a one-token prefill, and after a seq_rm rewind,
which the new bookkeeping makes common.
Every architecture other than qwen4exp builds llama_memory_hybrid rather than
llama_memory_hybrid_idx, has no PLE table and never asks for a history, so
nothing about its graph, its sequence operations or its state bytes changes.
(cherry picked from commit de170364c052c68fcf63285cc0028095edb9f23c)
* qwen4exp: tidy comments and simplify image token read
Rewrite the comments this series adds to the AGENTS.md rules: one or two lines,
no prose hard-wrapped mid-sentence, no narrative or history, and no comment that
only restates the code. Net 146 fewer comment lines, no code change.
Correct the PLE image comment: mtmd does not consume the placeholder ids. An
image is decoded as an embeddings-only batch, so ubatch->token is null and the
per-position ids never exist here. gemma3n and gemma4 hit the same case and
stand in row 0 of per_layer_token_embd; qwen4exp stands in the configured image
token id instead.
Read image_token_id straight from self.hparams in the converter. base.py merges
text_config into the root of hparams, and the key sits at the root of
config.json, so the config.json re-read was redundant.
(cherry picked from commit 205840c12169057da3e8d2f65ec4ceec3e18b980)
* qwen4exp: support a quantized KV cache in the QSA attention path
(cherry picked from commit 4c30574f81dc1115d08078c47b6cf8c789c0a842)
* llama: give qwen4exp a large-graph node budget
(cherry picked from commit 37c8c194e6a30e4c46ac29bee3fb264f091596ef)
* qwen4exp: drop an unused variable that breaks -Werror builds
(cherry picked from commit 528d032b51fa3cf935ed3ef6e0fb1c7401df53b5)
* quantize: dequantize and quantize large tensors in row bands
f32_conv_buf held the whole dequantized tensor, which is 204.8 GB for
per_layer_token_embd alone and dies with std::bad_alloc long before the
work buffer is reached. Dequantize and quantize in bands of whole rows
instead, capping the f32 staging at 1 GiB per band.
Rows are independent and the imatrix is indexed by column, so band
boundaries cannot change any output byte. Bands nest inside the existing
per-expert loop so each expert slice keeps its own imatrix, and a band is
kept to at least one quantization chunk per worker thread so the existing
multithreading still has work. F32 sources still stage nothing and are
banded by pointer arithmetic into the tensor.
llama_tensor_dequantize_impl now takes a first element offset; the single
caller is updated.
(cherry picked from commit 658c22549613555dbce57a772be4de8509eba3ee)
* llama: segment the qwen4exp fused QKV for tensor split
qwen4exp was missing from the gated delta net branch of get_split_segments,
so its attn_qkv.weight, shaped {n_embd, 2*key_dim + value_dim}, fell through
to the generic fused QKV rule and tripped
GGML_ASSERT(tensor->ne[axis] == n_embd + 2*n_embd_gqa) while loading with
--split-mode tensor. --split-mode layer was unaffected.
qwen4exp broadcasts K to the V heads by tiling, k_conv is grown with a plain
ggml_repeat_4d over the head axis so that v head j pairs with k head
j % n_k_heads. That is the Qwen 3.5 pattern, not the repeat interleave that
Qwen 3 Next builds explicitly, so qwen4exp takes the else branch and its V is
segmented on the scale of K.
Reported by benklop.
(cherry picked from commit 353d753f595dc81634ae6130188b31f06018f5ae)
* llama: fix the qwen4exp PLE history seq_rm(-1) iterator invalidation and the fatal-warning build
ple_hist_rm recursed over ple_hist with a range-based for and the recursive call
erases the entry it is iterating when the whole sequence is removed (p0 <= 0,
p1 < 0), so the loop then increments an invalidated iterator. It is unreachable
today only because llama_memory_recurrent::seq_rm rejects seq_id < 0 before
llama_memory_hybrid_idx::seq_rm reaches the history, which is a guard in another
class. Advance past the entry before recursing.
Two smaller things in the same area:
- the n_toks sanity bound in ple_hist_state_read was the literal 64, which is
the value of LLAMA_MAX_PLE_HEADS, not of the quantity being checked. The
window is at most ple_ngram_size - 1 tokens, so the bound is
LLAMA_MAX_PLE_NGRAM - 1, eight times tighter.
- build_conv_state_at left mem_size unused, so -DLLAMA_FATAL_WARNINGS=ON does
not compile. Predates this series; drop the line.
(cherry picked from commit 6eba44a89d5f328eb4859b844e1d28fb564cbe3e)
* qwen4exp: include llama-impl.h explicitly for llama_mul_mat_hadamard
(cherry picked from commit b634fd4d250d181ef82bf78bd00c1ae3b96a7af6)
* convert: fix the qwen4exp lint and type-check failures
flake8 flagged an unused MmprojModel import, and ty flagged seven errors in
the PLE streaming path: eos_token_id can be absent, and _ple_map, _ple_path,
_ple_row_dim and _ple_rows_per_shard are all Optional at the declaration but
were dereferenced without narrowing.
The map is opened and the stride fixed before the first shard is written, and
_finish_ple_table only runs once every shard has landed, so the invariants
hold. Assert them so the checker can see it. A missing eos_token_id now raises
with the reason instead of a TypeError from int(None).
* llama: give the qwen4exp indexer cache its own tensor names
The indexer KV cache and the attention KV cache both named their tensors
cache_k_l%d, so the Meta backend matched the indexer cache against the
attention split pattern and aborted in handle_set_rows. Tag the names
instead, and mirror the indexer cache: it has one key head and its
projections are mirrored.
(cherry picked from commit a1cdc8181134659766763a17762545a1f0e5db7b)
* qwen4exp: double the Q split granularity for tensor parallelism
qwen4exp fuses the attention gate into attn_q.weight the same way qwen3next
and qwen 3.5 do, so a device boundary must fall on a whole q+gate pair or the
Q heads stop lining up with the K/V heads and attn_output rows.
(cherry picked from commit 6c9a592f0a425a459ab6efae3b897cf68460e244)
* qwen4exp: keep the indexer cache in step across server slots
The QSA indexer keeps a side cache addressed by the cells of the attention
cache, so cell j has to hold the same token in both: the top-k indices it
produces are applied to the attention KQ mask. init_batch already hands the
indexer the attention cache's slot layout rather than letting it look for its
own, but the restore path did not. state_read called llama_kv_cache::state_read
on the two caches in turn and each ran its own find_slot over its own occupancy.
That agrees only for as long as nothing has already pushed the two caches apart,
which is the property a restore is supposed to re-establish rather than one it
can lean on.
The failure path was the worse half, and it is reachable from the public API
with nothing more than a short buffer. Truncating a good blob at 35 offsets and
feeding it to llama_state_seq_set_data left the two caches disagreeing at 5 of
them, and every one of 23 truncations of a whole-context blob did. Four of those
five land inside the attention section, so the attention cache drops the
sequence and the indexer keeps it; only the cut that lands in the indexer
section gives the opposite direction. llama_kv_cache::state_read cleans up its
own cache and rethrows, so whichever way it falls, nothing is left to bring the
two back together. The server papers over this by clearing the slot when a
prompt cache load fails; a caller of llama_state_seq_set_data that does not is
left with an indexer addressing cells that no longer mean what it thinks.
llama_kv_cache::state_read_sinfo reports the cells a restore landed in, or takes
a copy of them, and state_read_meta uses a supplied layout in place of find_slot
once it has checked that those cells are free here too. The indexer now adopts
the attention cache's restored layout by construction instead of reproducing it
by coincidence, and a layout that does not fit fails the read rather than being
applied over cells that already drifted. The hybrid restore is wrapped so that
any failure drops the sequence, or for a whole-context restore the context, from
all three caches at once, which is a state they do agree on.
* kv-cache: clear the cache once when restoring a whole context
state_read walks the streams of the cache in turn, and for a whole-context restore
each stream went through state_read_meta, which starts by calling clear(). clear()
resets every stream at once, so each stream after the first threw away the streams
already restored, and the K/V buffers with them. A non-unified cache holds one
stream per sequence, so a context saved with N sequences in it came back with only
the sequence in the last stream that carried any cells - the highest sequence id.
A unified cache has one stream and never showed it.
The cache is now emptied once, before the loop, which is what a whole-context
restore means. A blob whose streams are all empty now empties the cache as well,
where before it left the old contents in place.
* kv-cache: check the mirrored slot layout on a whole-context restore too
state_read_meta only looked at the layout it was given on the single-sequence path.
A whole-context restore lays the cells out from 0 in both caches, so they agree as
long as they restore the same number of cells, but nothing checked that they did: an
indexer section belonging to some other context was read over cells the attention
cache had filled from a different one, which is the state the indexer must never be
left in.
* qwen4exp: give the PLE conv history its own mirrored recurrent row
n_embd_r() reserved n_conv + ple_conv_state() so that one cache_r_l row could
carry both the delta-net conv state and the PLE dilated conv history, but the
QWEN4EXP arm of get_split_segments only described n_conv. Under -sm tensor the
segment sum came up short by ple_conv_state() and llama_memory_recurrent
construction aborted in ggml_backend_meta_alloc_ctx_tensors_from_buft.
Widening the segment list is not the fix. The Meta backend propagates a view's
split descriptor from its parent unchanged, so a view of one sub-range of a
split axis is sized as the whole row on every device; declaring the PLE tail as
a second segment merely moves the abort to "shape mismatch for VIEW" at graph
allocation. The two histories also want opposite policies: the delta-net state
is split by head to match wqkv and ssm_conv1d, while per_layer_tok_embd,
ple_conv1d and ple_norm_conv are all mirrored, so every device computes the
whole dilated conv and needs the whole history. One tensor cannot be both, and
the split state has no per-segment mirroring.
Move the PLE history into its own cache_ple_r_l%d row, mark it MIRRORED, and
return n_embd_r() to n_conv. The row is allocated only on layers where is_ple
holds, so mirroring one 92160-element row per device replaces a 92160-element
tail on all 36 recurrent rows: the recurrent R footprint drops rather than
grows. build_conv_state_at now takes its width from the tensor it was handed
and keys its gather on that tensor, which also drops a cont of a strided view.
* no more ple_hist (use master version)
* llama: give the qwen4exp full memory context its indexer cache
graph_reserve() walks a full memory context, and qwen4exp builds its
sparse attention only when the context exposes an indexer cache. the
full-context constructor left ctx_idx null, so the reserved worst case
was the dense fallback: a smaller graph than the one decode executes.
ggml-alloc then had to grow the compute buffer on the first decode,
past the size reported at load.
with -np 4 -c 32768 -fa on -ctk q8_0 -ctv q8_0 on an IQ1_S qwen4exp,
the reserved CUDA0 buffer was 217.00 MiB against 275.71 MiB actually
used, and CUDA_Host 42.31 MiB against 191.14 MiB. reserving the sparse
graph makes both match exactly, in unified and non-unified cache mode.
Co-authored-by: Pascal <[email protected]>
Assisted-by: Claude
* qwen4exp: shrink the PLE hparams storage
llama_hparams is held by value inside llm_graph_params and every llm_graph_input_*,
and llm_graph_params is a stack local in graph_reserve and process_ubatch, so its
width is paid on every worker thread stack.
is_ple_impl spent 2048 bytes carrying 512 bits. It is the one per-layer flag that is
not moved through the loader's uint32 array templates, so a bitset costs nothing in
call sites and also removes the uninitialized read that non-qwen4exp archs had, since
nothing filled the array for them.
The PLE head offsets and vocab sizes are token-space indices; the gather that consumes
them already truncates to int32, so 64-bit storage was never reachable. The gguf arrays
stay uint64 for file compatibility and are narrowed on load.
sizeof(llama_hparams) 34440 -> 31944, sizeof(llm_graph_params) 34872 -> 32376.
* llama: opt-in random-access mmap advice for host-resident gather tables
qwen4exp keeps per_layer_token_embd on the host: 26.8 GiB at IQ4_NL, read
by ggml_get_rows as 16 gathers of ~90-170 bytes per token, spread across
16 head regions ~20M rows apart. Measured over 4.75M gathers, no two
consecutive gathers land on the same 4 KiB page, so the readahead the
loader asks for buys nothing here and the whole table ends up cached to
serve about 4% of itself.
llama_mmap applies POSIX_FADV_SEQUENTIAL, MAP_POPULATE and a whole-file
POSIX_MADV_WILLNEED unconditionally. Those are right for streaming the
file once into buffers and wrong for whatever stays mapped afterwards.
Under LLAMA_MMAP_RANDOM the eager pull-in is skipped and the mapping is
advised random once every tensor has been read, so the load itself keeps
its sequential readahead. That alone drops the table to 4.4% resident but
serializes one NVMe latency per gather.
The second half is what pays for it: the PLE input already computes every
row index for the ubatch before the graph runs, so the pages those rows
fall on are handed to the kernel in one batch and the reads overlap.
POSIX_MADV_WILLNEED on POSIX, PrefetchVirtualMemory on Windows, which
takes the discontiguous ranges in a single call.
Off by default and off for every other model: the batched prefetch keys
off "this mapping was advised random", which nothing sets unless the user
opts in.
-c 512 --chunks 60, cold, IQ1_S, mean of 3:
default 35.3 s 26.82 GiB resident (100%)
advice only 104.5 s 1.19 GiB resident (4.4%)
advice + prefetch 34.2 s 1.19 GiB resident (4.4%)
PPL 4.2346 +/- 0.07862 in all three. IQ1_S KLD is unchanged in every
field, including Mean KLD 0.396070 +/- 0.001931 and Same top p 77.325%.
* llama: narrow the random-access mmap advice to the gather table
The advice was applied per mapping: every mapping the model kept got
POSIX_MADV_RANDOM plus a whole-file POSIX_FADV_RANDOM, and the eager
pull-in was skipped for every file. On qwen4exp that also hit
token_embd.weight, which sits 0.33 GiB past the PLE table in the same
shard and is read densely, not by sparse gathers. Measured over
-c 512 --chunks 60 on IQ1_S it fell to 8.45% resident, against 100% with
the feature off.
A model now nominates its gather tables (qwen4exp: per_layer_tok_embd)
and only those byte ranges are advised. The range is rounded out to
whole pages, which on this model takes in 832 bytes before and 192
after. token_embd goes back to 86.55% resident and the PLE table still
drops to 4.44%; smaps shows one VM_RAND_READ VMA of exactly the table
instead of one over all 27.16 GiB that stays mapped.
posix_fadvise is dropped from the narrowed path. POSIX_FADV_RANDOM
ignores its offset and length and marks the whole open file, and the
FMODE_RANDOM it sets is only read by page_cache_sync_ra() on the read()
path, which a fault on a MADV_RANDOM vma never reaches. POSIX_FADV_
DONTNEED does take a range, so the drop mode keeps it.
The eager pull-in is now skipped only for the files holding a nominated
table, and re-issued as WILLNEED over the rest of such a file, so other
shards load exactly as before.
prefetch_rows() keys off the tensor being nominated rather than off a
mapping-level flag, so the batched readahead lands only where the advice
did.
-c 512 --chunks 60, cold, IQ1_S, mean of 3, total wall:
default 32.50 s
whole mapping 30.05 s
narrowed 30.35 s
PPL 4.2061 in all three. IQ1_S KLD is bit-identical with the feature on
and off, including Mean KLD 0.396070 +/- 0.001931 and Same top p
77.325%. tg128 73.65 +/- 0.33 narrowed against 73.49 +/- 0.34 whole.
Assisted-by: Claude
* llama: fold the random-access prefetch into its own feature flag
LLAMA_MMAP_RANDOM_PREFETCH existed to measure the two halves of the feature
apart, and the measurement is done: on a cold cache over the same wikitext
run, MADV_RANDOM without the batched readahead takes 94.4 s against 36.7 s
for an untouched mapping, while the pair together take 34.1 s. Suppressing
the kernel's readahead only pays if we replace it, so the split let a user
select a 2.6x regression through a documented switch.
Keep the accessor, since the call site reads better than a mode comparison,
but derive it from the mode alone.
* FACP (Fewer Acronym Classes Please)
* qwen4exp: bias the QSA selection per block, not per cell
The QSA bias is a graph input, so it is pinned on the host and uploaded every
decode, and at -c 32768 -np 4 its twelve copies were 768 of the 815 MiB of
reserved host compute buffer.
Only one half of it needs a cell: whether the cell sits in the always-visible
tail, and whether its block was pooled. Both are properties of the block. The
other half - empty, other sequence, or in the future - is the plain visible/not
test the attention mask already carries over the same cells, so add that mask
instead of repeating it. The bias then holds one value per block.
A block sits wholly inside or wholly outside the tail because the tail starts on
a block boundary, so one value per block is exact. Cells no block covers keep
their -inf from the mask.
The mask is F16 and the bias F32, and a mixed ggml_add reinterprets the F16
buffer as float rather than converting it, so the cast is required.
reserved host compute buffer at -c 32768 -np 4:
--kv-unified 814.86 -> 238.86 MiB, CUDA0 721.07 -> 421.07 MiB
--no-kv-unified 214.86 -> 70.86 MiB, CUDA0 317.07 -> 265.07 MiB
Selection is unchanged: over 8192 tokens, four times the budget, every QSA
layer returns identical top-k indices and the logprobs are bitwise equal.
Two things a reviewer should know. A cell whose position divides past the last
block is guarded by an assert rather than handled, because no run reached it.
And the mask's same-position M-RoPE rule cannot fire for text and was never
exercised for images, so the 2D case is unverified.
* clean up code comments
* clean up new comments
* revert LLAMA_MMAP_RANDOM
* nits
* replace some changes with #27795
* improve the m-rope image for get_prev_tokens
* LazyChunkedTensor
* fix lint
* add some validations
* reduce input nodes
* trim output tokens
* nits
* some more sanity checks
* fix llm_graph_input_ple reuse
* exclude from webgpu test
---------
Co-authored-by: danielhanchen <[email protected]>
Co-authored-by: danielhanchen <[email protected]>
Co-authored-by: Xuan Son Nguyen <[email protected]>
Co-authored-by: Pascal <[email protected]>
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* spec : add DFlash2 support (local convolution + candidate selector) (#27342)
* support DFlash2
* Add p_min in DFlash2
Assisted-by: Claude Opus 5
* Revert unnecessary changes
Assisted-by: Claude Opus 5
* Revert draft sampling in rejection sampling
Assisted-by: Claude Opus 5
* Refactor code structure
Assisted-by: Claude Opus 5
* Delete embedding scaling
Assisted-by: Claude Opus 5
* Gate output transforms on DFlash2
Assisted-by: Claude Opus 5
* Optimize Dflash 2 cost
Assisted-by: Claude Opus 5
* Avoid using atoi
Assisted-by: Claude Opus 5
* Modify comments
Assisted-by: Claude Opus 5
* Move llama_model_dflash_selector_top_k to llama-ext.h
Assisted-by: Claude Opus 5
* Formatting
Assisted-by: Claude Opus 5
* Apply patch to fix the mrope bug
Assisted-by: Claude Opus 5
* fix ci
Assisted-by: Claude Opus 5
* Fix graph number calculation
Assisted-by: Claude Opus 5
* rename hid and unary
Assisted-by: Claude Opus 5
---------
Co-authored-by: Jian Chen <[email protected]>
Co-authored-by: Xuan-Son Nguyen <[email protected]>
* revert top-k.cu changes
---------
Co-authored-by: Zihan Zhang <[email protected]>
Co-authored-by: Jian Chen <[email protected]>
* ui : strip trailing container-format segments from parsed model names
* ui : show reasoning and modality icons on model options and search by modality
* ui : keep reasoning submenu visible regardless of model state
* ui : add show-org-name-in-trigger display setting
* ui : move model list into a submenu within the model selector
* ui : make model option hover and focus highlight override the active state
* ui : add raw model id tooltip to model selector options
* feat: Enable microphone input as default for audio models
* ui : fix eslint issues in chat form and model selector
* ui: show modality icons instead of file submenu in chat add menu
Assisted-by: pi
* chore: Format
* chore: Format
* ui: add ModelCapability enum and shared modality/capability icon constants
Assisted by: pi:GLM-5.3-Flash
* ui: derive modality badge icons and labels from shared constants
Assisted by: pi:GLM-5.3-Flash
* ui: split model option icons into capabilities and modalities
Replace the supportsThinking flag on ModelId with a capabilities object
keyed like ModelModalities, so future capabilities (tool calls, etc.)
slot in alongside reasoning. Icons and labels now come from the shared
CAPABILITY_ICONS/MODALITY_ICONS constants.
Assisted by: pi:GLM-5.3-Flash
* ui: replace per-conversation MCP overrides with per-conversation tool policy
MCP server enabled state is now global (server.enabled); per-conversation
control moves to disabled tool keys and categories seeded into each new
conversation. Aligns the add sheet with the dropdown options and flattens
MCP tool groups in the tools submenu.
Assisted-by: pi
* ui: keep tool policy migration running when defaults parse fails
A corrupt disabledToolKeys localStorage entry no longer aborts the
migration; it falls through with empty defaults so legacy MCP server
overrides still get converted.
Assisted-by: pi
* ui: fall back to global defaults when agentic flow has no tool policy
Passing empty disabled sets bypassed the global defaults and could
enable tools for callers that do not pass a policy yet.
Assisted-by: pi
* ui: align preferences section headers with their methods
The Reasoning Effort and Working Directory headers sat above tool
policy methods; move them above setCwd and setReasoningEffort. Also
clarify the disabled tools JSDoc: existing rows with an unset field
have an empty policy, defaults apply only when there is no active
conversation.
Assisted-by: pi
* ui: gate MCP server avatars on conversation tool policy
Servers whose tools are disabled for the current conversation (MCP
category or server-scoped key) no longer show as enabled for the chat.
Assisted-by: pi
* ui: drop unused MCP category toggle from tools panel hook
Per-conversation MCP control is server-granular; no component renders
a whole-category toggle, so remove the dead API.
Assisted-by: pi
* ui: skip MCP init when flow policy disables the MCP category
Resolve the effective tool policy before deciding whether to
initialize MCP so flows that will not send any MCP tools skip the
init work. Callers without a policy keep falling back to global
defaults.
Assisted-by: pi
* chore: format
* ui: restore reasoning section in mobile add sheet
The sheet rewrite dropped it; the desktop dropdown still has it.
MCP Prompts and Resources stay out of the sheet on purpose.
Assisted-by: pi
* ui: clear MCP server group key in enableAllToolsForServer
The group key disables every tool of the server regardless of
per-tool keys, so re-enabling a server from Settings did nothing
while it was set.
Assisted-by: pi
* ui: skip MCP init when no policy-enabled server remains
Extends the category-level check: the flow also skips MCP init when
every globally-enabled server has its server-scoped group key
disabled in the tool policy.
Assisted-by: pi
* ui: make Settings tools tab edit defaults with category toggles
Adds per-category checkboxes and a caption stating the tab applies
to new conversations; tool picks inside a chat only affect that
chat.
Assisted-by: pi
* ui: gate cwd picker and mention picker on effective tool policy
Both checked the global disabled set directly, so a conversation
that disabled file_search still showed search as available.
Assisted-by: pi
* ui: clean up tool key helpers and store docs
Documents getEnabledToolsForLLM properly, unstacks the JSDoc at
isEntryEnabled, makes setToolEnabled persist like setCategoryEnabled
(toggleTool now delegates to it), and routes the serverId-less MCP
branch of toolKey through getMcpServerToolsKey so both key formats
come from one place. Preferences banner comments become plain
comments so they no longer read as class member docs.
Assisted-by: pi
* ui: indeterminate group checkboxes and inert grayed rows
A category that is on with nothing enabled under it now shows the
mixed checkbox state instead of a checked box next to 0/N. Rows
grayed out by a disabled parent no longer stay clickable behind
opacity.
Assisted-by: pi
* ui: gate MCP prompt and resource capabilities on tool policy
hasPromptsCapability and hasResourcesCapability accept an optional
set of usable server ids; ChatFormActions resolves it from global
enablement minus the active conversation's policy. Restores the
per-chat gating the old mcpServerOverrides provided; callers without
arguments keep global behavior.
Assisted-by: pi
* ui: remove unmounted MCP submenu component
Never rendered anywhere; its entries are duplicates (prompts and
resources live in the attachment menu, servers in the add menu and
sheet) that would need capability wiring maintained for nothing.
Assisted-by: pi
* ui: fix model information dialog width on all screen sizes
The dialog sets container-type: inline-size, so auto width ignores
its contents and collapses to padding. Give it an explicit viewport
width on mobile and cap at 60rem on desktop.
Assisted-by: pi
* ui: scroll wide chat template in model information dialog
Long unbreakable Jinja tokens blew out the table and dialog width;
the block now scrolls horizontally instead of stretching.
Assisted-by: pi
* ui: use fixed table layout in model information dialog
Auto table layout sizes columns to content min-content, so the chat
template's long lines kept inflating the dialog despite the scroll
wrapper. Fixed layout pins the first column and gives the value
column a definite width the wrapper can scroll within. min-w-0 on
the grid item guards the same path on the grid side.
Assisted-by: pi
* ui: make model information dialog full-screen on mobile
Matches the settings dialog pattern: full viewport below md,
calc-sized and capped at 60rem on desktop.
Assisted-by: pi
* ui: stack chat template row in model information dialog
Label above the block in a single full-width cell, so the template
gets the whole table width and its horizontal scroll is usable on
narrow screens.
Assisted-by: pi
* ui: scroll model information header with the content
The base dialog header is sticky; this dialog overrides it to
relative so the title and description scroll away with the body.
relative keeps the header as the close button's containing block.
Assisted-by: pi
* ui: replace literal comment text in sheet group snippet
A // line inside the Svelte snippet rendered as visible text; use an
HTML comment.
Assisted-by: pi
* ui: let indeterminate state win over checked in group checkboxes
The checkbox indicator snippet renders the check icon whenever
checked, so the mixed state never showed. Pass the checked prop
as false while indeterminate.
Assisted-by: pi
* ui: initialize only policy-enabled MCP servers for a flow
ensureInitialized accepts an optional server id set; the agentic
flow passes the servers its tool policy leaves usable, so servers
disabled for the conversation no longer get connected. Callers
without arguments keep the global behavior.
Assisted-by: pi
* ui: derive group checkbox state in useToolsPanel
Moves the mixed-state derivation out of the submenu and sheet
snippets into one getGroupCheckState accessor; the snippets just
consume checked and indeterminate.
Assisted-by: pi
* ui: gate /prompt command on the conversation tool policy
The slash command's availability now follows the same rule as the
agentic flow instead of the global capability check, so it disables
itself when the conversation's policy leaves no usable MCP server.
Assisted-by: pi
* ui: remove dead MCP prompt menu trigger chain
The /prompt slash command is the surviving trigger; the menu-button
path (onMcpPromptClick, hasMcpPromptsSupport, showMcpPromptButton,
the MCP_PROMPT attachment item and its unrendered item arrays) has
no consumer left. Message display for inserted prompts is untouched.
Assisted-by: pi
* ui: render dash for mixed-state group checkboxes
The accessor refactor dropped the checked-and-not-indeterminate
guard, so the category-on flag won and the dash never showed. The
tooltip keeps using the raw parent flag since clicking a mixed
group still disables it.
Assisted-by: pi
* ui: fix group checkbox sticking checked after disable
Clicking a mixed-state group box let bits-ui optimistically flip
its internal checked flag; the derived checked prop did not change
across the transition (both mixed and off map to checked=false),
so Svelte never applied the settled value and the check icon stuck
while the count already read 0/7.
Pass the parent flag as checked and the mix as indeterminate, so
every group toggle changes checked; render the dash on top of a
checked box for the mixed state.
Assisted-by: pi
* fix: UI for Model Information dialog
* ui: keep MCP connections stable across policy switches
ensureInitialized folds the policy into its config signature, so
alternating two conversations with different policies tore down and
reconnected every server with health checks included. Tool collection
already filters by the flow policy, so initialize every
settings-enabled server instead and never pass a policy into the MCP
config. The duplicated policy-server check becomes one accessor on
ConversationPreferences.
Assisted-by: pi
* ui: remove dead MCP resources menu trigger chain
Same shape as the earlier prompt trigger cleanup: nothing renders the
MCP resources menu button, and the only live entry into resource
browsing is Settings > MCP Servers plus the attachment resource
picker. Drop onMcpResourcesClick, hasMcpResourcesSupport,
MCP_RESOURCES_CLICK, the AttachmentItemVisibleWhen enum and
hasResourcesCapability; the resources display, browser and picker
components are untouched.
Assisted-by: pi
* Add benchmark-only synthetic speculative acceptance to llama-server and llama-cli
* Address review comments
* Address review comments
* Add some comments in the code
* common : dedupe --n-cpu-moe / --spec-draft-n-cpu-moe override loops
* common : add --n-cpu-ffn to CPU-offload dense FFN weights of first N layers
* common : generalize llm_ffn_block_regex over the FFN regex, drop TODO
* hexagon: use non-host bufs by default and make the backend fully async
* hex-hb: remove optional hostbuf support and fix async copy
* hex-unary: relax supported unary check
* hex-bufs: use same get_alignment for host bufs
* snapdragon: bump android_platform to 34
* hex-rows: super hacky get/set rows for q8_0
* hex-get-rows: fix q8_0
* hex-get-rows: supprot for f16 and cleanup for q8_0
* hex-get-rows: generic macros and specialized thread funcs
* hex-get-rows: add DMA pipeline, vtcm_layout and kernel params
* hex-set-rows: fix q8_0 support, add dma and tracing
* hex-tests: override nmse threshold for HTP of Q8_0 quants
* hex-fa: add support for Q8_0 with inplace dequantizers
* hex-get-rows: simplify type dispatch
* hex-rows: simplify GET/SET_ROWS DMA pipeline
* hex-async: add events, set/get-tensor-async and rest of the async api support
* hex-repack: use slice instead of expert in repack functions
* hex-cpy: update event/async-cpy logging
* hex-set-rows: optimize smaller tensors
* hex-geglu: fix perf regression with larger tensors
* hex-get-rows: add missing header
* hex-set-rows: add missing header
* hex-bufs: ressurect GGML_HEXAGON_HOSTBUF but disable it by default
* hexagon: do not reject ops with non-heaxon buffers
* hex-get-rows: apply >=32 restriction only for q8_0
* hex-res: bump vtcm acquire timeout to 10 seconds
* hex-bufs: add support for cloning buffers between sessions to speed up tensor copies
* hex-async: rework event recording and batch flushing and integrate with meta backend
* hex-bufs: improved handling of repacked tensors
* hex-repack: handle get_tensor_2d offsets
* hex-dev: add support for devices with multiple NPUs
* hex-sync: add support for sync tokens to synchronize npu devices for async splits
* hex-mmap: cleanup mmap calls and add a retry for robustness
* hex-sync: add failsafe if sync wait gets stuck
* hex-sync: use sync_seq to check for completed events
* hex-sync: rotate tokens for extra robustness
* hex-devs: add supprot for legacy device names for now
* hex-bufs: add support for auto-cloning buffers from diff sessions
* hex-fusion: simplify and optimize htp-opnode fusion handling
* hex-sync: override opnode name so that it shows up in the profiles
* hex-trace: update scripts to handle multiple devices
* hex-sync: bump the size of the opbatch queue and number of sync tokens
* hex-cpy-sync: do not explicitly flush opbatches in cpy_tensor_async and add support for cpy-dma
* hex-sync: add graph-flush threshold to avoid single op batches
* hex-sync: add sync_peer so that we can flush peers we depend on during cross-device ops
* hex-bufs: introduce tensor->extra and shadow_bufs for repacking
* hex-l2: flush tiny tensors inline
* hex-sync: use explicit l2flush for sync tokens
* hex-extra: track weight flags via tensor extra
* hex-fence: rename sync to fence
* hex-repack: proper handling of set-tensor-2d in the shadow_buf
* hex-trace: remove obsolete opstage mask that we used for profiling
* hex-env: remove obsolete use_hmx variable
* hexagon: new unified run.py and build.py and updated docs
* snapdragon: update run script to auto-escapt test-backend-op -p argument
* hex-scripts: fix trailing spaces
* hex-scripts: fix flake8 warnings
* snapdragon: cleanup dst lib/bin dirs before copying new build
* hex-ops: add support for allreduce
* hex-ar: improved allreduce with dma pipeline
* hex-ar: align macros
* hex-ar: consistent use of fence_seq
* hex-ar: add AR_SELECT env var to select ALLREDUCE kernel or fallback
* hex-ar: add proper synchronize handling for ALLREDUCE
* hex-opbatch: looks like we now just rely on backend.synchronise to flush the batches, no need to flush them by threshold
* hex-ar: bump block size to improve dma efficiency
* hex-ar: fused ALLREDUCE+ADD
* hex-ar: cleaner fence buffer management
* hex-ar: futher allreduce tweaking to remove race conditions
* hex-ar: add simple solver and remove non-dma kernels
* hex-ar: add row-broadcast to fuse with bias ADD
* hex-fence: pass seq numbers via op_params
* hex-ar: allow for both entry/exit seq for completing entry wait
* hex-ar: align macros
* hex-ar: do not refetch broadcast row
* hex-fusion: move all fusion into opbatch::add_op for consistency with ALLREDUCE and things
* hex-fusion: fix incorrect MUL_MAT reordering
* hex-mm: make fused 2x and 3x matmuls more generic
* hex-fusion: move tensor fusion tagging to graph_compute
* hexagon: make sure to copy tensor->extra by value
* hex-get-rows: fix offset calc with row-chunking
* hex-repack: get_tensor_2d fixes for non-zero offsets
* snapdragon: make profile/trace scripts more robust and donot mix stdout/stderr by default
* hex-devices: use legacy device nameing by default to ease the transition
* hex-devices: hardcode CDSP domain IDs for current devices for now
* hex-optrace: improve multi-NPU timestamp alignment and overall handling of cycle values
* hex-optrace: more robust handling of the fence events
* ui : open MCP servers in a dialog from the chat form
Replace the MCP servers submenu with a single "MCP Servers" item that opens
a new DialogMcpServers dialog instead of navigating to the /mcp-servers route.
Assisted-by: pi
* ui : browse MCP resources from the server card
Make the Resources capability badge clickable so it opens the MCP resources
browser dialog, and drop the page-only chrome from SettingsMcpServers.
Assisted-by: pi
* ui : remove mcp-servers route and sidebar entry
MCP servers are now managed in a dialog, so drop the dedicated route and the
sidebar icon that navigated to it.
Assisted-by: pi
* ui : remove unused MCP servers submenu component
The submenu was replaced by the MCP servers dialog, so delete the component
and its export.
Assisted-by: pi
* feat(ui): add DialogSettingsChat dialog
* refactor(ui): switch SettingsChat to in-app section navigation
* feat(ui): open settings as dialog from sidebar
* refactor(ui): remove settings route and URL-based settings navigation
* fix(ui): adjust MCP dialogs for new base sizing
* chore: Formatting & linting
* feat(ui): make base dialog responsive and support sticky headers
* ui: move dialog close button to the sticky header
Assisted-by: pi
* chore: Formatting & linting
* ci : inline UI version resolution into ui-build.yml
* ci : build UI once and reuse the artifact in release jobs
Server jobs now extract the ui-build artifact into tools/ui/dist instead of npm-building the UI. Also removes the get-version job and the no-op -DHF_UI_VERSION flags.
Assisted-by: pi:Kimi-K3
* ui : disable the npm UI build by default (LLAMA_BUILD_UI=OFF)
The flag now only controls building the UI from source via npm. The UI
is still embedded by default from local tools/ui/dist or the prebuilt
download (LLAMA_USE_PREBUILT_UI=ON). CI jobs no longer npm-build the
UI; server-sanitize does not need node anymore.
Assisted-by: pi:Kimi-K3
* ci : rename the ui-build artifact to llama-ui.zip
Consistent with the other artifact names in the Actions summary.
Assisted-by: pi:Kimi-K3
* ci : clarify the windows artifact merge in release.yml
The windows-cuda/vulkan/sycl jobs build only the backend library;
llama-server (with the embedded UI) is injected into their zips from
the windows-cpu package during the release. State this in the job
comments and use accurate wording in the merge step.
Assisted-by: pi:Kimi-K3
* cuda: unblock mmq for MoE on sm_60
* cuda: duplicate mmq-config-pascal for dp4a and older
* cuda: reduce occupancy on non-dp4a pascal for Q2_K, Q4_K, Q5_K, Q6_K
This commit contains a fix for the conversion of NVIDIA Nemotron 3.5
Lightning which currently incorrectly converts when using a transformers
version later than 5.5.1.
When converting using [convert](https://github.com/ggml-org/convert) the
transformers version is 5.13.1 and this produces the following:
```console
WARNING:gguf.gguf_writer:Duplicated key name 'nemotron_h_moe.attention.head_count_kv', overwriting it with new value [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] of type ARRAY
```
This does not happen with transformers 5.5.1. The reason seems to be
that the configuration is different in later versions, for example when
using 5.13.1 the configuration block looks like this:
```console
transformers 5.13.1
raw has layers_block_type: True
autoconfig has layers_block_type: True
autoconfig layers_block_type: [
'linear_attention',
'moe',
'linear_attention',
'moe',
'linear_attention',
'full_attention',
'moe',
...
]
```
And with 5.5.1 we get:
```console
transformers 5.5.1
raw has layers_block_type: True
autoconfig has layers_block_type: True
autoconfig layers_block_type: [
'mamba',
'moe',
'mamba',
'moe',
'mamba',
'attention',
'moe'
...
]
```
In our conversion script we only match for attention, not full attention
which is causing this issue.
With the changes in this commit the output with transformers 5.13.1 will
be:
```console
(venv) $ gguf-dump models/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-BF16.gguf | grep head_count_kv
INFO:gguf-dump:* Loading: models/NVIDIA-Nemotron-3.5-Lightning-30B-A3B-BF16.gguf
29: [INT32] | 52 | nemotron_h_moe.attention.head_count_kv = [0, 0, 0, 0, 0, 2, ...]
```
Resolves: https://github.com/ggml-org/llama.cpp/issues/27718
Refs: https://github.com/ggml-org/convert/actions/runs/32949047680/job/98116096069#step:5:2391
* metal: WIP chunked SSD SSM_SCAN kernels for multi-token prefill
* metal: drop scalar SSD path; MMA + sequential tail
* drop WIP ssm scan test noise
* remove state_from_dst and rename CS and NSG constants
* remove unrelated added whitespace padding
* added clarity to mma_tokens calculation
* added clarity to use_mma bool checks
* added comments to metal ssd op constants for clarity
* reserve K tokens for sequential kernel rollback snapshots
* reset concurrency between mma and seq tail
* remove print args no longer used
* fixed comment to no longer point to specific line
* add FC_SSM_SCAN so seq path skips token offlset unless it's mma tail
* added changes to new ssm.metal for rebase after ggml-metal.metal refactor
* specialize ssm_scan tail with a template instead of a function constant
---------
Co-authored-by: dpantaleoni <[email protected]>
Co-authored-by: forforever73 <[email protected]>
* rpc: support apple RDMA as an RPC transport
* remove set_tensor micro optimization, rpc socket pinning per CR
* remove transparent reconnect
* trigger apple builds on RPC changes
---------
Co-authored-by: Ryan Churaman <[email protected]>
* devops: use GGML_NATIVE=OFF for OpenVINO
Same as in other Dockerfiles.
Should fix#23100
* enable backend dl and cpu all variants
---------
Co-authored-by: Sigbjørn Skjæret <[email protected]>
* server: fix tool calls getting silently stripped with --prefill-assistant
Last assistant carries tool_calls + --prefill-assistant is on → request
flips into continuation mode, add_generation_prompt forced off, tail
rebuilt from reasoning_content + content only. Tool calls just vanish.
- Auto-continuation now skips trailing assistant msgs that have tool calls
- continue_final_message on those throws a clear error instead of
silently corrupting the prompt
- Regression tests included, red before / green after
Fixes#27588
Developed with AI assistance, disclosed per the contribution policy.
* server : address review: fail on prefill-assistant + trailing tool_calls
Move validation into oaicompat_chat_params_parse (next to the existing
two-or-more-assistant check) and remove it from common_chat_templates_apply,
which has no precedent for validation. Drop the regression tests.
Per review: --prefill-assistant with a trailing assistant message
containing tool calls is not supported and should fail loudly.
* add ccache-buckets action
* use ccache-buckets
* only save on master
* install python3-venv for hip
* add jq and python3 for cuda
* only delete caches older than 5 minutes
* metal : null-check ggml_metal_buffer_init result to avoid OOM crash
ggml_backend_metal_buffer_type_alloc_buffer used the result of
ggml_metal_buffer_init without checking for NULL. ggml_metal_buffer_init
returns NULL when the underlying Metal allocation fails (e.g. an
out-of-memory condition), and the following ggml_metal_buffer_is_shared(res)
call dereferences it, turning a recoverable allocation failure into a hard
crash (EXC_BAD_ACCESS). This is easy to hit on memory-constrained devices
such as iOS when a model/context exceeds the available Metal budget.
Log the failure using the existing GGML_LOG_ERROR convention and return
NULL so the allocator surfaces a diagnosable error up the stack instead of
crashing.
* cont : fix log
---------
Co-authored-by: Georgi Gerganov <[email protected]>
* grammar : accept "\-" escape in character classes
gbnf_escape_char_class() escapes '-' as "\-" but parse_char() rejected
that escape, so generated tool-call grammars failed to parse.
Assisted-by: Claude Code <[email protected]>
* tests : add parser test for "\-" in char classes
Assisted-by: Claude Code <[email protected]>
* tests : add integration test for "\-" in char classes
Assisted-by: Claude Code <[email protected]>
* tests : drop integration and parser tests
* metal : per-device tuned (Q, NE) for flash-attn vec (#25750)
* rebase Q-generic FA vec body from 01dc93607 (#23114)
* add 53 f16 (Q,NE) flash-attn vec instantiations (vec 80 -> 133)
* add FA vec (Q,NE) tuning table + dispatch wiring + SMEM cap fallback
* add FA vec (Q,NE) perf sweep
* fill tuning result
* fold family table into a per-family representative SKU
* refactor tuning result format
* extend FA vec tuning to quantized KV caches
* sync fa vec tuner bucketing with runtime, use pointwise tuning regret
* update tuned table
* format and cleanup
* prefix fa_vec tuning procs with ggml_backend_metal_tuning_, drop unused fa_vec_override_active
* add device id -> token lookup for the offline tuning tool
* add ggml-metal-tuning skeleton
* add op-agnostic perf cell + median timing for the tuner
* add FA-vec graph build + tensor init to the tuner
* tools : add FA-vec (Q,NE) sweep, compression and table emit
* cool down and re-measure the dirty window on thermal drift
* test-backend-ops : replace the FA vec tune mode with a bounded (Q,NE) slice
* tools : document the Metal tuner, point the table comment at it
* abort on unknown KV type, single-source fa_vec_legal_ne
* cleanup
* honor -o in the FA vec (Q,NE) slice
* retune FA-vec (Q, NE) under a pointwise no-harm gate
* cont : add fa-vec tunings for M1 Pro, M2 Ultra, M5 Max
---------
Co-authored-by: Georgi Gerganov <[email protected]>
* metal : per-op source split + parallel compile (#24021)
* preliminary extract common header
* op source split
* split metallib into 8 libs && load in parallel
* derive kernel->library routing from functionNames
* x-macro lib list + underscore filenames, dedup QK_NL, MRC fixes
* op source split 8 to 20
* improve robustness of source fallback
* clean up
* change bool -> atomic_bool
* only prepend headers that source actually includes
* no semaphore, use GCD global queue
* dedup library compile path, fix NSError lifetime, rename gla
* relocate upstream concat/rope_back/repeat kernel changes into split files
* move ggml-common.h from common.h into dequantize.h to shrink binary size
---------
Co-authored-by: lvyichen <[email protected]>
* metal: add col2im_1d op (f32/f16/bf16) (#25176)
* metal : add set_rows with src0 f16 (#25434)
* metal : add CONV_2D_DW (depthwise convolution) support (#21565)
* metal : add Q2_0 support (#25419)
* metal: fuse snake activation (mul, sin, sqr, mul, add) (#25459)
* ggml-metal: FWHT kernel for metal backend (#25924)
* metal : port new kernels into the split sources
Move the kernels added on master after the split (lightning indexer,
DSv4 hyper-connections, silu_back, f16 bin ops, TQ2_0, the flash-attn KV
dequantization pass, rope offset/inplace, ssm_scan rollback, packed q8_0
dequantization and the tensor-API mat-mat K clamp) into the corresponding
kernels/*.metal sources. Copied verbatim, no functional change.
---------
Co-authored-by: lvyichen <[email protected]>
Co-authored-by: Georgi Gerganov <[email protected]>
`repetition_penalty` is standard HF key for repetion penalty.
Currently, only `penalty_repeat` is mapped, read `repetition_penalty`
and map it to `metadata.sampling_penalty_repeat`.
* ci : apply ccache-clear with older/min/dry-run to all ccache jobs
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : install gh in ccache-clear if missing (container jobs)
The ccache-clear action relies on the gh CLI, which is not present in
container-based jobs. Install it on demand so those jobs can clear caches.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : install gh via apt repo in ccache-clear
The install.sh script used previously is no longer served (404). Switch to
the official GitHub CLI apt repository, which is still available.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : pass --repo to gh cache commands in ccache-clear
In container jobs gh cannot auto-detect the repository from git, so
gh cache list/delete fail with 'failed to run git: not a git repository'.
Pass the repository explicitly via --repo using GITHUB_REPOSITORY.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : drop -new suffix from vulkan ccache key
The -new suffix was only needed to force a fresh cache. With
ccache-clear now evicting stale caches, the original key can be used
again. The old ccache-vulkan-ubuntu-24.04-arm-new entries still match
the ccache-clear key prefix and are cleaned up automatically.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : fix ccache-clear date parsing on macOS (BSD date)
macOS ships BSD date, which has no -d option. The older cutoff check
was silently disabled there: 'date: illegal option -- d' errors in the
log and the loop was only stopped by the min limit, risking deletion
of caches not older than the cutoff (e.g. saved by a concurrent job).
Parse the ISO-8601 timestamps with GNU date when available and fall
back to BSD date otherwise (TZ=UTC, fractional seconds dropped).
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : extract ccache-clear logic into scripts/ccache-clear.sh
The composite action now consists of a dedicated step that installs the
GitHub CLI when missing (e.g. in container jobs) and a thin step that
calls the new script. The script follows the make-release-checks.sh
conventions (usage/env header, set -euo pipefail, CLI flags) and only
checks that gh is available. The action inputs are unchanged, so the
workflow steps are untouched.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* ci : remove unused apple ccaches
description:"Corresponds to the ccache --evict-older-than AGE option, where AGE is the number of seconds or days followed by the 's' or 'd' suffix respectively."
# Graphics devices are leaked by Metal in Apple code sometimes, so we ignore those leaks
OBJC_DEBUG_MISSING_POOLS=YES "${cmd[@]}" 2>&1 | awk '{ print } index($0, "autoreleased with no pool in place") && !/class [a-zA-Z0-9]+Device autoreleased/ { found = 1 } END { exit found }'
@@ -20,8 +20,8 @@ If AI is used to generate any portion of the code, contributors must adhere to t
1. Explicitly disclose the manner in which AI was employed.
2. Check for an existing PR addressing the same change; if one exists, comment there to work with its author instead of opening a duplicate.
3. Perform a comprehensive manual review prior to submitting the pull request.
4. Be prepared to explain every line of code they submitted when asked about it by a maintainer.
3. Perform a comprehensive manual review prior to submitting the pull request. A proper code review usually takes something like one hour per 200-400 LOC and you should be spending **at least that much time on code review alone**.
4. Be prepared to explain every line of code you submit when asked about it by a maintainer.
5. It is strictly prohibited to use AI to write your posts for you (bug reports, feature requests, pull request descriptions, Github discussions, responding to humans, ...).
For more info, please refer to the [AGENTS.md](AGENTS.md) file.
LOG_WRN("DEPRECATED: `--load-mode` and `--mlock`/`--mmap`/`--direct-io` should not be combined; only the last flag on the command line will take effect\n");
}
};
// parse all CLI args now, so that -hf is available below for remote preset resolution
"JSON schema to constrain generations (https://json-schema.org/), e.g. `{}` for any JSON object\nFor schemas w/ external $refs, use --grammar + example/json_schema_to_grammar.py instead",
"JSON schema to constrain generations (https://json-schema.org/), e.g. `{\"type\": \"object\"}` for any JSON object",
"File containing a JSON schema to constrain generations (https://json-schema.org/), e.g. `{}` for any JSON object\nFor schemas w/ external $refs, use --grammar + example/json_schema_to_grammar.py instead",
"File containing a JSON schema to constrain generations (https://json-schema.org/), e.g. `{\"type\": \"object\"}` for any JSON object",
string_format("ip address to listen, or bind to an UNIX socket if the address ends with .sock (default: %s)",params.hostname.c_str()),
string_format("IP addresses to listen on, comma-separated, or UNIX socket paths ending in .sock; with multiple TCP addresses, :: binds IPv6 only; overlapping addresses result in undefined behavior (default: %s)",params.hostnames[0].c_str()),
[](common_params¶ms,conststd::string&value){
params.hostname=value;
params.hostnames.clear();
for(auto&host:parse_csv_row(value)){
host=string_strip(host);
if(!host.empty()){
params.hostnames.push_back(host);
}
}
if(params.hostnames.empty()){
throwstd::invalid_argument("--host requires at least one address");
Some files were not shown because too many files have changed in this diff
Show More
Reference in New Issue
Block a user
Blocking a user prevents them from interacting with repositories, such as opening or commenting on pull requests or issues. Learn more about blocking a user.