【CANN训练营学习笔记】从0到1掌握Ascend C算子工程开发全流程

训练营简介

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

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

📖 本文思维导图

mindmap
  root((工程开发全流程))
    开发模式对比
      快速开发
        直接编写kernel
        快速验证
      标准开发
        工程化
        产品化
        多种调用方式
    HOST与Device
      HOST侧
        CPU
        主控逻辑
        任务调度
      Device侧
        AI Core
        计算逻辑
        高性能计算
      数据流向
    Tiling机制
      核心概念
        Tiling实现
        Tiling块
        Tiling算法
        Tiling结构体
        Tiling函数
      三个关键环节
        定义结构体
        实现函数
        Kernel使用
    Shape支持
      固定Shape
        简单易用
        局限性大
      动态Shape
        参数传递
        灵活通用
      Shape推导
        错误检查
        内存优化
    工程结构
      msopgen工具
      目录职责
        op_proto
        op_host
        op_kernel
    编译部署
      CMake配置
      编译步骤
      部署流程

一、从快速开发到标准开发的进阶之路

今天学习了Ascend C算子的标准工程开发方式，这可以说是算子开发从"玩具"到"产品"的关键一步！之前我只用过快速开发模式——直接写个kernel函数，编译运行看结果。虽然简单快捷，但有很大的局限性。今天老师讲的标准开发模式让我眼前一亮，这才是真正的工程化开发！

1.1 两种开发方式的对比

老师一开始就给我们明确了两种开发方式的区别：

快速开发模式：

• 直接编写kernel函数
• 通过一个CPP程序调用和编译
• 适合快速验证算法
• 局限性：只能处理固定的数据规模

标准开发模式：

• 更侧重于工程化和产品化
• 实现过程与调用过程清晰分离
• 具备更高的通用性和可支持性
• 能够直接集成到大型模型中

1.2 标准开发的核心优势

老师特别强调了几个关键优势，我深有感触：

优势1：清晰的职责分离

将算子的实现过程与调用过程分离，这是软件工程的基本原则。实现者专注于算法优化，调用者只需要关心接口使用。

优势2：多种调用方式支持

标准开发的算子可以通过多种方式调用：

• NN接口（神经网络框架接口）
• 单算子模型
• PyTorch等主流深度学习框架

这意味着一次开发，到处使用！

优势3：支持动态Shape

这个太重要了！实际应用中，输入数据的大小往往是变化的。标准开发模式天然支持动态Shape，而快速开发模式只能处理固定大小的数据。

二、理解HOST与Device的架构设计

2.1 一个生动的类比

老师用了一个特别形象的比喻来解释HOST和Device的关系：

HOST侧：就像电脑主机（CPU）

• 负责执行主要应用程序
• 进行任务调度
• 管理整体流程

Device侧：就像专用显卡

• 负责执行专项任务
• 专注于高性能计算
• 由HOST调度

这个比喻让我一下就理解了两者的关系！

2.2 硬件对应关系

老师进一步解释了硬件层面的对应关系：

HOST：

• x86或Arm架构的服务器CPU
• Host Memory（主机内存）
• 运行主控程序

Device：

• 安装了昇腾硬件的AI加速卡
• 包含多个AI Core
• 独立的Global Memory
• AI Core内部有高速的Local Memory

2.3 数据流向理解

理解数据的流向对于算子开发至关重要，我画了一个简单的流程图：

Host Memory (数据源)
    ↓ 拷贝
Device Global Memory (设备全局内存)
    ↓ 搬运(MT2)
AI Core Local Memory (核心局部内存)
    ↓ 计算
AI Core Local Memory (计算结果)
    ↓ 搬运(MT3)
Device Global Memory (结果存储)
    ↓ 拷贝
Host Memory (结果返回)

这个流程清晰地展示了数据如何从主机到设备，再从设备的全局内存到AI Core的局部内存，最后又如何返回的整个过程。

三、Tiling机制：算子开发的核心

3.1 Tiling的概念体系

老师花了很多时间讲Tiling，我觉得这确实是算子开发中最重要也最难理解的部分。

Tiling实现：

• 完整的数据切分策略
• 参数计算过程

Tiling块：

• 每次搬运的一部分数据块
• 大小由Tiling算法决定

Tiling算法：

• 根据算子不同输入形状决定每个输入块大小
• 考虑Local Memory容量限制
• 平衡计算效率和内存使用

Tiling结构体：

• 统一管理算法所需参数
• 从Host侧传递到Device侧

Tiling函数：

• 实现tiling算法
• 将计算出的参数填入tiling结构体

3.2 为什么需要Tiling？

老师举了一个例子让我豁然开朗：

假设我们要处理一个1GB的数据，但AI Core的Local Memory只有几百KB。怎么办？

答案：分块处理！

就像吃一个大西瓜，你不能一口吞下去，而是切成小块一块块吃。Tiling就是这个"切西瓜"的策略。

3.3 Tiling实现的三个关键环节

老师总结了Tiling实现的三个关键步骤：

步骤1：定义Tiling结构体

位置：Host侧的头文件中（如op_host/*_tiling.h）

BEGIN_TILING_DATA_DEF(TilingData)
    TILING_DATA_FIELD_DEF(uint32_t, totalLength);    // 数据总长度
    TILING_DATA_FIELD_DEF(uint32_t, tileNum);        // 分块数量
    TILING_DATA_FIELD_DEF(uint32_t, blockSize);      // 每块大小
END_TILING_DATA_DEF()

这个结构体定义了所有需要传递给Kernel侧的参数。

步骤2：实现Tiling函数

位置：Host侧的CPP文件中（如op_host/*.cpp）

核心任务：为定义好的tiling结构体的每个字段赋予符合要求的分块策略值

实现流程：

1. 定义tiling结构体变量
2. 从Context中获取输入信息
3. 根据分块算法计算每个字段的值
4. 为结构体的每个字段赋值
5. 调用tiling.SetBuffer()保存结构体

步骤3：Kernel侧使用Tiling信息

在Kernel函数中：

__global__ __aicore__ void kernel_func(
    GM_ADDR x, 
    GM_ADDR y, 
    GM_ADDR z,
    GM_ADDR workspace,  // 新增：工作空间
    GM_ADDR tiling      // 新增：tiling数据
)

使用GetTilingData宏解析tiling数据：

TilingData tilingData;
GetTilingData(tilingData, tiling);

// 访问tiling参数
uint32_t totalLen = tilingData.totalLength;
uint32_t tileNum = tilingData.tileNum;

3.4 一个实际的Tiling算法示例

老师给了一个简单的Tiling算法示例，我把它整理出来：

// 假设：
// - 数据总长度：8192
// - Local Memory可用空间：2048
// - 需要两个输入和一个输出，共3个buffer

// 计算每个buffer的大小
uint32_t bufferSize = 2048 / 3;  // 每个buffer约682

// 计算需要多少个tile
uint32_t tileNum = (8192 + bufferSize - 1) / bufferSize;  // 向上取整

// 计算每个tile的实际大小
uint32_t tileSize = (8192 + tileNum - 1) / tileNum;  // 尽量均分

// 填充tiling结构体
tilingData.totalLength = 8192;
tilingData.tileNum = tileNum;
tilingData.blockSize = tileSize;

这个算法考虑了：

• Local Memory的容量限制
• 数据的均匀分配
• 边界情况的处理

四、固定Shape vs 动态Shape

4.1 固定Shape开发

老师首先讲了固定Shape的开发方式：

优势：

• 实现简单
• 开发者只需专注在已知数据规模下设计分块策略
• 调试方便

局限性：

• 只能处理单一类型的数据
• 当数据Shape变化时无法正常工作
• 不适合实际应用

示例：

// 固定Shape：只能处理长度为8192的数据
const uint32_t FIXED_LENGTH = 8192;
const uint32_t TILE_SIZE = 1024;
const uint32_t TILE_NUM = 8;

4.2 动态Shape开发

核心思路：

• 将实际数据的Shape信息通过参数动态传入
• 每次调用算子时，算子内部实时计算分块策略

实现方式：

通过tiling结构体传递关键参数：

• 数据总长度totalLength
• AI Core数量coreNum
• 每个核上的分块数量tileNum

示例：

// 动态计算tiling参数
void ComputeTiling(OpTilingContext& context) {
    // 获取输入shape
    auto inputShape = context.GetInputShape(0);
    uint32_t totalLength = inputShape[0];
    
    // 获取可用核心数
    uint32_t coreNum = context.GetCoreNum();
    
    // 计算每个核处理的数据量
    uint32_t lengthPerCore = (totalLength + coreNum - 1) / coreNum;
    
    // 计算每个核的分块数
    uint32_t tileNum = (lengthPerCore + TILE_SIZE - 1) / TILE_SIZE;
    
    // 填充tiling数据
    tilingData.totalLength = totalLength;
    tilingData.coreNum = coreNum;
    tilingData.tileNum = tileNum;
}

4.3 动态Shape的优势

老师用一个实际例子说明了动态Shape的重要性：

场景：一个推荐系统模型

• 批次1：100个用户，每个用户10个特征 → Shape: [100, 10]
• 批次2：50个用户，每个用户10个特征 → Shape: [50, 10]
• 批次3：200个用户，每个用户10个特征 → Shape: [200, 10]

如果用固定Shape，要么只能处理100个用户（浪费计算资源），要么就无法处理！

用动态Shape，一个算子就能适应所有情况！

五、Shape推导：智能的输出预测

5.1 Shape推导的核心任务

定义：根据算子输入Tensor的Shape信息，自动推导出输出Tensor的Shape

这听起来简单，但作用非常大！

5.2 Shape推导的两大作用

作用1：模型构图阶段的校验

在模型构建时，框架会根据每层的输出Shape推导规则，检查层与层之间的连接是否正确。

例子：

# 假设：
layer1_output: [batch, 128]
layer2_input:  需要[batch, 256]

# Shape推导会发现：128 ≠ 256
# 在构图阶段就报错，而不是等到运行时才发现！

这可以提前发现错误，避免浪费时间！

作用2：内存预分配优化

框架能够精确知道要为算子的输出分配多大的内存空间。

传统方式：

运行算子 → 发现输出大小 → 分配内存 → 重新运行

有Shape推导：

推导输出Shape → 预分配内存 → 运行算子（一次完成）

效率提升明显！

5.3 Shape推导的实现

老师给了一个向量加法算子的Shape推导示例：

// 向量加法：z = x + y
// 输入：x[N], y[N]
// 输出：z[N]

void InferShape(const OpDescriptor& desc, 
                std::vector<TensorDesc>& outputShapes) {
    // 获取输入shape
    auto inputShape1 = desc.GetInputShape(0);
    auto inputShape2 = desc.GetInputShape(1);
    
    // 校验：两个输入shape必须相同
    if (inputShape1 != inputShape2) {
        throw ShapeError("Input shapes must match!");
    }
    
    // 推导输出shape：与输入相同
    outputShapes[0] = inputShape1;
}

对于更复杂的算子（如矩阵乘法），Shape推导会更复杂：

// 矩阵乘法：C = A × B
// 输入：A[M, K], B[K, N]
// 输出：C[M, N]

void InferShape(const OpDescriptor& desc,
                std::vector<TensorDesc>& outputShapes) {
    auto shapeA = desc.GetInputShape(0);  // [M, K]
    auto shapeB = desc.GetInputShape(1);  // [K, N]
    
    // 校验：A的列数必须等于B的行数
    if (shapeA[1] != shapeB[0]) {
        throw ShapeError("Matrix dimensions incompatible!");
    }
    
    // 推导输出shape
    outputShapes[0] = {shapeA[0], shapeB[1]};  // [M, N]
}

六、工程创建与项目结构

6.1 使用msopgen创建算子工程

老师介绍了使用msopgen工具创建算子原型：

命令示例：

msopgen gen -i AddCustom.json -f pytorch -c ai_core-ascend910 -o ./

生成的文件结构：

AddCustom/
├── op_proto/
│   └── add_custom.py          # 算子原型定义
├── op_host/
│   ├── add_custom_tiling.h    # Tiling结构体定义
│   ├── add_custom.cpp         # Host侧实现
│   └── shape_infer.cpp        # Shape推导
├── op_kernel/
│   └── add_custom.cpp         # Kernel侧实现
├── CMakeLists.txt             # 编译配置
└── build.sh                   # 构建脚本

6.2 标准工程的目录职责

老师详细讲解了每个目录的职责：

op_proto目录：

• 算子的原型定义
• 输入输出参数定义
• 算子属性定义

op_host目录：

• Host侧的实现代码
• Tiling策略实现
• Shape推导实现
• 数据类型推导

op_kernel目录：

• Kernel侧的核心实现
• 实际的计算逻辑
• 运行在AI Core上的代码

七、编译与部署流程

7.1 编译配置

老师重点讲解了CMakeLists.txt的配置：

关键参数：ASCEND_PACKAGE_PATH

这个参数指定Ascend软件包的安装路径，默认是：

set(ASCEND_PACKAGE_PATH "/usr/local/Ascend/ascend-toolkit/latest")

如果你的安装路径不同，需要修改这个配置。

7.2 编译步骤

老师演示的编译流程：

# 1. 进入算子工程根目录
cd AddCustom/

# 2. 创建并进入build目录
mkdir build && cd build

# 3. 运行CMake配置
cmake ..

# 4. 编译
make

# 5. 打包
make package

编译成功后：

• 在build目录下生成custom_op_package.run安装包
• 这个就是可以部署的算子包

7.3 部署过程

部署步骤：

# 1. 进入build目录
cd build/

# 2. 运行安装包
./custom_op_package.run

# 3. 安装完成后，算子会被注册到系统中
# 可以通过NN接口、PyTorch等调用

验证安装：

import torch
import torch_npu

# 创建NPU设备上的张量
x = torch.randn(1024).npu()
y = torch.randn(1024).npu()

# 调用自定义算子
z = torch_npu.npu_add_custom(x, y)

print("算子调用成功！")

八、中级认证要求与实践

8.1 认证任务：实现Sigmoid算子

老师最后讲了中级认证的要求：实现Sigmoid算子

Sigmoid函数：
$$\text{sigmoid}(x)=\frac{1}{1+e^{-x}}$$

具体要求：

1. 补充Kernel侧和Host侧的完整代码
2. 在Host侧实现tiling结构体的变量定义
3. 算子必须支持Float16类型
4. 实现合理的tiling策略

8.2 实现注意事项

老师特别提醒了几个容易踩的坑：

坑1：精度问题

Sigmoid函数涉及指数计算，需要注意：

• Float16的表示范围有限
• 需要处理数值溢出
• 可能需要使用高精度API

坑2：API选择

昇腾提供了多个指数函数API：

• Exp：标准指数函数
• FastExp：快速指数函数（精度略低）
• Vexp：向量指数函数

需要根据精度要求选择合适的API。

坑3：Tiling策略

需要考虑：

• 临时变量的内存占用
• 中间结果的存储
• 边界情况的处理

8.3 我的实现思路

我根据老师的讲解，整理了一个实现思路：

第一步：定义Tiling结构体

BEGIN_TILING_DATA_DEF(SigmoidTilingData)
    TILING_DATA_FIELD_DEF(uint32_t, totalLength);
    TILING_DATA_FIELD_DEF(uint32_t, tileNum);
    TILING_DATA_FIELD_DEF(uint32_t, tileSize);
END_TILING_DATA_DEF()

第二步：实现Tiling函数

void ComputeTiling(OpTilingContext& context) {
    // 获取输入shape
    auto shape = context.GetInputShape(0);
    uint32_t totalLength = GetTotalSize(shape);
    
    // 计算tile大小（考虑临时变量）
    // 需要：input buffer + output buffer + temp buffer
    uint32_t availableSize = LOCAL_MEMORY_SIZE / 3;
    uint32_t tileSize = min(availableSize, totalLength);
    
    // 计算tile数量
    uint32_t tileNum = (totalLength + tileSize - 1) / tileSize;
    
    // 填充tiling数据
    tilingData.totalLength = totalLength;
    tilingData.tileNum = tileNum;
    tilingData.tileSize = tileSize;
}

第三步：实现Kernel函数

__global__ __aicore__ void sigmoid_kernel(
    GM_ADDR input,
    GM_ADDR output,
    GM_ADDR workspace,
    GM_ADDR tiling
) {
    // 1. 获取tiling数据
    SigmoidTilingData tilingData;
    GetTilingData(tilingData, tiling);
    
    // 2. 分配Local Memory
    LocalTensor<half> inputLocal = ...;
    LocalTensor<half> outputLocal = ...;
    LocalTensor<half> tempLocal = ...;
    
    // 3. 循环处理每个tile
    for (uint32_t i = 0; i < tilingData.tileNum; i++) {
        // 3.1 CopyIn：搬入数据
        DataCopy(inputLocal, input + offset, tileSize);
        
        // 3.2 Compute：计算sigmoid
        // temp = -input
        Muls(tempLocal, inputLocal, -1.0);
        
        // temp = exp(-input)
        Exp(tempLocal, tempLocal);
        
        // temp = 1 + exp(-input)
        Adds(tempLocal, tempLocal, 1.0);
        
        // output = 1 / (1 + exp(-input))
        Reciprocal(outputLocal, tempLocal);
        
        // 3.3 CopyOut：搬出结果
        DataCopy(output + offset, outputLocal, tileSize);
    }
}

九、我的学习心得

9.1 工程化思维的重要性

这次学习让我深刻认识到，算子开发不只是写代码，更重要的是工程化思维：

• 职责分离：Host和Device各司其职
• 参数化设计：通过Tiling支持动态Shape
• 错误检查：通过Shape推导提前发现问题
• 可维护性：清晰的目录结构和代码组织

9.2 Tiling是核心难点

Tiling机制是整个标准开发模式的核心，也是最难理解的部分：

• 需要深入理解内存管理
• 需要权衡计算效率和内存使用
• 需要考虑各种边界情况

建议多看示例代码，多动手实践。

9.3 工具链的熟练使用

标准开发涉及多个工具：

• msopgen：工程生成工具
• CMake：构建工具
• Profiling工具：性能分析工具

熟练掌握这些工具能大大提高开发效率。

9.4 实践是最好的老师

听课只能理解概念，只有自己动手实现一个完整的算子，才能真正掌握标准开发流程。

我计划接下来：

1. 完成Sigmoid算子的实现
2. 尝试实现更复杂的算子（如Softmax）
3. 学习性能优化技巧

十、总结

这次学习让我系统地掌握了Ascend C算子的标准工程开发方式。从HOST与Device的架构设计，到Tiling机制的深入理解，从固定Shape到动态Shape的进阶，从Shape推导的智能优化，到完整的编译部署流程，每一个环节都是工程化开发的必备知识。

核心要点回顾：

1. 标准开发 vs 快速开发：工程化 vs 原型验证
2. HOST与Device分离：主控逻辑 vs 计算逻辑
3. Tiling机制：数据切分的艺术
4. 动态Shape支持：真实应用的必备
5. Shape推导：智能的错误检查和内存优化
6. 完整工程结构：op_proto、op_host、op_kernel三位一体

对于想要学习Ascend C标准开发的同学，我的建议是：

1. 先理解整体架构
2. 重点掌握Tiling机制
3. 从简单算子开始实践
4. 逐步学习性能优化技巧
5. 多参考官方示例代码