Eval bug: Multi-GPU SYCL decode crash with false OOM if output tensor is on device with non-zero index #28765

Name and Version

0.4.0-dev 41fc758
built with Clang 24.0.0 for Linux x86_64

Operating systems

Linux

GGML backends

SYCL

Hardware

Intel Arc Pro B70 Graphics
Intel UHD Graphics 770

Models

any model, e.g. Llama-3.2-1B-Instruct-Q4_K_M

Problem description & steps to reproduce

SYCL backend crashes with OOM error at the first decode of any model only if output tensor of the model is located on SYCL device having index different than 0.
There has to be more than one SYCL device visible.

Steps to reproduce

llama-bench -m <model 1B> -fa 1 --device SYCL1 -ngl -1

Backtrace is like

#4  ggml_sycl_error (func="ggml_backend_sycl_get_tensor_async", line=5793)
#5  ggml_backend_sycl_get_tensor_async (...) at ggml/src/ggml-sycl/ggml-sycl.cpp:5792
#6  llama_context::decode (...) at src/llama-context.cpp:1884
#7  llama_decode (...)

Root cause analysis

The function ggml_backend_sycl_host_malloc added in #26789 hardcodes the device index to 0.

static void * ggml_backend_sycl_host_malloc(size_t size) {
    void * ptr = nullptr;
    try {
        ggml_check_sycl();
        // USM host memory is page-locked and device-accessible by construction
        auto & q = dpct::dev_mgr::instance().get_device(0).default_queue();
        ptr = sycl::malloc_host(size, q, sycl::property_list{});

Therefore llama always pins buffers to device 0 and then tries to copy using different device's queue. Memory belongs to a different context (in my environment also different platforms - this may be a required condition too). The pointer's value is unknown to the context of the queue that performs the copy.

Unfortunately error value is UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY. This is misleading - invalid context or invalid pointer should be used (a separate issue in compute-runtime).

Precise call stack:

1. llama-context.cpp asks for buffer from output tensor's device auto * output_dev = model.dev_output(); auto * output_dev_host_buft = output_dev ? ggml_backend_dev_host_buffer_type(output_dev) : nullptr;
2. ggml_backend_sycl_device_get_host_buffer_type ignores device GGML_UNUSED(dev)
3. allocation pins memory host on device 0 auto & q = dpct::dev_mgr::instance().get_device(0).default_queue(); ptr = sycl::malloc_host(size, q, sycl::property_list{});

Analysis correctness cross check
The issue is not related to which GPU card is the second one: ONEAPI_DEVICE_SELECTOR=level_zero:1,0 swaps result (card previously working crashes, card previously crashing - works, just index matters).
The issue is not related to memory size. Model crashes were observed at 1B, 7B and 3B sizes (my card has 32GB).
Not related to driver version - checked on NEO 26.22/oneAPI/icpx and NEO 26.36/clang - same behavior.

Related issues

1. OOM report from @arthw in sycl: remove unused flash attention kv buffers #27689 is caused by this issue
2. Eval bug: SYCL multi-GPU crash with Intel Arc Pro B50 + Arc A770 #27888 the same symptoms and environment (two devices, OOM error), possibly the same reason
3. Misc. bug: SYCL new host-pinned memory high CPU utilization for large allocations #27038 may be related, as code path referred there is also the path from [SYCL] Support host pinned mem to improve SYCL Host-to-Device Memory Access #26789

Workaround
Set GGML_SYCL_ENABLE_HOST_PINNED_MEM=0, that is use normal memory from heap.

Fix
I have a working fix and will open PR soon. Fix is based on using the device passed to get_host_buffer_type instead of 0.

First Bad Commit

a97123e
(confirmed it works on parent 2606220)

Relevant log output

level_zero backend failed with error: 39 (UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY)
Exception caught at file:ggml/src/ggml-sycl/ggml-sycl.cpp, line:5793, func:operator()
SYCL error: CHECK_TRY_ERROR((stream)->memcpy( data, (const char *)tensor->data + offset, size)): Exception caught in this line of code.
  in function ggml_backend_sycl_get_tensor_async at ggml/src/ggml-sycl/ggml-sycl.cpp:5793
ggml/src/ggml-sycl/common.hpp:145: SYCL error