NVIDIA A100 和 H100 硬件架构学习

NVIDIA A100 和 H100 硬件架构学习

引用

CSDN

1.

https://blog.csdn.net/qq_27763933/article/details/140904201

NVIDIA GPU架构代号概述

NVIDIA GPU有多个代号和架构，这些架构对应不同的世代和硬件特性。以下是NVIDIA主要GPU架构及其计算能力（Compute Capability）代号的简要概述：

1. Tesla架构

• G80、GT200
• Compute Capability: sm_10, sm_11, sm_12, sm_13

2. Fermi架构

• GF100, GF104, GF110
• Compute Capability: sm_20, sm_21

3. Kepler架构

• GK104, GK110
• Compute Capability: sm_30, sm_32, sm_35, sm_37

4. Maxwell架构

• GM107, GM204, GM206
• Compute Capability: sm_50, sm_52

5. Pascal架构

• GP100, GP102, GP104, GP106
• Compute Capability: sm_60, sm_61, sm_62

6. Volta架构

• GV100
• Compute Capability: sm_70, sm_72

7. Turing架构

• TU102, TU104, TU106
• Compute Capability: sm_75

8. Ampere架构

• GA100, GA102, GA104, GA106
• Compute Capability: sm_80, sm_86

9. Hopper架构

• H100
• Compute Capability: sm_90, sm_90a

这只是一个简要概述，具体的GPU型号可能会包含多种不同的子配置和强化特性，例如更多的CUDA核心、更高的内存带宽、更强的NVLink支持等。详细的功能和特性可以通过NVIDIA的最新文档和白皮书来获得。

举例说明：

• Tesla G80 和 GT200: 最早的GPU架构，主要用在基础的并行计算。
• Fermi: 引入了新的指令集架构和硬件功能，例如ECC内存支持。
• Kepler: 提升了能效，广泛应用于高性能计算和科学计算。
• Maxwell: 进一步优化了能效并改善了执行效率。
• Pascal: 引入了NVLink和统一内存，显著提高了深度学习的性能。
• Volta: 包含全新的Tensor Cores，用于加速深度学习任务。
• Turing: 包含了Ray Tracing Cores和改进的Tensor Cores，针对实时渲染和深度学习进行了优化。
• Ampere: 进一步增强了Tensor Cores性能，改善的memory和计算效率。
• Hopper: 最新的架构，进一步提升AI和数据中心计算的效率。

编译CUDA程序时，可以选择适合你的NVIDIA GPU架构的-arch参数。例如，如果你有一块Volta GPU，你可以这样编译程序：

nvcc -arch=sm_70 your_program.cu -o your_program

Hopper架构新增硬件特性

了解FlashAttentionV3的优化需要先了解Hopper的主要技术（Hopper White Paper概述）：

• NVIDIA Hopper Architecture In-Depth
• NVIDIA Ampere Architecture In-Depth

1. 新的第四代Tensor Core，整体加速了6x，单SM上的加速，SM数量的增加，频率升高，在同等数据类型上，张量内核的MMA（矩阵乘积）计算速度是A100 SM的2倍；同时支持了fp8的数据类型，与A100 fp16数据类型相比tensor core性能提升了4倍；
2. 新的DPX指令，相比A100在动态规划算法上加速7x；
3. IEEE FP64和F32相比A100加速3x，其中硬件计算单元提升2x，SM数量增加，频率升高；
4. 新增Thread block cluster的特性，编程层次变为：threads，thread blocks，thread block clusters，and grids。clusters使多个thread blocks能够在多个SMs上并发运行，同步，协同获取和交换数据；
5. Distributed shared memory，实现SM-to-SM的通信，用于跨多个SM共享内存块的加载、存储和原子操作。
   
6. 新的异步执行的特性，包括Tensor Memory Accelerator（TMA）单元。可以将大的数据块从GMEM高效的传输到SMEM，同时支持同一个cluster内不同的Thread blocks间，异步copy数据。
7. 新的Transformer Engine（硬件+软件），可以实现Fp16和Fp8的自动切换，训练加速9x，推理加速30x。
8. HBM3 memory subsystem提升2x的bandwidth。
9. 50MB的L2 cache的架构。
10. 第二代MIG（Multi-Instance GPU）每个GPU Instance增加3x的计算能力和2x的bandwidth；
11. 可信计算支持，保护用户数据；
12. 第四代NVIDIA NVLink，3x bandwidth在allreduce操作上。和50%的通用bandwidth提升的支持；
13. 第三代NVSwitch，总的switch throughput从7.2Tbits/sec提升到13.6Tbits/sec；
14. 新的NVLink Switch system；
15. PCIe Gen5支持128GB/sec的双向bandwidth（64GB/sec的单向带宽）。

疑问：TF32并没有增加芯片的峰值算力，为什么不直接将tensor core设计成支持fp32的类型？（降低能耗？）

Hopper架构优化的pipeline效果

核心思想：减少data_load、cuda core、tensor core对寄存器资源的竞争关系，加大pipeline hide latency效果

疑问：根据register file大小，理论上每个thread最多可以访问到512*32bit的registers(为什么文档说最多是256个registers?Flashattention3中register分配数量超过了256，达到264个)

TMA硬件单元

TMA的引入解放了load数据和计算，TMA不再和计算单元抢占register/thread资源，hide load数据的latency；
（类似biren br104 TDA硬件单元）

说明：
（1）通过copy descriptor的方式只需一次issue就可以完成global memory到share memory之间的async copy；
(2) TMA（只用到一个thread）解放了thread和register资源，去做其他independent工作；
(3) 支持一种全新的更高效的异步事务屏障（asynchronous transaction barrier）来处理数据copy和exchange，cluster内不同SM之间的数据通信也是基于这种新特性。

WGMMA指令

WGMMA指令的引入，合并SM里面的4个tensor core效果类似于一个大的tensor core，减少load tensor次数（A/B tensor共用），同时支持Tcore core的inputs来源于share memory（A100架构及之前的架构，inputs必须from registers），具体的WGMMA指令inputA from registers or share memory，inputB must from share memoryPTX ISA 8.5；减少了register的抢占，更有利于cuda core pipeline并行计算，hide cuda core计算的latency；
（类似biren br104 cwarps/Tmode概念）

setmaxnreg指令

setmaxnreg指令的引入，支持动态重新分配每个warp group可用register数量（from register pool）；
说明：Hopper架构新特性的指令大部分都是在PTX ISA version 8.0引入的PTX ISA 8.5
（类似biren br104 cwarps/Tmode下手都分配register用法）

FP8 Tensor Core

Hopper整体上支持FP8, FP16, BF16, TF32, FP64这些dtype类型的tensor core的计算，相比Ampere，fp8是新增加的数据类型：

FP8 Tensor Core支持FP32 and FP16两种类型的累加器, 并且支持两种FP8的输入类型:

• E4M3 with 4 exponent bits, 3 mantissa bits, and 1 sign bit（范围较小，精度较高）
• E5M2, with 5 exponent bits, 2 mantissa bits, and 1 sign bit（范围较大，精度较低）

flash attention3论文上也提到两点关于FP8在flash attention3上使能的工程细节：
（1）A,B tensor必须在K维度连续(V in-kernel transpose)；
（2）FP32 accumulator register layout is different from operand A FP8 operand register layout（QK结果permute）

H100 FP8相比A100 FP16提升了6x的吞吐量
横向对比tensor core计算，H100相比A100都有3x吞吐量的提升