NVIDIA A100 和 H100 硬件架构学习

1. 目前位置NV各种架构代号：

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

1. Tesla 架构

• G80、GT200

• Compute Capability: sm_10, sm_11, sm_12, sm_13

1. Fermi 架构

• GF100, GF104, GF110

• Compute Capability: sm_20, sm_21

1. Kepler 架构

• GK104, GK110

• Compute Capability: sm_30, sm_32, sm_35, sm_37

1. 4. Maxwell 架构

• GM107, GM204, GM206

• Compute Capability: sm_50, sm_52

1. Pascal 架构

• GP100, GP102, GP104, GP106

• Compute Capability: sm_60, sm_61, sm_62

1. Volta 架构

• GV100

• Compute Capability: sm_70, sm_72

1. Turing 架构

• TU102, TU104, TU106

• Compute Capability: sm_75

1. Ampere 架构

• GA100, GA102, GA104, GA106

• Compute Capability: sm_80, sm_86

1. 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 程序

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

nvcc -arch=sm_70 your_program.cu -o your_program

1. Hopper 相比 Ampere 新增硬件特性

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

https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/

https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/

• 新的第四代的Tensor Core，整体加速了6x，单SM上的加速，SM数量的增加，频率升高，在同等数据类型上，张量内核的 MMA（矩阵乘积）计算速度是 A100 SM 的 2 倍；同时支持了fp8的数据类型，与A100 fp16数据类型相比 tensor core 性能提升了 4 倍；

• 新的DPX指令，相比A100在动态规划算法上加速7x @黄明晓 应用场景调研；

• IEEE FP64和F32相比A100加速3x，其中硬件计算单元提升2x，SM数量增加，频率升高；

• 新增Thread block cluster的特性，编程层次变为：threads，thread blocks，thread block clusters， and grids。clusters 使多个thread blocks能够在多个 SMs 上并发运行，同步，协同获取和交换数据；

• Distributed shared memory，实现SM-to-SM的通信，用于跨多个 SM 共享内存块的加载、存储和原子操作。

• 新的异步执行的特性，包括Tensor Memory Accelerator（TMA）单元。可以将大的数据块从GMEM高效的传输到SMEM，同时支持同一个cluster内不同的Thread blocks间，异步copy数据。

• 新的Transformer Engine（硬件+软件）, 可以实现Fp16和Fp8的自动切换，训练加速9x，推理加速30x。

• HBM3 memory subsystem提升2x的bandwidth。

• 50MB的L2 cache的架构。

• 第二代MIG（Multi-Instance GPU）每个GPU Instance增加3x的计算能力和2x的bandwidth；

• 可信计算支持，保护用户数据；

• 第四代NVIDIA NVLink，3x bandwidth 在 allreduce操作上。和50%的通用bandwidth提升的支持；

• 第三代NVSwitch，总的switch throughput从7.2Tbits/sec提升到13.6Tbits/sec；

• 新的NVLink Switch system；

• PCIe Gen5支持128GB/sec的双向bandwidth（64GB/sec的单向带宽）。

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

1. Hopper 更优的pipeline效果

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

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

1. 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之间的数据通信也是基于这种新特性。

1. 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 memory PTX ISA 8.5；减少了register的抢占，更有利于cuda core pipeline 并行计算，hide cuda core 计算的latency；

（类似biren br104 cwarps/Tmode 概念）

1. setmaxnreg指令

setmaxnreg指令的引入，支持动态重新分配每个warp group 可用register数量（from register pool）；

说明：Hopper架构新特性的指令大部分都是在PTX ISA version 8.0引入的 PTX ISA 8.5

（类似biren br104 cwarps/Tmode下手都分配register用法）

1. 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吞吐量的提升