前言

在昇腾NPU上做AI应用开发，离不开CANN。模型要真正跑到芯片上，需要经过编译、调试、性能分析三个环节。CANN生态里负责这些工作的正是asc-tools这个仓库。它把Ascend编译器、算子调试器和性能分析工具打包在一起，目标就是让开发者能把PyTorch或MindSpore的模型，或者自己写的Ascend C算子，顺滑地转换成可以在昇腾AI处理器上执行的二进制。如果你刚拿到一台Atlas服务器，面对一堆ascend-toolkit目录不知道该点哪个命令，这篇文章就是为你写的。我们不讲抽象架构，只讲一个真实流程：从一段简单的算子代码开始，用asc-tools完成编译、上板调试、抓取性能数据，再依据数据回改代码。这个流程本身并不复杂，但每个环节都有容易踩的小坑，比如编译选项里漏了芯片型号、调试时数据搬移方向写反、性能数据里Cube单元和Vector单元分不清。我会把这些坑也放在流程里一起讲，让你读完就能直接上手。

asc-tools在CANN生态里扮演什么角色

CANN是昇腾异构计算架构，内部可以分成五层。最上面是AscendCL应用开发层，再往下是算子库和调优引擎，第三层就是计算编译层。asc-tools属于这一层，全称叫Ascend Compiler Tools。它上面接框架，下面接Runtime，主要干的活是把高层的模型描述或者算子源码，翻译成昇腾达芬奇架构能理解的指令序列。这个工具集里不是只有一个命令，而是一组互相配合的组件。编译器负责把算子或子图变成.o或.air文件，调试器负责在运行时看中间结果，性能分析器负责把芯片上的Cube、Vector、Scalar单元的占用情况拉下来。对于普通开发者来说，最常用的是三类入口：模型编译、算子编译、性能分析。模型编译通常把ONNX或PyTorch脚本转成离线模型，算子编译则处理Ascend C或自定义算子，性能分析覆盖整个推理或训练过程。这三类入口在asc-tools仓库里有对应的脚本和工具链，打开AtomGit上的asc-tools页面，能看到bin目录、scripts目录、docs目录、tests目录。bin下面放的是可执行命令和驱动脚本，scripts下面放的是批量编译或回归测试脚本，docs下面是使用说明，tests下面是功能测试用例。这个目录结构本身就说明，asc-tools不是单个二进制，而是一个围绕编译和调试构建的工程化套件。

准备一段可以在昇腾NPU上跑的算子代码

为了把asc-tools跑起来，我们得先有一个目标。这里选一个简单的向量加法算子，因为它逻辑清晰，容易验证，但又能暴露编译和调试里的真实问题。代码用Ascend C写，文件名叫add_custom.cpp。Ascend C是CANN提供的算子编程语言，语法接近C++，但里面有一些昇腾特有的关键字和API，比如GM（Global Memory）、UB（Unified Buffer）、DataCopy、Pipe等。一个Ascend C算子通常包含三个部分：输入输出张量定义、计算核函数、算子类封装。输入输出张量定义在GM上，核函数从GM把数据搬入UB，在UB里做计算，再把结果搬回GM。这个搬入搬出的过程是昇腾NPU编程里最容易出错的地方，因为芯片上有多个内存层级，数据在哪一层、怎么搬、搬多少，都会直接影响性能和正确性。

下面这段代码演示了一个最基础的向量加法核函数。代码里用x和y表示两个输入，out表示输出，长度用n表示。核函数先把x和y从GM复制到UB，然后在UB里逐元素相加，再把结果写回GM。这里故意用了比较原始的写法，没有做任何融合或向量化，目的就是为了先跑通编译和调试流程，后面再优化。

#include "ascendc.h"

class AddKernel {
public:
    __aicore__ inline AddKernel() {}

    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR out, uint32_t n) {
        this->n = n;
        xGm.SetGlobalBuffer((__gm__ half*)x);
        yGm.SetGlobalBuffer((__gm__ half*)y);
        outGm.SetGlobalBuffer((__gm__ half*)out);
        pipe.InitBuffer(inQueueX, 1, 256 * sizeof(half));
        pipe.InitBuffer(inQueueY, 1, 256 * sizeof(half));
        pipe.InitBuffer(outQueue, 1, 256 * sizeof(half));
    }

    __aicore__ inline void Process() {
        uint32_t loopCount = n / 256;
        for (uint32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(uint32_t i) {
        LocalTensor<half> x = inQueueX.AllocTensor<half>();
        LocalTensor<half> y = inQueueY.AllocTensor<half>();
        DataCopy(x, xGm[i * 256], 256);
        DataCopy(y, yGm[i * 256], 256);
        inQueueX.EnQue(x);
        inQueueY.EnQue(y);
    }

    __aicore__ inline void Compute(uint32_t i) {
        LocalTensor<half> x = inQueueX.DeQue<half>();
        LocalTensor<half> y = inQueueY.DeQue<half>();
        LocalTensor<half> out = outQueue.AllocTensor<half>();
        Add(out, x, y, 256);
        outQueue.EnQue(out);
        inQueueX.FreeTensor(x);
        inQueueY.FreeTensor(y);
    }

    __aicore__ inline void CopyOut(uint32_t i) {
        LocalTensor<half> out = outQueue.DeQue<half>();
        DataCopy(outGm[i * 256], out, 256);
        outQueue.FreeTensor(out);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, 1> inQueueX;
    TQue<QuePosition::VECIN, 1> inQueueY;
    TQue<QuePosition::VECOUT, 1> outQueue;
    GlobalTensor<half> xGm, yGm, outGm;
    uint32_t n;
};

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR out, uint32_t n) {
    AddKernel op;
    op.Init(x, y, out, n);
    op.Process();
}

这段代码写完之后，不能直接跑，因为Ascend C代码需要经过编译器翻译成达芬奇架构的核函数二进制。编译的过程会检查语法、数据类型、内存对齐、流水线依赖等。如果在GM地址类型上写错了，或者UB队列深度配得不合理，编译器会报错。比如上面的inQueueX深度是1，这意味着一次只缓存一个tile。如果后面改成双缓冲，深度就要改成2，否则没法掩盖搬入延迟。这些细节我们会在性能分析环节再回头看。

用asc-tools把算子编译成可执行文件

写好算子代码之后，接着进入编译阶段。asc-tools的编译入口通常是一个顶层驱动脚本，它会根据你指定的芯片型号、算子类型、输出格式，调用下面的编译器、链接器、打包器。在昇腾NPU上，不同芯片的指令集和内存布局有差异，所以编译时必须指定目标架构。常用的选项包括--framework指定框架类型，--soc_version指定芯片版本，--output指定输出目录。这里我们不罗列全部参数，只演示一个典型的自定义算子编译调用。实际使用时，你可以根据asc-tools的help信息或者仓库docs里的说明，找到自己需要的选项组合。

假设当前目录有add_custom.cpp和同目录的CMake配置，编译命令大致如下。这个命令把Ascend C源码编译成.o目标文件，再链接成算子库.so，同时生成算子信息文件.json供上层框架注册算子。

bash asc-tools/bin/op_compiler.sh \
  --kernel_name=add_custom \
  --source_path=./add_custom.cpp \
  --soc_version=Ascend910 \
  --output_path=./output \
  --compile_type=op

执行这条命令后，会在output目录下看到几个文件。add_custom.o是编译后的核函数目标文件，add_custom.so是打包后的算子库，add_custom.json是算子注册信息。如果这一步报错，最常见的原因有几个：soc_version写错了，比如把Ascend910写成Ascend910B；source_path里文件路径不对；kernel_name和函数名不一致；CMake配置里找不到Ascend C的头文件路径。遇到这种情况，不要急着改代码，先检查asc-tools的日志输出。日志里通常会打印出调用的子命令和报错位置，定位到具体哪一行后再回头改算子代码。这个排查过程比在CPU上编译复杂一点，因为asc-tools的编译链路比较长，从C++前端到中间表示，再到达芬奇指令，中间任何一层出问题都会以不同的错误形式抛出来。多看几遍日志，比反复重试更有效。

上板调试：确认算子在昇腾NPU上结果正确

编译通过后，算子还只是半成品，必须拿到真实的昇腾NPU设备上跑一遍。调试环节主要验证两件事：结果对不对，性能有没有明显异常。结果不对通常是因为数据搬移方向反了、数据类型转换错了、边界处理漏了。性能异常通常是因为流水打不满、内存搬移太多、或者核函数启动次数过多。asc-tools的调试工具可以帮我们抓取中间张量的值，也可以把算子运行时的硬件事件记录下来。对于向量加法这种简单算子，最常见的调试方式是构造一组输入，对比昇腾NPU输出和CPU参考实现的结果。如果两者不一致，再逐步缩小范围，看是哪一次DataCopy或哪一次Add出了问题。

下面这段代码演示了一个典型的上板调试脚本。它构造两个half类型的输入张量，调用刚刚编译出来的add_custom算子，然后把结果拉回主机内存，和NumPy的参考结果做对比。脚本里的device和host的区分很重要，因为昇腾NPU有自己的Device内存，数据不在同一个地址空间里。

import numpy as np
import acl

def test_add_custom():
    acl.init()
    device_id = 0
    acl.rt.set_device(device_id)

    n = 1024
    x = np.ones(n, dtype=np.float16)
    y = np.ones(n, dtype=np.float16)
    golden = x + y

    x_dev, _ = acl.rt.malloc(n * 2, 0)
    y_dev, _ = acl.rt.malloc(n * 2, 0)
    out_dev, _ = acl.rt.malloc(n * 2, 0)

    acl.rt.memcpy(x_dev, x.nbytes, x, x.nbytes, 1)
    acl.rt.memcpy(y_dev, y.nbytes, y, y.nbytes, 1)

    add_custom(x_dev, y_dev, out_dev, n)

    out = np.zeros(n, dtype=np.float16)
    acl.rt.memcpy(out, out.nbytes, out_dev, out.nbytes, 2)

    if np.allclose(out, golden, atol=1e-3):
        print("check passed")
    else:
        print("check failed")

    acl.rt.free(x_dev)
    acl.rt.free(y_dev)
    acl.rt.free(out_dev)
    acl.rt.reset_device(device_id)
    acl.finalize()

这个脚本跑通之后，说明算子逻辑是对的。但在真实场景里，输入往往不止1024个元素，形状也不是一维的。当数据量变大时，边界处理就变得重要。比如上面的核函数假设n是256的整数倍，如果实际长度不是整数倍，末尾的tile会读越界。解决这个问题的办法是在核函数里加一个tail处理分支，把余数部分单独CopyIn、Compute、CopyOut。这个改动很小，但如果不做，遇到任意长度输入就会崩溃。另外，调试时还要注意数据类型对齐。half类型是16位，Vector单元处理时通常要求地址按16字节或32字节对齐。如果输入地址没有对齐，DataCopy可能会触发硬件异常。这些细节在CPU调试里很少见，但在昇腾NPU上是家常便饭。

性能分析：找到真正的瓶颈

结果正确只是第一步。把算子放到网络里，如果它拖慢了整体推理速度，就需要做性能分析。asc-tools里带的性能分析工具可以采集昇腾NPU上每个核函数的运行时间、内存带宽、流水线利用率、指令发射情况等指标。采集方式通常是在运行程序时加上环境变量或命令前缀，把性能数据写到某个目录，然后用可视化工具或命令行工具打开。分析的重点不是看单个数字，而是看各个单元的负载是否均衡。达芬奇架构里有Cube单元负责矩阵乘，Vector单元负责向量和逐元素操作，Scalar单元负责控制流。向量加法主要用Vector单元，如果Vector单元利用率很低，说明核函数可能被内存搬移或同步阻塞了。

采集性能数据的命令通常是在执行测试脚本之前设置一个环境变量，指向asc-tools的性能分析入口。运行结束后，会生成一个目录，里面有timeline文件和summary文件。打开summary文件，可以看到每个算子的耗时占比。打开timeline文件，可以看到时间轴上的事件序列，比如Cube计算、Vector计算、内存拷贝、同步等待等。下面这个命令演示了一个典型的采集方式。

export PROFILING_MODE=ascend
export PROFILING_DIR=./prof_data
python test_add_custom.py
asc-tools/bin/profiling_parser.sh --input=./prof_data --output=./report

拿到报告后，先看向量加法算子的耗时占比。如果它在整个推理图中只占了很小比例，说明优化优先级不高，可以先不管。如果它占比很高，再细看它的内部指标。对于这个简单算子，最常见的瓶颈是GM到UB的搬移。因为每次只算256个元素，大部分时间可能在等数据，而不是在做加法。这时候可以考虑几种优化方向：增大tile尺寸，让一次搬入更多数据；使用双缓冲，让下一次搬入和上一次计算重叠；或者把多个算子融合，减少中间结果的读写。这些优化思路不是凭空想的，而是看timeline上有没有大片空白或频繁同步。如果timeline显示Vector单元一直在忙，但等待时间很短，那说明算子本身已经吃得比较满，瓶颈可能在别处。如果timeline显示很多空白，说明核函数启动间隔太大，或者数据供给跟不上。

使用前后的效率对比

做完一遍完整流程后，我们来看优化前后的对比。下面的表格从几个维度对比了手工写CPU循环和用asc-tools在昇腾NPU上编译运行的差异。这些描述是概括性的，没有具体数字，因为实际性能受数据规模、芯片型号、内存带宽、框架版本影响，不能用一个固定数字概括。但方向是稳定的：在昇腾NPU上跑合适的算子，通常比纯CPU实现快很多，而且asc-tools让整个流程可维护、可复现。

维度	使用前	使用后	差异来源
计算位置	在CPU上循环执行加法	在昇腾NPU的Vector单元上执行	达芬奇架构的Vector单元提供并行硬件加速
数据搬移	输入输出在主机内存与CPU缓存之间传递	输入输出在Device内存和UB之间按tile搬移	数据不离开NPU，减少跨PCIe拷贝
编译方式	手写Python或C++循环，无专门编译优化	用asc-tools调用编译器做指令级优化	编译器针对昇腾芯片做流水线排布和内存对齐
调试方式	用通用调试器逐行跟踪	用asc-tools的算子调试和性能采集工具	工具链知道昇腾NPU的内存层级和事件语义
代码可维护性	每个算子手写循环，逻辑分散	统一用Ascend C和asc-tools工具链管理	算子、编译、调试、分析都在同一个仓库体系里
性能调优依据	凭经验猜测或加日志	看timeline和summary里的硬件事件	性能数据来自真实芯片运行，而不是模拟

这个表格没有列出具体毫秒或百分比，但已经能说明问题。使用前的方案在CPU上跑，通用性强但缺乏针对昇腾NPU的优化；使用后的方案把整个计算搬到芯片上，并且用asc-tools做编译和性能分析。差异来源不是某一条命令，而是工具链对硬件的理解。比如编译器知道Cube和Vector单元的指令延迟，会把指令排成流水线；性能分析器知道GM和UB的带宽，会指出数据搬移是不是瓶颈。这些能力是纯CPU方案没有的。

从调优回到代码：一个典型改进

拿到性能数据后，我们通常会回到算子代码做第二轮修改。对于上面的向量加法，一个典型改进是引入双缓冲。双缓冲的意思是在UB里同时准备两个tile，一个tile在计算时，另一个tile在搬入数据。这样Vector单元不用等DataCopy完成，可以连续工作。实现双缓冲只需要把队列深度从1改成2，并且在循环里交错使用两个LocalTensor。这个改动很小，但对流水线的改善很明显。另一个改进是增大tile尺寸。如果输入长度很大，每次处理1024个元素而不是256个，可以减少循环次数和同步开销。但tile尺寸不能无限增大，因为UB容量有限，太大装不下。如果还要融合ReLU或Scale等操作，可以一次性把多个操作合并到一个核函数里，省掉中间结果的GM读写。这些调优动作都需要以性能数据为依据，而不是凭感觉改。改完后再用asc-tools编译一次，再上板跑一遍，对比新的timeline。这种迭代方式比一次性写完美代码更可靠，因为芯片上的真实行为很难完全预测。

把asc-tools融入日常开发节奏

真正用asc-tools不是跑一次就够了，而是把它融入项目开发节奏。小团队可以做一个Makefile或者CMake脚本，把算子编译、单元测试、性能采集串成一条命令。每次修改算子代码后，先跑本地编译，再上板跑测试和性能分析，再把报告归档。这样做的好处是，算子改动不会悄无声息地引入性能回退。很多团队在前期只关注结果正确性，等模型上线后才发现某个算子比预期慢了很多，再回头定位成本就高了。如果每次提交都顺手跑一遍asc-tools的性能采集，问题在早期就能暴露。

在持续集成里，可以做一个轻量流水线：代码合入后自动触发编译检查，生成.o和.so，确认没有编译错误；然后触发上板测试，验证结果正确；再触发性能回归，把当前版本的timeline和上一个版本的timeline做对比。如果某次改动让某个算子的Vector单元利用率明显下降，流水线就报警。这样的流程不需要每次跑完整模型，只跑几个代表性的算子就够了。asc-tools的脚本化能力很强，bin目录下的命令大多支持非交互式调用，输出也可以定向到文件。把输出文件名带上版本号或commit号，历史数据就自然形成了趋势图。

---

仓库地址：https://atomgit.com/cann/asc-tools