去年，NVIDIA (NV) 发布了Amper新架构的GPU，NVIDIA GPU架构的发展类似Intel的CPU，针对不同场景和技术革新，经历了不同架构的演进。
内容转自：智能计算芯世界
《NVIDIA A100 Tensor Core GPU技术白皮书》《NVIDIA Kepler GK110-GK210架构白皮书》《NVIDIA Kepler GK110-GK210架构白皮书》《NVIDIA Kepler GK110架构白皮书》《NVIDIA Tesla V100 GPU架构白皮书》- Kepler架构里，FP64单元和FP32单元的比例是1:3或者1:24；K80。
- Maxwell架构里，这个比例下降到了只有1:32；型号M10/M40。
- Pascal架构里，这个比例又提高到了1:2(P100)但低端型号里仍然保持为1:32，型号Tesla P40、GTX 1080TI/Titan XP、Quadro GP100/P6000/P5000
- Votal架构里，FP64单元和FP32单元的比例是1:2；型号有Tesla V100、GeForceTiTan V、Quadro GV100专业卡。
- Turing架构里，一个SM中拥有64个半精度，64个单精度，8个Tensor core，1个RT core。
- Ampere架构的设计突破，在8代GPU架构中提供了该公司迄今为止最大的性能飞跃，统一了AI培训和推理，并将性能提高了20倍。A100是通用的工作负载加速器，还用于数据分析，科学计算和云图形。
从Kelper、Maxwell、Pascal、Volta，GPU架构的更新体现在SM、TPC的增加，最终体现在GPU浮点计算能力的提升。NVIDIA GPU架构相关文章：
相比Pascal架构，Volta GPU架构的显著特征是它的Tensor Core，新的Tensor Core是专门为深度学习设计的，有助于提高训练神经网络所需的性能。Tensor Core推动卷积和矩阵运算。集成了Tensor Core也是NVIDIA新一代GPU架构Turing(图灵)的一个重要特性。关于GPU架构，NV 很有意思的是会用一些历史上杰出的科学家的名字来命名自己的硬件架构。总体上，NV GPU 用到的 SIMT 基本编程模型都是一致的，每一代相对前代基本都会在 SM 数量、SM 内部各个处理单元的流水线结构等等方面有一些升级和改动。这篇暂时不涉及到渲染管线相关的部分，其他诸如多少 nm 工艺、内存频率提升等等也都先略过，只关注计算相关的硬件架构演进。关于初代 GPU 的架构，Tesla可以找到的资料不太多，基本上都是从 Fermi 开始的。1、Fermi
Compute Capability：2.0，2.1
Fermi架构SM- 2 个 Warp Scheduler/Dispatch Unit
- 32 个 CUDA Core（分在两条 lane 上，每条分别是 16 个）
- 每个 CUDA Core 里面是 1 个单精浮点单元（FPU）和 1 个整数单元（ALU），可以直接做 FMA 的乘累加
我的理解是做一个双精 FMA 需要用到两个 CUDA Core？所以是 32 / 2 = 16
2、Kepler
Compute Capability：3.0,，3.2，3.5， 3.7；这一代 SM 整体结构上跟之前是一致的，只不过升级完了以后又往里面塞进去了更多的运算单元，其他部分也没有做太大的改动。- 4 个 Warp Scheduler，8 个 Dispatch Unit
- CUDA Core 增加到 192 个（4 * 3 * 16，每条 lane 上还是 16 个）
- 单独分出来 64 个（4 * 16，每条 lane 上 16 个）双精运算单元。
- SFU 和 LD/ST Unit 分别也都增加到 32 个
Kepler 是附近几代在硬件上直接有双精运算单元的架构，不用通过单精单元去做双精运算了，所以对比前后几代的双精浮点的性能话会发现 Kepler 要高出一截。3、Maxwell
Compute Capability：5.0， 5.2， 5.3可能是觉得 Kepler 往一个 SM 里面塞了太多东西，其实最终效率也并没有那么高，这一代的 SM 开始做减法了，每个 SM（SMM）中包含：- 4 个 Warp Scheduler，8 个 Dispatch Unit
- 32 个 SFU 和 LD/ST Unit（4 * 8）
Kepler 里面 192 这个数字也被诟病了（不是 2 的倍数）。这些硬件单元的流水线分布也不再是像 Kepler 那样大锅炖了，而是有点像是把 4 个差不多像是 Fermi 的 SM 拼在一起组成一个 SM：
- 1 个 Warp Scheduler 和 2 个 Dispatch Unit
图上没有看到之前 lane 的标记，不过我猜应该也还是 4 条，两条 CUDA Core 的 lane，1条 SFU，1条 LD/ST Unit。应该是工艺和频率的提升，Maxwell 每个 CUDA Core 的性能相比 Kepler 提升了 1.4 倍，每瓦性能提升了 2 倍。对 CUDA Core 的详细结构没有再介绍，姑且认为从 Fermi 开始一直到以后 CUDA Core 内部的结构都没有什么改变。另外一点是，前面说到的双精单元在这一代上也移除了。也许是觉得认为只有少数 HPC 科学计算才用的上的双精单元在这代上不太有必要吧。
4、Pascal
Compute Capability：6.0, 6.1, 6.2；这一代可以说是有了质的飞跃，还是先从 SM 开始：
可以看到一个 SM 内的部分作了进一步的精简，整体思路是 SM 内部包含的东西越来越少，但是总体的片上 SM 数量每一代都在不断增加，每个 SM 中包含：- 2 个 Warp Scheduler，4 个 Dispatch Unit
- 16 个 SFU 和 LD/ST Unit（2 * 8）
一个 SM 里面包含的 Process Block 数量减少到了 2 个，每个 Process Block 内部的结构倒是 Maxwell 差不多：- 1 个 Warp Scheduler 和 2 个 Dispatch Unit
单个 Process Block 的流水线增加到 6 条 lane 了？
- 面向 Deep Learning 做了一些专门的定制（CuDNN 等等）
- 除了 PCIE 以外，P100 还有 NVLink 版，单机卡间通信带宽逆天了，多机之间也能通过 Infiniband 进一步扩展 NVLink（GPUDirect）
然后 NV 现在已经把 Infiniband 行业的龙头 Mellanox 给收购了…… 说不定那时候就已经有这个想法了呢
- P100 上把 GDDR5 换成了 HBM2，Global Memory 的带宽涨了一个数量级
- 16nm FinFET 工艺，性能提升一大截，功耗还能控制住不怎么增加
- Unified Memory，支持把 GPU 的显存和 CPU 的内存统一到一个相同的地址空间，驱动层自己会做好 DtoH 和 HtoD 的内存拷贝，编程模型上更加友好了
CUDA Core 在这一代也终于有了升级，现在硬件上直接支持 FP16 的半精计算了，半精性能是单精的 2 倍，猜测应该是一个单精单元用来算两个半精的计算。5、Volta
Compute Capability：7.0， 7.2；又一个针对深度学习的质变 Feature，Tensor Core！看到 SM 的时候我们会发现这一代除了多出了一个额外的 Tensor Core 的单元以外，怎么 SM 的体积看起来好像又加回去了，每个 SM 中包含：- 4 个 Warp Scheduler，4 个 Dispatch Unit（发现不需要配 2 个 Dispatch 给每个 Scheduler 了？白皮书里面倒是没有对这个的解释）
事实上相比 Pascal 而言，单个 SM 中的单精运算单元数量是一致的，相当于把 Pascal 中的每个 Process Block 进一步地又拆成了 2 个，每个 Process Block 中包含：- 1 个 Warp Scheduler，1 个 Dispatch Unit
这里把原本的 CUDA Core 给拆开了，FP32 和 INT32 的两组运算单元现在是独立出现在流水线 lane 里面了，这一设计的好处是在前几代架构中 CUDA Core 同时只能处理一种类型的运算，而现在每个 cycle 都可以同时有 FP32 和 INT32 的指令在一起跑了。Pascal 中需要 6 个 cycles 来做一组 FMA，现在在 Volta 中只需要 4 个 cycles。另外每个 Warp Scheduler 还有了自己的 L0 指令 cache。这一代还改进了一下MPS，现在从硬件上直接支持对资源的隔离，方便多任务共享 GPU。Tensor Core：最重大的改动不用说也知道是 Tensor Core 了。
Tensor Core 的思路从系统设计上还是相当直接的，目前深度学习的 workload 中最主要的计算量都在矩阵的乘加上，因此为了专门去高效地支持这些 workload，就增加一些专用于矩阵运算的专用部件进去。这个也是常见的 AI ASIC（比如 Google 的 TPU、其他厂商的各种 xPU 等等）通常采用的思路，只不过 ASIC 可以从一开始就是针对特定的 workload 去的，因此设计上可以更直接更激进一些，直接上大量的 MMU（Matrix Multiply Unit），然后采用例如脉冲阵列这种设计去最大化它的 throughput。而 NV 的 GPU 毕竟还要用作其他一些通用的运算，所以只能往原本的 SM 流水线里面插进去一些额外的专用部件 lane 了。开个脑洞，要是哪一天发现除了 FMA 以外还有其他另外一种形式的运算有大量的需求，未来的 GPU 设计里面说不定也会出现其他 Core。好在 FMA 除了深度学习以外在 HPC 的 workload 里面也是挺常见的，这个设计以后还是比较有用的。Tensor Core 4x4 Matrix Multiply and AccumulateMixed Precision Multiply and Accumulate in Tensor CoreTensor Core 这个部件直接从 SM 的寄存器里面取两个 FP16 的矩阵作为输入，进行全精度的矩阵乘之后得到的结果可以是 FP16 或者 FP32 的，然后累加到 FP16/FP32 的 accumulator 里面去。数据类型选择 FP16 作为输入然后输出 FP32 猜测可能是为了保证结果不溢出，然后在加速部件设计等等方面做了一些 trade off。
所以 FP16 in -> FP16 out 和 FP16 in -> FP32 out 哪一个性能更好呢…
我没有测过，但是猜测可能默认结果是 FP32 out 更快？反而是输出 FP16 需要从 FP32 再转一次？
接下来道理我们都懂了，那 Tensor Core 要怎么用呢？这个部件的编程模型在一开始接触的时候可能会有一些坑。我们知道常规的 CUDA 代码需要制定 grid 的结构、block 的结构，然后其实我们写的 kernel 代码都是针对每一个单独的 thread 的，可以认为是 thread level 的编程。对一个子矩阵的 FMA 运算存在比较多的数据重用机会，这时候如果只是一个 thread 算一个矩阵块的 FMA 就比较浪费了，因此 Tensor Core 的设计是用一整个 warp 去共同完成一个 FMA 运算，一个 warp 中的 32 个 thread 可以复用寄存器里面的数据。CUDA 对 Tensor Core 的指南里面把这个叫做 “WMMA warp-wide macro-instructions”。所以 Tensor Core 的编程模型直接就是针对一整个 warp 写的。事实上，Tensor Core 的代码写起来还是有相当多的限制的，CUDA 给 Tensor Core 提供了C的API：首先用来做乘加的矩阵都需要放在这个叫 wmma::fragment 的变量里面，这个本质上就是定义了一个要放在 SM 寄存器上的存储空间，但是需要提供详细的 FMA 参数：- 第一个参数 Use 是这个 fragment 在 FMA 运算里面的角色，可选项有：matrix_a、matrix_b和 accumulator，含义就是字面意思，也没什么需要再解释的了。
- m，n，k，T 是这一个 warp 里面要处理的的 FMA 子矩阵的形状以及数据类型，不同的 Capability 能够支持的组合还不太一样，比如最基础的就是 a、b 都是 __half，accumulator 是 float，然后 m、n、k 都是 16。
m、n、k 的组合不是任意的，能支持的种类跟 Capability 直接相关，比如 V100 和后来出的 T4 能够支持的就不一样，具体可以在 Programming Guide 里面查。 - 最后这个 Layout 可选项有两个 row_major 和 col_major，代表这个 fragment 在内存里面实际存储的行列主序情况。
load_matrix_sync 和 store_matrix_sync 分别是把数据写到 fragment 空间里面和从这里面取出来写到别的地方去。fill_fragment 对 fragment 初始化。mma_sync 就是对整个 warp 调用 Tensor Core 去跑完这一个 FMA 运算了。常规的写法也是先把矩阵 A、B 都 load 到 shared_memory 上，然后再从 shared_memory 里面取对应 FMA 块大小的数据到 fragment 里面，mma_sync 跑完，最后从 fragment 里面把结果写到外面去。这里的注意点是上面这些代码（包括 fragment 定义以及下面几个函数的调用）都是针对 warp 的，即我们在写代码的一开始就需要考虑到每个 block 里面的 thread 结构，保证一个 warp 的 32 个 thread 执行的代码是完全相同的。相应地，对矩阵的分块也是需要在写代码的时候就考虑清楚，我们要保证每个 warp 处理的 a、b 矩阵的大小刚好是这个地方设定好的 m、n、k。看起来确实相当麻烦，不过想想可能好像也还好，本来如果要写出性能很好的 CUDA 代码来，每个 warp 要算多少东西也是需要精细考虑清楚的。
Volta 这一代对 SIMT 的编程模型也做了改变。在之前的 SIMT 流水线中，如果一个 warp 的指令里面出现了分支，这些分支块是不能被同时执行的。所以一直以来写 CUDA 代码都会要有一个原则是不要在一个 warp 里面出现不同的分支，要不需要花费两倍的时间去处理。这一代开始把 PC 和调用栈做成了每个线程独立的：Volta Warp with Per-Thread Program Counter and Call Stack现在呢，每个分支里面的指令可以在更细粒度的层面上进行混合调度了，也可以手动插入一些在 warp 层面同步的指令进去：Programs use Explicit Synchronization to Reconverge Threads in a Warp白皮书后面给了一个可以从这个改动上得到收益的 Starvation-Free Algorithms 的示例，修改带锁的双向链表的时候，不同 thread 可能会被 block 在锁上，以前的架构应该基本上不太可能能处理得了这种 case，新架构就保证了即使有些 thread 还在等待锁，另外的 thread 也有可能先拉出来跑。可能也是因为这样所以 1 个 Dispatch Unit 配 1 个 Warp Scheduler 了？因为线程指令的实现事实上更加复杂了。
所以其实最后还是同时只能执行一个分支里面的一部分，这个 upgrade 我暂时还没有想到具体的应用场景会有多常出现（上面这个带锁双向链表我觉得写在 CUDA 里面就很不常见啊…），以及会具体有多少性能收益，说不定还是原本的那种简单的设计更直接更高效一些呢。（期待一下未来的硬件里面会不会把这个恢复回去……）
以前 CUDA 编程原则里面不要写分支的那条在新架构下我觉得还是适用的，不写分支就不会有这么多额外的麻烦要考虑了。
另外有一个 Cooperative Group 的新设计倒是看起来感觉更有用一些。原本的 __syncthreads( ) 是针对一个 block 里面的所有 thread 做同步的，现在可以对不同 block 的不同 thread 单独定义同步组了，CUDA launch 的时候会把同一个组的一起 launch 上去，同步可以在一个更加细粒度的层面上完成。6、Turing
Compute Capability：7.5；TU102 GPU包含6个图像处理集群（GPC）、36个纹理处理集群（TPC）和72个流式多元处理器（SM）。
每个GPC均包含一个专用的光栅化引擎和6个TPC，且每个TPC均包含两个SM。每个SM包含：且我们可根据计算或图形工作负载将这些内存设置为不同容量。每个SM中的全新RT核心处理引擎负责执行光线追踪加速。Turing架构采用全新SM设计，每个TPC均包含两个SM，每个SM共有64个FP32核心和64个INT32核心。Turing SM支持并行执行FP32与INT32运算，每个Turing SM还拥有8个混合精度Turing Tensor核心和1个RT核心。
7、Ampere
Compute Capability：8.0；NVIDIA A100在AI训练（半/单精度操作，FP16/32）和推理（8位整数操作，INT8）方面，GPU比Volta GPU强大20倍。在高性能计算（双精度运算，FP64）方面，NVIDIA表示GPU的速度将提高2.5倍。
GA100 Full GPU with 128 SMs (A100 Tensor Core GPU has 108 SMs)A100采用是GA100 GPU的缩减版GA100架构和规格，GA100满配版和A100配置如下。NVIDIA Ampere GA100 GPU架构满配如下：- 8 TPCs/GPC, 2 SMs/TPC, 16 SMs/GPC, 128 SMs per full GPU
- 64 FP32 CUDA Cores/SM, 8192 FP32 CUDA Cores per full GPU
- 4第三代Tensor Cores/SM, 512第三代Tensor Cores per full GPU
- 6 HBM2 stacks, 12 512bit 内存控制器
NVIDIA Ampere GA100 GPU架构配置如下:- 2 SMs/TPC, up to 16 SMs/GPC, 108 SMs
- 64 FP32 CUDA Cores/SM, 6912 FP32 CUDA Cores
- 4第三代Tensor Cores/SM, 432第三代Tensor Cores
- 5 HBM2 stacks，10 512bit 内存控制器
Ampere GA100是迄今为止设计的最大的7nm GPU。GPU完全针对HPC市场而设计，具有科学研究，人工智能，深度神经网络和AI推理等应用程序。NVIDIA A100 是一项技术设计突破，在五项关键技术领域得到创新和突破：- NVIDIA Ampere架构 — A100的核心是NVIDIA Ampere GPU架构，其中包含超过540亿个晶体管，使其成为世界上最大的7纳米处理器。
- 基于TF32的第三代张量核(Tensor Core)： Tensor核心的应用使得GPU更加灵活，更快，更易于使用。TF32包括针对AI的扩展，无需进行任何代码更改即可使FP32精度的AI性能提高20倍。此外， TensorCore 现在支持FP64，相比上一代，HPC应用程序可提供多达2.5倍的计算量。
- 多实例（Multi-Instance）GPU — MIG是一项新技术功能，可将单个A100GPU划分为多达七个独立的GPU，因此它可以为不同大小的作业提供不同程度的计算，从而提供最佳利用率。
- 第三代NVIDIA NVLink —使GPU之间的高速连接速度加倍，可在服务器中提供有效的性能扩展。
- 结构稀疏性—这项新的效率技术利用了AI数学固有的稀疏特性来使性能提高一倍。
NVIDIA A100基于7nm Ampere GA100 GPU，具有6912 CUDA内核和432 Tensor Core，540亿个晶体管数，108个流式多处理器。采用第三代NVLINK，GPU和服务器双向带宽为4.8 TB/s，GPU间的互连速度为600 GB/s。另外，Tesla A100在5120条内存总线上的HBM2内存可达40GB。
NVIDIA GPU架构白皮书系列，下载链接：NVIDIA GPU架构白皮书
《NVIDIA A100 Tensor Core GPU技术白皮书》
《NVIDIA Kepler GK110-GK210架构白皮书》
《NVIDIA Kepler GK110-GK210架构白皮书》
《NVIDIA Kepler GK110架构白皮书》
《NVIDIA Tesla P100技术白皮书》
《NVIDIA Tesla V100 GPU架构白皮书》
《英伟达Turing GPU 架构白皮书》
参考链接：https://jcf94.com/2020/05/24/2020-05-24-nvidia-arch/https://wccftech.com/nvidia-ampere-ga100-gpu-7nm-architecture-specifications-performance-deep-dive/
转载申明：转载本号文章请注明作者和来源，本号发布文章若存在版权等问题，请留言联系处理，谢谢。
推荐阅读
更多架构相关技术知识总结请参考“架构师全店铺技术资料打包”相关电子书(35本技术资料打包汇总详情可通过“阅读原文”获取)。
全店内容持续更新，现下单“全店铺技术资料打包(全)”，后续可享全店内容更新“免费”赠阅，价格仅收188元(原总价290元)。
温馨提示：
扫描二维码关注公众号，点击阅读原文链接获取“架构师技术全店资料打包汇总(全)”电子书资料详情。
