企业官网建设流程全解析

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本文以 CANN asc-devkit 仓库中的向量加法入门样例为蓝本，完整剖析 Ascend C 算子开发的两大核心编程范式：基于LocalMemAllocator的静态 Tensor 实现（add）与基于 TPipe/TQue 队列机制的实现（add_tpipe_tque）。通过阅读本文，你将掌握昇腾 AI Core 上"加载—计算—存储"三段式流水结构、GM/UB 两级存储体系、多核按block_idx切分数据的并行方法，以及从编译运行到 printf/DumpTensor 调试、msOpProf 性能分析的完整实战链路。

样例总览：一个加法，两种内存与同步管理机制

向量计算样例基于 Ascend C 实现，用于演示两种不同的内存与同步管理机制。它们在功能上完全等价——都是完成两个同形状向量的逐元素相加，但在"如何管理片上内存、如何保证流水线阶段间同步"上给出了两种典型答案。

两个样例均在仓库 examples/01_simd_cpp_api/00_introduction/01_add 目录下，都通过 8 个 AI Core 并行完成计算，每个核处理 2048 个元素（blockLength = 2048），输入输出均为float类型、形状[8, 2048]、ND 数据排布，元素总数为 8×2048 = 16384 个。

样例一：静态 Tensor 实现（add）

算子功能与运行参数

Add 算子实现两个向量的逐元素相加，计算公式为：

$$ z_i = x_i + y_i $$

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

该样例使用 8 个核完成计算，每个核处理 2048 个元素，合计 16384 个 float 元素。目录结构如下：

├── add │ ├── CMakeLists.txt // 编译工程文件 │ ├── add.asc // Ascend C 样例实现与调用示例 │ └── README.md // 样例文档

三阶段流水结构："加载—计算—存储"

Add 算子的计算逻辑遵循昇腾 AI Core 上经典的"load-compute-store"三段式流水结构：

1. 通过 DataCopy 将输入数据 x、y 从 GM（Global Memory，芯片外部全局内存）搬入 UB（Unified Buffer，向量计算专用片上缓存）；
2. 在 UB 上对 xLocal、yLocal 执行向量加法，结果写入 zLocal；
3. 将计算结果从 UB 搬回 GM。

前置知识：

• GM（Global Memory）：AI Core 外部的全局存储，通过 GlobalTensor 访问，容量大但访问速度慢；
• UB（Unified Buffer）：AI Core 内部向量计算专用片上缓存，通过 LocalTensor 访问，容量有限但访问速度快；
• DataCopy：GM 与 UB 之间数据搬运的接口，搬运方向由参数顺序决定；
• PipeBarrier：流水线同步屏障，确保数据搬运完成后再进行后续操作，避免读写冲突；
• block_idx：内建变量，表示当前核的索引（等价于 GetBlockIdx()），用于多核并行计算中的数据切分。

核心代码逐段剖析

完整实现位于 add.asc，核心 Kernel 代码如下：

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>(); // 确保搬出完成 }

这段代码透露出几个值得注意的实现细节：

• 多核数据切分：每个核通过x + block_idx * blockLength计算自己在 GM 中的起始地址，配合 SetGlobalBuffer 绑定长度为blockLength的连续数据段，实现互不重叠的并行切分；
• UB 内存管理：这里采用LocalMemAllocator<AscendC::Hardware::UB>直接为三个 LocalTensor 分配 UB 空间，这是与 TPipe/TQue 方案最大的区别——内存由分配器统一管理，同步则全部依赖PipeBarrier；
• 同步时机：三次PipeBarrier<PIPE_ALL>()分别保证了"搬入完成 → 计算"、"计算完成 → 搬出"、"搬出完成"三个依赖关系。

调用方式与 Host 侧流程

调用方式为使用 Kernel 调用操作符<<<numBlocks, 0, stream>>>调用核函数，其中numBlocks = 8指定 8 个核并行执行（第二个参数0表示 blockDim 场景，由系统推导）。Host 侧的完整调用链（见 add.asc）遵循标准 ACL 流程：

aclInit(nullptr); aclrtSetDevice(deviceId); aclrtCreateStream(&stream); // aclrtMalloc 申请设备内存 + aclrtMallocHost 申请主机内存 // aclrtMemcpy 将 x、y 从 Host 拷贝到 Device（HOST_TO_DEVICE） add_custom<blockLength><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice); aclrtSynchronizeStream(stream); // aclrtMemcpy 将 z 从 Device 拷回 Host（DEVICE_TO_HOST） // VerifyResult 对比输出与 golden 数据，输出 "test pass!" aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize();

Host 侧通过std::equal将 Kernel 输出与逐元素相加的 golden 结果逐位比对，一致则打印test pass!。main函数中测试数据构造为x[i] = i * 0.1f、y[i] = i * 0.2f，golden 为x[i] + y[i]。

实现流程分析

样例二：TPipe/TQue 队列式实现（add_tpipe_tque）

目录结构与规格

├── add_tpipe_tque │ ├── scripts │ │ ├── gen_data.py // 输入数据与 golden 数据生成脚本 │ │ └── verify_result.py // 输出数据与 golden 数据一致性校验脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读写函数 │ ├── add_tpipe_tque.asc // Ascend C 样例实现（TPipe 与 TQue 管理内存和同步）及样例调用 │ └── README.md // 样例文档

处理流程

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

队列说明：该样例使用TPipe和TQue演示基本的队列式编程方法。EnQue用于将已搬到 UB 的LocalTensor入队，DeQue用于在后续阶段从队列中取出张量继续处理。其中 TPipe 负责统一管理 Device 端内存等资源——一个核函数必须且只能初始化一个 TPipe 对象，它通过InitBuffer为 TQue 和 TBuf 分配内存；TQue 则用于执行队列相关操作、管理相关资源。

核心代码逐段剖析

完整实现位于 add_tpipe_tque.asc：

__global__ __vector__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, uint32_t totalLength) { AscendC::TPipe pipe; AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX; AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueY; AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ; AscendC::GlobalTensor<float> xGm; AscendC::GlobalTensor<float> yGm; AscendC::GlobalTensor<float> zGm; // totalLength 表示全局输入长度；GetBlockNum() 返回本次启动的核数， // 这里按核数均分每个 block 的处理区间。 uint32_t blockLength = totalLength / AscendC::GetBlockNum(); // 根据 block_idx 计算当前核在 GM 中的起始地址，确保各核访问互不重叠的数据分片。 xGm.SetGlobalBuffer((__gm__ float*)x + blockLength * AscendC::GetBlockIdx(), blockLength); yGm.SetGlobalBuffer((__gm__ float*)y + blockLength * AscendC::GetBlockIdx(), blockLength); zGm.SetGlobalBuffer((__gm__ float*)z + blockLength * AscendC::GetBlockIdx(), blockLength); // 为输入和输出队列申请与 blockLength 对应的 UB 缓冲，保持 TPipe/TQue 的编程范式。 pipe.InitBuffer(inQueueX, 1, blockLength * sizeof(float)); pipe.InitBuffer(inQueueY, 1, blockLength * sizeof(float)); pipe.InitBuffer(outQueueZ, 1, blockLength * sizeof(float)); // 使用 DataCopy 将输入从 GM 搬运到 UB，并通过 EnQue 将 LocalTensor 入队，供后续计算阶段 DeQue 取用。 AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>(); AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>(); AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); // DeQue 取出输入张量，在 UB 内执行 Add，并将结果 EnQue 到输出队列，供后续写回 GM。 xLocal = inQueueX.DeQue<float>(); yLocal = inQueueY.DeQue<float>(); AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>(); AscendC::Add(zLocal, xLocal, yLocal, blockLength); outQueueZ.EnQue<float>(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); // 从输出队列取出结果，并写回当前核负责的 GM 分片。 zLocal = outQueueZ.DeQue<float>(); AscendC::DataCopy(zGm, zLocal, blockLength); outQueueZ.FreeTensor(zLocal); }

队列模板参数与深度选择的工程要点

TQue<TPosition::VECIN, 1>中第一个模板参数pos表示队列的逻辑位置（此处 VECIN 表示向量搬入流水、VECOUT 表示向量搬出流水），第二个参数depth表示队列深度。从 TQue 官方文档 可以提炼出以下工程要点：

• 深度语义：队列深度表示该队列可以连续入队/出队的次数。若对同一队列连续 n 次EnQue（中间没有DeQue），深度需设为 n；
• 深度与 double buffer 无关：队列机制用于实现流水线并行，double buffer 是在此基础上进一步提高流水线利用率。即使队列深度为 1，仍可开启 double buffer；
• 推荐深度为 1：非 Tensor 原地操作场景下，深度设为 1 时编译器对该场景做了特殊优化，性能通常更好，推荐设置为 1；Tensor 原地操作场景下需设为 0。样例中每个队列每次仅入队一次、出队一次，因此深度取 1，与文档推荐一致。

两种实现的对比与选型参考

两种实现的计算结果完全一致，选型取决于场景：学习流水线与同步机制选静态实现；工程化算子开发建议从 TPipe/TQue 范式起步，它与昇腾 算子开发编程指南 中的流水设计思路一脉相承。

编译与运行

环境变量配置

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

source ${install_path}/cann/set_env.sh

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

静态实现（add）的编译执行

在样例目录下依次执行：

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。

队列实现（add_tpipe_tque）的编译执行

该样例额外包含数据生成与结果校验环节，需要 Python3 与 numpy 环境：

mkdir -p build && cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程（默认 npu 模式） python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译产物运行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 校验输出结果是否正确

其中 gen_data.py 使用np.random.uniform(1, 10, [8, 2048])生成随机输入，并将逐元素相加的 golden 写入output/golden.bin；verify_result.py 采用np.isclose做容差比对（相对容差 1e-4、绝对容差 1e-5），错误元素占比不超过 1e-4 即判定通过，并打印test pass!。输入/输出数据通过 data_utils.h 中的ReadFile/WriteFile以二进制格式读写。

编译选项说明

从两个样例的 CMakeLists.txt 可以看到，工程通过find_package(ASC REQUIRED)引入 Ascend C 工具链，以project(kernel_samples LANGUAGES ASC CXX)声明 ASC 语言，add_executable(demo add.asc)将.asc源文件直接编译为可执行文件，并通过--npu-arch=${CMAKE_ASC_ARCHITECTURES}指定目标 NPU 架构。

运行结果

执行完成后输出如下，表示精度比对成功：

test pass!

功能调试：printf 与 DumpTensor

printf 格式化输出

printf 接口提供 CPU/NPU 域调试场景的格式化输出功能，在算子核侧实现代码中需要输出日志信息的位置调用即可：

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

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

DumpTensor 张量内容导出

对于基于算子工程开发的算子，DumpTensor 接口可用于导出指定 LocalTensor 的内容，并支持打印自定义附加信息（仅支持 uint32_t 类型信息，如打印当前行号）。在算子核侧实现代码中需要打印 Tensor 数据的位置调用，例如：

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

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

值得一提的是，add.asc 中预置了一段被#if 0屏蔽的调试代码，开启后可在计算完成后依次导出 xLocal、yLocal、zLocal 的前 32 个元素，是学习该接口用法的现成范例。

性能调试：msOpProf 工具

msOpProf 是单算子性能分析工具，提供msopprof与msopprof simulator两种使用模式，可帮助用户识别算子在内存、代码和指令方面的异常，用于全面的算子调优。它支持对不同运行模式（真机或仿真）和文件类型（可执行程序或算子二进制.o文件）进行性能数据采集与自动解析。

真机性能采集直接测量算子在昇腾 AI 处理器上的执行时间，适合在真机环境快速定位算子性能问题。对demo可执行程序执行算子调优：

msopprof ./demo

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

├──dump # 原始性能数据，用户无需查看 ├──ArithmeticUtilization.csv # Cube/Vector 指令周期占比 ├──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

msOpProf 的完整使用说明（含仿真模式、.o文件分析等进阶用法）收录在算子调优（msOpProf）对应的官方用户指南中，可结合 CANN 工具链文档查阅。

可优化方向与进阶路径

上述方向中，"搬运粒度"还与 DataCopy 的底层约束直接相关：count * sizeof(T)需要 32 字节对齐，若未对齐，搬运量会向下取整到 32 字节对齐。样例中每次搬运 2048 个 float（8192 字节），恰好满足 32 字节对齐要求，这也是切分粒度设计时需要始终牢记的硬性约束。

完整的性能调优过程可参考 Add 高性能调优样例，该样例是入门版本地到双缓冲流水、多核动态分配等优化手法的进阶续篇。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend