[Bug]: XPU TP=2 on dual Intel Arc Pro B70 (Battlemage): GP fault + xe BCS engine reset reproduces in intel/vllm:0.17.0-xpu on Ubuntu 24.04 HWE 6.17

[Zumbasam]

Your current environment

The output of python collect_env.py

Collecting environment information...
==============================
        System Info
==============================
OS                           : Ubuntu 24.04.4 LTS (x86_64)
GCC version                  : (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
Clang version                : Could not collect
CMake version                : version 4.2.3
Libc version                 : glibc-2.39

==============================
       PyTorch Info
==============================
PyTorch version              : 2.10.0+xpu
Is debug build               : False
CUDA used to build PyTorch   : None
ROCM used to build PyTorch   : N/A
XPU used to build PyTorch    : 20250301

==============================
      Python Environment
==============================
Python version               : 3.12.3 (main, Jan 22 2026, 20:57:42) [GCC 13.3.0] (64-bit runtime)
Python platform              : Linux-6.17.0-23-generic-x86_64-with-glibc2.39

==============================
      Intel XPU / GPU Info
==============================
Is XPU available             : True
XPU runtime version          : 20250301
Intel GPU models             :
GPU 0: Intel(R) Graphics [0xe223]
GPU 1: Intel(R) Graphics [0xe223]


--Compile time--
oneAPI compiler version      : 2025.3.2
SYCL compiler build          : 2025.3.2.20260112
oneCCL version               : 2021.17.2-5

--Runtime--
Intel Graphics Compiler (IGC): 2.24.8
Intel GMM (libigdgmm)        : 22.8.2
Level Zero loader version    : 1.26.0
Level Zero driver version    : 25.48.36300.8-0
vLLM XPU kernels version     : 0.1.4

==============================
          CPU Info
==============================
Architecture:                            x86_64
CPU op-mode(s):                          32-bit, 64-bit
Address sizes:                           46 bits physical, 48 bits virtual
Byte Order:                              Little Endian
CPU(s):                                  24
On-line CPU(s) list:                     0-23
Vendor ID:                               GenuineIntel
Model name:                              Intel(R) Core(TM) Ultra 9 285K
CPU family:                              6
Model:                                   198
Thread(s) per core:                      1
Core(s) per socket:                      24
Socket(s):                               1
Stepping:                                2
CPU(s) scaling MHz:                      26%
CPU max MHz:                             5700.0000
CPU min MHz:                             800.0000
BogoMIPS:                                7372.80
Flags:                                   fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rdt_a rdseed adx smap clflushopt clwb intel_pt sha_ni xsaveopt xsavec xgetbv1 xsaves split_lock_detect user_shstk avx_vnni lam wbnoinvd dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp hwp_pkg_req hfi vnmi umip pku ospke waitpkg gfni vaes vpclmulqdq rdpid bus_lock_detect movdiri movdir64b fsrm md_clear serialize arch_lbr ibt flush_l1d arch_capabilities
Virtualization:                          VT-x
L1d cache:                               768 KiB (20 instances)
L1i cache:                               1.3 MiB (20 instances)
L2 cache:                                40 MiB (12 instances)
L3 cache:                                36 MiB (1 instance)
NUMA node(s):                            1
NUMA node0 CPU(s):                       0-23
Vulnerability Gather data sampling:      Not affected
Vulnerability Ghostwrite:                Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit:             Not affected
Vulnerability L1tf:                      Not affected
Vulnerability Mds:                       Not affected
Vulnerability Meltdown:                  Not affected
Vulnerability Mmio stale data:           Not affected
Vulnerability Old microcode:             Not affected
Vulnerability Reg file data sampling:    Not affected
Vulnerability Retbleed:                  Not affected
Vulnerability Spec rstack overflow:      Not affected
Vulnerability Spec store bypass:         Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1:                Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:                Mitigation; Enhanced / Automatic IBRS; IBPB conditional; PBRSB-eIBRS Not affected; BHI BHI_DIS_S
Vulnerability Srbds:                     Not affected
Vulnerability Tsa:                       Not affected
Vulnerability Tsx async abort:           Not affected
Vulnerability Vmscape:                   Mitigation; IBPB before exit to userspace

==============================
Versions of relevant libraries
==============================

[pip3] flake8==7.3.0
[pip3] mypy==1.19.1
[pip3] mypy_extensions==1.1.0
[pip3] numpy==2.2.6
[pip3] pyzmq==27.1.0
[pip3] torch==2.10.0+xpu
[pip3] torchaudio==2.10.0+xpu
[pip3] torchvision==0.25.0+xpu
[pip3] transformers==4.57.6
[pip3] triton-xpu==3.6.0
[conda] Could not collect

==============================
         vLLM Info
==============================
ROCM Version                 : Could not collect
vLLM Version                 : 0.1.dev14456+gde3f7fe65 (git sha: de3f7fe65)
vLLM Build Flags:
  CUDA Archs: Not Set; ROCm: Disabled; XPU: Enabled
GPU Topology:
  Could not collect

==============================
     Environment Variables
==============================
CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0
CCL_ROOT=/opt/intel/oneapi/ccl/2021.17
I_MPI_ROOT=/opt/intel/oneapi/mpi/2021.17
CCL_CONFIGURATION=cpu_gpu_dpcpp
VLLM_WORKER_MULTIPROC_METHOD=spawn
CCL_CONFIGURATION_PATH=
ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE
CCL_ATL_TRANSPORT=ofi
ZE_AFFINITY_MASK=0,1
CCL_ZE_IPC_EXCHANGE=pidfd
VLLM_TARGET_DEVICE=xpu
LD_LIBRARY_PATH=/tmp/ucx_install/lib:/opt/intel/oneapi/tcm/1.4/lib:/opt/intel/oneapi/umf/1.0/lib:/opt/intel/oneapi/tbb/2022.3/env/../lib/intel64/gcc4.8:/opt/intel/oneapi/pti/0.16/lib:/opt/intel/oneapi/mpi/2021.17/opt/mpi/libfabric/lib:/opt/intel/oneapi/mpi/2021.17/lib:/opt/intel/oneapi/mkl/2025.3/lib:/opt/intel/oneapi/dnnl/2025.3/lib:/opt/intel/oneapi/debugger/2025.3/opt/debugger/lib:/opt/intel/oneapi/compiler/2025.3/opt/compiler/lib:/opt/intel/oneapi/compiler/2025.3/lib:/opt/intel/oneapi/ccl/2021.17/lib/:/usr/local/lib/
PYTORCH_NVML_BASED_CUDA_CHECK=1
TORCHINDUCTOR_COMPILE_THREADS=1
TORCHINDUCTOR_CACHE_DIR=/tmp/torchinductor_root

🐛 Describe the bug

Summary

intel/vllm:0.17.0-xpu reproduces a VLLM::Worker general protection fault followed by xe BCS engine resets on a dual Intel Arc Pro B70 system when running Qwen3-30B-A3B with TP=2 and dynamic FP8 on Ubuntu 24.04.4 HWE kernel 6.17.

The same failure pattern also reproduces with our local source-built vLLM XPU image. This suggests the issue is not caused by our source build alone, but likely involves an interaction between the host stack, xe driver/firmware, PCIe topology, and vLLM XPU TP=2 worker initialization / ProcessGroupXCCL path.

A standalone XCCL/SYCL collective sweep passes on the same host, so this does not appear to be a generic oneCCL/SYCL collective failure.

Hardware

• 2× Intel Arc Pro B70 32GB, Battlemage G31, device 8086:e223
• Z890 motherboard with PCIe switch topology
• No XeLink between GPUs
• p2p_access:0
• Each B70 reports PCIe 5.0 x8 host link
• ZE_AFFINITY_MASK=0,1 used for the two B70 GPUs

Host stack

• OS: Ubuntu 24.04.4 LTS
• Kernel: 6.17.0-23-generic HWE
• GPU driver: in-kernel xe
• GuC firmware: 70.44.1
• GuC firmware version 70.49.4 has been mentioned in upstream contexts as a newer candidate; we have not yet tested an upgrade.

Containers tested

Intel pre-built image

• intel/vllm:0.17.0-xpu
• Tested digest prefix: sha256:e961d08135a6...

Intel's published 0.17.0-xpu validation BOM differs from our host stack. It lists Ubuntu 25.04, KMD 6.14.0, PyTorch 2.10, vllm-xpu-kernels 0.1.4, and oneCCL 2021.15.7.8, but collect_env.py inside the tested intel/vllm:0.17.0-xpu container reports oneCCL 2021.17.2-5 and CCL_ROOT=/opt/intel/oneapi/ccl/2021.17.

Local source-built image

Built from vLLM Dockerfile.xpu:

• vLLM main
• PyTorch 2.11 + XPU
• vllm-xpu-kernels 0.1.7
• oneCCL 2021.17.2
• oneAPI 2025.3

Model / vLLM setup

• Model: Qwen/Qwen3-30B-A3B
• Local BF16 model path
• Dynamic FP8 quantization via --quantization fp8
• Tensor parallelism: --tensor-parallel-size 2

Common args:

--dtype float16
--quantization fp8
--tensor-parallel-size 2
--gpu-memory-utilization 0.88
--max-model-len 8192
--max-num-batched-tokens 8192
--max-num-seqs 64
--block-size 64
--trust-remote-code

Minimal reproducer

docker pull intel/vllm:0.17.0-xpu

docker run -d --name vllm-tp2-repro \
  --device /dev/dri \
  --group-add render \
  --group-add video \
  --network host \
  --shm-size 16g \
  --ipc host \
  -v /path/to/Qwen3-30B-A3B:/model:ro \
  -e CCL_ATL_TRANSPORT=ofi \
  -e CCL_ZE_IPC_EXCHANGE=pidfd \
  -e CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0 \
  -e ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE \
  -e ZE_AFFINITY_MASK=0,1 \
  -e ONEAPI_DEVICE_SELECTOR=level_zero:* \
  -e VLLM_WORKER_MULTIPROC_METHOD=spawn \
  intel/vllm:0.17.0-xpu \
  python3 -u -m vllm.entrypoints.openai.api_server \
  --model /model \
  --dtype float16 \
  --quantization fp8 \
  --tensor-parallel-size 2 \
  --gpu-memory-utilization 0.88 \
  --max-model-len 8192 \
  --max-num-batched-tokens 8192 \
  --max-num-seqs 64 \
  --block-size 64 \
  --enforce-eager \
  --trust-remote-code \
  --host 0.0.0.0 \
  --port 8000

Observed result

traps: VLLM::Worker[...] general protection fault
traps: VLLM::Worker[...] general protection fault

xe 0000:03:00.0: [drm] GT0: Engine reset: engine_class=bcs
xe 0000:07:00.0: [drm] GT0: Engine reset: engine_class=bcs

In one run, the GP fault was reported in libc.so.6 from the worker processes. We do not know whether this indicates a pthread/context issue, memory corruption surfacing in libc, or another native runtime issue.

Test matrix

Test	Stack	Key config	Result	BCS reset
Profile A	local source-build	CCL_ENABLE_SYCL_KERNELS=0	Stable, HEALTH=200, ~362 tok/s @ 50 concurrency	0
Profile B	local source-build	CCL_ENABLE_SYCL_KERNELS=1	XCCL / worker-init crash	yes
Test F	local source-build	CCL_ENABLE_SYCL_KERNELS=1, UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1	ProcessGroupXCCL::allreduce_impl crash	0
Test G	local source-build	above + CCL_ALLREDUCE=ring	Passed allreduce-init, later crashed during Dynamo/Inductor compile	0
C4 v2	local source-build	above + --enforce-eager, VLLM_XPU_ENABLE_XPU_GRAPH=0	HEALTH=200, inference works, but ~0.5 tok/s	0
C8	intel/vllm:0.17.0-xpu	Intel image, TP=2 FP8	VLLM::Worker GP fault + BCS reset	yes
C8b	intel/vllm:0.17.0-xpu	above + CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0	same GP fault + BCS reset on both GPUs	yes

Important control: standalone XCCL/SYCL sweep passes

Using the same host and relevant XPU/CCL environment as the failing vLLM path:

CCL_ENABLE_SYCL_KERNELS=1
torchrun --nproc-per-node=2 xccl_sweep.py

Results:

Collective	fp32	fp16	bf16	int32	Max size
all_reduce	OK	OK	OK	OK	33M elements / 134 MB
all_gather	OK	OK	OK	OK	33M elements / 134 MB
reduce_scatter	OK	OK	OK	not tested	33M elements / 134 MB

This suggests the hardware and base oneCCL/SYCL collective path are functional in isolation. The crash appears to require the vLLM TP=2 worker-init / ProcessGroupXCCL context.

Failure domains observed

We observed three separate failure domains while trying to make the source-built image start with CCL_ENABLE_SYCL_KERNELS=1.

1. BCS engine reset / Level Zero V2 copy-offload side effect

UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1 avoided the BCS engine reset in one worker-init canary, but did not fix the underlying vLLM/XCCL crash.

2. oneCCL topo allreduce / ProcessGroupXCCL path

The source-build crashed in:

c10d::ProcessGroupXCCL::allreduce_impl
c10d::ProcessGroupXCCL::collective
c10d::ops::allreduce_XPU

CCL_ALLREDUCE=ring got past this stage, but performance collapsed and the next failure occurred later during compile.

oneCCL documents that for GPU buffers, topo is the default algorithm, and that algorithms other than topo copy GPU-buffer data to host and follow the specified CPU algorithm. That likely explains the large performance drop with CCL_ALLREDUCE=ring.

3. vLLM compile / XPU graph path

After bypassing the allreduce failure with UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1 and CCL_ALLREDUCE=ring, the next failure occurred during Dynamo/Inductor compile.

Adding:

--enforce-eager
VLLM_XPU_ENABLE_XPU_GRAPH=0

allowed vLLM to reach HEALTH=200 and serve inference, but throughput was only ~0.5 tok/s in our smoke test. This is a debug path, not a usable workaround.

Stable fallback on our current host

The only stable and usable config we found on this Ubuntu 24.04.4 / HWE 6.17 host is:

CCL_ENABLE_SYCL_KERNELS=0
CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0
SYCL_UR_USE_LEVEL_ZERO_V2=0
CCL_ATL_TRANSPORT=ofi
CCL_ZE_IPC_EXCHANGE=pidfd
ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE
ZE_AFFINITY_MASK=0,1
ONEAPI_DEVICE_SELECTOR=level_zero:*
VLLM_WORKER_MULTIPROC_METHOD=spawn
VLLM_XPU_ENABLE_XPU_GRAPH=1

Observed performance:

• ~29 tok/s single-stream
• ~240 tok/s at 16 concurrency
• ~362 tok/s at 50 concurrency

Why this looks host-stack related

We are not claiming the exact root cause is known. However:

1. Intel's pre-built intel/vllm:0.17.0-xpu image reproduces the same failure pattern.
2. Adding CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0 removes the CCL topology warning but does not prevent the GP fault + BCS reset.
3. Standalone torchrun XCCL/SYCL collective sweeps pass on the same host.
4. The failure requires the vLLM TP=2 worker initialization / ProcessGroupXCCL context.
5. Our host differs from Intel's published 0.17.0-xpu validation BOM, especially OS/kernel/KMD/firmware.

A public external report claims that a similar 2× Intel Arc Pro B70 + Qwen3-30B-A3B + TP=2 + FP8 setup using intel/vllm:0.17.0-xpu on Ubuntu 25.10 reaches high throughput. We have not reproduced that environment yet, so we treat this only as motivation for a host-stack A/B test.

We are preparing a separate Ubuntu 25.10 spare-disk test and will report back.

Hardware behavior and safety

We observed BCS engine resets across multiple failed test runs. After each:

• The xe driver recovered the engine.
• GPU devices remained accessible.
• Temperatures returned to idle.
• No cascading errors were observed.
• A full reboot cleared all state cleanly.

Because BCS engine resets are not normal operation, we stopped repeated default Profile B testing on the current host stack.

Related upstream / external signals

• Intel's 0.17.0-xpu release notes say the container is validated on Intel Arc Pro B-Series Graphics and list the validation BOM as Ubuntu 25.04 + KMD 6.14.0 + oneAPI 2025.3.2 hotfix + PyTorch 2.10 + vllm-xpu-kernels 0.1.4 + oneCCL 2021.15.7.8.
• oneCCL docs state that CCL_ALLREDUCE=topo is the default for GPU buffers, and that non-topo algorithms copy GPU-buffer data to host and use the CPU algorithm.
• Intel SYCL runtime docs state that the Level Zero V2 adapter is used by default on Xe2/Battlemage, and that SYCL_UR_USE_LEVEL_ZERO_V2=0 selects the legacy adapter.
• vllm-xpu-kernels tracks ongoing work to make torch.compile the default execution path for XPU. This may be relevant to the later Dynamo/Inductor failure domain, but the primary C8/C8b failure happens earlier during worker initialization.
• There are similar multi-GPU Intel XPU reports in adjacent projects where single-GPU works but multi-GPU fails in SYCL/Level Zero cross-device paths. This may not be the same bug, but it suggests this class of issue is worth investigating.

Questions

1. Is dual Arc Pro B70 + Qwen3-30B-A3B + TP=2 + FP8 expected to work with intel/vllm:0.17.0-xpu on Ubuntu 24.04.4 HWE 6.17?
2. Is Ubuntu 24.04.4 / HWE 6.17 a supported host stack for intel/vllm:0.17.0-xpu on Intel Arc Pro B70?
3. Which host BOM should users target for dual Arc Pro B70 + Qwen3-30B-A3B + TP=2 + dynamic FP8?
4. Should this be routed to vLLM, vllm-xpu-kernels, oneCCL, Intel compute-runtime / Level Zero, or the Linux xe KMD?
5. What additional logs would be useful? We can provide redacted:
   • full dmesg -T
   • journalctl -k -b
   • docker logs
   • XCCL sweep script and output
   • lspci -nn
   • sycl-ls
   • xpu-smi output
   • exact container command lines

Before submitting a new issue...

• Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.

[Zumbasam]

Update: Ubuntu 25.10 A/B test attempted, but currently inconclusive.

The 25.10 install boots the same kernel version string as our 24.04 HWE install: 6.17.0-23-generic. However, the 25.10 run hit a host-side kernel/mm soft-lockup in the mmap/PTE-locking/THP path before vLLM produced Python logs (docker logs was empty).

This does not reproduce the original VLLM::Worker GP fault + xe BCS reset pattern and should not be interpreted as evidence for or against the original B70/XCCL issue. We are treating 25.10 as blocked for now and will continue using 24.04 as the stable baseline.

[Zumbasam]

Follow-up after 6 days of additional engineering — we now have a 3rd datapoint for the hierarchy this issue documents, and it's better than expected. Posting in case it helps anyone else debugging this stack.

Setup (delta from original issue body)
Same hardware (2× Intel Arc Pro B70, Z890 cascaded switch, PCIe 5.0 x8) and OS (Ubuntu 24.04.4 + HWE kernel 6.17.0-23). Stack changes since 2026-05-04:

vLLM main, commit 0c99629 (post-v0.20.0, 2026-05-01)
vllm-xpu-kernels v0.1.7 (2026-04-27 release)
Container: custom-built from vLLM main + vllm-xpu-kernels v0.1.7 (PyTorch 2.11)
Profile A baseline (CCL_ENABLE_SYCL_KERNELS=0) production-stable since 2026-05-06 — 444 tok/s output sustained @ 64 conc, 0/2048 failures in 37 min sustained bench, peak 67°C
Test
Goal: re-test CCL_ENABLE_SYCL_KERNELS=1 with all currently-known fixes stacked, to see which row of the hierarchy in the issue body we land in.

Env (Profile A baseline + the one change):

ZE_AFFINITY_MASK=0,1
VLLM_WORKER_MULTIPROC_METHOD=spawn
CCL_ENABLE_SYCL_KERNELS=1 # ← changed from 0
CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0
TORCH_LLM_ALLREDUCE=1
VLLM_USE_V1=1
UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1
CCL_ALLREDUCE=ring
VLLM_XPU_ENABLE_XPU_GRAPH=1
UR_L0_USE_IMMEDIATE_COMMANDLISTS=0 # libze fix, added 2026-05-06
Server: Qwen3-30B-A3B BF16 + --quantization fp8 dynamic, TP=2, --max-num-batched-tokens 2048 --enable-chunked-prefill --max-num-seqs 96 --gpu-memory-utilization 0.92 --max-model-len 8192.

Bench: vllm bench serve --random-input-len 1500 --random-output-len 500 --num-prompts 256 --max-concurrency 64 --endpoint /v1/completions.

Result
WorkerProc init now passes — no sycl::ProgramManager::getOrCreateURProgram crash like we saw on 2026-05-03. So adding UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1 + CCL_ALLREDUCE=ring + UR_L0_USE_IMMEDIATE_COMMANDLISTS=0 on top of CCL_ENABLE_SYCL_KERNELS=1 does get past the init-time crash described in the issue body.

5-minute canary at 64 conc:

======= Serving Benchmark Result ======

Successful requests: 256
Failed requests: 0
Maximum request concurrency: 64
Benchmark duration (s): 285.47
Total input tokens: 384000
Total generated tokens: 128000
Output token throughput (tok/s): 448.39
Peak output token throughput (tok/s): 896.00
Total token throughput (tok/s): 1793.56
---------------Time to First Token----------------
Mean TTFT (ms): 6692.48
Median TTFT (ms): 3078.03
P99 TTFT (ms): 34230.35
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms): 128.99
Median TPOT (ms): 136.44
P99 TPOT (ms): 140.41

===========================

vs Profile A baseline (CCL_ENABLE_SYCL_KERNELS=0, same stack): 444 tok/s sustained. So +1%, within noise. 0 BCS resets, peak GPU temp 66°C (well under our 85°C watchdog limit).

Crucially: no compile-time errors in 5 minutes of sustained load, no recompile warnings in the vllm logs.

Follow-up datapoint 2026-05-11 — rebuild with post-v0.1.7 kernels
To rule out "stale kernels" as an explanation for our paritet (vs the 912 tok/s claim some users have reported on similar hardware), we rebuilt with vllm-xpu-kernels==0.1.8.dev15+gaef85fb (HEAD aef85fb, 15 dev commits past v0.1.7 release), including:

aef85fb — OneDNN 3.12 + fp8 block gemm (vllm-project/vllm-xpu-kernels#173)
befa720 — upgrade fp8 moe gemm policy (vllm-project/vllm-xpu-kernels#333)
3d4f369 — refactor XpuFusedMoe — reduce host overhead (vllm-project/vllm-xpu-kernels#326)
060c542 — support uncontiguous Q in decode attn (vllm-project/vllm-xpu-kernels#318)
Same Profile A canary on the new build:

===== Serving Benchmark Result (new kernels) =====

Successful requests: 256
Failed requests: 0
Output token throughput (tok/s): 447.75
Peak output token throughput (tok/s): 896.00
Total token throughput (tok/s): 1791.01
Mean TTFT (ms): 7706.24
Median TPOT (ms): 134.93

=====================================
Side-by-side:

Metric v0.1.7 baseline post-v0.1.7 (dev15) Δ
Output throughput 444 tok/s 447.75 tok/s +0.8% (paritet)
Total throughput ~1795 1791.01 -0.2%
Median TPOT ~135 ms 134.93 ms identical
Failed requests 0 0 ✅
BCS resets 0 0 ✅
Compile errors 0 0 ✅
Peak GPU temp 67°C 65°C ✅
So the post-release kernel commits (XpuFusedMoe host-overhead refactor, OneDNN 3.12, fp8 MoE gemm policy upgrade) do not change Profile A throughput within noise — but they also do not regress it. New kernels are backwards-compatible.

We did not re-run the CCL_ENABLE_SYCL_KERNELS=1 configuration on the new build because (a) our hypothesis that newer kernels would close the gap to the 912 tok/s claim was already tested via a separate workload (Qwen3.6-27B Triton with --max-num-seqs 8) and rejected (-35% gap remains identical, 87.7 → 86.7 tok/s warm). The "stale kernels" hypothesis for both your reported regression and the community's high numbers appears to be ruled out for this particular axis.

Where this lands in the hierarchy
Mapping to the rows in the original issue body:

Fix layer Predicted outcome Our 2026-05-10 result
Default WorkerProc init crash (matches original report)

• UR_L0_V2_FORCE_DISABLE_COPY_OFFLOAD=1 BCS reset gone, allreduce crash remains (matches)
• CCL_ALLREDUCE=ring "comes past init but with big perf hit and later compile errors" init OK, +1% paritet, no compile errors in 5 min
  So our datapoint is better than the issue body's prediction for the third row. Two possible reasons:

The accumulated stack (libze fix UR_L0_USE_IMMEDIATE_COMMANDLISTS=0 from 2026-05-06 + post-v0.20.0 vLLM scheduler changes + vllm-xpu-kernels v0.1.7 with onednn 3.11.2) may have closed the perf gap and the "later compile errors" path
The original report's "big perf hit" and "later compile errors" may have been config- or workload-specific (different bench setup, different vLLM version, different oneCCL version)
Either way it's a positive update — for this exact stack, Profile B is safe to enable without a perf cliff. But we did not see the 2-2.5× speedup that a level1techs forum post claimed for similar hardware (912 tok/s @ 50 conc). That number does not replicate here — likely either burst rather than sustained, or some env we're still missing.

Caveats on our test
Only 5-minute canary, not 30-min sustained. The "later compile errors" prediction in the issue body might still show up at longer wallclock — we didn't run that long because there's no perf win to justify the production downtime
We have not tried CCL_ENABLE_SYCL_KERNELS=1 without CCL_ALLREDUCE=ring on the current stack. If anyone wants to bisect which of the three fixes is doing the work here, that'd be useful data
Net-net for our production deployment
We're staying on Profile A (CCL_ENABLE_SYCL_KERNELS=0) since there's no measurable win from flipping it. But the init-time crash this issue described is fixed on the current main+v0.1.7-kernels stack with the env we have above, at least for this workload.

Anyone else who has run dual-B70 TP=2 on the current stack and seen different numbers (especially the 912 tok/s claim from level1techs) — would love a comparison.

[chriswagner-ai]

Reproduction with different host stack — TP=2 GP fault confirmed, PP=2 surfaces a distinct phantom-OOM at first XPU all_reduce

Disclaimer & context

To be upfront, I am a Linux newbie and I use Claude to help me setup things and make things work.

Who I am: SEO consultant, long-time Windows user moving to Linux for AI workloads. Not a software engineer. Bought a dual Intel Arc Pro B70 setup both for productive work (local LLMs for client projects) and as a learning vehicle for the AI/GPU stack. Plan is to have 4 soon.

Setup journey: My current host stack is not the Intel-validated combination. The shortest version:

1. I initially attempted a bleeding-edge stack on advice from another AI assistant: PyTorch 2.11 + Python 3.14 + GCC 16 + Kernel 7.0.6. This produced segfaults at zeInit due to ABI mismatches and never reached a working state.
2. After researching Intel's published validation matrix for the B-Series (Ubuntu 25.04 / KMD 6.14.0 / oneAPI 2025.3.2 / PyTorch 2.10 / Python 3.12), I stepped back toward it: switched to Python 3.12, kept PyTorch 2.11.0+xpu (the newest stable Python 3.12 wheel currently in the PyTorch-XPU index).
3. My current bare-metal source-built stack is therefore closer to validated but still diverges on kernel (7.0.6 vs 6.14), distro (CachyOS rolling vs Ubuntu 25.04), GCC (16.1.1 vs distro default), and PyTorch (2.11 vs 2.10).

I am aware I run outside the validated BOM. I report this because the failures I reproduce match the original report despite the host stack differences, which is itself a useful data point.

Production of report: Together with Claude, which helped me read logs, propose hypotheses, design minimal tests, and write this comment. Commands were executed on my actual hardware, all log output is real and reproducible. Authorship note at the end.

---

Environment

Output of python -m torch.utils.collect_env — production state (TP=1, ZE_AFFINITY_MASK=0)

PyTorch version: 2.11.0+xpu
Is debug build: False
CUDA used to build PyTorch: None
ROCM used to build PyTorch: N/A

OS: CachyOS (x86_64)
GCC version: (GCC) 16.1.1 20260430
Clang version: 22.1.5
CMake version: version 4.3.2
Libc version: glibc-2.43

Python version: 3.12.13 (main, May 13 2026, 19:47:40) [GCC 16.1.1 20260430] (64-bit runtime)
Python platform: Linux-7.0.6-1-cachyos-x86_64-with-glibc2.43
Is CUDA available: False
Is XPU available: True
XPU used to build PyTorch: 20250302
Intel GPU driver version:

Intel GPU models onboard:
N/A
Intel GPU models detected:
* [0] _XpuDeviceProperties(name='Intel(R) Arc(TM) Pro B70 Graphics',
    platform_name='Intel(R) oneAPI Unified Runtime over Level-Zero',
    type='gpu', device_id=0xE223,
    uuid=868023e2-0000-0000-0300-000000000000,
    driver_version='1.15.38308', total_memory=32656MB,
    local_mem_size=128KB, max_compute_units=256, gpu_eu_count=256,
    gpu_subslice_count=32, max_work_group_size=1024,
    max_num_sub_groups=64, sub_group_sizes=[16 32],
    has_fp16=1, has_fp64=1, has_atomic64=1)

CPU:
AMD Ryzen Threadripper PRO 9965WX 24-Cores (48 threads)
L3 Cache: 128 MiB (4 instances)
NUMA: 1 node, CPU(s) 0-47

Versions of relevant libraries:
[pip3] dpcpp-cpp-rt==2025.3.2
[pip3] impi-rt==2021.17.2
[pip3] intel-cmplr-lib-rt==2025.3.2
[pip3] intel-cmplr-lib-ur==2025.3.2
[pip3] intel-opencl-rt==2025.3.2
[pip3] intel-openmp==2025.3.2
[pip3] intel-pti==0.16.0
[pip3] intel-sycl-rt==2025.3.2
[pip3] mkl==2025.3.1
[pip3] numpy==2.3.5
[pip3] oneccl==2021.17.2
[pip3] oneccl-devel==2021.17.2
[pip3] onemkl-license==2025.3.1
[pip3] onemkl-sycl-blas==2025.3.1
[pip3] onemkl-sycl-dft==2025.3.1
[pip3] onemkl-sycl-lapack==2025.3.1
[pip3] onemkl-sycl-rng==2025.3.1
[pip3] onemkl-sycl-sparse==2025.3.1
[pip3] pytorch-triton-xpu==3.6.0+git225cdbde
[pip3] tbb==2022.3.1
[pip3] tcmlib==1.4.1
[pip3] torch==2.11.0+xpu
[pip3] torchaudio==2.11.0+xpu
[pip3] torchvision==0.26.0+xpu
[pip3] triton-xpu==3.7.0
[pip3] umf==1.0.3

Only one GPU is visible in this state because ZE_AFFINITY_MASK=0 is set persistently in my shell config (pins TP=1 to the compute card by default). PCI 03:00.0 (UUID prefix 0300) is the compute card; PCI f3:00.0 (UUID prefix f300) is the display card carrying KDE Plasma 6.6 / Wayland.

Output of python -m torch.utils.collect_env — with Multi-GPU env triple (both cards visible)

After:

export ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE
export ZE_AFFINITY_MASK=0,1
export ONEAPI_DEVICE_SELECTOR=level_zero:*

python -c "import torch; print('device_count:', torch.xpu.device_count())"
# → device_count: 2

The relevant diff in collect_env is in the Intel GPU models detected block:

Intel GPU models detected:
* [0] _XpuDeviceProperties(name='Intel(R) Arc(TM) Pro B70 Graphics', ...
    uuid=868023e2-0000-0000-0300-000000000000, ...)   ← PCI 03:00.0 (compute)
* [1] _XpuDeviceProperties(name='Intel(R) Arc(TM) Pro B70 Graphics', ...
    uuid=868023e2-0000-0000-f300-000000000000, ...)   ← PCI f3:00.0 (display)

All other fields are identical to the production-state output above.

Note: Without this env triple, torch.xpu.device_count() returns 1 on this host even though both cards are present, healthy, and listed by lspci and sycl-ls. This is consistent with the upthread observation that ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE + ZE_AFFINITY_MASK=0,1 + ONEAPI_DEVICE_SELECTOR=level_zero:* is required for PyTorch to enumerate both Battlemage cards.

Additional context not in collect_env:

• 2× Intel Arc Pro B70 32GB, Battlemage G31, device 8086:e223
• ASUS Pro WS WRX90E SAGE SE motherboard
• PCIe 5.0 x16 native per card (different from original report's x8)
• 128 GiB system RAM
• GuC firmware: bmg_guc_70.bin v70.60.0
• HuC firmware: bmg_huc.bin v8.2.10
• DMC firmware: bmg_dmc.bin v2.6
• intel-compute-runtime: 26.18.38308.1 (system package, not in pip)
• level-zero-loader/headers: 1.28.2 (system package, not in pip)
• vLLM: 0.20.2rc1.dev345+g768f4a6f2 (source-built from main, not Intel pre-built image)
• KDE Plasma 6.6 / Wayland on the display GPU

Model used in failing TP=2 / PP=2 tests: RedHatAI/gemma-4-26B-A4B-it-FP8-Dynamic (Gemma 4 26B MoE, pre-quantized FP8) — different MoE architecture than the original report's Qwen3-30B-A3B.

---

Summary of findings

1. Device discovery requires the env triple. Without ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE + ZE_AFFINITY_MASK=0,1 + ONEAPI_DEVICE_SELECTOR=level_zero:*, torch.xpu.device_count() returns 1 even with both cards healthy. With the triple, it correctly returns 2. The collect_env diff above shows this directly.

2. TP=2 fails as in the original report — same GP fault + xe BCS engine reset signature. I did not re-stress the hardware once the symptom was confirmed.

3. PP=2 reaches further than TP=2. With the env triple, both workers initialize successfully (world_size=2, rank=0/1, backend=xccl), TCPStore rendezvous completes, eight ProcessGroupXCCL instances are created, four CPU-side collectives (BROADCAST × 2, BARRIER, ALLREDUCE) all run cleanly. The crash happens at vllm/v1/worker/xpu_worker.py:105 on the first XPU memory touch via torch.distributed.all_reduce(torch.zeros(1).xpu()).

4. The failure is a phantom OOM. PyTorch reports:

   torch.OutOfMemoryError: XPU out of memory. Tried to allocate 2.00 MiB.
   GPU 0 has a total capacity of 31.89 GiB of which 31.82 GiB is free.
   Of the allocated memory 0 bytes is allocated by PyTorch,
   and 0 bytes is reserved by PyTorch but unallocated.
   

   This is not memory pressure — the allocator has never allocated anything yet, and both cards are effectively idle.

5. The bug reproduces without vLLM in a 15-line minimal script. Using just torchrun --nproc-per-node=2 and dist.all_reduce(torch.zeros(1).xpu()) after cpu:gloo,xpu:xccl backend init produces the same OOM on first XPU touch.

6. After enough failed multi-GPU initialization attempts, the bug reproduces in a single-process script (no torchrun, no distributed) — python -c "import torch; torch.zeros(1).xpu()" crashes with the same phantom OOM. Recovery requires a full reboot. No software-level reset I tried (process kills, torch.xpu.empty_cache(), killing xpumd, clearing /tmp/torchinductor_* and ~/.triton/cache) recovered the card state. A simple reboot fully restores the working baseline (TP=1 inference at ~15–28 tok/s warm).

---

What I systematically ruled out

To narrow the root cause, the following hypotheses were tested and falsified:

Hypothesis	Test	Result
PyTorch XPU caching allocator is the cause	PYTORCH_NO_XPU_MEMORY_CACHING=1	Same crash, byte-identical trace
XCCL backend is the cause	Same code with cpu:gloo,xpu:xccl dual backend; gloo CPU allreduce passes (result = 2.0); XPU tensor creation alone (no allreduce) still crashes	XCCL is not involved in the crash; failure happens at first XPU memory allocation, before any distributed op
Multi-process / process group race is the cause	Single-process script (no torchrun) reproduces same OOM (after damaged state)	Bug exists outside multi-process context
vLLM-specific init sequence is missing	Replicated full xpu_worker.py:init_device() sequence in standalone script: torch.accelerator.device_count(), torch.accelerator.set_device_index(), torch.accelerator.empty_cache(), torch.xpu.get_device_properties().total_memory	Same crash at the subsequent torch.zeros(1).xpu()
Intel xpumd daemon memory mapping conflicts	fuser -v /dev/dri/renderD129 showed xpumd had F...m (memory-mapped) on the compute card; sudo kill -9 <xpumd pid> to release it	Same crash, card still unusable
Triton/PyTorch compile caches are stale	rm -rf /tmp/torchinductor_* ~/.triton/cache ~/.cache/torch	Same crash
KDE Plasma / display GPU interference	Tested allocation on the compute GPU (no display)	Same crash
Hypothesized 15-minute mei_lb rebind cycle in dmesg indicated a driver recovery loop	Compared mei_lb event timestamps with snapper snapshot timestamps and my own desk-presence pattern	Falsified — the mei_lb rebinds correlate with Plasma's suspend/resume cycles during idle, not with a driver loop. This was a wrong hypothesis on Claude's side; I caught it from the timestamp pattern.

What likely remains

By elimination, the root cause appears to be in one of:

• PyTorch XPU memory manager below the caching allocator layer
• Level Zero V2 runtime state corruption after failed XPU allocations
• intel-compute-runtime internal bookkeeping after engine reset events
• xe kernel driver persistent state after GuC firmware engine resets

The fact that only a full reboot recovers the working state — combined with dmesg showing xe engine reset events from kwin_wayland even during single-card baseline runs — suggests the issue is in a layer below PyTorch where state cannot be reset from userspace.

---

Reproduction trace (PP=2 with TORCH_DISTRIBUTED_DEBUG=DETAIL)

(EngineCore) DP group leader: world_size=2, local_world_size=2
(Worker rank=0) world_size=2 rank=0 local_rank=0 backend=xccl
(Worker rank=1) world_size=2 rank=1 local_rank=1 backend=xccl
[c10d] TCPStore server on port 58757
[c10d] Both clients connected (ports 50764, 50776)
ProcessGroupXCCL PG ID 0 GUID 0 (size=2) initialized on both ranks
ProcessGroupXCCL PG ID 1 GUID 1 (size=2) initialized on both ranks
[Rank 0/1] Running CPU BROADCAST [1] Long           ✓
[Rank 0/1] Running CPU BROADCAST [27] Byte           ✓
[Rank 0/1] Running CPU BARRIER                       ✓
[Rank 0/1] Running CPU ALLREDUCE [2] Int             ✓
ProcessGroupXCCL PG ID 2–7 (mix size=1 and size=2) initialized
[Rank 0] assigned as DP rank 0, PP rank 0, TP rank 0, EP rank 0
[Rank 0] crash at xpu_worker.py:105
  torch.distributed.all_reduce(torch.zeros(1).xpu())
  → torch.OutOfMemoryError: XPU out of memory.
    Tried to allocate 2.00 MiB.
    GPU 0 has a total capacity of 31.89 GiB of which 31.82 GiB is free.
[Rank 1] identical crash on GPU 1 (29.03 GiB free)

Minimal repro without vLLM

# reduce_all_test.py
import os, torch, torch.distributed as dist

rank = int(os.environ["RANK"])
local_rank = int(os.environ["LOCAL_RANK"])
world_size = int(os.environ["WORLD_SIZE"])

torch.xpu.set_device(local_rank)
dist.init_process_group(
    backend="cpu:gloo,xpu:xccl",
    rank=rank, world_size=world_size,
)

# CPU allreduce via gloo — works
cpu_t = torch.ones(1)
dist.all_reduce(cpu_t)   # ✓ result = 2.0

# XPU tensor creation — crashes (in damaged state)
xpu_t = torch.zeros(1).xpu()   # ✗ OutOfMemoryError, 2 MiB / 30 GiB free

Run with:

export ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE
export ZE_AFFINITY_MASK=0,1
export ONEAPI_DEVICE_SELECTOR=level_zero:*
torchrun --nproc-per-node=2 reduce_all_test.py

Single-process repro (in damaged state)

import torch
print(torch.xpu.device_count())   # 2
torch.xpu.set_device(0)
torch.xpu.empty_cache()
torch.xpu.synchronize()
t = torch.zeros(1).xpu()   # ✗ OutOfMemoryError

---

Working configuration on this host

• TP=1 only, with vllm serve ... --max-model-len 16384 --gpu-memory-utilization 0.90 --kv-cache-dtype fp8 --enforce-eager --attention-backend TRITON_ATTN
• Stable at ~15–28 tok/s warm on Gemma 4 26B FP8 (MoE, 128 experts top-8 active)
• Multi-GPU experiments require a reboot afterward to restore card usability

---

Files available on request

I can provide:

• Full vLLM startup logs for all PP=2 / TP=2 / single-process attempts
• Complete dmesg -T spanning the test session
• lspci -nn and lspci -vvv for both cards
• sycl-ls and device-property dumps
• pip freeze of the venv
• The minimal repro scripts (reduce_all_test.py, verification.py, replicate_vllm_init.py)

---

A note on this comment's authorship

AI and human collaboration:

• Me (the human): Hardware setup, all command execution, providing logs, deciding which hypotheses were worth pursuing. Also providing context that mattered — for example, when Claude proposed that recurring mei_lb rebind events in dmesg might indicate a driver recovery loop, I pointed out they correlated with the times I was away from the desk, suggesting Plasma's suspend/resume cycle as the more plausible cause. We verified the hypothesis against the data and it held.
• Claude: Helping parse stack traces, suggesting which environment variables and tests would isolate each layer of the stack, drafting this comment.

As said: Commands and observations from my machine. Methods reproductible. It's my first report here on GH and I hope that it will help narrowing down and fixing issues.

If anything here is unclear or you'd like a specific additional test, I'm happy to run it. Thank you for maintaining vLLM XPU support — Battlemage on a non-Ubuntu host is early-adopter territory and I really appreciate the work of everyone involved! :).

[Neppord]

Im having similar issues building the docker image from main, can only run TP=1, have two b70.

[F3zz1k]

Edit: correction, rebooting got it working and I ended up with the same 101t/s with TP=2... Seems the drivers get messed up after you hit an error. I'll post my working config below.

Original post:

Figured I'd drop this here as an extra data point:

Same-die TP=2 crashes on 4x Arc Pro B70 sharing one PCIe root complex on 1950x
Repro on 2× B70 (Battlemage), intel/llm-scaler-vllm:0.14.0-b8.2.1, env stack from this thread,
Qwen3.6-27B-AWQ, TP=2:

• Different root complexes: 101 t/s @ N=32, 0 errors. Very slow but works
• Same root complex: crash at apply_top_k_top_p → logits.sort:
• RuntimeError: level_zero backend failed with error: 2147483646 (UR_RESULT_ERROR_UNKNOWN)
• Both cards in dmesg at the same instant:
• xe NNNN:NN:00.0: [drm] VM worker error: -16
• xe NNNN:NN:00.0: [drm] VM worker error: -16

Ruled out: host RAM (capped via UR_L0_USM_ALLOCATOR, ~46 GiB free at crash, would OOM at >64GB otherwise), Xwayland (gdm stopped), PCIe link errors but for unrelated device. Single-card works fine at 218.19 t/s aggregate at N=32 but limited to 93k context.
CCL_TOPO_P2P_ACCESS=0 shifts the crash earlier but doesn't fix it. The [drm] VM worker error site
matches intel/compute-runtime#922.

[F3zz1k]

My working config with B70s:

Hardware: 2× Intel Arc Pro B70, Linux ≥ 7.1 with xe driver, intel-compute-runtime 26.18.38308.1,
Docker, ≥64 GiB RAM.

docker pull intel/llm-scaler-vllm:0.14.0-b8.2.1

launch_b70_tp2.sh:

#!/usr/bin/env bash
set -euo pipefail
ZE_AFFINITY="${ZE_AFFINITY:-0,1}"  # two SYCL device indices
PORT="${PORT:-9999}"

SED=$(base64 -w0 <<<'sed -i "s|image_processor.max_pixels or
image_processor.size\[\"longest_edge\"\]|getattr(image_processor, \"max_pixels\", None) or
image_processor.size[\"longest_edge\"]|"
/usr/local/lib/python3.12/dist-packages/vllm/model_executor/models/qwen2_vl.py')

docker rm -f vllm-b70-tp2 2>/dev/null || true
exec docker run --rm --name=vllm-b70-tp2 --privileged \
  --group-add="$(getent group video|cut -d: -f3)" --group-add="$(getent group render|cut -d: -f3)" \
  --device=/dev/dri:/dev/dri --pid=host --uts=host --userns=host --ipc=host --network=host \
  -v /dev/dri/by-path:/dev/dri/by-path -v /dev/shm:/dev/shm --shm-size=32g \
  -v "$HOME/.cache/huggingface":/home/sdp/.cache/huggingface -e HF_HOME=/home/sdp/.cache/huggingface \
  -e NEOReadDebugKeys=1 -e RenderCompressedBuffersEnabled=0 \
  -e EnableHostUsmAllocationPool=0 -e EnableUsmPoolLazyInit=1 \
  -e CCL_ZE_TMP_BUF_SIZE=134217728 \
  -e 'UR_L0_USM_ALLOCATOR=1;256M;host:1M,4,64K;shared:1M,4,64K;device:1M,4,64K' \
  -e ZE_AFFINITY_MASK="$ZE_AFFINITY" -e ZE_FLAT_DEVICE_HIERARCHY=COMPOSITE \
  -e SYCL_UR_USE_LEVEL_ZERO_V2=0 -e CCL_ENABLE_SYCL_KERNELS=1 \
  -e CCL_ALLREDUCE=ring -e CCL_TOPO_FABRIC_VERTEX_CONNECTION_CHECK=0 \
  -e TORCH_LLM_ALLREDUCE=1 -e VLLM_USE_V1=1 -e VLLM_XPU_ENABLE_XPU_GRAPH=1 \
  -e UR_L0_USE_IMMEDIATE_COMMANDLISTS=0 -e SED="$SED" \
  --entrypoint bash intel/llm-scaler-vllm:0.14.0-b8.2.1 \
  -c 'echo "$SED"|base64 -d|bash && exec vllm serve QuantTrio/Qwen3.6-27B-AWQ \
      --port '"$PORT"' --served-model-name qwen3.6-27b-awq-tp2 \
      --tensor-parallel-size 2 --max-model-len 16384 \
      --max-num-batched-tokens 2096 --max-num-seqs 32 \
      --gpu-memory-utilization 0.75 --enforce-eager --kv-cache-dtype fp8 --swap-space 0 \
      --reasoning-parser qwen3 --enable-auto-tool-choice --tool-call-parser hermes \
      --allow-deprecated-quantization'

Run + smoke test:

  ./launch_b70_tp2.sh &
  until curl -sf http://127.0.0.1:9999/v1/models >/dev/null; do sleep 5; done
  curl -sS -X POST http://127.0.0.1:9999/v1/completions -H "Content-Type: application/json" \
    -d '{"model":"qwen3.6-27b-awq-tp2","prompt":"def add(a,b):","max_tokens":40}'

Expected: ~4 min init, ~5 t/s sequential, ~101 t/s aggregate at N=32, ~13 GiB host RAM peak.

Caveat: if init fails with VM worker error: -16 / UR_RESULT_ERROR_UNKNOWN after retries — reboot. xe
driver state accumulates across crashed TP=2 starts.

[yisustc]

Hi @Zumbasam, we validate your tests with latest vllm docker image from vllm ci build, all are working well, please using this image public.ecr.aws/q9t5s3a7/vllm-ci-test-repo:b7c5baf63d5f95d18493a7fae3c75fc4853732a7-xpu and docker pull it directly to try again. And here is the environment we are testing, you can take as a reference.

Collecting environment information...

   System Info

==============================
OS : Ubuntu 24.04.3 LTS (x86_64)
GCC version : (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0
Clang version : Could not collect
CMake version : version 4.3.2
Libc version : glibc-2.39

==============================
PyTorch Info

PyTorch version : 2.11.0+xpu
Is debug build : False
CUDA used to build PyTorch : None
ROCM used to build PyTorch : N/A
XPU used to build PyTorch : 20250302

==============================
Python Environment

Python version : 3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0] (64-bit runtime)
Python platform : Linux-6.17.0-22-generic-x86_64-with-glibc2.39

==============================
Intel XPU / GPU Info

Is XPU available : True
XPU runtime version : 20250302
Intel GPU models :
GPU 0: Intel(R) Graphics [0xe223]
GPU 1: Intel(R) Graphics [0xe223]

--Compile time--
oneAPI compiler version : 2025.3.3
SYCL compiler build : 2025.3.3.20260319
oneCCL version : 2021.17.2-5

--Runtime--
Intel Graphics Compiler (IGC): 2.24.8
Intel GMM (libigdgmm) : 22.8.2
Level Zero loader version : 1.26.0
Level Zero driver version : 25.48.36300.8-0
vLLM XPU kernels version : 0.1.9

==============================
CPU Info

Architecture: x86_64
CPU op-mode(s): 32-bit, 64-bit
Address sizes: 46 bits physical, 57 bits virtual
Byte Order: Little Endian
CPU(s): 192
On-line CPU(s) list: 0-191
Vendor ID: GenuineIntel
Model name: Intel(R) Xeon(R) Platinum 8468V
CPU family: 6
Model: 143
Thread(s) per core: 2
Core(s) per socket: 48
Socket(s): 2
Stepping: 8
CPU(s) scaling MHz: 29%
CPU max MHz: 3800.0000
CPU min MHz: 800.0000
BogoMIPS: 4800.00
Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid dca sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb cat_l3 cat_l2 cdp_l3 intel_ppin cdp_l2 ssbd mba ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb intel_pt avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local split_lock_detect user_shstk avx_vnni avx512_bf16 wbnoinvd dtherm ida arat pln pts hfi vnmi avx512vbmi umip pku ospke waitpkg avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq la57 rdpid bus_lock_detect cldemote movdiri movdir64b enqcmd fsrm md_clear serialize tsxldtrk pconfig arch_lbr ibt amx_bf16 avx512_fp16 amx_tile amx_int8 flush_l1d arch_capabilities
Virtualization: VT-x
L1d cache: 4.5 MiB (96 instances)
L1i cache: 3 MiB (96 instances)
L2 cache: 192 MiB (96 instances)
L3 cache: 195 MiB (2 instances)
NUMA node(s): 2
NUMA node0 CPU(s): 0-47,96-143
NUMA node1 CPU(s): 48-95,144-191
Vulnerability Gather data sampling: Not affected
Vulnerability Ghostwrite: Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit: Not affected
Vulnerability L1tf: Not affected
Vulnerability Mds: Not affected
Vulnerability Meltdown: Not affected
Vulnerability Mmio stale data: Not affected
Vulnerability Old microcode: Vulnerable
Vulnerability Reg file data sampling: Not affected
Vulnerability Retbleed: Not affected
Vulnerability Spec rstack overflow: Not affected
Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1: Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2: Mitigation; Enhanced / Automatic IBRS; IBPB conditional; PBRSB-eIBRS SW sequence; BHI BHI_DIS_S
Vulnerability Srbds: Not affected
Vulnerability Tsa: Not affected
Vulnerability Tsx async abort: Not affected
Vulnerability Vmscape: Mitigation; IBPB before exit to userspace

==============================
Versions of relevant libraries

[pip3] flake8==7.3.0
[pip3] mypy==2.1.0
[pip3] mypy_extensions==1.1.0
[pip3] numpy==2.2.6
[pip3] pyzmq==27.1.0
[pip3] sentence-transformers==5.3.0
[pip3] torch==2.11.0+xpu
[pip3] torchaudio==2.11.0+xpu
[pip3] torchvision==0.26.0+xpu
[pip3] transformers==5.5.3
[pip3] triton-xpu==3.7.0
[conda] Could not collect

==============================
vLLM Info

ROCM Version : Could not collect
vLLM Version : 0.22.1rc1.dev190+gb7c5baf63 (git sha: b7c5baf)
vLLM Build Flags:
CUDA Archs: Not Set; ROCm: Disabled; XPU: Enabled
GPU Topology:
Could not collect

[chriswagner-ai]

Got it working bare metal - have a look, might be of help to someone: https://github.com/chriswagner-ai/intel-arc-b70-vllm-multi-gpu