[GSD-12973] Arc Pro B70 4-GPU: SYCL/Level Zero peer memcpy corrupts data and causes DEVICE_LOST

[arnoldthepillow-ui]

Arc Pro B70 4-GPU: SYCL/Level Zero peer memcpy corrupts data and causes DEVICE_LOST

Summary

On a 4x Intel Arc Pro B70 system, direct SYCL/Level Zero device-to-device copies between GPUs corrupt data. The runtime reports peer access is available via ext_oneapi_can_access_peer(...), but actual queue.memcpy(dst_device_ptr, src_device_ptr, size) returns incorrect bytes. Larger copies can trigger UR_RESULT_ERROR_DEVICE_LOST.

Host-staged copies (GPU -> host RAM -> GPU) work correctly.

This breaks llama.cpp multi-GPU inference unless cross-device copies are forced through host staging. With direct dGPU-to-dGPU copies enabled, model output becomes gibberish.

Hardware / Topology

• 4x Intel Arc Pro B70 GPUs
• Ubuntu 26.04 LTS
• Kernel: 7.0.0-27-generic
• Driver: xe
• PCIe fabric / switch topology with four B70s attached under the same root complex

PCIe GPU mapping:

0f:00.0 Intel Battlemage G31 / Arc Pro B70 -> renderD128
13:00.0 Intel Battlemage G31 / Arc Pro B70 -> renderD129
19:00.0 Intel Battlemage G31 / Arc Pro B70 -> renderD130
1d:00.0 Intel Battlemage G31 / Arc Pro B70 -> renderD131

Relevant topology excerpt:

03.1-[09-1f]----00.0-[0a-1f]
  +-00.0-[0b-14]
  | +-10.0-[0d-10] -> 0f:00.0 Intel Battlemage G31
  | +-18.0-[11-14] -> 13:00.0 Intel Battlemage G31
  +-04.0-[15-1e]
  | +-00.0-[17-1a] -> 19:00.0 Intel Battlemage G31
  | +-08.0-[1b-1e] -> 1d:00.0 Intel Battlemage G31

Full lspci -tv output is available in lspci-tv-b70-bug.txt.

Software Versions

intel-igc-core-2:    2.36.3
intel-igc-opencl-2:  2.36.3
intel-ocloc:         26.22.38646.4-0
intel-opencl-icd:    26.22.38646.4-0
libigdgmm12:         22.10.0
libze-intel-gpu1:    26.22.38646.4-0
Level Zero driver:   1.15.38646+4

sycl-ls sees all four GPUs:

[level_zero:gpu][level_zero:0] Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:1] Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:2] Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:3] Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]

Reproduction

I used a minimal SYCL repro that:

1. Enumerates Level Zero GPU devices.
2. Allocates a source USM buffer on GPU A.
3. Allocates a destination USM buffer on GPU B.
4. Fills the source with deterministic bytes.
5. Runs direct queue.memcpy(dst, src, size) on the destination queue.
6. Copies destination back to host and compares bytes.

Run command:

source /opt/intel/oneapi/setvars.sh
ONEAPI_DEVICE_SELECTOR=level_zero:gpu \
ZES_ENABLE_SYSMAN=1 \
UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 \
./sycl_peer_copy_probe

Failure example:

level_zero_gpu_count=4
gpu[0]=Intel(R) Arc(TM) Pro B70 Graphics
gpu[1]=Intel(R) Arc(TM) Pro B70 Graphics
gpu[2]=Intel(R) Arc(TM) Pro B70 Graphics
gpu[3]=Intel(R) Arc(TM) Pro B70 Graphics
pair 0->1 can_access=1
FAIL 0->1 size=1 first_bad=0 expected=0xaf actual=0xff
FAIL 0->1 size=2 first_bad=0 expected=0xc1 actual=0xff
FAIL 0->1 size=3 first_bad=0 expected=0xce actual=0xff
FAIL 0->1 size=4 first_bad=0 expected=0x62 actual=0xff
FAIL 0->1 size=7 first_bad=0 expected=0x21 actual=0xff
FAIL 0->1 size=16 first_bad=0 expected=0x9f actual=0xff
FAIL 0->1 size=31 first_bad=0 expected=0x13 actual=0xff
FAIL 0->1 size=64 first_bad=0 expected=0x35 actual=0xff
FAIL 0->1 size=255 first_bad=0 expected=0x29 actual=0xff
FAIL 0->1 size=4096 first_bad=0 expected=0x8d actual=0xff
FAIL 0->1 size=65536 first_bad=0 expected=0x76 actual=0xff
FAIL 0->1 size=1048576 first_bad=0 expected=0x6c actual=0xff
FAIL 0->1 size=16777216 exception=level_zero backend failed with error: 20 (UR_RESULT_ERROR_DEVICE_LOST)
FAIL 0->1 size=16777216 first_bad=0 expected=0xa6 actual=0x00
Abort was called at 284 line in file:
../../neo/shared/source/os_interface/linux/drm_neo.cpp

After the DEVICE_LOST failure, sycl-ls may temporarily show only 3 of the 4 Level Zero GPUs until reboot.

Expected Behavior

If ext_oneapi_can_access_peer(...) returns true, direct peer copies between those GPU devices should either:

• copy bytes correctly, or
• fail cleanly before corrupting data.

Actual Behavior

Direct peer copies return corrupted bytes. Larger copies can cause UR_RESULT_ERROR_DEVICE_LOST.

Workaround

Host-staged copies work:

GPU A -> host RAM -> GPU B

In llama.cpp, forcing host-staged cross-device copies avoids gibberish output:

GGML_SYCL_DEV2DEV_MEMCPY=2

Impact

This prevents reliable multi-GPU inference on 4x Arc Pro B70 when using direct dGPU-to-dGPU copies. It removes the expected benefit of the PCIe fabric for peer transfers and requires slower host-staged copies for correctness.

Attachments / Extra Outputs

I can provide:

• sycl_peer_copy_probe.cpp
• sycl_peer_copy_shared_context_probe.cpp
• sycl-ls-b70-bug.txt
• llama-ls-sycl-device-b70-bug.txt
• lspci-tv-b70-bug.txt
• packages-b70-bug.txt
• uname-b70-bug.txt

[arnoldthepillow-ui]

Adding the repro sources and collected command outputs proactively.

sycl_peer_copy_probe.cpp

#include <sycl/sycl.hpp>

#include <cstdint>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <vector>

static uint64_t mix(uint64_t x) {
    x ^= x >> 30;
    x *= 0xbf58476d1ce4e5b9ULL;
    x ^= x >> 27;
    x *= 0x94d049bb133111ebULL;
    x ^= x >> 31;
    return x;
}

static void fill_pattern(std::vector<uint8_t> & v, int src, int dst, size_t size) {
    const uint64_t seed = 0x9e3779b97f4a7c15ULL ^ ((uint64_t) src << 32) ^ (uint64_t) dst ^ size;
    for (size_t i = 0; i < v.size(); ++i) {
        v[i] = (uint8_t) (mix(seed + i) & 0xff);
    }
}

int main() {
    std::cout.setf(std::ios::unitbuf);

    std::vector<sycl::device> devices;
    for (const auto & platform : sycl::platform::get_platforms()) {
        if (platform.get_backend() != sycl::backend::ext_oneapi_level_zero) {
            continue;
        }
        for (const auto & dev : platform.get_devices(sycl::info::device_type::gpu)) {
            devices.push_back(dev);
        }
    }

    std::cout << "level_zero_gpu_count=" << devices.size() << "\n";
    for (size_t i = 0; i < devices.size(); ++i) {
        std::cout << "gpu[" << i << "]=" << devices[i].get_info<sycl::info::device::name>() << "\n";
    }

    std::vector<std::unique_ptr<sycl::queue>> queues;
    for (const auto & dev : devices) {
        queues.emplace_back(new sycl::queue(dev));
    }

    const std::vector<size_t> sizes = {1, 2, 3, 4, 7, 16, 31, 64, 255, 4096, 65536, 1048576, 16 * 1048576};
    bool failed = false;

    for (size_t src_i = 0; src_i < devices.size(); ++src_i) {
        for (size_t dst_i = 0; dst_i < devices.size(); ++dst_i) {
            if (src_i == dst_i) {
                continue;
            }

            const bool can_access = devices[dst_i].ext_oneapi_can_access_peer(
                devices[src_i], sycl::ext::oneapi::peer_access::access_supported);
            std::cout << "pair " << src_i << "->" << dst_i << " can_access=" << can_access << "\n";

            for (const size_t size : sizes) {
                std::vector<uint8_t> expected(size);
                std::vector<uint8_t> actual(size, 0);
                fill_pattern(expected, (int) src_i, (int) dst_i, size);

                void * src = sycl::malloc_device(size, *queues[src_i]);
                void * dst = sycl::malloc_device(size, *queues[dst_i]);
                if (!src || !dst) {
                    std::cerr << "allocation failed size=" << size << "\n";
                    return 3;
                }

                queues[src_i]->memcpy(src, expected.data(), size).wait();
                queues[dst_i]->memset(dst, 0xa5, size).wait();

                try {
                    queues[dst_i]->memcpy(dst, src, size).wait();
                    queues[dst_i]->memcpy(actual.data(), dst, size).wait();
                } catch (const std::exception & e) {
                    std::cout << "FAIL " << src_i << "->" << dst_i << " size=" << size
                              << " exception=" << e.what() << "\n";
                    failed = true;
                }

                if (std::memcmp(expected.data(), actual.data(), size) != 0) {
                    size_t first = 0;
                    while (first < size && expected[first] == actual[first]) {
                        ++first;
                    }
                    std::cout << "FAIL " << src_i << "->" << dst_i << " size=" << size
                              << " first_bad=" << first
                              << " expected=0x" << std::hex << std::setw(2) << std::setfill('0') << (int) expected[first]
                              << " actual=0x" << std::setw(2) << (int) actual[first]
                              << std::dec << "\n";
                    failed = true;
                } else {
                    std::cout << "OK " << src_i << "->" << dst_i << " size=" << size << "\n";
                }

                sycl::free(src, *queues[src_i]);
                sycl::free(dst, *queues[dst_i]);
            }
        }
    }

    return failed ? 1 : 0;
}

sycl_peer_copy_shared_context_probe.cpp

#include <sycl/sycl.hpp>

#include <cstdint>
#include <cstring>
#include <iomanip>
#include <iostream>
#include <vector>

static uint64_t mix(uint64_t x) {
    x ^= x >> 30;
    x *= 0xbf58476d1ce4e5b9ULL;
    x ^= x >> 27;
    x *= 0x94d049bb133111ebULL;
    x ^= x >> 31;
    return x;
}

static void fill_pattern(std::vector<uint8_t> & v, int src, int dst, size_t size) {
    const uint64_t seed = 0x9e3779b97f4a7c15ULL ^ ((uint64_t) src << 32) ^ (uint64_t) dst ^ size;
    for (size_t i = 0; i < v.size(); ++i) {
        v[i] = (uint8_t) (mix(seed + i) & 0xff);
    }
}

int main() {
    std::cout.setf(std::ios::unitbuf);

    std::vector<sycl::device> devices;
    for (const auto & platform : sycl::platform::get_platforms()) {
        if (platform.get_backend() != sycl::backend::ext_oneapi_level_zero) {
            continue;
        }
        for (const auto & dev : platform.get_devices(sycl::info::device_type::gpu)) {
            devices.push_back(dev);
        }
    }

    std::cout << "level_zero_gpu_count=" << devices.size() << "\n";
    if (devices.size() < 2) {
        return 2;
    }

    sycl::context shared_context(devices);
    std::vector<std::unique_ptr<sycl::queue>> queues;
    for (const auto & dev : devices) {
        queues.emplace_back(new sycl::queue(shared_context, dev));
    }

    const std::vector<size_t> sizes = {1, 2, 3, 4, 7, 16, 31, 64, 255, 4096, 65536, 1048576, 16 * 1048576};
    bool failed = false;

    for (size_t src_i = 0; src_i < devices.size(); ++src_i) {
        for (size_t dst_i = 0; dst_i < devices.size(); ++dst_i) {
            if (src_i == dst_i) {
                continue;
            }

            const bool can_access = devices[dst_i].ext_oneapi_can_access_peer(
                devices[src_i], sycl::ext::oneapi::peer_access::access_supported);
            std::cout << "pair " << src_i << "->" << dst_i << " can_access=" << can_access << "\n";

            for (const size_t size : sizes) {
                std::vector<uint8_t> expected(size);
                std::vector<uint8_t> actual(size, 0);
                fill_pattern(expected, (int) src_i, (int) dst_i, size);

                void * src = sycl::malloc_device(size, devices[src_i], shared_context);
                void * dst = sycl::malloc_device(size, devices[dst_i], shared_context);
                if (!src || !dst) {
                    std::cerr << "allocation failed size=" << size << "\n";
                    return 3;
                }

                queues[src_i]->memcpy(src, expected.data(), size).wait();
                queues[dst_i]->memset(dst, 0xa5, size).wait();
                queues[dst_i]->memcpy(dst, src, size).wait();
                queues[dst_i]->memcpy(actual.data(), dst, size).wait();

                if (std::memcmp(expected.data(), actual.data(), size) != 0) {
                    size_t first = 0;
                    while (first < size && expected[first] == actual[first]) {
                        ++first;
                    }
                    std::cout << "FAIL " << src_i << "->" << dst_i << " size=" << size
                              << " first_bad=" << first
                              << " expected=0x" << std::hex << std::setw(2) << std::setfill('0') << (int) expected[first]
                              << " actual=0x" << std::setw(2) << (int) actual[first]
                              << std::dec << "\n";
                    failed = true;
                } else {
                    std::cout << "OK " << src_i << "->" << dst_i << " size=" << size << "\n";
                }

                sycl::free(src, shared_context);
                sycl::free(dst, shared_context);
            }
        }
    }

    return failed ? 1 : 0;
}

sycl-ls-b70-bug.txt

[level_zero:gpu][level_zero:0] Intel(R) oneAPI Unified Runtime over Level-Zero V2, Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:1] Intel(R) oneAPI Unified Runtime over Level-Zero V2, Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:2] Intel(R) oneAPI Unified Runtime over Level-Zero V2, Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[level_zero:gpu][level_zero:3] Intel(R) oneAPI Unified Runtime over Level-Zero V2, Intel(R) Arc(TM) Pro B70 Graphics 20.2.0 [1.15.38646+4]
[opencl:cpu][opencl:0] Intel(R) OpenCL, AMD Ryzen 5 5600 6-Core Processor               OpenCL 3.0 (Build 0) [2026.21.3.0.31_160000]
[opencl:gpu][opencl:1] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) Pro B70 Graphics OpenCL 3.0 NEO  [26.22.38646.4]
[opencl:gpu][opencl:2] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) Pro B70 Graphics OpenCL 3.0 NEO  [26.22.38646.4]
[opencl:gpu][opencl:3] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) Pro B70 Graphics OpenCL 3.0 NEO  [26.22.38646.4]
[opencl:gpu][opencl:4] Intel(R) OpenCL Graphics, Intel(R) Arc(TM) Pro B70 Graphics OpenCL 3.0 NEO  [26.22.38646.4]

llama-ls-sycl-device-b70-bug.txt

Found 4 SYCL devices:
|  |                   |                                       |       |Max    |        |Max  |Global |                     |
|  |                   |                                       |       |compute|Max work|sub  |mem    |                     |
|ID|        Device Type|                                   Name|Version|units  |group   |group|size   |       Driver version|
|--|-------------------|---------------------------------------|-------|-------|--------|-----|-------|---------------------|
| 0| [level_zero:gpu:0]|             Intel Arc Pro B70 Graphics|   20.2|    256|    1024|   32| 34242M|         1.15.38646+4|
| 1| [level_zero:gpu:1]|             Intel Arc Pro B70 Graphics|   20.2|    256|    1024|   32| 34242M|         1.15.38646+4|
| 2| [level_zero:gpu:2]|             Intel Arc Pro B70 Graphics|   20.2|    256|    1024|   32| 34242M|         1.15.38646+4|
| 3| [level_zero:gpu:3]|             Intel Arc Pro B70 Graphics|   20.2|    256|    1024|   32| 34242M|         1.15.38646+4|
SYCL Optimization Feature:
|ID|        Device Type|Reorder|
|--|-------------------|-------|
| 0| [level_zero:gpu:0]|      Y|
| 1| [level_zero:gpu:1]|      Y|
| 2| [level_zero:gpu:2]|      Y|
| 3| [level_zero:gpu:3]|      Y|

lspci-tv-b70-bug.txt

-[0000:00]-+-00.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse Root Complex
           +-00.2  Advanced Micro Devices, Inc. [AMD] Starship/Matisse IOMMU
           +-01.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-01.1-[01]----00.0  Sandisk Corp WD_BLACK SN7100/WD PC SN7100S M.2 2280 NVMe SSD (DRAM-less)
           +-01.2-[02-08]--+-00.0  Advanced Micro Devices, Inc. [AMD] 500 Series Chipset USB 3.1 XHCI Controller
           |               +-00.1  Advanced Micro Devices, Inc. [AMD] 500 Series Chipset SATA Controller
           |               \-00.2-[03-08]--+-00.0-[04-06]----00.0-[05-06]----00.0-[06]--+-00.0  Advanced Micro Devices, Inc. [AMD/ATI] Navi 32 [Radeon RX 7700 XT / 7800 XT]
           |                               |                                            \-00.1  Advanced Micro Devices, Inc. [AMD/ATI] Navi 31 HDMI/DP Audio
           |                               +-08.0-[07]----00.0  Realtek Semiconductor Co., Ltd. RTL8852BE PCIe 802.11ax Wireless Network Controller
           |                               \-09.0-[08]----00.0  Realtek Semiconductor Co., Ltd. RTL8125 2.5GbE Controller
           +-02.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-03.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-03.1-[09-1f]----00.0-[0a-1f]--+-00.0-[0b-14]----00.0-[0c-14]--+-10.0-[0d-10]----00.0-[0e-10]--+-01.0-[0f]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]
           |                               |                               |                               \-02.0-[10]----00.0  Intel Corporation Device e2f7
           |                               |                               \-18.0-[11-14]----00.0-[12-14]--+-01.0-[13]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]
           |                               |                                                               \-02.0-[14]----00.0  Intel Corporation Device e2f7
           |                               +-04.0-[15-1e]----00.0-[16-1e]--+-00.0-[17-1a]----00.0-[18-1a]--+-01.0-[19]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]
           |                               |                               |                               \-02.0-[1a]----00.0  Intel Corporation Device e2f7
           |                               |                               \-08.0-[1b-1e]----00.0-[1c-1e]--+-01.0-[1d]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]
           |                               |                                                               \-02.0-[1e]----00.0  Intel Corporation Device e2f7
           |                               \-1c.0-[1f]----00.0  Broadcom / LSI PEX880xx PCIe Gen 4 Switch
           +-04.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-05.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-07.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-07.1-[20]----00.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Function
           +-08.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse PCIe Dummy Host Bridge
           +-08.1-[21]--+-00.0  Advanced Micro Devices, Inc. [AMD] Starship/Matisse Reserved SPP
           |            +-00.1  Advanced Micro Devices, Inc. [AMD] Starship/Matisse Cryptographic Coprocessor PSPCPP
           |            +-00.3  Advanced Micro Devices, Inc. [AMD] Matisse USB 3.0 Host Controller
           |            \-00.4  Advanced Micro Devices, Inc. [AMD] Starship/Matisse HD Audio Controller
           +-14.0  Advanced Micro Devices, Inc. [AMD] FCH SMBus Controller
           +-14.3  Advanced Micro Devices, Inc. [AMD] FCH LPC Bridge
           +-18.0  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 0
           +-18.1  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 1
           +-18.2  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 2
           +-18.3  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 3
           +-18.4  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 4
           +-18.5  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 5
           +-18.6  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 6
           \-18.7  Advanced Micro Devices, Inc. [AMD] Matisse/Vermeer Data Fabric: Device 18h; Function 7

packages-b70-bug.txt

+-- Desired=Unknown/Install/Remove/Purge/Hold
|+- Status=Not/Inst/Conf-files/Unpacked/halF-conf/Half-inst/trig-aWait/Trig-pend
||+ Err?=(none)/Reinst-required (Status,Err: uppercase=bad)
||| Name                          Version                       Architecture Description
+++-=============================-=============================-============-==========================================================
ii  intel-igc-core-2              2.36.3                        amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-igc-opencl-2            2.36.3                        amd64        Intel(R) Graphics Compiler for OpenCL(TM)
ii  intel-ocloc                   26.22.38646.4-0               amd64        Tool for managing Intel Compute GPU device binary format
ii  intel-opencl-icd              26.22.38646.4-0               amd64        Intel graphics compute runtime for OpenCL
ii  libigdgmm12:amd64             22.10.0                       amd64        Intel Graphics Memory Management Library -- shared library
ii  libze-intel-gpu1              26.22.38646.4-0               amd64        Intel(R) Graphics Compute Runtime for oneAPI Level Zero.
ii  linux-firmware                20260319.git217ca6e4.1ubuntu  all          Firmware for Linux kernel drivers (metapackage)
ii  linux-firmware-intel-graphics 20260319.git217ca6e4-0ubuntu2 all          Firmware for Intel graphics and IPU and VSC processors
ii  linux-image-generic           7.0.0-27.27                   amd64        Generic Linux kernel image

uname-b70-bug.txt

Linux shockwave 7.0.0-27-generic #27-Ubuntu SMP PREEMPT_DYNAMIC Thu Jun 18 19:13:49 UTC 2026 x86_64 GNU/Linux