1.3 昇腾存储层次与数据流

上一章我们知道了 AI Core 有 Cube/Vector/Scalar 三个计算单元。但算力再强，数据喂不上去也没用。这一章我们聚焦昇腾的存储层次（Memory Hierarchy）——它决定了算子和训练能跑多快，也是理解 Ascend C 中「数据搬运」概念的前提。

📑 目录

1. 昇腾存储层次总览
2. HBM：片外大容量高带宽内存
3. L2 Cache：片上共享缓存
4. AI Core 内部存储：UB / L0A / L0B / L0C
5. 数据流：一条数据从 HBM 到计算单元的旅程
6. 与 CUDA 内存模型的映射
7. 存储层次对性能优化的启示
总结
延伸阅读

1. 昇腾存储层次总览

昇腾的存储体系可以用一句话概括：

「片外 HBM 装大块数据，片上 L2 做中转缓存，AI Core 内部 Buffer 是计算的直接工作台。」

        HBM（片外，容量大、带宽高）
          │  数据搬运（DataCopy）
          ▼
        L2 Cache（片上共享）
          │
          ▼
   ┌───────────────────────────┐
   │   AI Core 内部             │
   │   ├─ Unified Buffer (UB)  │  ← 向量/通用计算的工作台
   │   ├─ L0A / L0B（Cube输入） │
   │   ├─ L0C（Cube输出）       │
   │   └─ 寄存器组（Register）  │
   └───────────────────────────┘

每一层都比上一层更快、但更小、更贵。优化昇腾算子的核心，就是让数据尽量停留在片上、减少 HBM 访问次数——这和 CUDA 里「减少全局内存访问、多用共享内存」是同一个道理。

2. HBM：片外大容量高带宽内存

HBM（High Bandwidth Memory，高带宽内存） 是昇腾片外的主内存，训练大模型时权重、梯度、激活、优化器状态全都住在这里。

关键认知：

容量：910 系列训练芯片的 HBM 容量远大于推理芯片（典型值随代际变化，请以官方规格为准）。
带宽：HBM 通过 3D 堆叠 + 宽接口获得极高带宽，是 GDDR 的数倍。Transformer 训练/推理中大量数据搬运都依赖 HBM 带宽。
对比：HBM 之于昇腾 ≈ 显存（VRAM）之于 GPU。你在 GPU 上理解过的「显存带宽瓶颈（memory-bound）」概念，在昇腾上同样成立。

在 Ascend C 里，HBM 被称为 Global Memory，算子代码中通过指针访问；开发者需要显式地把数据从 Global Memory 搬到片上（DataCopy），这正是昇腾「显式数据搬运」风格的体现。

3. L2 Cache：片上共享缓存

L2 Cache 是所有 AI Core 共享的片上二级缓存，位于 HBM 与 AI Core 之间：

作用：缓存频繁访问的数据（如权重的分块），减少对 HBM 的重复访问。
共享性：多个 AI Core 共享 L2，因此 L2 也承担了「核间数据暂存与通信」的一部分角色。
容量：远小于 HBM，但远大于 AI Core 内部 Buffer（数量级差异，以具体芯片为准）。

可以把它类比为 GPU 的 L2 Cache：GPU 里也有一块所有 SM 共享的 L2，用来缓存全局内存访问。

4. AI Core 内部存储：UB / L0A / L0B / L0C

这是昇腾最独特的地方——AI Core 内部有专为计算单元设计的几块 Buffer：

存储	全称/职责	主要使用方	类比
UB	Unified Buffer，统一缓冲区，AI Core 的「通用工作台」	向量计算、数据中转	近似 GPU 的共享内存（SMEM）
L0A	Cube 的输入 A 缓冲	AI Cube 矩阵乘左矩阵	近似 Tensor Core 的寄存器输入
L0B	Cube 的输入 B 缓冲	AI Cube 矩阵乘右矩阵	同上
L0C	Cube 的输出/累积缓冲	AI Cube 结果累积	近似 Tensor Core 的累积寄存器
寄存器	通用寄存器组	Scalar / 向量指令	通用寄存器

关键点：

UB 是 Ascend C 编程的「主场」：向量类算子（Add、Softmax 等）的数据通常先在 UB 里就位，然后让 AI Vector 计算。
L0A/L0B/L0C 是 Cube 的专用通路：矩阵乘数据从 HBM/L2 直接灌入 L0A/L0B，算完累积在 L0C，这条路径几乎不经过通用寄存器。
容量都有限：UB/L0 每块容量都只有几十 KB 到几百 KB 级别，因此「数据太大放不下」时必须 Tiling（分块）——把大张量切成能塞进 Buffer 的小块循环处理。

5. 数据流：一条数据从 HBM 到计算单元的旅程

以一次向量 Add 算子为例，看数据的完整旅程：

1. 数据 a、b 在 HBM（Global Memory）中
2. Ascend C 代码中执行 DataCopy：
   HBM ──DataCopy──▶ UB（a、b 被搬进 Unified Buffer）
3. AI Vector 执行 Add： UB 中逐元素相加 → 结果写入 UB
4. 结果在 UB 中就位后，执行 DataCopy：
   UB ──DataCopy──▶ HBM（结果写回 Global Memory）

如果是矩阵乘（GEMM），旅程略有不同：

1. HBM ──DataCopy──▶ L0A / L0B（Cube 输入）
2. AI Cube：L0A × L0B ──▶ L0C（FP32 累积）
3. 若后续需要激活：L0C ──▶ 向量路径 ──▶ AI Vector 处理
4. 结果 ──▶ UB ──DataCopy──▶ HBM

可以看到：昇腾算子性能的关键 = 搬运次数与计算时间的匹配。如果数据搬运速度跟不上计算单元的速度，算力再强也会「饿着」（memory-bound）。

6. 与 CUDA 内存模型的映射

如果你有 CUDA 基础，下面这张映射表能帮你快速迁移直觉：

CUDA / GPU	昇腾	说明
Global Memory（全局内存）	Global Memory（HBM）	显式搬运、带宽宝贵
L2 Cache	L2 Cache	核间共享缓存
Shared Memory（共享内存）	Unified Buffer（UB）	都是片上可编程的「手工管理」Buffer
寄存器	L0A/L0B/L0C + 寄存器组	Cube 专用输入输出缓冲对应 Tensor Core 的寄存器路径
`__syncthreads()`	数据同步/队列同步（EnQue/DeQue）	昇腾用队列/同步原语管理 Buffer 生命周期

一个显著差异要特别提醒：CUDA 里数据从全局内存到共享内存的搬运是「隐式+手工并存」的，而昇腾 Ascend C 中 GM→UB 的 DataCopy 几乎是必经之路。所以昇腾算子代码里「搬运」部分占比很高，这既是负担（要写更多代码），也是机会（更可控、更利于极致优化）。

7. 存储层次对性能优化的启示

理解存储层次后，昇腾算子优化立刻就有了方向：

减少 HBM 访问：能留在片上（UB/L2）的数据尽量复用，避免反复搬运。
提升搬运效率：DataCopy 尽量连续、对齐，一次搬大块数据，避免零散访问。
算搬重叠（流水线）：让数据搬运与计算并行（双缓冲/多缓冲），隐藏搬运延迟。
合理 Tiling：Tiling 大小要让数据恰好填满 UB/L0，既不多（放不下）也不少（浪费带宽）。

这些原则和 CUDA 优化中的「访存合并、共享内存复用、双缓冲」完全同构——性能优化的底层逻辑，跨硬件是相通的。

总结

昇腾存储层次：HBM（Global Memory）→ L2 → AI Core 内部（UB / L0A / L0B / L0C）。
AI Core 内部存储分工明确：UB 是向量计算的主场，L0A/L0B/L0C 是矩阵计算（Cube）的专用通路。
数据流 = 显式 DataCopy（GM↔UB）+ 计算单元执行 + 写回，搬运效率直接决定算子性能。
与 CUDA 内存模型高度同构，GPU 经验可以迁移；昇腾的特殊性在于「显式搬运」的权重更高。
优化的核心始终是：减少 HBM 访问、提高搬运效率、算搬重叠、合理 Tiling。

搜索