* 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>
208 lines
7.8 KiB
CMake
208 lines
7.8 KiB
CMake
message(STATUS "GGML_SYCL_TARGET=${GGML_SYCL_TARGET}")
|
|
|
|
if (NOT GGML_SYCL_TARGET MATCHES "^(INTEL)$")
|
|
message(FATAL_ERROR "GGML_SYCL_TARGET: Invalid target, the supported options are [INTEL]")
|
|
endif()
|
|
|
|
check_cxx_compiler_flag("-fsycl" SUPPORTS_SYCL)
|
|
|
|
if (DEFINED ENV{ONEAPI_ROOT})
|
|
message(STATUS "Using oneAPI Release SYCL compiler (icpx).")
|
|
elseif(SUPPORTS_SYCL)
|
|
message(WARNING "Using open-source SYCL compiler (clang++). Didn't detect ENV {ONEAPI_ROOT}.
|
|
If you expected the oneAPI Release compiler, please install oneAPI & source it, like:
|
|
source /opt/intel/oneapi/setvars.sh")
|
|
else()
|
|
message(FATAL_ERROR "C++ compiler lacks SYCL support.")
|
|
endif()
|
|
message(STATUS "SYCL found")
|
|
#todo: AOT
|
|
|
|
ggml_add_backend_library(ggml-sycl
|
|
ggml-sycl.cpp
|
|
../../include/ggml-sycl.h
|
|
)
|
|
|
|
file(GLOB GGML_HEADERS_SYCL "*.hpp")
|
|
file(GLOB GGML_SOURCES_SYCL "*.cpp")
|
|
file(GLOB SRCS "template-instances/fattn-tile*.cpp")
|
|
list(APPEND GGML_SOURCES_SYCL ${SRCS})
|
|
file(GLOB SRCS "template-instances/fattn-vec*.cpp")
|
|
list(APPEND GGML_SOURCES_SYCL ${SRCS})
|
|
|
|
target_sources(ggml-sycl PRIVATE ${GGML_HEADERS_SYCL} ${GGML_SOURCES_SYCL})
|
|
|
|
if (WIN32)
|
|
# To generate a Visual Studio solution, using Intel C++ Compiler for ggml-sycl is mandatory
|
|
if( ${CMAKE_GENERATOR} MATCHES "Visual Studio" AND NOT (${CMAKE_GENERATOR_TOOLSET} MATCHES "Intel C"))
|
|
set_target_properties(ggml-sycl PROPERTIES VS_PLATFORM_TOOLSET "Intel C++ Compiler 2025")
|
|
set(CMAKE_CXX_COMPILER "icx")
|
|
set(CMAKE_CXX_COMPILER_ID "IntelLLVM")
|
|
endif()
|
|
# Level Zero SDK path for Windows (only when GGML_SYCL_SUPPORT_LEVEL_ZERO is enabled)
|
|
if(GGML_SYCL_SUPPORT_LEVEL_ZERO)
|
|
if(DEFINED ENV{LEVEL_ZERO_V1_SDK_PATH})
|
|
set(LEVEL_ZERO_V1_SDK_PATH $ENV{LEVEL_ZERO_V1_SDK_PATH})
|
|
if(EXISTS "${LEVEL_ZERO_V1_SDK_PATH}")
|
|
target_include_directories(ggml-sycl PRIVATE "${LEVEL_ZERO_V1_SDK_PATH}/include")
|
|
set(LEVEL_ZERO_V1_SDK_LIB_PATH "${LEVEL_ZERO_V1_SDK_PATH}/lib")
|
|
else()
|
|
message(WARNING "LEVEL_ZERO_V1_SDK_PATH set but folder not found: ${LEVEL_ZERO_V1_SDK_PATH}")
|
|
endif()
|
|
endif()
|
|
endif()
|
|
endif()
|
|
|
|
macro(detect_and_find_package package_name)
|
|
set(test_source "
|
|
cmake_minimum_required(VERSION ${CMAKE_VERSION})
|
|
project(check_package LANGUAGES CXX)
|
|
find_package(${package_name} QUIET)
|
|
")
|
|
|
|
set(test_dir "${CMAKE_CURRENT_BINARY_DIR}/check_package_${package_name}")
|
|
file(WRITE "${test_dir}/CMakeLists.txt" "${test_source}")
|
|
|
|
set(cmake_args "")
|
|
if(CMAKE_GENERATOR)
|
|
list(APPEND cmake_args "-G" "${CMAKE_GENERATOR}")
|
|
endif()
|
|
if(CMAKE_GENERATOR_PLATFORM)
|
|
list(APPEND cmake_args "-A" "${CMAKE_GENERATOR_PLATFORM}")
|
|
endif()
|
|
if(CMAKE_GENERATOR_TOOLSET)
|
|
list(APPEND cmake_args "-T" "${CMAKE_GENERATOR_TOOLSET}")
|
|
endif()
|
|
if(CMAKE_CXX_COMPILER)
|
|
list(APPEND cmake_args "-DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}")
|
|
endif()
|
|
|
|
execute_process(
|
|
COMMAND ${CMAKE_COMMAND} ${cmake_args} .
|
|
WORKING_DIRECTORY "${test_dir}"
|
|
RESULT_VARIABLE result
|
|
OUTPUT_QUIET
|
|
ERROR_QUIET
|
|
)
|
|
|
|
if(result EQUAL 0)
|
|
find_package(${package_name} ${ARGN})
|
|
else()
|
|
message(WARNING "Detection of ${package_name} failed. The package might be broken or incompatible.")
|
|
set(${package_name}_FOUND FALSE)
|
|
endif()
|
|
endmacro()
|
|
|
|
detect_and_find_package(IntelSYCL)
|
|
if (IntelSYCL_FOUND)
|
|
# Use oneAPI CMake when possible
|
|
target_link_libraries(ggml-sycl PRIVATE IntelSYCL::SYCL_CXX)
|
|
else()
|
|
# Fallback to the simplest way of enabling SYCL when using intel/llvm nightly for instance
|
|
target_compile_options(ggml-sycl PRIVATE "-fsycl")
|
|
target_link_options(ggml-sycl PRIVATE "-fsycl")
|
|
endif()
|
|
|
|
target_compile_options(ggml-sycl PRIVATE "-Wno-narrowing")
|
|
|
|
message(STATUS "GGML_SYCL_SUPPORT_LEVEL_ZERO ${GGML_SYCL_SUPPORT_LEVEL_ZERO}")
|
|
if (GGML_SYCL_SUPPORT_LEVEL_ZERO)
|
|
# Link against Level Zero loader for direct device memory allocation.
|
|
# Avoids sycl::malloc_device triggering DMA-buf/TTM system RAM staging
|
|
# in the xe kernel driver during multi-GPU inference.
|
|
find_path(LEVEL_ZERO_INCLUDE_DIR level_zero/ze_api.h HINTS ${ONEAPI_ROOT}/include ${LEVEL_ZERO_V1_SDK_PATH}/include)
|
|
find_library(ZE_LOADER_LIB ze_loader HINTS ${ONEAPI_ROOT}/lib ${LEVEL_ZERO_V1_SDK_LIB_PATH} ENV LD_LIBRARY_PATH)
|
|
if(ZE_LOADER_LIB AND LEVEL_ZERO_INCLUDE_DIR)
|
|
target_link_libraries(ggml-sycl PRIVATE ${ZE_LOADER_LIB})
|
|
target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_SUPPORT_LEVEL_ZERO)
|
|
message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}")
|
|
message(STATUS "Level Zero headers found: ${LEVEL_ZERO_INCLUDE_DIR}")
|
|
else()
|
|
message(WARNING "Level Zero loader or headers not found, Level Zero support disabled")
|
|
endif()
|
|
endif()
|
|
|
|
# Link against oneDNN
|
|
set(GGML_SYCL_DNNL 0)
|
|
if(GGML_SYCL_DNN)
|
|
find_package(DNNL)
|
|
if(DNNL_FOUND)
|
|
if (NOT DEFINED DNNL_GPU_VENDOR)
|
|
# default to intel target
|
|
set(DNNL_GPU_VENDOR "INTEL")
|
|
if(NOT "${GGML_SYCL_TARGET}" STREQUAL "INTEL")
|
|
message(WARNING "oneDNN builds bundled with oneapi release only support INTEL target")
|
|
endif()
|
|
endif()
|
|
|
|
# Verify oneDNN was compiled for the same target as llama
|
|
if("${GGML_SYCL_TARGET}" STREQUAL "${DNNL_GPU_VENDOR}")
|
|
target_link_libraries(ggml-sycl PRIVATE DNNL::dnnl)
|
|
set(GGML_SYCL_DNNL 1)
|
|
get_target_property(CONFIGS DNNL::dnnl IMPORTED_CONFIGURATIONS)
|
|
foreach(CONFIG ${CONFIGS})
|
|
get_target_property(DNNL_LIB DNNL::dnnl IMPORTED_LOCATION_${CONFIG})
|
|
message(STATUS "Found oneDNN: ${DNNL_LIB}")
|
|
endforeach()
|
|
else()
|
|
message(WARNING
|
|
"oneDNN must be compiled for the same target as llama.cpp.
|
|
llama.cpp: ${GGML_SYCL_TARGET}, oneDNN: ${DNNL_GPU_VENDOR}.
|
|
Disabling oneDNN support.")
|
|
endif()
|
|
else()
|
|
message(STATUS "oneDNN not found, disabling oneDNN support")
|
|
endif()
|
|
else()
|
|
message(STATUS "oneDNN support disabled by the user")
|
|
endif()
|
|
target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_DNNL=${GGML_SYCL_DNNL})
|
|
|
|
if (GGML_SYCL_F16)
|
|
add_compile_definitions(GGML_SYCL_F16)
|
|
endif()
|
|
|
|
if (GGML_SYCL_TARGET STREQUAL "INTEL")
|
|
add_compile_definitions(GGML_SYCL_WARP_SIZE=16)
|
|
if (NOT GGML_SYCL_DEVICE_ARCH)
|
|
target_link_options(ggml-sycl PRIVATE -Xs -ze-intel-greater-than-4GB-buffer-required)
|
|
else()
|
|
message(STATUS "Skipping -ze-intel-greater-than-4GB-buffer-required for spir64_gen AOT")
|
|
endif()
|
|
|
|
# Link against Intel oneMKL
|
|
if (CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
|
|
set(SYCL_COMPILER ON)
|
|
endif()
|
|
find_package(MKL REQUIRED)
|
|
target_link_libraries(ggml-sycl PRIVATE MKL::MKL_SYCL::BLAS)
|
|
else()
|
|
# default for other target
|
|
message(FATAL_ERROR "GGML_SYCL_TARGET is not supported")
|
|
add_compile_definitions(GGML_SYCL_WARP_SIZE=32)
|
|
endif()
|
|
|
|
if (GGML_SYCL_GRAPH)
|
|
message(STATUS "find GGML_SYCL_GRAPH")
|
|
target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_GRAPH)
|
|
endif()
|
|
|
|
if (GGML_SYCL_HOST_MEM_FALLBACK)
|
|
message(STATUS "find GGML_SYCL_HOST_MEM_FALLBACK")
|
|
target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_HOST_MEM_FALLBACK)
|
|
endif()
|
|
|
|
if (GGML_SYCL_DEVICE_ARCH)
|
|
message(STATUS "GGML_SYCL_DEVICE_ARCH=${GGML_SYCL_DEVICE_ARCH} (AOT via spir64_gen)")
|
|
target_compile_options(
|
|
ggml-sycl PRIVATE
|
|
-fsycl-targets=spir64_gen
|
|
"SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\""
|
|
)
|
|
target_link_options(
|
|
ggml-sycl PRIVATE
|
|
-fsycl-targets=spir64_gen
|
|
"SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\""
|
|
)
|
|
endif()
|