fi_xpu(3)

📦 version 2.7.0, in arch rolling

✨ NAME ✨

fi_xpu — the OFI XPU API, and this is the fi_xpu(3) page from Libfabric v2.7.0 💖

🌸 SYNOPSIS 🌸

#include <rdma/fi_xpu.h>

🦋 OVERVIEW 🦋

okay so the FI_XPU capability lets a specified XPU (or device) access a given libfabric-allocated resource for data path operations, while control operations stay managed by host (CPU) software. A resource — endpoint, completion queue, counter — that has been configured with this capability may only be accessed by the given XPU for any data path operation, like submitting data transfers or reading completions.

As a general rule bestie: control path operations may only be invoked by the host CPU, while data path functions may only be accessed by a specified XPU. The following objects may be exported to an XPU: EP, CQ, and counters. AV and MR lookups produce raw data usable by the XPU, but the objects themselves remain host-only. Detailed XPU behavior is documented in the corresponding man pages (fi_av(3), fi_endpoint(3), fi_cq(3), fi_cntr(3), fi_mr(3)).

🎀 XPU CONTEXT 🎀

Related EP, CQ, and counter resources created with FI_XPU must all target the same XPU. The fid_xpu_ctx object groups these resources together and binds them to a specific device. An EP created with FI_XPU may also be bound to a host (non-XPU) CQ for error handling. It also provides the avenue for querying provider-specific sizes needed to interact with AV and MR data from the XPU.

An XPU context is created from a domain:

int fi_xpu_ctx(struct fid_domain *domain, struct fi_xpu_attr *attr,
               struct fid_xpu_ctx **ctx, void *context);

The number of XPU contexts that a domain supports is indicated by the max_xpu_ctx_cnt field in fi_domain_attr. A value of 0 indicates no support for XPU-initiated communication. A value of 1 indicates a 1:1 mapping between a domain and an XPU.

The XPU context is passed to EP, CQ, and counter creation calls via a pointer in their respective attribute structures (fi_ep_attr->xpu_ctx, fi_cq_attr->xpu_ctx, fi_cntr_attr->xpu_ctx). XPU resources bound to the same EP must share the same XPU context. A host CQ (without xpu_ctx set) may also be bound to an XPU EP for error handling.

The XPU context is also passed to fi_av_lookup2 and fi_mr_get_xpu_desc so that the provider can return data in the appropriate format for the target XPU. This allows AV and MR objects to be shared across multiple XPUs — the same AV or MR can be queried with different XPU contexts to get device-specific representations.

🧚 fi_xpu_attr 🧚

The fi_xpu_attr structure is an input argument that identifies the target XPU device and provides optional memory management callbacks. It is passed to fi_xpu_ctx() during context creation.

struct fi_xpu_ops {
    size_t size;
    int (*alloc)(uint64_t device, uint64_t size,
                 uint64_t alignment, uint64_t flags,
                 void **addr, int *fd, uint64_t *offset);
    int (*import)(uint64_t device, void *host_addr,
                  uint64_t size, uint64_t flags,
                  void **dev_addr);
    void (*free)(uint64_t device, void *addr);
};

struct fi_xpu_attr {
    int                  iface;    /* enum fi_hmem_iface */
    uint64_t             device;
    struct fi_xpu_ops    *ops;
};
  • iface — the heterogeneous memory interface type (e.g., FI_HMEM_CUDA, FI_HMEM_ZE).
  • device — the XPU device ordinal.
  • ops — optional pointer to a fi_xpu_ops structure containing memory management callbacks. If NULL, the provider uses its default mechanisms.

💜 fi_xpu_ops 💜

The fi_xpu_ops structure groups memory management callbacks for XPU device memory. The provider calls these when it needs to allocate, import, or free device memory on behalf of the XPU context.

  • size — must be set to sizeof(struct fi_xpu_ops). This allows future expansion of the structure while maintaining backward compatibility. The provider uses this field to determine which callback fields are present.
  • alloc — allocate XPU memory. The provider calls this when it needs device memory (e.g., for hardware queue buffers). If the FI_XPU_ALLOC_DMABUF flag is set, the consumer must also export a DMA-BUF file descriptor so the provider can DMA directly to the memory.
  • import — map a host virtual address into the XPU address space. The provider calls this when it has a host-side address (BAR MMIO or host RAM) that the XPU kernel needs to access. Flags indicate the memory type: FI_XPU_IMPORT_IOMEMORY for PCIe BAR MMIO regions, FI_XPU_IMPORT_DEVICEMAP for addresses that must be accessible from XPU kernels.
  • free — release memory previously allocated via alloc.

💅 Memory Callback Flags 💅

  • FI_XPU_ALLOC_DMABUF — passed to the alloc callback. Indicates the allocation must be exportable as a DMA-BUF fd for provider access.
  • FI_XPU_IMPORT_IOMEMORY — passed to the import callback. Indicates the host address points to PCIe BAR MMIO (device I/O memory).
  • FI_XPU_IMPORT_DEVICEMAP — passed to the import callback. Indicates the resulting pointer must be accessible from XPU kernel code.

fi_xpu_ctx_attr 🎀

Okay so fi_xpu_ctx_attr is the little struct that fi_xpu_ctx_query() hands back to you — it's basically the provider's vibe check for your XPU context, full of provider-specific output parameters. ✨

#define FI_XPU_CAP_EP      (1ULL << 0)
#define FI_XPU_CAP_CQ      (1ULL << 1)
#define FI_XPU_CAP_CNTR    (1ULL << 2)

struct fi_xpu_ctx_attr {
    uint64_t     caps;
    size_t       av_addr_size;
    size_t       mr_desc_size;
};
  • caps — a bitmask of the XPU capabilities this provider has going for it. Bestie, check these flags before you try to create XPU resources — if you go making an unsupported resource type anyway, you get -FI_ENOSYS thrown right back at you. 💀
    • FI_XPU_CAP_EP — provider supports XPU endpoints. That covers: device-side data transfer ops (post/dispatch) for the capabilities fi_getinfo returned (like FI_MSG, FI_TAGGED, FI_RMA, FI_ATOMIC), fi_ep_export_xpu to export the endpoint for device access, fi_av_lookup2 with FI_XPU to pull raw AV addresses, and fi_mr_get_xpu_desc with FI_XPU to pull raw MR descriptors. Full package. 👑
    • FI_XPU_CAP_CQ — provider supports XPU completion queues. That means fi_cq_export_xpu to export the CQ for device access, plus the device-side CQ functions (fi_xpu_cq_read, fi_xpu_cq_readfrom, fi_xpu_cq_readerr, fi_xpu_cq_sread, fi_xpu_cq_sreadfrom). 🦋
    • FI_XPU_CAP_CNTR — provider supports XPU counters. That's fi_cntr_export_xpu to export the counter for device access, plus the device-side counter functions (fi_xpu_cntr_read, fi_xpu_cntr_readerr, fi_xpu_cntr_wait, fi_xpu_cntr_add, fi_xpu_cntr_set, fi_xpu_cntr_adderr, fi_xpu_cntr_seterr). 🌸
  • av_addr_size — size in bytes of the raw address that fi_av_lookup2 gives you back (when it's called with the FI_XPU flag) for each AV entry. Every AV entry in a given context has the same size, so no surprises. 💅
  • mr_desc_size — size in bytes of the raw descriptor that fi_mr_get_xpu_desc returns (again, when called with the FI_XPU flag). All descriptors for a given context are the same size too. 🫶

fi_xpu_ctx_query 🔍

int fi_xpu_ctx_query(struct fid_xpu_ctx *ctx,
                     struct fi_xpu_ctx_attr *attr);

This is your ask-the-provider-about-XPU-context-parameters function. The caps field it returns tells you which XPU objects the provider actually supports (FI_XPU_CAP_EP, FI_XPU_CAP_CQ, FI_XPU_CAP_CNTR) — and yeah, check those flags before creating XPU resources, we said it once and we'll say it again. 😌

The av_addr_size and mr_desc_size fields it returns are exactly what you use to allocate correctly-sized buffers for your fi_av_lookup2 and fi_mr_get_xpu_desc calls. And heads up: different XPU contexts (like, ones targeting different devices) may report different sizes, so don't assume one size fits all. 🙈

EP 🌟

An XPU EP gets created by calling fi_endpoint2 with FI_XPU in the flags parameter, and fi_ep_attr->xpu_ctx pointing at an open XPU context. That's the whole recipe. 🍓

An EP created with FI_XPU can be bound to XPU CQs and XPU counters that were created with the same XPU context. Binding it to a host (non-XPU) CQ or counter? Allowed — but behavior is provider-specific, so iykyk, don't @ me. 💀 Control functions like fi_ep_bind, fi_enable, fi_setopt, and CM operations are only available to the host CPU. Host stays main character for control. 👑

Once it's bound and enabled, fi_ep_export_xpu exports the EP for device access. You (the caller) hand over a struct fid_xpu_ep, and the provider fills it with everything needed for the device-side data transfer functions. 🎀

Data transfer operations on the exported EP are only available to that specific XPU — this covers fi_msg(3), fi_rma(3), fi_tagged(3), fi_atomic(3). See fi_endpoint(3) for the details. ✨

CQ 🦋

An XPU CQ gets created by setting FI_XPU in fi_cq_attr->flags and fi_cq_attr->xpu_ctx to an open XPU context before you call fi_cq_open. Control operations (fi_cq_open, fi_close, fi_control) stay CPU-only. Like, that's just the rules. 💅

fi_cq_export_xpu exports the CQ for device access — you provide a struct fid_xpu_cq and the provider fills it with what the device-side completion functions need (fi_xpu_cq_read, fi_xpu_cq_readerr). But heads up: host-side read operations (fi_cq_read, fi_cq_sread) are NOT available on an exported CQ. 🚫 See fi_cq(3) for details.

Counters 🌈

An XPU counter gets created by setting FI_XPU in fi_cntr_attr->flags and fi_cntr_attr->xpu_ctx to an open XPU context before calling fi_cntr_open. Control operations (fi_cntr_open, fi_close) remain CPU-only. 🖥️

fi_cntr_export_xpu exports the counter for device access — you hand over a struct fid_xpu_cntr and the provider fills it with what the device-side counter functions need (fi_xpu_cntr_read, fi_xpu_cntr_wait). And again: host-side read operations (fi_cntr_read, fi_cntr_wait) are not available on an exported counter. 😭 See fi_cntr(3) for details.

AV 🧚

Okay, plot twist: the AV is not bound to an XPU context — it's a shared domain-level resource. Iconic, independent, does its own thing. 💜

Here's the flow, bestie: you query av_addr_size from fi_xpu_ctx_query(), allocate a buffer of that size, then call fi_av_lookup2 with the FI_XPU flag and the XPU context for each entry to retrieve the raw address, and copy the results into device-accessible memory. The raw address is a provider-specific representation that the XPU can actually use when posting work requests. ✨

fi_av_lookup2 is the extended version of fi_av_lookup — it accepts flags and an XPU context, which means the same AV can be queried for different XPU devices. See fi_av(3) for details. 🎀

MR 🔥

Same energy as the AV: the MR is not bound to an XPU context — shared domain-level resource. And the same MR can be queried with different XPU contexts. 🫶

So: query mr_desc_size from fi_xpu_ctx_query(), allocate a buffer, call fi_mr_get_xpu_desc with the FI_XPU flag and the XPU context to retrieve the raw descriptor (like, say, the hardware lkey), and copy the result to device-accessible memory. That raw descriptor replaces the desc parameter in data transfer operations initiated by the XPU. 💅

fi_mr_get_xpu_desc is an extended descriptor query that accepts flags and an XPU context, so the same MR can be queried for different XPU devices. See fi_mr(3) for details. 🌟

💻 DEVICE-SIDE API 💻

The device-side API is the set of communication functions you can call straight from XPU kernels. Include this header:

#include <rdma/fi_xpu_device.h>

This header gets compiled with a single XPU kernel compiler at a time. It supports XPU programming environments that give you a C/C++ interface:

  • NVIDIA CUDA (nvcc)
  • AMD ROCm HIP (hipcc)
  • Intel oneAPI Level Zero / SYCL (icpx -fsycl)

The same header covers all of the above — the FI_XPU_FUNC macro adapts the function qualifier (__device__, static inline, etc.) based on whichever compiler got detected at build time. One header to rule them all. 👑✨

Provider Dispatch Model 🎯

The device-side header uses a provider-identifier based dispatch model. Every exported XPU handle (fid_xpu_ep, fid_xpu_cq, fid_xpu_cntr) embeds struct fid_xpu as its first member. The provider fills in fid_xpu.prov_id during the export call (fi_ep_export_xpu, fi_cq_export_xpu, fi_cntr_export_xpu) with its assigned value from enum fi_xpu_provider:

enum fi_xpu_provider {
    FI_XPU_PROV_EFA = 1,
};

struct fid_xpu {
    uint32_t fclass;    /* FI_CLASS_EP, _CQ, _CNTR */
    uint32_t prov_id;   /* enum fi_xpu_provider */
    uint64_t prov_ctx;  /* provider-internal context */
};

struct fid_xpu_ep {
    struct fid_xpu fid;
};

struct fid_xpu_cq {
    struct fid_xpu fid;
};

struct fid_xpu_cntr {
    struct fid_xpu fid;
};

The generic dispatch functions take the typed handle (struct fid_xpu_ep *, fid_xpu_cq *, or fid_xpu_cntr *) and switch on fid.prov_id to route to the right provider-specific implementation. 💫

The prov_ctx field is how the provider locates all the state associated with the resource — it stores everything in device memory allocated via the XPU ops callbacks, and prov_ctx holds the address where that state lives. 🧠

And the tight range of provider IDs (starting at 1) lets the compiler generate an efficient jump table instead of a chain of comparisons. Slay, compiler optimization. ⚡

Scope 🌸

The scope parameter specifies the concurrency of the operation. All threads in the scope are required to issue the same operation — that's what lets the implementation optimize. No cap. 💅

  • FI_XPU_WORK_ITEM — CUDA: Thread · SYCL: Work item
  • FI_XPU_SUBGROUP — CUDA: Warp · SYCL: Subgroup
  • FI_XPU_WORK_GROUP — CUDA: Thread block · SYCL: Work group
  • FI_XPU_DEVICE — CUDA: Device · SYCL: Device

Data Transfer Operations 🚀

Device-side data transfer operations correspond to the standard libfabric APIs. Peep these man pages for the function signatures and semantics:

  • Message operations (fi_xpu_send, fi_xpu_recv): fi_msg(3)
  • Tagged operations (fi_xpu_tsend, fi_xpu_trecv): fi_tagged(3)
  • RMA operations (fi_xpu_write, fi_xpu_read): fi_rma(3)
  • Atomic operations (fi_xpu_atomic, fi_xpu_fetch_atomic, fi_xpu_compare_atomic): fi_atomic(3)

Completion Functions 🎀

Device-side completion functions operate on exported CQ and counter handles. Again, the man pages have the signatures and semantics:

  • Counter operations (fi_xpu_cntr_read, fi_xpu_cntr_wait, etc.): fi_cntr(3)
  • CQ operations (fi_xpu_cq_read, fi_xpu_cq_readerr, etc.): fi_cq(3)

💖 EXAMPLE 💖

okay bestie, this is the typical host-side setup flow for XPU-initiated RDMA communication using the XPU context. It’s giving “recipe for the whole vibe,” and we are not skipping steps. Here’s the exact code, preserved with love:

              /* 1. Query provider for FI_XPU support */
              struct fi_info *hints = fi_allocinfo();
              hints->caps = FI_MSG | FI_RMA | FI_XPU;
              fi_getinfo(FI_VERSION(2,7), NULL, NULL, 0, hints, &info);
              /* info->domain_attr->max_xpu_ctx_cnt > 0 confirms support */

              /* 2. Open domain */
              fi_domain(fabric, info, &domain, NULL);

              /* 3. Create XPU context for GPU 0 (with memory callbacks) */
              struct fi_xpu_ops my_ops = {
                  .size = sizeof(struct fi_xpu_ops),
                  .alloc = my_cuda_alloc,
                  .import = my_cuda_import,
                  .free = my_cuda_free,
              };
              struct fi_xpu_attr xpu_attr = {
                  .iface = FI_HMEM_CUDA,
                  .device = 0,
                  .ops = &my_ops,
              };
              struct fid_xpu_ctx *xpu_ctx;
              fi_xpu_ctx(domain, &xpu_attr, &xpu_ctx, NULL);

              /* 4. Query context for sizes and capabilities */
              struct fi_xpu_ctx_attr ctx_attr;
              fi_xpu_ctx_query(xpu_ctx, &ctx_attr);
              /* ctx_attr.caps indicates supported objects (FI_XPU_CAP_EP, etc.) */
              size_t av_entry_size = ctx_attr.av_addr_size;
              size_t desc_size = ctx_attr.mr_desc_size;

              /* 5. Create AV (domain-level, no xpu_ctx binding) */
              struct fi_av_attr av_attr = { .type = FI_AV_TABLE };
              fi_av_open(domain, &av_attr, &av, NULL);
              fi_av_insert(av, peer_addr, 1, &fi_addr, 0, NULL);

              /* 6. Create CQ and counter with xpu_ctx */
              struct fi_cq_attr cq_attr = { .format = FI_CQ_FORMAT_DATA,
                                            .xpu_ctx = xpu_ctx };
              fi_cq_open(domain, &cq_attr, &cq, NULL);

              struct fi_cntr_attr cntr_attr = { .events = FI_CNTR_EVENTS_COMP,
                                                .xpu_ctx = xpu_ctx };
              fi_cntr_open(domain, &cntr_attr, &cntr, NULL);

              /* 7. Create EP with xpu_ctx */
              info->ep_attr->xpu_ctx = xpu_ctx;
              fi_endpoint2(domain, info, &ep, FI_XPU, NULL);

              /* 8. Register MR (domain-level, no special flags) */
              struct fi_mr_attr mr_attr = { .mr_iov = &iov, .iov_count = 1,
                                            .access = FI_SEND | FI_RECV };
              fi_mr_regattr(domain, &mr_attr, 0, &mr);

              /* 9. Bind and enable EP */
              fi_ep_bind(ep, &av->fid, 0);
              fi_ep_bind(ep, &cq->fid, FI_TRANSMIT | FI_RECV);
              fi_ep_bind(ep, &cntr->fid, FI_SEND);
              fi_enable(ep);

              /* 10. Export EP/CQ/counter for device-side use */
              struct fid_xpu_ep xpu_ep;
              struct fid_xpu_cq xpu_cq;
              struct fid_xpu_cntr xpu_cntr;
              fi_ep_export_xpu(ep, 0, &xpu_ep);
              fi_cq_export_xpu(cq, 0, &xpu_cq);
              fi_cntr_export_xpu(cntr, 0, &xpu_cntr);

              /* 11. Get raw AV addr and MR desc for device-side use */
              void *raw_addr = malloc(av_entry_size);
              size_t len = av_entry_size;
              fi_av_lookup2(av, fi_addr, raw_addr, &len, FI_XPU, xpu_ctx);

              void *raw_desc = malloc(desc_size);
              len = desc_size;
              fi_mr_get_xpu_desc(mr, raw_desc, &len, FI_XPU, xpu_ctx);

              /* 12. Copy raw AV addr, MR desc, and exported handles to device memory */
              copy_to_device(gpu_addr, raw_addr, av_entry_size);
              copy_to_device(gpu_desc, raw_desc, desc_size);
              copy_to_device(gpu_xpu_ep, &xpu_ep, sizeof(xpu_ep));
              copy_to_device(gpu_xpu_cntr, &xpu_cntr, sizeof(xpu_cntr));

              /* 13. Launch device kernel — posts operations and polls completions
               *     using exported handles (already device-accessible):
               *
               *     __global__ void my_kernel(struct fid_xpu_ep *ep,
               *                               struct fid_xpu_cntr *cntr,
               *                               void *peer, void *desc, void *buf) {
               *         uint64_t prev = fi_xpu_cntr_read(cntr, FI_XPU_WORK_ITEM);
               *
               *         fi_xpu_send(ep, buf, 64, desc, 0, peer, NULL,
               *                     0, FI_XPU_WORK_ITEM);
               *
               *         fi_xpu_cntr_wait(cntr, prev + 1, -1, FI_XPU_WORK_ITEM);
               *     }
               */
              launch_kernel(my_kernel, gpu_xpu_ep, gpu_xpu_cntr, gpu_addr, gpu_desc, gpu_buf);

              /* 14. Cleanup */
              fi_close(&ep->fid);
              fi_close(&cq->fid);
              fi_close(&cntr->fid);
              fi_close(&mr->fid);
              fi_close(&av->fid);
              fi_close(&xpu_ctx->fid);
              fi_close(&domain->fid);

the comments inside are literally little sticky notes doing a lot of the explaining, so read them, bestie 💅✨

🔍 SEE ALSO 🔍

if you’re hopping around the docs, these are your next stops — every single one, no cap:

👑 AUTHORS 👑

credit where it’s due: OpenFabrics. 🫶

💖 that’s the whole page, bestie — from the Libfabric Programmer’s Manual, dated 2026-09-18, and it’s fi_xpu(3) 💅✨

📦 where this page comes from