* 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>
* switch ubuntu-latest to ubuntu-slim
* Fix the path for upload so CI doesn't fail
* Update .github/workflows/build-and-test-snapdragon.yml
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Use -slim image for key check and consistent naming for artifact dir
Signed-off-by: Max Krasnyansky <maxk@qti.qualcomm.com>
* Remove check-secret extra job
* move QDC key check for Run QDC jobs step specifically
* add a step before to check the secret for qdc jobs
---------
Signed-off-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Add the tests that we want to run on external CI
* remove extra files
* Fixes python issues, reove the deadlock on CI
* remove unecessary changes
* use override to ty.toml
* fix pre-commit and try tests with secret in external repo not upstream
* skip if key is unavailable
* Fix feedback
* switch hexagon to snapdragon
* cleanup
* fix secrets
* remove the copyrights at the top of the files
* 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>
* Update workflows to remove dependence on llvmpipe
* Try setting Dawn_DIR
* remove c++20 initializers
* Move to proper guid
* Try avoiding segfaults on vulkan backend process exit
* Remove compiler warnings on parameter casting
* Fix soft_max and update reg_tile accumulation to f32 for better precision
* Refactor flash_attn a bit
* remove c++20 initializers and format
* Increase div precision for NVIDIA
* revert div precision and comment out ggml-ci node for now
* Formatting
* Try debugging on a failing CI node
* Revert "Try debugging on a failing CI node"
This reverts commit 1971e33cba.
* vulkan: Programmatically add RoundingModeRTE to all shaders when the device supports it
* use FetchContent to get SPIRV-Headers
* Fetch spirv-headers unconditionally
* remove fetchcontent, rely on installed headers
* fix ubuntu job
* Update docs/build.md
* experimenting CI
* Experimenting CI fix for MinGW
* experimenting CI on Windows
* modified script for integration with VisualStudio
* added proxy handling
* adding python version for Windows execution
* fix iterator::end() dereference
* fixed proxy handling
* Fix errors occurring on Windows
* fixed ci script
* Reverted to master
* Stripping test items to simplify Windows test
* adjusting script for windows testing
* Changed shell
* Fixed shell
* Fixed shell
* Fix CI setting
* Fix CI setting
* Fix CI setting
* Experimenting ci fix
* Experimenting ci fix
* Experimenting ci fix
* Experimenting ci fix
* experimenting fix for unit test error
* Changed to use BUILD_LOW_PERF to skip python tests
* Fix CI
* Added option to specify Ninja generator
* Reverted proxy related changes
* cann: update docker images to 8.5.0
- bump CANN base image from 8.3.rc2 to 8.5.0
- bump ASCEND_VERSION from 8.1.RC1.alpha001 to 8.5.0
Move to newer stable releases.
* cann: update CANN.md
* Update CANN.md to include BF16 support
Added BF16 support information to the CANN documentation and corrected formatting for the installation instructions.
* Fix formatting issues in CANN.md
Fix 234: Trailing whitespace
* scripts: hip: gcn-cdna-vgpr-check: fix parsing of vgpr counts when an amdclang Remark block is interlieved with another from a different process
* Return warning ignore
* obay pep8 inline double space before inline commets
* add # noqa: NP100 for other prints too
* Add script changes to cause autotrigger
* Remove make dependency
* Added option to specify Ninja generator
* use ninja-build as default for several CI
* Revert "use ninja-build as default for several CI"
This reverts commit f552c4559b.
* changed use plain string rather than arrays
* Enabled ninja build by default for experimentation
* ci: add run.sh to test conditions to trigger GitHub CI and self-hosted runners
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* Enabled ninja build by default on self-hosted envs for experimentation
* ci: revert generator to ninja instead of ninja multi-config
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: install ninja-build for self-hosted workflows
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: revert ninja from self-hosted runners
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: missed one self-hosted step
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* ci: fix windows ci errors from an errenous revert
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* Added explicit build types for Ninja
Also reverted some needless change
* ci: use ninja multi-config for vulkan-x64 build
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
* added time command to measure build time
* Keeping some configs to use Ninja which show improvement
* minor fix based on review
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
* ci: rm `time` from custom containers
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
---------
Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Aaron Teo <aaron.teo1@ibm.com>
Co-authored-by: Aaron Teo <taronaeo@gmail.com>
