Eval bug: [SYCL] Qwen3-Next architecture: layer-split tensor allocation broken on dual Intel Arc Pro B70 (Battlemage/Xe2) on master b9093

[david-bue]

load_attempt.log

[SYCL] Qwen3-Next architecture: layer-split tensor allocation broken on dual Intel Arc Pro B70 (Battlemage/Xe2) on master b9093

Summary

Loading Qwen3-Coder-Next 80B-A3B (arch = qwen3next) on dual Intel Arc Pro B70 with --split-mode layer --tensor-split 50,50 requests a single 25.4 GB (25,371,093,248 byte) allocation on SYCL0, ignoring the tensor-split distribution across the two cards. Each card has 32 GB free at load time, so the model (46.20 GiB total) should fit comfortably split 50/50 with KV cache headroom.

The same model loaded successfully under Windows on the official prebuilt SYCL release from a few weeks earlier (asymmetric distribution: SYCL0 ~30.4 GB, SYCL1 ~31.3 GB). PMZFX's public benchmark of this exact model on dual B70 reports 43.4 t/s on commit ec6f7a6a5c (build b8840-12, 2026-04-21).

Reproduction

Hardware

• 2× Intel Arc Pro B70 (BMG-G31, full Big Battlemage die)
• 32 GB GDDR6 ECC per card, PCIe Gen 4 x8 each
• AMD Ryzen 5 4500, 64 GB system RAM
• ASRock B550M-A motherboard

Software

• Fedora 44, kernel 6.19.10-300.fc44.x86_64
• Intel Compute Runtime 26.09.37435.12
• oneAPI Level Zero 1.28.2 (V1, V2 disabled via SYCL_UR_USE_LEVEL_ZERO_V2=0)
• Intel oneAPI 2026.0.0 (DPC++ compiler 2026.0.0.20260331)
• llama.cpp commit 1e5ad35d5 (build 9093, master HEAD as of 2026-05-09)

Build

cmake -B build \
  -DGGML_SYCL=ON \
  -DGGML_SYCL_F16=ON \
  -DGGML_SYCL_TARGET=INTEL \
  -DCMAKE_C_COMPILER=icx \
  -DCMAKE_CXX_COMPILER=icpx \
  -DGGML_NATIVE=ON \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_C_FLAGS_RELEASE="-O3 -DNDEBUG" \
  -DCMAKE_CXX_FLAGS_RELEASE="-O3 -DNDEBUG" \
  -G Ninja
cmake --build build --config Release -j 4

Build completes cleanly. llama-server --version confirms IntelLLVM 2026.0.0.

Environment

export SYCL_UR_USE_LEVEL_ZERO_V2=0
export GGML_SYCL_DISABLE_OPT=1
export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1

Launch command

./build/bin/llama-server \
  -m Qwen3-Coder-Next-UD-Q4_K_XL.gguf \
  -ngl 99 \
  --split-mode layer \
  --tensor-split 50,50 \
  -c 32768 \
  -np 1 \
  --host 127.0.0.1 \
  --port 8080

Model: unsloth/Qwen3-Coder-Next-GGUF UD-Q4_K_XL quantization, 46.20 GiB on disk.

Result

Both devices are detected:

llama_prepare_model_devices: using device SYCL0 (Intel(R) Graphics [0xe223]) - 32568 MiB free
llama_prepare_model_devices: using device SYCL1 (Intel(R) Graphics [0xe223]) - 31790 MiB free

But tensor allocation immediately fails on a single 25.4 GB request to SYCL0:

load_tensors: loading model tensors, this can take a while... (mmap = true, direct_io = false)
ggml_backend_sycl_buffer_type_alloc_buffer: can't allocate 25371093248 Bytes of memory on device
alloc_tensor_range: failed to allocate SYCL0 buffer of size 25371093248
llama_model_load: error loading model: unable to allocate SYCL0 buffer
llama_model_load_from_file_impl: failed to load model

The 25.4 GB allocation pattern suggests the entire MoE expert tensor block (or a substantial fraction of model weights) is being allocated as a single contiguous SYCL buffer to SYCL0 before --tensor-split distribution kicks in. With 32 GB available, the allocation is just over the per-card limit but well within the dual-card combined budget.

sycl::ext::intel::experimental::level_zero per-allocation cap with UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 should allow allocations >4 GB; that env var is set and earlier smaller failures (~8 GB) on a different model also bypassed the 4 GB limit, so the 25 GB request reaches the underlying memory pool, which simply doesn't have that much contiguous space on a single device.

Expected behavior

--split-mode layer --tensor-split 50,50 should distribute model weights across both SYCL devices so that no single allocation exceeds either card's free VRAM. With 32 GB free per card, a 46 GiB model split should produce roughly 23 GB per card.

Comparison with known-good build

PMZFX's public benchmarks (https://github.com/PMZFX/intel-arc-pro-b70-benchmarks) report this exact model loading and running at 43.4 t/s tg128 on commit ec6f7a6a5c (b8840-12, 2026-04-21) using --split-mode layer on the same hardware class. Their hardware.md confirms --split-mode layer works on dual B70 (and --split-mode row segfaults, separately).

The Windows official prebuilt SYCL release from approximately the same era (early-to-mid April 2026) also loaded this model successfully on the same physical hardware, with asymmetric distribution:

• SYCL0: ~30.4 GB committed
• SYCL1: ~31.3 GB committed

Strongly suggests a regression in the b8840 → b9093 window affecting dual-card allocation for qwen3next-architecture models.

Additional notes

• --split-mode none and -ngl 99 on a single card cannot fit this 46 GiB model on a single 32 GB B70, so single-card fallback is not an option.
• --split-mode row is reported separately as crashing on dual B70 (PMZFX hardware.md, not yet filed upstream).
• A separate (likely related) failure occurs on Qwen3.6-27B (arch = qwen35, also hybrid attention/SSM) where the same dual-B70 SYCL build cannot allocate an 8.1 GB tensor regardless of context size. May share a root cause with this issue.

Full load attempt log

[Attached: load_attempt.log showing complete output from launch through failure]

llama.cpp version

$ ./build/bin/llama-server --version
version: 9093 (1e5ad35d5)
built with IntelLLVM 2026.0.0 for Linux x86_64

[NeoZhangJianyu]

@david-bue
1.
Following are the environment variables in your case:

export SYCL_UR_USE_LEVEL_ZERO_V2=0
export GGML_SYCL_DISABLE_OPT=1
export UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1

Please keep UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS only.
Test again.

Could run following code to test the malloc 4GB*4 memory on your PC?

icpx -fsycl malloc_4gb.cpp
./a.out

malloc_4gb.cpp:

#include <sycl/sycl.hpp>
#include <cstdint>
#include <iostream>

int main() {
  try {
    // Pick Intel GPU
    sycl::queue q{sycl::gpu_selector_v};
    auto dev = q.get_device();

    std::cout << "Device: "
              << dev.get_info<sycl::info::device::name>() << "\n";

    constexpr std::size_t bytes = 4ULL * 1024ULL * 1024ULL * 1024ULL; // 4 GiB
    constexpr std::size_t n = bytes / sizeof(std::uint32_t);

    // USM device allocation on GPU
    std::uint32_t* ptr = sycl::malloc_device<std::uint32_t>(n, q);
    if (!ptr) {
      std::cerr << "malloc_device returned nullptr for 4 GiB\n";
      return 1;
    }

    // Touch memory from a kernel so allocation is actually used
    q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) {
       ptr[i] = static_cast<std::uint32_t>(i[0]);
    }).wait();

    std::cout << "Allocated and initialized 4 GiB on GPU successfully.\n";

    sycl::free(ptr, q);
    return 0;
  } catch (const sycl::exception& e) {
    std::cerr << "SYCL exception: " << e.what() << "\n";
    return 2;
  }
}

[david-bue]

Findings from a full debug session

After ~7 hours of investigation on Fedora 44 with 2× Intel Arc Pro B70 (BMG-G31, 32 GB each), I can confirm the failure is upstream in intel-compute-runtime (NEO), not in llama.cpp. This is a regression in NEO somewhere between releases 25.40.35563.10 and 26.05.37020.3. The pool-manager hypothesis discussed in [intel/compute-runtime #916](intel/compute-runtime#916) is not the actual cause. The failure point is deeper.

Symptom

llama-cli and llama-server fail to load any model the moment the SYCL backend constructs a Level Zero context spanning both GPUs. Every USM device allocation returns UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY, regardless of size, despite 32 GB free per card.

ggml_backend_sycl_buffer_type_alloc_buffer: can't allocate 427970304 Bytes of memory on device
alloc_tensor_range: failed to allocate SYCL0 buffer of size 427970304
llama_model_load: error loading model: unable to allocate SYCL0 buffer

The size is irrelevant. Sub-2 MB allocations fail too. With ONEAPI_DEVICE_SELECTOR=level_zero:0 (single-device context), everything works.

What I tried, and what each result ruled out

1. Built NEO 26.14.37833.4 with commit 7bb2d32 reverted (the change blamed in #916, which flipped useUsmPoolManager from false to true). Same failure. The pool wasn't the actual problem.

2. Built NEO 25.40.35563.10 (the issue #916 reporter's "last known good"). Multi-device contexts work — Qwen3-Coder-Next 80B-A3B Q4_K_M loads across both cards. 25.40 → 26.x is confirmed as the regression window.

3. Built NEO 26.05.37020.3 with 7bb2d32 reverted. Same failure. Confirms the bug is not the named commit.

4. Installed Ubuntu 26.04's actual libze-intel-gpu1_26.05.37020.3-1_amd64v3.deb in a sibling prefix and LD_LIBRARY_PATH-prepended it. Verified with LD_DEBUG=libs that Ubuntu's NEO + L0 loader + gmmlib were the actually-loaded libraries. Same failure. Ubuntu's literal binary fails identically on this Fedora system. So the contents of NEO are not what makes Ubuntu setups work.

5. Set UR_L0_DISABLE_USM_ALLOCATOR=1 and SYCL_PI_LEVEL_ZERO_DISABLE_USM_ALLOCATOR=1 to fully disable UR's pool manager (real env vars, see unified-runtime/source/adapters/level_zero/v2/usm.cpp:46). Same failure. Rules out the UR pool layer too.

6. UR trace from SYCL_UR_TRACE=2:

urContextCreate(.DeviceCount = 2, .phDevices = {0x..., 0x...}, ...) -> UR_RESULT_SUCCESS;
urUSMDeviceAlloc(.hContext = ctx2, .hDevice = dev0, ..., .pool = nullptr, .size = 427970304, ...)
  -> UR_RESULT_ERROR_OUT_OF_DEVICE_MEMORY;

Note .pool = nullptr . Nno UR pool involved. The allocation fails at a layer below the pool.

7. Traced into NEO. The failing path is in level_zero/core/source/context/context_imp.cpp (ContextImp::allocDeviceMem) at line ~322. After the pool branches are skipped (which we've forced via patches and env vars), execution falls through to:

void *usmPtr = this->driverHandle->svmAllocsManager->createUnifiedMemoryAllocation(...);
if (usmPtr == nullptr) { return ZE_RESULT_ERROR_OUT_OF_DEVICE_MEMORY; }

SVMAllocsManager::createUnifiedMemoryAllocation() at shared/source/memory_manager/unified_memory_manager.cpp:498 returns nullptr in multi-device contexts, despite plenty of VRAM. The call into memoryManager->allocateGraphicsMemoryWithProperties(unifiedMemoryProperties, externalPtr) is what fails. This is the actual failure point.

Bisect range

The regression is in intel/compute-runtime between tags 25.40.35563.10 and 26.05.37020.3. After filtering out clang-format / refactor / revert noise, eight semantically interesting commits touched shared/source/memory_manager/unified_memory_manager.cpp in that window:

7c701fd69f feature:  Add multi-GPU capability for memAdvise           (Jan 19 2026)
dc0cb255c3 performance: Optimize compression condition
bb15deca7e fix: do not wait on tag when releasing in-place
dc85f33e49 fix: Disallow USM compression on peer access               (Nov 12 2025)
435c43d142 performance: Use 2MB heap for cl_buffers if possible
00ec4c82b2 fix: blocking usm reuse cleanup
6d6715bebf performance: Lock svm allocs mutex once in appendLaunchKernelWithArgs
17199556ca performance: lazy start usm reuse cleaner thread

Strongest candidate by smell is dc85f33e49. The only commit that explicitly changes behavior based on multi-device peer access. Reading the diff, it disables compression rather than returning null, so it shouldn't fail allocation directly, but it touches the same function we've isolated as the failure point. A git bisect between these two tags would identify the actual commit in ~5 build-and-test iterations.

Stack

• Fedora 44, kernel 7.0.4-200.fc44.x86_64
• xe driver (mandatory for BMG-G31), GuC v70.60.0, HuC v8.2.10, DMC v2.6
• intel-igc 2.30.3, intel-gmmlib 22.10.0, oneapi-level-zero 1.28.2
• Intel oneAPI Base Toolkit 2026.0 (icpx 2026.0.0)
• llama.cpp master commit 8e1f9d083 (version 9112)
• Build: SYCL backend, GGML_SYCL_TARGET=INTEL, GGML_SYCL_DEVICE_ARCH=bmg_g31, AOT, NDEBUG release
• Hardware: ASRock X570 Taichi (modded BIOS) + Ryzen 5 4500, 2× Intel-reference Arc Pro B70 32 GB at PCIe 4.0 x8 each, Above 4G + ReBAR on

Reproduction

$ ~/llama.cpp/build/bin/llama-cli --list-devices
Available devices:
  SYCL0: Intel(R) Graphics [0xe223] (32656 MiB, 32611 MiB free)
  SYCL1: Intel(R) Graphics [0xe223] (32656 MiB, 32020 MiB free)

# Fails. Multi-device context, hits the regression:
$ ONEAPI_DEVICE_SELECTOR='level_zero:0,1' ./llama-cli -m model.gguf -ngl 999 -p Hi -n 8
ggml_backend_sycl_buffer_type_alloc_buffer: can't allocate ... Bytes of memory on device
llama_model_load: error loading model: unable to allocate SYCL0 buffer

# Works. Single-device context:
$ ONEAPI_DEVICE_SELECTOR='level_zero:0' ./llama-cli -m model.gguf -ngl 999 -p Hi -n 8
# (tokens generate normally)

Workarounds for llama.cpp users until NEO ships a fix

1. Downgrade intel-compute-runtime to 25.40.35563.x. Last known-good. I built 25.40.35563.10 from source on Fedora 44 with three small GCC 15 patches (one <cstdint> include, one <climits> include, one -Wno-error=sfinae-incomplete flag for opencl/source/mem_obj/buffer.h). On dual-B70 this gives ~26 t/s on Qwen3-Coder-Next 80B-A3B Q4_K_M. The MoE kernels in 25.40 are months older than 26.x, so absolute throughput is below what 26.x would deliver if it worked. (For comparison: same hardware on Ubuntu 26.04 with intel-compute-runtime 26.05.37020.3-1 is reported to hit ~50 t/s on this exact model. The Ubuntu setup somehow avoids this bug. Likely via a different oneAPI/UR adapter combination that doesn't construct the failing code path.)

2. One llama-server per GPU, different ports, each with ONEAPI_DEVICE_SELECTOR=level_zero:<n> for a single-device context. Works on every 26.x NEO. Models must fit in one card.

3. Build llama.cpp with -DGGML_RPC=ON and run one rpc-server per GPU. Splits one model across both cards while each rpc-server stays in its own single-device L0 context. ~1-2% overhead.

4. Vulkan backend (-DGGML_VULKAN=ON) Different driver stack, doesn't touch NEO's USM allocator. ~70-90% of SYCL performance on dense models, often on par for MoE.

Suggested action for llama.cpp

When the SYCL backend detects a multi-device L0 context on intel-compute-runtime 26.05+, surface a clear startup warning pointing users to upstream NEO investigation and listing the workarounds above. The current alloc_tensor_range: failed to allocate SYCL0 buffer of size N reads like an out-of-memory error and sends people to check ReBAR / Above 4G / VRAM size, none of which are the actual problem.

Happy to attach full UR/L0 traces, the GCC 15 patch set, or the LD_DEBUG=libs resolution chain confirming Ubuntu binaries fail identically on Fedora.

[NeoZhangJianyu]

Yes, we find there is similar issue: #21747.
The root cause is 26.0 compute-runtime.

Thank you!

[NeoZhangJianyu]

@david-bue
Refer to the solution in PR:#22968

[github-actions]

This issue was closed because it has been inactive for 14 days since being marked as stale.