GPU 与 NPU 硬件组成详解
阅读提示：本文与《让深度学习跑得飞快：从第一性原理理解性能》配套。那篇从算力（Compute）/ 带宽（Bandwidth）/ 开销（Overhead）三个维度讲性能瓶颈，本文则从硬件角度回答「这三个维度对应芯片里的哪些部件」。
现代深度学习加速器虽然品牌、架构各异，但几乎都由五大部分构成：计算单元、存储层次、片上互连、控制调度、I/O 接口。理解了这五部分的组成与协作方式，就能明白为什么性能优化时矩阵算力、内存带宽和 CPU 开销这三个瓶颈会反复出现。
本文以 NVIDIA GPU（A100/H100）与华为昇腾 NPU（910B）为主要示例。
1. 计算单元（Compute）——干活的工厂
计算单元是芯片执行浮点运算的地方，也是我们最想喂饱的部分。现代加速器普遍采用异构计算设计：为矩阵乘法专门定制超强算力，同时保留通用单元处理其他算子。
1.1 GPU：流式多处理器（SM）与计算核心
SM（Streaming Multiprocessor，流式多处理器）是 NVIDIA GPU 组织计算资源的基本单元。一颗 GPU 由几十到上百个 SM（流式多处理器）组成（A100 SXM 为 108 个，H100 SXM 为 132 个），每个 SM（流式多处理器）内部包含：
- CUDA Core：通用浮点/整数运算单元，执行逐元素、激活、归约等操作。A100 每个 SM（流式多处理器）有 64 个 FP32 CUDA Core，其算力比 Tensor Core 慢约一个数量级。
- Tensor Core：专为矩阵乘法（matmul）设计的单元，贡献峰值算力的大头。A100 的 Tensor Core 在 TF32 稀疏/FP16 稠密两种主流口径下约 312 TFLOPS；H100 在此基础上引入 FP8，峰值进一步跃升（SXM 稠密口径约 1,979 TFLOPS）。
- 寄存器堆（Register File）：SM（流式多处理器）内最快的存储，供当前驻留的线程直接使用。
- 共享内存（Shared Memory / SRAM）：块内可编程缓存，用于块内线程协作与数据复用。A100 每个 SM（流式多处理器）最多约 164 KB。
- L1 缓存：与共享内存在物理上同属一块片上 SRAM，可按需划分。
- Warp Scheduler + 分发单元：以线程束（warp，32 个线程）为单位调度并发射指令。A100 每个 SM（流式多处理器）有 4 个 Warp Scheduler。
在 CUDA 编程模型中，kernel 启动后会生成大量线程，线程按 warp（32 个）和 block 两级组织，block 被调度到某个 SM（流式多处理器）上执行；一个 SM（流式多处理器）可以同时驻留多个 block（上限由寄存器与共享内存容量决定）。因此 SM（流式多处理器）数量 × 每 SM（流式多处理器）的寄存器/共享内存容量共同决定了 GPU 的并发能力（能同时驻留多少线程），这是内存带宽之外又一个影响实际吞吐的关键因素。
对应关系：昇腾 NPU 中的「AI Core」在定位上与 GPU 的「SM（流式多处理器）」类似——都是芯片内并行执行计算的基本单元，只是内部组织不同（NPU 是 Cube/Vector/Scalar 三分结构，GPU 是 CUDA Core + Tensor Core + SFU）。
1.2 NPU：Cube + Vector + Scalar
昇腾 NPU 采用类似的三级异构结构，但命名和分工不同。以昇腾 AI Core 为例，它由三个计算单元加一个片上缓冲组成：
| 单元 | 对应 GPU | 职责 |
|---|---|---|
| Cube（方阵） | Tensor Core | 矩阵乘法，AI Core 的核心算力来源 |
| Vector（向量） | CUDA Core | 逐元素运算、激活函数、池化 |
| Scalar（标量） | SFU / 控制逻辑 | 地址计算、循环控制、分支跳转、指令调度 |
| Unified Buffer（统一缓冲区） | 共享内存 / L1 | 片上数据暂存，供 Cube 与 Vector 消费 |
1.3 关键结论：矩阵单元决定了峰值算力
无论 GPU 还是 NPU，矩阵单元都贡献了 绝大部分峰值 FLOPS，非 matmul 算子在总 FLOPS 中只是「四舍五入误差」（如 BERT 中仅占约 0.2%）。因此：
- 你的模型 是不是 matmul 密集，直接决定能否贴近峰值算力。
- 不使用矩阵单元时，A100 的通用 FP32 算力只有约 19.5 TFLOPS（而 Tensor Core 在 TF32 下是 312 TFLOPS，差 16 倍）；注意这已是按 FMA（每条指令算一次乘加 = 2 FLOP）统计的结果。昇腾 NPU 同理，非矩阵运算的 Vector 单元算力远低于 Cube。
2. 存储层次（Memory）——仓库体系
算力单元再快，数据喂不进去也是白搭。存储层次解决「数据放哪、怎么流动」的问题。
2.1 从寄存器到 HBM
存储层次自顶向下依次为（以 GPU 为例，NPU 结构类似）：
寄存器堆（Register File） 容量极小，速度最快，算子内部直接使用
↓
共享内存 / L1 Cache（SRAM） 每 SM（流式多处理器）几十 KB ~ 几百 KB，块内共享
↓
L2 Cache 芯片级共享，约几十 MB（A100 为 40 MB）
↓
全局内存 DRAM（HBM） 几十 GB ~ 上百 GB，即 `nvidia-smi` 里那块
↓
主机内存（CPU RAM） 跨设备搬运的最终来源- 数据在寄存器堆 ↔ HBM 之间的流动，是带宽与延迟成本的来源。
- L2 Cache 承担跨 SM（流式多处理器）的数据复用与聚合，对大模型推理的注意力计算尤为关键。
- 数据搬运由两类机制完成，职责不同：
- 访存指令（LD/ST）：kernel 内部、线程执行到访存指令时，把数据从 HBM 经 L2 载入寄存器/共享内存。这是内存带宽受限区间的主要开销来源。
- 搬运引擎（Copy Engine / DMA）：独立于计算核的硬件单元，负责主机↔设备（H2D/D2H）、以及设备内存之间的批量搬运，可与计算重叠执行。
2.2 GPU 与 NPU 的存储分工
- GPU：每个 SM（流式多处理器）有自己的寄存器堆与共享内存（SRAM，与 L1 共享一块片上存储）；数据通过全局内存（HBM）进出，靠 L2 Cache 做中转。kernel 内数据的载入由访存指令（LD/ST）完成；此外有独立的 Copy Engine（DMA）负责 H2D/D2H 等主机侧搬运，可与计算重叠。正是「每次 kernel 都要从 DRAM 进出」造就了内存带宽受限区间。
- NPU：昇腾的 Unified Buffer（统一缓冲区）是片上 SRAM，数据由 DMA 搬运引擎在 HBM 与 UB 之间搬进搬出，再由 Cube 高效消费。NPU 对「数据一定要先落片上再计算」的约束比 GPU 更强，因此向量化搬运（DMA）与数据排布（format）管理对性能影响极大。
2.3 SRAM 与 HBM：片上高速缓存 vs 片外大容量内存
存储层次里最核心的一对对比，是片上 SRAM与片外 HBM。
SRAM（Static Random Access Memory，静态随机存取存储器）
- 位置：片上，即芯片内部，靠近计算单元。
- 特点：速度快（几个时钟周期即可访问）、容量小（KB ~ 几 MB 级）、价格昂贵。
- 作用：作为高速暂存区，让数据在进入矩阵单元前「就近等待」。GPU 的寄存器堆、共享内存、L1/L2 Cache 本质上都是 SRAM；NPU 的 Unified Buffer 同样是 SRAM。
- 性能影响：数据能否留在 SRAM 里反复复用，直接决定能否避开带宽瓶颈。算子融合的本质，就是让中间结果留在片上 SRAM 而不是反复往返 HBM。
HBM（High Bandwidth Memory，高带宽内存）
- 位置：片外，与芯片封装在一起（通过硅中介层垂直堆叠），但仍然在芯片外部。
- 特点：容量大（几十 ~ 上百 GB）、带宽高（A100 约 1.6~2 TB/s、H100 约 3.35 TB/s）、延迟比 SRAM 高约两个数量级（HBM 访问延迟约 100~200 ns，片上 SRAM 仅几纳秒）。
- 作用：充当「大仓库」，存放模型权重、KV Cache、激活和通信 Buffer。
nvidia-smi里显示的显存、昇腾里的内存，指的就是这部分。 - 性能影响：HBM 带宽是内存带宽受限区间的硬上限。算力每代翻倍远超带宽翻倍，「带宽追不上算力」是长期趋势，所以让数据尽量在 SRAM 里复用、减少 HBM 往返，是几乎所有推理优化（算子融合、KV Cache 复用、量化减字节）的共同目标。
| 特性 | SRAM（片上） | HBM（片外） |
|---|---|---|
| 位置 | 芯片内部 | 芯片外部，与芯片封装在一起 |
| 容量 | 每 SM（流式多处理器）几十 ~ 几百 KB，全局几 MB ~ 几十 MB | 几十 GB ~ 上百 GB |
| 访问延迟 | 几纳秒（约 2~5 个时钟周期） | 约 100~200 ns（比 SRAM 高约 2 个数量级） |
| 带宽 | 极高（紧邻计算单元，聚合带宽可达 TB/s 级，远高于 HBM） | 高（A100 约 2 TB/s、H100 约 3.35 TB/s），但比 SRAM 低一个数量级 |
| 典型成员 | 寄存器堆、共享内存、L1/L2 Cache、Unified Buffer | 全局内存（显存） |
一句话：SRAM 是「工作台」，HBM 是「仓库」。性能优化的核心，就是把东西尽量放在工作台上用（复用、融合），少往仓库来回搬。
带宽是硬约束：HBM 带宽决定内存带宽受限区间的上限（A100 约 2 TB/s、H100 约 3.35 TB/s）。算力每代翻倍的速度远超带宽翻倍速度，因此「带宽跟不上算力」是长期趋势——这正是《brrr》一文里强调的核心矛盾。
3. 片上互连（Interconnect）——工厂之间的传送带
计算单元之间、芯片之间、节点之间的数据搬运，靠互连完成。
| 层级 | GPU | NPU（昇腾） |
|---|---|---|
| 片内 | SM（流式多处理器）之间、SM（流式多处理器）与 L2 之间的片上网络（NoC） | AI Core 之间、与 HBM 控制器之间的 NoC |
| 片间 | NVLink（每链路约 50~100 GB/s 级）+ NVSwitch（全互联交换） | HCCS（华为高速互联）/ RoCE |
| 跨节点 | InfiniBand / 以太网（RoCEv2） | RoCE / 以太网 |
- GPU 集群中，NVSwitch 实现多 GPU 之间的全互联（All-to-All），是张量并行（TP）高效通信的硬件基础（如 DGX A100 8 卡、DGX H100 8 卡）。
- 昇腾集群中，HCCS 负责片间高速互联，多机场景再叠加 RoCE 网络。
- 片间互连带宽直接影响张量并行（TP）和模型并行（MP）的通信开销，是分布式推理与训练能否扩展的关键。
4. 控制调度（Overhead）——指令分发
4.0 先理解 kernel：让硬件干活的最小单元
kernel（内核/核函数）是「在 GPU/NPU 上执行的一段并行计算函数」，也是我们让加速器干活的最小单元。CPU 上一个普通的加法函数是串行循环：
void add(float* a, float* b, float* c, int n) {
for (int i = 0; i < n; i++) {
c[i] = a[i] + b[i]; // 串行循环
}
}同样的「逐元素相加」，写成 CUDA kernel 后，成千上万个线程会并行执行它：
__global__ void add(float* a, float* b, float* c) {
int i = blockIdx.x * blockDim.x + threadIdx.x; // 每个线程算自己的下标
c[i] = a[i] + b[i]; // 每个线程只处理一个元素
}__global__是声明 kernel 的关键字；调用add<<<grid, block>>>(a, b, c)时，CPU 只负责「启动」，真正的计算在 SM（流式多处理器）上并行完成。- kernel 启动是异步的：CPU 发出启动指令后不等它完成就继续执行，这正是 CPU 能「跑在 GPU 前面」的原因。
- 启动有固定开销：每次启动一个 kernel 都要付出微秒级的 CPU 派发成本（CUDA 单次启动约 3~10 μs）。所以小 kernel 越多越慢——算子融合与 CUDA Graphs / 图模式（可把启动开销压到亚微秒级）正是为消除这一成本而生。
- kernel 之间有依赖需要同步：如果 kernel B 需要 kernel A 的结果，必须显式等待（如
cudaDeviceSynchronize()或 Stream 间的依赖），否则会读到未完成的数据。
kernel 把「计算单元、存储层次、控制调度」三个硬件部分串了起来：线程被分派到 SM（流式多处理器）/ AI Core（计算单元），执行时数据在 HBM 与 SRAM 之间搬运（存储层次），而每次启动本身要消耗 CPU 派发时间（控制调度/开销）。
4.1 谁来「发号施令」
- 命令处理器（Command Processor）：接收 CPU 下发的 kernel 启动指令，排队调度。GPU 中对应 GigaThread Engine（全局线程调度器），负责把线程块分发到各 SM（流式多处理器）。
- 分发逻辑：把 kernel 分配到各个 SM（流式多处理器）/ AI Core。
- GPU 的每个 SM（流式多处理器）内部还有 Warp Scheduler（线程束调度器）：CUDA 以 warp（32 个线程）为最小调度单位，调度器为每个 warp 选择可发射的指令并交给执行单元。A100 每个 SM（流式多处理器）有 4 个 Warp Scheduler，每个时钟周期可发射一条 warp 指令。
- Stream 与异步执行：CUDA 的 Stream（流）让多个 kernel 可并行或重叠执行，GigaThread Engine 负责跨 Stream 的调度。
- NPU 没有 warp 概念，由标量单元解析指令、驱动矢量/矩阵单元流水线，配合核间同步与事件机制完成多核调度。
Warp 调度的意义：同一 warp 内的线程执行相同的指令（SIMT 模型），如果线程分支发散（divergence），warp 需要串行执行各分支，浪费吞吐。因此 kernel 中应尽量避免 warp 内分支，这是 GPU 优化的一条常用经验。
4.2 这是「开销」的硬件根源
CPU 派发一条 kernel 的时间如果超过 GPU 执行它的时间，GPU 就会闲置。这就是《brrr》一文所说的 Overhead-Bound（开销受限）区间：
- kernel 太小 → 派发开销 >> 执行时间 → GPU 变成「昂贵的镇纸」。
- 解法：算子融合、CUDA Graphs / 图模式、trace、真正的 JIT——本质都是减少「派发-执行」的次数。
昇腾的图模式（ACL Graph）与 GPU 的 CUDA Graphs 异曲同工，都是把一串 kernel 预编译成一张图，一次性提交，大幅降低派发开销。
5. I/O 与外部接口
- PCIe：CPU ↔ 加速器的主数据传输通道（影响数据搬运与模型加载）。PCIe Gen5 x16 单向约 63 GB/s。
- Copy Engine / DMA：负责主机与设备之间的异步数据搬运，可与计算重叠，减少 CPU 等待。
- 视频编解码单元：GPU 的 NVDEC/NVENC，昇腾也有集成的视频处理单元，多模态场景（图像/视频输入）会用到。
- 网络接口：GPU 通常外接网卡（InfiniBand HCA 或 RoCE NIC），部分昇腾板卡将 RoCE 网络控制器集成在芯片/模组上，减少外部依赖、降低数据拷贝开销。
- 其余外围：HBM 控制器、PCIe 控制器等，GPU 与 NPU 都具备，只是实现细节不同。
6. 总结：五大部分与三大瓶颈的对应关系
| 硬件组成 | 对应瓶颈 | 性能优化的核心手段 |
|---|---|---|
| 计算单元（Tensor Core / Cube） | 算力受限（Compute-Bound） | 用 Tensor Core / Cube、提高 matmul 占比 |
| 存储层次（SRAM ↔ HBM） | 带宽受限（Bandwidth-Bound） | 算子融合、数据排布、减少全局内存往返 |
| 控制调度（Command Processor / 派发） | 开销受限（Overhead-Bound） | CUDA Graphs / 图模式、trace、JIT |
| 片上互连（NVLink / NVSwitch / HCCS / RoCE） | 通信受限（分布式场景） | 张量并行、通信优化、PD 分离 |
| I/O 接口（PCIe / Copy Engine / 网络） | 数据搬运（非计算路径） | 异步传输、零拷贝、NVMe 直连 |
核心思想：让矩阵单元持续「吃饱」——数据搬运（带宽）和指令派发（开销）都不能成为瓶颈。判断当前处于哪个区间，再决定优化手段，比盲目套用技巧有效得多。