跳转至

Ascend C 入门

约 1162 个字 67 行代码 1 张图片 预计阅读时间 5 分钟

Ascend C 是 CANN 为算子开发推出的编程语言，原生遵循 C/C++ 标准，可以理解为“昇腾上的 CUDA C”。本章不追求覆盖全部细节，目标是让你看懂 Ascend C 代码的并行结构，并能照着官方样例写出第一个算子。

设备结构：为什么 Ascend C 长得和 CUDA 不太一样

回忆一下 CUDA：GPU 上有大量通用核心（SM），每个核心用 warp 调度 SIMT 线程。昇腾的 AICore 则是另一种思路——把“矩阵”和“向量”两类计算拆给两套专用单元：

  • Cube 单元：专做矩阵乘累加（对应张量核心），数据从 L0A/L0B 读入、累加结果写入 L0C；
  • Vector 单元：做各种向量/标量运算，工作在 Unified Buffer（UB）上。

数据都从 Global Memory（GM）出发，经 DMA 搬到片上（L1 → L0 / UB），算完再写回 GM。因此一个典型的 Ascend C 算子 = 搬运（DataCopy）+ 计算（Vector/Cube 指令）+ 搬运 的流水，而不是 CUDA 那种“线程各自读全局内存直接算”。

AICore 数据通路示意图

其他关键概念与 CUDA 的对应关系：

CUDA Ascend C
<<<grid, block>>> 启动多个线程块 一个 kernel 由多个 AICore SPMD 执行，GetBlockIdx() 取核号
shared memory / register L1 / UB / L0 等多级片上存储
线程块级同步 __syncthreads() 队列与事件同步（TQue、SetFlag/WaitFlag）
手写全局内存索引 DataCopy 接口成块搬运
nvcc CANN 工具链（基于毕昇编译器）

小试牛刀：add_custom

下面是官方经典样例 AddCustom 的简化版：对两个 half 向量求和。每个 AICore 处理数据的一段，处理过程中数据先被 DMA 进 UB，算完再写回。

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
#include "kernel_operator.h"
using namespace AscendC;
constexpr int TOTAL_NUM = 2 * 1024;          // 所有核合计要处理的元素总个数
constexpr int USE_CORE_NUM = 8;              // 使用的核数
constexpr int BLOCK_NUM = TOTAL_NUM / USE_CORE_NUM;   // 每个核负责的元素个数

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    // 初始化：设置全局内存地址，并在统一缓冲上开输入/输出队列
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
        xGm.SetGlobalBuffer((__gm__ half*)x + GetBlockIdx() * BLOCK_NUM, BLOCK_NUM);
        yGm.SetGlobalBuffer((__gm__ half*)y + GetBlockIdx() * BLOCK_NUM, BLOCK_NUM);
        zGm.SetGlobalBuffer((__gm__ half*)z + GetBlockIdx() * BLOCK_NUM, BLOCK_NUM);
        pipe.InitBuffer(inQueueX, 2, BLOCK_NUM * sizeof(half));  // 双缓冲
        pipe.InitBuffer(inQueueY, 2, BLOCK_NUM * sizeof(half));
        pipe.InitBuffer(outQueueZ, 2, BLOCK_NUM * sizeof(half));
    }
    // 主体：进——算——出的流水
    __aicore__ inline void Process() {
        CopyIn();
        Compute();
        CopyOut();
    }
private:
    __aicore__ inline void CopyIn() {
        LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();
        DataCopy(xLocal, xGm, BLOCK_NUM);    // GM -> UB
        DataCopy(yLocal, yGm, BLOCK_NUM);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }
    __aicore__ inline void Compute() {
        LocalTensor<half> xLocal = inQueueX.DeQue<half>();   // 等数据到位
        LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();
        Add(zLocal, xLocal, yLocal, BLOCK_NUM);              // Vector 指令，UB 上完成
        outQueueZ.EnQue<half>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    __aicore__ inline void CopyOut() {
        LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        DataCopy(zGm, zLocal, BLOCK_NUM);    // UB -> GM
        outQueueZ.FreeTensor(zLocal);
    }
    TPipe pipe;
    TQue<QuePosition::VECIN, 2> inQueueX, inQueueY;
    TQue<QuePosition::VECOUT, 2> outQueueZ;
    GlobalTensor<half> xGm, yGm, zGm;
};

// 设备侧入口，类比 CUDA 的 __global__ 函数
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

几个首次出现的记号

  • GM_ADDR：CANN 提供的宏，把入参展开为指向全局内存（GM）的指针；
  • __gm__ / __aicore__：地址空间与函数修饰符，分别表示“数据在 GM 上”“函数在 AICore 上执行”，类比 CUDA 的 __device__；
  • half：16 位半精度浮点（fp16），昇腾上最常用的计算精度之一；
  • TQue<..., 2> 中的 2 是队列深度，配合两块缓冲区实现“双缓冲”。

重点体会三件事：

  1. GetBlockIdx() 让每个 AICore 只处理自己那一段数据，这就是 SPMD 并行，和 CUDA 中用 blockIdx 划分数据一模一样；
  2. TPipe/TQue 管理着片上存储与同步：AllocTensor/EnQue/DeQue 的队列语义替你处理了“数据搬进 UB 之后才能算”的依赖关系，双缓冲让搬运与计算可以重叠；
  3. 计算发生在 UB 上：Add(zLocal, ...) 操作的是 LocalTensor，而不是像 CUDA 那样直接从全局内存 load/store。

Host 侧的调用有两条路线：一是轻量的 Kernel 直调（用 aclrt 接口申请设备内存、加载 kernel 后直接 launch）；二是工程上更常见的 标准算子工程——用 CANN 提供的工程模板组织 shape 推导、tiling（切分数据到各核）与 kernel，用算子工程编译脚本打包成自定义算子库（opp 包）再安装加载。至于 ATC 工具，它负责把 ONNX 等模型转换成 om 推理模型，属于推理部署路线，与算子开发是两条不同的路径。

新手建议先把官方 samples 仓库里的 Kernel 直调样例跑通（仓库已迁移至 GitCode：cann/samples，GitHub 镜像为 Ascend/samples）。AddCustom 直调样例位于 cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op/kernel_invocation/Add，一条命令完成编译与运行（脚本只接受 ascend910 / ascend310p 两种 SOC_VERSION；手上是 910B 等其他机型或暂时没有 NPU 时，先用 cpu 模拟模式跑通功能）：

git clone https://github.com/Ascend/samples.git
cd samples/cplusplus/level1_single_api/4_op_dev/6_ascendc_custom_op/kernel_invocation/Add

# 真机运行
bash run.sh add_custom ascend910 AiCore npu

# 没有 NPU？先用 CPU 模拟模式跑通功能
bash run.sh add_custom ascend910 AiCore cpu

跑通直调版之后，再学标准算子工程（样例见同级的 acl_invocation 部分，含 op_host / op_kernel / tiling 的完整工程结构）。

学习&拓展：

  1. 按上文命令把 add_custom Kernel 直调样例跑通（NPU 或 CPU 模拟模式均可），再浏览昇腾社区 Ascend C 算子开发文档核对样例中每个接口的语义。
  2. 把 add_custom 改成 float32 版本，并尝试把 USE_CORE_NUM 改大/改小，观察性能变化。
  3. 了解 tiling 概念：为什么大矩阵必须切块循环（L1/UB 装不下），以及“缓冲区大小 × 核数”如何决定切分方案。
  4. 阅读 cann-ops 仓库 中一个激活函数算子的源码，找出它的 Cube/Vector 用法与 add_custom 的差异。
  5. 了解大矩阵乘在昇腾上如何用 Cube 单元实现：GM → L1 → L0A/L0B → L0C 的数据通路。

一些链接

昇腾官方文档中心（CANN 文档入口）

Ascend C 算子开发官方文档

CANN 样例仓库（GitCode，含 Ascend C 算子样例）

cann-ops：官方 Ascend C 基础算子库源码

B站《Ascend C 算子开发入门课程》