CANN Runtime 实践：基于 aclrtMemcpyAsync 的 Device 到 Host 异步内存复制详解

CANN Runtime 实践基于 aclrtMemcpyAsync 的 Device 到 Host 异步内存复制详解【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章围绕 CANN Runtime 开源仓库中的4_d2h_async_memory_copy样例完整讲解如何通过aclrtMemcpyAsync接口实现 Device 到 HostD2H的异步内存复制。你将掌握从aclInit初始化、Device/Stream 管理、Host/Device 内存申请到异步拷贝与流同步、资源回收的完整调用链并理解异步复制与同步复制的本质区别、aclrtMemcpyKind枚举语义以及样例工程run.shCMakeLists.txt的编译运行机制。读完本文你可以直接在本仓库的样例基础上独立实现自己的 D2H 异步数据传输场景。样例概览与适用产品4_d2h_async_memory_copy位于 example/1_basic_features/memory/4_d2h_async_memory_copy属于仓库中基础特性 - 内存管理的一组样例之一同目录下还包含 H2H、H2D、D2D 等同步/异步复制样例可对照阅读 example/1_basic_features/memory/README.md。根据原文档该样例支持以下产品产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√运行前请确认当前硬件型号在上述列表内否则可能出现接口或行为差异。核心接口与调用流程原文档将样例涉及的关键功能点归纳为四类接口结合 main.cpp 源码完整调用序列如下初始化aclInit(nullptr)完成运行时初始化配置程序末尾aclFinalize()去初始化。Device 管理aclrtSetDevice(deviceId)指定运算 DeviceaclrtResetDeviceForce(deviceId)强制复位 Device 并回收资源。Stream 管理aclrtCreateStream(stream)创建 StreamaclrtSynchronizeStream(stream)阻塞等待 Stream 上所有任务完成aclrtDestroyStreamForce(stream)强制销毁 Stream。内存管理aclrtMallocHost申请 Host 端页锁定内存aclrtMalloc申请 Device 端内存aclrtFreeHost/aclrtFree分别释放。数据传输aclrtMemcpyAsync(hostPtrA, size, devPtrB, size, ACL_MEMCPY_DEVICE_TO_HOST, stream)发起 D2H 异步复制。整体数据流为在 Device 内存devPtrB中写入值123通过内核函数完成→ 异步复制到 Host 内存hostPtrA→ 同步 Stream 后读取hostPtrA中的值并打印验证。关键接口原型与参数说明aclrtMemcpyAsync的声明位于 include/external/acl/acl_rt.haclError aclrtMemcpyAsync( void* dst, // 目的地址指针 size_t destMax, // 目的地址内存的最大长度字节 const void* src, // 源地址指针 size_t count, // 待复制字节数 aclrtMemcpyKind kind, // 复制类型 aclrtStream stream); // 异步任务所在的 Stream结合头文件注释acl_rt.h与本样例调用方式需要注意异步语义接口将复制任务下发到指定 Stream 后立即返回调用方必须调用aclrtSynchronizeStream确保复制任务真正完成后再读取目的内存。原文档与头文件注释均明确强调这一点。kind取值aclrtMemcpyKind枚举定义在 acl_rt.h本样例使用ACL_MEMCPY_DEVICE_TO_HOSTDevice 到 Host。同一枚举还包含ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_DEVICE等样例目录下的1_h2d_async_memory_copy、5_d2d_async_memory_copy等即为其他方向的使用示例。对齐约束头文件注释指出片上 Device-to-Device 复制要求源、目的地址 64 字节对齐本样例为 D2H 场景但保持良好的对齐习惯仍然必要。Host 内存要求aclrtMallocHost申请的是可用于异步复制的 Host 内存接口注释明确说明该内存不能直接在 Device 中使用需要显式复制到 Device且必须通过aclrtFreeHost释放acl_rt.h。这正是 D2H 异步复制对 Host 侧缓冲区的标准要求。样例源码逐段解读main.cpp 核心逻辑约 45 行业务代码如下int32_t main() { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream stream nullptr; aclrtCreateStream(stream); // Allocate memory on the host and device uint64_t size 1 * 1024 * 1024; // 1 MB int* hostPtrA; int* devPtrB; CHECK_ERROR(aclrtMallocHost((void**)hostPtrA, size)); CHECK_ERROR(aclrtMalloc((void**)devPtrB, size, ACL_MEM_MALLOC_HUGE_FIRST)); // Write the data 123 to the virtual address devPtrB constexpr uint32_t blockDim 1; int writeValue 123; WriteDo(blockDim, stream, devPtrB, writeValue); // Copy memory from devPtrB to hostPtrA asynchronously CHECK_ERROR(aclrtMemcpyAsync(hostPtrA, size, devPtrB, size, ACL_MEMCPY_DEVICE_TO_HOST, stream)); // Read the value at address hostPtrA after stream synchronization aclrtSynchronizeStream(stream); int readValue *hostPtrA; INFO_LOG(Destination data: %d, readValue); // Release resources aclrtDestroyStreamForce(stream); aclrtFreeHost(hostPtrA); aclrtFree(devPtrB); aclrtResetDeviceForce(deviceId); aclFinalize(); return 0; }几个值得注意的实现细节Device 端写入不是简单赋值由于devPtrB是 Device 内存Host 无法直接访问样例通过WriteDo启动一个 AICore 内核把值123写入devPtrB。该内核定义在 example/kernel_func/write_read_value.cppextern C __global__ __aicore__ void DeviceWrite(__gm__ int* devPtr, int value) { int32_t idx block_idx; devPtr[idx] value; AscendC::printf(Source data: %d\n, value); }WriteDo通过DeviceWriteblockDim, nullptr, stream(devPtr, value)以单 block 在指定 Stream 上启动内核write_read_value.cpp。这说明 D2H 复制之前数据源必须已经由 Device 侧任务在 Stream 上按序产生——这也再次印证了同一 Stream 内任务按提交顺序执行的语义。错误处理宏每个 ACL 接口调用都用 example/utils.h 中定义的CHECK_ERROR宏包裹任一调用返回非ACL_SUCCESS即打印错误并退出#define CHECK_ERROR(call) \ do { \ aclError __ret (call); \ if (__ret ! ACL_SUCCESS) { \ ERROR_LOG(Operation failed: %s returned error code %d, #call, static_castint32_t(__ret)); \ return -1; \ } \ } while (0)资源回收顺序先销毁 Stream再依次释放 Host/Device 内存最后复位 Device 并调用aclFinalize与初始化的顺序严格对称。为什么这里必须先写 Device、再异步拷回、最后同步读取从源码可以看出样例刻意构造了一条完整的时间线内核写123到devPtrB→aclrtMemcpyAsync把数据拷到hostPtrA→aclrtSynchronizeStream等待整个 Stream 排空 → 读取hostPtrA校验。这演示了异步编程的核心纪律在异步 Stream 上任务的完成是最终一致的Host 侧必须在同步点之后才能安全消费复制结果。若省略aclrtSynchronizeStreamhostPtrA中的数据可能是未初始化的旧值。更完整的同步语义可参考 docs/zh/api_ref/06_stream_management.md 与 docs/zh/api_ref/11-03_memory_copy_and_set.md。编译与运行步骤一进入样例目录cd ${git_clone_path}/example/1_basic_features/memory/4_d2h_async_memory_copy其中${git_clone_path}为本仓库克隆到本地的根目录。步骤二设置环境变量# ${install_root} 替换为 CANN 安装根目录默认安装在 /usr/local/Ascend source ${install_root}/cann/set_env.sh # 自动识别 SOC_VERSION 和 ASCENDC_CMAKE_DIR source ${git_clone_path}/example/set_sample_env.shset_sample_env.shexample/set_sample_env.sh会完成三件事通过resolve_cann_env.sh解析 CANN 安装路径并导出ASCEND_INSTALL_PATH/ASCEND_HOME_PATH编译并运行 example/tools/get_soc_version/get_soc_version.cpp 小工具借助aclrtGetSocName自动探测当前硬件的SOC_VERSION在 CANN 包内按宿主架构x86_64 / aarch64自动定位ascendc.cmake所在目录并导出ASCENDC_CMAKE_DIR。脚本对 include/lib 布局有候选回退机制兼容多种 CANN 包目录结构。步骤三运行样例bash run.shrun.sh 的行为可以拆解为三个阶段环境自检校验ASCEND_HOME_PATH是否已设置即是否已sourceCANN 的set_env.sh随后加载${ASCEND_CANN_PATH}/bin/setenv.bash并打印当前编译用的SOC_VERSION。构建执行cmake -B build -DASCEND_CANN_PACKAGE_PATH${_ASCEND_CANN_PATH}配置、cmake --build build -j编译、cmake --install build安装随后运行./build/main并把输出通过tee写入output_msg.txt。自动校验用awk从输出中分别提取Source data:与Destination data:的值并比较一致则打印[SUCCESS] Memory copy successfully.不一致则打印[FAILURE]并以非零码退出。因此run.sh本身就是一次端到端的数据一致性断言。构建配置说明CMakeLists.txt 展示了样例的构建要点通过include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)引入 AscendC 内核构建工具链并用ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp)把 DeviceWrite 内核编译为静态库kernels头文件搜索路径包含${ASCEND_CANN_PACKAGE_PATH}/include与样例公共目录${CMAKE_CURRENT_SOURCE_DIR}/../../..即 example 根目录用于引入utils.h和kernel_func/kernel_ops.h最终add_executable(main main.cpp)并target_link_libraries(main PRIVATE ascendcl kernels)链接 ACL 运行时库与内核库。若想手动编译而不依赖run.sh在完成环境变量设置后执行cmake -B build -DASCEND_CANN_PACKAGE_PATH${ASCEND_INSTALL_PATH} cmake --build build -j ./build/main示例输出与结果解读原文档给出了样例的标准输出形态[INFO] Allocate memory on the host memory 0x... successfully [INFO] Allocate memory on the device memory 0x... successfully [INFO] Write the data 123 to the virtual memory 0x... [INFO] Copy memory from memory 0x... to memory 0x... [INFO] Destination data: 123 Source data: 123对照源码可以确定各条日志的来源前两行来自aclrtMallocHost/aclrtMalloc成功后的INFO_LOGmain.cpp第三行是内核下发写数据任务后的日志main.cpp第四行是aclrtMemcpyAsync调用成功后的日志main.cppSource data: 123由 Device 内核内部的AscendC::printf打印write_read_value.cppDestination data: 123由 Host 端同步后读取hostPtrA打印main.cpp。最终Destination data: 123与Source data: 123一致即证明 D2H 异步复制链路完整、数据无失真。常见问题与调优提示读到旧值 / 随机值多半是跳过了aclrtSynchronizeStream就读取hostPtrA。异步接口只保证任务被提交到 Stream不保证立即完成。run.sh报ASCEND_HOME_PATH is not set说明没有先执行source ${install_root}/cann/set_env.sh或ASCEND_HOME_PATH未被导出到当前 shell注意用source而非直接执行。复制大块数据count与destMax建议保持一致destMax不应小于实际复制长度Host 端使用aclrtMallocHost申请的页锁定内存更利于异步 DMA 传输。多方向复制对照若需实现 Host 到 Device 或 Device 到 Device 的异步复制仅需把kind换成ACL_MEMCPY_HOST_TO_DEVICE/ACL_MEMCPY_DEVICE_TO_DEVICE并参照同目录下 1_h2d_async_memory_copy、6_d2d_async_memory_copy 等样例。延伸阅读同组内存复制样例example/1_basic_features/memoryH2H/H2D/D2D、同步/异步、多 Stream 同步复制等接口参考acl_rt.haclrtMemcpyAsync、aclrtMemcpyKind、aclrtMallocHost等声明与注释文档 docs/zh/api_ref/11-03_memory_copy_and_set.md、docs/zh/api_ref/06_stream_management.md内核定义example/kernel_func/write_read_value.cpp 与 example/kernel_func/kernel_ops.h环境探测脚本example/set_sample_env.sh 与 example/tools/get_soc_version/get_soc_version.cpp【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考