0. Motivation

实在不知道博客更点什么内容比较好，干脆以Review的形式更新一些读论文的感受，主要集中在ISCA，ASPLOS，MICRO，SC，SOSP，OSDI，PLDI，HPCA等会议的已发表论文上(所以只是提出个人的看法，本身的价值已经由Accept这个事实体现了)。一切review代表个人观点，欢迎交流。

1. Template

Review Log分为以下板块：

• Summary
• Strengths and Weaknesses
• Question and Comments to Authors
• Paper Writting
• Overall Merit (1-5, Strong Reject (1), Weak Reject, Boardline, Weak Accept, Strong Accept (5))
• Confidence (1-4, No familarity (1), Some familarity (2), Knowledgeable (3), Expert (4))

2. Plan

1. Review Log: QIG: Quantifying the Importance and Interaction of GPGPU Architecture Parameters (TCAD VOL.37, NO.6, 2018)

NVIDIA Hopper 架构GPU浅析

原文地址(2022.03.22)

0. 引言

2022年GTC大会，老黄介绍了基于NVIDIA最新Hopper架构的NVIDIA H100 Tensor Core GPU。本文将带你深入Hopper架构并介绍Hopper架构的一些重要新特性。

1. NVIDIA H100 Tensor Core GPU简介

NVIDIA H100 Tensor Core GPU是NVIDIA发行的第九代数据中心GPU，该GPU被设计用来在大规模人工智能应用以及高性能计算应用中提供相较于前代商品NVIDIA A100 Tensor Core GPU数量级级别的提升。H100继承了A100上为提升人工智能及高性能计算性能的设计特点，且显著提高了架构效率。图1展示了基于新SXM5板卡的NVIDIA H100 Tensor Core GPU。

对于如今主流的人工智能以及高性能计算任务，配备有InfiniBand互联技术的H100相比A100能提供最多30倍的性能提升（译者：可能类似AMD Infinity Fabric，致力于提升设备、节点间的带宽）。最新的NVLink互联系统主要着眼于如今最具挑战性的、需要大量GPU计算节点提供加速进行模型并行的任务。基于InifniBand的加持，这些任务在H100上的性能又一次得到了飞跃。图2展示了A100, H100, H100+InfiniBand在高性能计算、AI训练与推理任务中的性能。

在本次GTC上，NVIDIA发布了新的NVIDIA Grace Hopper Superchip产品。 NVIDIA Hopper H100 Tensor Core GPU将为NVIDIA Grace Hopper Superchip CPU+GPU架构提供动力，该架构专为TB级加速计算而打造，并能在大型人工智能模型以及高性能计算任务重提供超10倍的性能。
NVIDIA Grace Hopper Superchip利用Arm架构的灵活性从头开始设计用于加速计算de CPU和服务器架构。 H100与NVIDIA Grace CPU通过NVIDIA的芯片互联技术连接，提供900 GB/s的总带宽，比PCIe Gen5快7倍。 与当今最快的服务器相比，这种创新设计可提供高达30倍的总带宽，并为使用巨量数据的应用提供超10倍的性能提升。

下面是NVIDIA H100 Tensor Core GPU新特性的简要总结：

• 新设计的流多处理器(SM)提供了许多性能及效率上的提升，具体如下：
  • 第四代张量核心：若将每个SM的提升、SM数量的提升以及频率的提升结合考虑，H100相比A100能提供超6倍的性能提升。在每个SM层面，H100相比A100在每种数据类型上能提供2倍的矩阵乘加(MMA)性能，在FP8数据类型上，相比于FP16数据类型则能提供4倍于A100的性能。而张量核心的稀疏特性能够利用深度学习中的细粒度结构化稀疏，提供2倍于稠密张量计算的性能。
  • 新的DPX指令：该指令能够加速动态规划算法，相比A100提升7倍的性能。例如基因组处理中的Smith-Waterman算法以及为机器人队伍寻找最优路径的Floyd-Warshall算法。
  • 3倍性能的IEEE标准浮点处理性能：相比A100，SM同频浮点性能为2倍。而得益于更高的频率以及更多的SM，总体性能为3倍。
  • 新的抽象层Thread block Cluster：这一特性能够使程序员能够显式地以大于运行在一个SM上的一个线程块Thread block的粒度控制程序的局部性。这一特性通过增加一个抽象层扩展了CUDA编程模型，现在的编程模型包括thread, thread block, thread block cluster, grid。Thread block cluster能够使多个线程块同时运行于多个SM，并支持同步、集合通信以及数据交换。
  • 分布式共享内存：这一机制允许SM之间点对点横跨多个SM的共享内存地进行包含load, store以及源自操作在内的通信。
  • 新的异步执行模型：通过张量存储加速器TMA能够在全剧内存以及共享内存间进行更高效的大规模数据通信。TMA同样还支持thread block cluster内的thread block进行异步数据拷贝。为保证数据原子移动以及同步，TMA还加入了新的异步事物屏障(asynchronous transaction barrier)。
• 新的Transformer引擎：该引擎结合张量核心以及软件特性专门加速Transformer模型的训练与推理。该引擎能够智能地管理计算过程，动态地在FP8以及FP16精度之间选择，自动处理层之间不同数据精度地重用和缩放。在大型语言模型上相比A100能获得9倍训练性能和30倍推理性能地提升。
• HBM3显存子系统：相比前代产品这一系统能提供将近2倍地带宽。H100 SXM5 GPU是世界上第一款配备了HBM3显存并提供领先的3 TB/s显存带宽地GPU。
• 50MB的L2缓存：缓存更多的数据、模型，以减少访存HBM3。
• 第二代多实例GPU技术(Multi-Instance GPU)：相比A100，每个GPU实例能提供3倍的计算能力以及2倍的访存带宽。且本代产品开始每个GPU实例能够提供加密保密计算。H100最多能够提供7个GPU实例，每个实例都有各自的图像视频解码单元(NVDEC,NVJPG)。现在每个GPU实例都有独立的性能监视器。
• 新的隐私计算支持：新机制能够保证在虚拟机或GPU实例中的用户数据具有更高的安全性，防止硬件、软件的攻击，更好地隔离资源。H100实现了世界上第一个原生的GPU配合CPU的加密计算以及高可信运行环境，并且完全匹配PCIe标准。
• 第四代NVLink：相比前代产品，H100在All-Reduce运算中能够提供3倍的带宽，而各种应用平均能提升50%的带宽。NVLink能提供7倍于PCIe Gen5的带宽(900 GB/s)。
• 第三代NVSwitch：这一技术同时包含节点内部以及节点外部用于互联在服务器、集群及数据中心中多个GPU的网络。节点内部每个NVSwitch提供64个NVLink端口以加速多GPU互联速率。相比前代产品，整个互联网络的吞吐量从7.2 Tbits/s提升到13.6 Tbits/s。新一代互联网络同时在硬件上加速了包含多播、归约等集合通信。
• 新的硬件互联系统：基于NVLink以及NVSwitch技术，新的互联系统支持32个节点或256个GPU以2:1胖树拓扑相连接。这一连接方式能够提供57.6 TB/s的互联带宽从而达到在FP8精度下1 exaFLOP的稀疏计算性能。
• PCIe Gen5：相比Gen4，能够提供双向共128 GB/s的带宽。PCIe Gen5使得H100能够以最高的性能与x86 CPU，SmartNICs以及DPU连接

除上文提到的信特性外，还有许多新细节能够提升性能、减小延迟和开销、提升编程便利性。

2. 深入NVIDIA H100 GPU架构

第一部分：飞行器概览

第一章：飞行器及其引擎

1. 飞行器概述

1.1 认识咆哮虫

E/A-18G咆哮虫是由麦克唐纳道格拉斯公司设计制造的舰载电子干扰飞行器（麦道现已属于波音公司的一部分）。飞行器的大致尺寸、三视图、座舱布局分别图1.1与图1.2中以及后续内容中展示。

为E/A-18G装备空中电子对抗系统（Airborn Electronic Attack，AEA）的目的旨在提高强电磁干扰作战环境下的生存能力以及电子攻击能力。在电子对抗系统AEA中包含起到控制器作用的电子攻击单元（Electronic Attack Unit，EAU）、ALQ-218(V)2接收机、ALQ-227B通讯反制套件（Communications Countermeasures Set，CCS）、多任务先进战术终端（Multimission Advanced Tactical Terminal Block 3，MATT）以及ALQ-99电子干扰吊舱。为满足E/A-18G的作战需求，数字存储设备（Digital Memory Device，DMD）以及干扰消除单元（Interference Cancellation Unit，INCANS）可被按需添加。

E/A-18G的作战任务是压制敌对搜索、捕获、跟踪以及制导雷达系统及一切可能对友方单位产生威胁的射频通道。这项任务包括在友方单位进出作战区域以及友方攻击机、战斗机作战时提供保护，这项任务通过电子攻击以及武器投放达成。这项任务的目的是在多种敌方威胁的情况下保护多种友方飞行器。该任务的电子攻击部分通过由任务计算机管控的具备空中电子攻击功能的载具人机界面达成，这一载具人机界面不仅包含经由任务计算机基于空中电子干扰威胁感知跟踪信息处理的威胁态势感知显示信息，还包含经由任务计算机通过显示器、驾驶舱控制按钮、以及手不离杆系统（HOTAS）上的开关状态获取的任务与指令界面。E/A-18G采用电子战战术作为进攻措施，其中攻势干扰是E/A-18G的主要功能。

E/A-18G由两台通用动力F414-GE-400涡轮风扇发动机提供动力。该型发动机配备有数字引擎控制（Full Authority Digital Engine Control，FADEC）。E/A-18G具有机翼前缘延伸（leading edge extensions，LEX，涡流发生器）以及可变中翼。E/A-18G的垂直尾翼有向外20度的角度。

在设计该型飞行器时放宽了静稳定性的限制已获得更好的操纵性以及更低的进近、着陆速度。E/A-18G通过数字电传飞控系统以液压的方式控制舵面（副翼、机翼双舵面、前缘襟翼、后缘襟翼、前缘延伸扰流板以及差速稳定器）。前缘襟翼包含一个能够通过增加机翼有效面积以增加进近、着陆时操控性能的机构。机翼上表面采用折叠机翼整流罩以及翼刀以增加大攻角跨音速段的操控性。飞行器的减速主要由各舵面的偏转负责。

加压座舱由电动座舱盖密封，辅助动力装置（Auxiliary Power Unit，APU）为引擎提供了自发启动时需要的电源。

1.2 飞行器质量

飞行器基础重量（包含除ALQ-99吊舱以外的空中电子攻击设备）为33733磅。请在DD-365文件中获取更为精确的飞行器质量。

第二章：飞行器系统

第三章：飞行任务及情况处置

第四章：飞行行动限制

0. 开始之前及总体感想

• 文章标题：Review Log: QIG: Quantifying the Importance and Interaction of GPGPU Architecture Parameters

• 期刊/会议：TCAD VOL.37, NO.6, 2018

• 这篇文章和我下一个想做的工作有重合，所以会以较为aggressive的态度来写review log（

• 本文主要介绍了一种用于GPU设计过程中设计空间探索（Design Space Explore，DSE）的系统，通过将设计参数与GPU Kernel的运行性能数据输入给一个可解释的机器学习模型（有点类似随机森林的决策树算法）来衡量不同任务下不同设计参数的优劣。当然这篇文章最重要的贡献我认为不在于模型准确率活着加速GPU架构设计等等，而在于其挖掘的GPU设计参数对于计算任务的影响，以及参数之间互相影响的模式。

• 本文采用可解释的机器学习模型，在获得高预测准确率的同时能够解释数据与硬件之间的关系。然而有些解释在我看来并不能真实反应硬件设计参数对性能的影响，也许是将相关性解释为了因果性，详见Review Log正文。

  1. Summary

  This work proposed an explainable performance model of GPGPUs, based on which the authors first predict the performance of a given kernel on giver hardware with high accuracy, then explain the relation between GPGPU behaviors and design parameters.

  2. Strength

• The paper is well organized and well-motivated.
• The performance model they use is highly explainable with considerable accuracy.
• Some important design insights are revealed via their model.
• The performance model can be easily generalized to other hardware and other performance metrics (from IPS to Power consumption).

  3. Weakness

• Some findings are misleading, which will be datailed in questions and comments.

  4. Questions and Comments to Authors

0. The problem this work target is emerging and important, and the authors explain their motivation well.
1. In the analysis of importance analysis, the core frequency is treated as the most important design factor. While this parameter is not the most challenging factor the hardware designers concerned. It’s trival that with higher frequency the performance and power of a GPGPU kernel will increase. So is the instruction replay overhead which is a factor related to the profiler itself. I think it’s necessary to filter out the influence of these parameters other wise some misleading findings will be given. (For example, execute all the kernels at a fixed core frequency or fixed power budget)
2. I think the benchmark amount and the problem size are limited, and the register per thread is always sufficient. While for some recent applications with very heavy computation burdens, there may exist other patterns. (这点其实有点时空警察，但是2018年应该也是有那些很大的模型、很大的计算任务了)
3. GPUWattch is available for power estimated with higher explainable.
4. Overall, the findings and hardware patterns proposed in this work is reasonable and insightful.

   5. Paper Writting (1-5)

   4

   6. Overall Merit (1-5)

   4 - Weak Accept

   7. Confidence (1-4)

   3 - Knowledgeable

This work is accepted by ASPLOS 2022.

Introduction

It’s universally acknowledged that deep learning applications are evolving fast and are becoming more and more complex. For example, In cloud service and auto-piloting, we need to provide service to several deep learning models. However, currently, the computation power of hardware also increases significantly, for example, Threadripper 3990X CPU has ~15 TFLOPS for FP32 while Radeon Instinct MI100 has>100 TFLOPS for matrix computation. This situation resulted in a problem that the performance of single hardware is far beyond the requirement of a single inference task of model: the QoS requirement of ResNet-50 in MLPerf is 15ms while the whole CPU use only 6ms to finish the job.
Under this circumstance, we need multi-tenant. In another word, we need to co-locate multiple tasks, or queries, onto one hardware.

Current Solutions

We can divide current solutions into hardware and software solutions. The hardware solutions mainly bypass the complex resource management of co-locating tasks by temporal sharing or physical isolation, like AI-MT for temporal sharing and NVIDIA MIG for resource isolation. For software solutions, there are some task co-locating solutions for traditional workloads including Parties, for deep learning tasks, DART proposed a cluster scale scheduling strategy. While our work mainly focuses on CPU scale strategy.

A straightforward solution

Of course, we have some simple, straightforward solutions. On CPUs, we can use MPI to serve multiple tasks and use CPU affinity and tools including res control, Intel CAT to limit resource usage. On GPUs, we have Multi-Process Service, Persistent Thread Block. These techniques support simply dumping multiple tasks onto single hardware. But how about the performance? It’s not good. Following list is an example of doing so.

1
2
3
4
5
6
7
8
9
10
11
12
13
14

if (MPI_CPU_ID == 0) {
    // CPUs
    sched_setaffinity(0, sizeof(cpu_set_t), 0x000f);
    task1.run();
    // GPUs
    task_kernel_1 <<<(1/4)_SM, ...>>>(...);
}
else if (MPI_CPU_ID == 1) {
    // CPUs
    sched_setaffinity(0, sizeof(cpu_set_t), 0xfff0);
    task2.run();
    // GPUs
    task_kernel_2 <<<(3/4)_SM, ...>>>(...);
}

What’s the limitations?

• The first is co-locating interference, according to our verification experiment, co-locating 4 identical GoogLeNet tasks on a CPU platform have 1.8x latency. If we inspect the system performance counter including LLC miss rate, LLC access, we observe server shared resources contention.
• The second is low QoS satisfaction rate. With the increase of QPS (Query Per Second), the satisfaction rate of the system drop significantly. Even if we apply fine-grained scheduling (schedule every layer/op of the network seperatly), the system with AMD Threadripper 3990X can only serve 50 queries with 99% satisfaction rate, while the theoritical performance of the CPU is far beyond this.