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

docs/backend/SYCL.md

@@ -796,6 +796,7 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
 | 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.|
+| GGML_SYCL_GET_MEM_API | 0 (default) or 1  | Set to get memory info (free, total) by Level Zero or SYCL API:<br>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.<br>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_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. |

ggml/src/ggml-sycl/base.hpp

@@ -0,0 +1,36 @@
+#ifndef GGML_SYCL_BASE_HPP
+#define GGML_SYCL_BASE_HPP
+
+/**
+ * Module: base
+ *
+ * Description:
+ * Provides zero-dependency, foundational primitives, core abstractions,
+ * and low-level system interfaces. This module acts as the lowest layer
+ * of the architecture and is consumed globally across all subsystems.
+ *
+ * Constraints:
+ * - STRICTLY zero upstream dependencies (leaf module).
+ * - High stability and backward compatibility required.
+ */
+
+#include <cstdio>
+
+extern int g_ggml_sycl_debug;
+
+#if defined(__clang__) && __has_builtin(__builtin_expect)
+// Hint the optimizer to pipeline the more likely following instruction in branches
+#    define LIKELY(expr)   __builtin_expect(expr, true)
+#    define UNLIKELY(expr) __builtin_expect(expr, false)
+#else
+#    define LIKELY(expr)   (expr)
+#    define UNLIKELY(expr) (expr)
+#endif
+
+#define GGML_SYCL_DEBUG(...)              \
+    do {                                  \
+        if (UNLIKELY(g_ggml_sycl_debug))  \
+            fprintf(stderr, __VA_ARGS__); \
+    } while (0)
+
+#endif  // GGML_SYCL_BASE_HPP

ggml/src/ggml-sycl/common.hpp

@@ -18,6 +18,7 @@
 #include <iostream>
 #include <string>
 
+#include "base.hpp"
 #include "dpct/helper.hpp"
 #include "ggml.h"
 #include "ggml-impl.h"
@@ -69,21 +70,6 @@ extern int g_ggml_sycl_fa_onednn;
 extern int g_ggml_sycl_fa_onednn_max_kv;
 
 
-#if defined(__clang__) && __has_builtin(__builtin_expect)
-// Hint the optimizer to pipeline the more likely following instruction in branches
-#    define LIKELY(expr)   __builtin_expect(expr, true)
-#    define UNLIKELY(expr) __builtin_expect(expr, false)
-#else
-#    define LIKELY(expr)   (expr)
-#    define UNLIKELY(expr) (expr)
-#endif
-
-#define GGML_SYCL_DEBUG(...)              \
-    do {                                  \
-        if (UNLIKELY(g_ggml_sycl_debug))  \
-            fprintf(stderr, __VA_ARGS__); \
-    } while (0)
-
 #define CHECK_TRY_ERROR(expr)                                            \
   [&]() {                                                                \
     try {                                                                \

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -35,6 +35,7 @@
 #include <sycl/backend.hpp>
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
 #include <level_zero/ze_api.h>
+#include <level_zero/zes_api.h>
 #endif
 #if defined(GGML_SYCL_GRAPH) && SYCL_EXT_ONEAPI_ASYNC_MEMORY_ALLOC
 #    include <sycl/ext/oneapi/experimental/async_alloc/async_alloc.hpp>
@@ -61,6 +62,7 @@
 #include "ggml-sycl/fwht.hpp"
 #include "ggml-sycl/gemm.hpp"
 #include "ggml-sycl/getrows.hpp"
+#include "ggml-sycl/mem.hpp"
 #include "ggml-sycl/norm.hpp"
 #include "ggml-sycl/presets.hpp"
 #include "ggml-sycl/quantize.hpp"
@@ -105,6 +107,8 @@ int g_ggml_sycl_enable_flash_attention = 1;
 int g_ggml_sycl_dev2dev_memcpy = DEV2DEV_MEMCPY_SYCL;
 int g_ggml_sycl_usm_system = 0;
 int g_ggml_sycl_enable_host_pinned_mem = 1;
+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 = {};
@@ -301,10 +305,27 @@ static const char* dev2dev_int2str(int dev2dev) {
     }
 }
 
+/*
+* There are several entry APIs to be called as first function in SYCL backend in different cases.
+* 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.
+*/
+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<int>(zes_init)
+                            << ". Sysman free-memory query may be unavailable.\n";
+    }
+#endif
+}
+
 static void ggml_check_sycl() try {
     static bool initialized = false;
 
     if (!initialized) {
+        initialize_sycl_begining();
+
         g_ggml_sycl_debug = ggml_sycl_get_env("GGML_SYCL_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);
@@ -317,8 +338,11 @@ static void ggml_check_sycl() try {
         g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
 
         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;
         }
 
 #ifdef SYCL_FLASH_ATTN
@@ -331,6 +355,7 @@ static void ggml_check_sycl() try {
         g_ggml_sycl_enable_host_pinned_mem =
             ggml_sycl_get_env("GGML_SYCL_ENABLE_HOST_PINNED_MEM", 1);
 
+
         GGML_SYCL_DEBUG("[SYCL] call ggml_check_sycl\n");
 
         GGML_LOG_INFO("Build with Macros:\n");
@@ -374,9 +399,12 @@ static void ggml_check_sycl() try {
 
 #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));
+        GGML_LOG_INFO("  GGML_SYCL_GET_MEM_API: %d (%s)\n", g_ggml_sycl_get_mem_api, mem_api_int2str(g_ggml_sycl_get_mem_api));
 #else
         GGML_LOG_INFO("  GGML_SYCL_DEV2DEV_MEMCPY: %d (%s), enable to SYCL API since missing GGML_SYCL_SUPPORT_LEVEL_ZERO_API\n",
                       g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy));
+        GGML_LOG_INFO("  GGML_SYCL_GET_MEM_API: %d (%s), enable to SYCL API since missing GGML_SYCL_SUPPORT_LEVEL_ZERO_API\n",
+                      g_ggml_sycl_get_mem_api, mem_api_int2str(g_ggml_sycl_get_mem_api));
 #endif
 
 #if defined(GGML_SYCL_DNNL)
@@ -5208,6 +5236,7 @@ catch (sycl::exception const &exc) {
 
 static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tensor * dst) try {
     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);
@@ -5590,17 +5619,15 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
-void ggml_backend_sycl_get_device_memory(int device, size_t *free,
+void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try {
-                                                   size_t *total) try {
     GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n");
-    ggml_sycl_set_device(device);
+    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total,
-
+                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    SYCL_CHECK(CHECK_TRY_ERROR(
+    if (!res) {
-        dpct::dev_mgr::instance().get_device(device).get_memory_info(*free, *total)));
+        GGML_ABORT("[%s] failed to get device memory size", __func__);
     }
-catch (sycl::exception const &exc) {
+} catch (const sycl::exception & exc) {
-  std::cerr << exc.what() << "Exception caught at file:" << __FILE__
+    std::cerr << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl;
-            << ", line:" << __LINE__ << std::endl;
     std::exit(1);
 }
 
@@ -6021,9 +6048,11 @@ 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;
-    ggml_sycl_set_device(ctx->device);
+    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total,
-    SYCL_CHECK(CHECK_TRY_ERROR(
+                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    dpct::dev_mgr::instance().get_device(ctx->device).get_memory_info(*free, *total)));
+    if (!res) {
+        GGML_ABORT("[%s] failed to get device memory size", __func__);
+    }
 }
 
 static enum ggml_backend_dev_type ggml_backend_sycl_device_get_type(ggml_backend_dev_t dev) {
@@ -6906,6 +6935,7 @@ ggml_backend_reg_t ggml_backend_sycl_reg() {
         static std::mutex mutex;
         std::lock_guard<std::mutex> lock(mutex);
         if (!initialized) {
+            initialize_sycl_begining();
             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;
 

ggml/src/ggml-sycl/mem.cpp

@@ -0,0 +1,162 @@
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+#include <sycl/sycl.hpp>
+
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+#include <level_zero/ze_api.h>
+#include <level_zero/zes_api.h>
+#endif
+
+#include "base.hpp"
+#include "mem.hpp"
+
+#include <cstdint>
+#include <iostream>
+#include <vector>
+
+const char * mem_api_int2str(int mem_api) {
+    if (mem_api == MEMORY_API_TYPE_SYCL) {
+        return "SYCL API";
+    } else if (mem_api == MEMORY_API_TYPE_LEVEL_ZERO) {
+        return "Level Zero API";
+    } else {
+        return "Unknown";
+    }
+}
+
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+    free_bytes  = 0;
+    total_bytes = 0;
+
+    uint32_t module_count = 0;
+
+#if defined(SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO)
+    constexpr sycl::backend kL0Backend = sycl::backend::ext_oneapi_level_zero;
+#else
+    constexpr sycl::backend kL0Backend = sycl::backend::level_zero;
+#endif
+
+    try {
+        ze_result_t zes_init = zesInit(0);
+        if (zes_init != ZE_RESULT_SUCCESS) {
+            std::cerr << "Warning: zesInit failed with code " << static_cast<int>(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<sycl::info::device::global_mem_size>();
+            free_bytes  = total_bytes;
+            return false;
+        }
+
+        ze_device_handle_t ze_dev = sycl::get_native<kL0Backend>(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<sycl::info::device::global_mem_size>();
+            free_bytes  = total_bytes;
+            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<sycl::info::device::global_mem_size>();
+            free_bytes  = total_bytes;
+            return false;
+        }
+
+        std::vector<zes_mem_handle_t> 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<sycl::info::device::global_mem_size>();
+            free_bytes  = total_bytes;
+            return false;
+        }
+
+        for (uint32_t i = 0; i < module_count; ++i) {
+            zes_mem_state_t state = {};
+            state.stype           = ZES_STRUCTURE_TYPE_MEM_STATE;
+            state.pNext           = nullptr;
+
+            r = zesMemoryGetState(modules[i], &state);
+            if (r != ZE_RESULT_SUCCESS) {
+                continue;
+            }
+
+            free_bytes += state.free;
+            total_bytes += state.size;
+        }
+
+        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<sycl::info::device::global_mem_size>();
+            free_bytes  = total_bytes;
+            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<sycl::info::device::global_mem_size>();
+        free_bytes  = total_bytes;
+        return false;
+    }
+}
+#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__);
+    total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
+
+#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.");
+            free_bytes = dev.get_info<sycl::ext::intel::info::device::free_memory>();
+            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;
+            return false;
+        }
+    } else {
+        GGML_SYCL_DEBUG(
+            "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory.");
+        free_bytes = total_bytes;
+    }
+#else
+    GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory.");
+    free_bytes = total_bytes;
+#endif
+    return true;
+}
+
+bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes, MemoryAPIType api_type) {
+    const auto name       = dev.get_info<sycl::info::device::name>();
+    const auto vendor     = dev.get_info<sycl::info::device::vendor>();
+    const auto global_mem = dev.get_info<sycl::info::device::global_mem_size>();
+
+    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<size_t>(global_mem));
+
+    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);
+        }
+        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;
+#endif
+    } else {  //MEMORY_API_TYPE_SYCL
+        return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes);
+    }
+}

ggml/src/ggml-sycl/mem.hpp

@@ -0,0 +1,16 @@
+#ifndef GGML_SYCL_MEM_HPP
+#define GGML_SYCL_MEM_HPP
+
+#include <sycl/sycl.hpp>
+
+enum MemoryAPIType {
+    MEMORY_API_TYPE_LEVEL_ZERO = 0,
+    MEMORY_API_TYPE_SYCL = 1,
+};
+
+const char* mem_api_int2str(int mem_api);
+
+bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes,
+    MemoryAPIType api_type);
+
+#endif  // GGML_SYCL_MEM_HPP