* Initial changes for Recurrent state rollback for nemotron for cpu and cuda
* Removing CPU RS rollback. Will enable it in subsequent PRs
* addition of test case
* Removing assert and calling runtime API to check if op is supported
* removing extra API and updating the call sites for K
* replace static cuda detection to runtime fused_op api
* address review comments and fallback when SSM rollback not supprted
* Adding changes for supporting RS-rollback in CPU. Also added test-backend-ops for cpu and cuda
* removing memory manipulation as rs rollback is now supported in CPU
* removing the static probe which is not needed now
* correcting the format
* address review comments
* enabling test for all the backends, unsupported backends will fallback to CPU
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* choose different graph based on the result of fused_ssm_op is supported or not and also handled memory->n_rs_seq >1 case incase of op is not supported
* Support K > 1 in ssm_scan for all backends
* Fix CI Issues
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Gaurav Garg <gaugarg@nvidia.com>
* OpenVINO backend: 1) enable gpt-oss moe on OV bk; 2) enable mxfp4 support
* OpenVINO backend: disable TOPK_MOE op test
* OpenVINO Backend: Add op FILL support
* OpenVINO backend: enable set rows with multi dims
* fix the name missmatch in setrow + view
* OpenVINO backend: enable op GGML_UNARY_OP_SIGMOID
* OpenVINO Backend: enable SQR & SQRT
* OpenVINO backend: 1) ensure unique node names for OpenVINO; 2) add org_src to recorde the src ggml tensor for OpenVINO dynamic shape infer
* OpenVINO backend: enable fallback for openVINO to CPU backend
* OpenVINO backend: fix accurace issue in gemma3n arch test
* fix mpt failed case
* OpenVINO backend: clean nodeinfo
* OpenVINO Backend: enable zero-size copy for view
* add concat ssm_conv in compute_dynamic_dim
enable qwen35
Fix after rebase
remove logging
* OpenVINO backend: disable EXP with FP32, which failed in op test. Root reason: the backend test initializes unary op inputs over a wide range, [-150, 150]. For FP32, exp(x) overflows around x ~= 88.7, so this test can randomly generate values right in or beyond the overflow region
* OpenVINO backend: fix CPY op test failed issue
* OpenVINO backend: fix GATED_DELTA_NET op test failed issue
* handle in-place op, handle qwen35 dynamic clearing of cache in cgraph
* handle qwen35 dynamic clearing of cache correctly
* Enable qwen35 dense multi seq
* Fix qwen35 9b gqa
* Fix after rebase
* Disable SOLVE_TRI
* openvino: fix NEOX RoPE accuracy on GPU stateful (mixed-rank Multiply)
In stateful mode the NEOX RoPE branch fed rank-3 data ([S, n_heads,
head_size]) into the Multiply against the rank-4 cos/sin tables
([1, S, 1, n_dims/2]). That mixed-rank broadcast is miscomputed by the
OpenVINO GPU plugin, corrupting the rotated Q/K and producing garbage
output (e.g. Phi-3-mini). Lift the data to rank-4 before the split/
Multiply so the operands are equal-rank, matching what the TYPE_NORMAL
branch already does. CPU and stateless paths are unaffected.
Phi-3-mini-Q4_K_M, wiki.test perplexity, GPU stateful:
before: PPL = 27120.43
after: PPL = 6.2263 (CPU reference: 6.2251)
* OpenVINO backend: 1) remove the unique name in llama.cpp; 2) add new ov name in ov bk; 3) fix issue in arch test & op test with latest code update
* OpenVINO Backenb: remove changes in llama.cpp
* Doc change (use x64 Native Tools Command Prompt for VS)
* Cleaner sentence
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
* OpenVINO Backend: cache key upgrade includes all src name
* OpenVINO Backend: enable llama arch test on ci
* OpenVINO Backend: move parameter node creating from decoder into translate
* OpenVINO Backend: create extra input ov node move from decoder to translate
* fix for op regression due to is_model_splitted
* openvino: fix CPY writeback for recurrent state rollback
Detect the rollback conv/gdn state writeback CPY nodes structurally
instead of by tensor name, since the rollback path in
build_conv_state does not call cb() and left the nodes unnamed. Add
per-node runtime offsets (rs_slot_begin_*, rs_src_begin_*) so the
cached IR handles any kv head, sequence count and snapshot slot for
both the conv state and the GDN state writeback.
Assisted-by: GitHub Copilot
* qwen35 moe
* optimize MoE expert aggregation with ReduceSum
* Skip GET_ROWS inaccurate test
* openvino: fallback dynamic MUL_MAT_ID shapes
* OpenVINO Backend: fix error in arch test model mpt
* fix error caused by cpy in arch test model kimi-linear
* OpenVINO Backend: fix error in arch test model minimax-m3
* openvino: fix GPU mul_mat_id op tests
* ggml-openvino: add GGML_OPENVINO_RELEASE_WEIGHTS to reclaim host weight RSS on GPU
The OpenVINO weight Constants are zero-copy views into host buffers
allocated by the backend (ggml_aligned_malloc, anonymous memory). On GPU
the plugin holds its own device copy after compile_model, so these host
pages are dead weight for inference. For a 1B Q4_K_M model this leaves
~850 MB of host RSS resident that the GPU path never reads again.
Add an opt-in GGML_OPENVINO_RELEASE_WEIGHTS mode that madvise(MADV_DONTNEED)s
the registered host weight buffers once the model is compiled, dropping
their resident pages while keeping the mappings valid (ggml still owns the
lifetime; tensors still point in). Measured steady-state RSS drops from
~1555 MB to ~710 MB on Llama-3.2-1B-Q4_K_M (Arc iGPU) with unchanged
throughput and correct output.
The GPU backend uses a single dynamic-shape model for both prefill and
decode, so a graph is compiled once and reused; the only event that forces
a recompile is clear_caches() on backend teardown. The change therefore:
- releases on the first cache-hit (model compiled, plugin has its copy);
- pins the compiled-model cache across backend teardown so a later
context reuses it instead of recompiling against the dropped pages;
- fails loud (GGML_ABORT) on a cache-miss recompile or on a second model
load, both of which would otherwise read zeroed weights or silently
reuse the wrong compiled graph.
Scope/limitations (all fail loud, never silently wrong): GPU only (the CPU
plugin reads the host Constants at inference time), one model per process,
and stable graph shapes. This reduces steady-state RSS, not the transient
compile-time peak. All changes are confined to the OpenVINO backend.
* ggml-openvino: stream weight requantization to cut the compile-time RSS peak
requantize_to_buffers() dequantized the entire tensor to a temporary
std::vector<float> of n_elements before requantizing. For token_embd.weight
(128256 x 2048) that transient is ~1 GB (1B model) / ~2 GB (8B), and it is
the single largest contributor to the OpenVINO compile-time memory peak --
it also fires twice for token_embd (once at load, once at graph build,
because token_embd is loaded via a CPU/mmap buffer and not cached as an OV
weight extra).
Stream the dequant instead: process a fixed window of complete rows
(CHUNK_ROWS=256) into a small scratch buffer and quantize/convert each chunk
straight into the output buffers. The transient F32 footprint is now
CHUNK_ROWS*ne0 floats regardless of tensor size.
quantize_q8_0/q8_1 gain an optional block_offset arg (default 0) so a chunk
writes its weights/scales/zp at the correct block. Streaming is applied to
the Q8_0_C / Q8_1_C / F16 targets (the large requant cases); the u4 (Q4_0)
path keeps the whole-array call because it packs two weights per byte with
running zp ORs, and a fallback handles any future target whose block size
does not divide a row.
Measured peak RSS (cold compile, GPU): 1B 2868 -> 1809 MB (-1.06 GB);
8B 11618 -> 9608 MB (-2.0 GB). Output verified unchanged
("capital of France is Paris"); throughput unchanged. Unlike
GGML_OPENVINO_RELEASE_WEIGHTS this reduces the transient peak, not just
steady-state, and needs no env flag. All changes confined to the OpenVINO
backend.
* ggml-openvino: avoid redundant token_embd requantization at compile
token_embd.weight is referenced twice in the graph path: as the GET_ROWS
embedding (a CPU/mmap-buffer tensor) it was re-extracted/re-requantized on
every weight-node build, and is_model_splitted() built a full (naive) set of
weight nodes just to test name membership — each requant is a ~1-2 GB F32
dequant of the 262M-element embedding.
Two changes:
- Add collect_weight_names(): a name-only collector for topology checks.
is_model_splitted() now uses it instead of create_weight_nodes(cgraph,
true), so the splitted-check no longer triggers any weight extraction.
- Memoize weight nodes built from non-OpenVINO buffers in a process-lifetime
cache keyed by tensor->data. These tensors have no OV buffer context to own
a cached extra, so without this they were rebuilt on every (re)compile;
prefill and decode graphs now share one build (verified: 2nd graph hits the
cache instead of re-requantizing).
Peak RSS is unchanged (the streaming-requant commit already removed the F32
transient); this removes redundant compile-time work. Output verified
unchanged ("capital of France is Paris"). Confined to the OpenVINO backend.
* ggml-openvino: gate compile-memory optimizations behind GGML_OPENVINO_REDUCE_COMPILE_MEM
The streaming requantization and the non-OpenVINO-buffer weight-node cache
(plus the name-only is_model_splitted path that pairs with it) are now opt-in
via GGML_OPENVINO_REDUCE_COMPILE_MEM. When unset, requantize_to_buffers()
fully materializes the F32 buffer and weights are rebuilt per compile exactly
as before; when set, the streaming path and the cross-compile weight cache
are used.
Default off keeps behavior identical to upstream unless explicitly enabled.
Verified: flag off -> peak RSS 2800 MB (original), flag on -> 1810 MB; output
"capital of France is Paris" in both modes. (GGML_OPENVINO_RELEASE_WEIGHTS,
added earlier, remains a separate opt-in for the steady-state release.)
* ggml-openvino: add frontend model cache (GGML_OPENVINO_MODEL_CACHE_DIR)
The plugin-level ov::cache_dir caches the compiled blob keyed by the OV
model, but producing that model still runs the full frontend every time:
weight requantization (incl. the large token_embd F32 transient) and the
ggml->OV graph conversion. This adds an opt-in frontend cache keyed off a
fingerprint computed directly from the ggml cgraph, so a hit imports a
previously exported CompiledModel and skips requant + convert + compile
entirely.
Key (model-cache.{h,cpp}) = 64-bit FNV-1a of: graph topology (n_nodes + per
node op/name), a sampled per-weight fingerprint (name/shape/type + bounded
head+tail byte sample), and blob-affecting config (device, flash-attn, rope
params, REDUCE_COMPILE_MEM/stateful flags, OpenVINO version). A sidecar
manifest stores every weight's fingerprint and is re-verified on load, so a
sampled-hash collision cannot cause a wrong-model hit (verified: two
different quantizations of the same model produce distinct cache entries).
Flow (dynamic single-model path only; split models defer to ov::cache_dir):
on a verified hit, core.import_model() restores the CompiledModel and a
lightweight decoder is built with a names-only weight map (membership is all
the decoder needs for I/O mapping; weights live in the imported model). On a
miss, compile as usual then export the blob (atomic temp+rename, manifest
written first). The frontend cache supersedes ov::cache_dir, so CACHE_DIR/
CACHE_MODE are stripped from the config used for the cached compile and the
import — a blob compiled with cache_dir set cannot be re-imported.
Measured 8B Q4_K_M (GPU): full requant+convert+compile 15.3s -> import 6.3s
(~2.4x faster compile phase). Output verified unchanged on cold and warm,
standalone and combined with REDUCE_COMPILE_MEM + RELEASE_WEIGHTS. Default
off; confined to the OpenVINO backend.
* ggml-openvino: harden frontend model cache correctness
The frontend model cache imports a previously exported CompiledModel keyed by a fingerprint of the ggml graph, weights, and blob-affecting config. The original key covered device, stateful execution, REDUCE_COMPILE_MEM, RoPE params, OpenVINO version, topology, and sampled weights, but missed runtime/frontend toggles that can change the lowered graph or the I/O binding contract. That made it possible to reuse a blob produced under a different OpenVINO backend configuration.
Add a small extra-config helper for the dynamic model-cache path and fold in the effective values of GGML_OPENVINO_DISABLE_KV_SLICE and GGML_OPENVINO_MANUAL_GQA_ATTN. MANUAL_GQA_ATTN is keyed by the behavior that actually takes effect: an explicit env value wins, otherwise GPU defaults to enabled and other devices default to disabled. This matches flash_attn_ext lowering and avoids unnecessary cache splits for equivalent configurations while separating genuinely different attention graphs.
DISABLE_KV_SLICE is also included because it changes the KV-cache tensor shape/output binding strategy used around imported models. Even when weights and graph topology are identical, switching this flag should not inherit a CompiledModel cache entry created for a different binding mode.
Also make cache artifact publication cleaner: write manifest.tmp and blob.tmp, publish the blob first, and publish the manifest last. Cache hits already require both blob and a verified manifest, so making the manifest the final visible artifact avoids leaving an apparently complete manifest for a failed or interrupted blob export. Temporary files are removed on the handled failure paths.
While touching this path, fix the indentation of the non-imported compile branch so the cache miss flow is easier to review. Behavior is otherwise unchanged: verified hits still import, misses still create weights, convert, compile, export, and create the infer request normally.
* ggml-openvino: add memory optimization umbrella switch
Add GGML_OPENVINO_MEMORY_OPTIMIZE as a single opt-in switch for the OpenVINO backend memory-saving paths. The existing fine-grained GGML_OPENVINO_REDUCE_COMPILE_MEM and GGML_OPENVINO_RELEASE_WEIGHTS variables remain supported and explicitly override the umbrella switch when set, so users can still bisect or disable one side of the optimization independently.
Centralize the policy in ggml_openvino_reduce_compile_mem_enabled() and ggml_openvino_release_weights_enabled(device). The umbrella switch enables compile-memory reductions everywhere REDUCE_COMPILE_MEM is used today: streaming requantization, non-OV weight-node caching, split-model weight-name collection, and the frontend model-cache fingerprint. On GPU it also enables host weight-buffer release unless GGML_OPENVINO_RELEASE_WEIGHTS is explicitly set.
Keep host weight release GPU-only because it relies on the plugin holding its own device copy after compile_model. Update the fail-fast diagnostic and comments to mention GGML_OPENVINO_MEMORY_OPTIMIZE, so users who enable the umbrella switch get accurate guidance if a later cache-miss recompile would read released host weight pages.
* ggml-openvino: rename compiled model cache env
Rename the frontend export/import cache environment variable from GGML_OPENVINO_MODEL_CACHE_DIR to GGML_OPENVINO_COMPILED_MODEL_CACHE_DIR. The cache stores blobs produced by ov::CompiledModel::export_model() and restores them with core.import_model(), so the new name distinguishes it from GGML_OPENVINO_CACHE_DIR, which configures OpenVINO plugin-level ov::cache_dir.
Update the registered env var, the cache-directory lookup, and comments around the frontend compiled-model cache. The old GGML_OPENVINO_MODEL_CACHE_DIR name is removed rather than kept as a fallback so there is a single spelling for the new option.
* docs: document OpenVINO memory optimization env vars
Add runtime configuration entries for the newly recognized OpenVINO environment variables.
Document GGML_OPENVINO_COMPILED_MODEL_CACHE_DIR as the frontend compiled-model cache used to export and import compiled blobs for matching single-graph models.
Document GGML_OPENVINO_MEMORY_OPTIMIZE as the umbrella switch, including how GGML_OPENVINO_REDUCE_COMPILE_MEM and the GPU-only GGML_OPENVINO_RELEASE_WEIGHTS override or inherit from it.
* ggml-openvino: fix Qwen3VL crash and deepstack correctness bug
1. GGML_OP_PAD was missing from compute_node_dynamic_dims(), causing a crash
on decode for models that pad the token embedding (n_embd -> n_embd_inp).
PAD never reorders/merges dims, so it keeps the same dynamic dim index as
its source.
2. process_view_input_new() chained VIEW inputs through src[0] (the
immediate op-graph parent) using offsets treated as relative to that
parent. But ggml_tensor::view_offs is always absolute from the true root
allocation (ggml collapses VIEW-of-VIEW chains internally). For the
per-layer deepstack view ("embd (view)", whose src[0] is "embd" - itself
an already-narrowed, zero-offset VIEW of the padded root, with the SAME
ggml shape as the deepstack view but a different absolute offset), this
caused an out-of-bounds re-slice that silently fell back to returning the
wrong (already-resolved sibling) tensor. In practice every deepstack ADD
ended up adding the real base token embedding into the residual stream
instead of zero, corrupting generation ("Hello my name is 1000000..."
instead of coherent text). Fixed by detecting this pattern (same shape as
the immediate src, different absolute offset) and re-slicing directly
from the untouched root tensor using the innermost view's absolute
offset.
Also adds a GGML_OPENVINO_DEBUG_NODE=<name1>,<name2>,... env var that attaches
extra debug Result nodes for arbitrary intermediate tensors, without binding
them to any ggml buffer (avoiding the risk of reading a ggml buffer that has
since been overwritten by a later in-place op). This was instrumental in
diagnosing bug #2 above and is left in as a general-purpose debugging aid.
* ggml-openvino: fix IMROPE inp_pos padding for NPU static shapes
IMROPE's inp_pos tensor packs 4 stacked t/h/w/e position planes into
ne[0] = 4*n_tokens instead of one value per token. On NPU's static-shape
path, inp_pos was padded/shaped as if it held a single plane, which
interleaved padding across the 4 planes and desynced later reshapes
from the rest of the (chunk_size-wide) graph.
- add GgmlOvDecoder::get_inp_pos_n_planes() to detect IMROPE's 4-plane layout
- get_graph_input_shape(): size inp_pos as n_planes * chunk_size (prefill)
or n_planes (decode) instead of assuming 1 value per token
- get_ov_input_tensor_static_prefill(): pad each plane to chunk_size
independently instead of one flat block
- get_ov_input_tensor_static_decode(): copy n_planes contiguous values
instead of asserting/copying a single scalar
* disable test-llama-archs tests.
* openvino: gate fallback with env var
* Revert changes in test-llama-archs
* Apply editor config
* reject CPY with quantized destination as unsupported
---------
Co-authored-by: Xuejun <Xuejun.Zhai@intel.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: suryasidd <surya.siddharth.pemmaraju@intel.com>
Co-authored-by: Mustafa Cavus <mustafacavus@intel.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
* metal: add TQ2_0 support
Add support for the GGML_TYPE_TQ2_0 (ternary, 2 bits per element) type in
the Metal backend.
Assisted-by: llama.cpp:DeepSeek-v4-Flash-0731
* cont : optimize mul_mv kernel
- float ops over integer ops
- precalculate sums
- hoist coef out of the inner loop
- contiguous y loads
llama.cpp:DeepSeek-v4-Flash-0731
* Add DMMV Q4_K and Q6_K ESIMD kernels
Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Refactor ESIMD kernels to share common code
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Move control of ESIMD from compile to runtime
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Use ESIMD by default when available
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Fix possible error when using ESIMD by default
While not an issue in the current version, this will become an
issue when additional QK ESIMD kernels are added (such as Q2_K).
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add explicit unroll to ESIMD kernels
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Tidy up ESIMD kernels a bit
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Add DMMV Q3_K ESIMD kernel
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
---------
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
It enables -fassociative-math, which reassociates FP reductions and can flip
greedy argmax on RDNA3.5 (e.g. MTP speculative decode diverging from the
non-speculative baseline). Drop it so HIP builds are IEEE-conformant.
Co-authored-by: Jim Wu <ywu@xilinx.com>
* gguf : harden loader against malformed tensor dims and metadata types
* gguf: address review on malformed-metadata hardening
- report the expected vs. actual type when general.alignment is not u32
- use ggml_nelements() > 0 for the zero-element guard and keep the
representability checks visually aligned
- add test-gguf cases for a wrong-typed alignment key and a zero-dim
tensor (both used to crash: assert-abort and SIGFPE respectively)
Ran tests/test-gguf: 164/164 pass. Used an AI assistant to help draft
these edits; reviewed and verified by me.
* cont : less comments
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* vulkan: TQ2_0 (ternary) support — dequant + dedicated mul_mat_vec + matmul via dequant_funcs
First Vulkan ternary type in ggml. Correctness: OM-125m TQ2_0 vs F16 top-12
logprobs identical to 4 decimals fully offloaded (float dequant path, no Q8_K
activation quant). Speed at 125m ~= F16 (overhead-bound at this scale); the
bandwidth win targets larger BitNet SKUs. MMQ/int-dot path intentionally not
wired yet.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* tests: enable TQ2_0 in backend-ops type lists
Vulkan now implements TQ2_0 (dequant, mul_mat_vec, mul_mm, get_rows); backends
without support skip via not-supported as usual. TQ1_0 stays disabled.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: Michael Trabalka <michael.trabalka@sqv.ai>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* llama: add new default load-mode auto which picks mmap unless a non-Metal iGPU is used
* Update ggml/src/ggml-hexagon/ggml-hexagon.cpp
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* set mmap_support to false on OpenCL backend
* fix order of load modes
* use -1 for auto
* resolve load mode auto earlier to correctly pick gpu host or cpu memory
* add load mode auto to llama-bench
* bump virtgpu api version, regenerate docs
---------
Co-authored-by: Piotr Wilkin (ilintar) <piotr.wilkin@syndatis.com>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* test new flash_attn test
* rebase and fix to disable subgrou matrices when max_kv_tile == 0
* delete log output
* Add i32 support to cpy and enables the all ops test
* restore the non target ci tests
* comment out of TODO of build-cpu.yml
* fix format
ggml_roll only asserts nb[0] == ggml_type_size, so a permuted src is a
valid input, but the CUDA and Metal roll kernels index by ne alone and
never read the nb strides. A non-contiguous src therefore produced
silently wrong results. Neither backend declared a contiguity
requirement in supports_op, so the scheduler did not fall back to the
CPU implementation, which does handle strides correctly.
Add the requirement to both backends, matching the existing
GGML_OP_ROPE guard, and add a permuted test_roll case.
ggml_metal_op_norm sized the threadgroup with
`nth = std::min(nth, args.ne00_t)`, which can leave nth not a multiple of
the simdgroup size. The kernels finish their row reduction with a
cross-simdgroup step where each lane of the last simdgroup reads one
per-simdgroup partial sum out of shmem_f32:
if (tiisg == 0) { shmem_f32[sgitg] = sumf; }
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
When the last simdgroup is partial it has fewer lanes than the
threadgroup has simdgroups, so the tail of the partial sums is never
read and the row sum is too small. For ne00_t = 33 nth becomes 33: two
simdgroups, but only one lane in the second, so one of the two partial
sums is dropped. The mean and variance are then wrong for the whole row.
Round ne00_t up to a whole number of simdgroups instead. Rounding up
rather than dropping the clamp keeps the threadgroup as small as
possible: deleting the line would raise nth to the next power of two
(ne00_t = 544 -> 1024 instead of 544), which costs idle lanes on 26 row
lengths below 8192 that were already correct, including 1536 and 3584.
GGML_OP_NORM is affected as well as GGML_OP_RMS_NORM - both dispatch
through ggml_metal_op_norm.
No mainstream LLM hidden size hits this: ne00_t is ne00/4 on the
vectorized path, so 4096, 8192, 2048 and friends all give a multiple of
32. It is reachable from other norm shapes, e.g. 320-channel norms.
Add NORM and RMS_NORM cases for ne0 = 33, 132 and 260 across the
existing eps values. 33 exercises the scalar path and 132/260 the
vectorized one, since only those divide by 4.
Before, on M3 Pro:
test-backend-ops test -b MTL0 -o NORM 25/50
test-backend-ops test -b MTL0 -o RMS_NORM 26/51
After:
test-backend-ops test -b MTL0 -o NORM 50/50
test-backend-ops test -b MTL0 -o RMS_NORM 51/51
test-backend-ops test -b MTL0 13943/13943
* cmake : add config version support (wip) [no ci]
This commit adds support for find_package using a version, for example:
```
find_package(ggml 0.19.0 REQUIRED)
```
examples/test-cmake has been updated to use this and build scripts have
been added to verify this manually. This is still a work in progress and
I'm not sure about the scripts and if we can find better ways to test
this but it might be useful to have for verification of changes to the
cmake build.
* cmake : add semver to ggml backends [no ci]
This commit adds a semver to the ggml backend modules files.
The motivation for this is that the backends are currently loaded just a
file extension, for example .so on linux. With the introduction of
semantic versioning installing a new version should just work but since
these files don't have a version they would get overwritten. Adding the
semver to the library names allows multiple version to be supported and
the correct one will be loaded by the code.
I've only tested this on linux and need to test on mac and win.
* Revert "cmake : add semver to ggml backends [no ci]"
This reverts commit 53a6c58a07591951324c891b9986b2cffe5c7972.
* examples : update build-install.sh and set GGML_BACKEND_DIR
* ggml : add fallback definitions for missing aarch64 HWCAP bits
* ggml : require HWCAP_ASIMDHP for the aarch64 fp16 cpu variants
Also rename has_fp16_va to has_fp16, the field gates the whole FEAT_FP16
extension, scalar and vector half-precision arithmetic together.
The NVFP4 quantization format stores a scaling factor for every group of
16 weights, packed into a single UE4M3 byte.
The SYCL GPU code was converting these scale values using the E4M3 path,
but that's *signed*, and these are unsigned values.
* tests: add SWIGLU perf cases
perf mode had no GLU coverage. Adds SWIGLU at 17408 columns, 512 and
2048 tokens, f16 and f32, with the operands both fused and split.
* sycl: consolidate fused-GLU kernels
They differed only in which op_* they called, so take the op as an argument and share a common launcher.
Their block sizes were all 256, so launch geometry is unchanged;
SYCL_GELU_BLOCK_SIZE and SYCL_SILU_BLOCK_SIZE lose their last users so are dropped.
* sycl: contiguous fast path for the fused GLU ops
o0 == n and o1 == n collapse the de-interleave index math to the
identity, so dispatch a flat kernel in that case. It fires for
ggml_glu_split with packed operands; a fused [gate|up] tensor keeps the
strided path. test-backend-ops perf -o SWIGLU on an Arc Pro B70: split
+14% f16 and +4% f32, fused unchanged.
- In MSL, declaring an array of matrix types like `threadgroup half4x4` causes
a 'no matching constructor' compilation error because MSL matrix types do not
have zero-argument default constructors and threadgroup variables cannot have
initializers.
- Fix this by declaring a POD `threadgroup half` array instead and casting
to `threadgroup half4x4 *` for matrix indexing.
Signed-off-by: JamePeng <jame_peng@sina.com>
* ci: prepare for amd rocm ci
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: fix editorconfig-checker
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: fix device not recognised
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: rename gpu-amd to gpu-hip
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: gpu-hip to gpu-rocm
haha
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* CUDA: allow integrated-GPU host output buffer in debug assert
On integrated GPUs (APUs), the scheduler can legitimately place a graph
node's output on the host-visible buffer, which ggml_cuda_compute_forward
already handles. The debug assert in ggml_cuda_graph_evaluate_and_capture
required every node output to be on the device buffer, so a debug build
aborts on such a node (e.g. attn_residual ADD -> ROCm_Host on RDNA3.5).
The source-tensor assert directly below already permits this via the
integrated + cuda_host exception; apply the same exception to the node's
own output buffer. Debug-only; no effect on release/compute.
Fixes test-recurrent-state-rollback on gfx1151 (Strix Halo).
* ci: enable unified memory for ROCm gfx1151 job
Work around a coherence issue on integrated RDNA3.5 (gfx1151) where GPU
kernels reading mmap-loaded weights can return incorrect output, which
makes test-llama-archs (and real inference) intermittently wrong.
GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 uses managed memory, which restores
coherence. Remove once the underlying ROCm/HIP issue is fixed.
* test-llama-archs: skip jamba on HIP backend
jamba produces incorrect output (~0.55 NMSE vs CPU) on the HIP backend on
RDNA3.5 (gfx1151); the SSM kernels need separate investigation. Skip it
for now, matching the existing per-backend carve-outs (WebGPU), so the
ROCm CI can run the test for the remaining architectures.
* ci: use HIP_LAUNCH_BLOCKING for ROCm gfx1151 job
The gfx1151 ROCm CI job produced incorrect inference output (qwen3 perplexity ~88 vs ~9.4) due to an async-execution correctness issue in the HIP path. Serializing kernel launches with HIP_LAUNCH_BLOCKING=1 restores correctness. This replaces the earlier GGML_CUDA_ENABLE_UNIFIED_MEMORY workaround, which did not fix batched inference.
* test-backend-sampler: skip top-k subtests on HIP backend
The ROCm backend does not support the TOP_K/ARGSORT op at vocab scale (no CUB; bitonic argsort is capped at ncols <= 1024), so top-k/top-p backend samplers cannot be offloaded. The penalties, set_sampler, mixed, and top_p subtests assert that offload happened, so they fail on HIP. Skip them until TOP_K is supported on the ROCm backend.
* Update tests/test-backend-sampler.cpp
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
* Update tests/test-backend-sampler.cpp
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---------
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Jim Wu <ywu@xilinx.com>
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* vulkan: add debug tooling to get more information about a DeviceLost error
* fix submission threshold applied too late
* use logging macros, throw instead of aborting
* clean up circular dependency
* ggml: add ggml_build_forward_order
ggml_build_forward_expand marks the tensor and all its ancestors for
compute, so using it as a pure ordering hint (keeping q, k and v
together) defeats ggml_build_forward_select: the unselected branch is
forced to run with inputs that were never uploaded. In the mtmd audio
graph this makes GEN_WAV calls execute the GEN_CODE branch with a
stale inp_code0, hitting the get_rows bound assert on CPU.
Add ggml_build_forward_order, which inserts nodes without the compute
flag; the flag is restored when the branch is actually selected.
Switch the q/k/v hints in clip_graph::build_attn to it.
* nit: reduce comments (AGENTS.md)
Georgi Gerganov
lnigam
Titaniumtown
Mendy Berger
Zijun Yu
Neo Zhang
jinzihao
Ian Faust
Todd Malsbary
Jim Wu
HarrisonSec
Jonathan Clohessy
Hongqiang Wang
michaeltrabalka-tech
lhez
Rafail Giavrimis
0
Ruben Ortlam
Masashi Yoshimura
Hitesh Chopra
Yash Raj Pandey
Hao-Chen2337
robertomeroni
Daniel Bevenius
Kilian Hu
Chris Lee
David Friehs
JamePeng
Pascal