图形 2.9 GPU硬件架构概述

GPU硬件架构概述

 

GPU是什么？

 

　　GPU全称是Graphics Processing Unit，图形处理单元。它的功能最初与名字一致，是专门用于绘制图像和处理图元数据的特定芯片，后来渐渐加入了其它很多功能。显卡就是由GPU、散热器、通讯元件、主板和显示器连接的各类插槽组成的。

 

 

GPU物理构架

 

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

　　当GPU结合散热风扇、PCI插槽、HDMI接口等部件之后，就组成了显卡，但显卡不能独立工作，需要装载在主板上，结合CPU、内存、显存、显示器等硬件设备，组成完整的PC机。

　　下面是历经多次变革的NVIDIA GPU架构：

 

  • 2008 - Tesla：最初给计算处理单元使用，应用于早期的CUDA系列显卡芯片中，但是并非真正意义上的图形处理芯片。
  • 2010 - Fermi ：第一个完整的GPU计算架构。首款可支持与共享存储结合纯cache层次的GPU架构，支持ECC的GPU架构。
  • 2012 - Kepler：相较于Fermi更快，效率更高，性能更好。
  • 2014 - Maxwell：其全新的立体像素全局光照 (VXGI) 技术首次让游戏 GPU 能够提供实时的动态全局光照效果。基于 Maxwell 架构的 GTX 980 和 970 GPU 采用了包括多帧采样抗锯齿 (MFAA)、动态超级分辨率 (DSR)、VR Direct 以及超节能设计在内的一系列新技术。
  • 2016 - Pascal ：将处理器和数据集成在同一个程序包内，以实现更高的计算效率。1080系列、1060系列都是基于Pascal架构。
  • 2017 - Volta：配备640 个Tensor 核心，每秒可提供超过100 兆次浮点运算(TFLOPS) 的深度学习效能，比前一代的Pascal 架构快5 倍以上。

除了图上所示的构架外，后续还有

  • 2018 - Turing ： 配备了名为 RT Core 的专用光线追踪处理器，能够以高达每秒 10 Giga Rays 的速度对光线和声音在 3D 环境中的传播进行加速计算。Turing 架构将实时光线追踪运算加速至上一代 NVIDIA Pascal™ 架构的 25 倍，并能以高出 CPU 30 多倍的速度进行电影效果的最终帧渲染。2060系列、2080系列显卡也是跳过了Volta直接选择了Turing架构。
  • 2020 - Ampere ： 同时配备第二代RT Core和第三代Tensor Core，将光线相交的处理性能提升了一倍，在渲染有动态模糊的影像时，比Turing快8倍。

 

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、系统内存交互的各种部件。

 

Fermi架构

 

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

 

Maxwell架构

 

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

 

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核（NVIDIA推出的统一计算架构）
  • 每个SM包含8 Tensor核（专为执行张量或矩阵运算而设计的专用执行单元）
  • 每个SM还包含256 KB寄存器文件

 

 

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的架构皆是以高度并行能力而设计的。

　　GPC包含着多组TPC，TPC包含多组SM，SM又包含着多组CORE。一组SM中可能包含了Poly Morph Engine（多边形引擎）、L1 Cache（L1缓存）、Shared Memory（共享内存）、Core（执行数学运算的核心）等组件。而一组CORE中又包含了ALU、FPU、Execution Context（执行上下文）、（Detch）、解码（Decode）等组件。

　　从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逻辑管线

 

  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的混合，深度和颜色的设置必须是原子操作，否则两个不同的三角形在同一个像素点就会有冲突和错误。

 

 

一些GPU技术

 

Early-Z

 

　　早期GPU的渲染管线的深度测试是在像素着色器之后才执行，这样会造成很多本不可见的像素执行了耗性能的像素着色器计算（具体渲染流程可跳转至渲染流水线章节）。后来，为了减少像素着色器的额外消耗，将深度测试提至像素着色器之前（下图），这就是Early-Z技术的由来。Early-Z技术可以将很多无效的像素提前剔除，避免它们进入耗时严重的像素着色器。Early-Z剔除的最小单位不是1像素，而是像素块（2*2）。

　　但是，以下情况会导致Early-Z失效： 

  • 开启Alpha Test：由于Alpha Test需要在像素着色器后面的Alpha Test阶段比较（DX的discard，OpenGL的clip），所以无法在像素着色器之前就决定该像素是否被剔除。
  • 开启Alpha Blend：启用了Alpha混合的像素很多需要与frame buffer做混合，无法执行深度测试，也就无法利用Early-Z技术。
  • 关闭深度测试。Early-Z是建立在深度测试开启的条件下，如果关闭了深度测试，也就无法启用Early-Z技术。
  • 开启Multi-Sampling：多采样会影响周边像素，而Early-Z阶段无法得知周边像素是否被裁剪，故无法提前剔除。
  • 以及其它任何导致需要混合后面颜色的操作。

 

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
延迟容忍度 低 高
并行目标 任务（Task） 数据（Data）
核心构架 多线程核心 SIMT核心
线程数量级别 10 10000
吞吐量 低 高
缓存需求量 高 低
线程独立性 低 高

 

 

 

 

 

 

 

 

 

 

CPU-GPU异构系统

 

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

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

　　二是耦合式架构（Couple），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资源管理模型（分离式架构）

 

 

  • MMIO（Memory Mapped IO）

    • CPU与GPU的交流就是通过MMIO进行的。CPU 通过 MMIO 访问 GPU 的寄存器状态。
    • DMA传输大量的数据就是通过MMIO进行命令控制的。
    • I/O端口可用于间接访问MMIO区域，像Nouveau等开源软件从来不访问它。

  • GPU Context

    • GPU Context代表了GPU计算的状态。
    • 在GPU中拥有自己的虚拟地址。
    • GPU 中可以并存多个活跃态下的Context。

  • GPU Channel

    • 任何命令都是由CPU发出。
    • 命令流（command stream）被提交到硬件单元，也就是GPU Channel。
    • 每个GPU Channel关联一个context，而一个GPU Context可以有多个GPU channel。
    • 每个GPU Context 包含相关channel的 GPU Channel Descriptors ， 每个 Descriptor 都是 GPU 内存中的一个对象。
    • 每个 GPU Channel Descriptor 存储了 Channel 的设置，其中就包括 Page Table 。
    • 每个 GPU Channel 在GPU内存中分配了唯一的命令缓存，这通过MMIO对CPU可见。
    • GPU Context Switching 和命令执行都在GPU硬件内部调度。

  • GPU Page Table

    • GPU Context在虚拟基地空间由Page Table隔离其它的Context 。
    • GPU Page Table隔离CPU Page Table，位于GPU内存中。
    • GPU Page Table的物理地址位于 GPU Channel Descriptor中。
    • GPU Page Table不仅仅将 GPU虚拟地址转换成GPU内存的物理地址，也可以转换成CPU的物理地址。因此，GPU Page Table可以将GPU虚拟地址和CPU内存地址统一到GPU统一虚拟地址空间来。

  • PCI-e BAR

    • GPU 设备通过PCI-e总线接入到主机上。 Base Address Registers(BARs) 是 MMIO的窗口，在GPU启动时候配置。
    • GPU的控制寄存器和内存都映射到了BARs中。
    • GPU设备内存通过映射的MMIO窗口去配置GPU和访问GPU内存。

  • PFIFO Engine

    • PFIFO是GPU命令提交通过的一个特殊的部件。
    • PFIFO维护了一些独立命令队列，也就是Channel。
    • 此命令队列是Ring Buffer，有PUT和GET的指针。
    • 所有访问Channel控制区域的执行指令都被PFIFO 拦截下来。
    • GPU驱动使用Channel Descriptor来存储相关的Channel设定。
    • PFIFO将读取的命令转交给PGRAPH Engine。

  • BO

    • Buffer Object (BO)，内存的一块(Block)，能够用于存储纹理（Texture）、渲染目标（Render Target）、着色代码（shader code）等等。

    • Nouveau和Gdev经常使用BO。

Nouveau是一个自由及开放源代码显卡驱动程序，是为NVidia的显卡所编写。

Gdev是一套丰富的开源软件，用于NVIDIA的GPGPU技术，包括设备驱动程序。

 

CPU-GPU数据流

 

　　下图是分离式架构的CPU-GPU的数据流程图：

 

  • 将主存的处理数据复制到显存中。
  • CPU指令驱动GPU。
  • GPU中的每个运算单元并行处理。此步会从显存存取数据。
  • GPU将显存结果传回主存。

 

 

Shader运行机制

 

　　Shader代码也跟传统的C++等语言类似，需要将面向人类的高级语言（GLSL、HLSL、CGSL）通过编译器转成面向机器的二进制指令，二进制指令可转译成汇编代码，以便技术人员查阅和调试。由高级语言编译成汇编指令的过程通常是在离线阶段执行，以减轻运行时的消耗。

　　在执行阶段，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）。

　　

 

漫反射例子说明

 

sampler mySamp;
Texture2D<float3> myTex;
float3 lightDir;
 
float4 diffuseShader(float3 norm, float2 uv)
{
    float3 kd;
    kd = myTex.Sample(mySamp, uv);
    kd *= clamp( dot(lightDir, norm), 0.0, 1.0);
    return float4(kd, 1.0);
}

经过编译后成为汇编代码：

<diffuseShader>:
sample r0, v4, t0, s0
mul    r3, v0, cb0[0]
madd   r3, v1, cb0[1], r3
madd   r3, v2, cb0[2], r3
clmp   r3, r3, l(0.0), l(1.0)
mul    o0, r0, r3
mul    o1, r1, r3
mul    o2, r2, r3
mov    o3, l(1.0)

 

　　在执行阶段，汇编代码会被GPU推送到执行上下文（Execution Context），然后ALU会逐条获取（Detch）、解码（Decode）汇编指令为二进制指令，并执行它们。

　　而对于SIMT架构的GPU，汇编指令有所不同，变成了SIMT特定指令代码：

<VEC8_diffuseShader>: 
VEC8_sample vec_r0, vec_v4, t0, vec_s0 
VEC8_mul    vec_r3, vec_v0, cb0[0] 
VEC8_madd   vec_r3, vec_v1, cb0[1], vec_r3 
VEC8_madd   vec_r3, vec_v2, cb0[2], vec_r3
VEC8_clmp   vec_r3, vec_r3, l(0.0), l(1.0) 
VEC8_mul    vec_o0, vec_r0, vec_r3 
VEC8_mul    vec_o1, vec_r1, vec_r3 
VEC8_mul    vec_o2, vec_r2, vec_r3 
VEC8_mov    o3, l(1.0)

 

　　并且Context以Core为单位组成共享的结构，同一个Core的多个ALU共享一组Context，如果有多个Core，就会有更多的ALU同时参与shader计算，每个Core执行的数据是不一样的，可能是顶点、图元、像素等任何数据：

 

GPU Context和延迟

 

　　由于SIMT技术的引入，导致很多同一个SM内的很多Core并不是独立的，当它们当中有部分Core需要访问到纹理、常量缓存和全局内存时，就会导致非常大的卡顿（Stall）。

　　如果有4组上下文（Context），它们共用同一组运算单元ALU。假设第一组Context需要访问缓存或内存，会导致2~3个周期的延迟，此时调度器会激活第二组Context以利用ALU。

　　当第二组Context访问缓存或内存又卡住，会依次激活第三、第四组Context，直到第一组Context恢复运行或所有都被激活。延迟的后果是每组Context的总体执行时间被拉长了，越多Context可用就越可以提升运算单元的吞吐量。

 

 

Geforce RTX 2060的扩展验证

 

　　NV shader thread group提供了OpenGL的扩展，可以查询GPU线程、Core、SM、Warp等硬件相关的属性。如果要开启次此扩展，需要满足以下条件：

  • OpenGL 4.3+；
  • GLSL 4.3+；
  • 支持OpenGL 4.3+的NV显卡；

 　　下面是具体的字段和代表的意义：

// 开启扩展
#extension GL_NV_shader_thread_group : require     (or enable)
 
WARP_SIZE_NV    // 单个线程束的线程数量
WARPS_PER_SM_NV // 单个SM的线程束数量
SM_COUNT_NV     // SM数量
 
uniform uint  gl_WarpSizeNV;    // 单个线程束的线程数量
uniform uint  gl_WarpsPerSMNV;  // 单个SM的线程束数量
uniform uint  gl_SMCountNV;     // SM数量
 
in uint  gl_WarpIDNV;       // 当前线程束id
in uint  gl_SMIDNV;         // 当前线程束所在的SM id，取值[0, gl_SMCountNV-1]
in uint  gl_ThreadInWarpNV; // 当前线程id，取值[0, gl_WarpSizeNV-1]
 
in uint  gl_ThreadEqMaskNV; // 是否等于当前线程id的位域掩码。
in uint  gl_ThreadGeMaskNV; // 是否大于等于当前线程id的位域掩码。
in uint  gl_ThreadGtMaskNV; // 是否大于当前线程id的位域掩码。
in uint  gl_ThreadLeMaskNV; // 是否小于等于当前线程id的位域掩码。
in uint  gl_ThreadLtMaskNV; // 是否小于当前线程id的位域掩码。
 
in bool  gl_HelperThreadNV; // 当前线程是否协助型线程。

 

　　利用以上字段，可以编写特殊shader代码转成颜色信息，可视化了顶点着色器、像素着色器的SM、Warp id，为我们查探GPU的工作机制和流程提供了途径，以便可视化窥探GPU的工作机制和流程。下面正式进入验证阶段，将以Geforce RTX 2060作为验证对象，加入扩展所需的代码，并修改颜色计算：

#version 430 core
#extension GL_NV_shader_thread_group : require
 
uniform uint  gl_WarpSizeNV;    // 单个线程束的线程数量
uniform uint  gl_WarpsPerSMNV;  // 单个SM的线程束数量
uniform uint  gl_SMCountNV;     // SM数量
 
in uint  gl_WarpIDNV;       // 当前线程束id
in uint  gl_SMIDNV;         // 当前线程所在的SM id，取值[0, gl_SMCountNV-1]
in uint  gl_ThreadInWarpNV; // 当前线程id，取值[0, gl_WarpSizeNV-1]
 
out vec4 FragColor;
 
void main()
{
    // SM id
    float lightness = gl_SMIDNV / gl_SMCountNV;
    FragColor = vec4(lightness);
}

由上面的代码渲染的画面如下：

 

从上面可分析出一些信息：

  • 画面共有32个亮度色阶，也就是Geforce RTX 2060有32个SM。
  • 单个SM每次渲染16x16为单位的像素块，也就是每个SM有256个Core。
  • SM之间不是顺序分配像素块，而是无序分配。
  • 不同三角形的接缝处出现断层，说明同一个像素块如果分属不同的三角形，就会分配到不同的SM进行处理。由此推断，相同面积的区域，如果所属的三角形越多，就会导致分配给SM的次数越多，消耗的渲染性能也越多。

 

接着修改片元着色器的颜色计算代码以显示Warp id：

// warp id
float lightness = gl_WarpIDNV / gl_WarpsPerSMNV;
FragColor = vec4(lightness);

 

得到如下画面：

 

由此可得出一些信息或推论：

  • 画面共有32个亮度色阶，也就是每个SM有32个Warp，每个Warp有8个Core。

  • 每个色块像素是4x8，由于每个Warp有8个Core，由此推断每个Core单次要处理2x2的最小单元像素块。

  • 也是无序分配像素块。

  • 三角形接缝处出现断层，同SM的推断一致。

再修改片元着色器的颜色计算代码以显示线程id：

// thread id
float lightness = gl_ThreadInWarpNV / gl_WarpSizeNV;
FragColor = vec4(lightness);

 

得到如下画面：

 

为了方便分析，用Photoshop对中间局部放大10倍，得到以下画面：

 

结合上面两幅图，也可以得出一些结论：

  • 相较SM、线程束，线程分布图比较规律。说明同一个Warp的线程分布是规律的。
  • 三角形接缝处出现紊乱，说明是不同的Warp造成了不同的线程。
  • 画面有32个色阶，说明单个Warp有32个线程。
  • 每个像素独占一个亮度色阶，与周边相邻像素都不同，说明每个线程只处理一个像素。

再次说明，以上画面和结论是基于Geforce RTX 2060，不同型号的GPU可能会不一样，得到的结果和推论也会有所不同。

 

 总结

 

　　通过前面介绍的逻辑管线层面和硬件执行层面，可以总结出：

  • 顶点着色器和像素着色都是在同一个单元中执行的（在原来的架构中vs和ps的确是分开的，后来nv把这个统一了）vs是按照三角形来并行处理的，ps是按照像素来并行处理的。
  • vs和ps中的数据是通过L1和L2缓存传递的。
  • warp和thread都是逻辑上的概念，sm和sp都是物理上的概念。线程数≠流处理器数。

 

扩展提问

 

1、GPU是如何与CPU协调工作的？

 

　　MMIO。CPU与GPU的交流就是通过MMIO进行的。CPU 通过 MMIO 访问 GPU 的寄存器状态。DMA传输大量的数据也是通过MMIO进行命令控制的。

在分离式架构中数据先从主存复制到显存中。CPU再向channel发动指令驱动GPU。

GPU中的每个运算单元并行处理，此步会从显存存取数据。最后将显存结果传回主存。

 

2、GPU也有缓存机制吗？有几层？它们的速度差异多少？

 

 如图所示，这个图其实表述的挺清楚了。这5层结构分别为：寄存器、L1缓存、L2缓存、GPU显存、系统显存。

 

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

 

3、GPU的渲染流程有哪些阶段？它们的功能分别是什么？

 

　　PC上会对每个三角形一次执行顶点着色器和片元着色器。移动平台更多为TBR/TBDR，会先把所有的三角形执行完顶点着色器，再执行片元着色器。

　　大体流程为：通过三角形索引取出数据，即Vertex Fetch。然后SM中以线程束(Warp)来调度处理顶点数据。warp完成所有指令后会被Viewport Transform模块处理，三角形被裁剪准备栅格化并决定将被分配到哪个光栅引擎。然后SM再分warp执行片元着色器。然后将数据移交给ROP(render output unit，渲染输入单元)中处理深度测试，和framebuffer的混合。

 

4、Early-Z技术是什么？发生在哪个阶段？这个阶段还会发生什么？会产生什么问题？如何解决？

 

　　Early-Z技术指将深度测试提前到片元着色前，提前将无用的像素提前剔除（剔除的不是1像素，而是2x2的像素块），避免大量无效的片元执行耗时严重的片元着色。

　　如果开启了AlphaTest活着AlphaBlend或者shader里有discard/clip指令以及关闭深度测试的情况就会使Early-Z失效。

　　Early-Z还会导致深度数据冲突，可以再写入深度值之前再次与frameBuffer的值做一次对比。

　　PC上可能会多执行一次Early-Z。但是TBR架构中会分阶段处理顶点着色和片元着色，对于移除看不见的片元其实更有优势，IOS就有HSR（Hidden Surface Removel）技术，安卓的高通芯片也有类似的技术，因此移动平台一般不用做Early-Z。

 

5、SIMD和SIMT是什么？它们的好处是什么？co-issue呢？

 

SIMD: Single Instrument Multiple Data，单指令多数据;

一条指令可以处理多维向量的数据，原来的多条指令用一条指令即可处理完。

SIMT: Single Instrument Multiple Threads，单指令多线程。

单个SM中的多个Core同时处理同一指令，并且每个Core存取的数据可以是不同的即a、b 、c的值可以不一样。

co-issue是为了尽可能充分利用SIMD，将低维向量合并成Vector4以提高ALU的利用率。

 

6、GPU是并行处理的么？若是，硬件层是如何设计和实现的？

 

是。

　　多个SM，每个SM有多个Warp，每个Warp又有多个Core，每个Core又有1个FPU和ALU等

有大量具有计算功能的Core和大量的线程调度得以掩盖IO延迟。

 

7、GPC、TPC、SM是什么？Warp又是什么？它们和Core、Thread之间的关系如何？

 

GPC: Graphics Processor Cluster，图形处理簇;

TPC: Texture Processor Cluster，纹理处理簇;

SM:Stream Multiprocessor，流多处理器;

Warp：线程束，GPU并行计算的最小粒度;

1个GPU可以有多个GPC，1个GPC可以有多个TPC，1个TPC可以有多个SM。

每个SM包含许多Core。他们由Warp Scheduler驱动，其Warp Scheduler管理一组32个Threads

 

8、顶点着色器（VS）和像素着色器（PS）可以是同一处理单元吗？为什么？

 

　　现在可，DirectX10引入了一种叫做统一着色器架构（Unified shader Architecture）的技术，Core（SP,流处理器）成了顶点处理单元和像素处理单元的统一。传统的顶点和像素分离渲染架构存在严重的资源分配不均的问题，两种单元渲染任务量不同，效率低下。而SP架构是统一结构，不再区分顶点和像素渲染，进行不同渲染任务时都能保证效率。

 

9、像素着色器（PS）的最小处理单位是1像素吗？为什么？会带来什么影响？

 

不是，是2x2的像素块。

在像素着色器中，会将相邻的四个像素作为不可分隔的一组，送入同一个SM内4个不同的Core。

能精简SM架构，减少硬件单元数量和尺寸，降低功耗提高效能比，同时也提供ddx ddy导数解决了mipmap等问题。

但可能会导致overdraw，即只有1个fragment需要绘制，但真正执行时也是4个fragment的消耗。对于三角形边缘的fragment，这种情况尤其明显。

 

10、Shader中的if、for等语句会降低渲染效率吗？为什么？

 

　　绝大多数情况会，由于SIMD的特性，每个ALU的数据不一样，导致if-else语句在某些ALU中执行的是true分支，有些ALU执行的是false分支，拉长了整个执行周期。for循环也是同理，for循环条件的不同也会导致线程的有效执行时间不同，可能最快的ALU执行完了，最慢的ALU才执行了1/N。但最快的那个得等最慢的执行完，才能继续执行下一组命令，整个warp的耗时都是以最长时间为准，造成了算力的浪费，降低了渲染效率。

　　那还有极少数情况是什么呢，就是32个线程都走到if或者else里面，当然，这很难发生。

 

11、如下图，渲染相同面积的图形，三角形数量少（左）的还是数量多（右）的效率更快？为什么？

 

三角形数量少的效率更快。

Vertex Fetch阶段中，更少的三角形意味着更少的vertex data。

而更多的三角形则可能意味着更多的overdraw。

即相同面积的区域，三角形越多消耗性能越多。

 

12、GPU Context是什么？有什么作用？

 

GPU Context代表了GPU计算的状态，含运行指令和数据状态等信息。

GPU能以Context将指令推送到各个Core，而多个Context可以轮流执行，当某个Context就比较大的卡顿时可以快速调度执行其他的Context，提高运算单元吞吐量，也提升了GPU算力的利用率。

 

13、造成渲染瓶颈的问题很可能有哪些？该如何避免或优化它们？

 

CPU、GPU交互效率限制

——合批，LOD，减少顶点数、三角形数，贴图优化，调整视锥，特效、动画避免在CPU端每帧修改提交数据。

 

overdraw

——确保Early-Z有效，确保HSR的效果，控制物体数量，尤其是数量多面积小的东西。

 

shader效率

——尽可能不适用分支循环，慎用像素裁剪之类可能会影响Early-Z或HSR的语句，减少消耗大的语句（如采样，复杂的数学函数）。

 

 

参考

 

图形学笔记(一) 底层知识背景

【技术美术百人计划】图形 2.7.2 GPU硬件架构概述

知乎-NVIDIA GPU 架构梳理-捏太阳

Data Transfer Matters for GPU Computing

OpenGL官网：NV extensions

 

跳转回百人合集

posted @ 2022-01-15 18:44  anesu  阅读(1736)  评论(0)  编辑  收藏  举报