news 2026/9/18 18:56:52

CANN Ascend C Add 向量加法算子入门实战：静态 Tensor 编程范式与多核流水实现解析

作者头像

张小明

前端开发工程师

1.2k 24
文章封面图
CANN Ascend C Add 向量加法算子入门实战：静态 Tensor 编程范式与多核流水实现解析

CANN Ascend C Add 向量加法算子入门实战：静态 Tensor 编程范式与多核流水实现解析

【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples

导读

本文围绕 CANN 开源仓库 cann-samples 中 Add 向量加法入门样例 展开，完整讲解基于 Ascend C SIMD C++ API 的静态 Tensor 编程方式：如何用"搬入—计算—搬出"三段式流水结构在 AI Core 上实现两个向量的逐元素加法，如何通过 8 个核并行分片提升吞吐，以及如何编译、运行、调试与性能剖析该算子。读完本文，你将掌握 GlobalTensor/LocalTensor/DataCopy/PipeBarrier 等核心编程要素的用法，并具备将同一套模板迁移到其他逐元素（Element-wise）算子的能力。

样例概述：做什么、怎么做

本样例位于 Samples/0_Introduction/01_simd_cpp_api/01_add/add，它演示了 Ascend C 向量加法的基本用法。算子的数学定义是逐元素加法：

$$z_i = x_i + y_i$$

  • x：输入张量，形状为[8, 2048]，数据类型为 float，数据排布格式为 ND；
  • y：输入张量，形状为[8, 2048]，数据类型为 float，数据排布格式为 ND；
  • z：输出张量，形状为[8, 2048]，数据类型为 float，数据排布格式为 ND。

样例运行参数：本样例使用 8 个核完成计算，每个核处理 2048 个元素（blockLength = 2048），数据总量为 8×2048 = 16384 个 float 元素。核函数启动时通过内核调用符<<<numBlocks, 0, stream>>>指定numBlocks = 8，即同时拉起 8 个 block（8 个核）并行执行。

支持的产品与 CANN 软件版本

产品CANN 软件版本
Ascend 950PR/Ascend 950DT>= CANN 9.1.0
Atlas A3 训练系列产品/Atlas A3 推理系列产品>= CANN 9.0.0
Atlas A2 训练系列产品/Atlas A2 推理系列产品>= CANN 9.0.0

目录结构

Samples/0_Introduction/01_simd_cpp_api/01_add/add ├── CMakeLists.txt // 编译工程文件 ├── add.asc // Ascend C 样例实现 & 调用样例 └── README.md // 样例说明文档

其中 add.asc 是单文件完整实现，同时包含核函数（kernel 侧）与宿主侧（host 侧）的调用、数据构造和精度校验逻辑，是理解"核函数怎么写、怎么调、怎么验"的最小闭环。

三个核心存储/同步概念：GM、UB、DataCopy、PipeBarrier

在阅读核函数代码之前，先建立四个基础概念（这也是后续所有 SIMD 向量算子通用的编程要素）：

  • GM（Global Memory）：AI Core 外部的全局存储，容量大但访问速度慢，通过GlobalTensor访问。它是算子的输入输出数据在设备侧的落脚点。
  • UB（Unified Buffer）：AI Core 内部的向量计算专用片上缓存，容量有限但访问速度快，通过LocalTensor访问。向量计算单元只能读取 UB 上的数据，因此输入数据必须先搬到 UB 才能参与计算。
  • DataCopy：在 GM 与 UB 之间搬运数据的 API，搬运方向由参数顺序决定：DataCopy(local, global, len)是 GM→UB 搬入，DataCopy(global, local, len)是 UB→GM 搬出。
  • PipeBarrier：流水线同步屏障，用于保证"数据搬运完成后再执行后续操作"，避免不同硬件流水（MTE 搬运单元与 Vector 计算单元）之间的读写冲突。

此外还有一个贯穿多核编程的内建变量：block_idx，它表示当前核的编号（等价于GetBlockIdx()），用于多核并行时的数据分片计算。

核函数实现：静态 Tensor 三段式流水

Add 算子的计算逻辑严格遵循"搬入—计算—搬出"三段式流水结构：

  1. 将输入数据 x 和 y 从 GM 搬运到 UB；
  2. 在 UB 上对xLocal、yLocal执行向量加法操作，结果存入zLocal；
  3. 将计算结果从 UB 搬运回 GM。

核心代码如下（与仓库中 add.asc 一致）：

template <uint32_t blockLength> __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { AscendC::InitSocState(); // Global Tensor：在GM上分配输入/输出缓冲区 AscendC::GlobalTensor<float> xGm, yGm, zGm; xGm.SetGlobalBuffer(x + block_idx * blockLength, blockLength); // 每个核按block_idx偏移处理各自的数据段 yGm.SetGlobalBuffer(y + block_idx * blockLength, blockLength); zGm.SetGlobalBuffer(z + block_idx * blockLength, blockLength); // Local Tensor：在UB上分配计算缓冲区 AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator; AscendC::LocalTensor<float> xLocal = ubAllocator.Alloc<float, blockLength>(); AscendC::LocalTensor<float> yLocal = ubAllocator.Alloc<float, blockLength>(); AscendC::LocalTensor<float> zLocal = ubAllocator.Alloc<float, blockLength>(); // GM -> UB: 搬入输入数据 AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); AscendC::PipeBarrier<PIPE_ALL>(); // 确保搬入完成后才进行计算 // 向量计算: z = x + y AscendC::Add(zLocal, xLocal, yLocal, blockLength); AscendC::PipeBarrier<PIPE_ALL>(); // 确保计算完成后才搬出 // UB -> GM: 搬出计算结果 AscendC::DataCopy(zGm, zLocal, blockLength); AscendC::PipeBarrier<PIPE_ALL>(); // 确保搬出完成 }

代码中的编程要素可以逐行拆解：

  • __vector__ __global__：声明这是一个在向量计算单元上执行、可从宿主侧启动的核函数；参数用__gm__ float*标记为 GM 地址空间。
  • InitSocState()：初始化 AI Core 硬件状态，为后续操作做准备，是所有核函数的第一行。
  • SetGlobalBuffer(addr, len)：将GlobalTensor绑定到 GM 上的某段连续内存。x + block_idx * blockLength的写法让每个核从自己的数据段起点开始访问，实现数据分片。
  • LocalMemAllocator<AscendC::Hardware::UB>：静态 Tensor 编程模式下在 UB 上申请内存的分配器，Alloc<float, blockLength>()为 float 类型的blockLength个元素分配一块连续空间。x、y、z 各分配一块，互不重叠。
  • AscendC::Add(dst, src0, src1, len)：向量加法指令，在 UB 上并行计算dst[i] = src0[i] + src1[i]。
  • 三处PipeBarrier<PIPE_ALL>()分别隔离"搬入→计算""计算→搬出""搬出→结束"，保证数据就绪后再被消费。

宿主侧调用与精度校验（源码级补充）

核函数本身不负责数据准备与结果验证，这些逻辑都在 add.asc 的宿主侧完成，整体链路为：

  1. 初始化运行环境：aclInit(nullptr)初始化 ACL 运行时，aclrtSetDevice(deviceId)选择设备（样例固定使用 device 0），aclrtCreateStream(&stream)创建执行流。
  2. 分配设备内存：通过aclrtMalloc为 x、y、z 各分配totalLength * sizeof(float)的设备侧内存（分配策略为ACL_MEM_MALLOC_HUGE_FIRST），再通过aclrtMallocHost分配一块宿主侧内存用于回拷结果。
  3. 数据上板：aclrtMemcpy把宿主侧的 x、y 数据以ACL_MEMCPY_HOST_TO_DEVICE方向拷入设备内存。
  4. 启动核函数：add_custom<blockLength><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice)，其中模板参数blockLength = 2048在编译期确定，运行时参数为三个设备地址；numBlocks = 8指定 8 个核并行。
  5. 同步与回拷：aclrtSynchronizeStream(stream)等待核函数执行完成，再用ACL_MEMCPY_DEVICE_TO_HOST将结果 z 拷回宿主。
  6. 资源释放：依次aclrtFree设备内存、aclrtFreeHost宿主内存、aclrtDestroyStream销毁流、aclrtResetDevice复位设备、aclFinalize结束 ACL 运行时。
  7. 精度校验：宿主侧用 CPU 直接计算 golden 结果golden[i] = x[i] + y[i]，再通过VerifyResult逐元素比对（std::equal）。比对通过打印test pass!并返回 0，否则打印test failed!返回 1。样例同时在标准输出打印输出与真值的前 20 个元素，便于人工核对。

main 函数中的数据生成方式为：x[i] = i * 0.1f、y[i] = i * 0.2f，共8 * 2048个元素。

实现流程解析：每个阶段在做什么、为什么

下表把核函数的执行过程按阶段拆解，明确每个动作的数据流动与设计意图：

阶段数据流动/行为实现目的/原因
初始化InitSocState()初始化 AI Core 硬件状态，为后续操作做准备
GM 地址分配SetGlobalBuffer(x + block_idx * blockLength, blockLength)每个核根据block_idx计算偏移量，处理不同的数据段，实现多核并行
UB 空间分配ubAllocator.Alloc<float, blockLength>()在 UB 上为 x、y、z 各分配一块连续内存，供向量计算使用
搬入（Stage 1）GM → UB：DataCopy(xLocal, xGm)、DataCopy(yLocal, yGm)将输入数据从 GM 搬运到 UB，因为向量计算单元只能访问 UB 上的数据
流水同步PipeBarrier<PIPE_ALL>()确保搬入完成后再开始计算，避免计算单元读取到未就绪的数据
计算（Stage 2）UB 上计算：Add(zLocal, xLocal, yLocal)在 UB 上执行向量加法，利用向量单元并行处理多个元素
流水同步PipeBarrier<PIPE_ALL>()确保计算完成后再开始搬出，避免搬出未完成的结果
搬出（Stage 3）UB → GM：DataCopy(zGm, zLocal)将计算结果从 UB 搬运回 GM，供后续使用或输出
流水同步PipeBarrier<PIPE_ALL>()确保搬出完成，保证数据一致性

可以看到：三段式结构中穿插的三次同步并不是冗余，而是分别守护"数据就绪""结果就绪""写回完成"三个关键时序点，这是保证多流水（MTE2 搬入、Vector 计算、MTE3 搬出）并发安全的最小同步骨架。

可优化方向分析：从能跑到跑得快

本样例是教学用的基础实现，刻意省略了性能优化。文档明确列出了四个可优化方向，理解它们有助于建立 Ascend C 性能优化的直觉：

序号可优化方向当前实现的问题预期优化收益
1多核动态分配固定使用 8 个核，未根据实际可用核数动态分配动态获取可用核数，充分利用多核并行能力，减少端到端耗时
2增大搬运粒度每次搬运 2048 个 float 元素（8KB），搬运粒度较小增大单次搬运数据量，减少搬运次数，摊薄启动开销，提升带宽利用率
3双缓冲流水线并行搬入、计算、搬出三个阶段严格串行执行，各硬件单元（MTE2/V/MTE3）无法同时工作采用 Ping-Pong 双缓冲机制，使搬入、计算、搬出可并行执行，隐藏搬运延迟
4L2 Cache bypassAdd 输入数据只读取一次，但默认经过 L2 Cache，增加了 Cache 污染对流式访问数据设置 L2 Cache bypass，减少不必要的 Cache 开销，提升搬运效率

这四个方向分别对应"核数利用率""搬运效率""流水并行度""缓存策略"四类通用优化手法。本仓库的 add_tpipe_tque 样例即给出了另一个维度的演进——用 TPipe/TQue 队列机制替代手写PipeBarrier，由框架自动管理内存分配与流水同步；而在 1_Features 与 2_Performance 中，还能找到双缓冲、多核动态分配等优化手法的完整落地案例（例如 softmax_regbase_story、rms_norm_quant_story 等演进式样例）。

功能调试：printf 与 DumpTensor

printf

printf接口提供 CPU 域 / NPU 域调试场景下的格式化输出功能。在算子 kernel 侧需要输出日志的位置直接调用即可，例如：

AscendC::printf("add blockIdx=%d\n", AscendC::GetBlockIdx());

注意：printf（PRINTF）接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP=0的方式关闭打印功能。

DumpTensor

DumpTensor用于 Dump 指定LocalTensor的内容，同时支持打印自定义的附加信息（仅支持uint32_t数据类型的信息，比如打印当前行号）。调用方式：

// 向量计算: z = x + y AscendC::Add(zLocal, xLocal, yLocal, blockLength); AscendC::DumpTensor(zLocal, 1, 32);

仓库的 add.asc 中预留了完整的调试代码示例，默认以#if 0关闭，开发者把#if 0改为#if 1即可依次 Dump xLocal、yLocal、zLocal 的内容（每段打印 32 个元素并附带自定义标识）：

#if 0 // Debug I/O. Set #if 1 to enable print. AscendC::printf("%s\n", "[ DumpTensor in xLocal]"); AscendC::DumpTensor(xLocal, 1, 32); AscendC::printf("%s\n", "[ DumpTensor in yLocal]"); AscendC::DumpTensor(yLocal, 2, 32); AscendC::printf("%s\n", "[ DumpTensor in zLocal]"); AscendC::DumpTensor(zLocal, 3, 32); #endif

注意：DumpTensor 接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP=0来关闭打印功能。

性能调试：msOpProf 单算子性能分析

msOpProf 是单算子性能分析工具，包含msopprof和msopprof simulator两种使用方式。该工具协助用户定位算子内存、算子代码以及算子指令的异常，实现全方位的算子调优，支持基于不同运行模式（上板或仿真）和不同文件形式（可执行文件或算子二进制 .o 文件）进行性能数据的采集和自动解析。

上板性能采集

上板性能采集可以直接测定算子在实际昇腾 AI 处理器上的运行时间，适合在板环境中快速定位算子性能问题。基于可执行文件 demo 执行：

msopprof ./demo

命令完成后，会在默认目录下生成以OPPROF_{timestamp}_XXX命名的文件夹，性能数据文件夹结构示例如下：

├──dump # 原始的性能数据，用户无需关注 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比 ├──L2Cache.csv # L2 Cache命中率，影响MTE2，建议合理规划数据搬运逻辑，增加命中率 ├──Memory.csv # UB，L1和主存储器读写带宽速率 ├──MemoryL0.csv # L0A，L0B，和L0C读写带宽速率 ├──MemoryUB.csv # Vector和Scalar到UB的读写带宽速率 ├──OpBasicInfo.csv # 算子基础信息 ├──PipeUtilization.csv # 采集计算单元和搬运单元耗时和占比 ├──ResourceConflictRatio.csv # UB上的bank group、bank conflict和资源冲突率在所有指令中的占比 └──visualize_data.bin # MindStudio Insight呈现文件

查看具体的性能分析结果：

# 查看Task Duration 以及各项数据 cat ./OPPROF_*/PipeUtilization.csv

对 Add 这样的向量算子而言，重点通常落在PipeUtilization.csv（看 Vector 计算单元与 MTE2/MTE3 搬运单元的耗时占比，判断是否存在搬运瓶颈）与L2Cache.csv（评估上一节提到的 L2 Cache bypass 优化空间）上。

编译运行全流程

配置环境变量

请根据当前环境上 CANN 开发套件包的安装方式配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

从仓库的 cmake/ascend.cmake 可以看出，编译系统会优先读取ASCEND_HOME_PATH环境变量定位工具链；未设置时，root 用户依次探测/usr/local/Ascend/ascend-toolkit/latest、/usr/local/Ascend/latest，非 root 用户探测$HOME/Ascend/ascend-toolkit/latest、$HOME/Ascend/latest，均找不到则报错要求显式设置。工具链编译器为${ASCEND_DIR}/${SYSTEM_PREFIX}/ccec_compiler/bin/bisheng，编译 ASC 语言工程时必须依赖该工具链。

编译与执行

在本样例目录下执行如下命令：

mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式 ./demo # 执行样例

使用 CPU 调试或 NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数即可，示例如：

cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # CPU调试模式 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

编译选项说明

选项可选值说明
CMAKE_ASC_RUN_MODEnpu（默认）、cpu、sim运行模式：NPU 运行、CPU 调试、NPU 仿真
CMAKE_ASC_ARCHITECTURESdav-2201（默认）、dav-3510NPU 架构：dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品，dav-3510 对应 Ascend 950PR/Ascend 950DT

架构选项在样例的 CMakeLists.txt 中通过find_package(ASC)引入 ASC 语言编译支持，并用target_compile_options将--npu-arch=${CMAKE_ASC_ARCHITECTURES}传给 ASC 编译器；仓库根目录的 CMakeLists.txt 进一步校验NPU_ARCH只能是dav-3510或dav-2201二者之一。因此，编译前请先确认目标设备的架构代号，选择与设备匹配的值。

执行结果

执行结果如下，说明精度对比成功：

test pass!

延伸：与 TPipe/TQue 队列式实现的对比

同目录下的兄弟样例 add_tpipe_tque 用完全相同的算法（z = x + y、形状[8, 2048]、float、ND）演示了另一种编程范式——基于TPipe和TQue的内存与同步管理机制。其核函数流程为：

  1. add_custom作为核入口接收totalLength；
  2. 通过GetBlockNum()计算当前 block 的数据长度，通过GetBlockIdx()计算当前核在 GM 中对应的数据起点；
  3. 用DataCopy把输入数据从 GM 搬到 UB，并通过EnQue将输入LocalTensor放入输入队列；
  4. 通过DeQue从输入队列取出输入张量，在 UB 中执行Add，再通过EnQue将结果LocalTensor放入输出队列；
  5. 通过DeQue从输出队列取出结果，并使用DataCopy写回当前核负责的 GM 分片。

两种范式对比可以直观看到静态 Tensor 与队列式编程的差异：本样例（静态 Tensor）手动调用LocalMemAllocator分配 UB 空间、手动插入PipeBarrier做同步，代码路径直观、适合理解底层流水；而 TPipe/TQue 版本把内存分配与同步交给框架管理，为后续引入多 buffer 流水并行（Ping-Pong 双缓冲）铺平了道路——这也是 可优化方向分析 中"双缓冲流水线并行"方向在框架层面的落地基础。

两个样例的工程结构也体现了从"单文件演示"到"脚本化验证"的演进：本样例把数据生成、真值计算内联在 main 函数中并直接打印test pass!；而 add_tpipe_tque 则拆分为scripts/gen_data.py（生成输入与 golden 数据）与scripts/verify_result.py（比对输出），更贴近算子工程的实际开发流程。

小结

通过 Add 入门样例，可以建立起 Ascend C 向量算子开发的最小知识闭环：

  • 核侧：InitSocState()→GlobalTensor/LocalTensor绑定与分配 →DataCopy搬入 →Add计算 →DataCopy搬出，配合PipeBarrier守护时序；
  • 宿主侧：ACL 初始化 → 设备内存分配与数据上板 →<<<>>>启动核函数 → 同步回拷 → 逐元素精度校验；
  • 调试与调优：printf/DumpTensor 做功能定位，msOpProf 做性能剖析，再沿多核动态分配、搬运粒度、双缓冲、L2 bypass 四个方向迭代优化。

掌握了这套模式，向 LeakyReLU、Gelu 等更复杂的逐元素算子迁移时，只需替换核函数中的向量计算指令与数据搬运细节，整体骨架可以完全复用。

【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples

创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

版权声明: 本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若内容造成侵权/违法违规/事实不符，请联系邮箱：809451989@qq.com进行投诉反馈，一经查实，立即删除！
网站建设 2026/9/18 18:56:43

IDEA代码提示慢？内存、索引、插件三管齐下，补全延迟压到50ms

你是不是也有过这种体验&#xff1a;项目打开以后&#xff0c;IDEA 底部一直显示 Indexing…&#xff0c;代码高亮正常&#xff0c;但敲代码的时候键盘按下去&#xff0c;补全列表要过一秒才弹出来。遇到大一点的接口&#xff0c;联想半天&#xff0c;偶尔连类名都提示不出来&a…

作者头像 李华
网站建设 2026/9/18 18:53:03

定压功放与定阻功放的区别、混接危害及广播系统配置排查指南

/* MD / 富文本中的 .toc（含博客园搬家等嵌套结构）；.toc-box 在侧栏，不受影响 */#content_views .toc,/* 编辑器常在目录前后插入空 p（:empty 仍占 20px），一并去掉避免顶空隙 */#content_views.markdown_views > p:empty:has(+ .toc),#content_views.markdown_views …

作者头像 李华
网站建设 2026/9/18 18:50:58

文件包含+任意文件上传组合链：从LFI到RCE应急加固

/* MD / 富文本中的 .toc（含博客园搬家等嵌套结构）；.toc-box 在侧栏，不受影响 */#content_views .toc,/* 编辑器常在目录前后插入空 p（:empty 仍占 20px），一并去掉避免顶空隙 */#content_views.markdown_views > p:empty:has(+ .toc),#content_views.markdown_views …

作者头像 李华
网站建设 2026/9/18 18:50:50

React Native热更新核心技术解析与实践

1. 为什么需要热更新能力在移动应用开发领域&#xff0c;传统发版模式存在几个致命痛点。每次功能迭代或问题修复都需要走完整的应用商店审核流程&#xff0c;iOS平台平均审核周期长达24-48小时&#xff0c;紧急情况下这个时间成本完全不可接受。更糟的是&#xff0c;用户设备上…

作者头像 李华