尧图网站建设 
CANN Runtime 样例解析使用 aclrtMemcpy 实现 Device 内同步内存复制D2D Sync Memory Copy【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章围绕 CANN runtime 仓库中的5_d2d_sync_memory_copy样例展开完整讲解如何在昇腾 AI 处理器的 Device 内部完成申请两块 Device 内存 → 由核函数写入源地址 → 通过aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_DEVICE方式同步复制 → 由核函数读取目的地址校验的完整链路。读完本文你将掌握aclrtMemcpy同步 D2D 拷贝的接口语义与参数约束、样例的编译运行流程、底层核函数DeviceWrite/DeviceRead如何配合验证拷贝结果以及run.sh如何自动化判定样例成功与否。1. 样例概述Device 内同步内存复制解决什么问题在 AI 推理与训练程序中数据常常先在 Device 上以aclrtMalloc申请经算子或核函数处理后需要在 Device 内存的另一个地址上复用例如特征图搬运、中间结果暂存、多 Buffer 轮转。此时并不需要经过 Host 中转直接调用内存复制接口即可完成 Device 内部的数据搬移。同步语义aclrtMemcpy是同步接口调用返回时内存复制任务已经完成Host 侧无需额外同步即可安全访问目的内存复制方向通过aclrtMemcpyKind枚举的ACL_MEMCPY_DEVICE_TO_DEVICE指定表示源、目的地址均位于 Device 侧样例价值作为 memory 系列样例中同步 D2D的基准示范与同目录下的 6_d2d_async_memory_copy异步版使用aclrtMemcpyAsync形成对照便于开发者理解同步与异步拷贝的差异。该样例是 example/1_basic_features/memory/README.md 中列出的第 5 个内存样例与0_h2h、1_h2d_sync、2_h2d_async、3_d2h_sync、4_d2h_async等样例共同覆盖了 Host/Device 间的各种数据传输方向。2. 产品支持情况根据 5_d2d_sync_memory_copy/README.md 的说明本样例支持以下产品产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√运行前需确认本机安装有对应版本的 CANN 软件包并完成 CANN 环境变量配置set_env.sh。3. 整体流程与数据流向从 main.cpp 的代码结构看整个样例的执行顺序可以归纳为初始化 → 资源创建 → 内存申请 → 写入 → 复制 → 读取校验 → 资源释放七个阶段aclInit 初始化 └─ aclrtSetDevice(0) 指定 Device └─ aclrtCreateStream 创建 Stream ├─ aclrtMalloc 申请 devPtrA / devPtrB各 1 MiB ├─ 核函数 DeviceWrite向 devPtrA 写入值 123并打印 Source data ├─ aclrtMemcpy(devPtrB, devPtrA, ACL_MEMCPY_DEVICE_TO_DEVICE) 同步复制 ├─ 核函数 DeviceRead读取 devPtrB 的值并打印 Destination data ├─ aclrtDestroyStreamForce 销毁 Stream ├─ aclrtFree 释放 devPtrA / devPtrB └─ aclrtResetDeviceForce 复位 Device aclFinalize 去初始化数据流向为内核写入 devPtrA → 同步 D2D 拷贝到 devPtrB → 内核读取 devPtrB通过对比源数据与目的数据是否均为123来验证拷贝正确性。4. 编译与运行4.1 切换到样例目录将仓库克隆到已安装 CANN 软件的环境中进入样例目录cd ${git_clone_path}/example/1_basic_features/memory/5_d2d_sync_memory_copy4.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其中set_sample_env.sh见 example/set_sample_env.sh负责通过解析 CANN 安装布局${cann_path}/include${cann_path}/lib64等候选组合定位 ACL 头文件与libacl_rt.so调用aclrtGetSocName自动探测当前环境的SOC_VERSION探测ascendc.cmake所在目录并导出ASCENDC_CMAKE_DIR供 CMake 构建内核库使用依次导出ASCEND_INSTALL_PATH、ASCEND_HOME_PATH、SOC_VERSION、ASCENDC_CMAKE_DIR四个变量。4.3 一键编译运行bash run.shrun.sh见 run.sh内部做了四件事校验ASCEND_HOME_PATH是否已设置未设置则报错提示先 sourceset_env.shsource ${ASCEND_HOME_PATH}/bin/setenv.bash补齐编译工具链环境以cmake -B build -DASCEND_CANN_PACKAGE_PATH...配置并执行构建、安装运行./build/main将输出同时写入终端与output_msg.txt再用awk分别提取Source data:与Destination data:的值并比较一致则输出[SUCCESS]否则输出[FAILURE]并返回非零退出码。5. 源码深度解读5.1 主程序main.cpp 逐段拆解样例主程序位于 main.cpp核心逻辑如下#include acl/acl.h #include kernel_func/kernel_ops.h #include utils.h #include cstdio int32_t main() { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream stream nullptr; aclrtCreateStream(stream); // Allocate memory on the device uint64_t size 1 * 1024 * 1024; int* devPtrA; int* devPtrB; CHECK_ERROR(aclrtMalloc((void**)devPtrA, size, ACL_MEM_MALLOC_HUGE_FIRST)); INFO_LOG(Allocate memory on the device memory %p successfully, devPtrA); CHECK_ERROR(aclrtMalloc((void**)devPtrB, size, ACL_MEM_MALLOC_HUGE_FIRST)); INFO_LOG(Allocate memory on the device memory %p successfully, devPtrB); // Write the data to the virtual address devPtrA constexpr uint32_t blockDim 1; int writeValue 123; WriteDo(blockDim, stream, devPtrA, writeValue); INFO_LOG(Write the data %d to the virtual memory %p, writeValue, devPtrA); // Copy memory from address devPtrA to address devPtrB synchronously CHECK_ERROR(aclrtMemcpy(devPtrB, size, devPtrA, size, ACL_MEMCPY_DEVICE_TO_DEVICE)); INFO_LOG(Copy memory from memory %p to memory %p, devPtrA, devPtrB); // Read the value at address devPtrB ReadDo(blockDim, stream, devPtrB); // Release resource of the device aclrtDestroyStreamForce(stream); aclrtFree(devPtrA); aclrtFree(devPtrB); aclrtResetDeviceForce(deviceId); aclFinalize(); return 0; }各步骤要点内存申请size 1 * 1024 * 10241 MiB使用ACL_MEM_MALLOC_HUGE_FIRST策略即优先申请大页内存失败时回退到普通页见下文第 6 节的枚举说明写入WriteDo以blockDim 1的核函数向devPtrA写入整数123同步拷贝aclrtMemcpy(devPtrB, size, devPtrA, size, ACL_MEMCPY_DEVICE_TO_DEVICE)destMax与count均传size方向为 Device→Device读取校验ReadDo读出devPtrB中的值用于人工/脚本校验与源数据一致资源回收依次销毁 Stream、释放内存、复位 Device、去初始化。注意示例中main()在 utils.h 定义的CHECK_ERROR宏作用下任何 ACL 调用返回非ACL_SUCCESS都会打印Operation failed: ... returned error code %d并提前return -1。这是样例统一的错误处理范式。5.2 辅助核函数write_read_value.cpp写入与读取动作由 example/kernel_func/write_read_value.cpp 中的两个__aicore__核函数完成extern 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); } 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); }DeviceWrite把value即 123写入devPtr[idx]idx block_idx配合blockDim 1即写入第 0 个元素并在 Device 侧打印Source dataDeviceRead从devPtr[idx]读出值并打印Destination data用于确认拷贝后的目的内存内容。WriteDo/ReadDo的宿主侧封装以blockDim, nullptr, stream语法下发核函数在 kernel_ops.h 中声明。该内核源码通过 CMake 的ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp)编译为静态库见 CMakeLists.txt最终与ascendcl一起链接进main可执行文件。5.3 构建脚本CMakeLists.txtinclude(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp) include_directories(${ASCEND_CANN_PACKAGE_PATH}/include) link_directories(${ASCEND_CANN_PACKAGE_PATH}/lib64) add_executable(main main.cpp) target_link_libraries(main PRIVATE ascendcl kernels)关键点依赖set_sample_env.sh导出的ASCENDC_CMAKE_DIR来include(ascendc.cmake)进而获得ascendc_library内核编译能力头文件目录指向 CANN 包内include链接目录指向lib64链接库为ascendclRuntime ACL 库与本地编译出的kernels内核库。6. 关键 CANN Runtime API 详解6.1 初始化与去初始化接口作用aclInit(nullptr)初始化 AscendCL 运行环境nullptr表示使用默认配置不加载配置文件aclFinalize()去初始化释放全局资源应与aclInit配对调用6.2 Device 管理接口作用aclrtSetDevice(deviceId)指定当前进程用于运算的 Device本例为 0 号 Device后续 Stream、内存申请均绑定该 DeviceaclrtResetDeviceForce(deviceId)强制复位 Device回收 Device 上资源Force语义下即使存在未完成任务也会直接复位6.3 Stream 管理接口作用aclrtCreateStream(stream)创建 Stream核函数与异步任务在该 Stream 上排队执行aclrtDestroyStreamForce(stream)强制销毁 Stream 并丢弃其上的所有任务同步拷贝本身不依赖 Stream 完成数据搬运但样例中的核函数写入/读取需要 Stream 来下发执行因此仍需创建 Stream。6.4 内存管理aclrtMalloc((void**)devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtFree(devPtr);aclrtMalloc用于申请 Device 侧内存第三个参数是aclrtMemMallocPolicy枚举定义于 include/external/acl/acl_rt.h常用取值包括枚举值语义ACL_MEM_MALLOC_HUGE_FIRST优先申请大页huge page内存失败回退普通页ACL_MEM_MALLOC_HUGE_ONLY仅申请大页内存ACL_MEM_MALLOC_NORMAL_ONLY仅申请普通页内存ACL_MEM_MALLOC_HUGE_FIRST_P2P/ACL_MEM_MALLOC_HUGE_ONLY_P2P/ACL_MEM_MALLOC_NORMAL_ONLY_P2P支持 P2P 访问的对应策略ACL_MEM_MALLOC_HUGE1G_ONLY仅申请 1G 大页内存样例选用ACL_MEM_MALLOC_HUGE_FIRST以兼顾大页性能与可用性。6.5 数据传输aclrtMemcpyaclrtMemcpy是本次样例的核心接口原型与语义在 include/external/acl/acl_rt.h 中有完整定义ACL_FUNC_VISIBILITY aclError aclrtMemcpy( void* dst, size_t destMax, const void* src, size_t count, aclrtMemcpyKind kind);参数含义dst目的地址指针本例为devPtrBDevice 地址destMax目的地址内存的最大长度字节必须不小于countsrc源地址指针本例为devPtrADevice 地址count待复制字节数本例为 1 MiBkind复制类型取aclrtMemcpyKind枚举aclrtMemcpyKind枚举同样定义于 acl_rt.h与本样例直接相关的是ACL_MEMCPY_DEVICE_TO_DEVICEDevice 到 Device 的复制即本样例使用的模式其他方向还包括ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_HOST、ACL_MEMCPY_DEFAULT以及ACL_MEMCPY_HOST_TO_BUF_TO_DEVICE、ACL_MEMCPY_INNER_DEVICE_TO_DEVICE等。同步语义aclrtMemcpy是同步复制接口函数返回时复制已完成Host 可以立即读取dst指向的内存例如随后直接在 Device 上下发DeviceRead核函数读取无需额外aclrtSynchronizeStream。这正是它与异步版aclrtMemcpyAsync需搭配 Stream 同步参见 6_d2d_async_memory_copy的核心区别。另外CANN 还在 include/external/acl/acl_rt_api.h 中提供了 C 模板重载可省去void*强转并自动推导参数template typename T, typename U static inline aclError aclrtMemcpy(T* dst, size_t destMax, const U* src, size_t count, aclrtMemcpyKind kind) { return ::aclrtMemcpy(static_castvoid*(dst), destMax, static_castconst void*(src), count, kind); }7. 输出结果与自动化校验7.1 预期输出正常运行时的输出如下地址为实际分配结果用0x...示意[INFO] Allocate memory on the device 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... Source data: 123 Destination data: 123其中Source data: 123由DeviceWrite核函数打印代表源内存devPtrA中的值Destination data: 123由DeviceRead核函数打印代表经同步 D2D 拷贝后目的内存devPtrB中的值。7.2 run.sh 的自动化判定run.sh在运行完./build/main后通过awk提取两个字段并比较source_value$(awk -F: /Source data:/ {gsub(/^ | $/, , $2); print $2; exit} output_msg.txt) destination_value$(awk -F: /Destination data:/ {gsub(/^ | $/, , $2); print $2; exit} output_msg.txt) 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. Value at source is ${source_value}, but value at destination is ${destination_value} exit 1 fi只有源值与目的值相等均为123时才会输出[SUCCESS]并以退出码 0 结束从而验证aclrtMemcpy的 Device→Device 同步复制确实把数据完整搬移到了目的地址。8. 使用要点与常见注意事项destMax必须不小于countaclrtMemcpy会校验目的内存最大长度防止越界写。样例中两块内存都按 1 MiB 申请destMax与count均为size是规范用法方向匹配源、目的均在 Device 侧时务必使用ACL_MEMCPY_DEVICE_TO_DEVICE混用 Host/Device 方向会导致拷贝失败或行为异常同步与异步的选择若拷贝后 Host 需要立即读取结果、或拷贝任务较少使用同步aclrtMemcpy更简单若需在拷贝期间让 Host 继续下发其他任务、或追求吞吐可参考 6_d2d_async_memory_copy 使用aclrtMemcpyAsync并配合aclrtSynchronizeStream错误处理所有 ACL 调用应检查返回值样例统一使用CHECK_ERROR宏utils.h在失败时打印错误码并退出便于快速定位问题资源释放顺序建议按销毁 Stream → 释放内存 → 复位 Device →aclFinalize的顺序回收资源避免句柄泄漏或残留任务导致的异常运行环境本样例依赖 CANN 环境变量set_env.sh、SOC_VERSION与ASCENDC_CMAKE_DIRset_sample_env.sh自动探测请确保在支持的产品Ascend 950 系列、Atlas A2/A3 系列上运行。9. 进一步探索对比学习异步 D2D 拷贝6_d2d_async_memory_copy同一系列的完整内存传输方向覆盖memory 样例目录Host↔Device 同步/异步拷贝、多 Stream 内存同步、IPC 共享等内存复制描述符方式单 Device 内异步复制并校验13_memcpy_descriptor接口声明与枚举定义acl_rt.h、C 模板重载见 acl_rt_api.h【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN Runtime 样例解析：使用 aclrtMemcpy 实现 Device 内同步内存复制（D2D Sync Memory Copy）
CANN Runtime 样例解析使用 aclrtMemcpy 实现 Device 内同步内存复制D2D Sync Memory Copy【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime导读本篇文章围绕 CANN runtime 仓库中的5_d2d_sync_memory_copy样例展开完整讲解如何在昇腾 AI 处理器的 Device 内部完成申请两块 Device 内存 → 由核函数写入源地址 → 通过aclrtMemcpy以ACL_MEMCPY_DEVICE_TO_DEVICE方式同步复制 → 由核函数读取目的地址校验的完整链路。读完本文你将掌握aclrtMemcpy同步 D2D 拷贝的接口语义与参数约束、样例的编译运行流程、底层核函数DeviceWrite/DeviceRead如何配合验证拷贝结果以及run.sh如何自动化判定样例成功与否。1. 样例概述Device 内同步内存复制解决什么问题在 AI 推理与训练程序中数据常常先在 Device 上以aclrtMalloc申请经算子或核函数处理后需要在 Device 内存的另一个地址上复用例如特征图搬运、中间结果暂存、多 Buffer 轮转。此时并不需要经过 Host 中转直接调用内存复制接口即可完成 Device 内部的数据搬移。同步语义aclrtMemcpy是同步接口调用返回时内存复制任务已经完成Host 侧无需额外同步即可安全访问目的内存复制方向通过aclrtMemcpyKind枚举的ACL_MEMCPY_DEVICE_TO_DEVICE指定表示源、目的地址均位于 Device 侧样例价值作为 memory 系列样例中同步 D2D的基准示范与同目录下的 6_d2d_async_memory_copy异步版使用aclrtMemcpyAsync形成对照便于开发者理解同步与异步拷贝的差异。该样例是 example/1_basic_features/memory/README.md 中列出的第 5 个内存样例与0_h2h、1_h2d_sync、2_h2d_async、3_d2h_sync、4_d2h_async等样例共同覆盖了 Host/Device 间的各种数据传输方向。2. 产品支持情况根据 5_d2d_sync_memory_copy/README.md 的说明本样例支持以下产品产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√运行前需确认本机安装有对应版本的 CANN 软件包并完成 CANN 环境变量配置set_env.sh。3. 整体流程与数据流向从 main.cpp 的代码结构看整个样例的执行顺序可以归纳为初始化 → 资源创建 → 内存申请 → 写入 → 复制 → 读取校验 → 资源释放七个阶段aclInit 初始化 └─ aclrtSetDevice(0) 指定 Device └─ aclrtCreateStream 创建 Stream ├─ aclrtMalloc 申请 devPtrA / devPtrB各 1 MiB ├─ 核函数 DeviceWrite向 devPtrA 写入值 123并打印 Source data ├─ aclrtMemcpy(devPtrB, devPtrA, ACL_MEMCPY_DEVICE_TO_DEVICE) 同步复制 ├─ 核函数 DeviceRead读取 devPtrB 的值并打印 Destination data ├─ aclrtDestroyStreamForce 销毁 Stream ├─ aclrtFree 释放 devPtrA / devPtrB └─ aclrtResetDeviceForce 复位 Device aclFinalize 去初始化数据流向为内核写入 devPtrA → 同步 D2D 拷贝到 devPtrB → 内核读取 devPtrB通过对比源数据与目的数据是否均为123来验证拷贝正确性。4. 编译与运行4.1 切换到样例目录将仓库克隆到已安装 CANN 软件的环境中进入样例目录cd ${git_clone_path}/example/1_basic_features/memory/5_d2d_sync_memory_copy4.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其中set_sample_env.sh见 example/set_sample_env.sh负责通过解析 CANN 安装布局${cann_path}/include${cann_path}/lib64等候选组合定位 ACL 头文件与libacl_rt.so调用aclrtGetSocName自动探测当前环境的SOC_VERSION探测ascendc.cmake所在目录并导出ASCENDC_CMAKE_DIR供 CMake 构建内核库使用依次导出ASCEND_INSTALL_PATH、ASCEND_HOME_PATH、SOC_VERSION、ASCENDC_CMAKE_DIR四个变量。4.3 一键编译运行bash run.shrun.sh见 run.sh内部做了四件事校验ASCEND_HOME_PATH是否已设置未设置则报错提示先 sourceset_env.shsource ${ASCEND_HOME_PATH}/bin/setenv.bash补齐编译工具链环境以cmake -B build -DASCEND_CANN_PACKAGE_PATH...配置并执行构建、安装运行./build/main将输出同时写入终端与output_msg.txt再用awk分别提取Source data:与Destination data:的值并比较一致则输出[SUCCESS]否则输出[FAILURE]并返回非零退出码。5. 源码深度解读5.1 主程序main.cpp 逐段拆解样例主程序位于 main.cpp核心逻辑如下#include acl/acl.h #include kernel_func/kernel_ops.h #include utils.h #include cstdio int32_t main() { aclInit(nullptr); int32_t deviceId 0; aclrtSetDevice(deviceId); aclrtStream stream nullptr; aclrtCreateStream(stream); // Allocate memory on the device uint64_t size 1 * 1024 * 1024; int* devPtrA; int* devPtrB; CHECK_ERROR(aclrtMalloc((void**)devPtrA, size, ACL_MEM_MALLOC_HUGE_FIRST)); INFO_LOG(Allocate memory on the device memory %p successfully, devPtrA); CHECK_ERROR(aclrtMalloc((void**)devPtrB, size, ACL_MEM_MALLOC_HUGE_FIRST)); INFO_LOG(Allocate memory on the device memory %p successfully, devPtrB); // Write the data to the virtual address devPtrA constexpr uint32_t blockDim 1; int writeValue 123; WriteDo(blockDim, stream, devPtrA, writeValue); INFO_LOG(Write the data %d to the virtual memory %p, writeValue, devPtrA); // Copy memory from address devPtrA to address devPtrB synchronously CHECK_ERROR(aclrtMemcpy(devPtrB, size, devPtrA, size, ACL_MEMCPY_DEVICE_TO_DEVICE)); INFO_LOG(Copy memory from memory %p to memory %p, devPtrA, devPtrB); // Read the value at address devPtrB ReadDo(blockDim, stream, devPtrB); // Release resource of the device aclrtDestroyStreamForce(stream); aclrtFree(devPtrA); aclrtFree(devPtrB); aclrtResetDeviceForce(deviceId); aclFinalize(); return 0; }各步骤要点内存申请size 1 * 1024 * 10241 MiB使用ACL_MEM_MALLOC_HUGE_FIRST策略即优先申请大页内存失败时回退到普通页见下文第 6 节的枚举说明写入WriteDo以blockDim 1的核函数向devPtrA写入整数123同步拷贝aclrtMemcpy(devPtrB, size, devPtrA, size, ACL_MEMCPY_DEVICE_TO_DEVICE)destMax与count均传size方向为 Device→Device读取校验ReadDo读出devPtrB中的值用于人工/脚本校验与源数据一致资源回收依次销毁 Stream、释放内存、复位 Device、去初始化。注意示例中main()在 utils.h 定义的CHECK_ERROR宏作用下任何 ACL 调用返回非ACL_SUCCESS都会打印Operation failed: ... returned error code %d并提前return -1。这是样例统一的错误处理范式。5.2 辅助核函数write_read_value.cpp写入与读取动作由 example/kernel_func/write_read_value.cpp 中的两个__aicore__核函数完成extern 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); } 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); }DeviceWrite把value即 123写入devPtr[idx]idx block_idx配合blockDim 1即写入第 0 个元素并在 Device 侧打印Source dataDeviceRead从devPtr[idx]读出值并打印Destination data用于确认拷贝后的目的内存内容。WriteDo/ReadDo的宿主侧封装以blockDim, nullptr, stream语法下发核函数在 kernel_ops.h 中声明。该内核源码通过 CMake 的ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp)编译为静态库见 CMakeLists.txt最终与ascendcl一起链接进main可执行文件。5.3 构建脚本CMakeLists.txtinclude(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC ../../../kernel_func/write_read_value.cpp) include_directories(${ASCEND_CANN_PACKAGE_PATH}/include) link_directories(${ASCEND_CANN_PACKAGE_PATH}/lib64) add_executable(main main.cpp) target_link_libraries(main PRIVATE ascendcl kernels)关键点依赖set_sample_env.sh导出的ASCENDC_CMAKE_DIR来include(ascendc.cmake)进而获得ascendc_library内核编译能力头文件目录指向 CANN 包内include链接目录指向lib64链接库为ascendclRuntime ACL 库与本地编译出的kernels内核库。6. 关键 CANN Runtime API 详解6.1 初始化与去初始化接口作用aclInit(nullptr)初始化 AscendCL 运行环境nullptr表示使用默认配置不加载配置文件aclFinalize()去初始化释放全局资源应与aclInit配对调用6.2 Device 管理接口作用aclrtSetDevice(deviceId)指定当前进程用于运算的 Device本例为 0 号 Device后续 Stream、内存申请均绑定该 DeviceaclrtResetDeviceForce(deviceId)强制复位 Device回收 Device 上资源Force语义下即使存在未完成任务也会直接复位6.3 Stream 管理接口作用aclrtCreateStream(stream)创建 Stream核函数与异步任务在该 Stream 上排队执行aclrtDestroyStreamForce(stream)强制销毁 Stream 并丢弃其上的所有任务同步拷贝本身不依赖 Stream 完成数据搬运但样例中的核函数写入/读取需要 Stream 来下发执行因此仍需创建 Stream。6.4 内存管理aclrtMalloc((void**)devPtr, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtFree(devPtr);aclrtMalloc用于申请 Device 侧内存第三个参数是aclrtMemMallocPolicy枚举定义于 include/external/acl/acl_rt.h常用取值包括枚举值语义ACL_MEM_MALLOC_HUGE_FIRST优先申请大页huge page内存失败回退普通页ACL_MEM_MALLOC_HUGE_ONLY仅申请大页内存ACL_MEM_MALLOC_NORMAL_ONLY仅申请普通页内存ACL_MEM_MALLOC_HUGE_FIRST_P2P/ACL_MEM_MALLOC_HUGE_ONLY_P2P/ACL_MEM_MALLOC_NORMAL_ONLY_P2P支持 P2P 访问的对应策略ACL_MEM_MALLOC_HUGE1G_ONLY仅申请 1G 大页内存样例选用ACL_MEM_MALLOC_HUGE_FIRST以兼顾大页性能与可用性。6.5 数据传输aclrtMemcpyaclrtMemcpy是本次样例的核心接口原型与语义在 include/external/acl/acl_rt.h 中有完整定义ACL_FUNC_VISIBILITY aclError aclrtMemcpy( void* dst, size_t destMax, const void* src, size_t count, aclrtMemcpyKind kind);参数含义dst目的地址指针本例为devPtrBDevice 地址destMax目的地址内存的最大长度字节必须不小于countsrc源地址指针本例为devPtrADevice 地址count待复制字节数本例为 1 MiBkind复制类型取aclrtMemcpyKind枚举aclrtMemcpyKind枚举同样定义于 acl_rt.h与本样例直接相关的是ACL_MEMCPY_DEVICE_TO_DEVICEDevice 到 Device 的复制即本样例使用的模式其他方向还包括ACL_MEMCPY_HOST_TO_HOST、ACL_MEMCPY_HOST_TO_DEVICE、ACL_MEMCPY_DEVICE_TO_HOST、ACL_MEMCPY_DEFAULT以及ACL_MEMCPY_HOST_TO_BUF_TO_DEVICE、ACL_MEMCPY_INNER_DEVICE_TO_DEVICE等。同步语义aclrtMemcpy是同步复制接口函数返回时复制已完成Host 可以立即读取dst指向的内存例如随后直接在 Device 上下发DeviceRead核函数读取无需额外aclrtSynchronizeStream。这正是它与异步版aclrtMemcpyAsync需搭配 Stream 同步参见 6_d2d_async_memory_copy的核心区别。另外CANN 还在 include/external/acl/acl_rt_api.h 中提供了 C 模板重载可省去void*强转并自动推导参数template typename T, typename U static inline aclError aclrtMemcpy(T* dst, size_t destMax, const U* src, size_t count, aclrtMemcpyKind kind) { return ::aclrtMemcpy(static_castvoid*(dst), destMax, static_castconst void*(src), count, kind); }7. 输出结果与自动化校验7.1 预期输出正常运行时的输出如下地址为实际分配结果用0x...示意[INFO] Allocate memory on the device 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... Source data: 123 Destination data: 123其中Source data: 123由DeviceWrite核函数打印代表源内存devPtrA中的值Destination data: 123由DeviceRead核函数打印代表经同步 D2D 拷贝后目的内存devPtrB中的值。7.2 run.sh 的自动化判定run.sh在运行完./build/main后通过awk提取两个字段并比较source_value$(awk -F: /Source data:/ {gsub(/^ | $/, , $2); print $2; exit} output_msg.txt) destination_value$(awk -F: /Destination data:/ {gsub(/^ | $/, , $2); print $2; exit} output_msg.txt) 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. Value at source is ${source_value}, but value at destination is ${destination_value} exit 1 fi只有源值与目的值相等均为123时才会输出[SUCCESS]并以退出码 0 结束从而验证aclrtMemcpy的 Device→Device 同步复制确实把数据完整搬移到了目的地址。8. 使用要点与常见注意事项destMax必须不小于countaclrtMemcpy会校验目的内存最大长度防止越界写。样例中两块内存都按 1 MiB 申请destMax与count均为size是规范用法方向匹配源、目的均在 Device 侧时务必使用ACL_MEMCPY_DEVICE_TO_DEVICE混用 Host/Device 方向会导致拷贝失败或行为异常同步与异步的选择若拷贝后 Host 需要立即读取结果、或拷贝任务较少使用同步aclrtMemcpy更简单若需在拷贝期间让 Host 继续下发其他任务、或追求吞吐可参考 6_d2d_async_memory_copy 使用aclrtMemcpyAsync并配合aclrtSynchronizeStream错误处理所有 ACL 调用应检查返回值样例统一使用CHECK_ERROR宏utils.h在失败时打印错误码并退出便于快速定位问题资源释放顺序建议按销毁 Stream → 释放内存 → 复位 Device →aclFinalize的顺序回收资源避免句柄泄漏或残留任务导致的异常运行环境本样例依赖 CANN 环境变量set_env.sh、SOC_VERSION与ASCENDC_CMAKE_DIRset_sample_env.sh自动探测请确保在支持的产品Ascend 950 系列、Atlas A2/A3 系列上运行。9. 进一步探索对比学习异步 D2D 拷贝6_d2d_async_memory_copy同一系列的完整内存传输方向覆盖memory 样例目录Host↔Device 同步/异步拷贝、多 Stream 内存同步、IPC 共享等内存复制描述符方式单 Device 内异步复制并校验13_memcpy_descriptor接口声明与枚举定义acl_rt.h、C 模板重载见 acl_rt_api.h【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考