写在前面

入门课的第一个自定义算子就是它，看着简单，却是新手最容易卡住的地方。这篇把流程和坑一次讲清。

一、先看全流程

昇腾算子开发固定五步：

算子分析 → 核函数定义 → 算子类实现 → 编译部署 → 调试验证

第一步别省。把输入输出个数、数据类型、计算逻辑写清楚：

算子类型：AddCustom
输入：x, y（float16）
输出：z
公式：z = x + y
并行策略：元素级并行

二、核函数

extern "C" __global__ __aicore__ void add_custom(
    __gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z,
    uint32_t totalLength)
{
    KernelAdd op;
    op.Init(x, y, z, totalLength);
    op.Process();
}

三个关键字是理解昇腾的密码：

• __global__ — 这段跑在 AI Core 上，不是 CPU
• __aicore__ — 运行在 AI Core 内部
• __gm__ — 指针指向全局内存（GM）

三、算子类：四个阶段

真正干活的是 KernelAdd，拆成 Init / CopyIn / Compute / CopyOut。

Init — 算清楚自己该处理哪一段

uint32_t blockLength = totalLength / GetBlockNum();
uint32_t offset = GetBlockIdx() * blockLength;

GetBlockIdx() 返回当前是第几个 AI Core。你只写一份代码，几十个核同时跑，每个核 offset 不同——这就是 SPMD 模型。

CopyIn / Compute / CopyOut — 搬进来、算、搬出去

// CopyIn
LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
DataCopy(xLocal, xGm[offset + i * tileLen], tileLen);
inQueueX.EnQue(xLocal);

// Compute
Add(zLocal, xLocal, yLocal, tileLen);

// CopyOut
outQueueZ.EnQue(zLocal);
LocalTensor<half> zOut = outQueueZ.DeQue<half>();
DataCopy(zGm[offset + i * tileLen], zOut, tileLen);

为什么要这么绕？ 因为 AI Core 内部只有一块几百 KB 的 UB，你传进来的指针指向的是几 GB 外的 GM。所以必须：

GM → UB → 计算 → UB → GM

一句话：CPU 编程关注"算什么"，NPU 编程关注"数据在哪、怎么搬"。

四、CPU / NPU 双模式验证

NPU 模式

CPU 报错就是逻辑错，NPU 报错多半是环境或对齐问题——分开验证能省很多时间。

# CPU 模式
export ASCENDC_CPU_DEBUG=1
./add_custom

# NPU 模式
./run.sh --soc_version=Ascend910B

五、六个坑，都踩过

1. DataCopy 必须 32 字节对齐
   单元素 float16 占 2 字节，一次搬 16 个才对齐。总长度不能整除时，尾块要单独处理，否则直接报错。
2. 多核均分有余数
   totalLength / blockNum 除不尽时最后一个核要补上余数，不处理就是精度错误。
3. BUFFER_NUM=2 是双缓冲
   开 2 份 buffer，计算和搬运能重叠，性能明显好于 1。这是最省事的一次优化。
4. 判断条件写 blockLength，别写 totalLength
   一个核算的是 blockLength，写错就整个崩。
5. 基本 API 和高级 API 别混着调
   用 Add 就一路 Add，用 Adds 不要又调 Add，混用容易出对齐问题。
6. CPU 通过 ≠ NPU 通过
   把浮点精度问题放到最后，先用整数数据验证功能。

六、总结
Add 算子值得亲手写一遍：它把异构计算、SPMD 并行、多级存储搬运、双缓冲流水全串起来了。后面写复杂算子，套路一模一样。

`