圖形 2.9 GPU硬件架構概述

本文轉載自查看原文 2022-01-15 18:44 1188 TA記錄/ TA/ 圖形

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邏輯管線

程序通過圖形API(DX、GL、WEBGL)發出drawcall指令，指令會被推送到驅動程序，驅動會檢查指令的合法性，然后會把指令放到GPU可以讀取的Pushbuffer中。
經過一段時間或者顯式調用flush指令后，驅動程序把Pushbuffer的內容發送給GPU，GPU通過主機接口（Host Interface）接受這些命令，並通過前端（Front End）處理這些命令。
在圖元分配器(Primitive Distributor)中開始工作分配，處理indexbuffer中的頂點產生三角形分成批次(batches)，然后發送給多個GPCs。這一步的理解就是提交上來n個三角形，分配給這幾個GPC同時處理。
在GPC中，每個SM中的Poly Morph Engine負責通過三角形索引(triangle indices)取出三角形的數據(vertex data)，即圖中的Vertex Fetch模塊。
在獲取數據之后，在SM中以32個線程為一組的線程束(Warp)來調度，來開始處理頂點數據。
SM的warp調度器會按照順序分發指令給整個warp，單個warp中的線程會鎖步(lock-step)執行各自的指令，如果線程碰到不激活執行的情況也會被遮掩(be masked out)
warp中的指令可以被一次完成，也可能經過多次調度，例如通常SM中的LD/ST(加載存取)單元數量明顯少於基礎數學操作單元。
由於某些指令比其他指令需要更長的時間才能完成，特別是內存加載，warp調度器可能會簡單地切換到另一個沒有內存等待的warp，這是GPU如何克服內存讀取延遲的關鍵，只是簡單地切換活動線程組。
一旦warp完成了vertex-shader的所有指令，運算結果會被Viewport Transform模塊處理，三角形會被裁剪然后准備柵格化，GPU會使用L1和L2緩存來進行vertex-shader和pixel-shader的數據通信。
接下來這些三角形將被分割，再分配給多個GPC，三角形的范圍決定着它將被分配到哪個光柵引擎(raster engines)，每個raster engines覆蓋了多個屏幕上的tile，這等於把三角形的渲染分配到多個tile上面。也就是像素階段就把按三角形划分變成了按顯示的像素划分了。
SM上的Attribute Setup保證了從vertex-shader來的數據經過插值后是pixel-shade是可讀的。
GPC上的光柵引擎(raster engines)在它接收到的三角形上工作，來負責這些這些三角形的像素信息的生成（同時會處理背面剔除和Early-Z剔除）。
32個像素線程將被分成一組，或者說8個2x2的像素塊，這是在像素着色器上面的最小工作單元，在這個像素線程內，如果沒有被三角形覆蓋就會被遮掩，SM中的warp調度器會管理像素着色器的任務。
接下來的階段就和vertex-shader中的邏輯步驟完全一樣，但是變成了在像素着色器線程中執行。由於不耗費任何性能可以獲取一個像素內的值，導致鎖步執行非常便利，所有的線程可以保證所有的指令可以在同一點。
最后一步，現在像素着色器已經完成了顏色的計算還有深度值的計算，在這個點上，我們必須考慮三角形的原始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、如下圖，渲染相同面積的圖形，三角形數量少（左）的還是數量多（右）的效率更快？為什么？