前一阵做模型功能迁移，一个自研模块需要把Y = (A + B) * scale这类基础计算落回加速卡执行。听起来就是一次普通算子替换，结果现成算子库翻遍了，要么只有分开的 Add 和 Mul，要么融合算子的输入输出限制对不上。用框架自带的两个算子前后拼一次，前向要多读写一遍全局内存，数据量一旦过百万，耗时和带宽都很难看。最后只能在 CANN-Ascend 平台上自己动手，用 C 算子开发把这个融合逻辑写成单个算子。

这篇文章就是我把整个过程重新梳理之后的完整记录。适合三类人看：一是模型迁移时发现算子库里总差“那么一个算子”的；二是想了解 CANN-Ascend 的 Ascend C 算子开发到底要过哪些关卡的；三是已经在写算子但经常被搬运、同步、对齐问题绊住脚的开发者。内容从为什么需要自研算子讲起，到工程结构、核心原理、完整代码示例、编译部署、性能优化，最后是踩坑记录，全程按我实际的开发顺序来写，不去绕那些空泛的概念。

1. 为什么自己写算子：现成算子覆盖不到的地方

1.1 从一次算子下沉经历说起

先还原我当时遇到的具体场景。模块里的计算很简单，对输入张量 A 和 B，先做逐元素相加，再统一乘一个标量 scale，得到输出 Y。框架里确实有 Add 和 Mul，拼起来也能跑。问题出在中间结果上：Add 运算完会把一个和输入同等大小的中间张量写回全局内存，Mul 再把它读出来，计算完再写回最终结果。

这一来一回多了一次完整的全局内存写和一次完整的全局内存读。对于带宽敏感型算子，多出来的开销往往比计算本身还大。我在一批 256K 个 float16 元素的张量上做过对比，两个算子串联的耗时大约是单个融合算子的 1.8 倍左右。原因是计算单元大部分时间都在等数据搬运，多绕一圈就多等一轮。

这就是自研算子的第一个价值：把多次 kernel 启动和中间张量的多次读写合并成一次。类似 FusedAddScale 这种简单场景，融合之后带宽占用大概能少一半。实际业务里更复杂的融合，比如把卷积、偏置、激活、量化全部打进一个算子，节省的时间和显存会更加可观。

1.2 什么信号说明你确实需要自定义算子

不是所有场景都值得自研算子，判断标准总结下来大概有这几条：

• 算子库没有对应实现，比如新论文里的激活函数、特殊的注意力变体、某些量化算子。
• 算子库有实现但融合粒度不够，多个算子串联导致过多的中间张量读写。
• 要实现层面对精度、内存布局有特殊要求，比如必须用 float32 做累加、必须输出特定格式化数据。
• 现有算子在特定 shape 下性能不达标，希望通过 tiling 和流水线优化拿到更高吞吐。

我把现成算子和自定义算子的差距整理成一张表，方便对照判断：

我当时的选择逻辑也很简单：先确认这是不是一次性需求，再确认计算逻辑是否足够稳定。FusedAddScale 看着简单，但它是我后续要扩展成一族融合算子的通用骨架，值得投入。如果只是临时验证一个 idea，直接用 Python 拼算子就行，没必要下沉。

2. 上手前需要想清楚的：Ascend C 到底在写什么

2.1 一个算子被拆成 Host 和 Device 两半

CANN-Ascend 的 Ascend C 算子开发，核心模型是把算子拆成 Host 侧和 Device 侧两部分。Host 侧跑在 CPU 上，负责任务切分，也就是常说的 tiling；Device 侧跑在 AI Core 上，负责真正的计算。两者不在同一个地址空间，通过参数和 tilingData 通信。

第一次接触这个模型的人最容易犯的错，是试图把 GPU 上写 CUDA Kernel 的习惯直接搬过来。在 CUDA 里，一个__global__函数同时承担切分和执行；在 Ascend C 里，切分逻辑放在 Host 侧的 tiling 函数里，Device 侧只拿到“我这一核处理哪一段数据”的指令，然后老老实实完成搬运、计算、写回。

我习惯用一个类比：Host 侧是工程总包，AI Core 是施工队。总包先看整体工程量，划分出几个施工段，告诉每个施工队干哪一段、用什么工具、多少材料；施工队自己不用关心整体计划，只管拿到图纸后按流程干完。这个“图纸”就是 tilingData。

2.2 一份算子工程的三件套

用工具链生成算子工程之后，目录里最核心的是三类文件，理解了它们就等于理解了整个开发框架：

• 算子原型描述文件：声明算子的输入、输出、属性，是算子在框架里的“身份证”。算子类型名、参数列表、形状推导规则都在这里。
• Host 侧实现文件：包含 tiling 函数，负责计算任务怎么切分、需要几个核、每个核处理多长数据。
• Device 侧实现文件：落在 AI Core 上执行的核心逻辑，包含数据搬运、计算、同步。

目录结构大致是这样：

FusedAddScale/ ├── op_info.json # 算子原型与编译配置 ├── host/ │ ├── fused_add_scale_tiling.cpp │ └── CMakeLists.txt ├── device/ │ ├── fused_add_scale_kernel.cpp │ └── CMakeLists.txt ├── scripts/ │ └── build.sh └── cmake/ └── ...

op_info.json决定了算子在网络中的名字和参数签名；host目录决定算子的并行策略；device目录决定单个 AI Core 上的执行逻辑。三样东西合起来，才构成一个能在训练或推理链路上被调用的完整算子。

2.3 工程从哪来：生成工具与编译链路

CANN-Ascend 工具链提供了一个叫算子生成器的命令行工具，可以从原型描述文件直接生成完整的算子工程骨架。常用形式是传入算子定义和芯片型号，生成 host、device 目录和构建脚本。

msopgen gen -i op_info.json -c ai_core-<soc_version> -out ./FusedAddScale

生成后再编译，算子会生成对应的二进制与算子描述，注册进算子库。后续通过 ACL 或框架侧接口就能拉起执行。整个流程从生成工程到跑通单算子调用，我第一次走大约花了半天，其中大部分时间花在环境变量和编译参数上。后面写熟了会发现，这套链路的核心就一句话：原型描述负责能不能被找到，Host 代码负责切得均不均匀，Device 代码负责跑得快不快。

3. 核心原理解读：数据在哪搬、怎么搬、怎么算

3.1 三层存储结构与显式搬运

写 Ascend C 算子绕不开存储层级。这里有一个极简化的三层模型：GM 是全局内存，容量大但访问慢；UB 是每个 AI Core 上的高速暂存区，容量有限但访问快；再往里有向量计算单元和矩阵计算单元，它们直接从 UB 取数计算。

计算流程通常是：先从 GM 把数据搬运到 UB，再由计算单元从 UB 读取数据计算，结果先放回 UB，最后由程序主动把结果搬回 GM。整个过程是显式的，搬运指令需要开发者自己写。这一点和 GPU 上默认“所有显存访存都自动完成”的体验非常不一样。

用厨房备菜的类比来说，GM 是仓库，UB 是案板，计算单元是灶台。算子开发不是只写“炒菜”这一步，而是要写一套备菜流程：今天要切多少菜、分几次拿、案板上同时放几盘、炒完一盘怎么撤下来。案板就那么大，菜放多了放不下，放少了又浪费等菜的时间。

3.2 数据队列 TQue 与缓冲区的循环利用

Ascend C 用队列抽象来管理 UB 上的缓冲区。每个队列绑定一块 UB 空间，开发者通过AllocTensor申请一块 LocalTensor，用DataCopy搬运数据进来，然后EnQue把数据块加入队列；计算侧通过DeQue从队列取出数据，算完再用FreeTensor释放。

这套队列机制解决了一个关键问题：UB 空间有限，不能一次性放下全部数据，必须分块复用。以 FusedAddScale 为例，如果输入是 100 万个 float16，UB 很可能放不下，那就按块循环：搬一块、算一块、写回一块，再把这块 buffer 腾出来装下一个数据块。队列帮我们管理了这块空间的申请和释放。

LocalTensor<half> x1Local = inQueueX1.AllocTensor<half>(); DataCopy(x1Local, x1Gm[curOffset], curLen); inQueueX1.EnQue(x1Local); LocalTensor<half> x1Compute = inQueueX1.DeQue<half>(); // 计算 outQueueY.EnQue(yLocal); inQueueX1.FreeTensor(x1Compute);

这段逻辑看起来平淡，但它是整个算子开发里最重要的循环结构。开发者不需要自己记账哪块 UB 空闲，队列替我们维护了状态。

3.3 同步不是可选项：SyncIO 与 SyncCompute

搬运和计算都是异步发出去的指令。DataCopy发起一个从 GM 到 UB 的搬运，它不会阻塞当前核；紧接着的乘法指令可能在数据还没到 UB 时就已经开始执行，拿到的就是脏数据。

所以 Ascend C 里到处是同步点。pipe.SyncIO()表示等待这一批搬运指令完成，pipe.SyncCompute()表示等待计算完成，确保数据写回 UB 后才能开始下一次搬运。漏加同步点的典型症状是结果“时好时坏”：小数据量时搬运快，侥幸没出错；数据量一大，偶发计算结果全乱。

同步点也是流水线设计的一部分。如果用双缓冲区（BUFFER_NUM = 2），可以在计算第 N 块的同时预取第 N+1 块数据，让搬运和计算重叠起来。没有同步点，双缓冲无从谈起；同步点加得太频繁，流水线又被切碎，性能上不去。这块是 Ascend C 编程最有“调优感”的地方。

3.4 Tiling：把任务切到每个核头上

tiling 的目标很简单：给定总数据量、可用核数、UB 大小，算出每个核处理多长数据、分几轮处理。Host 侧算出来的结果放进一个结构体，Device 侧拿到之后照着执行。

uint32_t coreNum = context->GetCoreNum(); uint32_t totalLength = context->GetInputShape(0)->GetShapeSize(); uint32_t blockLength = (totalLength + coreNum - 1) / coreNum; // 向上对齐到 16 的倍数，保证后续搬运长度满足 32 字节粒度 blockLength = ((blockLength + 15) / 16) * 16; uint32_t blockDim = (totalLength + blockLength - 1) / blockLength;

这里有个细节值得说明：为什么要做对齐取整？因为 GM 与 UB 之间的搬运通常要求按 32 字节粒度对齐，float16 一个元素 2 字节，16 个元素正好 32 字节。如果分配给每个核的长度不是 16 的整数倍，最后一段数据搬运的边界会变得很难处理。tiling 阶段把长度对齐，能省掉 Device 侧大量边界判断。

4. 实战：从零手写一个 FusedAddScale 融合算子

4.1 需求与实现约定

这个算子的语义：Y = (A + B) * scale。A、B、Y 是一维 float16 张量，长度相同且任意；scale 是一个标量属性。先限定输入总长度由外部确保不为 0，并且单核可处理长度足够大，方便演示。

选择这个例子有三个原因：逻辑足够简单，可以专注理解工程骨架；它包含了加法和标量乘两种典型向量运算；这个计算模式在归一化、残差连接、量化缩放里很常见，扩展性很好。

4.2 第一步：算子原型描述

算子原型描述文件定义了算子的输入、输出和属性。这里我用一个通用 JSON 风格的描述来表达，实际字段名以所选生成工具的版本为准，核心信息就是三类：输入张量、输出张量、标量属性。

{ "op_type": "FusedAddScale", "inputs": [ { "name": "x1", "param_type": "REQUIRED", "tensor_type": { "type": "float16", "format": "ND" } }, { "name": "x2", "param_type": "REQUIRED", "tensor_type": { "type": "float16", "format": "ND" } } ], "outputs": [ { "name": "y", "param_type": "REQUIRED", "tensor_type": { "type": "float16", "format": "ND" } } ], "attrs": [ { "name": "scale", "param_type": "REQUIRED", "type": "float", "value": 1.0 } ] }

写原型描述要特别注意算子名别写错，因为它直接对应后续框架调用和算子查找的名字。当时我踩过一次坑：核函数名内部写的是 FusedAddScale，原型描述里写成了 FusedAddScala，结果算子加载后怎么都找不到，排查了半天才发现是少了一个字母。

4.3 第二步：Host 侧 tiling 实现

tiling 函数的核心输出有三个：总长度、每个核处理的块长度、启动的核数。Device 侧根据这三个参数就能定位自己要处理的数据区间。

#include "register/op_def_registry.h" #include "tiling/tiling_api.h" namespace optiling { constexpr uint32_t kAlign = 16; static ge::graphStatus TilingFusedAddScale(gert::TilingContext* context) { auto shape = context->GetInputShape(0); uint32_t totalLength = 1; for (size_t i = 0; i < shape->GetDimNum(); ++i) { totalLength *= shape->GetDim(i); } uint32_t coreNum = context->GetCoreNum(); uint32_t blockLength = (totalLength + coreNum - 1) / coreNum; blockLength = ((blockLength + kAlign - 1) / kAlign) * kAlign; uint32_t blockDim = (totalLength + blockLength - 1) / blockLength; blockDim = blockDim > coreNum ? coreNum : blockDim; context->SetBlockDim(blockDim); auto tilingData = context->GetTilingData(); uint32_t* data = reinterpret_cast<uint32_t*>(tilingData); data[0] = totalLength; data[1] = blockLength; return ge::GRAPH_SUCCESS; } } // namespace optiling

这个实现里做了两个防御性处理：一是 blockLength 对 16 向上取整，避免出现非对齐长度；二是 blockDim 做了截断，防止总长度很小的时候启动过多空转核。实际工程中还要考虑每个核处理长度不能超过 UB 容量，这块可以在定义 tiling 结构时再加一个循环次数限制，我这里省略了循环拆分的细节。

4.4 第三步：Device 侧 Kernel 实现

Device 侧是真正跑在 AI Core 上的代码。核心逻辑：每个核根据 blockIndex 算出自己的起始偏移，循环处理属于自己的数据区间，每次搬运一块到 UB，计算完写回。

#include "kernel_operator.h" using namespace AscendC; constexpr int32_t BUFFER_NUM = 2; class KernelFusedAddScale { public: __aicore__ inline KernelFusedAddScale(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, uint32_t totalLength, uint32_t blockLength, float scale) : totalLength(totalLength), blockLength(blockLength), scale(scale) { x1Gm.SetGlobalBuffer((__gm__ half*)x1, totalLength); x2Gm.SetGlobalBuffer((__gm__ half*)x2, totalLength); yGm.SetGlobalBuffer((__gm__ half*)y, totalLength); pipe.InitBuffer(inQueueX1, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(inQueueX2, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(outQueueY, BUFFER_NUM, blockLength * sizeof(half)); } __aicore__ inline void Process() { uint32_t coreIdx = GetBlockIdx(); uint32_t coreStart = coreIdx * blockLength; uint32_t coreEnd = coreStart + blockLength; if (coreEnd > totalLength) { coreEnd = totalLength; } uint32_t step = blockLength; for (uint32_t offset = coreStart; offset < coreEnd; offset += step) { uint32_t curLen = (offset + step) < coreEnd ? step : (coreEnd - offset); CopyIn(offset, curLen); pipe.SyncIO(); Compute(offset, curLen); pipe.SyncCompute(); CopyOut(offset, curLen); } pipe.SyncIO(); } private: __aicore__ inline void CopyIn(uint32_t offset, uint32_t length) { LocalTensor<half> x1Local = inQueueX1.AllocTensor<half>(); LocalTensor<half> x2Local = inQueueX2.AllocTensor<half>(); DataCopy(x1Local, x1Gm[offset], length); DataCopy(x2Local, x2Gm[offset], length); inQueueX1.EnQue(x1Local); inQueueX2.EnQue(x2Local); } __aicore__ inline void Compute(uint32_t offset, uint32_t length) { LocalTensor<half> x1Local = inQueueX1.DeQue<half>(); LocalTensor<half> x2Local = inQueueX2.DeQue<half>(); LocalTensor<half> yLocal = outQueueY.AllocTensor<half>(); Add(x1Local, x2Local, yLocal, length); Muls(yLocal, (half)scale, yLocal, length); outQueueY.EnQue(yLocal); inQueueX1.FreeTensor(x1Local); inQueueX2.FreeTensor(x2Local); } __aicore__ inline void CopyOut(uint32_t offset, uint32_t length) { LocalTensor<half> yLocal = outQueueY.DeQue<half>(); DataCopy(yGm[offset], yLocal, length); outQueueY.FreeTensor(yLocal); } private: GlobalTensor<half> x1Gm, x2Gm, yGm; TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX1, inQueueX2; TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY; TPipe pipe; uint32_t totalLength; uint32_t blockLength; float scale; }; extern "C" __global__ __aicore__ void FusedAddScale(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, uint32_t totalLength, uint32_t blockLength, float scale) { KernelFusedAddScale op(x1, x2, y, totalLength, blockLength, scale); op.Process(); }

代码里的几个关键点：Process中每个核处理一段连续区间，这样可以避免多个核访问同一段数据的竞争；SyncIO在搬运之后同步，保证数据到达 UB 才开始计算；双缓冲区BUFFER_NUM = 2让下一块数据的搬运可以与当前块的计算重叠。受篇幅限制，这个示例假设数据长度已经满足对齐要求，实际处理尾部不对齐数据的方法放在第 6 节讲。

4.5 第四步：编译与在 ACL 侧拉起验证

编译算子使用生成好的工程目录，进入后执行构建脚本即可，通常会产出算子的动态库和描述文件。拿到算子之后，最直接的单算子验证方式是通过 ACL 接口拉起。核心流程分为五步：

1. 初始化 ACL 环境，设置当前设备。
2. 创建 stream。
3. 在设备侧申请输入输出内存，把 Host 侧输入数据拷贝到设备侧。
4. 准备算子描述和 tilingData。
5. launch 算子，等待执行完成后把输出拷贝回 Host。

aclInit(nullptr); aclrtSetDevice(0); aclrtCreateStream(&stream); void* devA = nullptr; aclrtMalloc(&devA, size, ACL_MEM_MALLOC_NORMAL_ONLY); aclrtMemcpy(devA, size, hostA, size, ACL_MEMCPY_HOST_TO_DEVICE); // 同理处理 devB、devY // 构造算子相关描述并 launch，接口名以当前版本为准 aclrtlaunchOp(desc, blockDim, tilingData, tilingSize, stream); aclrtSynchronizeStream(stream); aclrtMemcpy(hostY, size, devY, size, ACL_MEMCPY_DEVICE_TO_HOST);

第一次跑通单算子时，最直接的结果就是输出张量等于(A + B) * scale。从写代码到拿到这个结果，我大概用了一天，中间大部分时间都花在编译环境变量和清理对齐问题上。

5. 性能优化：只写对还不够，要写快

5.1 先量化，再谈优化

很多新人写算子时容易凭感觉优化，看到代码有循环就觉得慢，看到同步点就想删。这是最容易翻车的做法。Ascend C 算子性能好不好，关键是看 Profile 数据里搬运和计算的占比，而不是看代码行数。

我习惯先用性能分析工具抓一遍数据。重点看三个指标：算子总耗时、向量计算单元利用率、搬运带宽占比。通常拿到这三项，优化方向基本就清晰了：如果搬运占比极高，说明瓶颈在 GM 与 UB 的数据交换，重点优化 tiling 粒度和流水深读；如果计算单元等待严重，说明同步点过多或分块太小；如果负载不均，则要调整核间数据切分。

5.2 优化从三处下手：tiling、流水、尾块

以 FusedAddScale 为例，优化的顺序和内容大概是这样：

• 调 tiling 粒度：单核处理长度太大，UB 放不下，必须多轮循环；单核处理长度太小，循环次数变多，同步和队列管理开销增大。找到一个让循环次数和数据搬运效率平衡的值，这是最基础也最提升最大的一项。
• 加深流水线：把BUFFER_NUM从 1 改成 2 或 3，让搬运、计算、写回三级操作尽量重叠。数据量大的时候，双缓冲的收益很明显。
• 处理尾部不对齐块：对齐块用向量指令高效计算，尾部小块单独处理，避免把整块逻辑拖进慢路径。

还有一个容易被忽略的点：多核并行度。如果数据总量很小，盲目开满所有核反而会因为线程调度和空转增加耗时。tiling 时要有一个下限阈值，数据不足一定量时只开少量核，甚至单核处理更快。

5.3 实测参考数据

我用 FusedAddScale 做过一组简单的对照测试，输入为 float16，数据量 256K 个元素，单核朴素实现、多核对齐块实现、多核双缓冲实现三者差距明显。多核对齐块相比单核朴素实现，耗时下降明显；再叠加双缓冲，整体耗时又进一步降低。

这组数不绝对，不同芯片型号、不同数据量、不同编译选项都会有差异，但它足够说明问题：Ascend C 算子上限的提升，主要来自搬运和计算的重叠，以及核间负载的平衡。这也解释了为什么很多人把同一个逻辑从框架层下沉后，性能并没有立刻翻倍——如果没有正确设计流水线和 tiling，自研算子可能只比拼接快一点，甚至更慢。

5.4 不值得优化的场景

写算子也要学会放弃。如果数据量只有几千个元素，启动算子的调度开销远大于计算本身，这时优化得再好都是白搭。这种场景正确的做法不是优化，而是从根上避免形成独立算子：把它打进前后算子里一起融合，或者干脆保持在 Host 侧做轻量计算。

判断标准也很简单：用 Profile 数据看 kernel 耗时中启动和调度的占比。如果占比超过一半，这个算子就不值得单独存在。我在实际开发里做过一个微型算子，单独测耗时不到 0.05ms，但放在完整链路里，调度开销却占比很高，最后把它融合进目标算子里，链路整体收益比单独优化这个小算子明显得多。

6. 踩坑记录：新手最容易卡住的几个问题

6.1 长度不对齐导致的数据越界

症状是搬运时偶发报错，或者输出张量尾部数据全是垃圾值。根因是输入长度不是 16 的整数倍，导致最后一块数据的搬运长度落在 32 字节对齐边界之外。

对策有两类：一类是在 tiling 阶段把每核长度向上对齐，然后把总长度补到一个安全边界；另一类是使用带边界处理语义的搬运接口处理最后的残余块，它能处理不对齐的拷贝。我实际用下来，推荐的做法是主体大块使用普通搬运，最后残余部分单独走边界处理。这样既能保住对齐数据的性能，又能兜底任意长度输入。

6.2 同步点丢失导致偶发结果错误

最迷惑人的问题就是偶发错误：同一个算子同一批数据，多跑几次，大部分时候正确，偶尔全错。第一次遇到时我一度怀疑是编译器问题，后来才发现是漏了同步点。

排查思路是这样的：在Compute之前必须有SyncIO，否则 GM 数据可能还没到 UB；在CopyOut之前必须有SyncCompute，否则计算结果可能还没写回 UB。如果使用双 buffer，还要注意同步点不能加在 EnQue 和 DeQue 中间，否则会阻塞流水线。同步的本质是保证数据生产者完成、数据消费者才能开始。我在代码里习惯把同步点写到注释里，明确标注“等待 GM 到 UB 搬运完成”“等待向量计算完成”，避免后续维护时改动顺序。

6.3 不方便调试时如何定位

Ascend C 的核函数里不能像 Host 代码一样随手printf打日志，真遇到结果不对，定位起来很痛苦。我的做法是三步走：

1. 在 Host 侧构造随机输入和期望输出，期望输出用 CPU 上的朴素实现计算。
2. 把算子输出取回来，和期望输出做逐元素比对。
3. 如果整体不一致，逐步缩小数据范围，先测小 size，再测对齐 size，定位是通段逻辑问题还是边界问题。

比对时要特别注意浮点精度。float16 本身精度有限，(A + B) * scale的中间结果有舍入，不能用绝对相等来判断，要用相对误差和绝对误差的混合阈值。我一般用allclose模式，rtol 设为 1e-2，atol 设为 1e-3，先把大问题暴露出来。

6.4 接口语义理解偏差

Ascend C 提供的向量接口很多，比如Add、Muls、DataCopy、DataCopyPad，它们各自的入参顺序、是否支持 in-place、是否要求长度对齐，在不同版本上还有细微差异。最保险的做法是写完之后专门翻一遍接口说明文档，而不是靠记忆写。

我自己踩过的一个具体问题是把Muls的入参理解反了，写成了对源张量做修改但传参顺序不对，编译能过，结果却始终不生效。后来对照文档逐条核对入参语义才定位到。接口问题通常不会报错，只会在结果上露出马脚，所以新手阶段宁可多花半小时翻接口定义，也不要急着跑大数据量。

最后再说几句

写 FusedAddScale 这个例子的过程，几乎把 Ascend C 算子开发的主要环节都过了一遍：从需求判断、工程生成、tiling 设计、核函数编码，到编译验证、性能分析、问题排查。回过头看，真正难的从来不是某个接口怎么调，而是搞清楚数据在哪个环节、什么时候到位、谁负责搬运、谁负责计算。把存储层级和同步关系理清楚了，再复杂的算子也是这套逻辑的延伸。

如果你也在算子库覆盖不到的场景里挣扎，建议从类似 FusedAddScale 这样的小算子开始，先跑通完整链路，再逐步加形状约束、动态 shape、多输入融合。后面如果再遇到输入带维度信息、需要处理多维 tensor 的场景，只需要在 tiling 阶段把多维索引换算成一维偏移，Device 侧的搬运计算逻辑骨架基本不用大改。动手写第一个算子的过程会很折腾，但跑通的那一刻，后面就是一条顺畅的路。