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

快速结论：当系统存在多个 SYCL 设备、并且你把模型放在索引不为 0 的 SYCL 设备上运行时，llama.cpp 的 SYCL 后端会在首次 decode 阶段因 host pinned 内存被错误地绑定到 device 0 而崩溃，报出具有误导性的显存不足错误。优先确认是否只用单卡（SYCL0）或改用环境变量绕过，再升级到包含修复的版本。

适用环境：llama.cpp 0.4.0-dev（commit 41fc7584f），Clang 24.0.0，Linux x86_64，SYCL 后端，硬件为 Intel Arc Pro B70 Graphics + Intel UHD Graphics 770，系统中存在多于一个可见 SYCL 设备。

最快修复方案：暂无确认的一步修复方案；Issue 中给出的临时绕过方法是设置 GGML_SYCL_ENABLE_HOST_PINNED_MEM=0，改用普通堆内存，避免 host pinned 内存被固定到 device 0。该问题最终由 PR #28895 修复。

注意事项：绕过方法会关闭 host pinned memory 优化，可能影响主机与设备间的拷贝性能。切换 ONEAPI_DEVICE_SELECTOR 的设备顺序只是让崩溃换一张卡出现，不能真正解决问题。

问题场景

用户在 Linux 下使用带 SYCL 后端的 llama.cpp（0.4.0-dev，commit 41fc7584f，Clang 24.0.0 编译），机器上同时存在两块 Intel GPU（Arc Pro B70 Graphics 与 UHD Graphics 770），即有多个 SYCL 设备可见。当通过 --device SYCL1 把模型放到索引不为 0 的设备上、并执行 decode 时必然崩溃。例如：

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

模型的输出张量位于非 0 索引的 SYCL 设备时即可复现，与模型大小（1B、7B、3B 均可）、驱动版本（NEO 26.22 / oneAPI / icpx 与 NEO 26.36 / clang 表现一致）无关。

报错原文

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

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.

#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 (...)

原因分析

根据 Issue 作者的分析，最可能的根因是 PR #26789 引入的 ggml_backend_sycl_host_malloc 把设备索引硬编码为 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{});

由此产生的调用路径是：llama-context.cpp 向输出张量所在设备请求 host buffer → ggml_backend_sycl_device_get_host_buffer_type 忽略传入的 dev 参数（GGML_UNUSED(dev)）→ 分配时总是把 host pinned 内存固定到 device 0。随后代码试图用非 0 设备的 queue 去拷贝这块属于 device 0（在作者环境中甚至属于不同 platform）的内存，指针对执行拷贝的 context 不可见，于是报错。

由于底层返回的是 UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY，表象具有误导性，看起来像显存不足，实际是无效 context 或无效指针，这被作者认为是 compute-runtime 侧的一个独立问题。

交叉验证表明：把 ONEAPI_DEVICE_SELECTOR=level_zero:1,0 交换设备顺序会让原本正常的卡崩溃、原本崩溃的卡恢复正常，说明关键只在设备索引；与显存大小、驱动版本无关。

环境排查

确认 llama.cpp 版本是否包含问题提交：0.4.0-dev 41fc7584f 可复现，首个坏提交为 a97123e497968f3440264c0464a7adc7c999c027（其父提交 2606220d9 正常）。
确认 SYCL 后端已启用，且系统中有多于一个可见 SYCL 设备。
确认运行命令是否用 --device 指定了非 0 索引的设备（如 SYCL1）。
确认 GGML_SYCL_ENABLE_HOST_PINNED_MEM 当前取值；默认启用 host pinned 内存时更易触发。
确认 ONEAPI_DEVICE_SELECTOR 是否被设置，用来排除设备顺序造成的现象变化。
确认编译器与驱动组合：Clang 24.0.0 编译，NEO 26.22 / oneAPI / icpx 与 NEO 26.36 / clang 均能复现，说明驱动版本不是变量。

解决步骤

先确认崩溃是否只在输出张量位于非 0 索引 SYCL 设备时出现，例如把 --device SYCL1 换成 --device SYCL0 观察是否恢复。
可优先尝试 Issue 中给出的绕过方法：设置环境变量 GGML_SYCL_ENABLE_HOST_PINNED_MEM=0，让 host 端改用普通堆内存而非 pinned 内存，再次运行相同的 llama-bench 命令。
如果只是想让当前机器先跑起来，也可以暂时把模型放到索引为 0 的 SYCL 设备上执行。
升级到包含 PR #28895 的 llama.cpp 版本，该 PR 按 Issue 讨论链中的评论摘要用于彻底修复本问题（修复思路是使用 get_host_buffer_type 传入的设备，而不是写死 0）。

验证方法

在原本崩溃的多设备环境中，保持 --device SYCL1（非 0 索引）重新运行 llama-bench，观察首次 decode 是否仍出现 UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY 或 ggml_backend_sycl_get_tensor_async 相关报错。若使用绕过方法，可在关闭 GGML_SYCL_ENABLE_HOST_PINNED_MEM 后确认同一命令能正常完成；若升级到修复版本，可再开启 host pinned 内存复测，确认非 0 设备也能正常 decode。也可通过交换 ONEAPI_DEVICE_SELECTOR 设备顺序，验证崩溃是否不再随设备索引转移。