CANN Runtime 跨进程物理内存共享实战：基于进程白名单校验的 Device 内存导出与导入

CANN Runtime 跨进程物理内存共享实战：基于进程白名单校验的 Device 内存导出与导入

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

导读

本篇文章以 CANN Runtime 开源仓库中的7_physical_memory_sharing_withpid样例为骨架，讲解在同一个 Device 上、两个独立进程之间共享物理内存的完整方案。与普通 IPC 内存共享不同，本样例在导出共享句柄后启用进程白名单校验：进程 A 申请物理内存并导出可共享句柄，进程 B 将自身进程 ID 交给进程 A，由进程 A 调用aclrtMemSetPidToShareableHandle将进程 B 加入白名单后，进程 B 才能导入句柄、映射虚拟地址并读取设备侧数据。读完本文，你将掌握 CANN Runtime 虚拟内存管理（VMM）体系中aclrtMallocPhysical、aclrtReserveMemAddress、aclrtMapMem、aclrtMemExportToShareableHandle、aclrtMemImportFromShareableHandle等一系列接口的完整调用链，以及跨进程临时文件同步的工程实现。

一、样例背景：为什么需要跨进程共享物理内存

在多进程分布式训练、多实例推理等场景中，多个进程往往需要访问同一块 Device 物理内存。CANN Runtime 提供了基于"物理内存句柄（DrvMemHandle）+ 虚拟地址预留/映射"的内存模型，使得一块物理内存可以被导出为共享句柄，再被其他进程导入使用。

本样例的特殊之处在于引入了进程白名单安全机制：

• 进程 A 持有物理内存，调用aclrtMemExportToShareableHandle导出共享句柄；
• 进程 B 通过aclrtDeviceGetBareTgid获取自身在物理设备上登记的进程 ID，并把它交给进程 A；
• 进程 A 调用aclrtMemSetPidToShareableHandle将进程 B 加入白名单；
• 只有白名单内的进程才能通过aclrtMemImportFromShareableHandle导入该句柄并使用内存。

从源码注释可以看到，aclrtMemExportToShareableHandle支持通过 flags 控制白名单校验行为：ACL_RT_VMM_EXPORT_FLAG_DEFAULT为默认行为（启用校验），ACL_RT_VMM_EXPORT_FLAG_DISABLE_PID_VALIDATION则移除 PID 白名单校验（见 acl_rt.h）。本文样例采用默认校验行为，即更安全的共享方式。

二、样例总览与运行环境

2.1 产品支持情况

2.2 样例目录结构

样例位于 example/1_basic_features/memory/7_physical_memory_sharing_withpid：

7_physical_memory_sharing_withpid/ ├── CMakeLists.txt # 构建配置：编译内核库与 proc_a/proc_b 两个可执行文件 ├── proc_a.cpp # 进程 A：分配物理内存、导出句柄、设置白名单 ├── proc_b.cpp # 进程 B：上报 PID、导入句柄、映射并读取数据 ├── run.sh # 一键编译、启动双进程并自动校验共享结果 ├── README.md / README_en.md

样例复用仓库公共工具与内核代码：

• file_ops.cpp 与 mem_utils.h：跨进程临时文件读写工具；
• write_read_value.cpp：AscendC 设备侧内核，用于把数值写入共享内存；
• set_sample_env.sh：自动识别SOC_VERSION与ASCENDC_CMAKE_DIR的公共环境配置脚本。

run.sh会同时启动proc_a与proc_b两个进程，二者通过file/目录下的临时文件交换进程 ID、共享句柄和完成标志。

三、编译与运行步骤

1. 将样例代码下载到已安装 CANN 软件的环境，并切换到样例目录：

cd ${git_clone_path}/example/1_basic_features/memory/7_physical_memory_sharing_withpid

2. 设置环境变量：

# 将 ${install_root} 替换为 CANN 安装根目录，默认安装在 /usr/local/Ascend source ${install_root}/cann/set_env.sh # 设置 SOC_VERSION 与 ASCENDC_CMAKE_DIR # - SOC_VERSION：Ascend AI 处理器型号，例如 Ascend910_9362、Ascend910B2 等 # - ASCENDC_CMAKE_DIR：样例涉及调用 AscendC 算子，需配置 AscendC 编译器的 ascendc.cmake 路径， # 例如 /usr/local/Ascend/cann/x86_64-linux/tikcpp/ascendc_kernel_cmake source ${git_clone_path}/example/set_sample_env.sh

其中 set_sample_env.sh 会根据当前机器架构（x86_64/aarch64）自动探测 CANN 的 include/lib 目录布局与 ascendc.cmake 位置，避免手工配置出错。

3. 运行样例：

bash run.sh

run.sh的执行逻辑（见 run.sh）：

• 校验ASCEND_HOME_PATH环境变量，随后 source CANN 的bin/setenv.bash；
• 创建build/目录，执行cmake -B build -DASCEND_CANN_PACKAGE_PATH=...并编译、安装；
• 创建file/临时文件交换目录；
• 后台并发启动proc_a与proc_b，各自输出重定向到output_msg_proc_a.txt/output_msg_proc_b.txt；
• 等待两个进程结束后，从输出文件中提取Source data:与Destination data:的数值进行比较，一致则打印[SUCCESS] Memory sharing successfully.。

四、核心机制：进程白名单校验下的内存共享工作流

整个样例可以划分为三个阶段，对应两个进程的职责分工：

进程 A（导出方） 进程 B（导入方） ───────────────────── ───────────────────── 1. aclrtMallocPhysical 分配物理内存 2. aclrtReserveMemAddress 预留虚拟内存 3. aclrtMapMem 建立映射 4. aclrtMemSetAccess 设置访问权限 5. WriteDo 内核写入数据 123 6. aclrtMemExportToShareableHandle 导出句柄 7. 读取 B 的 PID ─────────► aclrtDeviceGetBareTgid 获取 PID 8. aclrtMemSetPidToShareableHandle 将 B 加入白名单 9. 写句柄到文件 ─────────► 读取句柄 aclrtMemImportFromShareableHandle 导入句柄 预留虚拟内存 + 映射 + 设置访问权限 aclrtMemcpy 拷贝到 Host 验证数据 aclrtUnmapMem + 写完成标志 10. 读取完成标志 11. 释放虚拟/物理内存

核心要点：白名单设置发生在句柄传递之前。进程 A 必须先通过aclrtMemSetPidToShareableHandle把进程 B 的 PID 登记到句柄上，再把句柄写入文件交给进程 B，进程 B 的导入操作才会成功。这一时序保证了共享句柄不会被未授权进程获取后直接使用。

五、关键 CANN Runtime API 详解

本样例覆盖了 CANN Runtime 的初始化、Device 管理、Stream 管理与内存管理四大类接口，完整列表如下：

• 初始化
  • aclInit：进行初始化配置。
  • aclFinalize：实现去初始化。
• Device 管理
  • aclrtSetDevice：指定用于运算的 Device。
  • aclrtResetDeviceForce：强制复位当前运算的 Device，回收 Device 上的资源。
• Stream 管理
  • aclrtCreateStream：创建 Stream。
  • aclrtDestroyStreamForce：强制销毁 Stream，丢弃所有任务。
• 内存管理
  • aclrtMemGetAllocationGranularity：查询内存申请粒度。
  • aclrtMallocPhysical：申请 Device 物理内存，并返回物理内存 handle。
  • aclrtReserveMemAddress：预留虚拟内存。
  • aclrtMapMem：将虚拟内存映射到物理内存。
  • aclrtMemSetAccess：设置虚拟内存访问权限。
  • aclrtMemExportToShareableHandle：导出物理内存共享句柄。
  • aclrtMemSetPidToShareableHandle：设置共享内存的进程白名单。
  • aclrtDeviceGetBareTgid：获取当前进程 ID。
  • aclrtMemImportFromShareableHandle：导入共享句柄并获取本进程可用的物理内存 handle。
  • aclrtUnmapMem：取消虚拟内存与物理内存之间的映射关系。
  • aclrtReleaseMemAddress：释放预留的虚拟内存。
  • aclrtFreePhysical：释放物理内存。
  • aclrtMallocHost：申请 Host 上的内存。
  • aclrtFreeHost：释放 Host 上的内存。
• 数据传输
  • aclrtMemcpy：通过内存复制的方式实现数据传输。

以下结合 acl_rt.h 中的接口声明，对共享相关的核心接口做进一步说明。

5.1 物理内存申请与虚拟内存管理

aclrtMallocPhysical按指定属性创建物理内存分配，返回句柄（acl_rt.h）：

aclError aclrtMallocPhysical( aclrtDrvMemHandle* handle, size_t size, const aclrtPhysicalMemProp* prop, uint64_t flags);

其中flags当前未使用、必须传 0；prop描述内存属性，样例中配置如下（见 proc_a.cpp）：

aclrtPhysicalMemProp prop = {}; prop.handleType = ACL_MEM_HANDLE_TYPE_NONE; // 物理内存句柄类型 prop.allocationType = ACL_MEM_ALLOCATION_TYPE_PINNED; // 固定分配类型 prop.location.type = ACL_MEM_LOCATION_TYPE_DEVICE; // 内存位置为 Device prop.location.id = 0; // Device ID prop.memAttr = ACL_HBM_MEM_NORMAL; // 普通 HBM 内存

aclrtReserveMemAddress预留一段虚拟地址范围（acl_rt.h）：

aclError aclrtReserveMemAddress( void** virPtr, size_t size, size_t alignment, void* expectPtr, uint64_t flags);

注意expectPtr必须传nullptr（由 Runtime 决定起始地址），flags为页类型标志。样例中直接以内存分配粒度作为预留大小。

aclrtMapMem将物理内存句柄映射到预留的虚拟地址范围（acl_rt.h）：

aclError aclrtMapMem( void* virPtr, size_t size, size_t offset, aclrtDrvMemHandle handle, uint64_t flags);

映射完成后还必须调用aclrtMemSetAccess(virPtr, size, desc, count)显式设置访问权限（ACL_RT_MEM_ACCESS_FLAGS_READWRITE），否则后续访问会失败——接口注释明确说明"此接口不授予访问权限，访问前需调用aclrtMemSetAccess"（见 acl_rt.h）。

5.2 共享句柄导出、白名单与导入

三个接口构成跨进程共享的核心三件套（acl_rt.h）：

// 导出：把本进程创建的物理内存句柄共享给其他进程 aclError aclrtMemExportToShareableHandle( aclrtDrvMemHandle handle, aclrtMemHandleType handleType, uint64_t flags, uint64_t* shareableHandle); // 设置白名单：仅白名单中的进程可使用该 shareableHandle aclError aclrtMemSetPidToShareableHandle(uint64_t shareableHandle, int32_t* pid, size_t pidNum); // 导入：从共享句柄导入，得到本进程可用的物理内存 handle aclError aclrtMemImportFromShareableHandle( uint64_t shareableHandle, int32_t deviceId, aclrtDrvMemHandle* handle);

参数要点：

• aclrtMemExportToShareableHandle的handleType为保留参数，必须传ACL_MEM_HANDLE_TYPE_NONE；flags可选ACL_RT_VMM_EXPORT_FLAG_DEFAULT（默认行为）或ACL_RT_VMM_EXPORT_FLAG_DISABLE_PID_VALIDATION（关闭 PID 白名单校验）；
• aclrtMemSetPidToShareableHandle的pid是待加入白名单的进程 ID 数组，pidNum为其个数，支持一次加入多个进程；
• aclrtMemImportFromShareableHandle的deviceId用于在指定 Device 上生成句柄。

进程 B 需要上报的 PID 通过aclrtDeviceGetBareTgid获取（acl_rt.h），返回的是"当前进程在物理设备上登记的 PID"，这正是白名单校验所依据的进程身份。

六、源码走读：进程 A（导出方）

proc_a.cpp 的完整执行流程如下。

① 初始化与对齐计算。调用aclInit、aclrtSetDevice(0)、aclrtCreateStream后，先查询内存分配粒度：

const size_t dataSize = 1024 * sizeof(float); aclrtPhysicalMemProp prop = { ... }; // 属性配置见 5.1 节 size_t granularity = 0UL; CHECK_ERROR(aclrtMemGetAllocationGranularity(&prop, ACL_RT_MEM_ALLOC_GRANULARITY_MINIMUM, &granularity)); // 按粒度向上取整对齐 size_t alignedSize = ((dataSize + granularity - 1U) / granularity) * granularity;

② 分配物理内存并建立虚拟地址映射。依次调用aclrtMallocPhysical、aclrtReserveMemAddress、aclrtMapMem、aclrtMemSetAccess，完成"物理内存句柄 + 虚拟地址 + 访问权限"的完整装配。

③ 写入数据。通过 AscendC 内核把数值 123 写入共享内存：

constexpr uint32_t blockDim = 1; int writeValue = 123; WriteDo(blockDim, stream, (int*)virPtr, writeValue);

WriteDo定义于 write_read_value.cpp，内部以DeviceWrite<<<blockDim, nullptr, stream>>>(devPtr, value)方式启动内核，将值写入__gm__全局内存并打印Source data: 123。这也是样例需要配置ASCENDC_CMAKE_DIR的原因——CMakeLists.txt通过ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp)将该内核编译进kernels静态库，供proc_a链接。

④ 导出句柄并设置白名单。这是本样例区别于普通共享的关键步骤：

uint64_t shareableHandle = 0ULL; CHECK_ERROR(aclrtMemExportToShareableHandle(handle, ACL_MEM_HANDLE_TYPE_NONE, 0, &shareableHandle)); // 读取进程 B 通过文件上报的 PID int32_t pid = 0; memory::ReadFile("file/pid.bin", "file/pid.bin.done", &pid, sizeof(pid)); // 将进程 B 加入白名单，然后才把句柄交给 B CHECK_ERROR(aclrtMemSetPidToShareableHandle(shareableHandle, &pid, 1)); memory::WriteFile("file/handle.bin", "file/handle.bin.done", &shareableHandle, sizeof(shareableHandle));

⑤ 等待完成标志并释放资源。进程 A 阻塞读取file/flag.bin的完成标志，待进程 B 确认数据读取成功后，依次aclrtUnmapMem、aclrtReleaseMemAddress、aclrtFreePhysical释放虚拟与物理内存，最后aclrtDestroyStreamForce、aclrtResetDeviceForce、aclFinalize收尾。

七、源码走读：进程 B（导入方）

proc_b.cpp 的执行流程如下。

① 上报自身 PID。先获取进程 ID 并写入临时文件，供进程 A 做白名单登记：

int32_t pid = 0; CHECK_ERROR(aclrtDeviceGetBareTgid(&pid)); memory::WriteFile("file/pid.bin", "file/pid.bin.done", &pid, sizeof(pid));

② 导入句柄。等待进程 A 写入句柄文件后读取，并导入为当前进程可用的物理内存 handle：

uint64_t shareableHandle = 0ULL; aclrtDrvMemHandle handle = nullptr; memory::ReadFile("file/handle.bin", "file/handle.bin.done", &shareableHandle, sizeof(shareableHandle)); CHECK_ERROR(aclrtMemImportFromShareableHandle(shareableHandle, deviceId, &handle));

③ 建立本地映射。进程 B 在自己的地址空间中重复"查粒度 → 预留虚拟内存 → 映射 → 设置访问权限"流程，将导入的物理内存映射到本地虚拟地址。注意：进程 B 也使用ACL_HBM_MEM_NORMAL、Device 0 等相同的aclrtPhysicalMemProp查询粒度。

④ 拷贝数据到 Host 并验证。通过aclrtMallocHost申请 Host 内存，再用aclrtMemcpy把共享内存中的设备数据拷贝到 Host：

int* hostPtrA; CHECK_ERROR(aclrtMallocHost(reinterpret_cast<void**>(&hostPtrA), granularity)); CHECK_ERROR(aclrtMemcpy(hostPtrA, granularity, virPtr, granularity, ACL_MEMCPY_DEVICE_TO_HOST)); int readValue = *hostPtrA; // 期望读到 123

⑤ 收尾。先aclrtUnmapMem解除映射，写入file/flag.bin完成标志通知进程 A，再依次aclrtReleaseMemAddress、aclrtFreePhysical、aclrtFreeHost释放全部资源。

八、跨进程文件同步机制

两个进程通过file/目录下的临时文件交换数据，其同步语义由 file_ops.cpp 实现，核心是"数据文件 + .done 完成标志文件"两阶段协议：

• WriteFile(filePath, doneFile, data, size)：先把数据写入filePath，再创建doneFile表示写入完成；
• ReadFile(filePath, doneFile, data, bufferSize)：以 1 秒间隔（waitTime = 1000000微秒）轮询doneFile是否存在，直到出现后读取数据文件，并对文件大小与缓冲区大小做越界保护。

本样例使用的三组文件及其语义：

这种实现方式将进程间握手显式化：进程 A 只有读到 PID 并完成白名单设置后才写句柄文件，进程 B 只有读到句柄文件才尝试导入，天然保证了共享时序安全。

九、构建配置与自动验证

9.1 CMake 构建配置

CMakeLists.txt 的关键配置：

# 内核编译：将设备侧写值内核编译为静态库 include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp) add_executable(proc_a proc_a.cpp ../file_ops.cpp) add_executable(proc_b proc_b.cpp ../file_ops.cpp) target_link_libraries(proc_a PRIVATE ascendcl kernels) target_link_libraries(proc_b PRIVATE ascendcl)

可见proc_a需要额外链接内核静态库kernels（因为它要启动DeviceWrite内核写入数据），而proc_b仅做aclrtMemcpy拷贝，只需链接ascendcl。两个进程共用../file_ops.cpp实现文件同步。

9.2 运行结果自动校验

run.sh在双进程结束后自动校验共享结果（见 run.sh）：从进程 A 输出中提取Source data:值、从进程 B 输出中提取Destination data:值，二者相等则打印成功信息，否则返回非零退出码并报告失败。

进程 A 典型输出：

[INFO] Process A: allocate physical memory successfully [INFO] Process A: reserve virtual memory successfully [INFO] Process A: export a shareable handle successfully, shareable handle = ... [INFO] Process A: add Process B to the whitelist successfully Source data: 123 [INFO] Process A: receive the completion signal from Process B, completion signal = 1 [INFO] Process A: release the virtual and physical memory successfully

进程 B 典型输出：

[INFO] Process B: get Process B's pid successfully [INFO] Process B: get a shareable handle successfully, shareable handle = ... [INFO] Process B: map virtual memory address to physical memory handle [INFO] Process B: copy memory from device address 0x... to host address 0x... Destination data: 123 [INFO] Process B: complete physical memory sharing [INFO] Process B: release the virtual and physical memory successfully

Source data: 123与Destination data: 123一致，证明进程 B 通过导入的共享句柄读取到了进程 A 写入设备内存的数据。

十、底层实现佐证

共享接口在 ACL 层的实现位于 src/acl/aclrt_impl/memory.cpp，可以看到每个接口都做了严格的入参校验并下沉到 RTS 层：

• aclrtMemExportToShareableHandleImpl（memory.cpp）：校验handle非空、handleType必须为ACL_MEM_HANDLE_TYPE_NONE、shareableHandle非空，再调用rtsMemExportToShareableHandle；
• aclrtMemImportFromShareableHandleImpl（memory.cpp）：校验输出handle非空后调用rtMemImportFromShareableHandle；
• aclrtMemSetPidToShareableHandleImpl（memory.cpp）：校验pid非空、pidNum为正整数，再调用rtMemSetPidToShareableHandle。

这些接口还通过ACL_PROFILING_REG接入 ACL 的 Profiling 埋点，相关的调用统计名称（如AclrtMemExportToShareableHandle、AclrtMemSetPidToShareableHandle、AclrtDeviceGetBareTgid）可在 src/acl/common/prof_reporter.h 与 profiling_manager.cpp 中查到，说明这些共享接口是受支持、可观测的标准能力。此外，src/runtime/cmake/arch5162_unsupported_acl_api.def 中对aclrtMemExportToShareableHandle、aclrtMemSetPidToShareableHandle等标注为RT_FEATURE_NOT_SUPPORT，意味着这些接口的可用性与具体芯片架构相关，跨架构移植时需以实际运行环境为准。

十一、注意事项与扩展阅读

11.1 使用要点

• 白名单设置必须先于句柄传递：进程 A 必须先在句柄上登记进程 B 的 PID，再把句柄交给进程 B，否则进程 B 的导入会因未授权而失败；
• 虚拟地址必须在各自进程内独立建立：进程 B 导入句柄后，仍需在自己的地址空间中执行"预留虚拟内存 → 映射 → 设置访问权限"三步，不能直接使用进程 A 的虚拟地址；
• 访问权限需显式设置：aclrtMapMem只建立映射关系，不授予访问权限，必须调用aclrtMemSetAccess才能访问；
• 内存申请需按粒度对齐：aclrtMallocPhysical的 size 应基于aclrtMemGetAllocationGranularity返回的粒度向上对齐；
• 句柄类型参数为保留参数：aclrtMemExportToShareableHandle的handleType必须传ACL_MEM_HANDLE_TYPE_NONE，aclrtMallocPhysical与aclrtMapMem的flags当前必须传 0。

11.2 与其他样例的对照

仓库中还提供了 8_physical_memory_sharing_withoutpid 样例，它演示的是不启用 PID 白名单校验的物理内存共享路径（对应ACL_RT_VMM_EXPORT_FLAG_DISABLE_PID_VALIDATION语义）。两相对照，可以直观理解白名单机制对跨进程共享安全性的约束：启用校验时进程必须显式互相信任并登记 PID，关闭校验时句柄一经导出即可被任意进程导入。建议根据实际部署的信任模型选择对应方案。

11.3 进程间 IPC 内存共享的通用主题

如需进一步了解 IPC 内存共享的页表对齐要求、多 Device 场景下的共享差异等话题，可参阅仓库文档 进程间IPC内存共享的页表对齐要求 与 跨DeviceP2P数据交互配置失败，以及内存管理章节 11-07_IPC_memory_sharing.md。

结语

7_physical_memory_sharing_withpid样例完整演示了 CANN Runtime 上"申请物理内存 → 导出共享句柄 → 进程白名单登记 → 跨进程导入 → 本地映射 → 数据读取 → 双向释放"的全链路，是理解 Runtime 虚拟内存管理模型与跨进程安全共享的最佳入门素材。本文结合 proc_a.cpp、proc_b.cpp、run.sh 及 memory.cpp 等源码，将样例行为与底层实现一一对应。开发者可直接基于该样例改造，例如扩展为多进程白名单（pidNum支持一次登记多个进程）、结合aclrtMemExportToShareableHandleV2使用更丰富的共享类型，或参考 8_physical_memory_sharing_withoutpid 关闭校验以适配可信内网环境。

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime