笔记13-图形 2.7.2 GPU硬件架构概述 | IronName

课前问题

1、GPU是如何与CPU协调工作的？
2、GPU也有缓存机制吗？有几层？它们的速度差异多少？
3、GPU的渲染流程有哪些阶段？它们的功能分别是什么？
4、Early-Z技术是什么？发生在哪个阶段？这个阶段还会发生什么？
5、SIMD和sIMT是什么？它们的好处是什么？co-issue呢？
6、GPU是并行处理的么？若是，硬件层是如何设计和实现的？
7、GPC、TPC、SM是什么？Warp又是什么？它们和Core、Thread之间的关系如何？
8、顶点着色器（VS）和像素着色器（PS）可以是同一处理单元吗？为什么？
9、像素着色器（PS）的最小处理单位是1像素吗？为什么？会带来什么影响？
10、Shader中的if、for等语句会降低渲染效率吗？为什么？
11、渲染相同面积的图形，三角形数量少的还是数量多的效率更快？为什么？
12、GPU Context是什么？有什么作用？
13、造成渲染瓶颈的问题很可能有哪些？该如何避免或优化它们？

GPU是什么

GPU全称是Graphics Processing Unit，图形处理单元。它的功能最初与名字一致，是专门用于绘制图像和处理图元数据的特定芯片，后来渐渐加入了其它很多功能。
我们日常讨论GPU和显卡时，经常混为一谈，严格来说是有所区别的。GPU是显卡（Videocard、Display card、Graphics card）最核心的部件,但除了GPU，显卡还有散热器、通讯元件、与主板和显示器连接的各类插槽。

GPU物理架构

由于纳米工艺的引入，GPU可以将数以亿记的晶体管和电子器件集成在一个小小的芯片内。从宏观物理结构上看，现代大多数桌面级GPU的大小跟数枚硬币同等大小，部分甚至比一枚硬币还小。
当GPU结合散热风扇、PCI插槽、HDMI接口等部件之后，就组成了显卡（上图）。
显卡不能独立工作，需要装载在主板上，结合CPU、内存、显存、显示器等硬件设备，组成完整的PC机。

微观物理结构

Nvidia Tesla 架构

拥有7组TPC（Texture/Processor Cluster，纹理处理簇)
每个TPc有两组SM（Stream Multiprocessor，流多处理器）
每个SM包含：
8个SP（Streaming Processor，流处理器）
2个SFU（Special Function Unit，特殊函数单元)
L1缓存、MTIssue（多线程指令获取）
C-Cache（常量缓存）、共享内存
除了TPC核心单元，还有与显存、CPU、系统内存交互的各种部件。

Nvidia Fermi架构

拥有16 个SM
2 个Warp Scheduler（线程束）
两组共 32个Core
16组加载存储单元（LD/ST）
4个特殊函数单元（SFU）
分发单元（Dispatch Unit）
每个Core: 1个FPU（浮点数单元）; 1个ALU（逻辑运算单元）

Nvidia Maxwell架构

采用了Maxwell的GM204，拥有4个GPC，每个GPC有4个SM，
对比Tesla架构来说，在处理单元上有了很大的提升。

Nvidia Turing架构

6 GPC（图形处理簇）, 36 TPC（纹理处理簇）, 72 SM（流多处理器）
每个GPC有6个TPC，每个TPC有2个SM
4,608 CUDA核; 72 RT核; 576 Tensor核; 288 纹理单元
12x32位 GDDR6内存控制器(共384位)

每个SM包含：

64 CUDA核 ( CUDA是NVIDIA推出的统一计算架构 )
8 Tensor核 (TensorCore是专为执行张量或矩阵运算而设计的专用执行单元 )
256KB 寄存器文件

GPU架构的共性

纵观上所有GPU架构，可以发现它们虽然有所差异，但存在着很多相同的概念和部件：

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

GPU为什么会有这么多层级且有这么多雷同的部件？答案是GPU的任务是天然并行的，现代GPU的架构皆是以高度并行能力而设计的

核心组件结构

包含关系 $GPC\owns TPC\owns SM\owns CORE$

SM 中包含Poly Morph Engine（多边形引擎）、L1Cache（L1缓存）、Shared Memory（共享内存）、Core（执行数学运算的核心）等
CORE 中包含ALU、FPU、Execution Context（执行上下文）、（Detch）、解码（Decode）

架构的运行

运行机制图，略

从 Fermi 开始NVIDIA使用类似的原理架构，使用一个GigaThreadEngine来管理所有正在进行的工作，GPU 被划分成多个GPCs(Graphics Processing Cluster)，每个GPC拥有多个SM（SMX、SMM）和一个光栅化引擎(RasterEngine)，它们其中有很多的连接，最显著的是Crossbar，它可以连接GPCs和其它功能性模块（例如ROP或其他子系统）。

程序员编写的 shader 是在 SM 上完成的。每个SM包含许多为线程执行数学运算的 core（核心）。例如，一个线程可以是顶点或像素着色器调用。这些core和其它单元由Warp Scheduler 驱动，WarpScheduler管理一组32个线程作为Warp（线程束）并将要执行的指令移交给 Dispatch Units。
GPU中实际有多少这些单元（每个GPC有多少个SM，多少个GPC…）取决于芯片配置本身

GPU逻辑管线

了解的部件和概念之后，可以深入阐述GPU的渲染过程和步骤。下面将以Fermi家族的SM为例，进行逻辑管线的详细说明。

1、程序通过图形API (DX、GL、WEBGL) 发出 drawcall 指令，指令会被推送到驱动程序，驱动会检查指令的合法性，然后会把指令放到 GPU 可以读取的 Pushbuffer 中。
2、经过一段时间或者显式调用 flush 指令后，驱动程序把 Pushbuffer 的内容发送给GPU，GPU通过主机接口（HostInterface）接受这些命令，并通过前端（Front End）处理这些命令。
3、在图元分配器 (Primitive Distributor) 中开始工作分配，处理 indexbuffer 中的顶点产生三角形分成批次(batches)，然后发送给多个GPCs。这一步的理解就是提交上来 n 个三角形，分配给这几个GPC同时处理。4、在GPC中，每个SM中的Poly Morph Engine负责通过三角形索引（triangle indices）取出三角形的数据(vertex data)，即图中的VertexFetch模块。
5、在获取数据之后，在SM中以32个线程为一组的线程束 (Warp) 来调度，来开始处理顶点数据

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

线程被遮掩是什么意思?

虽然被遮掩的线程没干活，但它们仍然占用了 warp 的一个指令周期（因为 warp 作为一个整体被调度）。如果分支频繁发散，会导致 warp 执行序列化，降低计算资源利用率（这就是为什么编程时要尽量避免 warp 内分支发散）。
“除非一个WARP的32个线程全都走到同一个分支，否则相当于所有分支都走一遍。”

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

8、由于某些指令比其他指令需要更长的时间才能完成，特别是内存加载，warp 调度器可能会简单地切换到另一个没有内存等待的 warp，这是 GPU 如何克服内存读取延迟的关键，只是简单地切换活动线程组

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

10、接下来这些三角形波分割，再分配给多个GPC，三角形的范围决定着它将被分配到哪个光栅引擎(rasterengines)，每个rasterengines覆盖了多个屏幕上的tile，这等于把三角形的渲染分配到多个tile上面。也就是像素阶段就把按三角形划分变成了按显示的像素划分了。
11、SM上的AttributeSetup保证了从vertex-shader来的数据经过插值后是pixel-shade是可读的。

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

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

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

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

Early-Z

早期GPU的渲染管线的深度测试是在像素着色器之后才执行，这样会造成很多本不可见的像素执行了耗性能的像素着色器计算。后来，为了减少像素着色器的额外消耗，将深度测试提至像素着色器之前，这就是Early-Z技术的由来。Early-Z技术可以将很多无效的像素提前剔除，避免它们进入耗时严重的像素着色器。Early-Z剔除的最小单位不是1像素，而是像素块（ $2*2$ ）
但是，以下情况会导致Early-Z失效：1、开启AlphaTest：由于 Alpha Test 需要在像素着色器后面的 Alpha Test 阶段比较（Dx 的 discard，OpenGL的clip），所以无法在像素着色器之前就决定该像素是否被剔除。2、开启AlphaBlend：启用了Alpha混合的像素很多需要与framebuffer做混合，无法执行深度测试，也就无法利用Early-Z技术。3、关闭深度测试。Early-Z是建立在深度测试开启的条件下，如果关闭了深度测试，也就无法启用Early-Z技术。4、开启Multi-Sampling：多采样会影响周边像素，而Early-Z阶段无法得知周边像素是否被裁剪，故无法提前剔除。5、以及其它任何导致需要混合后面颜色的操作。

SIMD和SIMT

SIMD（Single Instruction Multiple Data）是单指令多数据，在GPU的ALU单元内，一条指令可以处理多维向量（一般是4D）的数据。比如，有以下shader指令：
float4c=a+b;// a,b都是float4类型
对于没有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技术，只需一条指令即可处理完：

1 2	SIMD_ADD c, a, b // for(i=0;i<n;++i) a[i]=b[i]+c[i];

SIMT （Single Instruction Multiple Threads，单指令多线程）是SIMD的升级版，可对GPU中单个SM中的多个Core同时处理同一指令，并且每个Core存取的数据可以是不同的。

1	SIMT_ADD c, a, b

上述指令会被同时送入在单个SM中被编组的所有Core中，同时执行运算，但a、b、c的值可以不一样：

_global_void add(float *a, float *b, float *c) {
int i = blockldx.x * blockDim.x + threadldx.x;
a[i]=b[i]+c[i]; //no loop!
}

co-issue

co-issue 是为了解决SIMD运算单元无法充分利用的问题。由于float数量的不同，ALU利用率从100%依次下降为75%、50%、25%。
为了解决着色器在低维向量的利用率低的问题，可以通过合并1D与3D或2D与2D的指令。例如下图，DP3指令用了3D数据，ADD指令只有1D数据，co-issue会自动将它们合并，在同一个ALU只需一个指令周期即可执行完。
但是，对于向量运算单元（Vector ALU），如果其中一个变量既是操作数又是存储数的情况，无法启用co-issue技术。

CPU与GPU

对比

CPU是一个具有多种功能的优秀领导者。它的优点在于调度、管理、协调能力强，但计算能力一般
GPU相当于一个接受CPU调度的“拥有大量计算能力”的员工。

架构

根据CPU和GPU是否共享内存，可分为两种类型的CPU-GPU架构：

一是分离式架构，CPU和GPU各自有独立的缓存和内存，它们通过PCI-e等总线通讯。这种结构的缺点在于PCI-e相对于两者具有低带宽和高延迟，数据的传输成了其中的性能瓶颈。目前使用非常广泛，如PC等。
二是耦合式架构，CPU和GPU共享内存和缓存。AMD的APU采用的就是这种结构，目前主要使用在游戏主机中，如PS4、智能手机。
在存储管理方面，分离式结构中CPU和GPU各自拥有独立的内存，两者共享一套虚拟地址空间，必要时会进行内存拷贝。对于耦合式结构，GPU没有独立的内存，与CPU共享系统内存，由MMU进行存储管理。

GPU资源机制

内存架构

GPU与CPU类似，也有多级缓存结构：寄存器、L1缓存、L2缓存、GPU显存、系统显存
它们的存取速度从寄存器到系统内存依次变慢
由此可见，shader直接访问寄存器、L1、L2缓存还是比较快的，但访问纹理、常量缓存和全局内存非常慢，会造成很高的延迟。

GPU 内存分布在在 RAM 存储芯片或者 GPU 芯片上，他们物理上所在的位置，决定了他们的速度、大小以及访问规则

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

CPU-GPU数据流

分离式架构的CPU-GPU的数据流程：
1、将主存的处理数据复制到显存中。
2、CPU指令驱动GPU。
3、GPU中的每个运算单元并行处理。此步会从显存存取数据。
4、GPU将显存结果传回主存。

Shader 运行机制

在执行阶段，CPU端将shader二进制指令经由PCI-e推送到GPU端，GPU在执行代码时，会用Context将指令分成若干Channel推送到各个Core的存储空间。
右图为一个假象的Core：一个GPU Core包含8个ALU，4组执行环境（Execution context），每组有8个Ctx。这样，一个Core可以并发（Concurrent but interleaved）执行4条指令流（InstructionStreams），32个并发程序片元（Fragment）。