[SYCL][Level Zero] sycl::malloc_device allocates a 1:1 host-side (GTT) mirror when the context contains more than one device

[TSUMUGI-XE]

Summary

On a multi-GPU Level Zero platform, every sycl::malloc_device allocation is accompanied by a
host-side allocation of the same size, accounted as GPU-driver memory (GTT). It appears as soon
as the SYCL context contains more than one device, even if only one device is ever used, and even
if nothing is copied between devices.

Allocating via Level Zero directly (zeMemAllocDevice) on the same machine, with the same two
devices in the ze_context, does not produce it. That places the behaviour in the SYCL
runtime's mapping onto L0 rather than in the driver.

The host allocation is invisible to VmRSS, ps and free's process accounting, which makes it
easy to miss.

Environment

• 2× Intel Arc Pro B70 (Battlemage / BMG G31)
• Linux 7.1.1 (xe), compute-runtime 26.22 / IGC 2.36.3, oneAPI 2026.1
• Reproduced through PyTorch XPU (SYCL backend); no application-specific code involved

Reproduction

Plain device allocations, ~30 s, no model or framework logic. Only ZE_AFFINITY_MASK differs.

1 GPU  (ZE_AFFINITY_MASK=0), 4 GiB allocated
  VmRSS         +0.6 MiB
  nr_gpu_active +0.00 GiB
  0000:0e:00.0   VRAM +4.00 GiB   GTT +0.00 GiB

2 GPUs (ZE_AFFINITY_MASK=0,1), 4 GiB allocated on each
  VmRSS         +5.0 MiB
  nr_gpu_active +8.07 GiB
  0000:0e:00.0   VRAM +4.00 GiB   GTT +4.00 GiB
  0000:12:00.0   VRAM +4.18 GiB   GTT +4.07 GiB

Each device ends up holding its own allocation in VRAM and a same-sized GTT entry mirroring the
other device's allocation. Total host cost equals the sum of all device allocations, 1:1. Freeing
the allocations releases it.

Measured with:

• /proc/<pid>/fdinfo/<drm fd> → drm-total-vram0, drm-total-gtt
• /sys/devices/system/node/node0/vmstat → nr_gpu_active

nr_foll_pin_* moved by 4 pages and Mlocked stayed at 3.2 MiB, so this is driver-owned system
memory rather than pinned or locked user pages.

Level Zero control (same box, SYCL bypassed)

zeMemAllocDevice called directly through ctypes, varying only the number of devices in the
ze_context:

zeMemAllocDevice(4 GiB on device 0) + zeContextMakeMemoryResident

  ze_context with 1 device :  0000:0e:00.0  VRAM 4.00 GiB  GTT 0.00 GiB   nr_gpu_active +0.00 GiB
  ze_context with 2 devices:  0000:0e:00.0  VRAM 4.00 GiB  GTT 0.00 GiB   nr_gpu_active +0.00 GiB

No mirror in either case, while the SYCL path on the same hardware produces 8.07 GiB.

Note: zeMemAllocDevice is lazily committed. Without zeContextMakeMemoryResident both cases
report VRAM 0.00 / GTT 0.00, which looks like a clean result but means nothing was committed.

Impact

Frameworks that create one SYCL context spanning all visible devices pay host RAM equal to their
total device allocation. In our case, a tensor-parallel inference server across 2 GPUs consumes host
RAM equal to the total VRAM pool:

--gpu-memory-utilization 0.90  ->  57.1 GiB host RAM
                          0.80  ->  50.7 GiB host RAM
   (57.1/50.7 = 1.126 vs 0.90/0.80 = 1.125 — proportional to the VRAM pool, not to model size)

The model itself is 45 GiB and resides entirely in VRAM; the host RAM is purely the mirror. On a
123 GiB host that is roughly half the machine, and it contributed to an out-of-memory condition that
was hard to diagnose precisely because the memory does not appear in the usual counters.

Restricting the context to one device removes it, but that is not an option for tensor-parallel
work. ONEAPI_DEVICE_SELECTOR=level_zero:N is equivalent to ZE_AFFINITY_MASK=N here — it only
removes the second device.

Related

• llama.cpp issue [libspirv] Use clc and vectorize binary relational built-ins #22116 reports the same behaviour from a different application, with a proposed
  fix of replacing sycl::malloc_device with zeMemAllocDevice (PR [UR][NFC] Add INTEL suffixes to Intel Vendor EXTs #21597, closed). The Level Zero
  control above is consistent with that fix addressing the cause.

Reproducer scripts (SYCL/PyTorch and the raw Level Zero control) can be provided; happy to run
further variants — different allocation sizes, more than two devices, or explicit per-device
contexts — if that helps.

[pbalcer]

This was identified as a compute-runtime level-zero issue. You need an updated version of the GPU driver. I recommend the latest version available here. Please see pytorch/pytorch#180145.

[TSUMUGI-XE]

Thanks — that's very helpful, and the PyTorch PR is exactly the missing piece.

One data point in case it is useful: the numbers above were taken on compute-runtime 26.18, and I
re-ran with 26.22.38646 loaded (verified via /proc/self/maps, not just installed) — the mirror
is still there, same 1:1 ratio:

4 GiB allocated per GPU, ZE_AFFINITY_MASK=0,1
  26.18.38308   VmRSS +5.0 MiB   nr_gpu_active +8.07 GiB
  26.22.38646   VmRSS +4.4 MiB   nr_gpu_active +8.07 GiB

Since pytorch/pytorch#180145 reports it clean on 26.09, that may be a regression, or something
specific to this setup (Arc Pro B70 / BMG G31, custom 7.1.1 xe kernel).

I have not been able to test 26.27 yet — dropping in only libze_intel_gpu.so aborts in NEO
(command_stream_receiver.cpp:1205), so it needs a full host driver upgrade, which I'll do next.

I'll verify on 26.27 and take it to intel/compute-runtime if it still reproduces. Feel free to close
this one — thanks again for the pointer.

[TSUMUGI-XE]

Following up with the driver-version data, since it turned out to be worth checking.

We still reproduce on the latest compute-runtime:

4 GiB allocated per GPU, ZE_AFFINITY_MASK=0,1, dual Arc Pro B70 (BMG G31)
  26.18.38308   VmRSS +5.0 MiB   nr_gpu_active +8.07 GiB
  26.22.38646   VmRSS +4.4 MiB   nr_gpu_active +8.07 GiB
  26.27.39122   VmRSS +4.4 MiB   nr_gpu_active +8.07 GiB    <- latest, measured on the host directly

Also relevant, from ggml-org/llama.cpp#22116: another dual-B70 user reproduces on 26.05 and
26.09, while a llama.cpp maintainer running dual B580 does not reproduce on those same driver
versions. Combined with the above, the behaviour does not track the compute-runtime version in
either direction.

And pytorch/pytorch#180145 was closed on the conclusion that a driver upgrade fixed it, but the
question "multi-GPU only, or single too?" was never answered there — the author had a single GPU,
and the non-repro was also single-GPU. On our dual-GPU box a single-device context never shows it and
a two-device context always does, on every driver we tried.

Finally, a raw zeMemAllocDevice + zeContextMakeMemoryResident reproducer with both devices in
the ze_context allocates VRAM with no GTT at all, while the SYCL path on the same machine and
the same driver produces the 1:1 mirror. #22116 reports the same split.

So this may belong here after all rather than in compute-runtime — it looks like the DPC++
multi-device USM allocation path rather than the driver. Happy to move it if you disagree; I just
wanted to make sure the "upgrade the driver" answer was tested before closing.

It does not appear to be used. My first guess was that this is a per-allocation host staging
area for cross-device copies, so I measured the copies. They run at full peer speed:

same-device   xpu:0 -> xpu:0    588.32 GB/s
cross-device  xpu:0 -> xpu:1     28.60 GB/s
cross-device  xpu:1 -> xpu:0     28.58 GB/s

On this machine ze_peer measures 28.4 GB/s cross-card with checksums passing both directions (both
B70 on downstream ports of one Broadcom PEX880xx Gen4 switch, ACS redirect disabled), and a
host-staged path measures ~3 GB/s. The SYCL cross-device copies match the peer number, not the
host-staged one — so the host allocation is not serving as a staging buffer for them.

As far as I can tell it is allocated 1:1 with device memory and then not used. On a 123 GiB host
running tensor-parallel inference that came to ~57 GiB.

Disclosure: our kernel is a custom 7.1.1 build carrying a PCI P2P whitelist patch, so not a stock
configuration and I can't rule it out as a factor. If anyone knows which kernels the B580
(non-repro) and the other B70 (repro) reports ran on, that would separate GPU model from kernel
version. I can arrange a build of the same kernel version without the patch if that turns out to be
the deciding control — I'd rather not confound it by booting an older stock kernel instead.

[pbalcer]

Finally, a raw zeMemAllocDevice + zeContextMakeMemoryResident reproducer with both devices in
the ze_context allocates VRAM with no GTT at all, while the SYCL path on the same machine and
the same driver produces the 1:1 mirror. #22116 reports the same split.

How does your reproducer work? malloc_device works roughly like this:

sycl::malloc_device(context, device, size) {
  ptr = zeMemAllocDevice(device, size)
  foreach(device in context.devices()) {
    zeContextMakeMemoryResident(ptr, device);
  }
  return ptr;
}

[TSUMUGI-XE]

You are right, and my control was wrong. Thank you for the pointer — it changes the conclusion.

My reproducer called zeContextMakeMemoryResident on devices[0] only, even when the context
contained two devices. So the "2-device context" arm never asked device 1 to make the allocation
resident, which is exactly the step sycl::malloc_device performs. It was not measuring the same
thing, and its "no mirror" result proves nothing.

Re-running with residency made explicit (1 GiB, same machine, compute-runtime 26.27.39122):

ctx devices | zeContextMakeMemoryResident on | device1 GTT | nr_gpu_active
------------+--------------------------------+-------------+---------------
     1      | dev0                           |   0.00 GiB  |     +0.00 GiB
     2      | dev0 only   (my original test)  |   0.00 GiB  |     +0.00 GiB
     2      | dev0 + dev1 (what SYCL does)    |   1.00 GiB  |     +1.00 GiB

Per-device /proc/<pid>/fdinfo at each step, allocation made on device 0:

after zeMemAllocDevice           0000:0e:00.0  VRAM 0.00  GTT 0.00
                                 0000:12:00.0  VRAM 0.00  GTT 0.00
after MakeMemoryResident(dev0)   0000:0e:00.0  VRAM 1.00  GTT 0.00
                                 0000:12:00.0  VRAM 0.00  GTT 0.00
after MakeMemoryResident(dev1)   0000:0e:00.0  VRAM 1.00  GTT 0.00
                                 0000:12:00.0  VRAM 0.00  GTT 1.00   <-- appears here

So the 1:1 host mapping reproduces in raw Level Zero with no SYCL involved. It is created by the
driver when an allocation owned by one device is made resident on another, and sycl::malloc_device
reaches it simply by doing that for every device in the context.

Retracting, then: my earlier statement that this is a SYCL-runtime mapping issue, and that
llama.cpp PR #21597's zeMemAllocDevice substitution addressed the cause. If that replacement still
performs the same per-device residency calls, I would not expect it to help. Sorry for the noise.

Two questions remain from my side, and I think only the first is yours-or-not-yours to route:

1. Why is peer residency backed by a same-sized host (GTT) allocation? On this machine the two
   GPUs sit under a PLX switch with ACS redirect off and P2P works — I measure 28.6 GB/s device-to-device
   with SYCL copies, and ze_peer agrees. Nothing is ever copied between the devices in the test above;
   the memory is only made resident. Paying host RAM equal to the allocation for that looks like a
   driver-side implementation choice rather than a necessity, so I will take this to
   intel/compute-runtime unless you would rather it stayed here.

2. Should SYCL make every allocation resident on every device in the context, eagerly? It does,
   and it is easy to see in isolation. Allocating only on device 0, never touching device 1:

ZE_AFFINITY_MASK   allocated on   device0 GTT   device1 GTT   nr_gpu_active
-----------------+--------------+-------------+-------------+---------------
       0          |   xpu:0      |   unchanged |      -      |    +0.00 GiB
      0,1         |   xpu:0      |   unchanged |  +1.00 GiB  |    +1.00 GiB

The second GPU pays host RAM equal to an allocation it never sees. That is the
amplification factor in practice: host cost ends up equal to the total device allocation across the
context, whether or not any allocation is ever touched by a second device. In a tensor-parallel
server this is utilization × card capacity × ranks of host RAM
(--gpu-memory-utilization 0.90 -> 57.1 GiB, 0.80 -> 50.7 GiB, model entirely in VRAM).
If lazy or on-first-cross-device-use residency is feasible, that would remove the cost for the
common case where allocations stay device-local.

Happy to run further variants here — different sizes, more devices, or an on-demand-residency
experiment if you can point me at the intended path.

[pbalcer]

Please file a compute-runtime issue. It's not expected that an equivalent host physical allocation is created for every GPU device allocation.

Our runtime makes memory resident on all devices to enable efficient p2p operations between GPUs in the system. This is going to change in the next release, and we will restrict this behavior to GPUs with p2p explicitly enabled. See #21889. For now this is gated by SYCL_UR_L0_RESTRICT_USM_RESIDENCY_TO_P2P env variable.

[TSUMUGI-XE]

Filed as intel/compute-runtime#968 — thanks again.

One result that came out of your explanation and may be worth knowing here: since the residency is there to enable P2P, I tried simply skipping zeContextMakeMemoryResident for non-owning devices (an LD_PRELOAD shim) on a 2-GPU tensor-parallel server. It starts and runs normally — host RAM peak drops from 55.12 GiB to 1.55 GiB, KV cache, startup time, xccl init and generated output are all unchanged, and device-to-device copies still run at the same 28.6 GB/s. One workload and one smoke test, so not a general claim, but on this box the host-side buffer did not appear to be load-bearing for the P2P it backs.

Also, for the record: SYCL_UR_L0_RESTRICT_USM_RESIDENCY_TO_P2P is not reachable from a released oneAPI DPC++ yet — #21889 is in the sycl branch and the nightlies, but strings over libur_adapter_level_zero*.so in our runtime finds no occurrence. Having read the PR I do expect it to help us once it ships, since it keys off explicit urUsmP2PEnablePeerAccessExp calls rather than hardware P2P capability.