USM pool: freeing a resident pooled chunk evicts the whole pool; indirect-access kernels then hang (DEVICE_LOST) on Arc B570, 26.27.39122.11

Start with the named pool residency paths in shared/source/memory_manager/unified_memory_pooling.cpp and unified_memory_pooling.h, then trace how shared/source/memory_manager/unified_memory_manager.cpp scans allocations for indirect access. Reproduce with the supplied fill.spv kernel and repro.cpp, comparing runs with and without the probe. Done when pooled live allocations remain accessible after the probe is freed and the reproducer returns 3.0.

A kernel that reaches a pooled device USM allocation only indirectly (the pointer is not a kernel argument; zeKernelSetIndirectAccess(ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE) is set) hangs the GPU with ZE_RESULT_ERROR_DEVICE_LOST once another context in the same process has made a pooled allocation resident with zeContextMakeMemoryResident and then freed it. UMF does exactly this when SYCL creates a context (query_min_page_size in src/provider/provider_level_zero.c: 1 byte zeMemAllocDevice, zeContextMakeMemoryResident on the provider's resident devices, zeMemFree), so a HIP program on chipStar that calls hipblasCreate (oneMKL through SYCL interop) and then runs a thrust kernel loses the device. chipStar tracking issue: https://github.com/CHIP-SPV/chipStar/issues/1962

GPU hardware: Intel Arc B570 (Battlemage, PCI 8086:e20c), discrete, the only device selected (NEOReadDebugKeys=1 FilterBdfPath=0000:03:00.0; the reproducer uses driver 0 device 0).

Driver version: compute-runtime 26.27.39122.11 (libze_intel_gpu.so.1.15.39122), IGC 2.38.2, gmmlib 22.10.0, Level Zero loader 1.28.2. Installed from the release .deb packages of https://github.com/intel/compute-runtime/releases/tag/26.27.39122.11.

Linux distribution: Ubuntu 24.04.2 LTS, kernel 6.17.0-14-generic, xe driver, boot parameters ro quiet splash.

Actual behavior: the kernel launch after the probe completes, then the next command on the same immediate command list returns 0x70000001 (ZE_RESULT_ERROR_DEVICE_LOST); the kernel log shows xe 0000:03:00.0: [drm] GT0: Engine reset: engine_class=ccs. The next process using the device sometimes hangs on its first submission until the engine is reset again.

Expected behavior: the kernel writes the allocation and the copy returns 3.0, as it does without the other context's probe. The Level Zero specification (Device Residency, scripts/core/PROG.rst) says "the driver must ensure that all pages that will be accessed by the kernel are resident before program execution" and that the indirect access flags are the way to declare such accesses, and zeContextEvictMemory says "the memory is implicitly evicted when freed", which concerns the freed memory, not other live allocations.

Reproduction rate: every run.

Steps to reproduce: compile the kernel with clang --target=spirv64 -cl-std=CL2.0 -O2 -c fill.cl -o fill.spv, build the host program with g++ repro.cpp -lze_loader -o repro, run ./repro 0 (passes) and ./repro 1 (device lost).

// fill.cl: the pointer is passed as an integer, so the driver sees no buffer argument.
__kernel void fill(ulong p, double v, int n) {
  for (int i = get_local_id(0); i < n; i += get_local_size(0)) ((__global double *)p)[i] = v;
}

// repro.cpp
#include <level_zero/ze_api.h>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <iterator>
#include <vector>

#define CHECK(x) do { ze_result_t r_ = (x); printf("%-60s 0x%x\n", #x, r_); fflush(stdout); if (r_) exit(1); } while (0)

static ze_device_handle_t dev;
static ze_command_list_handle_t cl;
static ze_kernel_handle_t k;

static void fill(void *p, double v) {
  uint64_t a = (uint64_t)p; int n = 256;
  CHECK(zeKernelSetArgumentValue(k, 0, sizeof(a), &a));
  CHECK(zeKernelSetArgumentValue(k, 1, sizeof(v), &v));
  CHECK(zeKernelSetArgumentValue(k, 2, sizeof(n), &n));
  ze_group_count_t g = {1, 1, 1};
  CHECK(zeCommandListAppendLaunchKernel(cl, k, &g, nullptr, 0, nullptr));
  CHECK(zeCommandListHostSynchronize(cl, UINT64_MAX));
}

int main(int argc, char **argv) {
  int probe = argc > 1 ? atoi(argv[1]) : 1;
  CHECK(zeInit(0));
  uint32_t n = 1; ze_driver_handle_t drv; CHECK(zeDriverGet(&n, &drv));
  n = 1; CHECK(zeDeviceGet(drv, &n, &dev));
  ze_context_desc_t cd = {ZE_STRUCTURE_TYPE_CONTEXT_DESC};
  ze_context_handle_t ctx; CHECK(zeContextCreate(drv, &cd, &ctx));
  ze_command_queue_desc_t qd = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
  CHECK(zeCommandListCreateImmediate(ctx, dev, &qd, &cl));
  std::ifstream f("fill.spv", std::ios::binary);
  std::vector<unsigned char> spv((std::istreambuf_iterator<char>(f)), {});
  ze_module_desc_t md = {ZE_STRUCTURE_TYPE_MODULE_DESC, nullptr, ZE_MODULE_FORMAT_IL_SPIRV, spv.size(), spv.data()};
  ze_module_handle_t mod; CHECK(zeModuleCreate(ctx, dev, &md, &mod, nullptr));
  ze_kernel_desc_t kd = {ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0, "fill"};
  CHECK(zeKernelCreate(mod, &kd, &k));
  CHECK(zeKernelSetIndirectAccess(k, ZE_KERNEL_INDIRECT_ACCESS_FLAG_DEVICE));
  CHECK(zeKernelSetGroupSize(k, 256, 1, 1));

  ze_device_mem_alloc_desc_t dd = {ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC};
  void *d1; CHECK(zeMemAllocDevice(ctx, &dd, 2048, 64, dev, &d1));
  fill(d1, 1.0);

  if (probe) {  // what UMF's query_min_page_size does in its own context
    ze_context_handle_t ctx2; CHECK(zeContextCreate(drv, &cd, &ctx2));
    void *p; CHECK(zeMemAllocDevice(ctx2, &dd, 1, 0, dev, &p));
    CHECK(zeContextMakeMemoryResident(ctx2, dev, p, 1));
    CHECK(zeMemFree(ctx2, p));
  }

  void *d2; CHECK(zeMemAllocDevice(ctx, &dd, 2048, 64, dev, &d2));
  printf("d1=%p d2=%p\n", d1, d2);
  fill(d2, 3.0);
  double h = 0;
  CHECK(zeCommandListAppendMemoryCopy(cl, &h, d2, sizeof(h), nullptr, 0, nullptr));
  CHECK(zeCommandListHostSynchronize(cl, UINT64_MAX));
  printf("d2[0]=%g %s\n", h, h == 3.0 ? "PASS" : "FAIL");
}

$ ./repro 0
...
d1=0xffffffffff9ff800 d2=0xffffffffff9ff000
d2[0]=3 PASS
$ ./repro 1
...
zeContextMakeMemoryResident(ctx2, dev, p, 1)                 0x0
d1=0xffffffffff9ff800 d2=0xffffffffff9ff000
zeCommandListHostSynchronize(cl, UINT64_MAX)                 0x0
zeCommandListAppendMemoryCopy(cl, &h, d2, sizeof(h), nullptr, 0, nullptr) 0x70000001
$ NEOReadDebugKeys=1 MakeIndirectAllocationsResidentAsPack=0 ./repro 1
d2[0]=3 PASS

The 1 byte probe in the second context lands in the same device USM pool as d1 and d2. Larger probes that are not pooled (64 KB) do not trigger it, and neither does the probe without the zeMemFree.

What the source shows at 26.27.39122.11: making a pooled chunk resident raises the pool's per device count and makes the whole backing allocation resident (residencyOperation, https://github.com/intel/compute-runtime/blob/26.27.39122.11/shared/source/memory_manager/unified_memory_pooling.h#L76-L93). Freeing that chunk drops the count and, at zero, evicts the whole backing allocation, although other live chunks (d1, d2) are still in it (UsmMemAllocPool::freeSVMAlloc, https://github.com/intel/compute-runtime/blob/26.27.39122.11/shared/source/memory_manager/unified_memory_pooling.cpp#L143-L166). SVMAllocsManager::makeIndirectAllocationsResident only visits allocation IDs newer than the last scan for that command stream receiver (https://github.com/intel/compute-runtime/blob/26.27.39122.11/shared/source/memory_manager/unified_memory_manager.cpp#L1173-L1200), and a new pool chunk gets no new ID for the backing allocation, so the evicted backing allocation is never made resident again for indirect launches. Disabling the pack path with MakeIndirectAllocationsResidentAsPack=0 makes the reproducer pass. Current master still evicts the pool when the last explicitly resident chunk is released (releaseChunk and dropPoolResidency, https://github.com/intel/compute-runtime/blob/d09d0dbfc3083cd367f780d92a3c12c55e8523ad/shared/source/memory_manager/unified_memory_pooling.cpp#L251-L263 and https://github.com/intel/compute-runtime/blob/d09d0dbfc3083cd367f780d92a3c12c55e8523ad/shared/source/memory_manager/unified_memory_pooling.cpp#L414-L423) and keeps the incremental scan (https://github.com/intel/compute-runtime/blob/d09d0dbfc3083cd367f780d92a3c12c55e8523ad/shared/source/memory_manager/unified_memory_manager.cpp#L1296-L1323); I have not run master.

Is this a regression: not known; the pool residency tracking it depends on is enabled by default for device USM pools in 26.27.

chipStar works around it by keeping its own small device allocations out of the pool until this is fixed.

This issue was generated using AI