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
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 <noreply@anthropic.com>
* cont : fix var scope
* cont : comment out tests temporarily
Comment out tess to not break CI temporarily
Assisted-by: Claude Fable 5.1
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* rpc : hash-cache only weights
ggml_backend_rpc_buffer_set_tensor and ggml_backend_rpc_set_tensor_async
hashed every transfer above HASH_THRESHOLD and let `rpc-server -c` serve it
from its file cache. The cache is meant for weights, but the activations
ggml_backend_sched copies between backends took the same path: with a
two-node split of Qwen3.8-Flash-Next every prefill ubatch above 10 MB was
hashed, written to the worker's cache directory (1.4 TB after a day) and
later served from there. Use the hash path only for tensors in buffers
marked GGML_BACKEND_BUFFER_USAGE_WEIGHTS.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
* rpc : save a cache entry only for the tensor that missed the hash check
With the client hashing weights only, the server still wrote every
SET_TENSOR above HASH_THRESHOLD to the cache directory, so the compute
data the scheduler sends kept filling the disk. Remember the hash of the
last SET_TENSOR_HASH that missed and save only the SET_TENSOR that
follows it with that hash - the weight the client is re-sending.
* rpc : signal the cache decision in the SET_TENSOR payload
Replace the server-side `pending_cache` state with a `cache_flag` byte
in the SET_TENSOR message: the client sets it when SET_TENSOR_HASH
reported a miss, the server saves a cache entry only when it is set.
Bump RPC_PROTO_MAJOR_VERSION since the wire format changes.
---------
Co-authored-by: Patrick Hoffmann <patrickhoffmann@MacBook-Pro-14-HOP.local>
Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
* cuda: support row-contiguous SUM_ROWS
* organize the code and add GGML_OP_MEAN to support row-contiguous tensors using the same shared kernel, and add a test to MEAN permute/slice
* Keep original comments and add if/else branch
* exclude GPU/NPU failing POOL_2D case
* Fix pool case
* ggml-openvino: fix stateful decode for Gemma-4 per-layer-type head sizes
* ggml-openvino: fix MSVC narrowing error in permute
* ggml-openvino: classify sliding-window layers structurally on interleaved-SWA models
* ggml-openvino: add GGML_OPENVINO_REQUANT_KQUANT to select a 4-bit requant target
* ggml-openvino: add GGML_OPENVINO_SPILL_DIR to spill weight buffers to disk
* Stateful Performance: Added pass::KVStateSeqAxis to change KV layout
* ggml-openvino: fix stateful decode past the sliding-window size
Assisted-by: Claude Sonnet
* ggml-openvino: refuse stateful decode that cannot resume from the KV state
The stateful path seeds its KV state from ggml's cache when the decode position
is ahead of what the state holds. That only works when ggml's cache is a plain
prefix, where cell i holds position i. A sliding-window layer keeps just the last
n_swa positions and drops the rest, so past the window cell i no longer holds
position i and the seeded state is wrong.
Slicing the state to the decode position also had no bounds check, so a position
past the end surfaced as a bare ov::Exception from the ROI constructor
(llama_decode ret = -3, with no reason given at default verbosity).
Refuse both cases with a clear message instead, and refuse on the compile path
too, where a new model starts with an empty state and so can only serve a
sequence from its beginning. Reproducible with llama-bench -d, which restores a
saved sequence state rather than recomputing the depth prefill.
Assisted-by: Claude Opus 5
* ggml-openvino: use the per-layer KV head count for the stateful KV state
The stateful path reinterprets ggml's KV buffer [1, 1, seq, n_heads_kv * head_size]
as [1, seq, n_heads_kv, head_size]. The head size is already taken from the
tensor's own combined dim, because gemma-4 varies it per layer type, but the head
count still came from a model-level scalar that compute_llm_params() overwrites
per attention node, so it ended up holding whatever the last layer said.
gemma-4 varies the head count per layer too: 12B has 8 x 256 sliding layers and
1 x 512 full layers, 31B has 16 x 256 and 4 x 512. So 40 of 12B's 48 layers were
split as 1 x 2048 instead of 8 x 256, and attention read the state with the wrong
head split - both models decoded garbage on CPU and GPU. E2B is unaffected, its
head count is 1 everywhere.
Record the count per layer instead and look it up by the cache_k_l<N> leaf name.
Key it by layer, not by layer type: the sliding/full classification comes from
cache extents, which tie at a small -c, while the head count does not.
The stateful state trim now derives its sequence axis per state for the same
reason, since pass::KVStateSeqAxis matches per state on the head count.
Assisted-by: Claude Opus 5
* ggml-openvino: apply the KV state relayout to any KV head count
pass::KVStateSeqAxis was limited to states with a single KV head, where moving
the sequence axis from dim 1 to dim 2 is a pure metadata change. The limit was
also based on a measurement showing no gain for a multi-head model, but that was
taken at depth 0, which is the one depth where this change does nothing.
With several heads the pass does more than move metadata: it drops the reader
side transpose of the whole accumulated state, which the graph otherwise redoes
every token at a cost that grows with the context length, and replaces it with a
transpose of the single new row. Measured on GPU, tg128, alternating arms:
gemma-4-12B 6.27 -> 9.11 t/s at depth 8192 (stateless is 7.69, so stateful now
wins at depth instead of losing), Llama-3.2-1B 47.8 -> 59.6 t/s. Both are within
noise at depth 0, which is why the earlier check saw nothing.
The state refill needs the rows copied rather than reinterpreted now: ggml stores
[seq][n_heads_kv * head_size], and a relayout state with several heads is a
different element order. Without that, a refill would seed wrong data - it is
reachable today through llama-bench -d.
Assisted-by: Claude Opus 5
* ggml-openvino : support ggml_rope_set_offset and simplify op support gating
* add more cpy cases
* reject BF16 cpy on NPU
* Remove mul_mat_id fallback, gate large mul_mat_id only for mxfp4
* ggml-openvino: fuse the MoE expert block into MOECompressed on GPU
* ggml-openvino: skip GPU MUL_MAT_ID for unbound expert tensors
* ggml-openvino: requantize grouped 8-bit MoE experts on GPU
* Enable special strided CPY for conv state writeback
* openvino: support cacheless encoder models on NPU
Packed QKV views used by mmBERT were rejected by the ROPE support check. This split Q/K RoPE onto CPU, prevented cacheless attention detection, and sent fragmented encoder graphs through the decoder-oriented NPUW path.
Accept packed QKV RoPE views, detect cacheless attention from its mask, and run these models as a single full-sequence prefill without NPUW or a decode graph. Also provide static mask, output index, and mean-pooling shapes and inputs.
* openvino: optimize norm and RoPE translation
Replace the decomposed mean/variance normalization graph with an opset6 MVN operation. This preserves the GGML epsilon placement while allowing OpenVINO plugins to compile normalization as one operation with fewer intermediate tensors.
Cache RoPE sine and cosine outputs in the graph-wide tensor map. Build the cache key from all RoPE parameters and the optional frequency-factor input so compatible Q/K and layer nodes share one subgraph without mixing different RoPE configurations.
Expose NodeContext::put_shared() to publish translator-created outputs for graph-level reuse.
* ggml-openvino : simplify op translators and enable IMROPE/NEOX RoPE fusion
* remove unnecessary include and clean up PAD
* fix mulmat bug
* use ov::as_type_ptr instead of std::dynamic_pointer_cast
* ggml-openvino: fix mixed-dtype ADD/SWIGLU_CLAMP, gate unsupported ROPE/SOFTPLUS cases
- translate_add: upcast mismatched operand types (e.g. f16/f32 in fused
ADD_ADD) to f32, add, then cast once to the output type. opset1::Add
requires matching input types and downcasting first lost precision.
- translate_glu_swiglu_clamp: same fix, f16 Swish/Clamp rounding was
drifting past the test tolerance.
- supports_op: reject ROPE with ne[3] > 1 (multi-sequence) since the
cos/sin tables only cover one sequence, and SOFTPLUS on GPU since the
OpenVINO GPU kernel overflows to inf for large inputs (CPU is fine).
- ci/run.sh: serialize test-backend-ops on OpenVINO GPU; running two
workers concurrently crashes the GPU plugin (CL_OUT_OF_RESOURCES).
* openvino: share compiled models with per-context inference state; fix thread-safety
* ggml-openvino: gate MoE expert-sum ReduceSum shortcut past 8 experts
The ReduceSum shortcut for the MoE expert-plane-sum ADD chain drifts past
the 1e-7 test tolerance for >8 experts (f32 accumulation order vs CPU
reference), intermittently, like the existing Q4_K/Q5_K NMSE case.
Expose is_moe_expert_sum_add() so supports_op can gate on expert count
and fall back to CPU for just that reduction op.
* ggml-openvino: gate degenerate m=1,n=1 MUL_MAT on GPU
CI hit ERR=1.8e-3 (> 5e-4 tolerance) for a scalar-output f32 dot product
(m=1,n=1,k=2048); didn't reproduce locally in 8 tries, so likely an
internal fp16 accumulation path the GPU plugin picks for this tiny
shape. m=1 output dim doesn't occur in real model weights, so gate it.
* ggml-openvino: make SoftPlus decomposition opt-in native
Assisted-by: Codex
---------
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com>
Co-authored-by: ravi9 <ravi.panchumarthy@intel.com>
* metal : add FA kernels for HSK=96, HSV=64 (MiniCPM3)
MiniCPM3 sets attention.key_length to 96 and does not set
attention.value_length, which defaults to n_embd / n_head = 64. Metal had no
(96, 64) instantiation, so -fa auto aborted on the missing
kernel_flash_attn_ext_vec_f16_dk96_dv64.
Instantiate the tile kernel at (96, 64) for every K/V type that already has
(96, 96), and the vec kernel for the NE=4 configurations. Of the NE values the
vec dispatch considers, only NE=4 works here, because NL = 32/NE has to divide
both DK/4 = 24 and DV/4 = 16.
* tests : avoid redundant FA vec slice coverage
* 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
* 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 <johannesg@5d6.de>
* 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 <titaniumtown@proton.me>
* 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 <jianyu.zhang@intel.com>
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
Co-authored-by: Neo Zhang <NA>
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.
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 <maxk@qti.qualcomm.com>
* 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 <alexlu@qti.qualcomm.com>
* ggml-webgpu: Update to a recent version of Dawn
* No module scanning
* Accept review suggestion to update comment
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
---------
Co-authored-by: Masashi Yoshimura <yoshimura.masashi.frbs@gmail.com>
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.
* 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 <johannesg@5d6.de>
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.
