[GSD-13429] p2p copy will crash if peer has blocks mapped non-contiguously.

[stolk]

Pre-submission Checklist

• I am using the latest GPU driver version (releases)
• I have searched for similar issues and found none

GPU Hardware

Intel Arc Pro B70 (x2)

DRI Devices Information

$ ls -ls /dev/dri/*
0 crw-rw---- 1 root video  226,   0 Sep 10 17:16 /dev/dri/card0
0 crw-rw---- 1 root video  226,   1 Sep  3 13:12 /dev/dri/card1
0 crw-rw---- 1 root video  226,   2 Sep  3 13:12 /dev/dri/card2
0 crw-rw---- 1 root render 226, 128 Sep  3 13:12 /dev/dri/renderD128
0 crw-rw---- 1 root render 226, 129 Sep  3 13:12 /dev/dri/renderD129

/dev/dri/by-path:
total 0
0 lrwxrwxrwx 1 root root  8 Sep  3 13:12 pci-0000:04:00.0-card -> ../card1
0 lrwxrwxrwx 1 root root  8 Sep  3 13:12 pci-0000:67:00.0-card -> ../card0
0 lrwxrwxrwx 1 root root 13 Sep  3 13:12 pci-0000:67:00.0-render -> ../renderD128
0 lrwxrwxrwx 1 root root  8 Sep  3 13:12 pci-0000:b5:00.0-card -> ../card2
0 lrwxrwxrwx 1 root root 13 Sep  3 13:12 pci-0000:b5:00.0-render -> ../renderD129

GPU Detailed Information (lspci output)

$ lspci -tv | grep Battlemage
 +-[0000:64]-+-00.0-[65-68]----00.0-[66-68]--+-01.0-[67]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]
 \-[0000:b2]-+-00.0-[b3-b6]----00.0-[b4-b6]--+-01.0-[b5]----00.0  Intel Corporation Battlemage G31 [Intel Graphics]

Driver Version

26.31.3995.13

Installed GPU Driver Packages

mesa built from source.

gmmlib built from source.

$ sudo dpkg --list | grep -iE "igc|gmm|opencl|level-zero|fc|level_zero|ocloc|libze"
[sudo: authenticate] Password:
ii bpfcc-tools 0.35.0+ds-1ubuntu2 all tools for BPF Compiler Collection (BCC)
ii clinfo 3.0.25.02.14-1build1 amd64 Query OpenCL system information
ii intel-igc-core-2 2.40.13 amd64 Intel(R) Graphics Compiler for OpenCL(TM)
ii intel-igc-opencl-2 2.40.13 amd64 Intel(R) Graphics Compiler for OpenCL(TM)
ii intel-ocloc 26.31.39395.13-0 amd64 Tool for managing Intel Compute GPU device binary format
ii intel-ocloc-dbgsym 26.31.39395.13-0 amd64 debug symbols for intel-ocloc
ii intel-opencl-icd 26.31.39395.13-0 amd64 Intel graphics compute runtime for OpenCL
ii intel-opencl-icd-dbgsym 26.31.39395.13-0 amd64 debug symbols for intel-opencl-icd
ii intltool-debian 0.35.0+20060710.6build1 all Help i18n of RFC822 compliant config files
ii libbpfcc:amd64 0.35.0+ds-1ubuntu2 amd64 shared library for BPF Compiler Collection (BCC)
ii libcbor0.10:amd64 0.10.2-2ubuntu3 amd64 library for parsing and generating CBOR (RFC 7049)
ii libdebconfclient0:amd64 0.280ubuntu1 amd64 Debian Configuration Management System (C-implementation library)
ii libigc2 2.38.5-126.04 amd64 Core libraries for Intel(R) Graphics Compiler for OpenCL(TM)
ii libigdfcl2 2.38.5-126.04 amd64 OpenCL library for Intel(R) Graphics Compiler for OpenCL(TM)
ii libigdgmm12:amd64 22.10.0-126.04ppa1 amd64 Intel Graphics Memory Management Library -- shared library
ii libref-array1t64:amd64 0.6.2-3build1 amd64 refcounted array for C
ii libze-dev:amd64 1.32.0-126.04ppa1 amd64 oneAPI Level Zero -- development files
ii libze-intel-gpu1 26.31.39395.13-0 amd64 Intel(R) Graphics Compute Runtime for oneAPI Level Zero.
ii libze-intel-gpu1-dbgsym 26.31.39395.13-0 amd64 debug symbols for libze-intel-gpu1
ii libze1:amd64 1.32.0-126.04ppa1 amd64 oneAPI Level Zero -- share libraries
ii linux-firmware-qlogic 20260319.git217ca6e4-0ubuntu1.1 all Firmware for QLogic SCSI, FC, and IB host bus and Ethernet adapters
ii ocl-icd-libopencl1:amd64 2.3.4-1 amd64 Generic OpenCL ICD Loader
ii python3-bpfcc 0.35.0+ds-1ubuntu2 all Python 3 wrappers for BPF Compiler Collection (BCC)

Driver Installation Details

Driver built from source. Current HEAD of master.

Mesa built from source.

Linux Distribution

Other (please specify below)

Other Linux Distribution

Ubuntu 26.04

Kernel Version & Boot Parameters

kernel 7.0.0-30-generic

Actual Behavior

When doing a p2p copy, the driver can crash, depending on the mapping of the blocks on the peer.

To reproduce, see attached test code.

Expected Behavior

Expected: copy without crashing.

Actual: GPU crash.

Reproduction Rate

Always reproduces - 100%

Steps to Reproduce

Build repro:

$ gcc -O2 -o ze_vmm_peer_hang ze_vmm_peer_hang.c -lze_loader

$ ./ze_vmm_peer_hang 0x300000 32 shuffle
dev0: Intel(R) Arc(TM) Pro B70 Graphics
dev1: Intel(R) Arc(TM) Pro B70 Graphics
driver version: 0x10399e3
dev0 can access dev1: 1
page size: 2097152, blocks: 4, copy size: 3145728 (CROSSES block boundary)
VA reservation: 8 MiB
block mapping order: in order
dev1 VA range: 0x7db9bfa00000 .. 0x7db9c0200000

dev1 local fill across all 4 blocks ... 
  queue group 0 (compute), 1 queues
  local fill OK

dev0 peer-reads 3145728 bytes from dev1 VA
  queue group 0 (compute), 1 queues
  peer copy DID NOT COMPLETE after 10 s -> 0x70000001
>>> REPRODUCED

Is this a regression?

• Yes, this is a regression - functionality that previously worked is now broken

Last Known Working Driver Version

No response

First Known Failing Driver Version

No response

API Call Logs

strace Logs

No response

System Logs / dmesg Output

[641944.916911] xe 0000:67:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT
[641944.916919] xe 0000:67:00.0: [drm] Tile0: GT0: 
                	ASID: 268
                	Faulted Address: 0x0000eaab55010000
                	FaultType: 0
                	AccessType: 0
                	FaultLevel: 1
                	EngineClass: 5 ccs
                	EngineInstance: 0
[641944.916922] xe 0000:67:00.0: [drm] Tile0: GT0: Fault response: Unsuccessful -ENOENT
[641945.010117] xe 0000:67:00.0: [drm] Xe device coredump has been created
[641945.010121] xe 0000:67:00.0: [drm] Check your /sys/class/drm/card0/device/devcoredump/data
[641946.900893] xe 0000:67:00.0: [drm] exec queue reset detected
[641947.401932] xe 0000:67:00.0: [drm] exec queue reset detected
[641948.403789] xe 0000:67:00.0: [drm] exec queue reset detected

Backtrace (if crash or hang occurred)

gpu-crash.log

Source Code / Reproducer

No response

Command Line / Application Details

// ze_vmm_peer_hang.c - Level Zero reproducer, no SYCL, no Unified Runtime.
//
// A peer copy hangs the engine when the range on the REMOTE device crosses from
// one zeVirtualMemMap()'d physical allocation into the next adjacent one inside
// a single zeVirtualMemReserve() range.
//
//   device 1: one reserved VA range, N separate zePhysicalMemCreate() blocks
//             mapped end to end at page-size granularity
//   device 0: plain zeMemAllocDevice() destination
//   device 0 appends zeCommandListAppendMemoryCopy(dst, dev1_va, bytes)
//
//   bytes <= page size -> read stays inside the first block  -> completes
//   bytes >  page size -> read crosses into the second block -> never completes
//
// Local access across the same boundary is fine: device 1 fills the whole range
// itself first, and that completes.
//
// Build: gcc -O2 -o ze_vmm_peer_hang ze_vmm_peer_hang.c -lze_loader
// Run:   ./ze_vmm_peer_hang [bytes] [compute|copy]

#include <level_zero/ze_api.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define CHECK(x) do {                                                        \
        ze_result_t _r = (x);                                                \
        if (_r != ZE_RESULT_SUCCESS) {                                       \
            fprintf(stderr, "%s:%d: %s -> 0x%x\n", __FILE__, __LINE__, #x, _r); \
            exit(1);                                                         \
        }                                                                    \
    } while (0)

#define SYNC_TIMEOUT_NS 10000000000ull   /* 10 s */

static uint32_t queue_group(ze_device_handle_t dev, int want_copy) {
    uint32_t n = 0;
    CHECK(zeDeviceGetCommandQueueGroupProperties(dev, &n, NULL));
    ze_command_queue_group_properties_t * p = calloc(n, sizeof(*p));
    for (uint32_t i = 0; i < n; ++i) p[i].stype = ZE_STRUCTURE_TYPE_COMMAND_QUEUE_GROUP_PROPERTIES;
    CHECK(zeDeviceGetCommandQueueGroupProperties(dev, &n, p));

    uint32_t pick = 0, found = 0;
    for (uint32_t i = 0; i < n; ++i) {
        int is_compute = (p[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) != 0;
        int is_copy    = (p[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COPY) != 0;
        if (want_copy ? (is_copy && !is_compute) : is_compute) { pick = i; found = 1; break; }
    }
    if (!found) { fprintf(stderr, "no %s queue group\n", want_copy ? "copy" : "compute"); exit(1); }
    printf("  queue group %u (%s), %u queues\n", pick, want_copy ? "copy" : "compute", p[pick].numQueues);
    free(p);
    return pick;
}

static ze_command_list_handle_t imm_list(ze_context_handle_t ctx, ze_device_handle_t dev, uint32_t ordinal) {
    ze_command_queue_desc_t d = { ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC, NULL, ordinal, 0, 0,
                                  ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, ZE_COMMAND_QUEUE_PRIORITY_NORMAL };
    ze_command_list_handle_t cl;
    CHECK(zeCommandListCreateImmediate(ctx, dev, &d, &cl));
    return cl;
}

int main(int argc, char ** argv) {
    const size_t want_bytes = (argc > 1) ? strtoull(argv[1], NULL, 0) : 0;
    const int    want_copy  = (argc > 2) && strcmp(argv[2], "copy") == 0;
    const size_t reserve_gib = (argc > 3) ? strtoull(argv[3], NULL, 0) : 0;
    const int    shuffle     = (argc > 4) && strcmp(argv[4], "shuffle") == 0;

    setvbuf(stdout, NULL, _IONBF, 0);
    CHECK(zeInit(0));

    uint32_t ndrv = 0;
    CHECK(zeDriverGet(&ndrv, NULL));
    ze_driver_handle_t * drv = calloc(ndrv, sizeof(*drv));
    CHECK(zeDriverGet(&ndrv, drv));

    ze_device_handle_t d0 = NULL, d1 = NULL;
    ze_driver_handle_t drv_used = NULL;
    for (uint32_t i = 0; i < ndrv && !d1; ++i) {
        uint32_t nd = 0;
        CHECK(zeDeviceGet(drv[i], &nd, NULL));
        ze_device_handle_t * devs = calloc(nd, sizeof(*devs));
        CHECK(zeDeviceGet(drv[i], &nd, devs));
        int got = 0;
        for (uint32_t j = 0; j < nd; ++j) {
            ze_device_properties_t pr = { ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES };
            CHECK(zeDeviceGetProperties(devs[j], &pr));
            if (pr.type != ZE_DEVICE_TYPE_GPU) continue;
            if (got == 0) { d0 = devs[j]; printf("dev0: %s\n", pr.name); got = 1; }
            else          { d1 = devs[j]; printf("dev1: %s\n", pr.name); drv_used = drv[i]; break; }
        }
        free(devs);
    }
    if (!d1) { fprintf(stderr, "need 2 GPUs on one driver\n"); return 1; }

    ze_driver_properties_t dp = { ZE_STRUCTURE_TYPE_DRIVER_PROPERTIES };
    CHECK(zeDriverGetProperties(drv_used, &dp));
    printf("driver version: 0x%x\n", dp.driverVersion);

    ze_bool_t can = 0;
    CHECK(zeDeviceCanAccessPeer(d0, d1, &can));
    printf("dev0 can access dev1: %d\n", (int) can);

    ze_context_desc_t cd = { ZE_STRUCTURE_TYPE_CONTEXT_DESC };
    ze_context_handle_t ctx;
    CHECK(zeContextCreate(drv_used, &cd, &ctx));

    // page size the virtual-memory API demands for this size class
    size_t pagesize = 0;
    CHECK(zeVirtualMemQueryPageSize(ctx, d1, 2ull << 20, &pagesize));
    const size_t bytes = want_bytes ? want_bytes : pagesize + 4096;
    const int    nblk  = 4;
    printf("page size: %zu, blocks: %d, copy size: %zu (%s block boundary)\n",
           pagesize, nblk, bytes, bytes > pagesize ? "CROSSES" : "inside one");

    // dev1: one VA reservation, nblk separate physical blocks mapped end to end
    void * va = NULL;
    const size_t reserve_sz = reserve_gib ? (reserve_gib << 30) : pagesize * nblk;
    printf("VA reservation: %zu MiB\n", reserve_sz >> 20);
    CHECK(zeVirtualMemReserve(ctx, NULL, reserve_sz, &va));
    // Create every physical block first, then map them into the VA range. With
    // "shuffle" the mapping order is reversed, so blocks that are ADJACENT in VA
    // are guaranteed NOT to be adjacent in physical VRAM - which is the normal
    // state of a long-lived pool, and what an idle GPU hides.
    ze_physical_mem_handle_t * ph = calloc(nblk, sizeof(*ph));
    for (int i = 0; i < nblk; ++i) {
        ze_physical_mem_desc_t pd = { ZE_STRUCTURE_TYPE_PHYSICAL_MEM_DESC, NULL, 0, pagesize };
        CHECK(zePhysicalMemCreate(ctx, d1, &pd, &ph[i]));
    }
    printf("block mapping order: %s\n", shuffle ? "REVERSED" : "in order");
    for (int i = 0; i < nblk; ++i) {
        int slot = shuffle ? (nblk - 1 - i) : i;
        CHECK(zeVirtualMemMap(ctx, (char *) va + (size_t) slot * pagesize, pagesize, ph[i], 0,
                              ZE_MEMORY_ACCESS_ATTRIBUTE_READWRITE));
    }
    printf("dev1 VA range: %p .. %p\n", va, (char *) va + pagesize * nblk);

    ze_event_pool_desc_t epd = { ZE_STRUCTURE_TYPE_EVENT_POOL_DESC, NULL,
                                 ZE_EVENT_POOL_FLAG_HOST_VISIBLE, 2 };
    ze_device_handle_t both[2] = { d0, d1 };
    ze_event_pool_handle_t pool;
    CHECK(zeEventPoolCreate(ctx, &epd, 2, both, &pool));
    ze_event_desc_t ed = { ZE_STRUCTURE_TYPE_EVENT_DESC, NULL, 0, 0, ZE_EVENT_SCOPE_FLAG_HOST };
    ze_event_handle_t ev0, ev1;
    CHECK(zeEventCreate(pool, &ed, &ev1)); ed.index = 1;
    CHECK(zeEventCreate(pool, &ed, &ev0));

    // control: dev1 writes its OWN whole range, across every block boundary
    printf("\ndev1 local fill across all %d blocks ... ", nblk); fflush(stdout);
    uint32_t ord1 = 0;
    { uint32_t n = 0; CHECK(zeDeviceGetCommandQueueGroupProperties(d1, &n, NULL)); }
    printf("\n"); ord1 = queue_group(d1, 0);
    ze_command_list_handle_t cl1 = imm_list(ctx, d1, ord1);
    uint8_t pat = 0x22;
    CHECK(zeCommandListAppendMemoryFill(cl1, va, &pat, 1, pagesize * nblk, ev1, 0, NULL));
    if (zeEventHostSynchronize(ev1, SYNC_TIMEOUT_NS) != ZE_RESULT_SUCCESS) {
        printf("  local fill DID NOT COMPLETE\n"); return 3;
    }
    printf("  local fill OK\n");

    // dev0: plain device allocation as the destination
    ze_device_mem_alloc_desc_t md = { ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC };
    void * dst = NULL;
    CHECK(zeMemAllocDevice(ctx, &md, pagesize * nblk, 4096, d0, &dst));

    printf("\ndev0 peer-reads %zu bytes from dev1 VA\n", bytes);
    uint32_t ord0 = queue_group(d0, want_copy);
    ze_command_list_handle_t cl0 = imm_list(ctx, d0, ord0);
    CHECK(zeCommandListAppendMemoryCopy(cl0, dst, va, bytes, ev0, 0, NULL));

    ze_result_t r = zeEventHostSynchronize(ev0, SYNC_TIMEOUT_NS);
    if (r == ZE_RESULT_SUCCESS) {
        printf("  peer copy COMPLETED\n");
        return 0;
    }
    printf("  peer copy DID NOT COMPLETE after %llu s -> 0x%x%s\n",
           SYNC_TIMEOUT_NS / 1000000000ull, r,
           r == ZE_RESULT_NOT_READY ? " (ZE_RESULT_NOT_READY: engine hung)" : "");
    printf(">>> REPRODUCED\n");
    return 2;
}

oneAPI Version (if applicable)

$ icpx --version
Intel(R) oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/2026.1/bin/compiler
Configuration file: /opt/intel/oneapi/compiler/2026.1/bin/compiler/../icpx.cfg

Screenshots / Video

No response

Additional Notes

I have a fix. PR will follow.

[stolk]

NOTE: the ze repro is not fully deterministic, but the SYCL version of the repro will crash 100%

// vmm_min.cpp <bytes>
//
// Minimal form. dev1 has a VMM pool built from 2 MiB physical_mem commits that
// are adjacent in one reserved VA range. dev0 peer-reads <bytes> from offset 0
// into a plain local malloc_device buffer.
//
//   bytes <= 2 MiB  -> read stays inside the first commit
//   bytes >  2 MiB  -> read crosses into the second commit
//
// icpx -fsycl -O2 -o vmm_min vmm_min.cpp

#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/virtual_mem/physical_mem.hpp>
#include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <string>

using namespace sycl;
namespace syclex = sycl::ext::oneapi::experimental;

static const size_t GRAN = 2ull << 20;

int main(int argc, char ** argv) {
    const size_t bytes = (argc > 1) ? strtoul(argv[1], nullptr, 0) : (GRAN + 4096);
    const size_t srcoff = (argc > 2) ? strtoul(argv[2], nullptr, 0) : 0;
    const bool   local  = (argc > 3) && std::string(argv[3]) == "local";

    std::vector<device> gpus;
    for (auto & p : platform::get_platforms())
        for (auto & d : p.get_devices())
            if (d.is_gpu() && d.get_backend() == backend::ext_oneapi_level_zero) gpus.push_back(d);
    if (gpus.size() < 2) return 1;

    device d0 = gpus[0], d1 = gpus[1];
    queue q0(d0, {property::queue::in_order()});
    queue q1(d1, {property::queue::in_order()});
    context ctx = q0.get_context();
    if (d0.ext_oneapi_can_access_peer(d1)) d0.ext_oneapi_enable_peer_access(d1);

    // dev1: reserved VA + 4 separate 2 MiB commits, laid end to end
    uintptr_t base = syclex::reserve_virtual_mem(1ull << 35, ctx);
    std::vector<syclex::physical_mem> phys;
    for (int i = 0; i < 4; ++i) {
        syclex::physical_mem p(d1, ctx, GRAN);
        p.map(base + i * GRAN, GRAN, syclex::address_access_mode::read_write);
        phys.push_back(std::move(p));
    }
    void * src = (void *) (base + srcoff);
    q1.memset((void *) base, 0x22, 4 * GRAN).wait();

    void * dst = malloc_device(8ull << 20, local ? d1 : d0, ctx);

    printf("%s read %zu bytes at off 0x%zx (blocks %zu..%zu) ... ",
           local ? "LOCAL" : "PEER ", bytes, srcoff, srcoff / GRAN, (srcoff + bytes - 1) / GRAN);
    fflush(stdout);
    try { (local ? q1 : q0).memcpy(dst, src, bytes).wait(); printf("SUCCESS\n"); }
    catch (const std::exception & e) { printf("EXCEPTION: %s\n", e.what()); return 2; }
    return 0;
}