AscendC 硬件抽象架构：MainScalar / MTE / Vector / Cube
适用范围：David 系列DaVinci AICore
1. 架构总览
AscendC 将 AI Core 抽象为 4 类功能单元，对应 4 条并行流水线（pipeline），软件通过统一的 API 驱动它们协作。
┌──────────────────────────── AICore ────────────────────────────┐
│ │
│ ┌──────────────┐ 指令发射 ┌──────────────────────────────┐ │
│ │ MainScalar │───────────►│ 指令队列 / Scoreboard │ │
│ │ (控制单元) │ │ │ │
│ └──────────────┘ └──────┬───────┬───────┬────────┘ │
│ │ │ │ │ │
│ │ ┌──────────────┘ │ │ │
│ │ ▼ ▼ ▼ │
│ │ ┌────────────┐ ┌──────────┐ ┌────────┐ │
│ │ │ MTE │ │ Vector │ │ Cube │ │
│ │ │ (搬运单元) │ │ (向量计算)│ │(矩阵计算)│ │
│ │ └─────┬──────┘ └────┬─────┘ └───┬────┘ │
│ │ │ │ │ │
│ ▼ ▼ ▼ ▼ │
│ ┌────────┐ ┌──────────────────────────────────────────────┐ │
│ │ UB │ │ L1 / L0A / L0B / L0C │ │
│ │(UBuf) │◄─┤ (Cube 专用寄存器/L1) │ │
│ └────┬───┘ └──────────────────────────────────────────────┘ │
│ │ │ │
└───────┼──────────────────────────┼─────────────────────────────┘
▼ ▼
┌─────────┐ ┌──────────┐
│ GM │◄──────────────│ L2 │
│ (HBM) │ │ (Cache) │
└─────────┘ └──────────┘
1.1 四类单元职责
| 单元 |
全称 |
职责 |
对应流水 |
| MainScalar |
Main Scalar |
控制/发射：配置参数、调用其他单元、流程控制 |
Scalar 流水 |
| MTE |
Memory Transfer Engine |
搬运：GM↔UB、UB↔L1、格式转换（ND2NZ 等） |
MTE 流水（mte2/mte3） |
| Vector |
Vector Unit |
向量计算：elementwise、规约、激活、类型转换 |
Vector 流水（SIMD/SIMT） |
| Cube |
Cube Unit |
矩阵计算：GEMM（Mmad）、wgmma |
Cube 流水 |
1.2 设计哲学
| 原则 |
体现 |
| 通算分离 |
搬运（MTE）与计算（Vector/Cube）解耦，可并行 |
| 流水分工 |
每条流水独立调度，通过 scoreboard/barrier 同步 |
| 软件驱动 |
MainScalar 显式编排，编译器/程序员控制并行度 |
| 内存层级显式 |
UB/L1/L0 各级容量可见，软件 tiling 适配 |
2. MainScalar——控制/发射单元
2.1 定位
MainScalar 是 AICore 的控制核心，类似 CPU 的前端：取指、译码、发射、流程控制。
┌──────────────── MainScalar ────────────────┐
│ 取指 (ICache) │
│ ▼ │
│ 译码 │
│ ▼ │
│ 寄存器分配 (Sreg/Preg/GPR) │
│ ▼ │
│ 发射到各流水队列 ──► MTE queue │
│ ──► Vector queue │
│ ──► Cube queue │
│ ▼ │
│ 流程控制（分支/循环/同步） │
│ ▼ │
│ 栈管理（UB 模拟栈空间） │
└────────────────────────────────────────────┘
2.2 核心能力
| 能力 |
说明 |
| 标量计算 |
整数/地址运算，循环计数、地址计算 |
| 指令发射 |
向 MTE/Vector/Cube 队列发射计算/搬运指令 |
| 参数配置 |
配置 DMA 描述符、计算参数（shape/stride） |
| 流程控制 |
分支、循环、函数调用（vf_call） |
| 同步管理 |
发出 barrier、wait、notify 指令 |
3. MTE——Memory Transfer Engine 搬运单元
3.1 定位
MTE 是 AICore 的数据搬运引擎，负责所有跨内存层级的数据移动与格式转换。
┌──── MTE2（读路径）────┐
│ GM → L2 → L1 │
│ GM → L2 → UB │
│ L1 → UB │
GM ◄───────────►│ UB → L0A / L0B │
│ ND2NZ 格式转换 │
└───────────────────────┘
┌──── MTE3（写路径）────┐
│ UB → GM │
│ L0C → UB │
│ UB → L1 │
│ NZ2ND 格式转换 │
└───────────────────────┘
3.2 核心能力
| 能力 |
说明 |
API 示例 |
| DMA 搬运 |
GM↔UB、GM↔L1 大块搬运 |
DataCopy(ub, gm, len) |
| 格式转换 |
ND→NZ（行主序→分块Z序）、NZ→ND |
DataCopy 隐含转换 |
| Gather/Scatter |
非连续地址收集/分散 |
Gather/Scatter API |
| Prefetch |
预取数据到 L2/L1.5 |
preload2L1.5(gm) |
| Fixpipe |
L0C→UB 后处理（量化/偏置） |
fixpipe 操作 |
3.3 MTE2 与 MTE3
| 子流水 |
方向 |
典型操作 |
| MTE2 |
外部→内部（读） |
GM→UB、GM→L1、L1→L0A/L0B、Buffer→UB |
| MTE3 |
内部→外部（写） |
UB→GM、L0C→UB |
3.4 格式转换（ND2NZ）
AI 框架数据多为 ND（行主序），Cube 计算需 NZ（分块Z序）：
注意：ND2NZ 保序会引入头阻塞（SOC 指令保序冲突）
4. Vector——向量计算单元
4.1 定位
Vector 是向量计算引擎，处理 elementwise、规约、激活等 SIMT/SIMD 计算。
4.2 核心能力
| 能力 |
说明 |
典型算子 |
| Elementwise |
逐元素运算 |
Add、Mul、Relu、Sigmoid |
| 规约 |
沿轴规约 |
ReduceSum、ReduceMax |
| 类型转换 |
精度转换 |
Cast FP32→FP16、BF16→FP8 |
| 激活 |
激活函数 |
GELU、Swish、Softmax |
| Copy |
UB 内拷贝 |
Copy(ub_dst, ub_src) |
| SIMT 计算 |
线程级并行 |
自定义复杂逻辑 |
4.3 SIMD vs SIMT
| 模式 |
全称 |
特点 |
适用 |
| SIMD VF |
Single Instruction Multiple Data Vector Fun |
定长向量（VL=256B） |
规整向量计算 |
| SIMT VF |
Single Instruction Multiple Thread Vector Fun |
多线程（最多 2048 thread），4 EU 并行 |
复杂控制流、自定义算子 |
5. Cube——矩阵计算单元
5.1 定位
Cube 是矩阵乘引擎，核心是 Mmad（Matrix Multiply-Accumulate）单元，专为 GEMM 优化。
┌──────────────── Cube Unit ──────────────────┐
│ │
│ L0A ──────┐ │
│ (A 矩阵) │ │
│ ├──► Mmad ──► L0C │
│ L0B ──────┘ (矩阵乘) (累加结果) │
│ (B 矩阵) │
│ │
│ 数据: L1 → L0A/L0B → Mmad → L0C │
│ 指令: mmad │
└──────────────────────────────────────────────┘
5.2 核心能力
| 能力 |
说明 |
API 示例 |
| GEMM |
矩阵乘 C = A×B |
mmad(l0c, l0a, l0b, M, K, N) |
| Fixpipe |
L0C 后处理（量化/偏置/激活） |
fixpipe 操作 |
| L0C→UB |
结果输出到 UB |
David 100 新增 |
5.3 数据通路
Cube 计算数据流（完整路径）:
GM ──MTE2──► L1 ──MTE2──► L0A (A 矩阵)
──► L0B (B 矩阵)
L0A ──┐
├──► Mmad ──► L0C ──fixpipe──► UB ──MTE3──► GM
L0B ──┘
6. 内存层级与数据通路
6.1 内存层级总览
┌─────────────── 软件可见内存层级 ───────────────┐
│ │
│ GM (Global Memory / HBM) │ ← 大容量，跨核共享
│ ▲ │
│ │ MTE2/MTE3 │
│ ▼ │
│ L2 Cache (David 新增持久化) │ ← 自动缓存
│ ▲ │
│ │ │
│ ▼ │
│ L1 (Cube 专用) / UB (Vector 通用) │ ← 核内私有
│ ▲ │
│ │ MTE2 │
│ ▼ │
│ L0A / L0B (Cube 输入) / L0C (Cube 输出) │ ← 寄存器级
│ │
└───────────────────────────────────────────────┘
6.3 数据通路矩阵
| 源 \ 目的 |
GM |
L1.5 |
L1 |
UB |
L0A/L0B |
L0C |
| GM |
— |
MTE2 |
MTE2 |
MTE2 |
— |
— |
| L1.5 |
MTE3 |
— |
— |
MTE2 |
— |
— |
| L1 |
MTE3 |
— |
— |
MTE2 |
MTE2 |
— |
| UB |
MTE3 |
MTE3 |
MTE3 |
— |
MTE2 |
— |
| L0A/L0B |
— |
— |
— |
— |
— |
Mmad |
| L0C |
— |
— |
— |
fixpipe |
— |
— |
"—"表示无直接通路，需经中间层级中转。
7. 四单元协作编程模型
7.1 经典 GEMM 算子流程
时间 ──────────────────────────────────────────────────────────►
MainScalar: [配置DMA_A][配置DMA_B][wait][配置MMAD][wait][配置Fix][配置DMA_C]
│ │ │ │ │
MTE2: [搬运A→L1] [搬运B→L1] [搬运C→GM]
│ │ │ ▲
│ │ │
Cube: [LoadData A→L0A] │
[LoadData B→L0B] │
│ │
[Mmad L0C] │
│ │
Fixpipe: [L0C→UB] │
│ │
Vector: [Add bias] ──────────────────────┘
│
MTE3: [UB→GM]
7.2 Double Buffer 流水
为隐藏搬运延迟，采用 double buffer（DB）流水：
MainScalar: [cfg_DMA0][cfg_DMA1][wait0][cfg_MMAD0][cfg_DMA2][wait1][cfg_MMAD1]...
MTE2: [搬运buf0][搬运buf1][搬运buf2] ...
Cube: [LoadData0][MMAD0] [LoadData1][MMAD1] ...
Vector: [Fix0] [Fix1] ...
MTE3: [Write0] [Write1] ...
理想: MTE 与 Cube/Vector 完全重叠，吞吐翻倍
7.3 AscendC 编程接口映射
| 单元 |
AscendC API 类 |
示例 |
| MainScalar |
参数计算、流程控制 |
int tile = ...; |
| MTE |
DataCopy、LoadData |
DataCopy(ub, gm, len) |
| Vector |
Adds、Mul、Cast、Reduce |
Adds(out, in, 1.0f, len) |
| Cube |
mmad、wgmma |
mmad(l0c, l0a, l0b, M, K, N) |
| 同步 |
SetFlag/WaitFlag |
SetFlag<HardEvent::MTE2_VECTOR>(0) |
8. Pipeline 并行与同步
8.1 四流水并行模型
┌─────────┐ 指令 ┌────────┐ 数据 ┌─────────┐
│ Scalar │────────►│ MTE │────────►│ Cube │
│ (控制) │ │ (搬运) │ │ (矩阵计算)│
└─────────┘ └────────┘ └─────────┘
│ 数据 │
▼ ▼
┌────────┐ 数据 ┌─────────┐
│ UB │◄───────│ Vector │
│ (缓冲) │ │ (向量计算)│
└────────┘ └─────────┘
8.2 同步机制
| 同步类型 |
机制 |
适用 |
| 事件标志（Flag） |
SetFlag/WaitFlag + HardEvent |
流水间同步（MTE2→Cube） |
| Barrier |
barrier() |
核内多线程同步 |
| cross-core |
counter/notify |
核间同步 |
| Cluster sync |
cluster.sync() |
Cluster 内跨 Block 同步 |
| wait_prev_task_done |
Early Start 接力同步 |
Task 间接力 |
8.3 HardEvent 事件对
| 事件对 |
含义 |
MTE2_CUBE |
MTE2 搬运完成 → Cube 可用 |
CUBE_FIX |
Cube 计算完成 → Fixpipe 可用 |
FIX_VECTOR |
Fixpipe 完成 → Vector 可用 |
VECTOR_MTE3 |
Vector 完成 → MTE3 可写回 |
MTE2_VECTOR |
MTE2 搬运完成 → Vector 可用 |
8.4 多流水并行的代价
| 问题 |
原因 |
度量 |
| 同步代码行增加 |
多 pipeline 需显式 Flag |
多 pipeline 并行度量 |
| pc 定位困难 |
并行后无法精确定位出错 pc |
断点准确性度量 |
| pipe 切换开销 |
stage 切换产生额外指令 |
多 pipeline 并行度量 |
| Scalar Bound |
MainScalar 串行发射指令 |
Scalar bound 度量 |
8.5 算子流水编排原则
| 原则 |
说明 |
| 通算分离 |
MTE 搬运与 Cube/Vector 计算尽量并行 |
| Double Buffer |
用 2 倍 UB/L1 空间隐藏搬运延迟 |
| Tile 大小适配 |
tile 大小匹配 UB/L1/L0 容量，避免溢出 |
| 同步最小化 |
只在数据依赖处插入 Flag，减少等待 |
| Scalar 精简 |
MainScalar 代码尽量少，避免 Scalar Bound |
| 延迟 Copy |
无需转换时只记录地址，到 L0 才真实 Copy |
3
0
已为社区贡献27条内容
所有评论(0)