CANN Runtime 同步 H2D 内存复制实战：基于 aclrtMemcpy 的 Host 到 Device 数据传输样例解析

CANN Runtime 同步 H2D 内存复制实战：基于 aclrtMemcpy 的 Host 到 Device 数据传输样例解析

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

本技术指南以 1_h2d_sync_memory_copy 样例为核心，系统讲解在 CANN Runtime 上实现 Host 到 Device（H2D）同步内存复制的完整流程：从aclInit初始化、Device/Stream 管理、Host/Device 内存申请，到使用aclrtMemcpy完成数据传输，再到用 AscendC 内核回读校验结果。读完本文，你将掌握 H2D 同步复制的 API 调用链、各接口参数语义、构建运行环境配置，以及一套可复用的源码级校验方法。

样例定位：最小化 H2D 同步复制范例

1_h2d_sync_memory_copy位于 example/1_basic_features/memory/ 目录下，是 CANN Runtime 基础功能（basic features）系列中"内存复制"主题的第一个样例，与同目录下的1_h2d_async_memory_copy（异步 H2D 复制）、2_h2d_async_memory_copy、3_d2h_sync_memory_copy等样例共同构成完整的数据搬运场景矩阵。

该样例的业务逻辑非常聚焦：在 Host 侧申请一块内存写入数值123，在 Device 侧申请一块等大小内存，通过同步接口aclrtMemcpy将数据一次性拷贝到 Device，再通过一个 Device 内核将目标地址的值读取出来并打印，以此证明数据真正到达了 Device 内存。

样例目录仅包含四个文件，职责清晰：

文件	作用
main.cpp	主程序：初始化 → 申请内存 → 写数据 → 同步复制 → 内核回读 → 资源释放
CMakeLists.txt	构建脚本：编译主程序并链接ascendcl与 AscendC 内核静态库
run.sh	一键脚本：配置环境、构建、运行，并自动比对源/目的数据校验结果
README / README_en	样例说明（中/英文）

产品支持与运行前置条件

根据 样例文档 的产品支持矩阵，该样例支持以下硬件：

产品	是否支持
Ascend 950PR / Ascend 950DT	是
Atlas A3 训练系列产品 / Atlas A3 推理系列产品	是
Atlas A2 训练系列产品 / Atlas A2 推理系列产品	是

运行前需要满足的环境条件：

• 已安装 CANN 软件包（默认安装根目录为/usr/local/Ascend），并已下载本仓库样例代码；
• 环境中具备可用的昇腾 AI 处理器（Device）；
• 构建工具链：CMake（要求 3.16.0 及以上，见 CMakeLists.txt 中cmake_minimum_required(VERSION 3.16.0)）、C++ 编译器；
• 样例中会调用 AscendC 算子内核（用于 Device 侧读写），因此需要配置 AscendC 编译工具链（ascendc.cmake）。

构建与运行步骤

1. 切换到样例目录

cd ${git_clone_path}/example/1_basic_features/memory/1_h2d_sync_memory_copy

其中${git_clone_path}为克隆 CANN Runtime 仓库到本地的路径。

2. 设置环境变量

# ${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.sh

第一条命令导入 CANN 的基础运行环境。第二条命令执行 set_sample_env.sh，它内部做了三件关键事情：

• 自动探测 SOC_VERSION：通过编译运行 tools/get_soc_version/get_soc_version.cpp 这个小工具（底层调用aclrtGetSocName），自动识别当前环境的芯片型号，无需手动填写类似Ascend910_9362、Ascend910B2这样的型号字符串；
• 自动定位 ASCENDC_CMAKE_DIR：按宿主机架构（x86_64-linux / aarch64-linux）逐层探测 CANN 安装目录下的tikcpp/ascendc_kernel_cmake目录，找到其中包含ascendc.cmake的路径；
• 导出构建所需环境变量：最终导出ASCEND_INSTALL_PATH、ASCEND_HOME_PATH、SOC_VERSION、ASCENDC_CMAKE_DIR四个变量供 CMake 构建使用。

如果自动探测失败，也可以按照样例文档的说明手动设置：SOC_VERSION指定昇腾 AI 处理器型号，ASCENDC_CMAKE_DIR指定 AscendC 编译器ascendc.cmake所在路径（例如/usr/local/Ascend/cann/x86_64-linux/tikcpp/ascendc_kernel_cmake）。

3. 运行样例

bash run.sh

run.sh 内部执行了完整的"构建—运行—校验"流水线：

1. 校验ASCEND_HOME_PATH是否已设置（未设置则报错提示先 source CANN 的 set_env.sh），随后加载${ASCEND_CANN_PATH}/bin/setenv.bash；
2. 创建build目录，执行cmake -B build -DASCEND_CANN_PACKAGE_PATH=...配置工程，再执行cmake --build build -j编译、cmake --install build安装；
3. 运行./build/main并将标准输出同时打印到终端和output_msg.txt文件；
4. 用awk从输出文件中分别提取Source data:和Destination data:的值并比对，相等则输出[SUCCESS]，否则输出[FAILURE]并以退出码 1 结束。

构建脚本解析：如何链接 Runtime 与 AscendC 内核

CMakeLists.txt 展示了 CANN 样例的标准构建组织方式，值得注意的几点：

set(SOC_VERSION $ENV{SOC_VERSION}) set(ASCENDC_CMAKE_DIR $ENV{ASCENDC_CMAKE_DIR}) set(RUN_MODE "npu" CACHE STRING "run mode: npu") # 引入 AscendC 编译工具链，将 Device 侧内核源文件编译为静态库 include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp) # 链接 CANN 头文件、公共头文件与运行时库 include_directories(${ASCEND_CANN_PACKAGE_PATH}/include) include_directories(${CMAKE_CURRENT_SOURCE_DIR}/../../..) link_directories(${ASCEND_CANN_PACKAGE_PATH}/lib64) add_executable(main main.cpp) target_link_libraries(main PRIVATE ascendcl kernels)

• 头文件从${ASCEND_CANN_PACKAGE_PATH}/include引入，即 CANN 安装目录下的 ACL 公共头文件目录（对应本仓库的 include/external/acl/ 中的声明）；
• 可执行文件main链接ascendcl动态库（libacl_rt.so）与kernels静态库；
• ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp)将 write_read_value.cpp 中的 Device 内核函数编译为静态库，供 Host 侧通过<<<>>>语法启动。

关键 API 全景：本样例涉及的 CANN Runtime 接口

依据 样例文档 的梳理，该样例覆盖了 CANN Runtime 五大类接口：

• 初始化
  • aclInit：进行初始化配置（样例中传入nullptr表示使用默认配置）；
  • aclFinalize：去初始化，释放 ACL 全局资源。
• Device 管理
  • aclrtSetDevice：指定用于运算的 Device；
  • aclrtResetDeviceForce：强制复位当前运算的 Device，回收 Device 上的资源。
• Stream 管理
  • aclrtCreateStream：创建 Stream（任务队列）；
  • aclrtDestroyStreamForce：强制销毁 Stream，丢弃所有尚未执行的任务。
• 内存管理
  • aclrtMallocHost：在 Host 侧申请锁页内存（pinned memory），这类内存可用于与 Device 直接进行 DMA 传输；
  • aclrtMalloc：在 Device 侧申请内存；
  • aclrtFreeHost/aclrtFree：分别释放 Host 与 Device 侧内存。
• 数据传输
  • aclrtMemcpy：以内存复制方式实现 Host-to-Device 数据传输（同步语义）。

核心接口深入：aclrtMemcpy 的语义与参数

aclrtMemcpy是本样例真正的主角，其声明位于 include/external/acl/acl_rt.h：

aclError aclrtMemcpy(void* dst, size_t destMax, const void* src, size_t count, aclrtMemcpyKind kind);

头文件注释明确将其定义为"synchronous memory replication between host and device"（Host 与 Device 之间的同步内存复制）。与aclrtMemcpyAsync（异步版本，依赖 Stream 调度）不同，aclrtMemcpy是阻塞式的：函数返回时，数据搬运已经完成，无需额外同步即可安全读取目的地址。

各参数含义：

参数	含义
dst	目的地址指针（本样例为 Device 侧内存地址devPtrB）
destMax	目的地址内存的最大长度（本样例为1 * 1024 * 1024字节）
src	源地址指针（本样例为 Host 侧内存地址hostPtrA）
count	待复制的字节数
kind	内存复制类型，见aclrtMemcpyKind枚举

aclrtMemcpyKind枚举完整定义于 include/external/acl/acl_rt.h，本样例使用ACL_MEMCPY_HOST_TO_DEVICE（Host 到 Device），全部可选值如下：

typedef enum aclrtMemcpyKind { ACL_MEMCPY_HOST_TO_HOST, // Host 到 Host ACL_MEMCPY_HOST_TO_DEVICE, // Host 到 Device（本样例使用） ACL_MEMCPY_DEVICE_TO_HOST, // Device 到 Host ACL_MEMCPY_DEVICE_TO_DEVICE, // Device 到 Device ACL_MEMCPY_DEFAULT, // 默认，根据地址自动判断方向 ACL_MEMCPY_HOST_TO_BUF_TO_DEVICE, // Host 经中间缓冲到 Device ACL_MEMCPY_INNER_DEVICE_TO_DEVICE,// Device 内部（同卡）复制 ACL_MEMCPY_INTER_DEVICE_TO_DEVICE,// Device 间（跨卡）复制 } aclrtMemcpyKind;

值得指出的是ACL_MEMCPY_DEFAULT模式：系统会根据传入的源、目的地址自动判断复制方向，在编写通用封装时可以简化调用方的复杂度。

Device 内存申请策略：aclrtMalloc 与分配策略枚举

Host 侧使用aclrtMallocHost申请内存（声明见 include/external/acl/acl_rt.h），Device 侧使用aclrtMalloc申请，其第三个参数为内存分配策略。样例代码使用的是：

aclrtMalloc((void**)&devPtrB, size, ACL_MEM_MALLOC_HUGE_FIRST);

aclrtMemMallocPolicy枚举定义于 include/external/acl/acl_rt.h：

typedef enum aclrtMemMallocPolicy { ACL_MEM_MALLOC_HUGE_FIRST, // 优先申请 huge page 内存（本样例使用） ACL_MEM_MALLOC_HUGE_ONLY, // 仅申请 huge page 内存 ACL_MEM_MALLOC_NORMAL_ONLY, // 仅申请普通页内存 ACL_MEM_MALLOC_HUGE_FIRST_P2P, // 优先申请支持 P2P 的 huge page 内存 ACL_MEM_MALLOC_HUGE_ONLY_P2P, ACL_MEM_MALLOC_NORMAL_ONLY_P2P, ACL_MEM_MALLOC_HUGE1G_ONLY, // 仅申请 1G huge page 内存 ACL_MEM_MALLOC_HUGE1G_ONLY_P2P, ... } aclrtMemMallocPolicy;

ACL_MEM_MALLOC_HUGE_FIRST表示优先尝试分配大页（huge page）内存，大页内存能减少 TLB 缺失、提升大块数据搬运时的性能；若大页不足则回退到普通页内存，是一种"尽力而为"的稳妥策略，适合本样例这种单次 1MB 的搬运场景。

源码级流程解析：main.cpp 的完整调用链

main.cpp 完整展示了 H2D 同步复制的标准生命周期，我们逐段剖析：

阶段一：初始化与运行环境准备

aclInit(nullptr); // 初始化 ACL，nullptr 表示默认配置 int32_t deviceId = 0; aclrtSetDevice(deviceId); // 指定使用 0 号 Device aclrtStream stream = nullptr; aclrtCreateStream(&stream); // 创建 Stream，供后续内核下发使用

阶段二：Host 与 Device 内存申请

uint64_t size = 1 * 1024 * 1024; // 复制 1MB 数据 int* hostPtrA; int* devPtrB; CHECK_ERROR(aclrtMallocHost((void**)&hostPtrA, size)); // Host 锁页内存 INFO_LOG("Allocate memory on the host memory %p successfully", hostPtrA); CHECK_ERROR(aclrtMalloc((void**)&devPtrB, size, ACL_MEM_MALLOC_HUGE_FIRST)); // Device 内存 INFO_LOG("Allocate memory on the device memory %p successfully", devPtrB);

这里引入的CHECK_ERROR宏定义在 example/utils.h：它将每个 ACL 接口的返回值与ACL_SUCCESS比对，失败时打印具体接口名与错误码并提前返回，是 CANN 样例中最通用的错误处理范式。

阶段三：写源数据并执行同步复制

int writeValue = 123; *hostPtrA = writeValue; // Host 内存可直接按普通指针写入 INFO_LOG("Write the data %d to the virtual memory %p", writeValue, hostPtrA); INFO_LOG("Source data: %d", writeValue); // 同步执行 Host -> Device 复制：将 hostPtrA 指向的 size 字节拷贝到 devPtrB CHECK_ERROR(aclrtMemcpy(devPtrB, size, hostPtrA, size, ACL_MEMCPY_HOST_TO_DEVICE)); INFO_LOG("Copy memory from memory %p to memory %p", hostPtrA, devPtrB);

aclrtMallocHost申请的是 Host 侧锁页内存，因此可以像普通内存一样通过指针直接写入数据；aclrtMemcpy返回即表示 1MB 数据已全部落到 Device 内存devPtrB中。

阶段四：Device 内核回读验证

constexpr uint32_t blockDim = 1; ReadDo(blockDim, stream, devPtrB); // 在 Device 上启动内核读取 devPtrB 的值

同步复制完成后，代码通过ReadDo启动一个 AscendC 内核在 Device 侧读取目标地址。内核实现位于 example/kernel_func/write_read_value.cpp：

extern "C" __global__ __aicore__ void DeviceRead(__gm__ int* devPtr) { int32_t idx = block_idx; int value = devPtr[idx]; AscendC::printf("Destination data: %d\n", value); }

DeviceRead以block_idx（block 索引）作为地址偏移，读取devPtr处的值并通过AscendC::printf打印到 Host 侧日志。由于aclrtMemcpy是同步接口，数据在ReadDo被调用前已就绪，因此能直接读到123。该内核的 Host 侧启动封装ReadDo声明在 example/kernel_func/kernel_ops.h，与样例中的WriteDo等内核统一管理。

阶段五：资源释放

aclrtDestroyStreamForce(stream); // 强制销毁 Stream，丢弃所有任务 aclrtFreeHost(hostPtrA); // 释放 Host 内存 aclrtFree(devPtrB); // 释放 Device 内存 aclrtResetDeviceForce(deviceId); // 强制复位 Device，回收资源 aclFinalize(); // ACL 去初始化

释放顺序与申请顺序严格对称：先销毁 Stream，再释放两块内存，复位 Device，最后aclFinalize。aclrtDestroyStreamForce与aclrtResetDeviceForce都是"强制"语义的接口，会忽略未完成任务直接回收资源，适合样例这类"跑完即清场"的场景。

结果校验机制：脚本级自动比对

run.sh 的末尾实现了一个轻量级的自动化断言：

file_path=output_msg.txt ./build/main | tee "${file_path}" source_value=$(awk -F':' '/Source data:/ {gsub(/^ +| +$/, "", $2); print $2; exit}' "${file_path}") destination_value=$(awk -F':' '/Destination data:/ {gsub(/^ +| +$/, "", $2); print $2; exit}' "${file_path}") if [[ -n "${source_value}" && "${source_value}" = "${destination_value}" ]]; then echo "[SUCCESS] Memory copy successfully. Values at source and destination are equal: ${source_value}" else echo "[FAILURE] Memory copy failed. ..." exit 1 fi

运行输出被同时tee到output_msg.txt，脚本用awk提取Source data:（Host 侧写入值）与Destination data:（Device 内核回读值）并比较。两者相等即证明同步复制链路（写入 → 复制 → 内核回读）整体正确，这是"端到端数据正确性"的闭环验证，比单纯依赖 API 返回码更严谨。

预期输出

运行成功后，终端输出如下：

[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] Source data: 123 [INFO] Copy memory from memory 0x... to memory 0x... Destination data: 123

同时脚本会打印[SUCCESS] Memory copy successfully. Values at source and destination are equal: 123。其中0x...为实际运行时分配的内存地址，每次运行可能不同；123为源数据值。

与其他内存复制样例的对照

理解 H2D 同步复制后，建议对照同目录下其他样例以建立完整认知：

• 同步 vs 异步：本样例使用aclrtMemcpy（同步，返回即完成）；对照 1_h2d_async_memory_copy 使用aclrtMemcpyAsync（异步，依赖 Stream 上的同步点）；
• 方向扩展：D2H 场景见 3_d2h_sync_memory_copy；
• 进阶机制：若关注更细粒度的搬运控制，可进一步阅读 13_memcpy_descriptor（memcpy 描述符）与 9_multistream_sync_memory（多 Stream 同步内存）。

更系统的数据搬运、内存管理原理可参阅仓库文档 docs/zh/dev_guide/02_memory_management.md 与 docs/zh/dev_guide/02-01_data_copy.md，API 完整参考见 docs/zh/api_ref/11-03_memory_copy_and_set.md。

小结

本样例以最小可运行的形式展示了 CANN Runtime 上 H2D 同步内存复制的完整闭环：环境准备与构建配置 → 生命周期管理（aclInit/aclFinalize、Device、Stream）→ 双端内存申请（aclrtMallocHost/aclrtMalloc）→ 同步数据搬运（aclrtMemcpy+ACL_MEMCPY_HOST_TO_DEVICE）→ Device 内核回读验证 → 对称式资源释放。其中aclrtMemcpy的同步语义、aclrtMemcpyKind方向枚举、aclrtMemMallocPolicy分配策略，以及 run.sh 的自动化断言机制，都是后续编写任何昇腾 AI 数据搬运代码可直接复用的基础能力。

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