--- name: "fi_xpu" section: "3" platform: "linux" package: "libfabric" version: "2.7.0-1" release: "Arch Linux 2026-11-01" platforms: ["linux"] url: "https://new.carta.tech/fi_xpu" --- # fi\_xpu(3) OFI XPU API ## NAME fi\_xpu - OFI XPU API ## SYNOPSIS - ``` #include ``` ## OVERVIEW The FI\_XPU capability allows for a specified XPU (or device) to access a given libfabric allocated resource for data path operations, with control operations being managed by host (CPU) software. A resource (endpoint, completion queue, counter) which has been configured with this capability may only be accessed by the given XPU for any data path operation, such as submitting data transfers or reading completions. As a general rule, 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_av/3.md), [fi\_endpoint(3)](/fi_endpoint/3.md), [fi\_cq(3)](/fi_cq/3.md), [fi\_cntr(3)](/fi_cntr/3.md), [fi\_mr(3)](/fi_mr/3.md)). ## 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 The fi\_xpu\_ctx\_attr structure is returned by fi\_xpu\_ctx\_query(). It contains provider-specific output parameters for the given XPU context. - ``` #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* Bitmask of XPU capabilities supported by this provider. The application should check these flags before attempting to create XPU resources. Attempting to create an unsupported resource type will return -FI\_ENOSYS. - • FI\_XPU\_CAP\_EP: Provider supports XPU endpoints. This includes: device-side data transfer operations (post/dispatch) for the capabilities returned by fi\_getinfo (e.g., 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 retrieve raw AV addresses, and fi\_mr\_get\_xpu\_desc with FI\_XPU to retrieve raw MR descriptors. - • FI\_XPU\_CAP\_CQ: Provider supports XPU completion queues. This includes: fi\_cq\_export\_xpu to export the CQ for device access, and 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. This includes: fi\_cntr\_export\_xpu to export the counter for device access, and 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 returned by fi\_av\_lookup2 (when called with FI\_XPU flag) for each AV entry. All AV entries for a given context have the same size. - *mr\_desc\_size* Size in bytes of the raw descriptor returned by fi\_mr\_get\_xpu\_desc (when called with FI\_XPU flag). All descriptors for a given context have the same size. ### fi\_xpu\_ctx\_query - ``` int fi_xpu_ctx_query(struct fid_xpu_ctx *ctx, struct fi_xpu_ctx_attr *attr); ``` Query the provider for XPU context parameters. The returned caps field indicates which XPU objects the provider supports (FI\_XPU\_CAP\_EP, FI\_XPU\_CAP\_CQ, FI\_XPU\_CAP\_CNTR). The application should check these flags before attempting to create XPU resources. The returned av\_addr\_size and mr\_desc\_size fields are used to allocate appropriately sized buffers for fi\_av\_lookup2 and fi\_mr\_get\_xpu\_desc calls. Different XPU contexts (targeting different devices) may report different sizes. ### EP An XPU EP is created by calling fi\_endpoint2 with FI\_XPU in the flags parameter and setting fi\_ep\_attr-\>xpu\_ctx to an open XPU context. An EP created with FI\_XPU may be bound to XPU CQs and XPU counters created with the same XPU context. Binding to a host (non-XPU) CQ or counter is allowed but behavior is provider-specific. Control functions such as fi\_ep\_bind, fi\_enable, fi\_setopt, and CM operations are available only to the host CPU. Once bound and enabled, fi\_ep\_export\_xpu exports the EP for device access. The caller provides a struct fid\_xpu\_ep which the provider fills with the information needed for device-side data transfer functions. Data transfer operations on the exported EP are only available to the given XPU. This includes [fi\_msg(3)](/fi_msg/3.md), [fi\_rma(3)](/fi_rma/3.md), [fi\_tagged(3)](/fi_tagged/3.md), [fi\_atomic(3)](/fi_atomic/3.md). See [fi\_endpoint(3)](/fi_endpoint/3.md) for details. ### CQ An XPU CQ is created by setting FI\_XPU in fi\_cq\_attr-\>flags and fi\_cq\_attr-\>xpu\_ctx to an open XPU context before calling fi\_cq\_open. Control operations (fi\_cq\_open, fi\_close, fi\_control) remain CPU-only. fi\_cq\_export\_xpu exports the CQ for device access. The caller provides a struct fid\_xpu\_cq which the provider fills with the information needed for the device-side completion functions (fi\_xpu\_cq\_read, fi\_xpu\_cq\_readerr). Host-side read operations (fi\_cq\_read, fi\_cq\_sread) are not available on an exported CQ. See [fi\_cq(3)](/fi_cq/3.md) for details. ### Counters An XPU counter is 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. The caller provides a struct fid\_xpu\_cntr which the provider fills with the information needed for the device-side counter functions (fi\_xpu\_cntr\_read, fi\_xpu\_cntr\_wait). Host-side read operations (fi\_cntr\_read, fi\_cntr\_wait) are not available on an exported counter. See [fi\_cntr(3)](/fi_cntr/3.md) for details. ### AV The AV is not bound to an XPU context — it is a shared domain-level resource. The application queries av\_addr\_size from fi\_xpu\_ctx\_query(), allocates a buffer of that size, calls fi\_av\_lookup2 with FI\_XPU flag and the XPU context for each entry to retrieve the raw address, and copies the results to device-accessible memory. The raw address is a provider-specific representation usable by the XPU when posting work requests. fi\_av\_lookup2 is an extended version of fi\_av\_lookup that accepts flags and an XPU context, allowing the same AV to be queried for different XPU devices. See [fi\_av(3)](/fi_av/3.md) for details. ### MR The MR is not bound to an XPU context — it is a shared domain-level resource. The same MR can be queried with different XPU contexts. The application queries mr\_desc\_size from fi\_xpu\_ctx\_query(), allocates a buffer, calls fi\_mr\_get\_xpu\_desc with the FI\_XPU flag and the XPU context to retrieve the raw descriptor (e.g., the hardware lkey), and copies the result to device-accessible memory. The 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, allowing the same MR to be queried for different XPU devices. See [fi\_mr(3)](/fi_mr/3.md) for details. ## DEVICE-SIDE API The device-side API provides communication functions callable from XPU kernels. Include the following header: - ``` #include ``` This header is compiled with a single XPU kernel compiler at a time. It supports XPU programming environments that provide 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 the compiler detected at build time. ### Provider Dispatch Model The device-side header uses a provider-identifier based dispatch model. Each exported XPU handle (fid\_xpu\_ep, fid\_xpu\_cq, fid\_xpu\_cntr) embeds struct fid\_xpu as its first member. The provider populates 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 appropriate provider-specific implementation. The prov\_ctx field provides a mechanism for the provider to locate all state associated with the resource — the provider stores everything in device memory allocated via the XPU ops callbacks, and prov\_ctx holds the address where that state resides. The tight range of provider IDs (starting at 1) allows the compiler to generate an efficient jump table rather than a chain of comparisons. ### Scope The scope parameter specifies the concurrency of the operation. All threads in the scope are required to issue the same operation. This allows the implementation to optimize. | | | | | --- | --- | --- | | Scope | CUDA | SYCL | | FI\_XPU\_WORK\_ITEM | Thread | Work item | | FI\_XPU\_SUBGROUP | Warp | Subgroup | | FI\_XPU\_WORK\_GROUP | Thread block | Work group | | FI\_XPU\_DEVICE | Device | Device | ### Data Transfer Operations Device-side data transfer operations correspond to the standard libfabric APIs. See the following man pages for function signatures and semantics: - • Message operations (fi\_xpu\_send, fi\_xpu\_recv): [fi\_msg(3)](/fi_msg/3.md) - • Tagged operations (fi\_xpu\_tsend, fi\_xpu\_trecv): [fi\_tagged(3)](/fi_tagged/3.md) - • RMA operations (fi\_xpu\_write, fi\_xpu\_read): [fi\_rma(3)](/fi_rma/3.md) - • Atomic operations (fi\_xpu\_atomic, fi\_xpu\_fetch\_atomic, fi\_xpu\_compare\_atomic): [fi\_atomic(3)](/fi_atomic/3.md) ### Completion Functions Device-side completion functions operate on exported CQ and counter handles. See the following man pages for function signatures and semantics: - • Counter operations (fi\_xpu\_cntr\_read, fi\_xpu\_cntr\_wait, etc.): [fi\_cntr(3)](/fi_cntr/3.md) - • CQ operations (fi\_xpu\_cq\_read, fi\_xpu\_cq\_readerr, etc.): [fi\_cq(3)](/fi_cq/3.md) ## EXAMPLE The following illustrates the typical host-side setup flow for XPU-initiated RDMA communication using the XPU context: - ``` /* 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); ``` ## SEE ALSO [fi\_getinfo(3)](/fi_getinfo/3.md), [fi\_endpoint(3)](/fi_endpoint/3.md), [fi\_cq(3)](/fi_cq/3.md), [fi\_cntr(3)](/fi_cntr/3.md), [fi\_mr(3)](/fi_mr/3.md), [fi\_av(3)](/fi_av/3.md), fi\_set\_ops(3) ## AUTHORS OpenFabrics.