sycl : Fix get mem error (#28227)

* fix for unsupport zes API

* optimize the code

* adjust the log level

* rm unused head files

* Update docs/backend/SYCL.md

Co-authored-by: Titaniumtown <titaniumtown@proton.me>

* fix the error to detect level zero SDK/dev package, stop build after detect the error

* update the message

* fix the build error when missed to install level zero dev package

* rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions

---------

Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
Co-authored-by: Neo Zhang <NA>

.github/workflows/code-style.yml

@@ -15,7 +15,7 @@ concurrency:
 
 jobs:
   model-naming:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
     steps:
       - uses: actions/checkout@v6
       - name: Check model naming conventions

.github/workflows/editorconfig.yml

@@ -15,7 +15,7 @@ concurrency:
 
 jobs:
   editorconfig:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
     steps:
       - uses: actions/checkout@v6
       - uses: editorconfig-checker/action-editorconfig-checker@840e866d93b8e032123c23bac69dece044d4d84c # v2.2.0

.github/workflows/server-sanitize.yml

@@ -116,7 +116,6 @@ jobs:
         run: |
           source .venv/bin/activate
           cd tools/server/tests
-          export ${{ matrix.extra_args }}
           PYTEST_WORKERS=1 ./tests.sh
 
       - name: Slow tests
@@ -125,5 +124,4 @@ jobs:
         run: |
           source .venv/bin/activate
           cd tools/server/tests
-          export ${{ matrix.extra_args }}
           PYTEST_WORKERS=1 SLOW_TESTS=1 ./tests.sh

.pi/gg/SYSTEM.md

@@ -6,6 +6,7 @@ General:
 - PR and commit titles format: `<module> : <title>`. Lookup recents for examples
 - Don't try to build or run the code unless you are explicitly asked to do so
 - Use the `gh` CLI tool when querying PRs, issues, or other GitHub resources
+- When [MODEL] is needed, first try to get it from the `PI_MODEL_NAME` env var before asking the user
 
 Coding:
 - When in doubt, always refer to the CONTRIBUTING.md file of the project
@@ -20,7 +21,7 @@ Pull requests (PRs):
 - Don't explicitly wrap lines in the PR description (each paragraph and bullet is a single line)
 - When creating a pull request, look for the repository's PR template and follow it
 - For the AI usage disclosure section, write "YES. pi:llama.cpp/[MODEL]"
-- Ask the user to tell you what model was used and write it in place of [MODEL]
+- If `PI_MODEL_NAME` env var is not set, ask the user to tell you what model was used and write it in place of [MODEL]
 - Always create the pull requests in draft mode
 
 Commits:

ci/run.sh

@@ -775,7 +775,11 @@ function gg_run_test_backend_ops {
 
     set -e
 
-    local args_extra="-j $(nproc)"
+    local n_jobs=$(nproc)
+    if [ "${n_jobs}" -gt 2 ]; then
+        n_jobs=2
+    fi
+    local args_extra="-j ${n_jobs}"
 
     # TODO: fix multi-threaded for ROCm
     #       https://github.com/ggml-org/llama.cpp/actions/runs/34576278519/job/103297889044?pr=28740#step:3:4865
@@ -789,10 +793,11 @@ function gg_run_test_backend_ops {
         args_extra=""
     fi
 
+    # TODO: reduce the test-backend-ops timeout to 1800s
     if [ ! -z ${GG_BUILD_HIGH_PERF} ]; then
-        (time ./bin/test-backend-ops ${args_extra} -b CPU) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
+        (time timeout 3600 ./bin/test-backend-ops ${args_extra} -b CPU) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
     else
-        (time ./bin/test-backend-ops ${args_extra}       ) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
+        (time timeout 3600 ./bin/test-backend-ops ${args_extra}       ) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
     fi
 
     set +e

common/parsers/qwen3-coder.cpp

@@ -104,9 +104,34 @@ common_chat_params common_chat_params_init_qwen3_coder(const common_chat_templat
 
                     auto arg_open = p.tool_arg_open("<parameter=" + p.tool_arg_name(p.literal(param.name)) + ">\n");
 
-                    auto arg_value = param.schema->may_be_string() ?
-                        arg_string :
-                        p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
+                    auto types = param.schema->value_types();
+
+                    auto arg_value = p.eps();
+                    if (!types.has(common_chat_schema::TYPE_STRING)) {
+                        arg_value = p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
+                    } else if (types.is_only(common_chat_schema::TYPE_STRING)) {
+                        arg_value = arg_string;
+                    } else {
+                        // The string alternative accepts any text, so the grammar only keeps the raw string
+                        // rule. The parser still tries the JSON alternatives first to type the value.
+                        auto json_value = p.choice();
+                        if (types.has(common_chat_schema::TYPE_OBJECT)) {
+                            json_value |= p.json_object();
+                        }
+                        if (types.has(common_chat_schema::TYPE_ARRAY)) {
+                            json_value |= p.json_array();
+                        }
+                        if (types.has(common_chat_schema::TYPE_NUMBER) || types.has(common_chat_schema::TYPE_INTEGER)) {
+                            json_value |= p.json_number();
+                        }
+                        if (types.has(common_chat_schema::TYPE_BOOLEAN)) {
+                            json_value |= p.json_bool();
+                        }
+                        if (types.has(common_chat_schema::TYPE_NULL)) {
+                            json_value |= p.json_null();
+                        }
+                        arg_value = p.gbnf(p.atomic(p.tool_arg_json_value(json_value) + arg_close) | arg_string, "xml-arg-string");
+                    }
 
                     auto arg_rule = p.rule(rule_name, p.tool_arg(arg_open + arg_value));
 

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 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.<br>Value: <br>*  0: SYCL API (default), only support dGPUs.<br>* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case. <br>* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).<br>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:<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_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 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. |

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -8302,9 +8302,20 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c
     bool threashold_ok = tensor->ne[0] >= threshold_ne0 && tensor->ne[1] >= threshold_ne1 &&
             tensor->ne[2] == 1 && tensor->ne[3] == 1;
 
-    // q6_K adreno kernels requires ne1 is multiple of 128
-    if (tensor->type == GGML_TYPE_Q6_K) {
-        return threashold_ok && tensor->ne[1] % 128 == 0;
+    // The noshuffle layout packs 2 rows per 32-bit texel and the GEMV reads it at an
+    // ne1/2 texel stride with an exact-cover dispatch, so it is only addressable when
+    // ne1 is a multiple of 64; an unaligned ne1 truncates the stride and the weight is
+    // read misaligned. That is a property of the layout, not of one quant -- q4_K, q5_K
+    // and q8_0 read the same packing as q6_K. The bound is 64, not 128: a q8_0 attention
+    // weight of ne1 = 2880 is a multiple of 64 but not 128 and is correct.
+    switch (tensor->type) {
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
+        case GGML_TYPE_Q6_K:
+        case GGML_TYPE_Q8_0:
+            return threashold_ok && tensor->ne[1] % 64 == 0;
+        default:
+            break;
     }
     return threashold_ok;
 }

ggml/src/ggml-sycl/CMakeLists.txt

@@ -110,15 +110,21 @@ 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 development headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}")
     else()
-        message(WARNING "Level Zero loader or headers not found, Level Zero support disabled")
+        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()
 

ggml/src/ggml-sycl/base.hpp

@@ -17,6 +17,7 @@
 #include <cstdio>
 
 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

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;
@@ -113,8 +114,8 @@ 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_DEBUG("[SYCL] call ggml_sycl_init\n");
     ggml_sycl_device_info info = {};
 
     // Do not hard crash when there exists no SYCL devices.
@@ -205,12 +206,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;
@@ -314,23 +312,40 @@ 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) {
+        GGML_SYCL_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n",
+                            __func__, (int) res);
+    }
+    return res;
+}
+
+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
+
 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";
-    }
+    //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) {
         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);
@@ -344,9 +359,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;
@@ -405,6 +424,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));
@@ -945,6 +965,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;
@@ -1464,10 +1485,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<std::mutex> 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<std::array<float, GGML_SYCL_MAX_DEVICES>, struct ggml_backend_buffer_type> buft_map;
@@ -1520,6 +1542,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();
@@ -5341,8 +5364,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);
@@ -5725,11 +5748,27 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
+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;
+#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,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    if (!sycl_get_mem_info(device, free, total)) {
         GGML_ABORT("[%s] failed to get device memory size", __func__);
     }
     ggml_sycl_memtrace_report_device("device memory query", device, *free, *total);
@@ -6177,12 +6216,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;
-    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    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);
     ggml_sycl_memtrace_report_device("device memory query (dev)", ctx->device, *free, *total);
 }
 
@@ -7061,6 +7100,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;
 
@@ -7068,7 +7108,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_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;
 

ggml/src/ggml-sycl/mem.cpp

@@ -6,13 +6,13 @@
 #include <level_zero/zes_api.h>
 #endif
 
-#include "base.hpp"
-#include "mem.hpp"
-
 #include <cstdint>
 #include <iostream>
 #include <vector>
 
+#include "base.hpp"
+#include "mem.hpp"
+
 const char * mem_api_int2str(int mem_api) {
     if (mem_api == MEMORY_API_TYPE_SYCL) {
         return "SYCL API";
@@ -24,7 +24,12 @@ 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) {
+    GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__);
+
     free_bytes  = 0;
     total_bytes = 0;
 
@@ -37,41 +42,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<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;
+            GGML_SYCL_DEBUG("Device backend is not Level Zero.\n");
             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;
+            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<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
             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;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
             return false;
         }
 
@@ -90,73 +82,70 @@ 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<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            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());
-        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__);
+    GGML_SYCL_DEBUG("[SYCL] call %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.");
+            GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.\n");
             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;
+                "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.");
-        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.");
-    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) {
-    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));
+    GGML_SYCL_DEBUG("[%s]GPU Name:          %s\n", __func__,
+        dev.get_info<sycl::info::device::name>().c_str());
+    GGML_SYCL_DEBUG("[%s]GPU Vendor:        %s\n", __func__,
+        dev.get_info<sycl::info::device::vendor>().c_str());
 
     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);
+        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;
         }
-        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;
 }

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -333,6 +333,7 @@ static void ggml_vk_print_device_lost_info(const vk_device& device);
 struct vk_queue_handle {
     vk::Queue queue;
     vk_device_ref device;
+    std::mutex * device_submit_mutex = nullptr;
     virtual void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) = 0;
     virtual void lock()   {}   // no-op by default (internally synchronized case)
     virtual void unlock() {}
@@ -342,6 +343,11 @@ struct vk_queue_handle {
 struct vk_queue_handle_synchronized : vk_queue_handle {
     std::mutex mutex;
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
+        // Workaround for NVIDIA driver bug
+        std::unique_lock<std::mutex> device_guard;
+        if (device_submit_mutex) {
+            device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
+        }
         std::lock_guard<std::mutex> guard(mutex);
         try {
             queue.submit(submits, fence);
@@ -356,9 +362,14 @@ struct vk_queue_handle_synchronized : vk_queue_handle {
     void unlock() override { mutex.unlock(); }
 };
 
+// Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
 struct vk_queue_handle_unsynchronized : vk_queue_handle {
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
-        // Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
+        // Workaround for NVIDIA driver bug
+        std::unique_lock<std::mutex> device_guard;
+        if (device_submit_mutex) {
+            device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
+        }
         try {
             queue.submit(submits, fence);
         } catch (vk::DeviceLostError &) {
@@ -835,6 +846,7 @@ static bool ggml_vk_lightning_indexer_k_type_supported(ggml_type type) {
 
 struct vk_device_struct {
     std::recursive_mutex mutex;
+    std::mutex queue_submit_mutex;
     mutable std::shared_mutex pinned_memory_mutex;
 
     // Guards compile_pending, all_pipelines, and the dynamic pipeline maps
@@ -3520,6 +3532,10 @@ static std::unique_ptr<vk_queue> ggml_vk_create_queue(vk_device& device, uint32_
 
     h->queue = device->device.getQueue2(queue_info2);
     h->device = device;
+    // Avoid concurrent submissions on NVIDIA due to driver bug.
+    if (device->vendor_id == VK_VENDOR_ID_NVIDIA) {
+        h->device_submit_mutex = &device->queue_submit_mutex;
+    }
     q->handle = h;
 
     q->cmd_pool.init(device, q.get());

tests/test-backend-ops.cpp

@@ -10678,9 +10678,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_flash_attn_ext(256, 256, 2, {16, 1}, 16384,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0));
 
     // MLA shape: the V cache is a sub-view of the K cache, with quantized KV
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1},  113,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024,  64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1},  113,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1}, 1024,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1}, 1024,  64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
 
     // Sparse mask hint: supported decode/prefill layouts and dense fallbacks.
     test_cases.emplace_back(new test_flash_attn_ext(512, 512, 1, { 8, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false,  512));

tests/test-chat.cpp

@@ -846,6 +846,25 @@ static common_chat_tool nullable_int_tool{
     })",
 };
 
+static common_chat_tool string_union_tool{
+    /* .name = */ "set_union",
+    /* .description = */ "Set values whose types are unions with string",
+    /* .parameters = */ R"({
+        "type": "object",
+        "properties": {
+            "value": {
+                "type": ["string", "object"],
+                "description": "A string or object value"
+            },
+            "amount": {
+                "type": ["string", "integer"],
+                "description": "A string or integer value"
+            }
+        },
+        "required": ["value", "amount"]
+    })",
+};
+
 static common_chat_tool enum_no_type_tool{
     /* .name = */ "set_unit",
     /* .description = */ "Set a temperature unit",
@@ -3805,6 +3824,46 @@ static void test_template_output_peg_parsers(bool detailed_debug) {
             })
             .run();
 
+        // nullable string given null - parses as JSON null, not the string "null"
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_nullable_str>\n"
+               "<parameter=name>\nnull\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ nullable_string_tool })
+            .expect_tool_calls({
+                { "set_nullable_str", R"({"name": null})", {} },
+            })
+            .run();
+
+        // unions with string - JSON values of the other types are typed, everything else is a string
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_union>\n"
+               "<parameter=value>\n{\"a\": 1}\n</parameter>\n"
+               "<parameter=amount>\n2 dollars\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ string_union_tool })
+            .expect_tool_calls({
+                { "set_union", R"({"value": {"a": 1}, "amount": "2 dollars"})", {} },
+            })
+            .run();
+
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_union>\n"
+               "<parameter=value>\n{not valid json\n</parameter>\n"
+               "<parameter=amount>\n42\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ string_union_tool })
+            .expect_tool_calls({
+                { "set_union", R"({"value": "{not valid json", "amount": 42})", {} },
+            })
+            .run();
+
         // enum without explicit type key - should infer string from enum values
         tst.test(
                "<tool_call>\n"