CANN算子开发核心概念全解析

学了Ascend C的语法，能看懂代码了，但写起算子来还是云里雾里。什么是Tiling？为什么要分块处理？Kernel函数到底是怎么运行的？这些问题困扰了我好几天。直到在训练营听老师系统讲解，并且自己动手实践了几个算子，才真正理解这些概念。今天就把这些核心概念梳理清楚，相信能帮大家少走不少弯路。算子（Operator）就是神经网络中的一个计算单元。

qq_41397792

9502人浏览 · 2025-10-24 12:19:26

qq_41397792 · 2025-10-24 12:19:26 发布

CANN算子开发核心概念全解析

昇腾训练营报名链接：https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro

训练营简介：2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机、平板、开发板等大奖。

前言

学了Ascend C的语法，能看懂代码了，但写起算子来还是云里雾里。什么是Tiling？为什么要分块处理？Kernel函数到底是怎么运行的？

这些问题困扰了我好几天。直到在训练营听老师系统讲解，并且自己动手实践了几个算子，才真正理解这些概念。今天就把这些核心概念梳理清楚，相信能帮大家少走不少弯路。

一、什么是算子？

1.1 算子的定义

算子（Operator） 就是神经网络中的一个计算单元。

回顾一下CANN的整体架构（来自官网 https://www.hiascend.com/software/cann）：

CANN架构图

从架构图可以看到，算子库是CANN异构计算架构的重要组成部分，包含大模型融合算子和NN/CV/Math基础算子。

比如：

# PyTorch中的算子调用
import torch
x = torch.tensor([1.0, 2.0, 3.0])
y = torch.tensor([4.0, 5.0, 6.0])
z = torch.add(x, y)  # 这里的add就是一个算子

在底层，这个torch.add最终会调用硬件上的加法算子实现。在GPU上可能调用CUDA Kernel，在昇腾NPU上就调用CANN算子。

1.2 算子的分类

根据功能，算子大致分为：

我们训练营主要学的是计算类算子，这也是性能优化的重点。

1.3 算子的输入输出

一个算子通常包括：

输入Tensor：一个或多个
输出Tensor：一个或多个
属性参数：如卷积的stride、padding等

举个例子：

// Add算子
// 输入: x, y (两个Tensor)
// 输出: z (一个Tensor)
// 属性: 无
void Add(Tensor x, Tensor y, Tensor z);

// Conv2D算子
// 输入: input, weight, bias
// 输出: output
// 属性: stride, padding, dilation, groups
void Conv2D(Tensor input, Tensor weight, Tensor bias, Tensor output,
            int stride, int padding, ...);

二、Kernel函数：算子的核心

2.1 什么是Kernel？

Kernel是运行在AI Core上的函数，可以类比CUDA Kernel。

// Kernel函数定义
extern "C" __global__ __aicore__ void MyKernel(GM_ADDR input, GM_ADDR output, uint32_t size) {
    // 这段代码在AI Core上并行执行
}

关键点：

extern "C"：C链接，方便外部调用
__global__：全局可见
__aicore__：运行在AI Core上

2.2 Kernel的执行模型

NPU可以启动多个AI Core同时执行同一个Kernel，每个Core处理不同的数据块：

每个Core通过GetBlockIdx()知道自己是哪个核，从而处理对应的数据块。

// 获取当前AI Core的索引
uint32_t blockIdx = GetBlockIdx();

// 获取总的AI Core数量
uint32_t blockNum = GetBlockNum();

// 计算当前Core要处理的数据范围
uint32_t blockSize = totalSize / blockNum;
uint32_t startIdx = blockIdx * blockSize;
uint32_t endIdx = startIdx + blockSize;

2.3 我的理解

可以把Kernel想象成一个"工作模板"，NPU就像一个有多个工人的工厂。启动Kernel时：

把这个"工作模板"分发给每个工人（AI Core）
每个工人处理一部分数据
所有工人并行工作
最后把结果汇总

这种模型叫SIMT（Single Instruction, Multiple Threads），和CUDA很类似。

三、Tiling策略：为什么要分块？

3.1 问题的由来

假设我们要处理一个1GB的数据，但NPU的Unified Buffer（UB）只有几MB。怎么办？

答案是：分块处理（Tiling）！

3.2 Tiling的基本思想

把大数据切成小块（Tile），每次只处理一块：

原始数据: [1000000个元素]

分块后:
Tile 0: [0-255]
Tile 1: [256-511]
Tile 2: [512-767]
...
Tile N: [...]

每次处理流程：

3.3 Tiling的代码实现

// 总数据量
uint32_t totalLength = 10000;

// Tile大小（通常选择32的倍数，FP16时）
constexpr uint32_t TILE_SIZE = 256;

// 需要多少个Tile
uint32_t tileNum = totalLength / TILE_SIZE;  // = 39

// 剩余不足一个Tile的数据
uint32_t remainLength = totalLength % TILE_SIZE;  // = 16

// 循环处理每个Tile
for (int32_t i = 0; i < tileNum; i++) {
    // 1. 拷贝第i个Tile到UB
    DataCopy(localBuffer, globalBuffer[i * TILE_SIZE], TILE_SIZE);
    
    // 2. 在UB上计算
    Compute(localBuffer, TILE_SIZE);
    
    // 3. 拷贝结果回GM
    DataCopy(outputBuffer[i * TILE_SIZE], localBuffer, TILE_SIZE);
}

// 处理剩余数据
if (remainLength > 0) {
    ProcessTail(remainLength);
}

3.4 选择合适的Tile Size

Tile大小的选择很关键，我的经验：

// 太小（如TILE_SIZE = 32）
// ❌ 问题：循环次数太多，开销大
// ❌ 问题：无法充分利用向量指令

// 太大（如TILE_SIZE = 10000）
// ❌ 问题：可能超过UB容量
// ❌ 问题：多AI Core负载不均

// 合适的大小（TILE_SIZE = 256/512/1024）
// ✅ 既能充分利用UB
// ✅ 又能保证向量化效率
// ✅ 多核负载均衡

常用的Tile大小：

FP16：256、512、1024（对应4KB、8KB、16KB）
FP32：128、256、512（对应2KB、4KB、8KB）

四、双缓冲与流水线

4.1 单缓冲的问题

如果只用一个Buffer，执行流程是串行的：

可以看到，在Compute的时候，数据搬运单元是空闲的，浪费了！

4.2 双缓冲优化

使用两个Buffer，可以让数据搬运和计算并行：

可以看到，从Tile 1开始，CopyIn、Compute、CopyOut三个阶段并行执行！

4.3 双缓冲代码实现

constexpr int32_t BUFFER_NUM = 2;

class KernelAdd {
public:
    __aicore__ inline void Init(...) {
        // 初始化两个输入队列，每个队列有2个Buffer
        pipe.InitBuffer(queueX, BUFFER_NUM, TILE_SIZE * sizeof(half));
        pipe.InitBuffer(queueY, BUFFER_NUM, TILE_SIZE * sizeof(half));
        pipe.InitBuffer(queueZ, BUFFER_NUM, TILE_SIZE * sizeof(half));
    }
    
    __aicore__ inline void Process() {
        // 总循环次数 = tileNum + BUFFER_NUM
        // 多出来的是为了把流水线排空
        int32_t loopCount = tileNum + BUFFER_NUM;
        
        for (int32_t i = 0; i < loopCount; i++) {
            // Stage 1: CopyIn
            if (i < tileNum) {
                CopyIn(i);
            }
            
            // Stage 2: Compute（延迟BUFFER_NUM个周期）
            if (i >= 1 && i < tileNum + 1) {
                Compute(i - 1);
            }
            
            // Stage 3: CopyOut（延迟BUFFER_NUM*2个周期）
            if (i >= 2) {
                CopyOut(i - 2);
            }
        }
    }
};

这里有个细节：为什么要延迟？

Compute延迟1个周期：保证数据已经CopyIn完成
CopyOut延迟2个周期：保证计算已经完成

我刚开始没理解这个，导致数据还没拷进来就开始计算，结果全是错的。调试了半天才发现问题。

五、内存层次与数据搬运

5.1 NPU的内存架构

NPU有多级内存，速度和容量不同：

内存类型	位置	容量	速度	用途
DDR/HBM	板卡外部	16GB-32GB	慢（~100GB/s）	存储模型和数据
Global Memory	芯片内	几GB	较快	中转存储
Unified Buffer	AI Core内	几MB	快（~1TB/s）	计算时的数据
L0 Buffer	计算单元内	几KB	最快	寄存器级存储

5.2 数据搬运路径

典型的数据流：

我们写算子时，主要关注的是GM <-> UB这一层的搬运，用DataCopy实现。

5.3 数据搬运优化

几个优化技巧：

技巧1：合并小的DataCopy

// ❌ 不好：多次小拷贝
for (int i = 0; i < 100; i++) {
    DataCopy(dst[i], src[i], 1);  // 每次拷1个元素
}

// ✅ 好：一次大拷贝
DataCopy(dst, src, 100);  // 一次拷100个元素

技巧2：对齐访问

// 数据地址要32字节对齐（FP16）
// TILE_SIZE要是32的倍数
constexpr uint32_t TILE_SIZE = 256;  // ✅
// constexpr uint32_t TILE_SIZE = 100;  // ❌

技巧3：使用DMA异步拷贝

// 异步拷贝：不阻塞，可以和计算并行
DataCopyPad(dstLocal, srcGm, TILE_SIZE);

// 同步点：确保拷贝完成
pipe.barrier();

六、向量化编程

6.1 为什么要向量化？

NPU的计算单元是SIMD（Single Instruction, Multiple Data），一条指令可以同时处理多个数据。

举个例子：

// 标量方式（串行，慢）
for (int i = 0; i < 256; i++) {
    z[i] = x[i] + y[i];  // 每次处理1个元素，需要256次
}

// 向量方式（并行，快）
Add(z, x, y, 256);  // 一条指令处理256个元素

向量化能让性能提升几十倍甚至上百倍！

6.2 向量化的实现

Ascend C提供了丰富的向量API：

// 向量加法
Add(dst, src0, src1, count);

// 向量乘法
Mul(dst, src0, src1, count);

// 向量乘加（Fused Multiply-Add）
Mla(dst, src0, src1, src2, count);  // dst = src0 * src1 + src2

// 向量ReLU
Relu(dst, src, count);  // dst = max(0, src)

6.3 向量宽度

不同数据类型的向量宽度不同：

// FP16: 一次处理16个（128 bits）
half16 vec;

// FP32: 一次处理8个（256 bits）
float8 vec;

// INT8: 一次处理32个（256 bits）
int8x32 vec;

所以TILE_SIZE最好是向量宽度的倍数：

// FP16: TILE_SIZE应该是16的倍数
constexpr uint32_t TILE_SIZE = 256;  // ✅ 256 = 16 * 16

// FP32: TILE_SIZE应该是8的倍数
constexpr uint32_t TILE_SIZE = 256;  // ✅ 256 = 8 * 32

七、算子开发的完整流程

把前面的概念串起来，算子开发的典型流程：

具体步骤

步骤1：需求分析

算子: ElementWise Add
输入: x (FP16), y (FP16)
输出: z (FP16)
功能: z = x + y（逐元素相加）

步骤2：算法设计

# 伪代码
for i in range(n):
    z[i] = x[i] + y[i]

步骤3：Tiling策略

Tile大小: 256个元素（512字节）
Tile数量: n / 256
缓冲策略: 双缓冲

步骤4：编写Kernel

// 实现CopyIn、Compute、CopyOut三个函数

步骤5：多核并行

uint32_t blockIdx = GetBlockIdx();
uint32_t offset = blockIdx * blockLength;

步骤6：向量化

Add(zLocal, xLocal, yLocal, TILE_SIZE);  // 使用向量Add

步骤7：测试

// 对比CPU结果
// 测试性能

八、我踩过的坑总结

坑1：忘记处理tail数据

// ❌ 错误：只处理了完整的Tile
for (int i = 0; i < tileNum; i++) {
    Process(i);
}
// 如果totalLength不是TILE_SIZE的倍数，会漏掉最后的数据！

// ✅ 正确：处理tail
for (int i = 0; i < tileNum; i++) {
    Process(i);
}
uint32_t remainLength = totalLength % TILE_SIZE;
if (remainLength > 0) {
    ProcessTail(remainLength);
}

这个坑导致我的结果总是不对，调试了一整天。

坑2：多核数据重叠

// ❌ 错误：每个核都处理全部数据
uint32_t offset = 0;  // 固定从0开始

// ✅ 正确：每个核处理不同部分
uint32_t blockIdx = GetBlockIdx();
uint32_t blockNum = GetBlockNum();
uint32_t blockSize = totalLength / blockNum;
uint32_t offset = blockIdx * blockSize;

坑3：Buffer大小不够

// ❌ 错误：InitBuffer的大小小于实际使用
pipe.InitBuffer(queue, 2, 128 * sizeof(half));  // 只分配128个元素
DataCopy(localTensor, gmTensor, 256);  // 但拷贝了256个元素！

// ✅ 正确：确保Buffer够大
pipe.InitBuffer(queue, 2, TILE_SIZE * sizeof(half));