Async. I/O on GPU memory via io_uring_dmabuf#

Warning

This needs a kernel that is not yet released. See Kernel.

The io_uring_dmabuf backend lets io_uring read and write GPU memory on a regular NVMe block device. A CUDA or HIP allocation is exported as a dma-buf and registered with io_uring as a fixed buffer. The NVMe controller then transfers directly to and from the GPU over PCIe peer-to-peer.

Unlike uPCIe CUDA, this uses the kernel NVMe driver, so the device is e.g. /dev/nvme0n1.

Usage#

Select the backend with opts.be or --be:

xnvmeperf run /dev/nvme0n1 --be io_uring_dmabuf --direct \
    --iopattern randread --iosize 4096 --qdepth 32 --runtime 10 --cpumask 0x1

Register GPU memory with xnvme_buf_register() and release it with xnvme_buf_unregister(). Buffers from xnvme_buf_alloc() are in host memory, IO on them is regular io_uring IO.

cudaMalloc(&buf, nbytes);
xnvme_buf_register(dev, buf, nbytes);

xnvme_buf_register() registers the allocation with the ring of every queue of the device, and a new queue registers the allocations already registered. IO on a registered buffer is submitted as IORING_OP_READ_FIXED / IORING_OP_WRITE_FIXED.

Limitations#

Target

A raw NVMe block device opened with O_DIRECT, opts.direct, which the open checks

Polling

Not supported, IORING_SETUP_IOPOLL fails queue initialization

Vectored

Not supported (xnvme_cmd_pass_iov())

Size

A dma-buf of at most 1 TiB, the kernel’s limit for a registered buffer

Count

At most 256 mapped allocations per device

Threads

xnvme_buf_register() and xnvme_buf_unregister() must not run while IO is in flight on the device

Kernel#

The io_uring and nvme-pci support is applied to the io_uring tree, but not yet in a released kernel:

The backend is built with liburing and CUDA or HIP. No liburing release has the UAPI of the series yet, so xNVMe defines the part it uses itself. On a kernel without the series, opening the device fails with -ENOTSUP. The xnvme-devenv guest runs a kernel with the series, and the CI runs the backend there, on an emulated AMD GPU.

IOMMU#

The dma-buf exported by CUDA and HIP holds physical addresses. When an IOMMU translates for the NVMe device, xnvme_buf_register() maps those addresses into the IOMMU domain of the device at IOVA == physical address.

This needs the dmabuf-import and iommu-map-pa kernel modules, see Kernel Module and Enforcing IOMMU. Without them xnvme_buf_register() fails. Alternatively, boot with iommu=pt, amd_iommu=off or intel_iommu=off. The IOMMU group is read from the controller of the namespace; where it cannot be found, as for the head of a native-multipath namespace, while IOMMU groups exist, xnvme_buf_register() fails with -ENODEV rather than map without the IOMMU.

Note

The dma-buf must be physically contiguous. See Enforcing IOMMU for caveats of iommu-map-pa.

On this page