Grace Hopper

资料

关键特性

  • 计算架构：sm90
  • H100 + InfiniBand性能是A100的30x；
  • H100 + NVLink性能是H100 + InfiniBand的3x；
  • 与A100相比，H100的第四代Tensor Core有较大的提升：6x 芯片间速度，更快的SM，更多的SM，更高的clocks；同样的数据类型、数据量下，H100 SM计算速率是A100 SM的2x；
  • New thread block cluster: 支持比单个SM上的单个thread block更大的粒度的编程局部性支持；扩展了CUDA编程模型：thread, thread block, thread block cluster, grid；
  • Distributed shared memory: 支持直接SM-to-SM通信：跨域多个SM SMEM的数据交互；
  • New asynchronous execution + TMA(Tensor Memory Accelerator): TMA允许直接在HBM与SMEM之间传输大块的数据，也支持在一个thread block cluster中异步thread block拷贝的数据；
  • New asynchronous barrier: 新的异步事务屏障；
  • HBM3: 相比于A100的HBM2e，带宽提升2x，达到3TB/s；
  • 60MB L2 cache (GH100), 50MB L2 (H100 SXM5主板), 50MB L2 (H100 PCIe Gen5)；
  • 第二代MIG技术；
  • 第四代NVLink技术：all-reduce操作带宽提升3x，900GB/s，是PCIe Gen5的7x；
  • PCIe Gen5: 提供128GB/s总带宽，单个方向64GB/s；
  • thread block: 编程抽象，代表一组cooperative threads，一个SM能执行多个thread blocks(32个)，一个thread blocks不能跨多个SM；
  • thread block cluster的thread block(1cluster最多16个)可以跨SM，因DSMEM的存在支持SM-to-SM通信（虚拟地址空间统一），所以cluster内的thread blocks允许访问同一个SM内的SMEM，也支持访问cluster内的其它SM中的SMEM；

异步

  • 真正的异步GPU：支持应用构建出E2E的异步pipeline：允许将数据into and off the chip，完全重叠和隐藏计算与数据移动；
  • 只要少量CUDA thread通过使用TMA来完全利用起全部的内存带宽，同时其它剩余的CUDA thread聚焦于纯粹的计算，如Tensor Core计算的前后数据处理；

数据

  • 8个GPC (GPU processing clusters)
  • GH100: 144个SM
  • H100+SXM5: 132个SM, 80GB HBM3
  • H100+PCIe5: 114个SM, 80GB HBM2e
  • 128个FP32 CUDA Cores per SM
  • 4个第四代Tensor Core per SM

GH100 with 144 SMs

H100 SM 架构

  • 下图即为一个H100 SM(streaming multiprocessor)的架构图：

H100 SM

H100 SM

  • chip-to-chip对比A100有6x提升 （3x chip-to-chip processing rate + 2x clock-for-clock）；
  • 同FP16来对比，Tensor Core相比A100提升2x，如采用FP8则是4倍；
  • Sparsity相比标准Tensor Core运算，性能又能提升2x；
  • asynchronous execution + TMA + asynchronous transaction barrier
  • thread block cluster
  • SM-to-SM SMEM直接通信；

4th Gen Tensor Core (TC)

  • Tensor Core: 硬件MMA（matrix multiply and accumulate）计算核心；
  • 首次在V100中引入；
  • 与A100相比，4th Gen Tensor Core为每个SM提升了2x的运算能力；
  • 支持数据类型：INT8, FP16, BF16, TF32, INT8 MMA；

4th Tensor Core

FP8

  • H100支持了FP8 Tensor Core：支持FP32, FP16 accumulators，支持两种类型的FP8格式，如下图所示：

FP8 Tensor Core

  • FP8 vs FP16: 吞吐翻倍，内存减半

Thread block clusters

  • CUDA编程模型长期因GPU架构的特点，复用grid来实现程序中多个thread block的局部性，thread block包含多个thread，这些thread在单个SM上并行运行，其中这些线程通常是什么fast barriers来实现synchronize，以达到通过SM SMEM来交换数据的目标；
  • 而随着GPU SM数量超过100个，计算程序的变得越来越复杂，thread block作为编程模型中的唯一局部性单元已不足最大限度的利用GPU；
  • 因此H100引入了一种全新的thread block cluster架构，该构架以大于单个SM上的单个thread block的粒度来提供局部性控制，该架构扩展了CUDA编程模型，在GPU的物理编程结构上增加了一个级别：thread, thread block, thread block cluster, grid；
  • cluster是一组thread block，这些block可以保证同时调度到一组SM上，其目标是用于实现多个SM之间的高效协作，实现GPC内的SM并发；
  • GPC是硬件层次的一组SM，他们在物理上靠近，clusters具有hardware acceelerator barriers和新的内存访问协作能力，GPC提供了专用了SM-to-SM网络，用于cluster内的thread实现数据传输/共享；

Grid vs Thread block cluster

Distribute shared memory(DSMEM)

  • DSMEM: 在clusters中，threads可以直接访问其它SM的SMEM数据；
  • DSMEM中的虚拟地址空间是在逻辑上跨域cluster上的多个thread block的；
  • DSMEM的存在，cluster上的SM之间的数据交互就不再需要先写入Global Mem再读写SMEM的过程，而是利用专用的SM-to-SM网络完成对远程数据的快速、低延迟访问（与访问Global Mem相比，速度提升了7x）；

SM-to-SM in cluster

cluster性能提升

Asynchronous execution

  • 基于TMA和asynchronous transaction barrier能力，能进一步提升异步计算能力，进一步优化计算与内存copy的重叠；
  • TMA操作是异步的，并是利用A100中引入的基于SMEM的异步屏障能力；
  • TMA编程模型是单线程的，因此只会从一个warp中的选一个线程来执行TMA操作cuda::memcpy_async，多个线程会等待在屏障cuda::barrier上等数据传输完成；

TMA address

  • TMA的另一个优势是可以释放出线程出来执行其它独立的工作。在A100中，异步内存拷贝是利用特殊的指令执行，因此这些线程需要负责生成所有的地址并循环拷贝数据，而在Hopper上，TMA会处理所有的工作：单个线程创建一个copy descriptor，之后的工作由全部由TMA完成；

Asynchronous memory copy with TMA on H100 vs. LDGSTS on A100

Asynchronous barrier

  • 异步屏障是由Ampere架构引入的，用于解决场景：一组线程并行产生数据，这些数据在全部产生后将被继续用于后续的线程计算，这里通过asynchronous barrier （如下左图）分为了两步来实现：
    1. 先生成完了自己部分数据的线程，会发出非阻塞Arrive信号，之后该线程可以继续去执行其它工作；
    2. 那些需要消费上述数据的线程，在执行Wait命令后该线程会被阻塞，直到所有线程都完成了Arrive；
  • 该功能是让那些在等待的线程可以去执行无关的工作，以提升资源利用率；
  • 在Hopper架构中，优化了Wait状态的线程，之前系列GPU中，该状态的线程是处于spin自旋等待状态中，Hopper构架下线程是处于sleep状态；
  • Hopper架构中，asynchronous barrier仍然是重要的编程模型，但同时Hopper引入了一种新的asynchronous transaction barrier（如下右图），其工作模式与asynchronous barrier相似，仍然是分了两步，但除了计数thread arrive外，还需要计数transactions；
  • asynchronous transaction barrier主要是为了解决SM-to-SM SMEM数据交互过程中的线程同步问题，transaction计数本质就是一个byte count。该barrier wait会阻塞线程，直到所有线程完成Arrive，并transaction counts达到目标值；

Asynchronous barrier in A100 vs. asynchronous transaction barrier in H100

Compute Capability 9.0

  • V100 vs A100 vs H100 计算能力对比：

Compute Capability 9.0