From 1d3f075e23748337fb1e144f859cd2c6d2eb8e9c Mon Sep 17 00:00:00 2001 From: Neo Date: Wed, 2 Sep 2026 00:11:10 +0800 Subject: [PATCH 1/9] fix for unsupport zes API --- ggml/src/ggml-sycl/ggml-sycl.cpp | 30 ++++++++++++++++------- ggml/src/ggml-sycl/mem.cpp | 41 ++++++++++---------------------- 2 files changed, 34 insertions(+), 37 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 626f0f3bff12..e07bd022e7a3 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -110,7 +110,6 @@ int g_ggml_sycl_enable_host_pinned_mem = 1; int g_ggml_sycl_host_pinned_mem_2g = 0; int g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_LEVEL_ZERO; - static ggml_sycl_device_info ggml_sycl_init() { ggml_sycl_device_info info = {}; @@ -311,13 +310,26 @@ static const char* dev2dev_int2str(int dev2dev) { * It's the first internal function to be called by them in SYCL backend. * This function is used to do initialize work for the SYCL backend and set the global variables. */ + +static ze_result_t init_zes() { + ze_result_t res = zesInit(0); + if (res != ZE_RESULT_SUCCESS) { + std::cerr << "Warning: [" << __func__ << "] zesInit failed with code " << static_cast(res) + << ". Sysman free-memory query may be unavailable.\n"; + } + return res; +} + + +const ze_result_t get_zes_init_res() { + static ze_result_t zes_init_res = init_zes(); + return zes_init_res; +} + void initialize_sycl_begining() { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API - ze_result_t zes_init = zesInit(0); - if (zes_init != ZE_RESULT_SUCCESS) { - std::cerr << "Warning: zesInit failed [ggml_check_sycl] with code " << static_cast(zes_init) - << ". Sysman free-memory query may be unavailable.\n"; - } + //must be called in initialization stage, before any other Level Zero API calls + get_zes_init_res(); #endif } @@ -5634,8 +5646,9 @@ catch (sycl::exception const &exc) { void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try { GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n"); + printf("zjy 2 g_zes_init=%d, g_ggml_sycl_get_mem_api=%d\n", get_zes_init_res(), g_ggml_sycl_get_mem_api); bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total, - (MemoryAPIType) g_ggml_sycl_get_mem_api); + get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); if (!res) { GGML_ABORT("[%s] failed to get device memory size", __func__); } @@ -6061,8 +6074,9 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) { ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context; + printf("zjy 1 g_zes_init=%d, g_ggml_sycl_get_mem_api=%d\n", get_zes_init_res(), g_ggml_sycl_get_mem_api); bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total, - (MemoryAPIType) g_ggml_sycl_get_mem_api); + get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); if (!res) { GGML_ABORT("[%s] failed to get device memory size", __func__); } diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp index 5ec466420e02..09756108aac1 100644 --- a/ggml/src/ggml-sycl/mem.cpp +++ b/ggml/src/ggml-sycl/mem.cpp @@ -24,6 +24,9 @@ const char * mem_api_int2str(int mem_api) { } #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API +/* +* Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail. +*/ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) { free_bytes = 0; total_bytes = 0; @@ -37,41 +40,28 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot #endif try { - ze_result_t zes_init = zesInit(0); - if (zes_init != ZE_RESULT_SUCCESS) { - std::cerr << "Warning: zesInit failed with code " << static_cast(zes_init) - << ". Sysman free-memory query may be unavailable.\n"; - } if (dev.get_platform().get_backend() != kL0Backend) { - GGML_SYCL_DEBUG("Device backend is not Level Zero; falling back to SYCL memory query.\n"); - total_bytes = dev.get_info(); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("Device backend is not Level Zero.\n"); return false; } ze_device_handle_t ze_dev = sycl::get_native(dev); if (ze_dev == nullptr) { - GGML_SYCL_DEBUG("Level Zero device handle is null; falling back to SYCL memory query.\n"); - total_bytes = dev.get_info(); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("Level Zero device handle is null.\n"); return false; } ze_result_t r = zesDeviceEnumMemoryModules(ze_dev, &module_count, nullptr); if (r != ZE_RESULT_SUCCESS || module_count == 0) { - GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n"); - total_bytes = dev.get_info(); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n"); return false; } std::vector modules(module_count); r = zesDeviceEnumMemoryModules(ze_dev, &module_count, modules.data()); if (r != ZE_RESULT_SUCCESS || module_count == 0) { - GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n"); - total_bytes = dev.get_info(); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n"); return false; } @@ -90,16 +80,12 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot } if (total_bytes == 0) { - GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes. Falling back to SYCL memory query.\n"); - total_bytes = dev.get_info(); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes.\n"); return false; } return true; } catch (const sycl::exception & e) { GGML_SYCL_DEBUG("Level Zero memory query failed: %s\n", e.what()); - total_bytes = dev.get_info(); - free_bytes = total_bytes; return false; } } @@ -112,23 +98,20 @@ bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105) if (dev.has(sycl::aspect::ext_intel_free_memory)) { try { - GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory."); + GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.\n"); free_bytes = dev.get_info(); - return true; } catch (const sycl::exception &) { GGML_SYCL_DEBUG( - "Failed to query free memory using SYCL aspect::ext_intel_free_memory. Using total memory as free " - "memory."); - free_bytes = total_bytes; + "Failed to query free memory using SYCL aspect::ext_intel_free_memory.\n"); return false; } } else { GGML_SYCL_DEBUG( - "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory."); + "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory.\n"); free_bytes = total_bytes; } #else - GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory."); + GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory.\n"); free_bytes = total_bytes; #endif return true; From db36812078a1dfa862fa13e286b32fff33242fb1 Mon Sep 17 00:00:00 2001 From: Neo Date: Wed, 2 Sep 2026 08:37:09 +0800 Subject: [PATCH 2/9] optimize the code --- ggml/src/ggml-sycl/mem.cpp | 43 ++++++++++++++++++++++---------------- 1 file changed, 25 insertions(+), 18 deletions(-) diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp index 09756108aac1..ea6ce6c62e10 100644 --- a/ggml/src/ggml-sycl/mem.cpp +++ b/ggml/src/ggml-sycl/mem.cpp @@ -6,12 +6,14 @@ #include #endif -#include "base.hpp" -#include "mem.hpp" - #include #include #include +#include +#include + +#include "base.hpp" +#include "mem.hpp" const char * mem_api_int2str(int mem_api) { if (mem_api == MEMORY_API_TYPE_SYCL) { @@ -100,6 +102,7 @@ bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & try { GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.\n"); free_bytes = dev.get_info(); + return true; } catch (const sycl::exception &) { GGML_SYCL_DEBUG( "Failed to query free memory using SYCL aspect::ext_intel_free_memory.\n"); @@ -107,14 +110,12 @@ bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & } } else { GGML_SYCL_DEBUG( - "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory.\n"); - free_bytes = total_bytes; + "Device does not support SYCL aspect::ext_intel_free_memory.\n"); } #else - GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory.\n"); - free_bytes = total_bytes; + GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105.\n"); #endif - return true; + return false; } bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes, MemoryAPIType api_type) { @@ -129,17 +130,23 @@ bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes if (api_type == MEMORY_API_TYPE_LEVEL_ZERO) { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API GGML_SYCL_DEBUG("[%s]Querying free memory using Level Zero API.\n", __func__); - if (!query_free_memory_by_ze(dev, free_bytes, total_bytes)) { - //fallback to SYCL API if Level Zero API fails - GGML_SYCL_DEBUG("[%s]Falling back to SYCL API for memory query.\n", __func__); - return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes); + if (query_free_memory_by_ze(dev, free_bytes, total_bytes)) { + return true; } - return true; -#else - GGML_SYCL_DEBUG("[%s]Level Zero API support is not enabled. Please enable it to use this feature.\n", __func__); - return false; + //fallback to SYCL API if Level Zero API fails + GGML_SYCL_DEBUG("[%s]Falling back to SYCL API for memory query.\n", __func__); #endif - } else { //MEMORY_API_TYPE_SYCL - return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes); } + + //MEMORY_API_TYPE_SYCL + if(get_memory_size_by_sycl_api(dev, free_bytes, total_bytes)){ + return true; + } + + //Todo, fallback to other methods to get free memory size, such as using OS-specific APIs (e.g., /proc/meminfo on Linux, GlobalMemoryStatusEx on Windows, etc.) + GGML_SYCL_DEBUG( + "[%s] Can't get free mem size by Level Zero and SYCL API. Using total memory as free memory.\n", __func__); + free_bytes = total_bytes; + + return true; } From 381e7b5f7438b9e6b05b4306f7fdf36139fdfabf Mon Sep 17 00:00:00 2001 From: Neo Zhang Jianyu Date: Wed, 2 Sep 2026 13:38:36 +0800 Subject: [PATCH 3/9] adjust the log level --- docs/backend/SYCL.md | 5 +++-- ggml/src/ggml-sycl/base.hpp | 7 +++++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 17 ++++++++--------- ggml/src/ggml-sycl/mem.cpp | 11 +++++++---- 4 files changed, 25 insertions(+), 15 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index b30e109ad408..6b8b3c0654c6 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -790,14 +790,15 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm | Name | Value | Function | |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------| -| GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG | +| GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function: GGML_SYCL_DEBUG() for common debug. | +| GGML_SYCL_DEV_DEBUG | 0 (default) or 1 | Enable log function: GGML_SYCL_DEV_DEBUG() for development by replacing GGML_SYCL_DEBUG() in special codes. Restore to GGML_SYCL_DEBUG() before commit code.| | GGML_SYCL_DEV2DEV_MEMCPY | 0 (default), 1, 2 | Choose the method of dev2dev memory copy.
Value:
* 0: SYCL API (default), only support dGPUs.
* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case.
* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).
SYCL & L0 API are easy to be impacted by Intel GPU driver issue. When you meet the garbled output or crash issues in multiple GPUs case, try with this debug flag to work around or check the issue.| | GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.| | GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) | | GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. | | GGML_SYCL_ENABLE_HOST_PINNED_MEM | 0 or 1 (default) | Enable host pinned memory to speed up copy data from host to device. When disable it, host memory will common malloc() on CPU. Disable it when use `--load-model mlock`.| | GGML_SYCL_HOST_PINNED_MEM_2G | 0 (default) or 1 | Limit the max memory allocation to be no more than 2GB when enable host pinned memory. USM allocations above 2 GiB take the relaxed/large-allocation path, which serializes H2D copies with compute and prevents copy/compute overlap. It will impact the startup time. Need more test. Depend on `GGML_SYCL_ENABLE_HOST_PINNED_MEM=1`.| -| GGML_SYCL_GET_MEM_API | 0 (default) or 1 | Set to get memory info (free, total) by Level Zero or SYCL API:
0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.
1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return total size for free size.| +| GGML_SYCL_GET_MEM_API | 0 (default) or 1 | Set to get memory info (free, total) by Level Zero or SYCL API:
0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.
1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return the free size as value of total size.| | GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).| | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. | | GGML_SYCL_FA_ONEDNN | 1 (default) or 0 | Enable the oneDNN fused SDPA (flash-attention) path on supported GPUs. Set to 0 to always use the native SYCL flash-attention kernel. | diff --git a/ggml/src/ggml-sycl/base.hpp b/ggml/src/ggml-sycl/base.hpp index 3afd57ccb2db..fe96c4ab855e 100644 --- a/ggml/src/ggml-sycl/base.hpp +++ b/ggml/src/ggml-sycl/base.hpp @@ -17,6 +17,7 @@ #include extern int g_ggml_sycl_debug; +extern int g_ggml_sycl_dev_debug; #if defined(__clang__) && __has_builtin(__builtin_expect) // Hint the optimizer to pipeline the more likely following instruction in branches @@ -33,4 +34,10 @@ extern int g_ggml_sycl_debug; fprintf(stderr, __VA_ARGS__); \ } while (0) +#define GGML_SYCL_DEV_DEBUG(...) \ + do { \ + if (UNLIKELY(g_ggml_sycl_dev_debug)) \ + fprintf(stderr, __VA_ARGS__); \ + } while (0) + #endif // GGML_SYCL_BASE_HPP diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index e07bd022e7a3..a339bed6a0cf 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -91,6 +91,7 @@ static bool g_sycl_loaded = false; int g_ggml_sycl_debug = 0; +int g_ggml_sycl_dev_debug = 0; int g_ggml_sycl_enable_optimize = 1; int g_ggml_sycl_enable_graph = 0; int g_ggml_sycl_enable_dnn = 1; @@ -310,18 +311,16 @@ static const char* dev2dev_int2str(int dev2dev) { * It's the first internal function to be called by them in SYCL backend. * This function is used to do initialize work for the SYCL backend and set the global variables. */ - static ze_result_t init_zes() { ze_result_t res = zesInit(0); - if (res != ZE_RESULT_SUCCESS) { - std::cerr << "Warning: [" << __func__ << "] zesInit failed with code " << static_cast(res) - << ". Sysman free-memory query may be unavailable.\n"; - } + if (res != ZE_RESULT_SUCCESS) { + GGML_SYCL_DEV_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n", + __func__, (int) res); + } return res; } - -const ze_result_t get_zes_init_res() { +ze_result_t get_zes_init_res() { static ze_result_t zes_init_res = init_zes(); return zes_init_res; } @@ -340,6 +339,7 @@ static void ggml_check_sycl() try { initialize_sycl_begining(); g_ggml_sycl_debug = ggml_sycl_get_env("GGML_SYCL_DEBUG", 0); + g_ggml_sycl_dev_debug = ggml_sycl_get_env("GGML_SYCL_DEV_DEBUG", 0); g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1); g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0); g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1); @@ -411,6 +411,7 @@ static void ggml_check_sycl() try { GGML_LOG_INFO("Running with Environment Variables:\n"); GGML_LOG_INFO(" GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug); + GGML_LOG_INFO(" GGML_SYCL_DEV_DEBUG: %d\n", g_ggml_sycl_dev_debug); #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API GGML_LOG_INFO(" GGML_SYCL_DEV2DEV_MEMCPY: %d (%s)\n", g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy)); @@ -5646,7 +5647,6 @@ catch (sycl::exception const &exc) { void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try { GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n"); - printf("zjy 2 g_zes_init=%d, g_ggml_sycl_get_mem_api=%d\n", get_zes_init_res(), g_ggml_sycl_get_mem_api); bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total, get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); if (!res) { @@ -6074,7 +6074,6 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) { ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context; - printf("zjy 1 g_zes_init=%d, g_ggml_sycl_get_mem_api=%d\n", get_zes_init_res(), g_ggml_sycl_get_mem_api); bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total, get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); if (!res) { diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp index ea6ce6c62e10..71451d477b60 100644 --- a/ggml/src/ggml-sycl/mem.cpp +++ b/ggml/src/ggml-sycl/mem.cpp @@ -30,6 +30,8 @@ const char * mem_api_int2str(int mem_api) { * Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail. */ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) { + GGML_SYCL_DEBUG("[%s] Querying free memory using Level Zero API.\n", __func__); + free_bytes = 0; total_bytes = 0; @@ -85,7 +87,8 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes.\n"); return false; } - return true; + return total_bytes >= free_bytes; + } catch (const sycl::exception & e) { GGML_SYCL_DEBUG("Level Zero memory query failed: %s\n", e.what()); return false; @@ -94,7 +97,7 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot #endif bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) { - GGML_SYCL_DEBUG("[%s]Querying free memory using SYCL API.\n", __func__); + GGML_SYCL_DEBUG("[%s] Querying free memory using SYCL API.\n", __func__); total_bytes = dev.get_info(); #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105) @@ -129,12 +132,12 @@ bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes if (api_type == MEMORY_API_TYPE_LEVEL_ZERO) { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API - GGML_SYCL_DEBUG("[%s]Querying free memory using Level Zero API.\n", __func__); + GGML_SYCL_DEBUG("[%s] Querying free memory using Level Zero API.\n", __func__); if (query_free_memory_by_ze(dev, free_bytes, total_bytes)) { return true; } //fallback to SYCL API if Level Zero API fails - GGML_SYCL_DEBUG("[%s]Falling back to SYCL API for memory query.\n", __func__); + GGML_SYCL_DEBUG("[%s] Falling back to SYCL API for memory query.\n", __func__); #endif } From cbdb6d525ae9fd7e7e526fb5cec896f915470215 Mon Sep 17 00:00:00 2001 From: Neo Zhang Jianyu Date: Wed, 2 Sep 2026 13:57:09 +0800 Subject: [PATCH 4/9] rm unused head files --- ggml/src/ggml-sycl/mem.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp index 71451d477b60..19b811add994 100644 --- a/ggml/src/ggml-sycl/mem.cpp +++ b/ggml/src/ggml-sycl/mem.cpp @@ -9,8 +9,6 @@ #include #include #include -#include -#include #include "base.hpp" #include "mem.hpp" From d2f23593ffae0024619af391cf0f16681d4c4be4 Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Wed, 2 Sep 2026 15:27:50 +0800 Subject: [PATCH 5/9] Update docs/backend/SYCL.md Co-authored-by: Titaniumtown --- docs/backend/SYCL.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 6b8b3c0654c6..29b006aa2cef 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -791,7 +791,7 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm | Name | Value | Function | |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------| | GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function: GGML_SYCL_DEBUG() for common debug. | -| GGML_SYCL_DEV_DEBUG | 0 (default) or 1 | Enable log function: GGML_SYCL_DEV_DEBUG() for development by replacing GGML_SYCL_DEBUG() in special codes. Restore to GGML_SYCL_DEBUG() before commit code.| +| GGML_SYCL_DEV_DEBUG | 0 (default) or 1 | Enable log function: GGML_SYCL_DEV_DEBUG() for developmental purposes by replacing GGML_SYCL_DEBUG() in special codes. Restore to GGML_SYCL_DEBUG() before committing code.| | GGML_SYCL_DEV2DEV_MEMCPY | 0 (default), 1, 2 | Choose the method of dev2dev memory copy.
Value:
* 0: SYCL API (default), only support dGPUs.
* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case.
* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).
SYCL & L0 API are easy to be impacted by Intel GPU driver issue. When you meet the garbled output or crash issues in multiple GPUs case, try with this debug flag to work around or check the issue.| | GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.| | GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) | From 9a0e49aad452190d7f2a652704da2457e1fdecd2 Mon Sep 17 00:00:00 2001 From: "zhang.jianyu" Date: Sun, 6 Sep 2026 23:23:35 +0800 Subject: [PATCH 6/9] fix the error to detect level zero SDK/dev package, stop build after detect the error --- ggml/src/ggml-sycl/CMakeLists.txt | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index a8d9c0d804bf..d335231dd46e 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -110,15 +110,19 @@ if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API) # 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_path(LEVEL_ZERO_DEV_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) + if(ZE_LOADER_LIB AND LEVEL_ZERO_DEV_INCLUDE_DIR) target_link_libraries(ggml-sycl PRIVATE ${ZE_LOADER_LIB}) target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_SUPPORT_LEVEL_ZERO_API) message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}") - message(STATUS "Level Zero headers found: ${LEVEL_ZERO_INCLUDE_DIR}") + message(STATUS "Level Zero headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}") else() - message(WARNING "Level Zero loader or headers not found, Level Zero support disabled") + message(FATAL_ERROR "Level Zero SDK/Development package not found! " + "Please install the Level Zero SDK/development package or " + "set -DGGML_SYCL_SUPPORT_LEVEL_ZERO_API=OFF in build command. " + "Level Zero API is not mandatory for SYCL backend, " + "but it is required by the special features for better performance on Intel GPUs.") endif() endif() From 8a53e0ad2fbb502083f4cff1344eea7f77de6a59 Mon Sep 17 00:00:00 2001 From: "zhang.jianyu" Date: Sun, 6 Sep 2026 23:28:40 +0800 Subject: [PATCH 7/9] update the message --- ggml/src/ggml-sycl/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index d335231dd46e..aa5a87551a95 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -116,7 +116,7 @@ if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API) target_link_libraries(ggml-sycl PRIVATE ${ZE_LOADER_LIB}) target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_SUPPORT_LEVEL_ZERO_API) message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}") - message(STATUS "Level Zero headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}") + message(STATUS "Level Zero development headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}") else() message(FATAL_ERROR "Level Zero SDK/Development package not found! " "Please install the Level Zero SDK/development package or " From a53449ab0fe846332166a73a448a43172760ca2c Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Mon, 7 Sep 2026 11:19:39 +0800 Subject: [PATCH 8/9] fix the build error when missed to install level zero dev package --- ggml/src/ggml-sycl/CMakeLists.txt | 12 +++++++----- ggml/src/ggml-sycl/ggml-sycl.cpp | 24 ++++++++++++++++++------ 2 files changed, 25 insertions(+), 11 deletions(-) diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index aa5a87551a95..d2196f74d565 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -118,11 +118,13 @@ if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API) message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}") message(STATUS "Level Zero development headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}") else() - message(FATAL_ERROR "Level Zero SDK/Development package not found! " - "Please install the Level Zero SDK/development package or " - "set -DGGML_SYCL_SUPPORT_LEVEL_ZERO_API=OFF in build command. " - "Level Zero API is not mandatory for SYCL backend, " - "but it is required by the special features for better performance on Intel GPUs.") + message(WARNING "Level Zero loader or development headers not found, " + "Level Zero API support disabled. " + "Please install the Level Zero SDK/development package " + "to support Level Zero API features. " + "Level Zero API is not mandatory for SYCL backend, " + "but it is required by the special features for better " + "function & performance on Intel GPUs.") endif() endif() diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index a339bed6a0cf..4d771611b1ed 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -311,6 +311,7 @@ static const char* dev2dev_int2str(int dev2dev) { * It's the first internal function to be called by them in SYCL backend. * This function is used to do initialize work for the SYCL backend and set the global variables. */ +#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API static ze_result_t init_zes() { ze_result_t res = zesInit(0); if (res != ZE_RESULT_SUCCESS) { @@ -324,6 +325,7 @@ ze_result_t get_zes_init_res() { static ze_result_t zes_init_res = init_zes(); return zes_init_res; } +#endif void initialize_sycl_begining() { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API @@ -5645,11 +5647,23 @@ catch (sycl::exception const &exc) { std::exit(1); } +bool sycl_get_mem_info(int device, size_t * free, size_t * total) { + MemoryAPIType mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL; +#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API + mem_api_type = get_zes_init_res() == ZE_RESULT_SUCCESS ? + (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL; +#else + mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL; +#endif + bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), + *free, *total, mem_api_type); + GGML_SYCL_DEBUG("[SYCL] [%s] total = %zu free = %zu\n", __func__, *total, *free); + return res; +} + void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try { GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n"); - bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total, - get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); - if (!res) { + if (!sycl_get_mem_info(device, free, total)) { GGML_ABORT("[%s] failed to get device memory size", __func__); } } catch (const sycl::exception & exc) { @@ -6074,9 +6088,7 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) { ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context; - bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total, - get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL); - if (!res) { + if (!sycl_get_mem_info(ctx->device, free, total)) { GGML_ABORT("[%s] failed to get device memory size", __func__); } } From 19b054cb1c59d53a5f629b1760175965c559a5c4 Mon Sep 17 00:00:00 2001 From: Neo Zhang Date: Mon, 7 Sep 2026 18:23:51 +0800 Subject: [PATCH 9/9] rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions --- ggml/src/ggml-sycl/ggml-sycl.cpp | 35 +++++++++++++++++++++++--------- ggml/src/ggml-sycl/mem.cpp | 14 ++++++------- 2 files changed, 31 insertions(+), 18 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 4d771611b1ed..6152da169a32 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -112,6 +112,7 @@ int g_ggml_sycl_host_pinned_mem_2g = 0; int g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_LEVEL_ZERO; static ggml_sycl_device_info ggml_sycl_init() { + GGML_SYCL_DEBUG("[SYCL] call ggml_sycl_init\n"); ggml_sycl_device_info info = {}; // Do not hard crash when there exists no SYCL devices. @@ -202,12 +203,9 @@ static ggml_sycl_device_info ggml_sycl_init() { } #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API - // Large buffers can be allocated before ggml_check_sycl() initializes other - // g_ggml_sycl_enable_* globals, so initialize this one as early as we can. + //update g_ggml_sycl_use_level_zero_api according to the device support g_ggml_sycl_use_level_zero_api = - info.ext_oneapi_level_zero && ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1); -#else - g_ggml_sycl_use_level_zero_api = 0; + info.ext_oneapi_level_zero && g_ggml_sycl_use_level_zero_api; #endif return info; @@ -315,7 +313,7 @@ static const char* dev2dev_int2str(int dev2dev) { static ze_result_t init_zes() { ze_result_t res = zesInit(0); if (res != ZE_RESULT_SUCCESS) { - GGML_SYCL_DEV_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n", + GGML_SYCL_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n", __func__, (int) res); } return res; @@ -323,6 +321,7 @@ static ze_result_t init_zes() { ze_result_t get_zes_init_res() { static ze_result_t zes_init_res = init_zes(); + GGML_SYCL_DEBUG("[SYCL] call %s: zesInit result: %d\n", __func__, (int) zes_init_res); return zes_init_res; } #endif @@ -330,11 +329,13 @@ ze_result_t get_zes_init_res() { void initialize_sycl_begining() { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API //must be called in initialization stage, before any other Level Zero API calls + GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__); get_zes_init_res(); #endif } static void ggml_check_sycl() try { + GGML_SYCL_DEBUG("[SYCL] ggml_check_sycl()\n"); static bool initialized = false; if (!initialized) { @@ -352,9 +353,13 @@ static void ggml_check_sycl() try { g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1); g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0); +#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API + g_ggml_sycl_use_level_zero_api = ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1); +#else + g_ggml_sycl_use_level_zero_api = 0; +#endif g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL); g_ggml_sycl_get_mem_api = ggml_sycl_get_env("GGML_SYCL_GET_MEM_API", MEMORY_API_TYPE_LEVEL_ZERO); - if (g_ggml_sycl_use_level_zero_api == 0) { g_ggml_sycl_dev2dev_memcpy = DEV2DEV_MEMCPY_SYCL; g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_SYCL; @@ -950,6 +955,7 @@ inline void * aligned_malloc_host(size_t alignment, size_t size) { static ggml_backend_buffer_t ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) try { + GGML_SYCL_DEBUG("[SYCL] call %s: size=%zu\n", __func__, size); ggml_check_sycl(); ggml_backend_sycl_buffer_type_context * buft_ctx = (ggml_backend_sycl_buffer_type_context *)buft->context; @@ -1469,10 +1475,11 @@ static ggml_backend_buffer_type_i ggml_backend_sycl_split_buffer_type_interface }; ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split) { + GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n"); + static std::mutex mutex; std::lock_guard lock(mutex); - GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n"); ggml_check_sycl(); // FIXME: this is not thread safe static std::map, struct ggml_backend_buffer_type> buft_map; @@ -1525,6 +1532,7 @@ static const char * ggml_backend_sycl_host_buffer_type_name(ggml_backend_buffer_ //host pinned memory static void * ggml_backend_sycl_host_malloc(size_t size) { + GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_malloc\n"); void * ptr = nullptr; try { ggml_check_sycl(); @@ -5263,8 +5271,8 @@ catch (sycl::exception const &exc) { } static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tensor * dst) try { + GGML_SYCL_DEBUG("[SYCL] ggml_sycl_compute_forward: dst=%s, op=%s\n", dst->name, ggml_op_name(dst->op)); if (!g_sycl_loaded) return false; - initialize_sycl_begining(); if (dst->src[0] != nullptr && ggml_backend_buffer_is_sycl_split(dst->src[0]->buffer)) { ggml_sycl_set_peer_access(dst->src[1]->ne[1], ctx.device); @@ -5648,7 +5656,11 @@ catch (sycl::exception const &exc) { } bool sycl_get_mem_info(int device, size_t * free, size_t * total) { + GGML_SYCL_DEBUG("[SYCL] [%s] g_ggml_sycl_get_mem_api=%d\n", + __func__, g_ggml_sycl_get_mem_api); + MemoryAPIType mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL; + #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API mem_api_type = get_zes_init_res() == ZE_RESULT_SUCCESS ? (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL; @@ -6087,10 +6099,12 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t } static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) { + GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__); ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context; if (!sycl_get_mem_info(ctx->device, free, total)) { GGML_ABORT("[%s] failed to get device memory size", __func__); } + GGML_SYCL_DEBUG("[SYCL] call %s total %zu free %zu\n", __func__, *total, *free); } static enum ggml_backend_dev_type ggml_backend_sycl_device_get_type(ggml_backend_dev_t dev) { @@ -6966,6 +6980,7 @@ static const ggml_backend_reg_i ggml_backend_sycl_reg_interface = { // backend registry ggml_backend_reg_t ggml_backend_sycl_reg() { + GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_reg\n"); static ggml_backend_reg reg; static bool initialized = false; @@ -6973,7 +6988,7 @@ ggml_backend_reg_t ggml_backend_sycl_reg() { static std::mutex mutex; std::lock_guard lock(mutex); if (!initialized) { - initialize_sycl_begining(); + ggml_check_sycl(); ggml_backend_sycl_reg_context * ctx = new ggml_backend_sycl_reg_context; const int min_batch_size = getenv("GGML_OP_OFFLOAD_MIN_BATCH") ? atoi(getenv("GGML_OP_OFFLOAD_MIN_BATCH")) : 32; diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp index 19b811add994..ad5bfe0ff1eb 100644 --- a/ggml/src/ggml-sycl/mem.cpp +++ b/ggml/src/ggml-sycl/mem.cpp @@ -28,7 +28,7 @@ const char * mem_api_int2str(int mem_api) { * Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail. */ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) { - GGML_SYCL_DEBUG("[%s] Querying free memory using Level Zero API.\n", __func__); + GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__); free_bytes = 0; total_bytes = 0; @@ -95,7 +95,7 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot #endif bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) { - GGML_SYCL_DEBUG("[%s] Querying free memory using SYCL API.\n", __func__); + GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using SYCL API.\n", __func__); total_bytes = dev.get_info(); #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105) @@ -120,13 +120,11 @@ bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & } bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes, MemoryAPIType api_type) { - const auto name = dev.get_info(); - const auto vendor = dev.get_info(); - const auto global_mem = dev.get_info(); - GGML_SYCL_DEBUG("[%s]GPU Name: %s\n", __func__, name.c_str()); - GGML_SYCL_DEBUG("[%s]GPU Vendor: %s\n", __func__, vendor.c_str()); - GGML_SYCL_DEBUG("[%s]GPU Global Memory: %zu bytes\n", __func__, static_cast(global_mem)); + GGML_SYCL_DEBUG("[%s]GPU Name: %s\n", __func__, + dev.get_info().c_str()); + GGML_SYCL_DEBUG("[%s]GPU Vendor: %s\n", __func__, + dev.get_info().c_str()); if (api_type == MEMORY_API_TYPE_LEVEL_ZERO) { #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API