Chenfan Blog

Do cool things that matter.

0%

NVIDIA GPU 架构演进

最近 NV 刚发布了新架构的 GPU,自之前理过 Intel CPU 的架构变化 后再来理一下 GPU 的。

硬/软件接口那篇有介绍过 GPU 的结构,当时也是以 Fermi 架构为例的,NV 很有意思的是会用一些历史上杰出的科学家的名字来命名自己的硬件架构。

总体上,NV GPU 用到的 SIMT 基本编程模型都是一致的,每一代相对前代基本都会在 SM 数量、SM 内部各个处理单元的流水线结构等等方面有一些升级和改动。这篇暂时不涉及到渲染管线相关的部分,其他诸如多少 nm 工艺、内存频率提升等等也都先略过,只关注计算相关的硬件架构演进。


Tesla

关于初代 GPU 的架构,找到的资料不太多,基本上都是从 Fermi 开始的。

Fermi

Compute Capability: 2.0, 2.1

Fermi 架构 SM

每个 SM 中包含:

  • 2 个 Warp Scheduler/Dispatch Unit
  • 32 个 CUDA Core(分在两条 lane 上,每条分别是 16 个)
    • 每个 CUDA Core 里面是 1 个单精浮点单元(FPU)和 1 个整数单元(ALU),可以直接做 FMA 的乘累加
    • 每个 cycle 可以跑 16 个双精的 FMA
  • 16 个 LD/ST Unit
  • 4 个 SFU

我的理解是做一个双精 FMA 需要用到两个 CUDA Core?所以是 32 / 2 = 16

Kepler

Compute Capability: 3.0, 3.2, 3.5, 3.7

这一代 SM 整体结构上跟之前是一致的,只不过升级完了以后又往里面塞进去了更多的运算单元,其他部分也没有做太大的改动。

Kepler 架构 SM

每个 SM(这里叫 SMX 了)中包含:

  • 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 要高出一截。

Maxwell

Compute Capability: 5.0, 5.2, 5.3

Maxwell 架构 SM

可能是觉得 Kepler 往一个 SM 里面塞了太多东西,其实最终效率也并没有那么高,这一代的 SM 开始做减法了,每个 SM(SMM)中包含:

  • 4 个 Warp Scheduler,8 个 Dispatch Unit
  • 128 个 CUDA Core(4 * 32)
  • 32 个 SFU 和 LD/ST Unit(4 * 8)

Kepler 里面 192 这个数字也被诟病了(不是 2 的倍数)。

这些硬件单元的流水线分布也不再是像 Kepler 那样大锅炖了,而是有点像是把 4 个差不多像是 Fermi 的 SM 拼在一起组成一个 SM:
每个 Process Block 里面是:

  • 1 个 Warp Scheduler 和 2 个 Dispatch Unit
  • 32 个 CUDA Core
  • 8 个 SFU 和 LD/ST 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 科学计算才用的上的双精单元在这代上不太有必要吧。

Pascal

Compute Capability: 6.0, 6.1, 6.2

这一代可以说是有了质的飞跃,还是先从 SM 开始:

Pascal 架构 SM

可以看到一个 SM 内的部分作了进一步的精简,整体思路是 SM 内部包含的东西越来越少,但是总体的片上 SM 数量每一代都在不断增加,每个 SM 中包含:

  • 2 个 Warp Scheduler,4 个 Dispatch Unit
  • 64 个 CUDA Core(2 * 32)
  • 32 个双精浮点单元(2 * 16,双精回来了!)
  • 16 个 SFU 和 LD/ST Unit(2 * 8)

一个 SM 里面包含的 Process Block 数量减少到了 2 个,每个 Process Block 内部的结构倒是 Maxwell 差不多:

  • 1 个 Warp Scheduler 和 2 个 Dispatch Unit
  • 32 个 CUDA Core
  • 多了 16 个 DP Unit
  • 8 个 SFU 和 LD/ST 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 倍,猜测应该是一个单精单元用来算两个半精的计算。

Volta

Compute Capability: 7.0, 7.2

又一个针对深度学习的质变 Feature,Tensor Core!

Volta 架构 SM

看到 SM 的时候我们会发现这一代除了多出了一个额外的 Tensor Core 的单元以外,怎么 SM 的体积看起来好像又加回去了,每个 SM 中包含:

  • 4 个 Warp Scheduler,4 个 Dispatch Unit(发现不需要配 2 个 Dispatch 给每个 Scheduler 了?白皮书里面倒是没有对这个的解释)
  • 64 个 FP32 Core(4 * 16)
  • 64 个 INT32 Core(4 * 16)
  • 32 个 FP64 Core(4 * 8)
  • 8 个 Tensor Core (4 * 2)
  • 32 个 LD/ST Unit(4 * 8)
  • 4 个 SFU(发现对特殊计算的需求减少了?)

事实上相比 Pascal 而言,单个 SM 中的单精运算单元数量是一致的,相当于把 Pascal 中的每个 Process Block 进一步地又拆成了 2 个,每个 Process Block 中包含:

  • 1 个 Warp Scheduler,1 个 Dispatch Unit
  • 16 个 FP32 Core
  • 16 个 INT32 Core
  • 8 个 FP64 Core
  • 2 个 Tensor Core
  • 8 个 LD/ST Unit
  • 1 个 SFU

这里把原本的 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 设计里面说不定也会出现其他 xx Core。好在 FMA 除了深度学习以外在 HPC 的 workload 里面也是挺常见的,这个设计以后还是比较有用的。

Tensor Core 4x4 Matrix Multiply and Accumulate

Mixed Precision Multiply and Accumulate in Tensor Core

Tensor 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:

1
2
3
4
5
6
7
template<typename Use, int m, int n, int k, typename T, typename Layout=void> class fragment;

void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm);
void load_matrix_sync(fragment<...> &a, const T* mptr, unsigned ldm, layout_t layout);
void store_matrix_sync(T* mptr, const fragment<...> &a, unsigned ldm, layout_t layout);
void fill_fragment(fragment<...> &a, const T& v);
void mma_sync(fragment<...> &d, const fragment<...> &a, const fragment<...> &b, const fragment<...> &c, bool satf=false);

PTX 的指令应该更多一些,不过我没有详细看过。

首先用来做乘加的矩阵都需要放在这个叫 wmma::fragment 的变量里面,这个本质上就是定义了一个要放在 SM 寄存器上的存储空间,但是需要提供详细的 FMA 参数:

  • 第一个参数 Use 是这个 fragment 在 FMA 运算里面的角色,可选项有:matrix_amatrix_baccumulator,含义就是字面意思,也没什么需要再解释的了。
  • 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_majorcol_major,代表这个 fragment 在内存里面实际存储的行列主序情况。

load_matrix_syncstore_matrix_sync 分别是把数据写到 fragment 空间里面和从这里面取出来写到别的地方去。fill_fragmentfragment 初始化。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 要算多少东西也是需要精细考虑清楚的。

SIMT Model Upgrade & COOPERATIVE GROUPS

Volta 这一代对 SIMT 的编程模型也做了改变。

在之前的 SIMT 流水线中,如果一个 warp 的指令里面出现了分支,这些分支块是不能被同时执行的。所以一直以来写 CUDA 代码都会要有一个原则是不要在一个 warp 里面出现不同的分支,要不需要花费两倍的时间去处理。

SIMT Warp Execution Model of Pascal and Earlier GPUs

这一代开始把 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 上去,同步可以在一个更加细粒度的层面上完成。

Turing

Compute Capability: 7.5

Ampere

Compute Capability: 8.0


To be continued.