2.9深入GPU硬件架构及运行机制

余光..

已于 2024-06-02 00:22:42 修改

阅读量865

点赞数 22

分类专栏：百人计划文章标签：技术美术图形渲染

于 2024-06-02 00:21:54 首次发布

本文链接：https://blog.csdn.net/qq_65461674/article/details/139381503

版权

百人计划专栏收录该内容

15 篇文章

订阅专栏

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是什么

NVIDIA GPU芯片实物图

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

二、Gpu物理架构

由于纳米工艺的引入，GPU可以将数以亿记的晶体管和电子器件集成在一个小小的芯片内。从宏观物理结构上看，现代大多数桌面级GPU的大小跟数枚硬币同等大小，部分甚至比一枚硬币还小（下图）。

当GPU结合散热风扇、PCI插槽、HDMI接口等部件之后，就组成了显卡（上图）。

显卡不能独立工作，需要装载在主板上，结合CPU、内存、显存、显示器等硬件设备，组成完整的PC机。

三、英伟达GPU架构的发展史

NVIDIA GPU架构历经多次变革，从起初的Tesla发展到最新的Turing架构，发展史可分为以下时间节点：

2008 - Tesla

Tesla最初是给计算处理单元使用的，应用于早期的CUDA系列显卡芯片中，并不是真正意义上的普通图形处理芯片。

2010 - Fermi

Fermi是第一个完整的GPU计算架构。首款可支持与共享存储结合纯cache层次的GPU架构，支持ECC的GPU架构。

2012 - Kepler

Kepler相较于Fermi更快，效率更高，性能更好。

2014 - Maxwell

其全新的立体像素全局光照 (VXGI) 技术首次让游戏 GPU 能够提供实时的动态全局光照效果。基于 Maxwell 架构的 GTX 980 和 970 GPU 采用了包括多帧采样抗锯齿 (MFAA)、动态超级分辨率 (DSR)、VR Direct 以及超节能设计在内的一系列新技术。

2016 - Pascal

Pascal 架构将处理器和数据集成在同一个程序包内，以实现更高的计算效率。1080系列、1060系列基于Pascal架构

2017 - Volta

Volta 配备640 个Tensor 核心，每秒可提供超过100 兆次浮点运算(TFLOPS) 的深度学习效能，比前一代的Pascal 架构快5 倍以上。

2018 - Turing

Turing 架构配备了名为 RT Core 的专用光线追踪处理器，能够以高达每秒 10 Giga Rays 的速度对光线和声音在 3D 环境中的传播进行加速计算。Turing 架构将实时光线追踪运算加速至上一代 NVIDIA Pascal™ 架构的 25 倍，并能以高出 CPU 30 多倍的速度进行电影效果的最终帧渲染。2060系列、2080系列显卡也是跳过了Volta直接选择了Turing架构。

四、Gpu微观物理结构

NVidia Tesla架构

Tesla微观架构总览图

下面将阐述它的特性和概念：

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

NVidia Fermi架构

Fermi微观架构总览图

下面将阐述它的特性和概念：

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

NVidia Maxwell架构

Maxwell微观架构总览图

下面将阐述它的特性和概念：

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

NVidia Turing架构

Turing微观架构总览图

下面将阐述它的特性和概念：

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

Turing单个SM的结构图

每个SM包含：

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

核心组件结构：

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

五、GPU架构的共性

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

GPC（图形处理簇）
TPC（纹理处理簇）
Thread（线程）
SM、SMX、SMM（Stream Multiprocessor，流多处理器）
Warp线程束、Warp Scheduler（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的架构皆是以高度并行能力而设计的

六、GPU渲染总览

Fermi架构的运行机制总览图

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

程序员编写的shader是在SM上完成的。

每个SM包含许多为线程执行数学运算的Core（核心）。

例如，一个线程可以是顶点或像素着色器调用。

这些Core和其它单元由Warp Scheduler驱动，Warp Scheduler管理一组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通过主机接口（Host Interface）接受这些命令，并通过前端（Front End）处理这些命令。

3、在图元分配器(Primitive Distributor)中开始工作分配，处理indexbuffer中的顶点产生三角形分成批次(batches)，然后发送给多个GPCs。这一步的理解就是提交上来n个三角形，分配给这几个GPC同时处理。

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如何克服内存读取延迟的关键，只是简单地切换活动线程组

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)在它接收到的三角形上工作，来负责这些这些三角形的像素信息的生成（同时会处理背面剔除和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、开启Alpha Test：由于Alpha Test需要在像素着色器后面的Alpha Test阶段比较（DX的discard，OpenGL的clip），所以无法在像素着色器之前就决定该像素是否被剔除。

2、开启Alpha Blend：启用了Alpha混合的像素很多需要与frame buffer做混合，无法执行深度测试，也就无法利用Early-Z技术。

3、关闭深度测试。Early-Z是建立在深度测试开启的条件下，如果关闭了深度测试，也就无法启用Early-Z技术。

4、开启Multi-Sampling：多采样会影响周边像素，而Early-Z阶段无法得知周边像素是否被裁剪，故无法提前剔除。

5、以及其它任何导致需要混合后面颜色的操作。

九、SIMD和SIMT

SIMD（Single Instruction Multiple Data）是单指令多数据，在GPU的ALU单元内，一条指令可以处理多维向量
（一般是4D）的数据。比如，有以下shader指令：
float4 c = 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技术，只需一条指令即可处理完：
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存取的数据可以是不同的。
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]; //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架构：

右图一是分离式架构，CPU和GPU各自有独立的缓存和内存，它们通过PCI-e等总线通讯。这种结构的缺点在于 PCI-e 相对于两者具有低带宽和高延迟，数据的传输成了其中的性能瓶颈。目前使用非常广泛，如PC等。

右图二是耦合式架构，CPU 和 GPU 共享内存和缓存。AMD 的 APU 采用的就是这种结构，目前主要使用在游戏主机中，如 PS4、智能手机。

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

十三、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 条指令流（Instruction Streams），32 个并发程序片元（Fragment）。

漫反射例子说明：