'\" t .\" Automatically generated by Pandoc 3.1.3 .\" .\" Define V font for inline verbatim, using C font in formats .\" that render this, and otherwise B font. .ie "\f[CB]x\f[]"x" \{\ . ftr V B . ftr VI BI . ftr VB B . ftr VBI BI .\} .el \{\ . ftr V CR . ftr VI CI . ftr VB CB . ftr VBI CBI .\} .TH "fi_xpu" "3" "2026\-09\-18" "Libfabric Programmer\[cq]s Manual" "Libfabric v2.7.0" .hy .SH NAME .PP fi_xpu - OFI XPU API .SH SYNOPSIS .IP .nf \f[C] #include \f[R] .fi .SH OVERVIEW .PP 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. .PP 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 (\f[V]fi_av\f[R](3), \f[V]fi_endpoint\f[R](3), \f[V]fi_cq\f[R](3), \f[V]fi_cntr\f[R](3), \f[V]fi_mr\f[R](3)). .SH XPU CONTEXT .PP Related EP, CQ, and counter resources created with FI_XPU must all target the same XPU. The \f[V]fid_xpu_ctx\f[R] 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. .PP An XPU context is created from a domain: .IP .nf \f[C] int fi_xpu_ctx(struct fid_domain *domain, struct fi_xpu_attr *attr, struct fid_xpu_ctx **ctx, void *context); \f[R] .fi .PP The number of XPU contexts that a domain supports is indicated by the \f[V]max_xpu_ctx_cnt\f[R] field in \f[V]fi_domain_attr\f[R]. 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. .PP The XPU context is passed to EP, CQ, and counter creation calls via a pointer in their respective attribute structures (\f[V]fi_ep_attr->xpu_ctx\f[R], \f[V]fi_cq_attr->xpu_ctx\f[R], \f[V]fi_cntr_attr->xpu_ctx\f[R]). 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. .PP The XPU context is also passed to \f[V]fi_av_lookup2\f[R] and \f[V]fi_mr_get_xpu_desc\f[R] 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 \[em] the same AV or MR can be queried with different XPU contexts to get device-specific representations. .SH fi_xpu_attr .PP The \f[V]fi_xpu_attr\f[R] structure is an input argument that identifies the target XPU device and provides optional memory management callbacks. It is passed to \f[V]fi_xpu_ctx()\f[R] during context creation. .IP .nf \f[C] 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; }; \f[R] .fi .TP \f[I]iface\f[R] The heterogeneous memory interface type (e.g., FI_HMEM_CUDA, FI_HMEM_ZE). .TP \f[I]device\f[R] The XPU device ordinal. .TP \f[I]ops\f[R] Optional pointer to a \f[V]fi_xpu_ops\f[R] structure containing memory management callbacks. If NULL, the provider uses its default mechanisms. .SS fi_xpu_ops .PP The \f[V]fi_xpu_ops\f[R] 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. .TP \f[I]size\f[R] Must be set to \f[V]sizeof(struct fi_xpu_ops)\f[R]. This allows future expansion of the structure while maintaining backward compatibility. The provider uses this field to determine which callback fields are present. .TP \f[I]alloc\f[R] Allocate XPU memory. The provider calls this when it needs device memory (e.g., for hardware queue buffers). If the \f[V]FI_XPU_ALLOC_DMABUF\f[R] flag is set, the consumer must also export a DMA-BUF file descriptor so the provider can DMA directly to the memory. .TP \f[I]import\f[R] 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: \f[V]FI_XPU_IMPORT_IOMEMORY\f[R] for PCIe BAR MMIO regions, \f[V]FI_XPU_IMPORT_DEVICEMAP\f[R] for addresses that must be accessible from XPU kernels. .TP \f[I]free\f[R] Release memory previously allocated via alloc. .SS Memory Callback Flags .TP \f[I]FI_XPU_ALLOC_DMABUF\f[R] Passed to the alloc callback. Indicates the allocation must be exportable as a DMA-BUF fd for provider access. .TP \f[I]FI_XPU_IMPORT_IOMEMORY\f[R] Passed to the import callback. Indicates the host address points to PCIe BAR MMIO (device I/O memory). .TP \f[I]FI_XPU_IMPORT_DEVICEMAP\f[R] Passed to the import callback. Indicates the resulting pointer must be accessible from XPU kernel code. .SH fi_xpu_ctx_attr .PP The \f[V]fi_xpu_ctx_attr\f[R] structure is returned by \f[V]fi_xpu_ctx_query()\f[R]. It contains provider-specific output parameters for the given XPU context. .IP .nf \f[C] #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; }; \f[R] .fi .TP \f[I]caps\f[R] 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. .IP \[bu] 2 \f[V]FI_XPU_CAP_EP\f[R]: Provider supports XPU endpoints. This includes: device-side data transfer operations (post/dispatch) for the capabilities returned by \f[V]fi_getinfo\f[R] (e.g., FI_MSG, FI_TAGGED, FI_RMA, FI_ATOMIC), \f[V]fi_ep_export_xpu\f[R] to export the endpoint for device access, \f[V]fi_av_lookup2\f[R] with FI_XPU to retrieve raw AV addresses, and \f[V]fi_mr_get_xpu_desc\f[R] with FI_XPU to retrieve raw MR descriptors. .IP \[bu] 2 \f[V]FI_XPU_CAP_CQ\f[R]: Provider supports XPU completion queues. This includes: \f[V]fi_cq_export_xpu\f[R] to export the CQ for device access, and the device-side CQ functions (\f[V]fi_xpu_cq_read\f[R], \f[V]fi_xpu_cq_readfrom\f[R], \f[V]fi_xpu_cq_readerr\f[R], \f[V]fi_xpu_cq_sread\f[R], \f[V]fi_xpu_cq_sreadfrom\f[R]). .IP \[bu] 2 \f[V]FI_XPU_CAP_CNTR\f[R]: Provider supports XPU counters. This includes: \f[V]fi_cntr_export_xpu\f[R] to export the counter for device access, and the device-side counter functions (\f[V]fi_xpu_cntr_read\f[R], \f[V]fi_xpu_cntr_readerr\f[R], \f[V]fi_xpu_cntr_wait\f[R], \f[V]fi_xpu_cntr_add\f[R], \f[V]fi_xpu_cntr_set\f[R], \f[V]fi_xpu_cntr_adderr\f[R], \f[V]fi_xpu_cntr_seterr\f[R]). .TP \f[I]av_addr_size\f[R] Size in bytes of the raw address returned by \f[V]fi_av_lookup2\f[R] (when called with \f[V]FI_XPU\f[R] flag) for each AV entry. All AV entries for a given context have the same size. .TP \f[I]mr_desc_size\f[R] Size in bytes of the raw descriptor returned by \f[V]fi_mr_get_xpu_desc\f[R] (when called with \f[V]FI_XPU\f[R] flag). All descriptors for a given context have the same size. .SS fi_xpu_ctx_query .IP .nf \f[C] int fi_xpu_ctx_query(struct fid_xpu_ctx *ctx, struct fi_xpu_ctx_attr *attr); \f[R] .fi .PP Query the provider for XPU context parameters. The returned \f[V]caps\f[R] 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 \f[V]av_addr_size\f[R] and \f[V]mr_desc_size\f[R] fields are used to allocate appropriately sized buffers for \f[V]fi_av_lookup2\f[R] and \f[V]fi_mr_get_xpu_desc\f[R] calls. Different XPU contexts (targeting different devices) may report different sizes. .SS EP .PP An XPU EP is created by calling \f[V]fi_endpoint2\f[R] with \f[V]FI_XPU\f[R] in the flags parameter and setting \f[V]fi_ep_attr->xpu_ctx\f[R] to an open XPU context. .PP 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 \f[V]fi_ep_bind\f[R], \f[V]fi_enable\f[R], \f[V]fi_setopt\f[R], and CM operations are available only to the host CPU. .PP Once bound and enabled, \f[V]fi_ep_export_xpu\f[R] exports the EP for device access. The caller provides a \f[V]struct fid_xpu_ep\f[R] which the provider fills with the information needed for device-side data transfer functions. .PP Data transfer operations on the exported EP are only available to the given XPU. This includes \f[V]fi_msg\f[R](3), \f[V]fi_rma\f[R](3), \f[V]fi_tagged\f[R](3), \f[V]fi_atomic\f[R](3). See \f[V]fi_endpoint\f[R](3) for details. .SS CQ .PP An XPU CQ is created by setting \f[V]FI_XPU\f[R] in \f[V]fi_cq_attr->flags\f[R] and \f[V]fi_cq_attr->xpu_ctx\f[R] to an open XPU context before calling \f[V]fi_cq_open\f[R]. Control operations (\f[V]fi_cq_open\f[R], \f[V]fi_close\f[R], \f[V]fi_control\f[R]) remain CPU-only. .PP \f[V]fi_cq_export_xpu\f[R] exports the CQ for device access. The caller provides a \f[V]struct fid_xpu_cq\f[R] which the provider fills with the information needed for the device-side completion functions (\f[V]fi_xpu_cq_read\f[R], \f[V]fi_xpu_cq_readerr\f[R]). Host-side read operations (\f[V]fi_cq_read\f[R], \f[V]fi_cq_sread\f[R]) are not available on an exported CQ. See \f[V]fi_cq\f[R](3) for details. .SS Counters .PP An XPU counter is created by setting \f[V]FI_XPU\f[R] in \f[V]fi_cntr_attr->flags\f[R] and \f[V]fi_cntr_attr->xpu_ctx\f[R] to an open XPU context before calling \f[V]fi_cntr_open\f[R]. Control operations (\f[V]fi_cntr_open\f[R], \f[V]fi_close\f[R]) remain CPU-only. .PP \f[V]fi_cntr_export_xpu\f[R] exports the counter for device access. The caller provides a \f[V]struct fid_xpu_cntr\f[R] which the provider fills with the information needed for the device-side counter functions (\f[V]fi_xpu_cntr_read\f[R], \f[V]fi_xpu_cntr_wait\f[R]). Host-side read operations (\f[V]fi_cntr_read\f[R], \f[V]fi_cntr_wait\f[R]) are not available on an exported counter. See \f[V]fi_cntr\f[R](3) for details. .SS AV .PP The AV is not bound to an XPU context \[em] it is a shared domain-level resource. .PP The application queries \f[V]av_addr_size\f[R] from \f[V]fi_xpu_ctx_query()\f[R], allocates a buffer of that size, calls \f[V]fi_av_lookup2\f[R] with \f[V]FI_XPU\f[R] 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. .PP \f[V]fi_av_lookup2\f[R] is an extended version of \f[V]fi_av_lookup\f[R] that accepts flags and an XPU context, allowing the same AV to be queried for different XPU devices. See \f[V]fi_av\f[R](3) for details. .SS MR .PP The MR is not bound to an XPU context \[em] it is a shared domain-level resource. The same MR can be queried with different XPU contexts. .PP The application queries \f[V]mr_desc_size\f[R] from \f[V]fi_xpu_ctx_query()\f[R], allocates a buffer, calls \f[V]fi_mr_get_xpu_desc\f[R] with the \f[V]FI_XPU\f[R] 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. .PP \f[V]fi_mr_get_xpu_desc\f[R] is an extended descriptor query that accepts flags and an XPU context, allowing the same MR to be queried for different XPU devices. See \f[V]fi_mr\f[R](3) for details. .SH DEVICE-SIDE API .PP The device-side API provides communication functions callable from XPU kernels. Include the following header: .IP .nf \f[C] #include \f[R] .fi .PP This header is compiled with a single XPU kernel compiler at a time. It supports XPU programming environments that provide a C/C++ interface: .IP \[bu] 2 NVIDIA CUDA (nvcc) .IP \[bu] 2 AMD ROCm HIP (hipcc) .IP \[bu] 2 Intel oneAPI Level Zero / SYCL (icpx -fsycl) .PP The same header covers all of the above \[em] the \f[V]FI_XPU_FUNC\f[R] macro adapts the function qualifier (\f[V]__device__\f[R], \f[V]static inline\f[R], etc.) based on the compiler detected at build time. .SS Provider Dispatch Model .PP The device-side header uses a provider-identifier based dispatch model. Each exported XPU handle (\f[V]fid_xpu_ep\f[R], \f[V]fid_xpu_cq\f[R], \f[V]fid_xpu_cntr\f[R]) embeds \f[V]struct fid_xpu\f[R] as its first member. The provider populates \f[V]fid_xpu.prov_id\f[R] during the export call (\f[V]fi_ep_export_xpu\f[R], \f[V]fi_cq_export_xpu\f[R], \f[V]fi_cntr_export_xpu\f[R]) with its assigned value from \f[V]enum fi_xpu_provider\f[R]: .IP .nf \f[C] 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; }; \f[R] .fi .PP The generic dispatch functions take the typed handle (\f[V]struct fid_xpu_ep *\f[R], \f[V]fid_xpu_cq *\f[R], or \f[V]fid_xpu_cntr *\f[R]) and switch on \f[V]fid.prov_id\f[R] to route to the appropriate provider-specific implementation. The \f[V]prov_ctx\f[R] field provides a mechanism for the provider to locate all state associated with the resource \[em] the provider stores everything in device memory allocated via the XPU ops callbacks, and \f[V]prov_ctx\f[R] holds the address where that state resides. .PP The tight range of provider IDs (starting at 1) allows the compiler to generate an efficient jump table rather than a chain of comparisons. .SS Scope .PP 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. .PP .TS tab(@); l l l. T{ Scope T}@T{ CUDA T}@T{ SYCL T} _ T{ FI_XPU_WORK_ITEM T}@T{ Thread T}@T{ Work item T} T{ FI_XPU_SUBGROUP T}@T{ Warp T}@T{ Subgroup T} T{ FI_XPU_WORK_GROUP T}@T{ Thread block T}@T{ Work group T} T{ FI_XPU_DEVICE T}@T{ Device T}@T{ Device T} .TE .SS Data Transfer Operations .PP Device-side data transfer operations correspond to the standard libfabric APIs. See the following man pages for function signatures and semantics: .IP \[bu] 2 Message operations (fi_xpu_send, fi_xpu_recv): \f[V]fi_msg\f[R](3) .IP \[bu] 2 Tagged operations (fi_xpu_tsend, fi_xpu_trecv): \f[V]fi_tagged\f[R](3) .IP \[bu] 2 RMA operations (fi_xpu_write, fi_xpu_read): \f[V]fi_rma\f[R](3) .IP \[bu] 2 Atomic operations (fi_xpu_atomic, fi_xpu_fetch_atomic, fi_xpu_compare_atomic): \f[V]fi_atomic\f[R](3) .SS Completion Functions .PP Device-side completion functions operate on exported CQ and counter handles. See the following man pages for function signatures and semantics: .IP \[bu] 2 Counter operations (fi_xpu_cntr_read, fi_xpu_cntr_wait, etc.): \f[V]fi_cntr\f[R](3) .IP \[bu] 2 CQ operations (fi_xpu_cq_read, fi_xpu_cq_readerr, etc.): \f[V]fi_cq\f[R](3) .SH EXAMPLE .PP The following illustrates the typical host-side setup flow for XPU-initiated RDMA communication using the XPU context: .IP .nf \f[C] /* 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 \[em] 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); \f[R] .fi .SH SEE ALSO .PP \f[V]fi_getinfo\f[R](3), \f[V]fi_endpoint\f[R](3), \f[V]fi_cq\f[R](3), \f[V]fi_cntr\f[R](3), \f[V]fi_mr\f[R](3), \f[V]fi_av\f[R](3), \f[V]fi_set_ops\f[R](3) .SH AUTHORS OpenFabrics.