* jinja : parse unary +/- before variables
Lexer already emits unary_operator for -n / +n, and runtime executes
unary -. Parse them at multiplicative precedence so slices like
items[:-n] and GigaChat indent[:-indent_factor] work.
* jinja : keep filters/tests outside unary operands
Unary +/- must bind only the primary/postfix operand so -n|abs is
(-n)|abs, not -(n|abs). Add unary + and filter/test regression coverage.
Signed-off-by: sinksilk <[email protected]>
---------
Signed-off-by: sinksilk <[email protected]>
The OpenAI chat completions API specifies content part type "video_url"
with a {"url": ...} object, and clients typically send data: URIs
(e.g. data:video/mp4;base64,...). The llama-server only accepted the
non-standard "input_video" type and rejected data: URIs for video
(accept_base64_uri=false), so any OpenAI-conformant client failed with
"unsupported content[].type" or "Invalid uri format".
- accept "video_url" as an alias of "input_video"
- read the media object from whichever key was used
- allow data: URIs for video (data:video/*), as already done for images
This commit adds a new recipe/target to the Makefile which allows the
logits verification to be run on pre-existing model outputs.
The motivation for this is that for large models it can take a long time
to run them models, and especially for the original model which seldom
changes this is very time consuming. With this change we can run the
original model one which will store the tokens and logits, and then
manually run the converted model and the run use this recipe to verify
them against the orignal model.
* dspark: add Gemma 4 draft support
Add GGUF conversion and runtime support for full-attention and SWA Gemma 4
DSpark drafts, including tied output weights and boolean backbone metadata.
Assisted-by: Codex
* dflash: infer Gemma draft features from metadata
* 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
# 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 }'
- name:Test
id:cmake_test
@@ -74,16 +100,6 @@ jobs:
cd build
ctest -L main -E "test-llama-archs" --verbose --timeout 900
@@ -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.
Sigbjørn Skjæret
Si Chen
YiChen Lv
calebrio02
Xie Wenxiang
Daniel Bevenius
Hrishith Thadicherla
Georgi Gerganov
Piotr Wilkin (ilintar)
Ruben Ortlam
Aman Gupta
Neo Zhang
Anant Shrivastava
Erik Winter
Max Krasnyansky
Felix Ye
shaofeiqi
Pascal
David Friehs
Empressia
Jiang, FishandLiu, Russell <[email protected]>
Xuan-Son Nguyen
Michael de Gansandyomaytk
Bartowski
David M. Rogers
Eric Rodrigues Pires
Emanuil RusevandXuan Son Nguyen
wendadawen
miyan
Nicolas Mowen
Yuri Khrustalev
Agustín Mista
Asahi-Prv
Adrien Gallouët
Foad Abo Dahood
lingyezhixing
leejet
cwriterandcwriter
Łukasz Ślusarczyk
ynankani
Mikolaj Kucharski
Nandan Vallamdasu
Mostafa
Silverside
pl752andCopilot Autofix powered by AI
Samriddha Sinha
Yangyu ChenandJohannes Gäßler
Niklas Wenzel
Aldehir Rojas
Aleksander GrygierandPascal
Andrei
bri-prism
Aparna M P
chiheb ben cheikh
dsproule
Todor Boinovski
Alexey Kopytko
Aaron Teo
z
Daniel Varga
Sait Furkan Teke
Nikita Gordeev
yanghongandSigbjørn Skjæret
drluoto
Pedro Cuenca
Kartik Gulia
Ravi PanchumarthyandMostafa Faheem
midagedevandSigbjørn Skjæret
Titaniumtown
lhez
Abir Deol
Michael Taylor
Jeff Bolz
Eve
Gaurav Garg
Marco ColomboandMax Krasnyansky
GeorgeandSigbjørn Skjæret
uvos
I3eg1nner
y198
WenqiangJia2026
SG-Amadeus
asbelin
Jhen-Jie Hong
Trivikram Reddy
Sandro Steeger
Hongqiang WangandLi He
Mohamed Elashri
Chris Peterson
Yanzhao Wang
shivamkumard-ctrl
Aman Karki
Abhiram
Apoorv Parle
Christian Kastner
Alex
thelittlefiremanandJohannes Gäßler
Clint Herron
fairydreamingandStanisław Szymczyk
Yaniss Amazouz
Bernard Ladenthin
MiaoMing Chen
Alessandro de Oliveira Faria (A.K.A.CABELO)
Ed Addario
Mendy BergerandMasashi Yoshimura
Rohanjames1997
Logan Chu
Jesus Gulfo
Iggy Jackson
Julian Pscheid
Cordell Blanchardandgabby-zy
Masato Nakasaka
linsen458-specandlinsen
WakeUpMorty
Sarah Wu
Pepper Gray
Ankit Khandelwal
Frank Dai
Zhaolun Yin
Kevin HopperandKevin Hopper
AuroraRAS
Pranesh Gonegandlaandpraneshgo
DevVexusandClaude Opus 5
Siavash Norouzi
Anjielon
PikaPikachu
Daniel HanandGeorgi Gerganov
Jingxin (Philip) Li
Nick Farrell
Isaac
nachobh
yzyyzyhhhandhappyyzy
Tom Tan
Evan Huus
Frosty40
Ozymandias_EBON
Sergey Sklyarov
kbenkhaled
Mads Marquart
Nathan Wilson
Tarek Dakhran
EurekaticandRaulAbejonDelgado
cqderek
Mario Limonciello
Aman Chadha(IVIXMMI)andAcmmi
Laurent ZuijdwijkandMarshall
Alan Tseng
anujj
Lukasz Stolcman
James Francis
Ludovic Henry
Jonathan Clohessy
Buğra Özgürsoy