3.1 Ascend C 编程模型：核函数、Tiling 与 UB

Ascend C 是昇腾的算子开发语言，可以理解成「面向昇腾 NPU 的 CUDA」。但它解决的不是「如何组织线程」，而是「如何组织数据流」。这一章建立 Ascend C 的三大核心心智模型：核函数、Tiling、两级内存视图。

📑 目录

1. Ascend C 到底是什么
2. 核心概念一：核函数（Kernel Function）
3. 核心概念二：Tiling（数据切分）
4. 核心概念三：两级内存视图
5. 一个最小 Add 算子的骨架
6. Ascend C 与 CUDA 的编程范式对比
7. 学习路径建议
总结
延伸阅读

1. Ascend C 到底是什么

Ascend C 是昇腾提供的、基于 C++ 的算子编程语言，开发者用它编写「在 AI Core 上运行的核函数」，并管理数据搬运、Tiling 与同步。

对标 CUDA：都是「用高层语言写底层算子」。
但范式不同：CUDA 是「线程 → 块 → 网格」的 SIMT 模型；Ascend C 是「核函数 + Tiling + 显式数据搬运」的数据流模型。
由 CANN 工具链编译：Ascend C 源码经编译器生成 AI Core 可执行代码。

为什么昇腾不用「线程模型」？因为昇腾的计算单元是「专用化」的（Cube 管矩阵、Vector 管向量、Scalar 管控制），而且 AI Core 内部的计算依赖「数据先在片上 Buffer 就位」。线程模型适合「灵活调度一堆同质线程」，而昇腾的模型更像「编排一台流水线机器」。

2. 核心概念一：核函数（Kernel Function）

核函数（Kernel Function）是跑在 AI Core 上的计算函数，是 Ascend C 算子的执行主体。

关键认知：

核函数用 __global__ 类似的关键字/宏标记（Ascend C 中通过特定关键字声明为核函数）。
一个核函数在多个 AI Core 上并行执行（类似 CUDA 的 grid 跨 SM 并行）。
核函数内部使用 Ascend C 提供的高级接口：数据搬运（DataCopy）、矢量计算（Add、Exp…）、矩阵计算（Matmul）等。
核函数由 Host 侧发起调用（通过 ACLRT_LAUNCH_KERNEL 之类的机制下发），执行是异步的。

Host 侧（CPU）
   │  下发核函数 + Tiling 参数
   ▼
AI Core 侧：多核并行执行核函数
   │  DataCopy 搬数据 → 计算 → DataCopy 写回
   ▼
Host 侧：同步等待完成

3. 核心概念二：Tiling（数据切分）

Tiling 是把大张量切成能装进片上 Buffer（UB/L0）的小块，逐块计算的过程。

为什么必须 Tiling？因为：

HBM 里可能放着几 GB 的张量；
而 UB 只有几十~几百 KB；
所以一次只能搬一小块上来算，算完搬走，再来下一块。

Tiling 是昇腾算子开发里最需要动脑子的部分：

Tiling 维度	要决策的事
块大小（Block）	每次搬多少数据进 UB，是否恰好填满 UB
循环组织	外层循环遍历所有块，内层处理单块
边界处理	张量尺寸不整除块大小时，最后一块怎么补/怎么截
多核切分	把整个张量分成几份分给不同的 AI Core

Tiling 参数通常在 Host 侧计算，再通过参数传给核函数——所以 Ascend C 算子代码里经常同时存在「Host 侧 Tiling 计算」和「Device 侧核函数」两部分。

4. 核心概念三：两级内存视图

Ascend C 的编程内存视图可以简化为两层：

Global Memory（GM）：对应 HBM，容量大、访问慢。数据住在这里。
Unified Buffer（UB）：AI Core 内的计算工作台，容量小、速度快。计算在这里发生。

数据流永远是：

GM ──DataCopy──▶ UB ──计算──▶ UB ──DataCopy──▶ GM

（矩阵计算则会经过 L0A/L0B/L0C，见 3.4。）

这条「先搬后算、算完搬回」的强制路径，是 Ascend C 与 CUDA 最大的体验差异。 CUDA 里写 a[i] + b[i] 编译器会处理缓存；Ascend C 里你得显式说「把数据搬到 UB，算，再搬回」。

5. 一个最小 Add 算子的骨架

先看整体骨架，细节在 3.2 展开：

// 核函数：在 AI Core 上执行
__global__ void addKernel(const __gm__ float *a,
                          const __gm__ float *b,
                          __gm__ float *c,
                          const int totalLen) {
    // 1. 确定本核处理的数据范围（多核切分）
    // 2. 循环处理 Tiling 后的每个 Block
    for (int i = 0; i < numBlocks; i++) {
        // 2.1 DataCopy: GM → UB（搬入 a、b 分块）
        // 2.2 矢量计算：Add(aUb, bUb, cUb, blockLen)   // 在 UB 上
        // 2.3 DataCopy: UB → GM（搬出结果）
    }
}

Host 侧部分负责：分配 GM 内存、计算 Tiling 参数、下发核函数、同步。

6. Ascend C 与 CUDA 的编程范式对比

维度	CUDA	Ascend C
并行模型	Thread/Block/Grid（SIMT）	核函数 + 多 AI Core
最小执行单元	Warp（32 线程）	单核处理的数据块
数据搬运	隐式缓存 + 可选手动共享内存	强制显式 DataCopy
访存视图	全局/共享/寄存器	GM / UB / L0A/L0B/L0C
分块	靠线程索引隐式划分	显式 Tiling 参数
同步	`__syncthreads()`	队列/缓冲区同步（EnQue/DeQue）
典型代码结构	一个 kernel 函数	Host Tiling + Device 核函数

核心差异一句话：CUDA 让「线程」去拿数据，Ascend C 让「数据」被主动搬到位。

7. 学习路径建议

先吃透 3.2 的矢量 Add 算子——它能覆盖 Tiling、DataCopy、矢量计算、写回的全部要素，是昇腾算子开发的「Hello World」。
再做 Softmax 类规约算子——引入块内规约与更复杂的 Tiling。
再做 GEMM（3.4）——切换到矩阵编程，理解 Cube 的 Tiling 逻辑。
最后学习 msprof（3.5）——用数据说话，验证每一步优化。

总结

Ascend C 是昇腾的算子编程语言，三大核心心智：核函数、Tiling、两级内存视图（GM/UB）。
数据流强制路径：GM → UB → 计算 → GM，搬运是显式且必不可少的。
Tiling 决定「数据怎么切、怎么循环」，是算子性能的最大杠杆。
与 CUDA 同构但不相同：理解了「线程模型」到「数据流模型」的转变，Ascend C 就入门了。

搜索