我最早接触 Ascend C 算子开发，是因为跑模型时被一个简单的 LayerNorm 卡住了性能。框架自带的算子能用，但融合度不够，读一遍数据做均值，再读一遍做方差，最后再来一遍归一化，Global Memory 被反复读写，耗时肉眼可见地涨。那时候我就意识到，如果想在 NPU 上拿到真正的性能，绕不开手动写算子这一步。这篇文章就从我自己的实战经历出发，完整讲一遍 Ascend C 算子开发从环境搭建、工程结构、双域编程模型，到写算子、部署验证、调优避坑的完整链路。适合刚接触算子开发、被文档和各种抽象概念绕晕的开发者，也适合已经在写但总踩坑、想系统梳理一遍的人。
1. 为什么框架自带算子不够用：算子开发解决的真实问题
1.1 哪些情况必须自己写算子
很多同学一开始都跟我一样有疑问：框架里nn.LayerNorm、nn.Softmax、nn.GELU什么都有，为什么还要自己折腾算子开发？
答案很简单：框架内置算子只是“能用”，不是“好用”。先说最常见的场景——算子缺失。模型里用了新出的激活函数、量化算子、自定义 Loss，或者某个特殊的归一化方式，框架里根本没有对应实现，这就必须自己写。PyTorch 社区里经常用一堆小算子组合拼出一个功能，在 GPU 上跑没问题，因为 GPU 的小算子启动开销还能忍受，但在 NPU 上，算子启动、数据搬运、内存分配的成本都更高，组合算子多跑一步就是多一份损耗。
第二类场景是算子融合。LayerNorm 看起来只是一个算子，实际上内部有三步逻辑：算均值、算方差、归一化。如果框架把它拆成两个或三个算子串起来，每个算子都要把中间结果从 Global Memory 读出来、写进去，数据量翻了好几倍。自己写一个融合算子，把中间结果放在片上高速存储里，一次数据搬运做完所有事情，性能直接上几个台阶。
第三类场景是特殊形状和逻辑。比如动态 shape、稀疏计算、布尔掩码、多维归约这类需求，框架实现往往不够灵活，或者边界条件处理得一塌糊涂。自己写算子时，Shape 推导、数据切分、逻辑分支全部握在自己手里，遇到什么形状都能兜住。
我当时就是被一个组合算子逼到走投无路，才开始认真学 Ascend C。学完之后回头看，算子开发并不是什么“底层硬核黑科技”，它其实是一套很成熟的工程范式，关键是把几个核心概念搞清楚：数据放在哪、计算单元长什么样、数据怎么搬运、并行怎么切分。搞懂这四件事，剩下的就是写代码。
1.2 写算子前先算一笔“性能账”
不是所有算子都值得手写。动手之前，我建议你先拿笔算一算：这个算子到底是计算密集还是访存密集。
拿矩阵乘法来说，计算量是 (2 \times M \times N \times K)，访存量是 (M \times K + K \times N + M \times N)。当 M、N、K 都很大时，计算量远大于访存量，这叫计算密集。这种算子的性能瓶颈在计算单元上，重点关注乘法单元、卷积单元利用率。
而 LayerNorm、Softmax、各种激活函数，计算量只有 (O(N))，访存量也是 (O(N))。看起来平衡，但对比一下算力：AI 处理器上的向量单元每秒能算几十到几百 T 的 FLOPs，而数据搬运带宽只有可怜的几百 GB/s 到几 TB/s。也就是说，计算单元大部分时间在“等饭”——等数据从内存搬过来。这类算子就是访存密集，瓶颈在搬运带宽上。
打个比方：计算单元是个厨子，切菜很快，但食材在仓库里，厨师每次做菜都要坐车去仓库取一次。如果你把一次做菜的流程拆成“取面粉”“取水”“取鸡蛋”三趟，每趟都要来回跑仓库，那当然慢。融合算子的思路就是，提前把所有食材一次性取到厨房的操作台上，厨子在台上切切炒炒，最后把成品端出去。
所以我的建议是：访存密集的算子，优先考虑融合和减少搬运次数；计算密集的算子，优先考虑切分、流水并行、用上专用计算单元。Ascend C 算子绝大多数场景都属于前者——矢量算子的主场。
2. Ascend C 的工程骨架：从目录结构到上板运行的完整链路
2.1 一个标准算子工程的目录结构长什么样
我第一次新建 Ascend C 算子工程时，最大的感受是“怎么这么多文件”。习惯了写 Python 一行行做实验，突然面对 C++ 头文件、算子信息库、Tiling 实现、Host 侧调用代码，确实有点懵。但搞清楚每个文件干什么之后，其实结构非常清晰，甚至有点固定的套路感。
一个典型的算子工程大概长这样：
layer_norm/ ├── CMakeLists.txt ├── op_host/ │ ├── layer_norm.cpp # host侧：Shape推导、Tiling数据准备 │ ├── layer_norm_tiling.cpp # 算子的切分策略计算 │ └── layer_norm_tiling.h ├── op_kernel/ │ ├── layer_norm.cpp # device侧：真正的NPU kernel代码 │ └── layer_norm_tiling.h # tiling入参数据结构 └── op_info.json # 算子信息描述：输入输出、属性等跟 GPU 上只需要写一个kernel.cu就完事不同，Ascend C 明确把算子开发拆成 Host 侧和 Device 侧两个部分。Host 侧负责计算“这一层数据怎么切块”，Device 侧负责“每一块怎么算”。中间通过 Tiling 数据结构沟通。
我记得第一次看到op_info.json时完全不知道它是干嘛的。其实它约等于算子的“户口本”，框架要靠它知道这个算子接收几个输入、输出几个 Tensor、每个输入是什么数据类型、有没有属性参数、需不需要动态 Shape 推导。文件里少了字段，算子加载时就会报错。
编译环节一般是 CMake + 自带的编译命令，环境变量准备好之后，命令行跑一下编译脚本，会自动生成算子动态库.so和 kernel 侧的可执行二进制。把所有产物拷贝到算子安装目录，框架就能通过aclnn_layer_norm这类接口来调用了。
2.2 从源码到上板：运行链路里最容易被忽略的环节
很多人写完算子源码，编译过了，却在调用阶段各种报错，原因往往在于没有真正理解运行链路。
算子从源码到在设备上跑通，要穿过的链路大致是：
- 编写 Host 侧代码和 Kernel 侧代码，分别编译，生成动态库。
- 应用进程启动，初始化设备，加载算子动态库。
- 框架算子调用层把输入 Tensor 的 shape、dtype 传给算子的 Shape 推导逻辑。
- 推导出输出 Tensor 的 shape 后，分配输出内存。
- 调用 Host 侧 Tiling，根据输入 shape 和片上资源计算切分参数。
- 启动 Kernel，把 Tiling 参数和输入输出地址传给设备侧。
- 设备侧根据 Tiling 参数，逐步把 Global Memory 里的数据搬到片上、计算、搬回。
这个链路里最容易出问题的，是第 5 步和第 6 步衔接的地方：Tiling 参数和 Kernel 里用的数据结构必须完全一致。我踩过的坑是，Host 侧结构体里定义了一个uint32_t blockDim，Kernel 侧定义了一个int32_t blockDim，两个文件各自编译，编译时不可能互相检查，运行时就数据错乱、维度对不上。这类问题在 CPU 编程里几乎不会发生，但在 Host/Device 分离编译的模型下特别容易踩。我的习惯是，把 Tiling 结构体定义放在一个单独的头文件里，Host 和 Kernel 都 include 同一个头文件，保证定义永远一致。
另一个被我忽略了很久的环节是 Shape 推导。Ascend C 里，Host 侧InferShape函数决定输出的 shape 和 dtype。有一次我写了归一化算子，输入是[batch, seq_len, hidden]，输出应该一模一样，却忘了给输出shape赋值，调用时报“输出 shape 为空”的错误。这种问题本身不难，但如果你没有把链路理解清楚，会被报错信息带偏。
3. Host 与 Device 双域模型：搞懂 CPU 侧和 AI 核心侧的分工
3.1 内存层级：GM、UB、LocalTensor 到底是什么
写 Ascend C 算子，离不开三个关键存储概念：GM（Global Memory）、UB（Unified Buffer）、LocalTensor。
GM 就是 NPU 上的全局内存，等效于 GPU 的显存，容量大，但访问速度慢。你能改写的原始输入数据、要输出的结果，都在 GM 里。在 Ascend C 代码里，你操作的 GM Tensor 通常长这样：
__global__ void layer_norm_kernel(GM tensor input_x, GM tensor input_gamma, GM tensor input_beta, GM tensor output_y)这里的GM tensor可以理解成“指向全局内存的一块数据视图”，支持按形状读取数据，也可以直接对里面的元素做简单操作。
UB 才是真正干活的地方，它是一块片上高速存储，容量很小，可能只有几百 KB。所有复杂的向量计算，都要先把数据从 GM 拷到 UB，计算完再拷回 GM。UB 上的数据通过LocalTensor访问。
我常说，GM 相当于仓库，UB 相当于厨房操作台。操作台只有这么大，不可能把整批货物都摆上去，所以每次只能切一部分，反复地搬、切、算。
这里有一个特别重要的点：在 UB 上声明 LocalTensor，你不只是“声明一块内存”，而是在申请一个明确的内存生命周期。用TPipe管理时还要跟队列（Queue）绑定。第一次学的时候我完全没意识到这一点，直接拿LocalTensor的地址乱算，结果数据互相覆盖，算出来的结果一塌糊涂。
Ascend C 还有个容易混淆的地方：很多 API 操作的是LocalTensor，但你要先通过 TQue 从队列里 DeQue 拿到这块 LocalTensor 的使用权，算完再 EnQue 交回去。
3.2 队列模型与 TPipe：让数据搬运和计算重叠起来
再进一步，Ascend C 的TPipe是一个二维的队列灌水模型。
拿流水线来说，CopyIn 就是把仓库的货搬到厨房台面（UB），Compute 是切菜，CopyOut 是把成品装车搬回仓库。如果只能搬一箱、切一箱、装一箱，整个流程永远是串行的。有了队列和多块缓冲（Double Buffer）之后，可以一边搬下一批，一边切当前批，一边装上一批，三者重叠起来了。
代码里的大体框架如下：
class LayerNormKernel { public: __aicore__ inline LayerNormKernel(GM tensor in_x, GM tensor in_gamma, GM tensor in_beta, GM tensor out_y, const LayerNormTilingData& tiling) : x_gm(in_x), gamma_gm(in_gamma), beta_gm(in_beta), y_gm(out_y), tiling(tiling) {} __aicore__ inline void Init(GM tensor in_x, GM tensor in_gamma, GM tensor in_beta, GM tensor out_y) { pipe.InitBuffer(tmp_buf, block_size * sizeof(float)); x_que.InitBuffer(block_size * sizeof(float)); // ... 其他队列初始化 } __aicore__ inline void Process() { int32_t loop = tiling.totalLength / block_size; for (int32_t i = 0; i < loop; i++) { CopyIn(i); Compute(i); CopyOut(i); } } private: __aicore__ inline void CopyIn(int32_t i) { LocalTensor<float> x_local = x_que.AllocTensor<float>(); DataCopy(x_local, x_gm[i * block_size], block_size); x_que.EnQue(x_local); // gamma、beta 同理 } __aicore__ inline void Compute(int32_t i) { LocalTensor<float> x_local = x_que.DeQue<float>(); // 这里对 x_local 做向量计算 // 比如 ReduceSum、Muls、Add 等 y_que.EnQue(x_local); } __aicore__ inline void CopyOut(int32_t i) { LocalTensor<float> y_local = y_que.DeQue<float>(); DataCopy(y_gm[i * block_size], y_local, block_size); y_que.FreeTensor(y_local); } private: TPipe pipe; TQue<QuePosition::VECIN, 1> x_que; TQue<QuePosition::VECOUT, 1> y_que; TBuf<QuePosition::VECCALC> tmp_buf; // GM tensor、tiling 等成员 };这段代码是简化版，但可以清楚地看出流水线的骨架：每个循环迭代里 CopyIn、Compute、CopyOut 分别操作不同队列，通过队列同步机制天然重叠——当第 i 次迭代在计算时，第 i+1 次迭代的数据搬运已经开始了。
新手最容易犯的错误是把这三个步骤写成顺序执行，或者没理解EnQue和DeQue的匹配关系。AllocTensor申请到空块后必须EnQue，DeQue拿到数据后算完一定要FreeTensor。队列深度的循环搞错，程序要么卡死，要么数据错位，而且这种问题很难通过读日志发现，因为卡死原因往往是“队列为空时去 DeQue”，就像去空银行柜台取号，永远等不到叫号。
4. 实战：手写 LayerNorm 算子的完整代码走读
4.1 数学拆解：把 LayerNorm 变成一串向量操作
在写代码之前，先用数学语言把 LayerNorm 拆干净。给定输入 (x \in R^{B\times S\times H})，对最后一维做归一化：
[ \mu_{b,s} = \frac{1}{H} \sum_{h=1}^{H} x_{b,s,h} ] [ \sigma^2_{b,s} = \frac{1}{H} \sum_{h=1}^{H} (x_{b,s,h} - \mu_{b,s})^2 ] [ y_{b,s,h} = \frac{x_{b,s,h} - \mu_{b,s}}{\sqrt{\sigma^2_{b,s} + \epsilon}} \cdot \gamma_h + \beta_h ]
计算流程可以转成向量指令的“三步走”：
- 对 H 维 ReduceSum，得到每个 (b, s) 行的和 (sum)，再乘 (1/H) 得到均值。
- 用均值广播算出方差：先把 x 拷贝一份，逐元素减均值，再逐元素平方，再 ReduceSum，乘 (1/H)。
- 归一化：逐元素减均值、乘方差的倒数、乘 gamma、加 beta。
可以看出，计算量不大，关键是减少对 GM 的读写。设计成融合算子，中间结果全放 UB，数据只需从 GM 读一次 x、gamma、beta，写一次 y。相比拆成三个算子，GM 访问量直接减少到三分之一以下。
4.2 Host 侧代码：Shape 推导与 Tiling 计算
我们先来看 Host 侧核心代码，它负责回答两大问题：输出长什么样、每个核处理多少数据。
#include "register/op_def_registry.h" #include "tiling/platform/platform_ascendc.h" namespace { constexpr int32_t BLOCK_SIZE = 256 * 1024; // 256KB 一个 block } void LayerNormTiling(const gert::TilingContext* ctx) { auto x_shape = ctx->GetInputShape(0)->GetStorageShape(); int32_t totalLength = 1; for (auto dim : x_shape) totalLength *= dim; LayerNormTilingData tiling; tiling.totalLength = totalLength; tiling.blockDim = 8; // 实际应根据核数决定 tiling.blockSize = BLOCK_SIZE / sizeof(float); ctx->SetTilingData(&tiling, sizeof(tiling)); ctx->SetBlockDim(tiling.blockDim); } REGISTER_OP("LayerNorm") .INPUT("x", TensorType({DT_FLOAT})) .INPUT("gamma", TensorType({DT_FLOAT})) .INPUT("beta", TensorType({DT_FLOAT})) .OUTPUT("y", TensorType({DT_FLOAT})) .OP_TYPE("LayerNorm") .SHAPE_INFER(shape_infer_func) .TILING(LayerNormTiling);按我的经验，Host 侧最容易被坑的地方是blockDim 的选择。它不是拍脑袋定的，要根据设备上的 AI 核数量、普通数据量来决定。数据量非常小的时候，开 8 个核反而会因为切分太碎而降低效率；数据量大的时候，blockDim 开少了又吃不满算力。稳妥做法是：先拿platform_ascendc.h里的核数接口取到实际核数，再结合totalLength / 每个核能处理的元素数来定。
Shape 推导函数这里也提醒一下：必须显式设置输出形状。算子输入[2, 4, 128]与[2, 128, 4]的维度含义不同，虽然名字都是x，Shape 推导代码必须跟 Kernel 内部的数据切分逻辑保持一致，否则调用时输出的内存布局就是错的。
4.3 Kernel 侧代码：归约、广播与逐元素运算
Kernel 侧代码是真正跟 NPU 矢量单元打交道的地方。简化后的核心逻辑长这样：
__global__ void layer_norm_kernel(GM tensor x, GM tensor gamma, GM tensor beta, GM tensor y, const LayerNormTilingData tiling) { int32_t blockIdx = GetBlockIdx(); // 每个 AI 核处理总数据的一段 int32_t perBlock = tiling.totalLength / tiling.blockDim; int32_t start = blockIdx * perBlock; TPipe pipe; TQue<QuePosition::VECIN, 2> inQueue; TQue<QuePosition::VECOUT, 1> outQueue; TBuf<QuePosition::VECCALC> tmpBuf; LocalTensor<float> xLocal = inQueue.AllocTensor<float>(); DataCopy(xLocal, x[start], perBlock); inQueue.EnQue(xLocal); LocalTensor<float> xForCompute = inQueue.DeQue<float>(); // 第一步：算各行的 sum // 这里按 H 维做归约，伪代码如下 // 需要把 perBlock 按照 H 拆成若干 row，逐 row 处理 // 第二步：算 variance // 先算临时向量 = (x - mean)^2，再归约 // 第三步：归一化 // y = (x - mean) * rsqrt(var + eps) * gamma + beta outQueue.EnQue(xForCompute); LocalTensor<float> yLocal = outQueue.DeQue<float>(); DataCopy(y[start], yLocal, perBlock); outQueue.FreeTensor(yLocal); }上面代码是为了让你看清结构，实际实现里我会把归约、方差、归一化拆成几个内联函数，并仔细处理每个 AI 核只负责自己那一段的边界。
这中间最有技术含量的一步是“按行归约”。如果把二维输入当成一维数组，每批就是一个 H 的区间，你需要在这个区间上做 ReduceSum。Ascend C 提供了ReduceSum之类的接口，可以直接在 LocalTensor 上做归约，配合Broadcast把均值广播回去。但要注意归约结果存在临时缓冲区，别和原始输入缓冲区重叠，否则你算着算着原始数据就没了。
我第一版就是在原地算方差时直接干掉了 x——步骤二是 x 减均值，中途用了临时缓冲临时算均值，结果下一个循环 x 已经不是原值了，怎么调数值都不对。最后老老实实把“临时均值缓冲”“临时方差缓冲”“归一化输出缓冲”分成三个独立内存块，才一步到位。先不说性能，正确性永远是第一步。
4.4 上板验证：数值对齐才是真的跑通
写完之后不代表结束，还要验证数值是否正确。我的验证套路是这样：
- 先用 CPU 上的 numpy 计算一个标准结果，保存成
.npy文件。 - 用随机初始化把输入数据填充成固定随机种子，同样生成
.npy。 - 调用刚编译的算子，输出结果。
- 两个结果做绝对误差比较，要求误差在 1e-3 量级。
因为浮点数运算顺序不同，CPU 和 NPU 的计算误差是必然的，所以不能要求完全一致。我习惯用allclose的宽松模式，rtol=1e-2, atol=1e-3。如果误差达到 1e-1，那基本就是逻辑错误或者数据错位，赶紧回头查 Tiling 和边界。
也可以参考社区里常用的调试手段：在 Kernel 里加临时逻辑，把中间结果写回 GM 的临时输出 tensor，再来和 CPU 参考值逐段对齐。比如我先只输出 sum，跟 numpy 的 sum 比，对上了再输出 variance，一路分段验证，定位非常快。这个方法比对着代码空想高效一百倍。
5. 调试链路与性能剖析：把问题从 NPU 里挖出来
5.1 为什么 NPU 调试比 CPU 痛苦：数据不可见与报错模糊
在 NPU 上调试算子的最大痛点是数据不可见。你在 Kernel 里写的代码运行在 AI 核上，不能随便 printf，也没有 IDE 可以打断点看变量。报错信息往往也很模糊，比如“数据越界”“队列溢出”“同步错误”，你根本不知道是哪一行触发的。
我的经验是把调试分成两个阶段：
第一阶段是 host 侧调试。shape 推导、tiling 计算这些逻辑在 CPU 上就能跑，直接断点调试、打印都行。很多初始化错误、结构体不一致、shape 推错问题，这个阶段就能抓到一大半。
第二阶段才是 device 侧调试。这时候我固定使用“三件套”：临时中间结果回写 GM、手动分段校验、以及二分定位。
具体来说，我在调试阶段会在 Kernel 里故意把某个中间结果写到一个额外的 GM 输出位置，然后 host 侧把它读出来跟参考值对比。举个例子，怀疑是均值算错了，那就只输出均值；怀疑是方差算错了，那就只输出方差。用“分而治之”的思路，把整个算子拆成多个独立的小函数，每个函数都单独验证，最后再拼起来。这个办法在 FKernel 代码里看起来有点笨，但配合外部 Python 脚本对比，几十行就能定位到根因。
5.2 性能剖析思路：先看时间占比，再决定优化方向
数值正确之后，下一步就是性能。跑 profiler 一看，最直观的指标就是算子总耗时。但总耗时会骗人，你必须拆解出三个时间组成部分。
一是计算时间，即 AI 核上执行向量指令的时间；二是搬运时间，即 GM 和 UB 之间数据存取的时间；三是同步与空闲时间，即队列等待、核间同步造成的空转。
有了三段式数据，优化方向就非常清晰：
- 如果计算时间占比高，重点查指令数量、for 循环开销，考虑能不能用更高阶的向量指令一次做更多操作。
- 如果搬运时间占比高，重点查数据量、访问模式，看能不能合并搬运、调整缓冲区大小、减少中间结果回写 GM。
- 如果空闲时间占比高，说明流水线没有重叠起来，可能是队列深度不够，或者各处 EnQue/DeQue 卡住了同步。
举个例子，我之前写的一个算子，profiler 显示整体耗时 800ms，但真正计算只花了 100ms，剩下全在等数据。数据搬运量是 1.2GB，而设备带宽按理可以 300GB/s 以上，理想搬运时间应该不到 50ms。问题出在 Kernel 里一维度地串行 CopyIn、Compute、CopyOut，完全没重叠。改成双缓冲并发之后，搬运时间缩短到 120ms，总耗时直接砍到原来的四分之一。这个优化没有改任何数学逻辑，纯粹是重新安排了流水线，但效果立竿见影。
性能剖析还有一个小技巧：先用简单的带宽维度估算上限。假设算子要处理 (N) 个浮点数，每个 4 字节，总数据量为 (4N) 字节，理论耗时下限就是 (4N / BW)。把实测时间和这个理论值一比，就知道离带宽上限还有多远。如果实测只有理论值的 20%，那一定还有大量的等待和搬运浪费可以压缩。
6. 那些文档没有写但实战一定会遇见的坑
6.1 对齐问题：UB 存储不是随便怎么放都可以
写 Ascend C 算子，稍微复杂一点的数据布局就会碰到对齐问题。UB 上的数据对齐是有要求的，通常是 32 字节甚至更高对齐，具体取决于指令类型和数据类型。我踩过最典型的坑是：结构体里塞了一个 float 数组，又塞了一个 int 字段，导致数组起始偏移不是 32 字节的整数倍，DataCopy 直接报错。
解决办法是：从 Buffer 分配时尽量用AlignUp对齐，或者干脆把多个小数组分别放到不同的 TBuf 里，避免联合布局带来的对齐灾难。数据对齐问题往往只在你把多个变量塞到同一块缓冲时才出现，单一大数据块反而不容易踩坑。
另一个对齐问题是元素个数。UB 处理数据时很多指令要求一次处理固定数量（比如 64 个 float）的整数倍，如果数据量不是指令粒度的整数倍，最后一笔怎么处理就得额外写边界逻辑。千万别贪图省事直接超处理几个元素，后面会导致内存越界写坏别的数据。
6.2 Tiling 与边界数据：最后一个 Block 到底归谁管
很多人第一次写 Kernel 都会假设totalLength / blockDim刚好整除。但真实世界的数据形状千奇百怪，大概率有余数。这个余数不处理好，要么丢数据，要么越界。
我的处理模板是：
int32_t perBlock = tiling.totalLength / tiling.blockDim; int32_t remain = tiling.totalLength % tiling.blockDim; int32_t start = blockIdx * perBlock; int32_t curLen = perBlock; if (blockIdx == tiling.blockDim - 1) { curLen += remain; }看似很简单，但很多人栽在这里，因为 Kernel 内部循环的循环次数和CopyIn的搬运长度都必须跟着curLen走，不能只用固定的blockSize。另外，如果remain本身很大，而且大于 UB 单次可处理的块大小，还要继续把curLen二次切成小片，循环处理。Tiling 的目标永远是“每个核都能正确处理自己的那一段”，而不是“所有核处理量完全一样”。
还有一点很多人忽略：gamma、beta 这类与最后一维等长的张量，在所有核之间其实是共享的。如果按一维平铺数据，第 0 号核处理前 1024 个元素时，gamma 要取前 128 个，第 1 号核处理 1024~2047 个元素时，gamma 又要重新从 GM 取一遍同一份 gamma。这个读法在逻辑上没问题，但会让 GM 访问量翻倍。优化方案是：先把 gamma、beta 各拷一份到 UB，每个核重复使用，或者如果片上放得下，就常驻在 UB 里。这个优化细节，文档里很少明确写，但实测效果很明显。
6.3 数值正确之后再谈优化：先对齐再提速
这是我不只一次重复的教训：先保证正确，再谈性能优化。很多优化手段（比如双缓冲、循环展开、手工内存排布重组）会大幅增加代码复杂度，一旦结果对不上，排查难上加难。
我的工作流是：
- 第一版先写朴素实现，保证数值正确，能通过测试用例。
- 用 numpy 参考做好回归测试脚本，后续每次改动都自动跑到。
- 跑 profiler 找出瓶颈，根据瓶颈做一轮优化。
- 优化完立刻跑回归，确认数值没变坏、误差还在容忍范围内。
- 保存好每个版本的代码，优化过程用 git 记录，方便随时回溯。
有了这套流程，即使后来我把代码改得比较花哨，也一直能保持结果稳定。某同事有一次直接把 double buffer 和循环展开一起上，本地过了，大规模测试就对不上了，最后查了两天，问题出在队列深度不够导致数据竞争。如果当时没有保存朴素版本做对照，可能崩溃得更久。
6.4 一个实用经验：把中间结果打印到 Host 侧
最后分享一个小技巧。Ascend C 的 Host 侧代码运行在 CPU 上，你有完整打印能力。我习惯在 Kernel 的临时 GM 输出里留一个 dump 槽位，调试时把中间结果分批次填进去，Host 侧统一打印成 CSV 或者直接转成 numpy 比较。
具体做法是在 TilingData 里加一个debugFlag字段，正常推理时置 0，调试时置 1。Kernel 里根据标志决定要不要把某个中间结果写到额外的 GM 空间。这个开关可以留着不删，也算是开发期的一个“后门”，测试数据下一跑就知道是哪个环节出了问题。
如果当前平台支持 host 侧打印，也可以直接把某个核的局部值传到 host，打印出来看。这比读一堆十六进制内存高效太多。不过要注意，这只能用来调试，千万别在生产路径里留打印，打印本身极其耗时，会严重拖慢算子性能。
我自己的经验是，算子开发的难度其实不在写代码本身，而在“看不见数据的世界里保持逻辑准确”。Ascend C 把 Host 和 Device 分得清清楚楚，逼着你把数据布局、切分策略、内存搬运全都显式地思考一遍。这个过程初看繁琐，但一旦你理解了队列模型和双域协作的逻辑，后面写新算子就是替换计算部分的事。下次遇到框架自带的算子在 NPU 上跑不满带宽，不妨自己动手写一个融合版本，你会看到性能数字带来的那种实实在在的爽感。
