尧图网页设计 
CANN 入门样例解析基于 AscendC 与 VectorCore 的 vector_add 向量加法实现【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples导读vector_add是 CANN 高性能实战样例仓库 cann-samples 中位于 Samples/0_Introduction 下的入门级样例它展示了如何在昇腾 AI 处理器的VectorCore向量核硬件单元上使用Ascend C 编程语言实现z x y的向量加法算子。通过本样例读者可以掌握单算子开发中最核心的三件事Tiling 切分与多核调度、TPipe/TQue 双缓冲流水并行、Host 侧 ACL 编程框架与精度校验同时还能了解基于aclprofTensor的算子性能分析元信息上报方法为后续阅读仓库中 1_Features 与 2_Performance 的调优实践打下基础。一、样例概述vector_add样例在昇腾 AI 处理器的 VectorCore 上完成一次完整的向量加法Host 侧通过 ACLAscend Computing LanguageRuntime 接口完成设备初始化、内存申请与搬运、Stream 创建与同步、结果回拷与逐元素校验Device 侧编写一个__global__ __aicore__ __vector__修饰的 Kernel 函数在多个 AI Vector Core 上并行执行每个核负责一段连续数据精度基准Host 侧以h_C[i] h_A[i] h_B[i]的严格逐元素比较作为标准 CPU 精度基准不依赖浮点容差样例数据由uniform_real_distributionfloat在[0.0, 10.0)区间随机生成加法结果可精确表示因此允许全等比较。样例核心源码位于 Samples/0_Introduction/vector_add/main.asc构建配置位于 Samples/0_Introduction/vector_add/CMakeLists.txt。二、关键特性特性说明流水并行使用TPipeTQueTPosition::VECIN/VECOUT, BUFFER_NUM构建 DoubleBuffer 双缓冲流水隐藏搬入/搬出与计算之间的等待参数可配totalLength向量长度可在运行入口调整Tiling 参数核数、块长、tile 大小由 Host 侧依据硬件资源自动推导精度对比提供标准 CPU 实现作为精度基准逐元素比对并打印执行结果性能分析元信息使用aclprofTensor/aclprofRangePushEx上报算子输入输出张量元信息供 Profiling 工具解析该 API 自 CANN 9.1.0 起提供三、核心实现原理从 Host 到 Kernel 的完整链路3.1 整体结构main.asc采用单文件单算子工程结构包含三层main()main.asc#L223-L238ACL 初始化、设置 Device、创建 Stream、调用run_vector_add(stream, 409600)后清理资源run_vector_add()main.asc#L139-L221Host 侧数据准备、Tiling 计算、Kernel 下发与结果校验add_kernelmain.asc#L85-L137Device 侧 Ascend C Kernel完成双缓冲流水向量加法。int main() { CHECK_ACL(aclInit(nullptr)); int32_t deviceId 0; CHECK_ACL(aclrtSetDevice(deviceId)); aclrtStream stream nullptr; CHECK_ACL(aclrtCreateStream(stream)); int result run_vector_add(stream, 409600); // 默认向量长度 409600 CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); return result; }所有 ACL 调用统一走CHECK_ACL宏main.asc#L30-L38做错误检查失败时打印错误码与文件行号并提前返回。设备内存通过aclrtMalloc分配、并由自定义删除器AclrtFreeDeletermain.asc#L41-L47配合std::unique_ptr自动释放避免异常路径下的内存泄漏。3.2 Tiling 计算多核切分与 Tile 尺寸推导calc_tiling_params()main.asc#L67-L83是整条数据通路的关键它根据硬件信息自动推导三个 Tiling 参数std::tupleint64_t, int64_t, int64_t calc_tiling_params(int64_t totalLength) { constexpr static int64_t MIN_ELEMS_PER_CORE 1024; // 每核最少处理的元素数 constexpr static int64_t BUFFER_NUM 2; // 双缓冲 constexpr static int64_t QUEUE_NUM 3; // 每队列分配 3 份缓冲区空间 auto ascendcPlatform platform_ascendc::PlatformAscendCManager::GetInstance(); uint64_t ubSize; ascendcPlatform-GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize); // 查询 UB 容量 int64_t coreNum ascendcPlatform-GetCoreNumAiv(); // 查询 VectorCore 核数 int64_t numBlocks std::min(coreNum, (totalLength MIN_ELEMS_PER_CORE - 1) / MIN_ELEMS_PER_CORE); numBlocks std::max(numBlocks, static_castint64_t(1)); int64_t blockLength (totalLength numBlocks - 1) / numBlocks; // 每个核处理的块长 int64_t tileSize ubSize / BUFFER_NUM / QUEUE_NUM; // 每个 tile 的字节数 constexpr int64_t ALIGN_SIZE 32; tileSize (tileSize / ALIGN_SIZE) * ALIGN_SIZE; // 32 字节对齐 return std::make_tuple(numBlocks, blockLength, tileSize); }推导逻辑可以概括为三步核数分配numBlocks取min(核心数, ceil(totalLength / 1024))即每个核至少处理 1024 个元素数据量不足以铺满所有核时只启用需要的核数且最少为 1 个核块长计算blockLength按ceil(totalLength / numBlocks)向上取整保证每个核的数据量均衡尾块由 Kernel 内部裁剪Tile 大小tileSize由 UB 总容量除以BUFFER_NUM(2)与QUEUE_NUM(3)得出这是为了给双缓冲的输入队列 X、输入队列 Y、输出队列 Z 各留出独立空间最后按 32 字节对齐对应 Vector 搬运的对齐要求。这里使用的platform_ascendc::PlatformAscendCManager来自 kernel_operator.h 配套的platform/platform_ascendc.h头文件通过GetCoreMemSize(CoreMemType::UB, ...)与GetCoreNumAiv()在运行时查询硬件实际资源使 Tiling 具备跨代际硬件的自适应性——这也是仓库中 CMakeLists.txt 同时声明支持dav-2201Ascend 910B/C与dav-3510Ascend 950的原因之一。3.3 Kernel 实现TPipe/TQue 双缓冲流水add_kernelmain.asc#L85-L137是 Ascend C 编程模型中典型的CopyIn → Compute → CopyOut三段式流水template typename T __global__ __aicore__ __vector__ void add_kernel( GM_ADDR x, GM_ADDR y, GM_ADDR z, int64_t totalLength, int64_t blockLength, uint32_t tileSize) { constexpr static int64_t BUFFER_NUM 2; AscendC::TPipe pipe; AscendC::GlobalTensorT xGm, yGm, zGm; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueX; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueY; AscendC::TQueAscendC::TPosition::VECOUT, BUFFER_NUM outQueueZ; pipe.InitBuffer(inQueueX, BUFFER_NUM, tileSize); pipe.InitBuffer(inQueueY, BUFFER_NUM, tileSize); pipe.InitBuffer(outQueueZ, BUFFER_NUM, tileSize); xGm.SetGlobalBuffer((__gm__ T *)x blockLength * AscendC::GetBlockIdx()); yGm.SetGlobalBuffer((__gm__ T *)y blockLength * AscendC::GetBlockIdx()); zGm.SetGlobalBuffer((__gm__ T *)z blockLength * AscendC::GetBlockIdx()); // ... 按 tile 循环DataCopyPad 搬入 - AscendC::Add 计算 - DataCopyPad 搬出 }实现要点全局缓冲区偏移通过AscendC::GetBlockIdx()获取当前核 ID将xGm/yGm/zGm的起始地址各自偏移blockLength * blockIdx实现核间数据不重叠的分块访问尾块处理currentBlockLength用totalLength - blockIdx * blockLength与blockLength取较小值处理最后一个核数据不足一个块长的情况tile 层同理tileElementNum取min(剩余元素, elementNumPerTile)对齐搬运AscendC::DataCopyPad配合DataCopyExtParamsblockCount1、blockLentileElementNum * sizeof(T)、srcStride/dstStride0完成 GM↔UB 搬运支持非对齐的尾部数据计算指令核心运算只有一行AscendC::Add(zLocal, xLocal, yLocal, tileElementNum)一条向量指令完成一批元素的加法双缓冲TQue..., 2声明每个队列包含 2 个 bufferAllocTensor/EnQue/DeQue/FreeTensor的交替使用让下一 tile 的 CopyIn 与当前 tile 的 Compute 重叠执行从而隐藏 MTE 搬运延迟。值得对照的是同目录下的 vector_function_addRegBase/VF 编程模型仅支持dav-3510与 vector_add_c_apiC 语言风格 APIasc_simd.h三者功能一致但编程模型不同适合对照学习 MemBase、RegBase 与 C API 三种写法的差异。3.4 Kernel 下发与性能分析元信息上报Kernel 以add_kernelfloatnumBlocks, nullptr, stream(...)的 CUDA 风格语法下发main.asc#L198。下发前后样例通过 ACL Profiling 接口上报了张量元信息aclprofTensor tensors[] { MakeVectorTensor(PROF_INPUT_TENSOR, numElements), // x MakeVectorTensor(PROF_INPUT_TENSOR, numElements), // y MakeVectorTensor(PROF_OUTPUT_TENSOR, numElements), // z }; aclprofTensorInfo tensorInfo { aclprofStr2Id(PROF_OP_NAME), // vector_add aclprofStr2Id(PROF_OP_TYPE), // VectorAdd 0, sizeof(tensors) / sizeof(aclprofTensor), 0, static_castuint32_t(numBlocks), stream, tensors }; aclprofEventAttributes attrs { ACL_PROF_EVENT_ATTR_VERSION, sizeof(aclprofEventAttributes::message), ACL_PROF_MESSAGE_TYPE_TENSOR_INFO, tensorInfo }; aclprofRangePushEx(attrs); // 压栈性能分析区间 add_kernelfloatnumBlocks, nullptr, stream(...); aclprofRangePop(); // 弹栈其中MakeVectorTensormain.asc#L56-L65将张量描述为ACL_FORMAT_NDACL_FLOAT 一维 shape 的元信息。这套aclprofTensor元信息上报 API 由acl/acl_prof.h提供自 CANN 9.1.0 起才可用这也是本样例环境要求必须不低于 CANN 9.1.0 weekly20260708的原因。上报后Profiling 工具即可在性能分析中还原出算子的输入输出 shape相关参考实现在仓库子模块 third_party/asc-devkit 的 examples 目录中。四、参数说明参数含义默认值/取值totalLength向量长度元素个数即每个输入向量的元素数run_vector_add(stream, 409600)中默认为 409600可自行调整Dtype 模板参数Kernel 数据类型支持FLOAT32add_kernelfloat可通过模板实例化扩展从源码结构看Kernel 已通过template typename T实现泛化main.asc#L85tileSize / sizeof(T)与tileElementNum * sizeof(T)均按类型推导因此扩展到其他数据类型只需新增模板实例与对应 Host 侧数据生成逻辑。五、支持架构与环境要求5.1 支持架构NPU ARCHdav-2201Ascend 910B/C、dav-3510Ascend 950架构约束在构建期由 cmake/sample_common.cmake 中的cann_sample_check_arch(dav-2201 dav-3510)宏强制检查当配置的NPU_ARCH不在期望列表时该样例会在 CMake 配置阶段被跳过并打印Skip sample ...: NPU_ARCH... is not supported不会进入构建也不会出现在--target help中。5.2 环境要求CANN 版本需 CANN 9.1.0 weekly20260708及以上。本样例使用了aclprofTensor等性能分析元信息上报 API见 main.asc#L176-L199旧版本缺少该接口将无法编译通过Toolkit 安装须安装完整社区版 CANN Toolkit 并执行source ${install_path}/ascend-toolkit/set_env.sh使环境变量生效详见 README.md 的环境部署一节构建依赖cmake 3.16.0、python 3.8.0、zip、git以及通过pip3 install -r requirements.txt安装的 Python 三方库环境自检建议构建前在仓库根目录运行python3 scripts/check_env.pyscripts/check_env.py脚本只检查不修改任何配置。六、编译与运行6.1 配置构建从仓库根目录初始化构建NPU_ARCH为必填参数# Ascend 950 cmake -S . -B build -DNPU_ARCHdav-3510 # Ascend 910B/C cmake -S . -B build -DNPU_ARCHdav-22016.2 编译 vector_add 目标cmake --build build --target vector_add构建由 Samples/0_Introduction/vector_add/CMakeLists.txt 驱动add_executable(vector_add main.asc)声明目标--npu-arch${NPU_ARCH}与-O3作为 Ascend C 编译选项链接m、dl、platform、tiling_api、msprofiler等运行时库其中msprofiler即支撑aclprofTensor上报的库产物安装到0_Introduction/vector_add目录。6.3 运行样例切换到可执行文件所在目录后直接运行cd ./build/Samples/0_Introduction/vector_add/ ./vector_add成功时打印Vector add completed successfully!若存在精度问题则打印Vector add failed!需要注意的是vector_add默认不接收命令行参数main()中固定传入 409600如需验证不同totalLength可直接修改 main.asc#L231 的入参后重新编译。七、结果验证与自动化测试样例内置了严格的 Host 侧校验将设备端结果h_C拷回后逐一检查h_C[i] h_A[i] h_B[i]任一元素不相等即判定失败并返回退出码 1main.asc#L206-L220。仓库同时提供了对应的 CI 用例tests/ci_functional_test.yaml它模拟全量构建安装后的运行方式- id: vector_add steps: - name: run_default cwd: build_out/0_Introduction/vector_add cmd: [./vector_add] pass_criteria: exit_code: 0 stdout_contains: [Vector add completed successfully!] stdout_not_contains: [Vector add failed!]也就是说通过cmake --build build --parallel全量构建并cmake --install build --prefix ./build_out安装后在build_out/0_Introduction/vector_add/目录执行./vector_add且退出码为 0、标准输出包含成功信息即视为功能测试通过。这也给开发者提供了一条编译产物 → 可执行文件 → 判定标准的完整验证路径。八、同类入门样例横向对照Samples/0_Introduction目录下共有三个向量加法语义相同的入门样例适合放在一起对照学习样例编程模型核心 API支持架构数据流vector_add本文MemBaseTQue/TPipeAscendC::Add操作 UB 上的 LocalTensordav-2201, dav-3510每步运算在 UB 上读写vector_add_c_apiC APIasc_simd.hasc_copy_gm2ub_sync/asc_add_sync/asc_copy_ub2gm_syncdav-2201同步搬入→同步计算→同步搬出vector_function_addRegBaseVector FunctionAscendC::Reg::LoadAlign/Add/StoreAlignUpdateMaskdav-3510Load 进寄存器→寄存器间运算→Store 回 UB其中 vector_add 的双缓冲流水与 C API 版本的单缓冲同步执行、RegBase 版本的寄存器直算形成渐进式对照前者体现以空间换时间的经典流水思想后两者分别体现接口简化与寄存器级数据驻留的优势。概念背景可进一步参考 vector_function_getting_started性能进阶可阅读 2_Performance 下的各调优 story。九、小结vector_add虽小却完整覆盖了 Ascend C 单算子开发的全部关键环节Tiling 三要素核数、块长、tile 大小由硬件资源查询自动推导兼顾多核并行与 UB 容量约束TPipe/TQue 双缓冲流水让搬运与计算重叠是后续所有性能调优样例的基石模式ACL Host 框架提供从初始化、内存管理、Stream 同步到精度校验的完整范式aclprofTensor 元信息上报展示了如何让算子被 Profiling 工具正确识别是 CANN 9.1.0 新特性的实战示例。建议读者在运行通过后尝试修改totalLength观察尾块处理逻辑或对照 vector_add_c_api 与 vector_function_add 体会三种编程模型的差异再进入 matmul 等更复杂的入门样例继续进阶。参考资料cann-samples 根 README环境部署与快速入门vector_add 样例源码vector_add 构建配置入门样例索引架构检查宏实现vector_add CI 功能测试用例asc-devkit 子模块性能分析参考样例所在【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN 入门样例解析：基于 AscendC 与 VectorCore 的 vector_add 向量加法实现
CANN 入门样例解析基于 AscendC 与 VectorCore 的 vector_add 向量加法实现【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples导读vector_add是 CANN 高性能实战样例仓库 cann-samples 中位于 Samples/0_Introduction 下的入门级样例它展示了如何在昇腾 AI 处理器的VectorCore向量核硬件单元上使用Ascend C 编程语言实现z x y的向量加法算子。通过本样例读者可以掌握单算子开发中最核心的三件事Tiling 切分与多核调度、TPipe/TQue 双缓冲流水并行、Host 侧 ACL 编程框架与精度校验同时还能了解基于aclprofTensor的算子性能分析元信息上报方法为后续阅读仓库中 1_Features 与 2_Performance 的调优实践打下基础。一、样例概述vector_add样例在昇腾 AI 处理器的 VectorCore 上完成一次完整的向量加法Host 侧通过 ACLAscend Computing LanguageRuntime 接口完成设备初始化、内存申请与搬运、Stream 创建与同步、结果回拷与逐元素校验Device 侧编写一个__global__ __aicore__ __vector__修饰的 Kernel 函数在多个 AI Vector Core 上并行执行每个核负责一段连续数据精度基准Host 侧以h_C[i] h_A[i] h_B[i]的严格逐元素比较作为标准 CPU 精度基准不依赖浮点容差样例数据由uniform_real_distributionfloat在[0.0, 10.0)区间随机生成加法结果可精确表示因此允许全等比较。样例核心源码位于 Samples/0_Introduction/vector_add/main.asc构建配置位于 Samples/0_Introduction/vector_add/CMakeLists.txt。二、关键特性特性说明流水并行使用TPipeTQueTPosition::VECIN/VECOUT, BUFFER_NUM构建 DoubleBuffer 双缓冲流水隐藏搬入/搬出与计算之间的等待参数可配totalLength向量长度可在运行入口调整Tiling 参数核数、块长、tile 大小由 Host 侧依据硬件资源自动推导精度对比提供标准 CPU 实现作为精度基准逐元素比对并打印执行结果性能分析元信息使用aclprofTensor/aclprofRangePushEx上报算子输入输出张量元信息供 Profiling 工具解析该 API 自 CANN 9.1.0 起提供三、核心实现原理从 Host 到 Kernel 的完整链路3.1 整体结构main.asc采用单文件单算子工程结构包含三层main()main.asc#L223-L238ACL 初始化、设置 Device、创建 Stream、调用run_vector_add(stream, 409600)后清理资源run_vector_add()main.asc#L139-L221Host 侧数据准备、Tiling 计算、Kernel 下发与结果校验add_kernelmain.asc#L85-L137Device 侧 Ascend C Kernel完成双缓冲流水向量加法。int main() { CHECK_ACL(aclInit(nullptr)); int32_t deviceId 0; CHECK_ACL(aclrtSetDevice(deviceId)); aclrtStream stream nullptr; CHECK_ACL(aclrtCreateStream(stream)); int result run_vector_add(stream, 409600); // 默认向量长度 409600 CHECK_ACL(aclrtDestroyStream(stream)); CHECK_ACL(aclrtResetDevice(deviceId)); CHECK_ACL(aclFinalize()); return result; }所有 ACL 调用统一走CHECK_ACL宏main.asc#L30-L38做错误检查失败时打印错误码与文件行号并提前返回。设备内存通过aclrtMalloc分配、并由自定义删除器AclrtFreeDeletermain.asc#L41-L47配合std::unique_ptr自动释放避免异常路径下的内存泄漏。3.2 Tiling 计算多核切分与 Tile 尺寸推导calc_tiling_params()main.asc#L67-L83是整条数据通路的关键它根据硬件信息自动推导三个 Tiling 参数std::tupleint64_t, int64_t, int64_t calc_tiling_params(int64_t totalLength) { constexpr static int64_t MIN_ELEMS_PER_CORE 1024; // 每核最少处理的元素数 constexpr static int64_t BUFFER_NUM 2; // 双缓冲 constexpr static int64_t QUEUE_NUM 3; // 每队列分配 3 份缓冲区空间 auto ascendcPlatform platform_ascendc::PlatformAscendCManager::GetInstance(); uint64_t ubSize; ascendcPlatform-GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize); // 查询 UB 容量 int64_t coreNum ascendcPlatform-GetCoreNumAiv(); // 查询 VectorCore 核数 int64_t numBlocks std::min(coreNum, (totalLength MIN_ELEMS_PER_CORE - 1) / MIN_ELEMS_PER_CORE); numBlocks std::max(numBlocks, static_castint64_t(1)); int64_t blockLength (totalLength numBlocks - 1) / numBlocks; // 每个核处理的块长 int64_t tileSize ubSize / BUFFER_NUM / QUEUE_NUM; // 每个 tile 的字节数 constexpr int64_t ALIGN_SIZE 32; tileSize (tileSize / ALIGN_SIZE) * ALIGN_SIZE; // 32 字节对齐 return std::make_tuple(numBlocks, blockLength, tileSize); }推导逻辑可以概括为三步核数分配numBlocks取min(核心数, ceil(totalLength / 1024))即每个核至少处理 1024 个元素数据量不足以铺满所有核时只启用需要的核数且最少为 1 个核块长计算blockLength按ceil(totalLength / numBlocks)向上取整保证每个核的数据量均衡尾块由 Kernel 内部裁剪Tile 大小tileSize由 UB 总容量除以BUFFER_NUM(2)与QUEUE_NUM(3)得出这是为了给双缓冲的输入队列 X、输入队列 Y、输出队列 Z 各留出独立空间最后按 32 字节对齐对应 Vector 搬运的对齐要求。这里使用的platform_ascendc::PlatformAscendCManager来自 kernel_operator.h 配套的platform/platform_ascendc.h头文件通过GetCoreMemSize(CoreMemType::UB, ...)与GetCoreNumAiv()在运行时查询硬件实际资源使 Tiling 具备跨代际硬件的自适应性——这也是仓库中 CMakeLists.txt 同时声明支持dav-2201Ascend 910B/C与dav-3510Ascend 950的原因之一。3.3 Kernel 实现TPipe/TQue 双缓冲流水add_kernelmain.asc#L85-L137是 Ascend C 编程模型中典型的CopyIn → Compute → CopyOut三段式流水template typename T __global__ __aicore__ __vector__ void add_kernel( GM_ADDR x, GM_ADDR y, GM_ADDR z, int64_t totalLength, int64_t blockLength, uint32_t tileSize) { constexpr static int64_t BUFFER_NUM 2; AscendC::TPipe pipe; AscendC::GlobalTensorT xGm, yGm, zGm; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueX; AscendC::TQueAscendC::TPosition::VECIN, BUFFER_NUM inQueueY; AscendC::TQueAscendC::TPosition::VECOUT, BUFFER_NUM outQueueZ; pipe.InitBuffer(inQueueX, BUFFER_NUM, tileSize); pipe.InitBuffer(inQueueY, BUFFER_NUM, tileSize); pipe.InitBuffer(outQueueZ, BUFFER_NUM, tileSize); xGm.SetGlobalBuffer((__gm__ T *)x blockLength * AscendC::GetBlockIdx()); yGm.SetGlobalBuffer((__gm__ T *)y blockLength * AscendC::GetBlockIdx()); zGm.SetGlobalBuffer((__gm__ T *)z blockLength * AscendC::GetBlockIdx()); // ... 按 tile 循环DataCopyPad 搬入 - AscendC::Add 计算 - DataCopyPad 搬出 }实现要点全局缓冲区偏移通过AscendC::GetBlockIdx()获取当前核 ID将xGm/yGm/zGm的起始地址各自偏移blockLength * blockIdx实现核间数据不重叠的分块访问尾块处理currentBlockLength用totalLength - blockIdx * blockLength与blockLength取较小值处理最后一个核数据不足一个块长的情况tile 层同理tileElementNum取min(剩余元素, elementNumPerTile)对齐搬运AscendC::DataCopyPad配合DataCopyExtParamsblockCount1、blockLentileElementNum * sizeof(T)、srcStride/dstStride0完成 GM↔UB 搬运支持非对齐的尾部数据计算指令核心运算只有一行AscendC::Add(zLocal, xLocal, yLocal, tileElementNum)一条向量指令完成一批元素的加法双缓冲TQue..., 2声明每个队列包含 2 个 bufferAllocTensor/EnQue/DeQue/FreeTensor的交替使用让下一 tile 的 CopyIn 与当前 tile 的 Compute 重叠执行从而隐藏 MTE 搬运延迟。值得对照的是同目录下的 vector_function_addRegBase/VF 编程模型仅支持dav-3510与 vector_add_c_apiC 语言风格 APIasc_simd.h三者功能一致但编程模型不同适合对照学习 MemBase、RegBase 与 C API 三种写法的差异。3.4 Kernel 下发与性能分析元信息上报Kernel 以add_kernelfloatnumBlocks, nullptr, stream(...)的 CUDA 风格语法下发main.asc#L198。下发前后样例通过 ACL Profiling 接口上报了张量元信息aclprofTensor tensors[] { MakeVectorTensor(PROF_INPUT_TENSOR, numElements), // x MakeVectorTensor(PROF_INPUT_TENSOR, numElements), // y MakeVectorTensor(PROF_OUTPUT_TENSOR, numElements), // z }; aclprofTensorInfo tensorInfo { aclprofStr2Id(PROF_OP_NAME), // vector_add aclprofStr2Id(PROF_OP_TYPE), // VectorAdd 0, sizeof(tensors) / sizeof(aclprofTensor), 0, static_castuint32_t(numBlocks), stream, tensors }; aclprofEventAttributes attrs { ACL_PROF_EVENT_ATTR_VERSION, sizeof(aclprofEventAttributes::message), ACL_PROF_MESSAGE_TYPE_TENSOR_INFO, tensorInfo }; aclprofRangePushEx(attrs); // 压栈性能分析区间 add_kernelfloatnumBlocks, nullptr, stream(...); aclprofRangePop(); // 弹栈其中MakeVectorTensormain.asc#L56-L65将张量描述为ACL_FORMAT_NDACL_FLOAT 一维 shape 的元信息。这套aclprofTensor元信息上报 API 由acl/acl_prof.h提供自 CANN 9.1.0 起才可用这也是本样例环境要求必须不低于 CANN 9.1.0 weekly20260708的原因。上报后Profiling 工具即可在性能分析中还原出算子的输入输出 shape相关参考实现在仓库子模块 third_party/asc-devkit 的 examples 目录中。四、参数说明参数含义默认值/取值totalLength向量长度元素个数即每个输入向量的元素数run_vector_add(stream, 409600)中默认为 409600可自行调整Dtype 模板参数Kernel 数据类型支持FLOAT32add_kernelfloat可通过模板实例化扩展从源码结构看Kernel 已通过template typename T实现泛化main.asc#L85tileSize / sizeof(T)与tileElementNum * sizeof(T)均按类型推导因此扩展到其他数据类型只需新增模板实例与对应 Host 侧数据生成逻辑。五、支持架构与环境要求5.1 支持架构NPU ARCHdav-2201Ascend 910B/C、dav-3510Ascend 950架构约束在构建期由 cmake/sample_common.cmake 中的cann_sample_check_arch(dav-2201 dav-3510)宏强制检查当配置的NPU_ARCH不在期望列表时该样例会在 CMake 配置阶段被跳过并打印Skip sample ...: NPU_ARCH... is not supported不会进入构建也不会出现在--target help中。5.2 环境要求CANN 版本需 CANN 9.1.0 weekly20260708及以上。本样例使用了aclprofTensor等性能分析元信息上报 API见 main.asc#L176-L199旧版本缺少该接口将无法编译通过Toolkit 安装须安装完整社区版 CANN Toolkit 并执行source ${install_path}/ascend-toolkit/set_env.sh使环境变量生效详见 README.md 的环境部署一节构建依赖cmake 3.16.0、python 3.8.0、zip、git以及通过pip3 install -r requirements.txt安装的 Python 三方库环境自检建议构建前在仓库根目录运行python3 scripts/check_env.pyscripts/check_env.py脚本只检查不修改任何配置。六、编译与运行6.1 配置构建从仓库根目录初始化构建NPU_ARCH为必填参数# Ascend 950 cmake -S . -B build -DNPU_ARCHdav-3510 # Ascend 910B/C cmake -S . -B build -DNPU_ARCHdav-22016.2 编译 vector_add 目标cmake --build build --target vector_add构建由 Samples/0_Introduction/vector_add/CMakeLists.txt 驱动add_executable(vector_add main.asc)声明目标--npu-arch${NPU_ARCH}与-O3作为 Ascend C 编译选项链接m、dl、platform、tiling_api、msprofiler等运行时库其中msprofiler即支撑aclprofTensor上报的库产物安装到0_Introduction/vector_add目录。6.3 运行样例切换到可执行文件所在目录后直接运行cd ./build/Samples/0_Introduction/vector_add/ ./vector_add成功时打印Vector add completed successfully!若存在精度问题则打印Vector add failed!需要注意的是vector_add默认不接收命令行参数main()中固定传入 409600如需验证不同totalLength可直接修改 main.asc#L231 的入参后重新编译。七、结果验证与自动化测试样例内置了严格的 Host 侧校验将设备端结果h_C拷回后逐一检查h_C[i] h_A[i] h_B[i]任一元素不相等即判定失败并返回退出码 1main.asc#L206-L220。仓库同时提供了对应的 CI 用例tests/ci_functional_test.yaml它模拟全量构建安装后的运行方式- id: vector_add steps: - name: run_default cwd: build_out/0_Introduction/vector_add cmd: [./vector_add] pass_criteria: exit_code: 0 stdout_contains: [Vector add completed successfully!] stdout_not_contains: [Vector add failed!]也就是说通过cmake --build build --parallel全量构建并cmake --install build --prefix ./build_out安装后在build_out/0_Introduction/vector_add/目录执行./vector_add且退出码为 0、标准输出包含成功信息即视为功能测试通过。这也给开发者提供了一条编译产物 → 可执行文件 → 判定标准的完整验证路径。八、同类入门样例横向对照Samples/0_Introduction目录下共有三个向量加法语义相同的入门样例适合放在一起对照学习样例编程模型核心 API支持架构数据流vector_add本文MemBaseTQue/TPipeAscendC::Add操作 UB 上的 LocalTensordav-2201, dav-3510每步运算在 UB 上读写vector_add_c_apiC APIasc_simd.hasc_copy_gm2ub_sync/asc_add_sync/asc_copy_ub2gm_syncdav-2201同步搬入→同步计算→同步搬出vector_function_addRegBaseVector FunctionAscendC::Reg::LoadAlign/Add/StoreAlignUpdateMaskdav-3510Load 进寄存器→寄存器间运算→Store 回 UB其中 vector_add 的双缓冲流水与 C API 版本的单缓冲同步执行、RegBase 版本的寄存器直算形成渐进式对照前者体现以空间换时间的经典流水思想后两者分别体现接口简化与寄存器级数据驻留的优势。概念背景可进一步参考 vector_function_getting_started性能进阶可阅读 2_Performance 下的各调优 story。九、小结vector_add虽小却完整覆盖了 Ascend C 单算子开发的全部关键环节Tiling 三要素核数、块长、tile 大小由硬件资源查询自动推导兼顾多核并行与 UB 容量约束TPipe/TQue 双缓冲流水让搬运与计算重叠是后续所有性能调优样例的基石模式ACL Host 框架提供从初始化、内存管理、Stream 同步到精度校验的完整范式aclprofTensor 元信息上报展示了如何让算子被 Profiling 工具正确识别是 CANN 9.1.0 新特性的实战示例。建议读者在运行通过后尝试修改totalLength观察尾块处理逻辑或对照 vector_add_c_api 与 vector_function_add 体会三种编程模型的差异再进入 matmul 等更复杂的入门样例继续进阶。参考资料cann-samples 根 README环境部署与快速入门vector_add 样例源码vector_add 构建配置入门样例索引架构检查宏实现vector_add CI 功能测试用例asc-devkit 子模块性能分析参考样例所在【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考