* rename GGML_SYCL_SUPPORT_LEVEL_ZERO to GGML_SYCL_SUPPORT_LEVEL_ZERO_API, and GGML_SYCL_ENABLE_LEVEL_ZERO to GGML_SYCL_USE_LEVEL_ZERO_API
* fix code format
* fix error when rebase
* add dev2dev memcpy by SYCL API
* mv GGML_SYCL_DEV2DEV_MEMCPY to runntime table
* update the detect method for p2p comm
* fix the erro created during fix confilct
---------
Co-authored-by: Neo Zhang <NA>
* Add interface is_model_splitted() to check the c-graph is splited or not
* Infer and propagate dynamic-dimension indices for all tensors in the GGML graph in api compute_model_outputs()
* Only do this for fallback sub graph
* Move dynamic dims compute in graph missmatch
* ggml-openvino: fix tensor data handling for PERMUTE/VIEW ops in split models
* ggml-openvino:add comments
* ggml-openvino: override VIEW op_case to 0 for split model inputs
* openvino backend: Handle unsupported VIEW shape-mismatch in OpenVINO backend
* Enable additional mul_mat tests and add tensor data saving function (#81)
* ggml-openvino: fix CONT/TRANSPOSE mapping and improve dynamic-dimension handling
* OpenVINO: add NORM/TANH support and rework SOFT_MAX translation
* ggml-openvino: extend VIEW handling
* Enable -fa off (#118)
* Enable --context-shift
* Fix llm param compute error for normal softmax not the softmax in attention
* OpenVINO backend: fix error for attention size compute in llm param
* use tensor->extra in infer_request i/o
* OpenVINO backend: refacter the compute_llm_params() func add get_attention_pattern_case to easy extand
* OpenVINO backend: clean unused code
* 1to1 match op update (#146)
* added translate_1to1_match_1_input function and updated gelu and tanh translations
* Remove unused translation function calls
---------
Co-authored-by: Mustafa Cavus <mustafacavus@intel.com>
* initial gemma4 support
* removed hardcoded names for kv cache slicing
* OpenVINO backend: Add new attention pattern for llm parameters compute
* flash attn Q shape static conversion
* Remove slice in permute translation when n_seq is 1
* return optional in extract_layer_from_name
* OpenVINO backend: refactor VIEW related operation (#148)
* OpenVINO backend: refactor VIEW related operation
* Enable VIEW handling in following ops
* OpenVINO backend does not support GGML_OP_NORM & GGML_OP_L2_NORM with VIEW input accuracy issue from OpenVINO
* OpenVINO backend: Add ops l2_norm & pad
* OpenVINO backend does not support CPY with non-contiguous data or mismatched types
* add op SSM_CONV GATED_DELTA_NET
* OpenVINO backend: fix error for bf16 in OV gpu plugin
* reverted static Q input shape for attention layer
* OpenVINO backend: remove hardcode name inp_tokens, which ignore some leaf case
* Disable remote tensor due to bug in ov gpu
* Disable n_token > 1 GATED_DELTA_NET on gpu
* OpenVINO backend: fix the view op dynamic handling issue in gemma4 & enable view + get_row
* OpenVINO backend: clean code
* OpenVINO backend: enable view + norm/rms_norm
* OpenVINO backend: concat op
* OpenVINO backend: argsort op
* OpenVINO backend: enable unary + view & GGML_UNARY_OP_SOFTPLUS
* Fix issue for test-backend-ops in TOPK_MOE, which compare VIEW ops result, VIEW node in OpenVINO no need compare, the whole graph result is correct
* OpenVINO backend: enable sum_rows
* OpenVINO backend: enable clamp
* OpenVINO backend: enable DIV
* OpenVINO backend: enable GGML_OP_MUL_MAT_ID
* OpenVINO backend: disable MUL_MAT_ID_FUSION case with large mem needed
* OpenVINO backend: Disable GGML_OP_ARGSORT, cause test_backend-ops failed
* OpenVINO backend: fix issue in mul_mat_id
* OpenVINO backend: Disable DIV with broadcast on GPU
* OpenVINO backend: update DIV
* use ov internal op GatedDeltaNet
* OpenVINO backend: enable llama erch test qwen3next
* OpenVINO backend: enable RMS_NORM + VIEW & remove op_case 2 for rope
* OpenVINO backend: fix error
* suggested changes, need review
* suggested changes, need review
* OpenVINO backend: clean unused code & fix build warning
* OpenVINO backend: enable minicpm3 for arch test
* Disable GDN op (#177)
* disable gated_delta_net
* update stateful_kv_size correctly in mismatch case
* OpenVINO backend: enable arch test for qwen3vl
* OpenVINO backend: enable cohere2 for arch test
* OpenVINO backend: enable t5 for arch test
* OpenVINO backend: enable jamba for arch test
* OpenVINO backend: remove warning for tmp
* OpenVINO backend: enable kimi-linear for arch test
* Remove unused
* Fix gpt-oss accuracy issue
* OpenVINO backend: enable arctic for arch test
* OpenVINO backend: enable grok for arch test
* Gemma4 initial npu support (#179)
* Initiall gemma4 npu support
* temp. fix for gemma4 accuracy bug on npu
* Remove hardcoded names for npu-fold handling
* revert static n tokens for cont translation as it is not needed
* removed unused variable
* ggml-openvino: add GGML_OPENVINO_ENABLE_CACHE env var to control decoder cache. Add environment variable GGML_OPENVINO_ENABLE_CACHE (default: YES). When set to NO, the decoder_cache is bypassed and models are rebuilt from the cgraph on every inference call in both dynamic and static compute paths. This is useful for debugging and verifying correctness without caching interference.
* Revert "Gemma4 initial npu support (#179)"
This reverts commit 0d29a9c4a52dc2c8aa52990f1a3854cfb01768ad.
* OpenVINO backend: disable debug log print
* Update TBB discovery. Delegated to OpenVINOs own config.
* OpenVINO backend: GGML_OPENVINO_ENABLE_CACHE YES -> 1
* OpenVINO backend: fallback FLASH_ATTN_EXT in gemma3n to CPU backend
* Add raw ov infer profiling metric
* Add OV raw infer time metric to static compute path
Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com>
* Modify precision of static profiling
* update to OV 2026.2, add OV windows CI
* fix editorconfig-checks
* Initiall gemma4 npu support
* temp. fix for gemma4 accuracy bug on npu
* Remove hardcoded names for npu-fold handling
* revert static n tokens for cont translation as it is not needed
* removed unused variable
* test-llama-archs fix
* Fix gemma4 flash_attn fallback
* support im2col
* fix code style
* disable add_rope_sin_cos optimization
* stateless boradcast and rope optimizations
* Enable manual gqa attn by default for stateless gpu
* manual gqa: fixed static batch
* gemma4 llama-bench ctx update fix
* Update OV win CI
* stateful rope fusion temp. fix
* OpenVINO backend: Conslolidate supported ops
* Exclude unsupported GGML_OP_SUB cases
* Exclude unsupported TOPK_MOE cases
* OpenVINO Backend: MUL_MAT enhancements
* Update OV CI
* support f16 mask input for npu
* Make GGML_OPENVINO_* env vars usage uniform
Standardize all GGML_OPENVINO_* env flags:
positive integers >0 to enable. Unset, empty, =0, or non-numeric values to disable.
This fixes cases where text values or empty strings enabled features.
* OpenVINO backend: Enhance envvar handling
* more cleanup
* move ggml_openvino_env_flag to appropriate place
* OpenVINO backend: add REPEAT translator, Q5_1 weights, and GLU view-input fix
* ggml-openvino: fix -Werror=cast-qual in extract_q5_1_data
* Update openvino.Dockerfile
Use BuildKit cache mounts for faster Docker rebuilds.
Use apt instead of dpkg, remove unused .ddeb downloads, add DLLAMA_BUILD_TESTS=OFF.
* ggml-openvino: centralize env var access via *getenv_str/getenv_int helpers
Replace getenv and legacy flags with _str and _int helpers.Minor cleanup, doc updates.
* OpenVINO backend: Enable GGML_OP_ADD_ID
* Uptade openvino backend clamg-format
* clang-format
* Update OPENVINO.md (#211)
* OpenVINO backend: fix accuracy issue for op CONCAT with i64 precision
* Remove strict concurrency for gpu-openvino-low-perf
* Update openvino CI keynames; add ccache-clear
* Apply suggestions from code review
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
* Fix formatting
---------
Co-authored-by: Xuejun Zhai <Xuejun.Zhai@intel.com>
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: Mustafa Cavus <mustafacavus@intel.com>
Co-authored-by: Xuejun <XuejunZhai@intel.com>
Co-authored-by: Wang Yang <yang4.wang@intel.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com>
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com>
Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>
This introduces an optional feature to allocate large GPU buffers (≥ 1GB)
using USM system allocations if supported by the device. It allows using
buffers from the system allocator then letting the system manage memory
migrations between host and device as necessary.
This feature is disabled by default and requires the GGML_SYCL_USM_SYSTEM
environment variable to enable. If USM system allocations are not supported
by the device or the system, we fallback to regular allocations.
This feature can allow VRAM overcommit. For example, the test below fails
on B580 due to lack of memory for allocation, but it passes when enabling
USM system allocations:
./examples/sycl/test.sh -m Qwen3.5-27B-Q3_K_M.gguf -lv 4
Signed-off-by: Francois Dugast <francois.dugast@intel.com>
* vulkan: support non-contig unary/glu ops
Change unary/glu ops to pass in all strides and use fastdiv for the index
calculation. Put all unary ops in one file, similar to glu, to share the
code. codex went ahead and added expm1 without me asking, but I had to
make it do a real precision analysis rather than just making stuff up.
unary.comp initially couldn't use generic_unary_head because there wasn't
space for xielu's additional constants. Fixing this required packing the
fastdiv 'L' values.
* attempt to workaround compiler bug
* resolve conflict from #23991
* use expm1
* Tidy up SYCL doc a bit
- Add explicit links to referenced items
- Fix spelling errors
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Correct documented default for GGML_SYCL_GRAPH
The default is ON, not OFF:
$ cmake -LAH -B build | grep GGML_SYCL_GRAPH
...
GGML_SYCL_GRAPH:BOOL=ON
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Move docker instructions from SYCL.md to docker.md
This makes them directly accesible from the Quick Start section
of the top-level README.md.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Refer to intel.Dockerfile for ARGs and their defaults
The defaults are always changing; this avoids accuracy errors
from duplicating the information.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* Remove mention of Nvidia in SYCL row of backend table
This support was removed in 2026.02 - refer to the SYCL.md News.
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
---------
Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
* add to support Q1_0, NVFP4, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ1_S, IQ1_M, IQ3_S, IQ4_NL, IQ4_XS, I32, MXFP4, Q2_K, Q3_K, Q5_K, and Q6_K in GET_ROWS OP
* correct the link
* fix(action): update SpacemiT toolchain URL and version
Change-Id: If4cc1c738a855274103f8c3ad52daa33528acd0c
* fix(action): add -L flag to curl command for URL redirection
Change-Id: I9b6c37390f0c7a733a36308c8fb53d22d234ab06
* snapdragon: update compiler flags to enable all CPU features
* snapdragon: update readme to point to toolchain v0.6
* snapdragon: bump toolchain docker to v0.6
* SYCL: fix multi-GPU system RAM exhaustion by using Level Zero allocations
Replace sycl::malloc_device with zeMemAllocDevice for GPU memory allocation
in the SYCL backend. sycl::malloc_device triggers the xe kernel driver's
DMA-buf/TTM path which mirrors every VRAM allocation 1:1 in system RAM.
zeMemAllocDevice uses the SVM/P2P path with no host staging.
On a dual Intel Arc Pro B70 system (64GB VRAM, 64GB RAM), a 15.6 GiB model
consumed 60 GiB of system RAM via sycl::malloc_device, causing OOM crashes.
With zeMemAllocDevice, the same workload uses ~6.7 GiB of system RAM with
no performance regression.
All Level Zero calls include automatic fallback to the original SYCL
allocation path if Level Zero interop is unavailable.
* SYCL: address review feedback - remove try/catch, check device types, deduplicate
- Remove try/catch from malloc/free/memcpy helpers, check backend and
device type upfront instead (ggml_sycl_is_level_zero, ggml_sycl_is_dgpu)
- Move shared helpers (is_level_zero, is_dgpu, free_device) to common.cpp
and declare in common.hpp to eliminate code duplication
- Use SYCL_CHECK(CHECK_TRY_ERROR()) for fallback sycl::free calls
- Guard dev2dev_memcpy L0 path to dGPU-to-dGPU only, preserving the
host-staged path for iGPU-to-dGPU transfers
- Add Windows Level Zero SDK path detection (LEVEL_ZERO_V1_SDK_PATH)
in CMakeLists.txt (co-authored with @arthw)
* SYCL: add build/runtime flags for Level Zero, address review feedback
Implements the architecture suggested by @arthw: compile-time and runtime
flags to cleanly separate Level Zero and SYCL memory API paths.
- Add GGML_SYCL_SUPPORT_LEVEL_ZERO cmake option (default ON). All Level
Zero code is wrapped in #ifdef so the build works on systems without
the Level Zero SDK installed (e.g. CPU-only CI servers). Both the
loader library and headers are checked before enabling.
- Add GGML_SYCL_ENABLE_LEVEL_ZERO runtime env var (default 1). Controls
whether Level Zero or SYCL memory APIs are used. Only one API style is
used per session, no mixing. If Level Zero is enabled but the devices
don't support the Level Zero backend, it auto-disables with a warning.
- Remove Level Zero code from dpct_malloc. It was unused (dpct::device_memory
is not called anywhere in the backend) and used try/catch for flow control.
- Update SYCL.md with documentation for both new parameters.
Tested on Intel Arc Pro B70 (32GB), single-GPU and dual-GPU, with both
GGML_SYCL_SUPPORT_LEVEL_ZERO=ON and OFF builds. AI-assisted development
(Claude). Code reviewed and tested on my hardware.
* SYCL: unify Level Zero malloc/free call sites, address review feedback
Move ggml_sycl_malloc_device to common.cpp alongside ggml_sycl_free_device.
Both functions are now unconditionally available — Level Zero code is
#ifdef'd inside the functions, not at call sites. All call sites use
uniform SYCL_CHECK(CHECK_TRY_ERROR()) wrapping with no #ifdef blocks.
Addresses arthw's review: wrap all malloc/free in SYCL_CHECK for stack
traces on failure, eliminate duplicated #ifdef/else patterns at 6 call
sites (-29 lines net).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: add Level Zero SDK to CI, fix device check and missed alloc paths
Add Level Zero SDK installation to Ubuntu and Windows SYCL CI jobs
so the Level Zero code path is compiled and tested in CI.
Fix two bugs found during extended dual-GPU testing (no
ONEAPI_DEVICE_SELECTOR set):
- The Level Zero backend check was iterating all SYCL devices
including CPU. The OpenCL CPU device caused Level Zero to be
disabled for the GPUs, defeating the fix on multi-GPU systems.
Added is_gpu() filter so only GPU devices are checked.
- sycl_ext_malloc_device/sycl_ext_free (tensor reorder temp buffers)
were still calling sycl::malloc/sycl::free directly, bypassing the
Level Zero path. Routed through ggml_sycl_malloc_device/free_device
for consistency with the other device memory call sites.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* SYCL: address arthw review feedback on Level Zero memory API structure
- Move ggml_sycl_malloc_device to static function in ggml-sycl.cpp;
only ggml_sycl_free_device (used by common.cpp) stays in common.cpp
- Switch both helpers to use g_ggml_sycl_enable_level_zero global
instead of per-call queue backend checks
- Remove #ifdef wrapper from global definition; always declare at 0,
add #else branch in init block so it stays 0 when L0 not compiled in
- Update init loop comment to explain GPU-only device check
- CMakeLists: message(STATUS) before the if block; align option wording
AI-assisted implementation. Reviewed and tested on dual Intel Arc Pro
B70 (32 GB each): test-backend-ops OK on both GPUs, single/dual-GPU
Q4_K_M and Q8_0 bench correct, zeMemAllocDevice GTT delta confirmed
<5 MiB per 4 GiB allocation (vs ~4 GiB shadow with sycl::malloc_device).
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
* SYCL: remove unused cstdio/cstdlib includes from common.cpp
Leftover from the deleted ggml_sycl_queue_supports_level_zero helper.
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
* Apply suggestions from code review
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* SYCL: preserve Level Zero allocation path during early malloc
* ci: fix Level Zero package conflict in Intel Docker build
* ci: find Level Zero loader in oneAPI package step
* ci: allow Windows SYCL package without Level Zero DLL
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-authored-by: Neo Zhang <zhang.jianyu@outlook.com>
* SYCL: reduce allocation overhead during flash attention
* tidy up whitespace
* add a note about the flag
* move ggml_sycl_fattn_* into fattn-buffers.hpp
* refactor implementation into fattn-buffers.cpp
* move new_fattn_kv_buffers back into ggml-sycl.cpp
2026-05-09 09:30:39 +03:00
Intel AI Get-to Market Customer Success and Solutions
* Write a readme on Multi-GPU usage in llama.cpp
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Address review comments
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* docs : update speculative decoding parameters after refactor (#22397)
Update docs/speculative.md to reflect the new parameter naming scheme
introduced in PR #22397:
- Replace --draft-max/--draft-min with --spec-draft-n-max/--spec-draft-n-min
- Replace --spec-ngram-size-n/m with per-implementation variants
- Add documentation for all new --spec-ngram-*- parameters
- Update all example commands
Assisted-by: llama.cpp:local pi
* pi : add rule to use gh CLI for GitHub resources
Assisted-by: llama.cpp:local pi
* docs : run llama-gen-docs
* arg : fix typo
* opt arc770 for Q4_0
* add for Q4_0
* update the script
* add help script for windows
* update guide
* fix format issue
* convert from dos to unix for format issue
* fix missed -sm parameter
* hexagon: restore HTP_OPMASK_QUEUE
* hexagon: honor OPMASK_SKIP_COMPUTE in hmx-matmul
* hex-prof: restore op profiling
* hex-prof: enable PMU
* hexagon: simplify and improve op-queuing with full profiling support
Add separate profile descriptors.
* hexagon: remove opsync and rename opmask into opstage
opsync is no longer needed since the profiler is fully async now.
opmask name was confusing and opstage is more accurate.
* hexagon: refactor opbatch queue handling
* hexagon: add iface hooks for enabling profiler from the host
Also move all the PMU setup stuff out of the hex-utils since it's not inteded for normal use.
* hexagon: make profiler mode configurable
On older devices getting PMU counters is expensive so it's now optional.
* hexagon: add support for setting profiler pmu events from env
* hexagon: simplify profiler output (no need to print buffs, etc)
* hexagon: simplify pmu counter formating
* hexagon: add a simple profile post-proc tool
* hex-prof: add support for reading logs from stdin
* hexagon: document GGML_HEXAGON_PROFILE
* hex-prof: update default width for dims field
* hex-prof: fix linter warnings and errors
* Update ggml/src/ggml-hexagon/htp/htp-ops.h
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Update scripts/snapdragon/ggml-hexagon-profile.py
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
---------
Co-authored-by: Trivikram Reddy <tamarnat@qti.qualcomm.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* upgrade oneAPI to 2025.3.3
* update
* seperate SYCL CI and support release binary package for ubuntu 24
* add dependence
* remove wrong copy lines
* add missed line
* remove other task to test the release for SYCL
* rm more for test release
* fix file name
* correct the error in running
* support build for fp32/fp16
* rm ubuntu-24-sycl-fp16 for duplicated
* refactor build setting
* update guide for ubuntu 24 release package, restore the release.yml for other backend
* user docker replace to install oneAPI
* use download installation package to replace docker
* use wget to download and install oneapi, replace the apt cmd
* enable ccache for oneAPI installation
* fix format error
* enable cache for oneAPI installation
* update guide
* Update .github/workflows/release.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Update .github/workflows/release.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Update .github/workflows/build-sycl.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Update .github/workflows/release.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
---------
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Thread safety per request only
* Fix ROPE yarn case
* Fix sticky stateful config
* Use i4/i8 directly for symmetric quant
* Use weightless caching
* Add WeightlessCacheAttribute to reduce NPU memory usage
* Gelu tanh support (#125)
* Imrope support (#126)
* fix(openvino): explicit ov::Tensor frees in ggml_backend_openvino_free
* add GPU,NPU support in OV Dockerfile
* add build-openvino.yml ci
* Fix sticky stateful config
* add concurrency to ov-gpu ci runs. Move OV CI to build-openvino.yml
* fix thread-safety of shared runtime context
* rope type abstraction for frontend translations
* fix editorconfig
---------
Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com>
Co-authored-by: Dan Hoffman <dhoff749@gmail.com>
Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com>
* nix: support unified apple-sdk
* Impl roll op for Metal
* Revert "nix: support unified apple-sdk"
This reverts commit abfa473360.
* update ops.md
* update op docs
Xuan-Son Nguyen