UBMEM and CUDA IPC

导言

进程 A 已经在加速卡上生成一块数据，进程 B 能不能直接使用？把指针发过去是否就够了？把 UBMEM 和 CUDA IPC 放在一起看，最有价值的切入点就是这个问题。

本文面向知道进程、指针和设备内存，但尚未写过跨进程设备通信的读者。先用 CUDA IPC 建立“共享句柄 → 本地映射 → 同步访问 → 有序释放”的直觉，再沿 HCOMM 的真实实现解释 UBMEM。它们的交集是建立可访问的内存视图；地址有效、数据就绪和访问成本，仍是三个需要分别回答的问题。

为什么不能直接传指针

假设进程 A 在 GPU 0 上存有四个 int32：[10, 20, 30, 40]，进程 B 想读其中第三个数。这里只用 16 字节描述逻辑数据，实际分配可以更大。

最容易想到的办法，是把 A 的设备指针当作一个整数，通过 socket 发给 B。但指针数值只在它所属的地址空间和访问上下文里有意义。B 收到这个数，并没有因此获得对应的设备映射。NVIDIA 的 IPC 编程指南明确区分了“本进程的设备指针”和“可以跨进程交换的句柄”。[^cuda-ipc]

CUDA 的 UVA（Unified Virtual Addressing，统一虚拟寻址）也没有消除这个边界：它统一的是同一个 OS 进程内的 host 与设备虚拟地址空间，不能让 A 的原始指针直接在 B 中生效。[^cuda-uva]

可以把它理解为两个人各自使用一份地图：A 地图上的坐标，不是 B 地图上的通行证。共享机制需要告诉驱动“要访问的是哪一个内存对象”，由驱动在 B 的地址空间中建立有效入口。

这个过程涉及三个不同对象：

• 物理内存：真正保存 [10, 20, 30, 40] 的存储区域。
• 共享句柄或 key：跨进程交换的资源描述。它不是数组内容，也不能直接解引用。
• 本地映射地址：导入成功后，B 用来访问该区域的设备指针。

下图只画地址关系，帮助分清“一份物理数据”和“两个访问入口”。

自绘示意图。蓝色线表示地址映射关系，紫色虚线表示句柄等元数据的交换；没有画出实际数据传输。图中的 16 字节是逻辑数据范围，不代表驱动实际分配或共享的粒度。

如果 A 的本地基址是 pA，B 导入得到 pB，那么第三个元素分别通过字节地址 pA + 8、pB + 8 定位。两个基址不必相同，它们可以指向同一个底层对象。这里的加法按字节计算；如果使用 int32_t*，相应表达式是 pA[2] 与 pB[2]。CUDA 明确不保证不同进程导入同一个 handle 后得到相同地址。[^cuda-device]

这也给 tensor 共享一个直接启示：交换内容通常还需要带上 allocation 内的偏移、有效字节数，以及上层理解数据所需的 dtype、shape、stride。IPC 负责内存对象，tensor 的语义仍由应用解释。

CUDA IPC 怎样共享设备内存

CUDA IPC 的 IPC 是 Interprocess Communication，即进程间通信。本文先讨论经典 cudaIpc* API，在同一台主机的两个独立 Linux 进程中共享设备内存。两个进程可以使用同一块 GPU，也可以在支持 peer access 的条件下使用不同 GPU。

导出、导入与关闭

经典路径的关键步骤很少：

1. A 用 cudaMalloc 分配内存，并用 cudaIpcGetMemHandle 从分配基址获得 cudaIpcMemHandle_t。
2. A 通过普通进程间通道，把 handle 和必要的应用元数据交给 B。
3. B 调用 cudaIpcOpenMemHandle，拿到对 B 有效的设备指针。
4. B 的设备任务完成后，用 cudaIpcCloseMemHandle 关闭自己的映射。
5. A 确认所有使用者都已经关闭映射，才释放原始分配。[^cuda-device]

此时 A 中的数据没有被 IPC 自动复制成 B 的一份私有数据。B 可以在满足设备访问条件时直接使用导入地址，也可以主动把内容复制到自己管理的 buffer。前者是共享访问，后者会产生新的物理副本。

跨 GPU 时，cudaIpcMemLazyEnablePeerAccess 允许导入操作尝试启用 peer access；应用可以先通过 cudaDeviceCanAccessPeer 检查能力。映射成功后的设备指针，也不意味着 CPU 可以把它当普通 host 指针读取。[^cuda-device]

能访问，还不代表可以开始读

假设 A 已导出了 handle，但产生数组的 kernel 仍在执行。B 即使成功导入，也可能读得太早。共享内存的建立，与生产者工作完成，没有自动的先后关系。

CUDA 还提供 IPC event。A 创建可跨进程共享的事件，把它记录在生产者 stream 上；B 导入事件，让自己的消费者 stream 等待它。事件需要同时带有 cudaEventInterprocess 与 cudaEventDisableTiming 标志，等待操作则使用 cudaStreamWaitEvent。[^cuda-device][^cuda-stream]

下面是本文编写的一次性交接伪代码，不是可直接编译的程序。两端已选好设备并创建 stream；省略了变量声明、返回值检查、socket 实现与异常清理。produce_four_ints 写入四个整数，consume_four_ints 读取它们；send/recv 是可靠的 host 消息通道，epoch 是本次交接编号。

// 进程 A：分配的所有者
cudaMalloc(&pA, 2 * 1024 * 1024);  // 共享分配为 2 MiB，逻辑数据仅 16 B
cudaEventCreateWithFlags(&ready,
    cudaEventInterprocess | cudaEventDisableTiming);
cudaIpcGetMemHandle(&mem_handle, pA);
cudaIpcGetEventHandle(&event_handle, ready);
produce_four_ints<<<1, 32, 0, streamA>>>(pA);
cudaEventRecord(ready, streamA);
send_to_B(mem_handle, event_handle, 16, epoch);  // Record 调用成功后再发
recv_closed_from_B(epoch);                     // 等 B 完成并关闭映射
cudaEventDestroy(ready);
cudaFree(pA);

// 进程 B：使用者；这一段在另一个进程中运行
recv_from_A(mem_handle, event_handle, valid_bytes, epoch);
cudaIpcOpenMemHandle(&pB, mem_handle, cudaIpcMemLazyEnablePeerAccess);
cudaIpcOpenEventHandle(&readyB, event_handle);
cudaStreamWaitEvent(streamB, readyB, 0);
consume_four_ints<<<1, 32, 0, streamB>>>(pB);
cudaStreamSynchronize(streamB);  // 成功返回后，消费者才真正完成
cudaEventDestroy(readyB);
cudaIpcCloseMemHandle(pB);
send_closed_to_A(epoch);

这里有两个容易省掉、却不能省掉的顺序。

A 必须先调用 Record，再允许 B 提交 Wait。CUDA event 在第一次记录前代表空工作集；等待操作使用的是调用当时已经捕获的状态，不会自动等待未来才发生的记录。消息在 Record 成功返回之后发送，保证 B 等待的是这一轮生产任务。Record 返回并不表示 kernel 已完成，真正的设备等待由 event 承担。[^cuda-event]

B 必须完成访问、关闭映射，再通知 A 释放。内存所有权留在 A，B 只拥有导入映射；B 不应对 pB 调用 cudaFree。A 在 B 关闭前释放原始内存属于未定义行为。[^cuda-device]

循环复用 buffer 时，还需要防止 A 在 B 读完前覆盖下一轮数据。上例用一次性交接避开了事件重录和多轮并发；把它改成流水线，必须补上轮次、槽位和消费完成协议。

分配方式决定共享接口

cudaIpcGetMemHandle 对应经典 cudaMalloc 分配，不能把所有 CUDA 指针都塞进去：

• **cudaMallocManaged**：经典 IPC 不支持这类分配。Unified Memory 管理 CPU/GPU 访问和数据驻留，不能因此推导出任意进程自动共享。[^cuda-ipc]
• cudaMallocAsync / 内存池：需要支持共享的显式 pool，以及 cudaMemPoolExportToShareableHandle、cudaMemPoolImportFromShareableHandle 和 allocation 的 export/import 配套接口。不能直接沿用经典 handle 路径。[^cuda-pool]
• CUDA VMM 分配：物理内存、虚拟地址保留、映射和访问权限分开管理，用的是 Driver API 的可共享 handle 路径。[^cuda-vmm]

上例把分配大小设成 2 MiB，也有文档依据：NVIDIA 提醒，cudaMalloc 可能从更大的底层块中划分小分配，IPC 可能暴露整个底层块，因此建议共享的分配大小按 2 MiB 对齐。这里说的是共享分配的隔离边界，不是“tensor 必须有 2 MiB”。[^cuda-ipc]

UBMEM 在昇腾栈中的位置

前面把共享机制拆开后，再看 UBMEM 就容易找到它的具体职责了。

本文的 UBMEM 指昇腾通信栈中的 UB 内存语义路径，具体分析 HCOMM 的 COMM_PROTOCOL_UB_MEM。它与 Ascend C 算子里简称 UB 的 Unified Buffer 不同；后者是片上存储，讨论的是算子局部数据驻留、搬运和 bank conflict。[^unified-buffer]

先看类型，再看名字

HCOMM 把 COMM_PROTOCOL_UB_MEM 放在 CommProtocol 中，与 UB_CTP、UB_RTP、RoCE 等协议类型并列；CommMemType 则单独区分 Device 和 Host 内存。这个类型划分说明：UBMEM 描述访问路径和通信语义，实际数据仍驻留在具体的设备或主机内存中。[^hcomm-types]

在本文核查的 HCOMM 提交 4ab526d49bac264c71eb4d435ac8c0d6418912b4 中，公开协议矩阵把 UB_MEM 列在 Ascend 950PR/950DT 的 COMM_ENGINE_AICPU_TS 和 COMM_ENGINE_AIV 下，A2/A3 矩阵未列出。这个结论限定于该版本的 HCOMM 协议支持；它不表示 A2/A3 没有设备内存 IPC。较早的 CANN IPC 文档就已列出 A2/A3 支持。[^hcomm-types][^cann-ipc]

这里也需要分清三个相邻名词：UnifiedBus 是互联体系；UBMEM 是这里研究的内存语义路径；UDMA 等搬运资源回答的是任务怎样执行。仅凭其中一个名字，无法判断一次访问究竟由哪个引擎发起、经过哪条链路或获得哪种完成保证。

沿源码找到共享对象

在这个固定版本里，AIV UBMEM 的资源建立可以沿以下调用关系阅读。它描述的是控制面和映射建立，没有省略成“建链后自动完成数据传输”。

1. 选择 Endpoint：UbMemEndpoint::Init 以 COMM_PROTOCOL_UB_MEM 建立对应内存管理资源。
2. 登记 buffer：UbMemRegedMemMgr::RegisterMemory 使用 LocalIpcRmaBuffer。普通构造路径先记录页对齐后的共享范围和偏移，再生成 IPC key。
3. 交换描述：AivUbMemTransport 通过 socket 交换 ExchangeIpcBufferDto，字段包含地址、大小、偏移、PID、共享名称和内存标签。
4. 建立对端视图：接收端构造 RemoteIpcRmaBuffer。跨 PID 分支导入 key，随后用“导入基址 + 偏移”得到可访问地址；同 PID 分支复用已有地址并执行相应预取处理。
5. 交给使用者：Channel 的 GetRemoteMems 返回远端内存描述，供后续设备侧操作使用。[^hcomm-endpoint][^hcomm-local][^hcomm-transport][^hcomm-remote]

这条路径最值得停一下看的地方，是底层确实出现了 CANN IPC。适配器中的两个实际调用分别是：

// HrtIpcSetMemoryName 中生成共享 key 的调用
aclError ret = aclrtIpcMemGetExportKey(ptr, ptrMaxLen, name, nameMaxLen, 1UL);
// HrtIpcOpenMemory 中导入 key 的调用，位于另一个函数
aclError ret = aclrtIpcMemImportByKey(&ptr, name, 0UL);

摘录自 orion_adapter_rts.cc 的第 614、644 行，省略各自函数的日志和错误处理，两个调用不是一段连续可执行程序。[^hcomm-adapter]

name 是这里交换的共享 key，ptrMaxLen 是导出范围；第二个调用通过输出参数 ptr 返回本进程可使用的地址。CANN 官方文档对 aclrtIpcMemImportByKey 的描述与这个用法相符。[^cann-ipc]

因此，与经典 CUDA IPC 更直接对应的是 CANN 的设备内存 IPC 接口；HCOMM UBMEM 则会在其具体实现里组合这些资源能力。这是依据两侧 API 与上述调用路径作出的机制比较，不代表两套 ABI、权限选项或硬件范围相同。

有 API 名称，不代表这条后端支持它

如果只看 HCOMM 接口目录，很容易顺手写出 HcommMemReg → HcommMemExport → HcommMemImport，再把它解释成 UBMEM 的标准路径。当前证据恰好不支持这种写法。

HcommMemExport 和 HcommMemImport 的文档明确说明：UB_MEM Endpoint 不支持这两个通用接口，调用不执行实际操作。相应的 UbMemRegedMemMgr 实现打印不支持信息并返回成功码；仅检查返回值，并不能证明已经拿到了有效描述。上述 AIV 路径通过 Channel 内部的 DTO 交换来完成资源建立。[^hcomm-export]

数据面也要单独核对。在该提交的 AivUbMemChannel 类中，Write、Read、NotifyRecord、NotifyWait 和 ChannelFence 方法返回 HCCL_E_NOT_SUPPORT。这个 Host 侧资源类的行为，不能外推成整个 UBMEM 没有读写能力；它说明不能把别的 Channel 后端调用方式原样套过来。实际读写、通知和完成顺序，应继续沿选定设备接口或算子实现核查。[^hcomm-channel]

这也是 UBMEM 与 CUDA IPC 比较时最需要保持的边界：映射建立的共同点已经有证据，后续执行和同步的等价性还需要逐条证明。

两者可以比较到哪里

如果问题是“能否让另一进程访问已经存在的数据”，两者确实有共同机制。若问题变成“谁更快”“谁能跨多少台服务器”，就必须补齐设备、互联域、分配方式和操作类型。

比较维度	经典 CUDA IPC	本文核查的 HCOMM UBMEM 路径
直接对象	cudaIpc* 内存与事件 API	COMM_PROTOCOL_UB_MEM Endpoint / Channel 资源
跨进程资源描述	cudaIpcMemHandle_t	Channel 交换的 IPC key 及 buffer 元数据
地址如何得到	cudaIpcOpenMemHandle 返回本地设备指针	跨 PID 分支由 CANN IPC 导入，再加 buffer 偏移
物理数据是否自动复制	导入建立映射，不自动生成私有副本	所示路径建立共享视图；后续是否复制取决于执行代码
数据何时可读	应用建立 event / stream 等同步关系	必须核查选定设备数据面与同步接口
范围怎样判断	本文示例为同机 Linux；跨 GPU 检查 peer access	结合 950 支持矩阵、CANN/驱动、拓扑和实际可访问域

表中 CUDA IPC 依据 NVIDIA API 文档，UBMEM 依据上述固定源码；最后一行刻意没有把“支持 UB”写成“任意两台服务器可以共享”。这是两类不同的证据强度。

跨节点需要更换比较对象

把经典 CUDA IPC 的 handle 通过 TCP 发到另一台普通服务器，不会凭空建立 GPU 到 GPU 的可访问路径。传递资源描述的通道，并不决定资源可访问的范围。

NVIDIA 的多节点 NVLink 场景使用 CUDA VMM 的 CU_MEM_HANDLE_TYPE_FABRIC 等机制，并要求受支持的平台与 IMEX 环境。VMM 把物理分配创建、共享 handle 导入导出、地址保留、映射和权限设置分开；典型接口包括 cuMemCreate、cuMemExportToShareableHandle、cuMemImportFromShareableHandle、cuMemAddressReserve、cuMemMap、cuMemSetAccess。[^cuda-vmm][^cuda-vmm-api]

所以，如果要讨论 UB 超节点与 NVLink 多节点系统，比较对象应该扩展为完整的 fabric 内存共享能力，同时检查寻址域、权限、故障和同步。只比较 UBMEM 与 cudaIpcGetMemHandle 两个名字，会把产品能力和单个 API 混在一起。

零拷贝需要说明省掉了什么

共享映射可以省掉一份显式的中间副本，例如避免“先复制到 host，再传递，再复制回设备”的路径。但 B 在另一张卡上读 A 的数据，数据仍要经过互联；导入映射也没有把远端 HBM 变成本地 HBM。

对于前面的数组，可以选择：

• 直接读共享映射：保留 A 的一份数据，B 按需访问；代价可能落在远端读取、访问粒度和重复访问上。
• 先复制到 B 再计算：多一份 B 的本地存储，支付一次显式传输；后续重用可以在本地完成。

这是由数据位置推导出的取舍，没有同条件测量，不能宣布其中一种一定更快。性能比较至少应分别记录映射建立耗时、稳定状态传输或访问耗时、同步耗时，以及每个 buffer 的复用次数。

实践时先验证什么

我更倾向于先把一个 buffer 的完整生命周期跑通，再讨论“换后端能省多少通信”。原因来自上面的机制：地址拿到了，却没等生产完成，或者读取尚未结束就释放，这些问题都发生在带宽成为瓶颈之前。

一个有明确检查结果的最小实验可以这样安排：

1. 确认分配身份：记录原始分配基址、大小、逻辑数据偏移、设备标识和分配 API。先用独立 buffer，避免框架内存池和 view 干扰判断。
2. 确认共享视图：两端打印各自的本地地址，允许它们不同；消费者实际校验四个整数，不能用“导入返回成功”代替数据检查。
3. 确认时序：让消费者在生产完成事件之后读取，并让所有者等待消费者结束再复用或释放。多轮实验加入轮次编号，检查读到的是本轮数据。
4. 确认释放边界：消费者完成设备工作后关闭映射，所有者收到关闭确认后释放。超时或对端异常时，不把超时本身当作设备已经停止访问的证明。
5. 最后比较成本：增加 buffer 大小和复用次数，分别测首次建立、复用访问及显式复制路径，保持设备、拓扑和消费算子一致。

证据边界

本文完成了官方文档核查和 HCOMM 固定提交的静态调用追踪，没有在 CUDA GPU 或 Ascend 950 上执行上述实验，因此不提供带宽、时延或硬件互通的实测结论。CUDA 平台支持的当前文档还存在口径差异：Programming Guide 写 Linux，Runtime API Reference 同时提及 Windows 兼容支持；本文只以 Linux 为示例范围，其他平台应核查目标工具链的具体说明。[^cuda-ipc][^cuda-device]

回到开头：B 要读到 A 的第三个整数，需要的不只是一个地址数值，而是正确的共享对象、对 B 有效的映射、生产完成的依赖，以及覆盖整个访问过程的内存生命周期。CUDA IPC 把其中一些步骤暴露为直接 API；HCOMM UBMEM 在所研究的路径中组合了 CANN IPC 与 Channel 资源。沿着这几件事读文档和源码，比从“都是共享内存，所以应该一样”出发，更容易找到真正可迁移的部分。

参考资料

核查日期为 2026 年 9 月 10 日。HCOMM 所有源码与仓内文档链接固定到提交 4ab526d49bac264c71eb4d435ac8c0d6418912b4；NVIDIA 在线文档按访问时内容核查，示例为本文编写。有关昇腾通信整体分层，可继续阅读本站 Ascend Communication Stack。

[^cuda-ipc]: NVIDIA，CUDA Programming Guide：Interprocess Communication，页面更新日期 2026-09-09；参见进程局部指针、经典 IPC、分配边界与多节点 handle 的区分。
[^cuda-uva]: NVIDIA，CUDA Programming Guide：Unified and System Memory，参见 Unified Virtual Address Space 的单进程边界。
[^cuda-device]: NVIDIA，CUDA Runtime API：Device Management，参见 cudaIpcGetMemHandle、cudaIpcOpenMemHandle、cudaIpcCloseMemHandle、cudaIpcGetEventHandle。
[^cuda-stream]: NVIDIA，CUDA Runtime API：Stream Management，参见 cudaStreamWaitEvent、cudaStreamSynchronize。
[^cuda-event]: NVIDIA，CUDA Runtime API：Event Management，参见 cudaEventRecord 对首次记录、重复记录与等待捕获状态的说明。
[^cuda-pool]: NVIDIA，Using the NVIDIA CUDA Stream-Ordered Memory Allocator, Part 2，参见跨进程内存池与 allocation 共享。
[^cuda-vmm]: NVIDIA，CUDA Programming Guide：Virtual Memory Management，参见 VMM、fabric handle、平台能力查询和 IMEX。
[^cuda-vmm-api]: NVIDIA，CUDA Driver API：Virtual Memory Management，参见创建、导入导出、映射与权限 API。
[^unified-buffer]: 昇腾社区，Avoiding Bank Conflicts in the Unified Buffer，用于区分片上 Unified Buffer 与通信语境 UBMEM。
[^cann-ipc]: CANN 8.3.RC1，aclrtIpcMemImportByKey。此来源支持 IPC API 含义和该历史版本的产品矩阵，不用于证明当前 HCOMM 的完整硬件范围。
[^hcomm-types]: HCOMM，CommProtocol 与 CommMemType。
[^hcomm-endpoint]: HCOMM，ub_mem_endpoint.cc 与 ub_mem_reged_mem_mgr.cc。
[^hcomm-local]: HCOMM，local_ipc_rma_buffer.cc 与 exchange_ipc_buffer_dto.h。
[^hcomm-transport]: HCOMM，aiv_ub_mem_transport.cc，重点为 BufferPack、SendMemInfo、RmtBufferUnpackProc、GetRemoteMems。
[^hcomm-remote]: HCOMM，remote_rma_buffer.cc，重点为 RemoteIpcRmaBuffer 构造与关闭路径。
[^hcomm-adapter]: HCOMM，orion_adapter_rts.cc，导出与导入适配见第 612–655 行。
[^hcomm-export]: HCOMM，HcommMemExport、HcommMemImport；空操作实现见 ub_mem_reged_mem_mgr.cc 第 53–66 行。
[^hcomm-channel]: HCOMM，aiv_ub_mem_channel.cc，区分资源管理方法与尚不支持的数据面方法。