mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-09-15 18:13:29 +02:00
master
400
Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
|
77d554b26d |
OpenVINO: optimize stateful decode and GPU MoE inference (#28638)
* 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>
|
||
|
Aman KarkiandGitHub
|
4c9233c034 |
cuda : enable i16 and i32 for DUP (#28897)
* cuda : enable i16 and i32 for DUP * docs : update ops table for DUP on CUDA |
||
|
|
b6b003d2cb |
sycl : Fix get mem error (#28227)
* 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> |
||
|
Aldehir RojasandGitHub
|
acecd56032 |
common : implement common_schema internal representation for JSON schemas (#28736)
* common : implement common_schema types
* common : implement a json schema optimizer
* common : reduce optimizations
* common : refactor json-schema-to-grammar to use common_schema
* common : use common_trie
* common/schema : implement type/kind resolution
* cont : cleanup
* cont : remove common_chat_tool_parameters
* cont : simplify schema resolution
* cont : pass common_schema through the json-schema-to-grammar builder
* cont : cleanup
* cont : move enums under common_schema and add type enum
* cont : reduce test cases
* cont : clean up
* cont : clean up
* refactor : rename common_schema_parse to common_schema_from_json
* tests : fix gcc dangling-reference warning in test-json-schema
* tests : take the schema label as const char * to satisfy gcc dangling-reference
* refactor : rename common_schema_builder parse_* methods to build_*
* cont : fix may_be_string
* cont : properly handle empty tool parameters
* cont : add tests for empty $ref
* cont : remove dead code
* cont : update docs
* cont : make "{}" mean any object for json_object as well
* cont : restore (min|max)Length to imply string type
* cont : rename common_schema to common_chat_schema
|
||
|
|
eafe15a5e3 |
hexagon: support for multi-device model split (aka row-split) (#28589)
* 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> |
||
|
|
d3146f2b56 |
ggml-webgpu: Update to a recent version of Dawn (#28683)
* 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> |
||
|
Daniel BeveniusandGitHub
|
3bcfeb700f |
cmake : add PCH and unity build to improve build times (#28091)
* 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
|
||
|
Aaron TeoandGitHub
|
f1b6fbf35c |
ggml-cpu(s390x): add Q1_0 vector intrinsic support (#28606)
* ggml-cpu: add `ggml_vec_dot_q1_0_q8_0` support Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * ggml-cpu: clean up variable naming for understanding Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> * docs: update support for Q1_0 Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> --------- Signed-off-by: Aaron Teo <aaron.teo1@ibm.com> |
||
|
|
5a4d0fecae |
CUDA: replace GGML_FA_ALL_QUANTS with GGML_FA_QUANTS, more control over what is compiled (#28079)
* CUDA: add configurable FA quant combinations Assisted-by: Codex * remove all flags but , add runtime fallback with warning for uncompiled combination * Update docs/build.md Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * apply code review comments --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> |
||
|
|
3ad1ba7336 |
[Model] Support for Spark2_5ForCausalLM implementation (#27868)
* Add Spark3 Model * rename spark3 -> spark2_5 Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> Co-authored-by: dongjiang <dongjiang2010@gmail.com> |
||
|
Nick FarrellandGitHub
|
cd8cdf397d |
sycl: attribute device allocations by site (GGML_SYCL_MEMTRACE) (#27631)
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 |
||
|
Adrien GallouëtandGitHub
|
24f5bf8a41 |
ggml : remove GGML_CUDA_PEER_MAX_BATCH_SIZE (#28177)
Signed-off-by: Adrien Gallouët <angt@huggingface.co> |
||
|
Frosty40andGitHub
|
6703d7894c |
sycl: fuse rms_norm+mul+add and add+add residual chains (#27610)
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. |
||
|
Todor BoinovskiandGitHub
|
e107984bcf | ops: add Hexagon to ops.md and update main README.md (#28263) | ||
|
Neo ZhangandGitHub
|
5eec3ad017 | sycl : support limit max alloc memory within 2GB for host-pinned memory (#27559) | ||
|
Jonathan ClohessyandGitHub
|
518b76236b |
kleidiai : Update KleidiAI Documentation (#26078)
Signed-off-by: Jonathan Clohessy <Jonathan.Clohessy@arm.com> |
||
|
Neo ZhangandGitHub
|
a32af33de2 |
sycl : Enhance to get the free memory of Intel GPU (#27968)
* enhance get mem info by l0 an SYCL API * remove debug code, format the code * update SYCL.md for GGML_SYCL_GET_MEM_API |
||
|
|
d7bd3bfcad |
snapdragon: python SDK setup (Windows) (#27903)
* port setup-build.ps1 to setup_sdk.py, to facilitate installation of Hexagon and OpenCL SDKs on Windows * rename setup_sdk.py -> setup-sdk.py * flake8 fix: print() -> logger.info() --------- Co-authored-by: Kristopher Urquhart <kurquhar@qti.qualcom.com> |
||
|
|
511f9c1379 |
OpenVINO: Update OV to 2026.3.1, whisper.cpp support, Qwen3.5 on NPU, and new ops (#27843)
* OpenVINO Backend: Fuse IM2COL + MatMul convolution into OpenVINO convolution * ci:ggml-ov: Skip recurrent state rollback tests * ci:ggml-ov: Skip recurrent state rollback tests * Update OPENVINO.md * ggml-openvino : add env-var gated op support debugging * Fix ggml_rope_set_offset case * OpenVINO backend: Support Whisper.cpp * Fix code style * openvino : enable qwen35 on NPU Static shapes: - get_graph_input_shape() left the s_copy / s_copy-leaf inputs dynamic ([1,1,1,-1]) even in static mode, which propagated a dynamic slot dim through GET_ROWS into the conv/GDN state, the state reshapes and the GDN output. - With -np 1 the s_copy defrag remainder gathers zero rows; short-circuit that CPY to the untouched cache instead of emitting a degenerate Slice/Concat, and skip binding its zero-byte ggml tensor as an output (the dynamic path already did the latter, the static path wrote the full cache over a 0-byte buffer). Token-count independence: - In static mode the compiled model's token count is the prefill chunk size or 1, not the captured cgraph's. Offsets derived from the captured count were therefore wrong. Anchor the GDN state slice at the end of the packed [attn | state] output and drop the rs_src_begin runtime inputs, and make VIEWs over the GDN output / conv_input pass through so the consumer does the slicing. - CONT could not identify its token axis when the graph was captured with a single token (every trailing dim has the same stride and size 1) and baked the captured shape into the prefill model. Chunked prefill: - The last chunk is padded with fabricated tokens. Attention masks them, but the recurrent path folded them into cache_r/cache_s permanently. Add a chunk_valid_len runtime input, use it to zero g and beta for padded steps (making the recurrence an exact identity) and to end the conv snapshot window at the last valid token, and disable the recurrent-cache reset after the first chunk so earlier chunks are not wiped. - get_is_prefill() and the chunk loop bound read inp_pos->ne[0] directly, but IMROPE stacks 4 position planes, so every decode step was run through the padded prefill model and the loop ran extra out-of-bounds chunks. cache_rs_reset_idx/len now stay runtime Parameters in static mode, since can_reuse_statically() does not invalidate the cached model on ComputeParams changes. Add GGML_OPENVINO_FORCE_STATIC to exercise the static path on CPU. * Update to OpenVINO 2026.3.1 * ggml-openvino: forward NPU compilation mode parameters Add GGML_OPENVINO_NPU_COMPILE_CONFIG to the backend's cached environment so callers can configure the NPU compiler without using the generic property escape hatch. When the value is non-empty, pass it to OpenVINO as NPU_COMPILATION_MODE_PARAMS. This enables settings such as optimization-level=3 for NPU compilation while preserving the existing behavior when the variable is unset and leaving CPU and GPU configuration unchanged. Document the variable, its NPU-only scope, and the optimization-level=3 example in the OpenVINO backend runtime configuration table. * ggml-openvino : support RELU, POOL_2D, QUICK_GEGLU, and ROLL ops * reorder op table * exclude GPU/NPU failing POOL_2D case * move op type detection to compute_op_case * Relax rope supported cases * Fix pool case * Update openvino doc, gpu driver in ov docker * openvino: remove unused static remote context branch * openvino: parallelize static model build * Apply editorconfig --------- Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com> Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com> Co-authored-by: zhaixuejun1993 <xuejun.zhai@intel.com> |
||
|
Gaurav GargandGitHub
|
2bb9bddafa |
spec: Add benchmark-only synthetic speculative acceptance options (#27711)
* Add benchmark-only synthetic speculative acceptance to llama-server and llama-cli * Address review comments * Address review comments * Add some comments in the code |
||
|
Max KrasnyanskyandGitHub
|
192067b72d |
hexagon: support for multi-NPU devices (IQ9, IQ10) and fully asynchronous backend (#26501)
* hexagon: use non-host bufs by default and make the backend fully async * hex-hb: remove optional hostbuf support and fix async copy * hex-unary: relax supported unary check * hex-bufs: use same get_alignment for host bufs * snapdragon: bump android_platform to 34 * hex-rows: super hacky get/set rows for q8_0 * hex-get-rows: fix q8_0 * hex-get-rows: supprot for f16 and cleanup for q8_0 * hex-get-rows: generic macros and specialized thread funcs * hex-get-rows: add DMA pipeline, vtcm_layout and kernel params * hex-set-rows: fix q8_0 support, add dma and tracing * hex-tests: override nmse threshold for HTP of Q8_0 quants * hex-fa: add support for Q8_0 with inplace dequantizers * hex-get-rows: simplify type dispatch * hex-rows: simplify GET/SET_ROWS DMA pipeline * hex-async: add events, set/get-tensor-async and rest of the async api support * hex-repack: use slice instead of expert in repack functions * hex-cpy: update event/async-cpy logging * hex-set-rows: optimize smaller tensors * hex-geglu: fix perf regression with larger tensors * hex-get-rows: add missing header * hex-set-rows: add missing header * hex-bufs: ressurect GGML_HEXAGON_HOSTBUF but disable it by default * hexagon: do not reject ops with non-heaxon buffers * hex-get-rows: apply >=32 restriction only for q8_0 * hex-res: bump vtcm acquire timeout to 10 seconds * hex-bufs: add support for cloning buffers between sessions to speed up tensor copies * hex-async: rework event recording and batch flushing and integrate with meta backend * hex-bufs: improved handling of repacked tensors * hex-repack: handle get_tensor_2d offsets * hex-dev: add support for devices with multiple NPUs * hex-sync: add support for sync tokens to synchronize npu devices for async splits * hex-mmap: cleanup mmap calls and add a retry for robustness * hex-sync: add failsafe if sync wait gets stuck * hex-sync: use sync_seq to check for completed events * hex-sync: rotate tokens for extra robustness * hex-devs: add supprot for legacy device names for now * hex-bufs: add support for auto-cloning buffers from diff sessions * hex-fusion: simplify and optimize htp-opnode fusion handling * hex-sync: override opnode name so that it shows up in the profiles * hex-trace: update scripts to handle multiple devices * hex-sync: bump the size of the opbatch queue and number of sync tokens * hex-cpy-sync: do not explicitly flush opbatches in cpy_tensor_async and add support for cpy-dma * hex-sync: add graph-flush threshold to avoid single op batches * hex-sync: add sync_peer so that we can flush peers we depend on during cross-device ops * hex-bufs: introduce tensor->extra and shadow_bufs for repacking * hex-l2: flush tiny tensors inline * hex-sync: use explicit l2flush for sync tokens * hex-extra: track weight flags via tensor extra * hex-fence: rename sync to fence * hex-repack: proper handling of set-tensor-2d in the shadow_buf * hex-trace: remove obsolete opstage mask that we used for profiling * hex-env: remove obsolete use_hmx variable * hexagon: new unified run.py and build.py and updated docs * snapdragon: update run script to auto-escapt test-backend-op -p argument * hex-scripts: fix trailing spaces * hex-scripts: fix flake8 warnings * snapdragon: cleanup dst lib/bin dirs before copying new build * hex-ops: add support for allreduce * hex-ar: improved allreduce with dma pipeline * hex-ar: align macros * hex-ar: consistent use of fence_seq * hex-ar: add AR_SELECT env var to select ALLREDUCE kernel or fallback * hex-ar: add proper synchronize handling for ALLREDUCE * hex-opbatch: looks like we now just rely on backend.synchronise to flush the batches, no need to flush them by threshold * hex-ar: bump block size to improve dma efficiency * hex-ar: fused ALLREDUCE+ADD * hex-ar: cleaner fence buffer management * hex-ar: futher allreduce tweaking to remove race conditions * hex-ar: add simple solver and remove non-dma kernels * hex-ar: add row-broadcast to fuse with bias ADD * hex-fence: pass seq numbers via op_params * hex-ar: allow for both entry/exit seq for completing entry wait * hex-ar: align macros * hex-ar: do not refetch broadcast row * hex-fusion: move all fusion into opbatch::add_op for consistency with ALLREDUCE and things * hex-fusion: fix incorrect MUL_MAT reordering * hex-mm: make fused 2x and 3x matmuls more generic * hex-fusion: move tensor fusion tagging to graph_compute * hexagon: make sure to copy tensor->extra by value * hex-get-rows: fix offset calc with row-chunking * hex-repack: get_tensor_2d fixes for non-zero offsets * snapdragon: make profile/trace scripts more robust and donot mix stdout/stderr by default * hex-devices: use legacy device nameing by default to ease the transition * hex-devices: hardcode CDSP domain IDs for current devices for now * hex-optrace: improve multi-NPU timestamp alignment and overall handling of cycle values * hex-optrace: more robust handling of the fence events |
||
|
Pranav UttarkarandGitHub
|
bf94216469 | Implemented vulkan cross_entropy_loss and cross_entropy_loss_back (#27216) | ||
|
Xuan-Son NguyenandGitHub
|
4a08fa2970 | test: move tools/parser to tests (#27548) | ||
|
Niklas WenzelandGitHub
|
9a286ac98d | docs: improve Windows build instructions (#27381) | ||
|
Xuan-Son NguyenandGitHub
|
873e5d8e39 |
model: use ggml_rope_set_offset() (#27382)
* model: use ggml_rope_set_offset() * partially apply to deepseek2 |
||
|
vkandGitHub
|
a298422da7 | docs: fix typos in ET.md (#27457) | ||
|
Oliver SimonsandGitHub
|
bf0040e15f |
CI: Use LLVM's OpenMP over MSVC_DEBUG_non_redist on Windows (#26678)
* CI: Use LLVM's OpenMP over MSFT_DEBUG_non_redist on Windows Currently, we ship the non-redist debug version of microsoft's libomp. This PR changes this to official LLVM's release, also packaging the license as needed. * Remove LLVM SHA from job name to increase legibility * Add temp validations to CI * Revert "Add temp validations to CI" This reverts commit |
||
|
|
9d77fa1725 |
ci : Update OpenVINO to 2026.3, skip nemotron-h rollback test (#27292)
* update to ov-2026.3, update device drivers * ci: skip nemotron-h rollback test on OpenVINO The OpenVINO backend does not support SSM_SCAN, so the Nemotron-H recurrent state rollback graph is split and cannot preserve the recurrent cache output shape. Keep the test enabled for other backends and retain the qwen35 OpenVINO rollback coverage. --------- Co-authored-by: ravi9 <ravi.panchumarthy@intel.com> |
||
|
|
9cd719af21 |
model: support speculators-format checkpoints for DSpark (#26275)
* dspark: support speculators-format checkpoints (SpecForge exports) Speculators-format DSpark drafts (e.g. SpecForge exports for the Gemma-4-26B-A4B target) differ from the dense DeepSpec checkpoints in three ways: - the config nests the backbone hparams under transformer_layer_config and gives the extract layers as aux_hidden_state_layer_ids - the block is the DFlash 1+N fill-in layout: the anchor slot is a bonus token, not a prediction slot. Written as dflash.bonus_anchor; such drafts build the block and read the mask positions exactly like DFlash (n_max drafts from a 1+n_max block), only the Markov/confidence sampling comes from DSpark - the draft output vocab may be reduced (draft_vocab_size < vocab_size) with a d2t remap table. The converter expands lm_head/markov_w2 back to the full vocab and synthesizes an lm_head bias of -1e9 on the rows the draft cannot produce, so the runtime needs no d2t remapping. Such drafts ship their own (now optional) token_embd/output tensors instead of sharing the target's Verified against gemma4-26b-a4b-dspark: greedy outputs are byte-identical with and without the draft; acceptance 0.46, mean draft len 3.7 (n_max 6). Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Fable 5 * dspark: fold the speculators draft class into DSparkModel One class now covers every DSpark variant. What used to pick the class is a single flag, because the arch name turns out to be the only thing that separates the two families: SpecForge also exports a flat schema that carries no speculators_* fields yet still uses the 1+N bonus-anchor block, so keying on those fields would silently mis-read its drafts. Also rename i0 to i_first_pred in the draft read loop and the Markov head, and give the head a real bonus_anchor bool instead of testing i0 > 0. Converting the Qwen3-8B DeepSpec draft and both gemma-4 speculators drafts produces byte-identical GGUFs. The one behaviour change is that the markov_head_type check now also covers the DeepSpec checkpoints, which previously skipped it. Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Opus 5 * dspark: address review comments - rename bonus_anchor to sample_from_anchor (GGUF key and code), matching the checkpoint config field; absent key still means anchor-first - rework the reduced draft vocab to match EAGLE3: d2t is written as I64 absolute target ids and the logits are scattered at runtime, instead of expanding lm_head/markov_w2 and synthesizing an output bias at conversion - move the t2d skip to modify_tensors, like EAGLE3 - drop _is_specforge: the arch name only picks the sample_from_anchor default, embed/lm_head sharing is decided by the draft vocab size - deduplicate the tok_embd create_tensor left behind by the rebase Verified with the RedHat gemma-4-31b speculator draft: greedy output is byte-identical with and without the draft; acceptance 0.26 (n_max 7). Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Fable 5 * dspark: fold the sample_from_anchor read into the block_size block * dspark: fix flake8 continuation indent * clean up * dspark: key the sample_from_anchor default off the export format Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Fable * dspark: drop t2d in filter_tensors Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Fable * dspark: map model.lm_head instead of bypassing the dflash prefix Co-authored-by: desovo7 <942845546@qq.com> Assisted-by: Claude Fable --------- Co-authored-by: desovo7 <942845546@qq.com> Co-authored-by: ruixiang63 <wangruixiang07@outlook.com> |
||
|
Georgi GerganovandGitHub
|
7c35571e5d |
ci : allow make-release to target a specific commit (#27234)
* ci : allow make-release to target a specific commit The make-release workflow now accepts an optional 'commit' input. When set, that commit is checked out and the release checks verify that it belongs to the branch selected in the "Run workflow" dialog and is not older than 3 days from the branch tip. The check is part of make-release-checks.sh (driven by the RELEASE_BRANCH env), so it follows the same dry-run semantics as the other checks. Assisted-by: pi:llama.cpp/Qwen3.8-27B * cont : scan latest 100 relase workflow runs * cont : do not check manually for release.yml success |
||
|
Xuan-Son NguyenandGitHub
|
fa88ae9368 |
convert: add @ModelBase.example (#27208)
* convert: add @ModelBase.example * add docs * add more variants * BailingMoeV3ForCausalLM * rm pocket-tts |
||
|
|
37a215c9e9 |
[SYCL] support OP OPT_STEP_ADAMW, OPT_STEP_SGD (#25268)
* fix conflict * fix conflict of ops.md * fix conflict of ops.md * update the ops.md --------- Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com> |
||
|
Fathi BoudraandGitHub
|
0177dcc730 |
common: migrate the deprecated --mmap/--no-mmap to --load-mode (#26934)
Replace the deprecated --mmap, --no-mmap, --mlock, and --direct-io flags with the unified --load-mode argument across scripts, examples, and documentation. Internal warning message and env var docs updated accordingly. Signed-off-by: Fathi Boudra <fathi.boudra@linaro.org> |
||
|
TitaniumtownandGitHub
|
6509138622 |
sycl: fuse mul_mat(gate) + mul_mat(up) + GLU for q4_K dense FFN (#26779)
Measured on Arc Pro B70 (Battlemage, Level Zero), llama-bench -r 20, two
interleaved rounds, tg128:
qwen2.5-3B-Instruct Q4_K_M 154.18 -> 158.53 t/s +2.8%
gemma-2-2b-it Q4_K_M 162.45 -> 165.62 t/s +2.0%
llama-batched-bench on qwen2.5-3B, S_TG by batch size:
B=1 142.72 -> 147.57 t/s +3.4%
B=2 243.72 -> 268.26 t/s +10.1%
B=4 359.58 -> 398.02 t/s +10.7%
B=8 449.75 -> 505.63 t/s +12.4%
|
||
|
|
aee56b3abf |
OpenVINO: Qwen3.5, memory optimization, and test-recurrent-state-rollback (#26952)
* 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> |
||
|
Neo ZhangandGitHub
|
a97123e497 |
[SYCL] Support host pinned mem to improve SYCL Host-to-Device Memory Access (#26789)
* support host pinned mem, ggml_backend_sycl_host_buffer_type_get_max_size, * fix the thread-safe issue |
||
|
Todd MalsbaryandGitHub
|
8efbf65dbd |
sycl : Add DMMV ESIMD Q3_K kernel (#26251)
* 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> |
||
|
|
8e7f22b67e |
common: add system-level config file (#26118)
* common: Add CLI > ENV > models-presets > INI precedence
1. CLI flags have the highest precedence
2. ENV vars have the second-highest precedence
3. System and User configs have the lowest precedence
- Linux/BSD/Mac
- /etc/llama.cpp/config.ini < ${XDG_CONFIG_HOME:-~/.config}/llama.cpp/config.ini
- Windows
- %PROGRAMDATA%\llama.cpp\config.ini < %APPDATA%\llama.cpp\config.ini
* fix UB
* use common_get_env
* ignore_unknown_keys
* nits
* add docs
---------
Co-authored-by: Xuan Son Nguyen <son@huggingface.co>
|
||
|
|
680a9ae63d |
cmake : introduce semantic versioning (#26839)
* cmake : introduce semantic versioning (wip)
This commit introduces semantic versioning to llama.cpp.
* squash! cmake : introduce semantic versioning (wip)
* cmake : update test-cmake README notes [no ci]
* include libmtmd in output so show its semversioned
* ci : add make-release workflow
* ci : fix build number check in build-cmake-pkg.yml
* examples : remove trailing whitespace
* ci : abort if upstream ggml version does not exist
* ci : extract step contents into scripts
* ci : add GGML_NATIVE=OFF to ubuntu job
* examples : remove CI build information from test-cmake [no ci]
This commit removes the nightly/release information that I added
previously to keep this focused only on using building and installing
llama.cpp with cmake and being able to quickly verify changes or
troubleshoot issues.
* ci : merge scripts into single script
* remove -dev-build_number support
This commit removes the incremental build number (versioning) support
that I added. This was incorrect and we should only use the semver for
the version. Releases will be tag a nightly build and package
maintainers/managers that build from source can use the tag and it is
therefor important that the correct version is reported. So a
nightly-build will report the semver without the build number. The build
number and commit as availble via cmake and test-cmake has been updated
to include an example of using them:
```console
$ ./build.sh
[test-cmake] version: 0.1.0, build: 10360 (
|
||
|
|
dd1ea52433 |
llama : support multi-output backend sampling (#25532)
* Enable backend sampling with token speculation * Clamp the mask sum before converting it into the sampled index * Add a numeric context parameter declaring the maximum outputs one sequence * More fixes * Don't reuse memory for output views. * Match dist between CPU and GPU * Fix CPU and backend sampling mismatches * Simpify some of the changes * Fix tests on Vulkan * More test fixes * Rebase changes * Rebase and address review comments * Address review comments * Address review comments * Update src/llama-sampler.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> |
||
|
Neo ZhangandGitHub
|
31558dbb76 |
sycl : Support DSv4 OPs: LIGHTNING_INDEXER,DSV4_HC_COMB,DSV4_HC_POST,DSV4_HC_PRE (#26568)
* support DSv4 OPs: LIGHTNING_INDEXER,DSV4_HC_COMB,DSV4_HC_POST,DSV4_HC_PREwq * update ops.md * fix format issue |
||
|
Neo ZhangandGitHub
|
c1f4109898 | sycl : update guide Q&A and script for device setting (#26442) | ||
|
|
0713275082 |
mtmd: support Qwen3-TTS (note: breaking change to llama-tts binary) (#26254)
* convert text model * main model load ok * convert encoder ok * speaker encoder loading ok * speaker enc graph * adapt vocab for backbone (with some tricks) * add suppress_tokens * poc new mtmd gen api * convert code_predictor to gguf * load gen_code model ok * add clip_encode * wire up * code gen cgraph init version Co-authored-by: Pascal <admin@serveurperso.com> * code2wav convert to gguf * code2wav graph ok * wire up in/out * (wip) subgraph * wire up * wip, correct code2wav * demo (to be removed) * code2wav preserve kv between calls * demo voice clone * llama: add llama_model_get_tok_embd * mtmd_helper_gen_audio API * fix clamp cold prefix Co-authored-by: Pascal <admin@serveurperso.com> * fuse snake op Co-authored-by: Pascal <admin@serveurperso.com> * demo: use proper sampling * update dev docs * polymorphism helper * revamp llama-tts binary * update docs * fix compile * fix lint * nits * add guide + docs * more timings info * clean up code comments * security fixes * update docs * use ggml_build_forward_select, clean up comments * fix ci * use ISO 639-1 language code * rename CODE2WAV --> GEN_WAV, update docs * clean up * clean up tts.cpp * add seq_id * add step_prompt() * mtmd_helper_model_can_chat * clean up comments --------- Co-authored-by: Pascal <admin@serveurperso.com> |
||
|
Pranav UttarkarandGitHub
|
f26efa02a7 |
vulkan backend ops: implemented GATED_LINEAR_ATTN (#25601)
* vulkan : add GATED_LINEAR_ATTN op * docs : update Vulkan ops * vulkan : remove unused GLA spec constant * Updated ops.md * ops.md update |
||
|
|
9d9a6d29f6 |
SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt proc… (#25025)
* SYCL: add oneMKL GEMM flash attention for XMX-accelerated prompt processing * fattn-mkl: fix interleaved dst layout in normalize kernel - Fix mkl_fa_normalize_head: use interleaved dst layout ((query * n_q_heads + head) * DV) matching TILE's flash_attn_combine_results. Previously used dense head-major layout which wrote head outputs to wrong addresses, corrupting attention for all models except Qwen3.6-27B (where GQA=6 heads were sparse enough to avoid visible overlap). - Remove 7 redundant stream->wait() calls — SYCL in-order queue already serializes pure SYCL kernel dependencies. Retain only the 4 MKL GEMM ↔ SYCL handshake barriers (oneMKL GEMM uses its own internal queue that does not respect SYCL in-order). - Remove unused dst_row_stride, diagnostic clutter, and dead K/V hex dump (fa_diag block in fattn-mkl.cpp). - Add MKL_FA_DISABLE=1 env var for A/B testing. - Add FA-DISP watchdog (MKL_FA_DEBUG=1) and FA-DIAG output fingerprint (MKL_FA_DIAG=1) in fattn.cpp. Tested: Gemma-4-26B, Gemma-4-31B, Qwen3.6-27B, Qwen3.6-35B-A3B Perf (B70/Battlemage, 32K, q8_0 KV): Gemma-4-26B: 1473 t/s MKL vs 746 TILE (1.97x) Qwen3.6-27B: 609 t/s MKL vs 330 TILE (1.85x) Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback: rename env vars, use GGML_LOG_INFO, document in SYCL.md Completed the following: - Rename MKL_FA_DISABLE → GGML_SYCL_ENABLE_MKL_FA (inverted: 0 to disable) - Rename MKL_FA_DEBUG → GGML_SYCL_MKL_FA_DEBUG - Rename MKL_FA_DIAG → GGML_SYCL_MKL_FA_DIAG - Replace fprintf(stderr, ...) / fflush(stderr) with GGML_LOG_INFO() macro - Document all three env vars in docs/backend/SYCL.md under Runtime - Add comment explaining MKL FA activation trigger (flash-attn + quantized KV cache + batch-size >= 1024 + n_kv >= 1024) Resolves review feedback from arthw. Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * Thank you for the review feedback round 2: use ggml_sycl_get_env, remove dup waits, gate perf macros - Replace raw getenv() with ggml_sycl_get_env() in all 4 env-var checks (fattn.cpp: GGML_SYCL_ENABLE_MKL_FA, GGML_SYCL_MKL_FA_DEBUG, GGML_SYCL_MKL_FA_DIAG; fattn-mkl.cpp: GGML_SYCL_MKL_FA_DEBUG) - Remove duplicated stream->wait() before ev.wait_and_throw() in GEMM KQ and GEMM VKQ — ev.wait_and_throw() already waits for completion - Gate MKL_ACCUM macro behind do_print so timing accumulators are no-ops in normal operation - Remove redundant MIT/Intel copyright header from fattn-mkl.cpp - Remove unused #include <cfloat> - Expand SYCL.md MKL FA docs with step-by-step activation trigger and example llama-cli command Again, thank you!!! Co-Authored-By: Claude Code on DeepSeek-v4-Pro * fattn-mkl: enable MKL FA for all KV cache types Remove the quantized-only restriction on MKL activation — the MKL kernel converts any non-F16 K/V to F16 via to_fp16_sycl before GEMM, so F16 (default), BF16, and F32 caches all benefit from XMX hardware acceleration. The type restriction was an unnecessary gate. Before (F16/BF16 default cache + FA on at 32K prefill): ~356 t/s (TILE path) After: ~670 t/s (MKL path, matching quantized-cache baseline) Minimal change: two conditions removed, one comment updated in fattn.cpp. No kernel or conversion code changes — the dequant pipeline already covers all types. * fattn-mkl: rename mkl_disable -> mkl_enable for clarity * fattn-mkl: refine MKL FA dispatch gates Three changes: 1. Remove quantized-only restriction - MKL FA activates for all KV cache types (F16 default, BF16, F32, quantized). The MKL kernel converts non-F16 K/V via to_fp16_sycl before GEMM. 2. Rename mkl_disable -> mkl_enable to match env var (GGML_SYCL_ENABLE_MKL_FA). 3. Replace batch-size threshold with Q->ne[1] >= 32 gate. Keeps TG (Q=1) and MTP drafts (Q=3-8) on VEC path where fused kernel beats MKL launch overhead. Routes all multi-token prefill through XMX-accelerated GEMM. Production data confirms Q patterns: 1-8 TG, 32-127 cache reuse, 128+ full reprocess. At 32K F16/BF16 FA-on: 356 -> 670 t/s. * ggml-sycl: fix F16 cache + MKL FA multi-turn corruption; add gate guards Two changes: 1. Always copy F16 K/V to dense row-major buffers before MKL GEMM. Previously F16 was read in-place with raw tensor strides. During multi-turn conversations, the accumulated KV cache had different stride properties than a fresh prefill, producing corrupted outputs. Now dense F16 gets a fast memcpy; interleaved (Gemma) gets a strided copy kernel. This matches what the quantized paths already did through to_fp16_sycl. 2. Gate MKL FA on unsupported op params (max_bias, logit_softcap, batch dim mismatch) and pathological F16 strides (nb[1] not a multiple of ne[0]*2). These conditions would previously crash inside the MKL kernel. Pathological strides (test-only) and ALiBi/softcap fall through to TILE/VEC which handle them correctly. The stride check uses modulo rather than equality, so both dense (nb1 == ne0*2) and interleaved (nb1 == H * ne0*2) pass — all real models use these layouts. Only test cases with overlapping rows (nb1=32 or nb1=75 for ne0=40) are blocked. Thanks to hmscider for the oneDNN FA PR (#25222) which surfaced the same insight: always normalize inputs to contiguous F16 before GEMM. Co-Authored-By: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> * fattn-mkl: fix quant+GQA KV strides, tighten MKL gate, add K>=1024 tests Adding K>=1024 flash-attn test cases surfaced several MKL bugs: - Quant K/V with a padded seq-view (real KV cache) used the wrong strides in the dequant path... only the true Gemma interleave layout should reconstruct strides. nb[2] vs ne[1]*nb[1] - Gate was firing on shapes the kernel doesn't handle: head_dim < 64 or not a multiple of 64, MHA, attention sinks, and bf16 decode... fell through to vec which no bf16 case. Gate MKL to the validated envelope: gqa>=2, head_dim 64 through 512 (has to be a multiple of 64) with matching K/V head size, mask, no sinks/alibi/softcap... everything else falls back to tile. Covers Qwen Dense/MoE and Gemma4 Dense/MoE Ran test-backend-ops -o FLASH_ATTN_EXT: 3641/3641 pass. Perplexity unchanged... 6.7267 MKL vs 6.7290 stock using Qwen 27b q5_k_xl * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * fattn-mkl: bound attention scratch so it doesn't grow with batch or context... also dropped the bf16 comment in fattn.cpp per arthw review. * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * Update ggml/src/ggml-sycl/fattn-mkl.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * apply arthw suggestions: enum for dequant modes, macro for wg_size, env-var one-liners --------- Co-authored-by: Claude Code using DeepSeek-V4-Pro <noreply@anthropic.com> Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> |
||
|
|
1c5b89ff63 |
sycl : support dev2dev memcpy by DEV2DEV_MEMCPY_FORWARD (#26234)
Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com> |
||
|
Georgi GerganovandGitHub
|
6b36c23056 |
readme : refresh (#26280)
* docs : center badges and links, remove Hot topics - Use <div align="center"> for GitHub-compatible centering - Add dev branches and compile times links - Remove Hot topics section Assisted-by: llama.cpp:Qwen3.6-27B * readme : remove sections * docs : center badges, remove Hot topics, extract sections, remove tools - Use <div align="center"> for GitHub-compatible centering - Add dev branches and compile times links - Add lib llama API and llama-server REST API links - Remove Hot topics section - Remove Recent API changes section - Extract XCFramework section into docs/xcframework.md - Extract Completions section into docs/completions.md - Extract Obtaining and quantizing models into docs/models.md - Remove tools usage sections (llama-cli, llama-server, etc.) - Move Contributing section to the end Assisted-by: llama.cpp:Qwen3.6-27B * cont : arrange links * cont : fix ws * cont : remove seminal papers * cont : change sample model * cont : trim-down contributing section * cont : sort backends alphabetically * cont : words * cont : add fig captions * docs : models words * readme : shorter caption * cont : fix typo * cont : add window frame to screenshot |
||
|
|
ad77bd31a6 |
docs: Adapt conda-forge package name (#26229)
Co-authored-by: dev-tinker <dev-tinker@users.noreply.github.com> |
||
|
|
84075273c8 |
spec: add DSpark speculative decoding (#25173)
* spec: add DSpark speculative decoding DSpark (DeepSpec, 2026) on top of the merged DFlash drafter. It reuses the DFlash encoder/decoder graph, target feature extraction and KV-cache injection, and the verify/accept path unchanged; the draft model is a new "dspark" arch adding a low-rank Markov head (markov_w1/w2) and an optional (unused here) confidence head. No new public APIs. The proposal is the only change: the block is anchor-first (position 0 already predicts the first draft) and the decoder graph applies a semi-autoregressive, previous-token conditioned logit bias in-graph, chained per block position: logits'(i) = logits(i) + markov_w2 . markov_w1[prev(i)] prev(0) = the block's anchor token, prev(i>0) = argmax(logits'(i-1)) vectorized across all blocks in the batch; the anchors are fed through a dedicated graph input (token 0 of every block). Greedy stays lossless (verify unchanged, same as DFlash). - new arch "dspark" (llama_model_dspark : llama_model_dflash, reuses the graph, loads the markov/confidence tensors; shares the target's embed/lm_head). - Qwen3DSparkModel converter. - new spec type "draft-dspark" (common_speculative_impl_draft_dspark : common_speculative_impl_draft_dflash, overrides draft() only: submits whole anchor-first blocks and greedily reads back the biased logits). * spec: read draft block size in the dflash impl * docs: add DSpark section to speculative.md * spec: keep dspark block size read in the dspark impl * dspark : add TODOs for incomplete parts - confidence head is loaded but not used yet - confidence-scheduled prefix pruning is not implemented - the in-graph Markov chain is greedy-only - only Qwen3 backbones are supported for now (also noted in docs) * spec: fold DSpark into the DFlash arch Address review: drop LLM_ARCH_DSPARK and the dspark.block_size / markov_rank GGUF keys. A DSpark draft now converts to a DFlash GGUF; the Markov head tensors are detected by presence (like eagle3 d2t), block_size is read from the existing dflash.block_size key, and the block anchors are taken as a strided view of the decoder's token input instead of a separate graph input. * spec: add confidence-based draft pruning for DSpark The DSpark confidence head predicts per-position acceptance of the drafted block. --spec-draft-conf-min truncates the block at the first position below the threshold (default 0 = disabled). * fold the dspark impl into dflash, selected by spec type * address review comments * dspark: clean up and improve naming * update readme * remove trailing whitespace * dflash: draft full n_max blocks, defer dp.n_max to the central truncation The DSpark markov head views the draft batch as a uniform [n_seqs x block] grid, but the per-seq dp.n_max clamp could produce blocks of different sizes, silently corrupting the strided views and the resulting logits. Drop the clamp and always draft the full n_max block for every sequence: dp.n_max is already enforced by the central truncation in common_speculative_draft(), the same way eagle3 handles it. Co-authored-by: Zaire404 <3147879462@qq.com> * dflash: assert the markov head block-uniformity invariant, require the conf head With the draft batch always submitting equal-size n_max blocks, a non-divisible token count can only mean the batch was split across ubatches or a caller broke the layout - fail loudly instead of silently dropping the markov bias. The block_drafts > block_size early return stays: worst-case graph reserve passes legitimately build with n_seq_tokens > block_size. Also make conf_proj required when the markov head is present: the confidence head is part of the DSpark checkpoint format, and a missing head would otherwise leave --spec-draft-conf-min silently reading stale embeddings instead of confidences. Co-authored-by: Zaire404 <3147879462@qq.com> * dspark: fold conf_min into p_min p_min and conf_min express the same thing - the minimum predicted survival probability for a drafted position - differing only in how the estimate is obtained: token probability for regular drafters, the trained confidence head for DSpark. The DSpark readback never used p_min, so reuse it for the confidence threshold and drop the separate --spec-draft-conf-min flag. Both defaulted to 0 (disabled), so behavior is unchanged. Co-authored-by: Zaire404 <3147879462@qq.com> * dflash: note the confidence broadcast workaround Requested in review: the ggml_repeat only adapts the [1, n_tok] confidences to the n_embd-wide embd_nextn transport so that llama_get_embeddings_nextn can be reused - not a placeholder. Co-authored-by: Zaire404 <3147879462@qq.com> * cont : clarify [no ci] --------- Co-authored-by: Ruixiang Wang <wangruixiang07@outlook.com> Co-authored-by: Zaire404 <3147879462@qq.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> |
||
|
|
d6b61ac0d3 |
sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path (#25880)
* sycl: fix use-after-return of the SDPA scale in the oneDNN flash-attention path The scale was uploaded with an async memcpy sourced from a stack local. On the in-order queue that copy is ordered behind the K/V staging kernels; once n_kv is large enough (>= ~26k observed on Arc Pro B70) the staging outlives the host stack frame and the copy reads recycled memory, feeding the SDPA a garbage scale. Output then collapses to a single repeated token and the KV cache is poisoned for the rest of the session. Short contexts win the race by accident, and test-backend-ops caps FLASH_ATTN_EXT at kv=1024, which is why CI never caught it. The previous device_count > 1 wait_and_throw() gate (and reverting it, PR #25741) fixes the symptom only by keeping the frame alive across the copy at the cost of a host sync on every FA call. Fix: cache one device scalar per (device, value) -- the scale is constant per model -- and upload it synchronously once. The single-device fast path (no per-call host sync) is then safe: every device-side hazard already serializes on the in-order queue. The multi-GPU conservative wait is kept unchanged. Also: - GGML_SYCL_FA_ONEDNN_MAX_KV env (0 = unlimited): optional n_kv ceiling that routes very long sequences to the native FA kernel. - test-backend-ops: FLASH_ATTN_EXT F16 cases up to kv=65536 (Qwen3.6-27B geometry hsk=hsv=256 GQA 6, and hsk=128 GQA 4), closing the kv=1024 blind spot. Note the race itself needs a live multi-op pipeline to reproduce; single-op runs pass even on broken builds. Verified on Arc Pro B70 (bmg_g31), Qwen3.6-27B Q4_K, -c 131072: output byte-identical at temp 0 to the native FA path through 32k-deep prefill, with prefill depth-flat at 820-840 t/s (vs 340-350 native at 32k depth). Assisted-by: Claude Fable 5 * sycl: handle GGML_SYCL_FA_ONEDNN_MAX_KV like the other runtime env vars and document it Review feedback on #25880: - read the variable once at backend init into g_ggml_sycl_fa_onednn_max_kv via ggml_sycl_get_env, and print it in the startup env listing (-lv 4 shows it) - document GGML_SYCL_FA_ONEDNN and GGML_SYCL_FA_ONEDNN_MAX_KV in the SYCL.md runtime table Also trim the added FLASH_ATTN_EXT cases to kv={4096,16384}: the 32768/65536 shapes exceed the legacy NMSE threshold on both the oneDNN and native kernels (long-sequence fp16 accumulation drift, present before this PR) and would fail CI for an unrelated reason. Assisted-by: Claude Fable 5 * sycl: clarify GGML_SYCL_FA_ONEDNN_MAX_KV default is disabled Assisted-by: Claude Fable 5 * sycl: state default behavior of GGML_SYCL_FA_ONEDNN_MAX_KV explicitly Assisted-by: Claude Fable 5 * Update ggml/src/ggml-sycl/fattn-onednn.cpp Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> * sycl: write the SDPA scale from a kernel instead of caching it The per-(device, value) scale cache was a function-local static unordered_map with no synchronization, so concurrent backend instances could access and rehash it at the same time. Write the scalar with a single_task instead. The value is captured into the command, so no host memory has to outlive the call -- which is what the use-after-return fix needed in the first place. That removes the shared container, the leaked device allocation and the string key, and it also closes the remaining async-memcpy-from-a-stack-local on the first flash-attention call. Ordering does not rely on timing: the queue is created with sycl::property::queue::in_order and the dnnl stream wraps that same queue, so the write completes before the SDPA reads the scalar. The multi-GPU wait_and_throw() branch is unchanged. Also drop the <cstdlib> include, which is unused. Assisted-by: Claude Opus 5 --------- Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com> |