AscendC 硬件抽象架构MainScalar / MTE / Vector / Cube适用范围David 系列DaVinci AICore1. 架构总览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 四类单元职责单元全称职责对应流水MainScalarMain Scalar控制/发射配置参数、调用其他单元、流程控制Scalar 流水MTEMemory Transfer Engine搬运GM↔UB、UB↔L1、格式转换ND2NZ 等MTE 流水mte2/mte3VectorVector Unit向量计算elementwise、规约、激活、类型转换Vector 流水SIMD/SIMTCubeCube Unit矩阵计算GEMMMmad、wgmmaCube 流水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→NDDataCopy隐含转换Gather/Scatter非连续地址收集/分散Gather/Scatter APIPrefetch预取数据到 L2/L1.5preload2L1.5(gm)FixpipeL0C→UB 后处理量化/偏置fixpipe 操作3.3 MTE2 与 MTE3子流水方向典型操作MTE2外部→内部读GM→UB、GM→L1、L1→L0A/L0B、Buffer→UBMTE3内部→外部写UB→GM、L0C→UB3.4 格式转换ND2NZAI 框架数据多为 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、SoftmaxCopyUB 内拷贝Copy(ub_dst, ub_src)SIMT 计算线程级并行自定义复杂逻辑4.3 SIMD vs SIMT模式全称特点适用SIMD VFSingle Instruction Multiple Data Vector Fun定长向量VL256B规整向量计算SIMT VFSingle Instruction Multiple Thread Vector Fun多线程最多 2048 thread4 EU 并行复杂控制流、自定义算子5. Cube——矩阵计算单元5.1 定位Cube 是矩阵乘引擎核心是 MmadMatrix Multiply-Accumulate单元专为 GEMM 优化。┌──────────────── Cube Unit ──────────────────┐ │ │ │ L0A ──────┐ │ │ (A 矩阵) │ │ │ ├──► Mmad ──► L0C │ │ L0B ──────┘ (矩阵乘) (累加结果) │ │ (B 矩阵) │ │ │ │ 数据: L1 → L0A/L0B → Mmad → L0C │ │ 指令: mmad │ └──────────────────────────────────────────────┘5.2 核心能力能力说明API 示例GEMM矩阵乘 C A×Bmmad(l0c, l0a, l0b, M, K, N)FixpipeL0C 后处理量化/偏置/激活fixpipe 操作L0C→UB结果输出到 UBDavid 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 数据通路矩阵源 \ 目的GML1.5L1UBL0A/L0BL0CGM—MTE2MTE2MTE2——L1.5MTE3——MTE2——L1MTE3——MTE2MTE2—UBMTE3MTE3MTE3—MTE2—L0A/L0B—————MmadL0C———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 bufferDB流水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 ...;MTEDataCopy、LoadDataDataCopy(ub, gm, len)VectorAdds、Mul、Cast、ReduceAdds(out, in, 1.0f, len)Cubemmad、wgmmammad(l0c, l0a, l0b, M, K, N)同步SetFlag/WaitFlagSetFlagHardEvent::MTE2_VECTOR(0)8. Pipeline 并行与同步8.1 四流水并行模型┌─────────┐ 指令 ┌────────┐ 数据 ┌─────────┐ │ Scalar │────────►│ MTE │────────►│ Cube │ │ (控制) │ │ (搬运) │ │ (矩阵计算)│ └─────────┘ └────────┘ └─────────┘ │ 数据 │ ▼ ▼ ┌────────┐ 数据 ┌─────────┐ │ UB │◄───────│ Vector │ │ (缓冲) │ │ (向量计算)│ └────────┘ └─────────┘8.2 同步机制同步类型机制适用事件标志FlagSetFlag/WaitFlag HardEvent流水间同步MTE2→CubeBarrierbarrier()核内多线程同步cross-corecounter/notify核间同步Cluster synccluster.sync()Cluster 内跨 Block 同步wait_prev_task_doneEarly Start 接力同步Task 间接力8.3 HardEvent 事件对事件对含义MTE2_CUBEMTE2 搬运完成 → Cube 可用CUBE_FIXCube 计算完成 → Fixpipe 可用FIX_VECTORFixpipe 完成 → Vector 可用VECTOR_MTE3Vector 完成 → MTE3 可写回MTE2_VECTORMTE2 搬运完成 → Vector 可用8.4 多流水并行的代价问题原因度量同步代码行增加多 pipeline 需显式 Flag多 pipeline 并行度量pc 定位困难并行后无法精确定位出错 pc断点准确性度量pipe 切换开销stage 切换产生额外指令多 pipeline 并行度量Scalar BoundMainScalar 串行发射指令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