CANN Runtime Host 内存零拷贝实战：基于 aclrtHostRegisterV2 映射与 AscendC Kernel 直接读写

CANN Runtime Host 内存零拷贝实战基于 aclrtHostRegisterV2 映射与 AscendC Kernel 直接读写【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章基于 0_simple_zero_copy 示例完整讲解 CANN Runtime 中Host 内存零拷贝的实现路径通过aclrtMallocHost申请页锁定pinnedHost 内存再用aclrtHostRegisterV2aclrtHostGetDevicePointer将 Host 内存映射为 Device 可访问地址并让 AscendC Kernel 直接读写该映射地址从而绕开显式的 Device 内存搬运。读完本文你将掌握 Host 内存注册/映射/反注册的完整调用链、AscendC Kernel 二进制加载与参数装配流程以及一个可直接编译运行的向量加法零拷贝示例的每个细节。示例解决什么问题在常规 CANN 编程模型中Host 与 Device 之间的数据交换依赖aclrtMemcpy等显式拷贝Host 数据先拷入 Device 内存Kernel 计算结果再拷回 Host。这带来了两次数据搬运的开销也会让代码里到处充满aclrtMemcpyAsync调用。0_simple_zero_copy展示的是另一种数据通路把页锁定的 Host 内存注册并映射为 Device 地址。Kernel 侧拿到的三个指针两个输入、一个输出直接指向映射后的 Host 内存Kernel 计算完成后结果直接写回 Host 可读缓冲区整个流程没有任何显式 Device 内存拷贝。示例在单个受支持 Device 上执行 FP16 向量加法申请三块页锁定 Host 内存注册并映射为 Device 可访问地址将这三个映射地址作为参数传给 AscendC KernelKernel 将两个确定性 FP16 输入相加并直接写入映射后的 Host 输出缓冲区Stream 同步后应用逐一校验全部 16,384 个结果。任何 API 调用失败、结果校验失败或资源清理失败都会打印ERROR并以非零值退出。产品支持示例 README 明确声明支持以下产品产品是否支持Ascend 950PR / Ascend 950DT是Atlas A3 训练系列产品 / Atlas A3 推理系列产品是Atlas A2 训练系列产品 / Atlas A2 推理系列产品是说明以当前仓库 README_en.md 的记录为准。示例运行时set_sample_env.sh会自动探测当前环境的SOC_VERSION例如示例输出中的Ascend910B3因此在实际支持型号上均可按后续步骤直接编译运行。整体工作流程从源码 simple_zero_copy.cpp 的RunSimpleZeroCopy()可以看到整个流程被组织为四个阶段if (InitializeRuntime(session) ! 0 || PrepareMappedBuffers(session) ! 0 || LoadKernel(session, function, args) ! 0 || ExecuteAndVerify(session, function, args) ! 0) { result -1; } ReleaseSession(session, result);初始化运行时aclInit→aclrtSetDevice→aclrtCreateStream准备映射缓冲区对三块缓冲区分别执行aclrtMallocHost→aclrtHostRegisterV2→aclrtHostGetDevicePointer并填充确定性输入数据加载 Kernel 并装配参数aclrtBinaryLoadFromFile→aclrtBinaryGetFunction→aclrtKernelArgsInit→aclrtKernelArgsAppend依次追加三个映射地址→aclrtKernelArgsFinalize执行并校验aclrtLaunchKernelWithConfig→aclrtSynchronizeStream→ 在 Host 侧直接逐元素校验输出缓冲区。整个过程中 Host 输出缓冲区无需回拷校验代码直接读取session.output.host[index]这正是零拷贝的核心体现。编译与运行1. 获取并进入示例目录将示例代码下载到已安装 CANN 的环境中并切换到示例目录cd ${git_clone_path}/example/3_memory_advanced/host_register/0_simple_zero_copy2. 设置环境变量# 将 ${install_root} 替换为 CANN 安装根目录 source ${install_root}/set_env.sh # 自动探测 SOC_VERSION 与 ASCENDC_CMAKE_DIR source ${git_clone_path}/example/set_sample_env.sh其中 set_sample_env.sh 会完成三项关键工作通过 get_soc_version 辅助程序调用 Runtime ACL APIaclrtGetSocName自动探测SOC_VERSION依据探测到的 CANN 安装路径在tikcpp/ascendc_kernel_cmake等候选目录中查找ascendc.cmake导出ASCENDC_CMAKE_DIR同时导出ASCEND_INSTALL_PATH与ASCEND_HOME_PATH。导出的变量随后会被 CMakeLists.txt 消费SOC_VERSION用于内核编译目标架构ASCENDC_CMAKE_DIR用于include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)。3. 编译并执行bash run.shrun.sh 使用set -euo pipefail保证错误即停其内部流程为cmake -B build -DASCEND_CANN_PACKAGE_PATH${ASCEND_HOME_PATH} cmake --build build -j cmake --install build ./build/main | tee output_msg.txt值得注意的构建细节CMakeLists.txt 中ascendc_fatbin_library(simple_zero_copy_kernel ${KERNEL_FILES})将 add_custom.cpp 编译为 fatbin产物路径为./out/fatbin/simple_zero_copy_kernel/simple_zero_copy_kernel.o——这正是源码中kKernelPath指向的文件可执行程序main链接${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so即 Runtime 侧 ACL 库编译选项为-O2 -stdc17 -D_GLIBCXX_USE_CXX11_ABI0 -Wall -Werror运行输出通过tee同时写入output_msg.txt文件便于后续分析。关键 API 与源码实现解析示例 README 将涉及接口按功能域归类本节结合源码逐一展开。初始化、设备管理与 Stream 管理对应 simple_zero_copy.cpp 的InitializeRuntime功能接口作用初始化aclInit(nullptr)完成 ACL 运行时初始化配置设备管理aclrtSetDevice(kDeviceId)选择用于计算的 Device示例固定为 0 号设备管理aclrtResetDeviceForce(kDeviceId)强制复位当前 Device 并回收 Device 资源Stream 管理aclrtCreateStream创建 StreamKernel 在指定 Stream 上排队执行Stream 管理aclrtSynchronizeStream等待 Stream 上 Kernel 全部完成之后才能安全地在 Host 侧读取结果Stream 管理aclrtDestroyStream销毁 StreamHost 内存管理零拷贝映射三步曲这是本示例最核心的部分对应AllocateMappedBuffersimple_zero_copy.cppint AllocateMappedBuffer(MappedBuffer buffer) { CHECK_ERROR(aclrtMallocHost(reinterpret_castvoid**(buffer.host), kBufferSize)); CHECK_ERROR(aclrtHostRegisterV2(buffer.host, kBufferSize, ACL_HOST_REG_MAPPED)); buffer.registered true; CHECK_ERROR(aclrtHostGetDevicePointer(buffer.host, buffer.device, 0)); if (buffer.device nullptr) { ERROR_LOG(Mapped Device address is null.); return -1; } return 0; }三个步骤缺一不可aclrtMallocHost申请页锁定pinned的 Host 内存。页锁定内存不会被操作系统换出到磁盘是后续 DMA/映射操作的前提声明见 acl_rt.h。aclrtHostRegisterV2(ptr, size, flag)将已有的页锁定 Host 内存注册为 Device 可映射内存。flag取ACL_HOST_REG_MAPPED在 acl_rt.h 中定义为0x2UL注释为 Map host memory to device address即要求把 Host 内存映射为 Device 地址接口声明见 acl_rt.h。aclrtHostGetDevicePointer(pHost, pDevice, 0)获取已注册 Host 内存对应的 Device 映射地址。第三个参数flag当前必须传 0接口注释明确 flag for extensions (must be 0 for now)见 acl_rt.h。示例额外对返回的device指针做了空指针校验。获取到的buffer.device地址就是后续传给 Kernel 的GM_ADDR 全局内存地址。示例为输入 A、输入 B、输出各准备一块映射缓冲区MappedBuffer inputA/inputB/output并通过std::fill_n填入确定性数据std::fill_n(session.inputA.host, kElementCount, kHalfOne); // 0x3C00即 1.0 std::fill_n(session.inputB.host, kElementCount, kHalfTwo); // 0x4000即 2.0 std::fill_n(session.output.host, kElementCount, 0); // 输出先清零数据规模相关常量simple_zero_copy.cpp常量值含义kDeviceId0使用的 Device 编号kBlockDim8Kernel 启动的核数kElementCount8 * 2048 16384元素总数FP16kBufferSizekElementCount * sizeof(uint16_t) 32 KB单块缓冲区字节数kHalfOne/kHalfTwo/kHalfThree0x3C00/0x4000/0x4200分别为 1.0、2.0、3.0 的 FP16 位模式由于校验期望值为kHalfThree输入确定为 1.0 2.0 3.0。Kernel 加载与参数装配对应 simple_zero_copy.cpp 的LoadKernel接口作用aclrtBinaryLoadFromFile(kKernelPath, nullptr, session.binary)从 fatbin 文件加载 AscendC Kernel 二进制aclrtBinaryGetFunction(session.binary, add_custom, function)获取 Kernel 函数句柄函数名对应 add_custom.cpp 中的extern C __global__ __aicore__ void add_custom(...)aclrtKernelArgsInit(function, args)初始化 Kernel 参数列表aclrtKernelArgsAppend(args, devicePointer, sizeof(uintptr_t), paramHandle)依次追加参数三个映射后的 Device 地址aclrtKernelArgsFinalize(args)完成参数装配之后的参数列表可直接用于启动参数追加的顺序至关重要必须与 Kernel 形参顺序一致AppendKernelPointer依次追加inputA.device、inputB.device、output.device对应 Kernel 入口add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)中的x / y / z。Kernel 启动与结果校验对应ExecuteAndVerifysimple_zero_copy.cppCHECK_ERROR(aclrtLaunchKernelWithConfig(function, kBlockDim, session.stream, nullptr, args, nullptr)); CHECK_ERROR(aclrtSynchronizeStream(session.stream)); for (size_t index 0; index kElementCount; index) { if (session.output.host[index] ! kHalfThree) { ... return -1; } }aclrtLaunchKernelWithConfig以 8 个核kBlockDim 8在指定 Stream 上启动add_customKernelaclrtSynchronizeStream阻塞等待 Kernel 完成——只有同步完成后Host 侧直接读取映射输出缓冲区才是安全的校验循环直接在 Host 侧读取output.host全程零显式拷贝。资源释放逆序清理示例在ReleaseSession中实现了与资源申请顺序相反的完整清理simple_zero_copy.cppaclrtBinaryUnLoad卸载 Kernel 二进制对每个缓冲区依次执行aclrtHostUnregister反注册→aclrtFreeHost释放页锁定内存aclrtDestroyStream销毁 StreamaclrtResetDeviceForce强制复位 Device、回收 Device 资源aclFinalize完成去初始化。每个清理调用都通过RecordCleanupError记录错误码确保即使清理阶段失败也会以非零值上报体现示例对资源泄漏与错误传播的严格处理。AscendC Kernel 侧实现零拷贝的另一头是 Kernel 对映射地址的直接访问。Kernel 实现在 add_custom.cpp采用经典的 AscendC 流水线结构入口函数extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)三个GM_ADDR参数即 Host 侧传入的三个映射地址数据划分TOTAL_LENGTH 8 * 2048USE_CORE_NUM 8每个核负责BLOCK_LENGTH 2048个元素通过AscendC::GetBlockIdx()定位本核起始偏移双缓冲流水每核把数据再拆成TILE_NUM 8个 Tile、BUFFER_NUM 2个缓冲区双缓冲TILE_LENGTH 128Process()循环内依次执行CopyIn → Compute → CopyOutCopyIn用AscendC::DataCopy从全局内存映射地址搬入LocalTensor并EnQueComputeAscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH)完成向量加法CopyOut把结果DeQue后DataCopy写回全局内存zGm——即写回映射的 Host 输出缓冲区。值得强调的是Kernel 内SetGlobalBuffer绑定的是 Host 侧传入的映射地址因此CopyOut写回的目的地就是 Host 可读缓冲区这就是Kernel 直接写 Host 内存的本质。示例输出正常运行时的输出如下[INFO]: Current compile soc version is Ascend910B3 ... [INFO] Start to run simple_zero_copy sample. [INFO] Registered three Host buffers and obtained their Device mapping addresses. [INFO] Verified 16384 FP16 additions through mapped Host memory. [INFO] Run the simple_zero_copy sample successfully.日志宏定义于 utils.hINFO_LOG输出[INFO]前缀ERROR_LOG输出[ERROR]前缀CHECK_ERROR宏会在任意 ACL 调用返回非ACL_SUCCESS时打印出错调用与错误码并立即返回失败。因此输出中出现[ERROR]即表示相应阶段失败程序最终返回非零退出码。小结0_simple_zero_copy完整演示了 CANN Runtime 的 Host 内存零拷贝数据通路aclrtMallocHost保证页锁定、aclrtHostRegisterV2(ACL_HOST_REG_MAPPED)完成注册、aclrtHostGetDevicePointer取得映射地址随后 Kernel 直接读写映射内存并在同步后于 Host 侧校验结果。示例还顺带示范了 AscendC Kernel 二进制的加载、参数装配与启动以及严格的逆序资源清理是一份适合作为零拷贝编程起点的完整可运行样例。若需进一步了解 Host 注册相关的通用封装接口如模板化的aclrtHostRegisterV2、aclrtHostGetDevicePointer重载可参考 acl_rt_api.h。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考