【笔记】【百人计划】图形2.7.2 GPU硬件架构概述

思考

image-20220822162918576

简述

GPU是Graphics Processing Unit（图形处理单元）的缩写。顾名思义，GPU最初就是用于更高效的图像的绘制和图元数据处理。和CPU相比，GPU简化了控制模块与Cache模块，但是具有大量的ALU（算术逻辑单元），并且可以进行并行运算。因此，CPU更擅长逻辑、控制、串行运算，GPU更易于运行并行程序。

GPU和通常所说的显卡严格来说是据有区别的，GPU是显卡（Video card、Display card、Graphics card）最核心的部件，但除了GPU，显卡还有扇热器、通讯元件、与主板和显示器连接的各类插槽。

image-20220822162950559

我们通常所说的显卡和GPU实际上具有一定区别。GPU是显卡的核心部分，再加上散热器，以及和主板、显示器的各种插槽等，就构成了显卡。显卡不能独立工作，需要装载到主板上，结合CPU、内存、显存、显示器等硬件组成完整的PC

image-20220822163026806

当前市场主流的GPU厂商有Nvidia、AMD、ARM、Imagination、Qualcomm、苹果。

目前在PC、手机、VR一体机等设备上，GPU都发挥着作用，为我们在设备上绘制出更快速、优质的画面奠定了基础。甚至除了绘制，在人工智能领域，近几年神经网络能够成为主流，GPU在其中也起到重要的作用。

发展历史

1999年，Nvidia的GeForce 256 发布，Nvidia将GPU定义为“集成了变换，照明，三角形设置/裁剪和渲染引擎的单芯片处理器，每秒能够处理至少1000万个多边形。GeForce 256被认为是第一块完成商业化的GPU，将GPU引入市场，是市场上第一款采用硬件加速T&L（Transformation&Lighting）的消费级卡。而在GeForce 256之前，也并非没有“具有完整T&L引擎的显卡”，但是它们从未引入市场。实际上，20世纪70年代就有了“GPU”的概念，索尼于1994年PS1推出时使用“GPU”的术语。该系统有一个32位索尼GPU（由东芝设计）。这个阶段的GPU也已经基本具备渲染、Z缓冲、Alpha混合、雾计算、模板缓冲、纹理映射、纹理过滤等功能。

从下面开始，以Nvidia的GPU发展过程为代表来讲述现代GPU和GPU架构的发展历史。

2001年的GeForce3 开始，引入了可编程的渲染管线（可编程顶点着色器，可配置的32位浮点片段着色管线）。

2006年推出了GeForce8 系列（G8x），也正是在这一代的GPU的基础上，开始有了GPU架构——Tesla架构。

2010年的Fermi架构是第一个完整的GPU计算架构，是NVIDIA GPU 架构自初代 G80 以来最重大的飞跃。后来的Kepler架构、Maxwell架构都是基于Fermi架构的基础。

2012年的Kepler架构显卡（600，700系列）和2014年的Maxwell架构显卡采用28nm工艺制造，使性能提高了20%。

2014年的Maxwell架构，立体像素全局光照VXGI技术首次让游戏GPU能够提供实时动态全局光照

2016年的GeForce GTX系列采用接替Maxwell的Pascal架构。Pascal 架构将处理器和数据集成在同一个程序包内，以实现更高的计算效率。1080系列、1060系列都是基于Pascal架构。

2017年发布的Volta 架构配备了640 个Tensor core，这是专为深度学习设计的内核，每秒可提供超过100 兆次浮点运算(TFLOPS) 的深度学习效能，比前一代的Pascal 架构快5 倍以上。使用Volta架构的Titan V显卡也是专用于AI和仿真。

从2018年开始的Turing架构，融合了Volta专用的Tensor Core，同时增加了RT core，这是光线追踪技术的核心。RT Core 能够以高达每秒 10 Giga Rays 的速度对光线和声音在 3D 环境中的传播进行加速计算。Turing 架构将实时光线追踪运算加速至上一代 NVIDIA Pascal架构的 25 倍，并能以高出 CPU 30 多倍的速度进行电影效果的最终帧渲染。为了体现对光线追踪的支持，NVIDIA在这之后的显卡前缀都采用了RTX。

2020年的RTX 30系列采用了最新的Ampere架构。Ampere架构的Tensor core采用新的Tensor Float32（TF32）与64位浮点（FP64）精度标准。RT core也进行了更新，运算能力达到第一代的两倍。

GPU架构组成

每一代的GPU架构虽然具有差异，但也有一些相同的概念和组成部分。

以Fermi架构为例，Fermi架构拥有16个SM （Stream Multiprocessor 流多处理器/计算单元）。SM支持并发执行多个thread。Kepler、Maxwell架构中甚至对SM升级为SMX、SMM。程序员编写的Shader是在SM上运行的。

每个SM包含：

  • 2个Warp Scheduler
  • 2个Dispatch Unit（分发单元）
  • 2个Warp（线程束）
  • 16组LD/ST（加载存储单元）
  • 4个SFU（特殊函数单元）
  • 128KB Register File（寄存器）
  • 64KB Shared Memory /L1 Cache
  • Uniform Cache（全局缓存）

每个Warp包含32个Core（也叫做SP——Stream Processor）；

每个Core包含1个FPU（浮点数单元），1个ALU（算术逻辑单元）

Warp Scheduler负责Warp调度，Warp Scheduler的指令通过Dispatch Units送到运算核心（Core）执行。

LD/ST(Load/Store)加载/存储模块用于辅助一个Warp从Shared Memory或显存加载或存储数据。

对于L1 Cache，部分GPU架构中是与Shared Memory共用的（如Fermi架构）。

Turing：

image-20220822163459068

共性：

  • GPC（图形处理簇）
  • TPC（纹理处理簇）
  • Thread（线程）
  • SM、SMX、SMM（流多处理器）
  • Warp线程束、Warp Scheduler（Warp编排器）
  • SP（Streaming Processor，流处理器）
  • Core（执行数学运算的核心）
  • ALU（逻辑运算单元）
  • FPU（浮点运算单元）
  • SFU（特殊函数单元）
  • ROP（render output unit，渲染输入单元）
  • Load/Store Unit（加载存储单元）
  • L1 Cache
  • L2 Cache
  • Shared Memory（共享内存）

微观物理架构

  • GPC
    • TPC
      • SM
        • Poly Morph Engine（多边形引擎）
        • L1 Cache
        • Shared Memory
        • Core
          • ALU
          • FPU
          • Execution Context（执行上下文）
            • 汇编代码会被GPU推送到执行上下文（Execution Context），然后ALU会逐条获取（Detch）、解码（Decode）汇编指令，并执行它们。

GPU存储架构

内存架构

GPU的存储架构分级由

  • Register File (Memory)寄存器
    • 位于每个SM中，访问速度最快的存储体，用于存放线程执行时所需的变量
  • Shared Memory共享内存
    • 位于每个SM中
  • L1 Cache
  • L2 Cache
  • Constant Memory常量内存
    • 位于每个SM中和片外的RAM存储器中
  • Texture Memory纹理内存
    • 位于每个SM中和片外的RAM存储器中
  • Global Memory全局内存
    • 位于片外存储体。容量大，访问延迟高、传输速度较慢，使用二级缓存L2 Cache做缓冲
  • Local Memory本地内存
    • 一般位于片内存储体，变量、数组、结构体等都存放在此处，但是有大数组、大结构体以至于寄存器区放不下他们，编译器在编译阶段就会将他们放到片外的DDR芯片中（最好的情况也会放到L2 Cache），且将他们标记为“Local”型

构成，它们的存取速度从寄存器到显存依次变慢。

他们物理上所在的位置，决定了他们的速度、大小以及访问规则

存储类型 访问周期
寄存器 1
共享内存 1~32
L1缓存 1~32
L2缓存 32~64
纹理、常量缓存 400~600
全局内存 400~600

在不同的平台（PC或移动设备），GPU的存储架构分为分离式和耦合式（是否共享内存）。

分离式架构中CPU和GPU各自有独立的缓存和内存，通过PCI-e总线通讯，缺点是PCI-e相对于两者具有低带宽和高延迟，数据的传输成为其中的性能瓶颈。一般用于PC；

耦合式架构中CPU与GPU共享内存和缓存。AMD的APU采用的就是这种架构。目前主要是用在游戏主机、移动设备中，如PS4，智能手机。

img

在储存管理方面，分离式架构中CPU与GPU各自拥有独立的内存，两者共享一套虚拟地址空间，必要时会进行内存拷贝。对于耦合式架构，GPU没有独立的内存，与CPU共享系统内存，由MMU进行存储管理。

image-20220822173521235

image-20220822173644684

GPU逻辑管线

我们根据GPU的渲染过程来了解每个部分在GPU中发挥的作用。以Fermi的SM为例

1、程序通过图形API(DX、GL、WEBGL)发出draw call指令，指令会被推送到驱动程序，驱动会检查指令的合法性，然后会把指令放到GPU可以读取的Push buffer中。

2、经过一段时间或者显式调用flush指令后，驱动程序把Push buffer的内容发送给GPU，GPU通过主机接口（Host Interface）接受这些命令，并通过前端（Front End）处理这些命令。

3、在图元分配器(Primitive Distributor)中开始工作分配，处理index buffer中的顶点产生三角形分成批次(batches)，然后发送给多个GPCs（Graphics Processing Cluster）。

4、在GPC中，每个SM中的Poly Morph Engine负责通过三角形索引(triangle indices)取出三角形的数据(vertex data)——Vertex Fetch。

5、在获取数据之后，在SM中以32个线程为一组的线程束(Warp)来调度，来开始处理顶点数据。

6、SM的warp调度器会按照顺序分发指令给整个warp，单个warp中的线程会锁步(lock-step)执行各自的指令，如果线程碰到不激活执行的情况也会被遮掩(be masked out)。

7、warp中的指令可以被一次完成，也可能经过多次调度，例如通常SM中的LD/ST(加载存取)单元数量明显少于基础数学操作单元。

8、由于某些指令比其他指令需要更长的时间才能完成，特别是内存加载，warp调度器可能会简单地切换到另一个没有内存等待的warp，这是GPU如何克服内存读取延迟的关键，只是简单地切换活动线程组。为了使这种切换非常快，调度器管理的所有warp在寄存器文件中都有自己的寄存器。这里就会有个矛盾产生，Shader需要越多的寄存器，就会给warp留下越少的空间，就会产生越少的warp，这时候在碰到内存延迟的时候就会只是等待，而没有可以运行的warp可以切换。

9、一旦warp完成了vertex shader的所有指令，运算结果会被Viewport Transform模块处理，三角形会被裁剪然后准备栅格化，GPU会使用L1和L2缓存来进行vertex shader和pixel shader的数据通信。

10、接下来这些三角形将被分割，再分配给多个GPC，三角形的范围决定着它将被分配到哪个光栅引擎(raster engines)，每个raster engines覆盖了多个屏幕上的tile，这等于把三角形的渲染分配到多个tile上面。也就是像素阶段就把按三角形划分变成了按显示的像素划分了。

11、SM上的Attribute Setup保证了从vertex shader来的数据经过插值后是pixel shade是可读的。

12、GPC上的光栅引擎(raster engines)在它接收到的三角形上工作，来负责这些这些三角形的像素信息的生成（同时会处理裁剪Clipping、背面剔除和Early-Z剔除）。

13、32个像素线程将被分成一组，或者说8个2x2的像素块，这是在像素着色器上面的最小工作单元，在这个像素线程内，如果没有被三角形覆盖就会被遮掩，SM中的warp调度器会管理像素着色器的任务。

14、接下来的阶段就和vertex shader中的逻辑步骤完全一样，但是变成了在像素着色器线程中执行。 由于不耗费任何性能可以获取一个像素内的值，导致锁步执行非常便利，所有的线程可以保证所有的指令可以在同一点。

15、最后一步，现在像素着色器已经完成了颜色的计算还有深度值的计算，在这个点上，我们必须考虑三角形的原始api顺序，然后才将数据移交给ROP(render output unit，渲染输入单元)，一个ROP内部有很多ROP单元，在ROP单元中处理深度测试，和framebuffer的混合，深度和颜色的设置必须是原子操作，否则两个不同的三角形在同一个像素点就会有冲突和错误。

SIMD和SIMT

  • SIMD（Single Instruction Multiple Data），单指令多数据，在GPU的ALU内，一条指令可以处理多维向量（一般是4D）的数据，比如以下shader指令
    • float4 c=a+b;
1
2
3
4
5
6
7
;对于没有SIMD的处理单元，需要4条指令将4个float数值相加，汇编如下
ADD c.x, a.x, b.x
ADD c.y, a.y, b.y
ADD c.z, a.z, b.z
ADD c.w, a.w, b.w
;有了SIMD技术，只需一条指令
SIMD_ADD c,a,b
  • SIMT（Single Instruction Multiple Threads），单指令多线程，是SIMD的升级版，可对GPU中单个SM中的多个Core同时处理同一指令，并且每个Core存取的数据可以不同
1
2
3
4
5
6
7
SIMT_ADD c,a,b
;上述指令会被同时送入在单个SM中被编组的所有Core中，同时执行运算，但a,b,c的值可以不同
;相当于
__global__ void add(float *a, float *b, float *c) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	a[i] = b[i] + c[i];
}

co-issue

co-issue是为了解决SIMD运算单元无法充分利用的问题。如下图，由于float的数量不同，ALU的利用率也不同

image-20220822171640666

为了解决着色器在低维向量的利用率低的问题，可以通过合并1D与3D或2D与2D的指令。例如下图，DP3指令用了3D数据，ADD指令只有1D数据，co-issue会自动将它们合并，在同一个ALU只需一个指令周期即可执行完

image-20220822171813268

但是对于向量运算单元（Vector ALU），如果其中一个变量既是操作数又是存储数的情况，无法使用co-issue

总结

  • vs和fs都是在同一个单元中执行的，vs按照三角形来并行处理，fs按照像素来并行处理。
  • vs和fs中的数据通过L1和L2缓存传递
  • warp和thread都是逻辑上的概念，sm和sp、core都是物理上的概念。线程束≠流处理器数

优化建议：

  1. 尽量使用自己扩展的几何实例化代替unity提供的静态合批，动态合批。前者将合并mesh增加额外的vbo内存占用。后者则会增加cpu端的耗时开销。
  2. 尽量减少顶点数与三角形面数。前者减少vs计算，另外可以减少gpu显存中frameData的内存存储。后者减少fs的消耗
  3. 避免每帧提交buffer数据，比如unity的CPU版本粒子系统。可使用GPU版本粒子系统，将修改数据移动到gpu端。另外特别提醒尽量避免大片的透明粒子特效，这将造成严重的overdraw
  4. 减少渲染状态的设置与获取。例如在update中获取设置shader的属性或公共变量。因为前面说到cpu通过MMIO获取寄存器数据，这将耗费更多的时间周期。
  5. 3D物件应使用LOD减少处理的顶点与面数消耗，开启mipmap减少贴图缓存命中的丢失
  6. 避免Alpha Test的使用，造成Early-Z失效
  7. 避免三角面过小，这会加剧过度绘制的情况。也就是一个三角形只占3个像素点，却使用了12个线程去计算像素值，然后遮蔽其余9个的计算机结果。
  8. 在寄存器数量与变体中寻找平衡，使用if变量达成静态分支，取代变体。一方面可以减少变体数量，一方面可以使得URP中的SRP Batch更高效合批
  9. 避免动态的判断分支。也就是Shader中if true和false都会走的情况
  10. 减少复杂函数调用。

参考资料

[1] https://www.cnblogs.com/timlly/p/11471507.html#22-gpu%E5%8E%86%E5%8F%B2

[2] https://www.nvidia.cn/data-center/ampere-architecture/

[3] 图形处理单元 - 维基百科，自由的百科全书 (wikipedia.org)

[4] https://www.bilibili.com/video/BV1aM4y1g75f

[5] 云渲染与分布式并行集群渲染技术 (daimajiaoliu.com)