论文阅读：Benchmarking and Dissecting the Nvidia Hopper GPU Architecture

论文阅读：Benchmarking and Dissecting the Nvidia Hopper GPU Architecture

原文链接：https://arxiv.org/pdf/2402.13499.pdf

Notes

• Firstly, we conduct conventional latency and throughput comparison benchmarks across the three most recent GPU architectures, namely Hopper, Ada, and Ampere.
• 我们深入研究了最新 Hopper 功能的全面讨论和基准测试，包括 Hopper DPX 动态编程 (DP) 指令集、分布式共享内存以及 FP8 张量核心的可用性
• 张量核心性能和编程指令集
  • the tensor core performance and programming instruction sets
• tensor cores and high-bandwidth memory
• 张量核心（TC）单元最初随 Volta 架构引入，专注于以 FP16 和 FP32 精度运算加速深度神经网络
• 后续的 Ampere 架构扩展了 TC 功能，包括稀疏性和更广泛的数据精度，例如 INT8、INT4、FP64、BF16 和 TF32
• Hopper 架构进一步扩展了这一点，引入了对 FP8 精度的支持，显着增强了 LLM 训练和推理加速。
• Hopper introduces innovative features:
  • Dynamic Programming X (DPX) instructions：DPX 指令可加速各种动态编程算法，通常涉及大量最小/最大运算来比较先前计算的解决方案
  • distributed shared memory (DSM)：DSM 支持 SM 到 SM 的直接通信，包括跨多个 SM 共享内存块的加载、存储和原子操作
  • an enhanced asynchronous execution mechanism (Tensor Memory Accelerator) for diverse scenarios：Hopper支持集群内线程块之间的异步复制，提高效率

1 本文贡献

• We conduct detailed instruction-level testing and analysis on memory architecture and tensor cores across three GPU generations with different architectures.
  • Our analysis highlights the unique advantages and potential of the Hopper architecture
• We compare AI performance across recent GPU generations, examining latency and throughput of tensor cores at the instruction level, transformer engines at the library level, and real LLM generation at the application level
• 我们的研究代表了对 Hopper 架构独特功能的首次探索，包括 DPX、异步内存操作和分布式共享内存。

2 相关工作

• 融合矩阵乘法累加 (MMA) 是人工智能中的一项关键运算，自 Volta 架构以来主要由 Nvidia GPU 中的张量核心 (TC) 加速
• 
• • 应用方面：
    • GPU 动态电压和频率调节 (DVFS) 对深度学习期间能耗和性能的影响测试各种 GPU 架构，DVFS 设置，和 DNN 配置。
• 必要性：
  • 尽管张量核和AI很流行，但另一个趋势是DPX的支持、异步操作支持和分布式共享内存。

3 METHODOLOGY

3.1 内存单元访问:延迟和吞吐量

• 如何测试这两个指标：

  • L1 Cache：
    • 延迟:我们首先使用 ca 修饰符将数据从全局内存加载到 L1 缓存。然后我们使用一个线程来访问这个L1缓存以获得延迟。
    • 吞吐量：首先使用 ca 修饰符将内存加载到 L1 缓存中。我们只发出一个具有 1024 个线程的块来重复访问 L1 缓存。我们记录消耗的时间和访问的数据量来计算L1缓存的带宽。
  • Shared memory：
    • Testing the shared memory is similar to testing the L1 cache
    • 唯一的区别是不需要指定修饰符来显式预热共享内存。
    • 我们使用一个线程来测试延迟，并使用一个具有 1024 个线程的块来测试带宽，就像测试 L1 缓存一样。
  • L2 Cache：
    • For the latency test, we use the same method as for L1 cache testing.
    • The only difference is that the cg modifier is used instead of ca, ensuring that the cache we load is L2
    • 测试吞吐量：我们首先使用 cg 修饰符将内存加载到 L2 缓存中。**由于L2缓存被所有SM共享，**我们使用大量的块来访问二级缓存。然后我们根据访问的数据量和消耗的时间来计算二级缓存的带宽

  Streaming Multiprocessors (SM)：这通常是指在图形处理单元 (GPU) 中的一种结构单元。SM 是 GPU 上用于并行计算的核心部件，它们能够执行程序中的指令。不同型号的 GPU 有不同数量的 SM。

  • Global memory：
    • 我们首先分配超过L2大小的全局内存，以避免L2预取，然后初始化全局内存
    • 初始化有两个目的。第一个是使测试能够以固定的步幅进行，第二个是预热TLB以避免冷遗漏的发生。
    • 我们连续启动了四个线程，每个线程负责读取8个字节，从而形成一个32字节的内存读取事务。
    • Finally, we can calculate the memory access latency of each thread.
    • 为了吞吐量测试，分配了比L2大得多的内存空间，将每个线程设置为使用向量化内存访问来读取 float4。每个线程读5次写1次，最后，根据消耗的时间和数据量计算内存带宽。

  TLB是Translation Lookaside Buffer的缩写，它是计算机体系结构中的一个重要部分**，用于存储虚拟地址到物理地址的转换。**TLB的预测是指对TLB中将要发生的转换进行预测，以提高地址转换的效率。

3.2 张量核心延迟和吞吐量

• Tensor Core’s Evolution：

  • 从开始只支持FP16（Volta Architecture）作为输入数据类型到逐渐支持更多的精度（BF16, TF32, FP64, INT8, INT4, Binary, and more）

  安培和艾达·洛夫莱斯两种GPU的特性。它们为用户提供了选择的灵活性，可以使用传统的C级别wmma（CUDA Warp Matrix Multiply）API，这是一种用于执行矩阵乘法运算的CUDA库函数，或者使用PTX级别的***mma（Matrix Multiply Accumulate）***指令。PTX（Parallel Thread Execution）是CUDA中的一种汇编语言，mma指令允许开发者更加底层地控制GPU上的矩阵乘法运算。因此，这句话强调了GPU用户可以根据需求和偏好选择合适的方法来执行矩阵乘法运算

  • wmma API 在充分利用 TC 的功能方面存在局限性，而 mma 指令可以利用自 Ampere 以来引入的高级稀疏矩阵乘法功能
  • Hopper 仍然支持 wmma 和 mma API
  • 同时Hopper TC 的全部潜力只能通过 wgmma 指令来实现。
  • mma指令计算D(m×n) = A(m×k)×B(k×n)+D(m×n)
    • mma计算时通过一个CUDA warp（例如32个threads）同步执行
    • wgmma计算时通过一个CUDA warp（例如4个CUDA warp）异步执行
    • The matrix shapes for mma instructions can be m16n8k16 or m16n8k8，while wgmma supports m64nN k16 where N can be 16, 32, 64, 128, 256
  • wgmma 的优点：直接从共享内存加载矩阵 A 和 B，mma 需要在执行之前将所有矩阵存储在寄存器文件中。此外，wgmma 还支持 mma 不需要的某些有用参数。
    

• Benchmarking Levels and Performance Metrics:

  • *strikes a suitable balance between granularity and complexity.*在粒度和复杂性之间取一个平衡
  • we disassemble PTX instructions to SASS codes to achieve a deeper understanding of the operations拆分PTX指令成SASS代码
  • 两个指标：延迟和吞吐量：
    • ·延迟表示所用的时间，以时钟周期为单位，从开始向执行管道发出指令开始，到结果可供后续使用时结束
    • 吞吐量被量化为Total OPS/Duration，其中OPS表示乘法或加法运算。利用总时钟周期来计算由于执行不同 TC 指令期间 GPU 频率的潜在变化而导致的吞吐量。

3.3 Transformer Engine

• Hopper architecture.它能够利用 Hopper 和 Ada 架构提供的 FP8 精度
• Linear Layer:
  • 在 Transformer 架构中，大部分计算开销来自线性层，特别是矩阵乘法，而 Transformer 引擎提供 te.Linear 实现，可以在 FP8 Tensor Core 上以更高的吞吐量执行矩阵乘法。
  • 当使用带有 te.Linear 的 Transformer Engine 进行 FP8 精度的矩阵乘法时，TE 将线性层中的输入和权重转换为 FP8。这个转换过程涉及数据变换和量化操作。
    • 由于 FP8 的动态范围可能不包含输入张量的最大值，因此 TE 将输入数据的最大绝对值识别为缩放因子（scaling factor）。
    • 然后使用 inp_fp8 = inp_fp16/scale 调整输入数据以适应 FP8 的表示范围，然后在 FP8 Tensor Core 中进行矩阵乘法 out_fp8 = inp_fp8 × w_fp8
    • 它使用 out_fp16 = out_fp8×scale 缩放结果。此操作会带来一些开销。
  • 当执行相对较小的矩阵乘法的时候转换明显大于 FP8 Tensor Core 中 GEMM 内核计算产生的转换结果。
  • 评估：
    • 测量两个相同矩阵 D(n × n) = A(n × n) × B(n × n) 的 te.Linear 的吞吐量 (GFLOPS)。
      
• TransformerLayer：
  • Transformer Engine (TE) 通过对Transformer层结构进行特定算子融合优化**，利用 FP8 提供的效率改进**
    • 例如，LayerNorm MLP将layernorm和MLP组合在transformer结构内，允许layernorm和后续MLP层之间的数据传输采用FP8格式
      • 这种方法不仅消除了数据格式转换开销，而且有效地利用了 FP8 内存传输优势
  • TE 提供 te.TransformerLayer 模块，其中包含变压器层结构的所有算子优化，通过调整其参数促进各种大型语言模型 (LLM) 结构的实现
    • some operators, such as Softmax and GeLU, have not been quantized to FP8 by TE,resulting in significant data format conversion overhead
    • 此外，DotProductAttention 运算符使用 flash-attention而不是 FP8 Tensor Core。
  • Transformer层的线性层的计算开销主要取决于隐藏层的大小:
    • 提出了一个问题：与 FP16 相比，哪种隐藏状态（嵌入维度）将为 FP8 的 TE 带来更好的性能?
      • 通过检查开源 LLM、Llama、将激活函数修改为 SwiGLU并将标准化为 RMSNorm来研究这一点
      • 根据隐藏状态的大小设置层结构参数，隐藏状态4096、5120和8192分别对应于Llama配置7b、13b和70b。
        
• 将输入固定为 (4, 512, hidden_size)，其中 4 是batch_size,512 是序列长度，注意力掩码设置为 None。然后计算了对单层进行一次编码所需的延迟（毫秒），重点关注单层的编码任务。
• LLM Generation：
  • 目前Transformer Engine尚未对主流decode-only任意语言模型提供最优支持
  • 为了测试推理性能：

    • 分别将原模型结构中的nn.Linear和RMSNorm替换为te.Linear和te.RMSNorm，以确保模型中的大部分模块都使用Transformer Engine

    • 为了评估 TE 在为 Llama 生成文本方面的有效性，使用 ShareGPT 数据集作为 LLM 的输入。

      • ShareGPT 数据集包含用户和 ChatGPT 之间的对话，已被用户分享。对这些数据集进行标记，并根据它们的输入和输出长度生成合成的客户端请求。

    • 使用吞吐量作为评估指标，它代表每秒可以处理的文本总长度：

      Throughput = (input_len + output_len)/time

3.4 New CUDA Programming Features

• DPX：
  • Nvidia 从 CUDA 12 开始提供 DPX 功能，以加速动态编程代码，提高编程的简易性。
    • 对于延迟评估，利用线程迭代发出 DPX 函数，计算他们的平均延迟。
    • 在吞吐量测试中，使用一个块重复发出DPX函数，确定每个SM的DPX指令吞吐量
    • 改变启动块的数量，并观察 DPX 吞吐量和启动块计数之间的关系。
• Asynchronous Data Movement:
  • 异步执行的引入是Ampere架构的一大亮点
  • 此功能允许使用 cuda::memcpy async 在 GPU 全局内存和共享内存之间进行非阻塞数据传输，避免数据移动期间的线程占用。它促进计算与数据传输的重叠，有效减少整体执行时间
  • Hopper 架构以Ampere的异步复制为基础，通过更先进的**张量内存加速器 (TMA)** 增强了这一点，以实现复杂的异步复制
    • 为了评估此功能的效率，使用来自官方 CUDA 样本的 globalToShmemAsyncCopy 应用程序进行实证研究
    • 该应用程序实现矩阵乘法，并利用从全局到共享内存的异步数据复制来实现 8.0 或更高的计算能力。
    • 比较了两种实现：“SyncShare”，采用同步复制到共享内存来进行传统的平铺tiled矩阵乘法；以及“AsyncPipe”，通过异步数据移动增强平铺。异步版本使用两级管道，共享内存缓冲区大小加倍，使计算和数据复制能够在不同的执行流之间重叠。
      • 矩阵A的宽度和矩阵B的高度设置为2048，决定了每个线程的计算量。
      • 将块大小从 8×8 更改为 32×32 来评估异步操作对 warp 并发的影响
      • 对不同的块数进行基准测试，通过调整矩阵 A 的高度和矩阵 B 的宽度来优化计算吞吐量。
• Distributed Shared Memory：
  • Hopper 架构的特点是集群内有一个直接的 SM 到 SM 通信网络，使一个线程块中的线程能够访问另一个块的共享内存，称为分布式共享内存 (DSM)
  • 根据官方文档，该网络可以将不同SM上的块之间的数据传输开销减少多达7倍
  • 此外，对于共享内存需求限制 SM 上的活动块数量的情况，DSM 可以对同一集群内的数据进行分区，从而减轻每个块的共享内存需求。
  • DSM 的可编程性通过 CUDA C 函数 cluster.map_shared_rank(SMEM, DST_BLOCK_RANK) 实现，返回目标块的共享内存地址
  • SMEM代表共享内存指针，DST_BLOCK_RANK是簇中的目标块等级。这被编译成PTX代码mapa，它映射目标块中共享变量的地址
  • We assess DSM using three benchmarks：
    • Latency Measurement
    • RBC（基于环的复制）吞吐量测量
    • Histogram Application with DSM

4 EXPERIMENTAL RESULTS

HBM2e HBM2e代表"High Bandwidth Memory 2e"，它是一种高带宽内存技术。HBM2e是HBM2（High Bandwidth Memory 2）的改进版本，旨在提供更高的内存带宽和更大的容量，以满足高性能计算和图形处理等领域的需求。

**HBM2e与传统的*GDDR（Graphics Double Data Rate）***内存相比，具有更高的数据传输速率和更低的功耗。这使得HBM2e成为高性能计算、人工智能、深度学习和图形处理等应用领域的理想选择。通过将存储器堆叠在处理器芯片上方，HBM2e技术可以显著提高内存带

GDDR6X是一种图形双数据速率（GDDR）存储技术，它于2020年推出。这种技术是由英特尔旗下的NVIDIA公司开发的，用于其最新一代图形处理器（GPU）。GDDR6X相比传统的GDDR6存储技术有更高的带宽和更快的数据传输速度，这使得GPU能够更有效地处理大规模的图形数据和复杂的计算任务。GDDR6X的推出为游戏玩家、图形设计师和科学计算等领域的用户提供了更好的性能和体验。

• throughput:
• Tensor Core Latencies and Throughputs
• SASS analysis
• • mma指令被编译(compilation)成SASS指令，命名约定遵循既定的模式：
    • HMMA（针对浮点类型）、IMMA（针对整数类型）和 BMMA（针对二进制类型）。值得注意的是 mma 中存在两种特殊类型：INT4 和 FP8
    • FP8（Ada 中引入的一种新数据类型）没有可用的 mma 指令。
• Unlike mma, wgmma instructions are compiled into the new GMMA SASS instructions.
  • wgmma 不提供对 INT4 Tensor Core 的支持
• mma results:
• • 表中的稀疏形状代表压缩形状。换句话说，实际指令修饰符的k是表中的两倍。
• 对于 A100 和 H800，具有较大形状的相同精度 mma 指令通常可以实现更好的吞吐量。
• 但这种现象在RTX4090上就消失了。稀疏和密集 mma 指令表现出相同的延迟，稀疏 mma 指令可实现更高的吞吐量。
• 稀疏和密集 mma 指令表现出相同的延迟，使用稀疏 mma 指令实现更高的吞吐量。
• wgmma results:
  • 作为一组专为 Hopper GPU 设计的扭曲组级 Tensor Core 指令，wgmma 指令是异步执行的先驱指令。
  • 在 H800PCIe GPU 上使用 Tensor Core 的用户在执行计算时应充分考虑功耗限制
  • 在“RS”和“SS”模式下，同一指令的延迟和吞吐量保持相对一致。
    • 原因：由于大量的计算工作量和进程的异步特性而有效隐藏了共享内存访问延迟。
      • 发现：在稀疏wgmma中，“SS”模式从大小为m×k的共享内存中检索数据，并在执行稀疏wgmma指令期间基于元数据执行2:4稀疏剪枝
      • “RS”模式直接从大小为 m × k/2 的修剪寄存器文件中访问数据
      • 从共享内存中获取指令的延迟导致“SS”模式下的稀疏 wgmma 指令无法达到预期的峰值性能
• wgmma results with different N values:
  • ***wgmma.m64nNk16.f32.f16.f16,***建议尽可能选择较大的 N 值 (>= 64)，以获得卓越的性能。
• Energy efficiency：
  • 可以发现H800的能效明显更高
  • • Transformer Engine Performance
• Te.Linear analysis：
  • 使用 FP8 张量核心加速线性层的矩阵乘法。
  • FP8 性能受到数据格式转换和量化运算符开销的影响
  • 这强调了 FP8 的高吞吐量潜力，但也强调需要特定条件才能获得最佳计算密度。
• Te.TransformerLayer analysis：
  • Transformer Engine 将整个 Transformer Layer 结构压缩为 te.TransformerLayer。
  • 随着计算密度的增加，H800的计算优势变得明显
• LLM Inference Throughput Results：
  • decode-only模型在推理过程中受内存限制，因此 FP8 Tensor Core 的计算优势并不显着
  • • 有可能当模型规模和输入数据长度增加时，并且有良好的算子融合支持，可以取得一定的提升

New Features of Hopper

• DPX:
  • 可以观察到的是，对于relu指令，H800的性能明显更好
  • 对于16位运算，H800也有显着的加速，高达13倍。
  • • 事实上，通过观察SASS代码，发现Hopper上使用了新指令（VIMNMX）。
• 与之前的IMNMX相比，性能似乎并没有明显提升,但总的来说，带有DPX硬件加速的Hopper架构比上一代架构具有更好的性能。
• 此外，__vibmax_s32 数据在 RTX4090 和 A100 上不可用。
  • 原因是编译优化将此函数优化为max指令。如果想阻止这种优化，吞吐量测量将受到很大影响。
• DPX 函数的吞吐量与块的数量成正比。
• 当区块数量刚好超过SM数量的整数倍时，吞吐量急剧下降，随着区块数量的增加逐渐恢复到最大水平。
• 当块的数量是 SM 数量的整数倍时，出现最大吞吐量。因此，有足够的理由推断DPX加速单元位于SM级别。
• Asynchronous Data Movement:
  • 在两个 GPU 上，AsyncPipe 通常优于具有较小块大小（例如 8×8 和 16×16）的 SyncShare。
  • 原因是在小块大小下，warp数量不足会阻碍隐藏同步共享内存复制延迟。
  • AsyncPipe 中的两级管道允许跨不同阶段同时进行数据移动和计算。
  • 随着块大小的增加，优势逐渐减弱即使 H800 上的块大小为 32×32，AsyncPipe 的吞吐量也常常比 SyncShare 差。较大的块大小会导致较高的warp并发性，从而有效隐藏共享内存复制延迟。
  • • 更大的块大小和更多的可并行指令会带来更高的吞吐量
• 随着集群中越来越多的块竞争 SM 到 SM 的带宽，整体吞吐量会变得越来越低。
• Balancing this tradeoff by selecting optimal block and cluster sizes is an important direction for exploration.
• 最佳簇大小对于不同的块大小是不同的（对于块大小128，CS=4，对于块大小512，CS=2）。增加块和簇大小可能会使 SM 到 SM 网络利用率饱和，从而可能因资源争用而降低整体性能。

5 CONCLUSION

• 本文使用指令级基准测试深入研究了最新三种 Nvidia GPU 架构的内存层次结构和张量核心性能。发现 hopper 架构在内存带宽和张量核心方面都表现出优势，这与官方说法一致。值得注意的是，在张量核心上，需要使用最新的wgmma指令来充分利用第四代张量核心的所有性能。
• 实验表明，当运算规模较大时，低精度数据类型会表现出更大的优势。此外，还探讨了 Hopper 架构的主要功能：DPX、异步数据移动和分布式共享内存。研究增强了对最新架构特征和性能的理解，有助于优化算法设计和应用程序性能。