昇腾自定义算子开发实战：从Ascend C到ATC部署全流程

这阵子接到一个迁移需求，要把一套基于三维重建的算法链路整体搬到昇腾设备上跑。模型本身倒还好，真正卡住我的是光栅化和高斯溅射那一段——原工程里有一段高度依赖CUDA的自定义kernel，昇腾预制算子根本覆盖不到，模型一跑到那个节点就报“算子不支持”。项目在这个问题上卡了将近一周，后来我决定自己动手写算子：从Ascend C的hello world开始，到把自研算子接进ATC转换链路，再到用profiling工具把单算子耗时从几十毫秒压到个位数毫秒。整个走下来，我对“昇腾自定义算子”这件事的最大感受是：门槛没有传说中那么高，但它确实有一套和CUDA不太一样的思考方式，尤其是tiling策略和数据流水线的组织。这篇就把完整链路捋一遍，覆盖版本配套、核函数编写、编译验证、性能剖析、部署接入和实战踩坑，适合正在做模型迁移、或者准备参加CAN算子开发相关赛事的工程师参考。

先说结论：昇腾设备上90%的模型需求，CANN预制算子都能覆盖，但剩下那10%的算法定制场景，恰恰决定了项目能不能真正落地。你越早掌握自定义算子的开发方式，迁移时就越主动。

1. 别急着写代码：昇腾算子开发的定位与版本配套

很多人拿到昇腾开发板或者Atlas服务器之后，第一反应是“赶紧装环境”，然后就开始照着文档写算子。这种思路不能说错，但容易在版本配套上栽跟头。昇腾这套软件栈和CUDA生态最大的区别就是分层相对封闭，驱动、CANN、框架插件三者的版本必须互相匹配，随便混搭很容易出现“驱动认不到卡”或者“CANN找不到算子”的诡异问题。所以我的建议是，动手写任何一行算子代码之前，先把底层的配套关系梳理清楚。

1.1 为什么需要自定义算子：预制算子覆盖不到的场景

CANN提供了数百个预制算子，覆盖了卷积、池化、归一化、矩阵乘、softmax等常见操作。但在实际项目中，总是会遇到这么几类情况：第一类是新算法结构，比如三维重建里的可微光栅化、高斯溅射的前向投影，这些步骤并没有现成的框架算子可以直接映射，CUDA实现往往是逐像素操作，昇腾这边预制算子根本没有对应项；第二类是算子融合需求，比如把激活函数融合进卷积、把一个elementwise链融合成一个算子，目的是减少中间张量的搬运开销；第三类是性能定制，用预制算子的组合方式能跑通，但中间结果反复在全局内存和片上缓存之间搬移，速度慢得没法接受，这时候就需要自己写一个融合算子把整条链路压缩到一次数据搬运里完成。

我的判断标准很简单：先用npu-smi info确认设备型号，用官方算子清单过一遍模型结构，只有确认预制算子确实覆盖不了，或者性能差到无法接受，才启动自定义算子开发。自定义算子不是目的，解决业务问题是目的，别为了炫技而重复造轮子。

1.2 算子开发方式选型：为什么使用Ascend C而不是TBE

昇腾自定义算子开发历史上主要有两条路线：早期的TBE（Tensor Boost Engine）和现在主推的Ascend C。TBE通过Python DSL描述计算逻辑，开发上手快，但它在编译优化粒度和调试能力上都有局限，复杂算子性能很难调上去。Ascend C则是基于C++的底层编程语言，提供类核函数的编程模型，程序员能直接控制数据搬运、缓冲区管理和流水线。

我强烈建议新项目直接选择Ascend C。虽然代码量大一些，但它有几个关键优势：一是对硬件细节的掌控力强，可以精确管理全局内存（GM）和片上统一缓冲区（UB）之间的数据流转；二是支持CPU模式调试，可以在没有昇腾设备的环境里先验证核函数逻辑；三是整个CANN生态后续迭代的重心在Ascend C上，TBE相关的文档和工具链基本处于维护状态。说白了，如果你只写一次算子无所谓，但你要是打算长期在这个生态里做开发，Ascend C是唯一值得投入的方向。

1.3 驱动、CANN、框架插件的版本配套关系

这个坑我踩得最惨。一开始我在一台Atlas设备上直接装了一个新版本的torch_npu，结果运行时一直报“ACL接口初始化失败”，排查了整整两天，最后发现是CANN版本和驱动版本不匹配。昇腾的版本配套关系是严格的三层结构：底层是NPU驱动（Driver），中间是CANN软件栈，上层是框架适配插件，比如torch_npu、mindspore。这三层只要有一层对不上，就会出现各种莫名其妙的问题。

我整理了一个当前阶段比较常用的配套思路，具体以官方发布配套表为准：

安装之前建议先执行npu-smi info查看驱动版本，再对照CANN官方文档的版本配套表选择CANN包。不要手动把CANN的安装目录从一个设备拷贝到另一个设备，也不要直接pip安装最新版torch_npu而不看它依赖的CANN版本号。配套版本这件事，五分钟能查明白，错了可能折腾你五天。

2. 实践一个Ascend C算子：从核函数逻辑到tiling切分

做好环境规划之后，就可以真正开始写算子。下面我用一个矢量加法算子作为实战案例，因为这个算子虽然简单，但包含了Ascend C开发中最核心的要素：核函数定义、缓冲区管理、队列同步和tiling策略。这套思路一旦你掌握了，换到矩阵运算、光栅化算子甚至融合算子上都是同一个套路。

2.1 算子结构设计：AI Core的执行视角

昇腾AI处理器的基本计算单元叫AI Core，每个AI Core内部有Cube单元（矩阵计算）、Vector单元（向量计算）和Scalar单元（标量计算）。作为一个Ascend C算子开发者，你关注的并不是这些单元的微架构细节，而是数据如何从全局内存（GM）搬到片上统一的缓冲区（UB），计算完成后又如何写回GM。

整个执行流程可以类比成一个流水线车间：GM是仓库，UB是车间里的操作台，Vector单元是工人。算子要做的事情就是反复执行“从仓库取货到操作台→工人在操作台上加工→把成品搬回仓库”这个循环。而Ascend C提供的关键抽象——TPipe、TQue、LocalTensor、GlobalTensor——全部是为了支撑这样一个搬运-计算-写回模型。

2.2 核函数侧实现：数据搬运、计算与队列管理

Ascend C的核函数需要使用特定的修饰符__global__ __aicore__，并在函数内部通过TPipe来管理UB缓冲。核心代码结构大致如下：

#include "kernel_operator.h" using namespace AscendC; constexpr int32_t BUFFER_NUM = 2; // 双缓冲 class KernelAdd { public: __aicore__ inline KernelAdd(TPipe* pipe, GM_ADDR x, GM_ADDR y, GM_ADDR z, int32_t totalLen, int32_t tileLen) : totalLen(totalLen), tileLen(tileLen) { xGm.SetGlobalBuffer((__gm__ half*)x, totalLen); yGm.SetGlobalBuffer((__gm__ half*)y, totalLen); zGm.SetGlobalBuffer((__gm__ half*)z, totalLen); pipe->InitBuffer(xQue, BUFFER_NUM, tileLen * sizeof(half)); pipe->InitBuffer(yQue, BUFFER_NUM, tileLen * sizeof(half)); pipe->InitBuffer(zQue, BUFFER_NUM, tileLen * sizeof(half)); } __aicore__ inline void Process() { int32_t blockIdx = GetBlockIdx(); int32_t blockNum = GetBlockNum(); int32_t perBlock = totalLen / blockNum; int32_t start = blockIdx * perBlock; int32_t curLen = (blockIdx == blockNum - 1) ? (totalLen - start) : perBlock; for (int32_t i = 0; i < curLen / tileLen; i++) { CopyIn(i, start); Compute(i); CopyOut(i, start); } // 尾块处理略 } private: __aicore__ inline void CopyIn(int32_t i, int32_t start) { LocalTensor<half> xLocal = xQue.AllocTensor<half>(); LocalTensor<half> yLocal = yQue.AllocTensor<half>(); DataCopy(xLocal, xGm[start + i * tileLen], tileLen); DataCopy(yLocal, yGm[start + i * tileLen], tileLen); xQue.EnQue(xLocal); yQue.EnQue(yLocal); } __aicore__ inline void Compute(int32_t i) { LocalTensor<half> xLocal = xQue.DeQue<half>(); LocalTensor<half> yLocal = yQue.DeQue<half>(); LocalTensor<half> zLocal = zQue.AllocTensor<half>(); Add(zLocal, xLocal, yLocal, tileLen); zQue.EnQue(zLocal); xQue.FreeTensor(xLocal); yQue.FreeTensor(yLocal); } __aicore__ inline void CopyOut(int32_t i, int32_t start) { LocalTensor<half> zLocal = zQue.DeQue<half>(); DataCopy(zGm[start + i * tileLen], zLocal, tileLen); zQue.FreeTensor(zLocal); } private: GlobalTensor<half> xGm, yGm, zGm; TQue<QuePosition::VECIN, BUFFER_NUM> xQue, yQue; TQue<QuePosition::VECOUT, BUFFER_NUM> zQue; int32_t totalLen, tileLen; }; extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, int32_t totalLen, int32_t tileLen) { TPipe pipe; KernelAdd op(&pipe, x, y, z, totalLen, tileLen); op.Process(); }

这段代码里有几个关键点需要特别理解。AllocTensor是从UB队列里申请一块缓冲区，EnQue是把填好数据的缓冲区交给下一级消费，DeQue是从队列里取出一块已经生产好的缓冲区。通过这种队列机制，Ascend C编译器才能自动识别出搬运和计算之间的依赖关系，并尝试在双缓冲场景下做流水线重叠，也就是数据搬运当前tile的同时，Vector单元正在计算上一个tile。你在写代码时不要去手动加同步等待，把队列关系表达清楚，编译器会自动编排。

2.3 Host侧逻辑与Tiling策略计算

核函数侧只处理“一个核在某一时刻计算什么”，而每个核计算多少数据、总的tile怎么切分、切多大，这些逻辑在Host侧（也就是CPU侧）完成。Ascend C的算子工程里，Host侧需要实现算子原型定义（输入输出、属性、数据校验）和Tiling结构计算。

以我上面这个add算子为例，Tiling计算要考虑几个参数：总数据长度totalLen、可用的核数blockNum、每个tile的长度tileLen。其中tileLen的确定是最关键的，它受限于UB容量。以昇腾910B上的AI Core为例，UB可用空间大约192KB。我们要同时放x、y、z三份数据，开了双缓冲意味着缓冲区数量再加倍，那么单份张量一个tile最多使用约 192KB / (3 * 2) = 32KB。使用half类型占2字节，也就是每个tile最多可以放16384个半精度数。再考虑对齐约束，16384个half恰好是32KB，天然满足32字节对齐。

Host侧Tiling结构可以定义如下：

struct AddTilingData { int32_t totalLen; int32_t tileLen; int32_t blockNum; };

这个结构体会被写入Tiling buffer，然后在kernel侧通过参数传入。Host侧计算逻辑里需要根据设备实际的AI Core数量来设置blockNum，并处理总长度不能被核数整除的情况——通常是让最后一个核处理余数部分。我建议在Host侧把所有边界条件都处理好，不要在kernel里做太多分支判断，kernel的每个分支都可能带来性能损失。

2.4 为什么tiling是整个算子的性能命门

刚开始写算子的人往往只关心功能正确，跑出正确结果就觉得很满足了。但tiling策略直接决定了算子能不能吃满AI Core的算力。如果tile切得太大，UB放不下，程序运行时会越界写坏数据，偶尔报错偶尔正确，非常难排查；如果tile切得太小，数据搬运的次数就会变多，而搬运本身是要消耗时间的，搬运时间占比一旦高起来，即便计算单元再快，整体耗时也会被拖住。

更合理的做法是尽量让tile块与硬件的数据位宽、缓存行对齐，同时利用多核并行。比如总长度是1024000个元素，设备有32个AI Core，每个核分到32000个元素，32000可以被tileLen=16000整除，每个核循环两次就能跑完。这类计算在Host侧可以通过对totalLen / blockNum做向上取整到tileLen的倍数来规划，保持所有核的负载尽量均衡。tiling的调优不是一次到位的，通常需要结合profiling数据反复调整，这个后面专门说。

3. 编译、跑板、剖析：性能瓶颈排查的完整链路

代码写完之后，接下来的流程是：用模板生成算子工程、编译、CPU模式验证、上板跑单算子、profiling分析性能。很多初学者喜欢一口气把算子写到工程里然后直接上板，结果编译报错一堆，调试效率极低。我的建议是把验证拆成几个阶段，每个阶段用最轻量的方式暴露问题。

3.1 用msopgen生成算子工程骨架

昇腾提供了算子工程生成工具msopgen，可以根据你指定的算子名称生成一套标准的Ascend C工程骨架。执行方式类似：

msopgen gen -i add_custom -c ai_core-ascend910b -out .

生成完毕后的工程目录大概是这样的结构：

add_custom/ ├── CMakeLists.txt ├── op_host/ │ ├── add_custom_tiling.h │ └── add_custom.cpp ├── op_kernel/ │ └── add_custom.cpp ├── build.sh └── scripts/

op_host里放的是Host侧逻辑，包括算子原型注册和tiling计算；op_kernel里放的是核函数实现。生成后，把核函数代码替换成自己的实现，再补充Host侧的原型定义和tiling计算，就可以编译了。msopgen的好处是你不需要手写CMakeLists和目录结构，避免在最容易出错的工程配置上花时间。

3.2 编译与CPU模式验证

编译时使用工程自带的build.sh，它会调用底层编译器把kernel代码编译成昇腾设备可执行的指令，Host侧代码编译成CPU可执行的so。特别要指出的是，Ascend C支持CPU模式，你可以在没有昇腾硬件的情况下先用CPU模拟运行kernel，这功能对逻辑调试特别有用。启动CPU模式时，可以在代码里加一个调用kernel的main函数，编译时加上CPU模式相关参数，然后在宿主机上直接运行。

CPU模式的定位是验证逻辑正确性，不是性能。你可以在CPU模式下一行一行地看数据变化，尤其是Add(zLocal, xLocal, yLocal, tileLen)这种计算的输入输出位置是不是搞反了、缓冲区索引是不是越界、尾块处理是否有遗漏。等CPU模式跑出的结果和参考实现完全一致，再上板。

3.3 上板单算子运行与结果比对

上板之后，建议先做一个单算子级别的二项式测试：构造随机输入数据，用自定义算子的输出和CPU/GPU参考输出对比，要求误差在允许范围内。昇腾设备上可以用AscendCL提供的接口加载算子包，直接下发单算子任务。如果你的算子包在单算子模式下性能很慢，也不要急着怀疑硬件，先检查是不是profiling工具配置没开、或者算子首次加载时的算子编译缓存还没有建立。

单算子运行通过后，再把它接进模型验证。如果模型推理结果不对，通常有三种可能：算子本身算错了、tiling在某个边界shape上算错了、或者接进模型时数据排布（format）和算子期望不一致。我之前遇到过一个看起来很诡异的问题：小shape对，大shape偶尔不对，浮点误差也正常。最后定位出来是tiling结构体里某字段在Host侧没有初始化，导致在特定shape下读到脏数据。所以我在tiling代码里都会先memset整个tiling buffer，再填充字段。

3.4 用msprof把性能瓶颈一个个挖出来

功能没问题之后，才轮到性能分析。昇腾上最常使用的profiling工具是msprof，可以采集AI Core利用率、Vector利用率、搬运耗时、任务同步等待等关键指标。采集命令大致是这样：

msprof --application="./add_main" --output=./prof_data

跑完之后分析prof_data目录下的csv文件，我一般重点看这几个指标：

以我的add算子为例，如果Vector利用率只有20%，而数据搬运耗时占比高达70%，说明整个算子几乎都在等数据，计算单元闲得发慌。这时候优先检查双缓冲是否生效，以及tileLen是否选得合理。如果tileLen过小，搬运启动次数过多，搬运效率自然上不去；如果tileLen过大，导致UB中同时只能存在一份缓冲，计算和搬运就无法重叠，也会拖慢整体吞吐。

3.5 双缓冲与流水线优化：让搬运和计算重叠

双缓冲的原理很简单：第一个tile还在计算时，第二个tile的数据已经从GM搬到UB里了。这样数据搬运和向量计算同时进行，AI Core的利用率自然提升。在Ascend C的队列模型里，双缓冲通过将队列深度设为2来实现，就是代码里TQue<QuePosition::VECIN, BUFFER_NUM>中的BUFFER_NUM=2。

但开了双缓冲不意味着一定有效，还需要确认代码写得不会串行化。实践中常见的问题是：CopyIn、Compute、CopyOut三个函数在循环里严格顺序执行，看起来逻辑没问题，但如果编译器分析出前后tile之间存在依赖，比如你反复用同一个全局tensor作为中间结果，双缓冲就会被自动退化成单缓冲。解决方法是确保不同tile之间没有数据竞争，每块UB buffer都通过AllocTensor和FreeTensor管理生命周期，让编译器清楚当前缓冲的使用状态。另外还可以考虑多stream并行，让不同AI Core在执行不同tile时进一步打散同步点。

4. 部署接入与上线调优：让自定义算子真正进入业务链路

算子本身性能和正确性都达标后，下一步就是把算子部署到真实链路里。昇腾的自定义算子部署，核心是把编译好的算子包让两个东西认识：一是模型转换工具ATC，二是运行时执行引擎。这一步如果配置不对，模型转换时会提示找不到算子，运行时会提示算子加载失败。

4.1 生成算子包并接入ATC转换

编译完成的算子工程会生成算子so，一般包括libcust_op.so和对应的算子原型描述文件。部署时需要把这几个文件放到CANN能扫描到的自定义算子目录，并设置环境变量：

export ASCEND_CUSTOM_OPP_PATH=/path/to/custom/opp

让ATC能识别自定义算子的方式，是保证算子的原型注册（op proto）和tiling实现都打进so，并且so文件路径包含在ASCEND_CUSTOM_OPP_PATH对应的vendor目录中。然后执行模型转换：

atc --model=model.onnx --framework=5 --output=model_ascend \ --soc_version=Ascend910B4 --output_type=FP32

如果模型里使用的算子名和你自定义算子的注册名不一致，有两种处理方式：一是修改模型中的算子名，让其匹配自定义算子的原型名；二是在ATC转换时通过算子的语义匹配让ATC自动替换。大部分情况下我推荐第一种，直接、可控、不容易误伤其他算子。改完模型后，ATC如果仍然报找不到算子，优先检查ASCEND_CUSTOM_OPP_PATH是否设置，以及so是否确实被编译成了当前CANN版本能识别的格式。

4.2 在PyTorch昇腾侧调用自定义算子

在训练场景下，自定义算子往往需要嵌入PyTorch计算图中。通过torch_npu框架，昇腾设备在PyTorch里被抽象成NPU设备，执行model.npu()后模型计算会被调度到昇腾设备。但自定义算子在训练链路里通常有两种接入方式：

第一种方式是把自定义算子导出到ONNX，利用ATC转换成om模型后走推理引擎。这条路适合推理部署。

第二种方式是使用AscendCL的aclnn接口，在Python侧直接封装一个PyTorch自定义autograd.Function，forward里调用aclnn接口，backward里也调用对应的反向算子。这样做的好处是算子可以参与反向传播和梯度计算，适合训练场景。但开发成本比推理链路高，通常还要实现反向算子的kernel，这时的重点是确保反向算子与正向算子使用同一套tiling策略，否则在梯度计算时shape变化会触发边界bug。

对于这个add算子，如果只是debug，最方便的方式是用torch_npu下的自定义算子调试接口直接跑通NPU张量的输入输出。我个人的建议是：不要一开始就搞完整的训练链路，先在一个Python脚本里构造固定shape的张量，调用自定义算子包的aclnn接口，确认结果与CPU一致，再把算子挂进训练循环。

4.3 部署中的显存、多卡并发与稳定运行策略

算子从demo变成服务，还要考虑资源规划。昇腾设备虽然叫NPU，但它也有类似显存的片外存储，多卡场景下要按卡分配数据，避免某张卡OOM。部署时我会先跑一遍全量数据的profiling，确认单算子的峰值内存占用，然后根据峰值内存和推理batch size计算一个卡上最多能放多少个并发实例。

多卡并发时，需要注意host侧的CPU线程不要成为瓶颈。算子的启动和同步都是通过AscendCL接口完成的，如果host侧用了Python的多线程，还要关注进程绑核和NUMA拓扑。实际场景里我常看到的现象是：模型本身很快，但多卡数据加载和预处理把CPU吃满了，导致NPU空转。这种时刻要用多进程而非多线程，并且把数据预处理和NPU计算分配到不同的CPU核心，避免共享同一个解释器锁。

4.4 算子版本管理与灰度发布

自定义算子一旦进入业务，就不再是“本地代码”，而是需要纳入版本管理的线上组件。我强烈建议把算子工程的源码、编译产物、依赖的CANN版本、测试向量和基准性能数据全部放到同一个仓库里，每次修改都记录一份性能对比报告。算子这种底层组件，改了一行tiling代码，可能让某类shape性能提升30%，同时让另一类shape性能退化15%。没有基准数据，你是发现不了这种回归的。

灰度发布的顺序我习惯是：先在单卡环境跑离线评测，再在测试环境跑在线评测，最后切小流量上线。每次发布都保留上一个算子包的版本，遇到问题立刻回退。自定义算子出问题通常不是编译期报错，而是运行期偶发的计算错误，这种问题最难查，所以版本管理和灰度发布不是流程上的形式主义，而是给你留一条退路。

5. 半年踩坑复盘：几个容易反复的细节问题

文章最后分享几个我在这类开发任务里踩过、也看同事踩过的坑。它们并不是什么高深原理，但每一个都能让人卡上好几天。

5.1 内存对齐与UB越界：最隐蔽的错误来源

昇腾的DataCopy指令对地址对齐有严格要求，特别是global memory到UB的搬运，通常要求32字节对齐。如果你的tileLen算出来不是对齐单位的整数倍，搬运接口不会直接报错，而是搬运到错误的地址，导致计算结果出现“随机性”错误。排查这类问题的方法是：在Host侧tiling代码里，对所有长度参数都做对齐处理，宁可多搬运几个无用元素，也不要让长度参数“非对齐”。

另外，UB缓冲区的大小是有限资源。有时候你在kernel里多开了一个临时tensor，或者tileLen算大了一些，编译和单次运行都正常，但多次循环之后就会溢出，直接踩坏相邻buffer的数据。这种问题的坑在于它的偶发性，最好在开发阶段就在关键位置加上边界检查断言，或者对超大shape做一轮专门的压测。

5.2 版本错配导致的“灵异现象”

我在文章开头说过版本配套关系，这里再强调一次细节：驱动、CANN、torch_npu三者错配时，现象往往不是直接报版本错误，而是“运行时找不到设备”“算子上板后跑不起来”“PyTorch导入torch_npu时崩溃”。因为这些组件之间通过底层接口通信，版本不匹配时接口符号可能对不上，但又不会给出清晰的提示。所以遇到这些灵异现象，先不要急着debug代码，先核对版本配套表，这个动作应该排在任何调试步骤之前。

我个人的工作流是在每个项目里建一个环境信息文件，记录驱动版本、CANN版本、Python版本、framework插件版本以及每个算子的编译参数。排查问题的时候，第一件事不是看代码，而是看这个文件。因为昇腾的软件栈更新较快，你三个月前编译的算子so，三个月后在新环境里可能就加载不了了，这时候重新编译往往比深挖代码更高效。

5.3 算子编译缓存带来的二次运行困惑

昇腾的ATC转换和算子运行时都有算子编译缓存机制。第一次运行某个shape的算子时，系统会生成对应的kernel指令并缓存下来；第二次再遇到相同shape，就直接复用缓存。这个机制本身很好，但如果你修改了算子代码，重新编译了算子包，却没有清缓存，运行时的表现可能仍然是旧kernel逻辑，让你误以为“代码改了没生效”。遇到这种情况，除了重新编译算子包，还要清理ATC和runtime的缓存目录，或者换一个输出目录重新生成om模型。

5.4 建立自己的算子性能基线库

最后一条建议，也算是我这几年开发算子沉淀下来的习惯：每写一个算子，都建一个性能基线记录。记录内容包括不同shape下的耗时、AI Core利用率、Vector利用率、搬运耗时占比、tiling参数。不要只记录最高性能的那组数据，要把调整过程中的关键中间版本也记下来。这样你不仅能知道当前版本跑得快不快，还能知道它为什么快、改坏了可以从哪个版本回退。

昇腾自定义算子开发这条路，最考验人的不是某个具体API怎么调用，而是能不能建立一套自己的排查和验证节奏。我先通过CPU模式验证逻辑，再上板验证正确性，再用profiling验证性能，最后用灰度发布保障上线。这套流程看起来朴素，但每一次都能帮我快速定位问题。如果你正在被某个昇腾算子问题卡住，不妨停下来，把版本、shape、tiling参数和profiling数据列成一张表，问题通常就会自己浮出水面。