uPCIe CUDA#

Warning

This backend requires an out-of-tree kernel module that is not yet upstream. See Kernel Module for details.

The upcie-cuda backend enables direct PCIe peer-to-peer (P2P) data transfers between the NVMe device and the GPU, with I/O buffers allocated from a heap backed by CUDA device memory.

Physical addresses for GPU memory are resolved once at initialization through the Linux dma-buf interface. The CUDA driver exports memory at 64 KiB granularity (one dma-buf page per 64 KiB of GPU memory), and these physical addresses are stored in a lookup table (LUT) indexed by GPU page. At submission time, PRP entries are built at 4 KiB granularity by computing sub-page offsets within each 64 KiB LUT entry. This matches the host page size used by NVMe for PRP construction and sector alignment.

Memory Architecture#

This backend uses a hybrid memory model:

Structure

Location

Reason

Data buffers (xnvme_buf_alloc)

GPU device memory (device heap, 1 GiB by default)

Transferred directly by the NVMe controller via PCIe P2P, bypassing host DRAM

SQ, CQ, PRP lists

Host hugepage memory (host heap, 1 GiB by default)

The CPU writes and the NVMe controller DMA-reads these structures; host-accessible memory is required

CQ, with XNVME_QUEUE_P2P_CQ_MIRROR

GPU device memory, mirrored into the host CQ by a resident kernel

Keeps the controller’s completion writes behind its data writes on one path; see Completions in device memory

SQ of a GPU-issued queue

GPU device memory; host memory with XNVME_QUEUE_SQ_HOSTMEM

The controller fetches entries; from DRAM that is faster than across the root complex, see GPU-Resident Queue API

The NVMe data path goes GPU ↔ NVMe without touching host DRAM. The control path (submission queue entries, completion queue entries, PRP lists) still flows through host memory. As a result, both the device heap and the host heap are initialized when the first upcie-cuda device is opened.

The heap is allocated in the CUDA context current on the thread opening that first device, or, when none is current, in the device’s primary context, which the backend retains and makes current. An application on the runtime API thus finds the heap in the context it already uses on every thread, and what it creates after the open is untouched by the last close.

Completions in device memory#

PCIe keeps a requester’s posted writes in order. With one destination that costs nothing; with two, the root complex holds each write until the one ahead of it has been accepted at the other. A controller that writes its payloads into the GPU and its completions into host memory changes destination at every command, and on a fast enough device that shows up as a ceiling on small I/O: on a drive that reaches 4.0M IOPS at 512 B with the buffers in host memory, the same run into device memory stops at 2.8M, from a queue depth of 64 upwards. Larger transfers are unaffected, and so are drives that never approach the rate. The other split, payloads in host memory and completions in the GPU, has not been measured; by this account it pays the same.

Opening a queue with XNVME_QUEUE_P2P_CQ_MIRROR places its completion queue in device memory as well, beside the data, so the controller has a single destination again. The CPU still owns the queue: it builds submissions, rings doorbells and polls for completions exactly as before, against the host-memory CQ the queue would have had anyway. What changes is who fills that CQ. A kernel resident on the GPU for as long as the process has such a queue watches each device-memory CQ with one warp, and copies completions into the host CQ as their phase bit turns. No command is issued from the GPU and the GPU never touches the controller’s registers, so this needs none of the setup that GPU-Resident Queue API does and works behind a translating IOMMU.

The mirror costs about a microsecond per completion, which is a loss below a queue depth of 32 and a gain above 64. On the drive above the ceiling is gone at every size measured, 512 B through 4 KiB, with host memory and device memory within a percent or two of each other. One kernel serves every queue in the process, up to 64 of them, and is launched with the first and retired with the last. It runs on a stream of its own, in the context that holds the heap: the caller’s, or the primary context as above. So while such a queue is open, cudaDeviceSynchronize() and cudaFree() in that context wait for the kernel and do not return: synchronize on streams instead, and free device memory once the last such queue is terminated. For the same reason a process does not hold such a queue and a GPU-issued one at once, since creating and destroying the latter synchronise the device: each is refused with EBUSY while the other is open. The --p2p-cq-mirror flag of xnvmeperf and of the test tools sets the flag. uPCIe HIP implements it as well. The upcie backend rejects the flag with ENOTSUP, having no GPU to put the CQ in; the other backends ignore it, as they do the polling flags they do not implement.

The host and the kernel share a table with a slot per queue: the addresses of the two completion queues, the depth, a generation and an acknowledgement. The host fills a slot and raises its generation to an odd value; the warp copies until the generation changes, then acknowledges the value it worked under, and the host frees nothing before it sees that. A warp finding an even generation acknowledges the odd value before it, since a queue detached before the warp took it up has nothing to drain. The kernel is the only CUDA device code in the library and is compiled by nvcc in a translation unit of its own; the rest of upcie-cuda is C against the driver API.

A controller served by homi takes the flag as well. The client cannot create the queue itself there, so it asks the server for one whose completion queue alone lies in the heap the client registered, naming it by offset; the submission queue and the PRP scratch stay in the server’s heap, as for any served queue, and the server creates the pair with the completion queue at the client’s address. What comes back names a completion queue in the server’s heap that the controller never writes, and the resident kernel keeps that one a copy of the queue it placed, so the submitting side is unchanged.

A caller can hand over device memory it allocated itself with xnvme_buf_register(), which registers the range through the same registry the heap uses, so such a buffer is usable exactly as one from xnvme_buf_alloc(). Registering a range twice is cheap, since what it covers is refcounted. On a controller a server holds, the range is exported as a dma-buf and registered with the server instead, and the description that comes back is adopted into the same table; that is a socket round trip per call, so a client registers its buffers before it submits, not per command. Either way the range is whole device pages, 64 KiB on CUDA, since that is the granule the export works in.

Choosing the GPU and Heap Sizes#

Three fields of struct xnvme_opts control where the heaps live and how large they are:

Option

CLI

Default

gpu_id

--gpu-id

0

device_heap_size

--device-heap-size

0, meaning 1 GiB

host_heap_size

--host-heap-size

0, meaning 1 GiB

gpu_id is the CUDA device ordinal the device heap is allocated on. When a CUDA context is current on the opening thread, it must belong to that GPU, or the open fails with EINVAL.

The device heap is rounded up to 64 KiB and the host heap to the hugepage size.

Each heap is allocated once per process, on the first open, so the values given to that open are the ones that take effect, and later opens leave the heaps as they are.

Not every tool takes these options on its command line. xnvmeperf takes --gpu-id and sizes both heaps from the workload it is given, and homi takes all three.

Kernel Module#

Physical address resolution for CUDA device memory relies on importing a dma-buf exported by the CUDA driver into the kernel’s DMA subsystem. The ioctls that do so are not upstream. They ship as dmabuf_import, an out-of-tree module serving /dev/dmabuf_import, in upcie-experimental-dkms. That DKMS package is attached to every uPCIe release and carries its version. Install it from the release that subprojects/upcie.wrap pins:

sudo apt install ./upcie-experimental-dkms_<version>_all.deb

These ioctls began as a patch to the in-tree udmabuf driver, which meant rebuilding the kernel to get them. That is no longer necessary, and the module supersedes the udmabuf-import package that carried the patched version.

Enforcing IOMMU#

Under uio_pci_generic the controller consumes physical addresses and the registry hands it those. Under vfio-pci an IOMMU translates, so every address it sees is an IOVA. Device memory cannot get one the way host memory does, since IOMMU_IOAS_MAP_FILE rejects the dma-bufs the CUDA driver exports.

Where the controller is on vfio-cdev, iommufd is asked first: it maps the heap once for the whole process and needs no out-of-tree module. It cannot always, and what follows is what happens when it cannot.

The iommu_map_pa module inserts the VRAM into the controller’s domain directly, and the registry’s table holds the resulting IOVAs. It ships in the same upcie-experimental-dkms package as dmabuf_import, installed as Kernel Module describes.

IOVA window#

The module maps at IOVAs of its own choosing without telling the domain’s owner, so a window is reserved for it: 64 GiB at 256 GiB, with the IOAS allowed every usable IOVA except that. Override where it does not fit:

XNVME_UPCIE_IOVA_BASE=0x2000000000  # 128 GiB
XNVME_UPCIE_IOVA_SIZE=0x400000000   # 16 GiB

A window outside every usable IOVA range fails xnvme_dev_open() with ERANGE.

Several controllers#

The window holds one set of addresses for the process. The first controller opened claims it and the heap is installed in it; every controller opened after that is attached to the same range, so all of them reach the heap, and anything registered with xnvme_buf_register(), at the same IOVAs through one translation table. Two limits follow, on the vfio-pci path only:

  • The heap and every registered buffer have to fit in the window. A full window fails the registration with ENOSPC.

  • The translation table is sized as uPCIe (host memory) describes. XNVME_UPCIE_VA_BITS bounds it.

A controller attaching later brings its own reserved regions. Where one overlaps the ranges the IOAS was told to allow, that attach fails with EADDRINUSE. Move the window with XNVME_UPCIE_IOVA_BASE.

uio_pci_generic is unaffected. Physical addresses read the same from every controller, so one table serves them all.

Registering buffers#

A buffer from xnvme_buf_alloc() comes off the heap, which every controller maps, so it works on any of them.

A buffer handed to xnvme_buf_register() is registered against the device passed to the call, so behind an IOMMU it resolves for that controller only. Register it with each device it is used from, and unregister it against each:

xnvme_buf_register(dev_a, buf, nbytes);
xnvme_buf_register(dev_b, buf, nbytes);

Under uio_pci_generic one table serves every controller, so a single registration covers them all.

What upstream would have to change#

None of the out-of-tree code here exists because the kernel cannot do these things. It exists because the interfaces that would are closed to a dma-buf that a GPU exported, or to memory that is not RAM. Three changes would retire it:

  • A way to resolve a dma-buf to physical addresses. This is what dmabuf_import does, and what a driver already does internally when it maps one for DMA. Nothing equivalent is exposed to user space.

  • IOMMU_IOAS_MAP_FILE accepting a GPU dma-buf. It refuses one today, which is why the mapping is installed from physical addresses instead.

  • A CUDA runtime that can map a BAR it is handed. cuMemHostRegister(..., CU_MEMHOSTREGISTER_IOMEMORY) puts a BAR in the GPU’s address space, which is what lets a kernel ring a doorbell. It takes a host mapping and resolves it, and for a mapping made through a vfio device it will not: only the first page of one is accepted, and the doorbells are never in it. Taking a dma-buf wrapping the BAR, which vfio can export, would make the mapping a thing the runtime is given rather than something it has to work out from a virtual address.

System Configuration#

Hardware Requirements#

P2P transfers require a GPU whose BAR1 window maps all of its device memory into the host PCIe address space, which in practice means resizable BAR enabled in the firmware. The window must be at least as large as the GPU’s total memory, whatever the size of the device heap. To compare the two:

nvidia-smi -q -d memory

If Total under BAR1 Memory Usage is smaller than Total under FB Memory Usage, opening a device fails with EOPNOTSUPP.

Hugepages#

In addition to the device heap, opening an upcie-cuda device also initializes the host heap (1 GiB by default) used to hold NVMe control structures. Follow the hugepage setup steps in uPCIe (host memory) before opening an upcie-cuda device.

GPU IOMMU domain#

Important

GPU-issued I/O behind a translating IOMMU needs the iommu_map_pa module. With it loaded the backend maps the controller’s BAR0 into the GPU’s domain itself and nothing else is needed; without it the queue is refused, and the fallback is the domain type, set as described below. Host-issued I/O, including Completions in device memory, needs none of this.

Needed only for GPU-resident queues, meaning xnvmeperf cuda-run, cuda-verify and anything else built on GPU-Resident Queue API. Host-driven I/O needs nothing here.

With a GPU-resident queue the GPU writes the doorbell register itself, rather than the CPU doing it, so the write is peer-to-peer traffic into the controller’s BAR0. It is translated by the GPU’s own domain, not the controller’s that Enforcing IOMMU sets up, and CUDA hands the GPU a physical address. With the GPU in a translating domain nothing has mapped it, so every write faults, which is what the identity mapping of the BAR into that domain prevents. The backend installs it, at the BAR’s own physical address, when it opens a controller for GPU-issued I/O and the GPU’s group is of type DMA or DMA-FQ, through the same iommu_map_pa module the heap relies on for the reverse direction; the throughput is then the one measured with the IOMMU off. The mapping is per controller and held for as long as the controller is open. The domain is the kernel’s, the one the GPU’s driver allocates its own DMA addresses from, and those come from the top of the 32-bit range downwards, so a BAR placed below 4 GiB can collide with them. Where the identity range is taken, GPU-issued I/O is declined for that controller, since a mapping the driver placed there looks the same as one an earlier opener of the controller installed. Setting XNVME_UPCIE_CUDA_DOORBELL_SHARE takes the collision to be the latter, for a second process opening the same controller for the same GPU. The mapping is then the first process’s and goes when it closes the controller, so the first has to outlive the second. What follows describes the situation without the module.

Recognising it#

The run completes while every I/O fails, taking the queue’s timeout rather than --runtime:

 Device                        IOPS      MiB/s   Failed
 0000:01:00.0                  2.30       0.01       32

The kernel log names the GPU, not the NVMe:

nvidia 0000:2b:00.0: AMD-Vi: IO_PAGE_FAULT domain=0x0012 address=0xee901010

0xee901010 is BAR0 plus 0x1010, the submission-queue doorbell of queue 2. An Intel host reports the equivalent under DMAR:.

Fixing it persistently#

Boot with the GPU in a passthrough domain. Add iommu.passthrough=1 to GRUB_CMDLINE_LINUX_DEFAULT in /etc/default/grub, then sudo update-grub && sudo reboot. This changes only the default domain type, so vfio-pci still attaches its own enforcing domain to the NVMe.

Fixing it without a reboot#

A group’s domain type can be switched only while no device in that group has a driver bound. Find the group and everything in it:

readlink -f /sys/bus/pci/devices/0000:2b:00.0/iommu_group
ls /sys/bus/pci/devices/0000:2b:00.0/iommu_group/devices/

The addresses and group number below are from one host. Substitute what those report, then release, unbind, switch, and rebind:

sudo fuser -v /dev/nvidia*   # an empty result is what you want

sudo systemctl stop nvidia-dcgm nvidia-persistenced
sudo rmmod nvidia_uvm   # plus nvidia_drm and nvidia_modeset where loaded

echo 0000:2b:00.1 | sudo tee /sys/bus/pci/drivers/snd_hda_intel/unbind
echo 0000:2b:00.0 | sudo tee /sys/bus/pci/drivers/nvidia/unbind

echo identity | sudo tee /sys/kernel/iommu_groups/16/type

echo 0000:2b:00.0 | sudo tee /sys/bus/pci/drivers/nvidia/bind
echo 0000:2b:00.1 | sudo tee /sys/bus/pci/drivers/snd_hda_intel/bind
sudo modprobe nvidia_uvm
sudo systemctl start nvidia-dcgm

The switch fails with EBUSY if anything is still bound, and does not survive a reboot.

Verifying the switch#

cat /sys/kernel/iommu_groups/16/type      # identity
sudo xnvmeperf cuda-verify --be upcie-cuda --iosize 4096 --qdepth 32 \
    --nqueues 2 0000:01:00.0
sudo dmesg | grep -iE 'AMD-Vi|DMAR'

cuda-verify should report 0 mismatches and dmesg nothing at all.

Note

This applies to any peer-to-peer write from a GPU into another device’s BAR, not just to xNVMe.

GPU-Resident Queue API#

The upcie-cuda backend supports GPU-resident NVMe queue pairs via the libxnvme_cuda API. See GPU-Resident Queue API for the full API reference, including host-side setup, CUDA kernel dispatch, and queue depth semantics.

GPU-resident queues behind a translating IOMMU need the iommu_map_pa module, through which the backend puts the controller’s registers where the GPU’s doorbell writes land; see GPU IOMMU domain for the mechanism and for the domain setting that stands in for the module when it is absent.

By default the whole queue pair lives in device memory: the kernel writes entries locally, the controller fetches them across PCIe, completes into device memory and moves payloads to and from it, so nothing of the I/O touches host DRAM. Passing XNVME_QUEUE_SQ_HOSTMEM to xnvme_cuda_queue_create() moves the submission queue alone into host memory, taken from the process’s own heap or, on a served controller, from the server’s, and mapped into the GPU’s address space; the kernel then writes entries across PCIe and the controller fetches them from DRAM. The reader decides: a controller fetching entries from device memory pays a PCIe round trip per fetch and keeps only so many in flight, which on a fast controller capped a queue pair below what the host can drive, whereas the GPU’s writes and the controller’s completion writes are posted and cost nothing of the sort. Measured with 512 B random reads, the host placement raised one queue pair from 1.65M to 3.8M IOPS and eight drives from 40M to 49M, at the same latency at a depth of one. What it costs is the entries themselves on the GPU’s transmit side, four 16-byte writes per command, which is idle traffic for reads but competes with the controller’s payload reads on writes; and the placement depends on host memory the GPU can map. The --sq-hostmem flag of xnvmeperf cuda-run and cuda-verify selects it.

A kernel writes an entry as four 16-byte stores cached at L2 only, so the L2 combines them into whole sectors before they leave the GPU. With narrower stores each would cross PCIe as a posted write of its own, and sixteen of them per entry bound the rate before the controller does.

Validation#

The cijoe task test-gpu.yaml exercises the upcie-cuda backend against a PCIe NVMe device. Devices must be tagged with the cuda label in your cijoe configuration:

[[devices]]
uri = "0000:xx:00.0"
nsid = 1
labels = ["dev", "pcie", "nvm", "cuda"]
driver_attachment = "userspace"

Run the task with:

cd cijoe && cijoe tasks/test-gpu.yaml --config configs/<your-config>.toml

Limitations#

  • 100 controllers per process, 16 where the heap goes in through iommu_map_pa. The first is the controller table (XNVME_BE_CREF_MAX_ENTRIES), the second the one importer a process attaches its controllers to (DMAMEM_IOMMU_MAP_PA_DEVS_MAX). Under an enforcing IOMMU the heap and every registered buffer must also fit in the IOVA window, see Enforcing IOMMU.

  • vfio-pci may need the mapping installed by hand. Describing the GPU heap by mapping its dma-buf is what IOMMU_IOAS_MAP_FILE would do, and it refuses one a GPU runtime exported. Where iommufd cannot map it, the mapping goes in from the physical addresses behind it instead. That happens either way: a server does it for a client, and a process holding the controller does it for itself. See What upstream would have to change.

  • Doorbells come from sysfs under vfio-pci. A queue the GPU submits on needs the doorbell page in the GPU’s address space, and the CUDA runtime will not take it from a vfio mapping, so it is taken from the BAR’s resource0 instead. Both name the same registers. This costs a served client something: everything else it needs arrives as a descriptor over the socket, and this it has to open for itself, which resource0 only permits to root. A client that submits from the host is unaffected.

  • I/O the GPU issues needs no server. A controller this process opened builds the same queue in the same device memory, and behind an IOMMU installs the mapping itself rather than being given one.

  • I/O the GPU issues behind a translating IOMMU needs the mapper module. See GPU IOMMU domain. What a client submits from the host works either way.

  • One GPU per process. The device heap lives on the GPU chosen by the first open. See Choosing the GPU and Heap Sizes.

  • No PRP list chaining. Each request has a single 4 KiB PRP list page (512 entries). Combined with PRP1, the maximum transfer size per command is 513 × 4 KiB ≈ 2 MiB.

  • No buf_realloc. Buffer reallocation is not implemented and returns ENOSYS.

  • No pseudo commands. Show registers, controller reset, subsystem reset, and namespace rescan all return ENOSYS.