作者 | leowwlwang

“你買的4090多少錢？”、“H100性能真厲害！” ，GPU的價格性能一直是大家樂于談論的話題，作者也經常可以在茶余飯后聽到這樣的討論。在熱火朝天地談論性能指標、價格以外，本著”知其然也要知其所以然“的道理，作者學習整理了GPU本身的工作原理，編程模型，架構設計，在這里將我的學習筆記整理成文與大家分享，希望與大家一起 “知其然也要知其所以然”。

一、引言

1. Why GPU?

為什么要使用GPU？很多同學的第一反應就是“快”，這當然沒錯。而一個更嚴謹的說法是，GPU兼顧了“通用性”與“高效性”，才使得其一步步成為高性能計算的首選。

針對計算性能，1974年Dennard等人提出了Dennard縮放比例定律（Dennard Scaling）。

Dennard 縮放比例定律 (Dennard Scaling) ：當晶體管特征尺寸縮小時，其功率密度保持恒定。具體表現為電壓隨特征尺寸線性下降，電流密度保持穩定，使得單位面積的功耗與晶體管尺寸成比例關系。

一言以蔽之：晶體管越小越省電。

推導到芯片設計領域：晶體管縮小，芯片能塞入的晶體管更多，同時保持整體能耗穩定，推動計算機性能持續提升。

在計算機發展的前四十年間，基于Dennard定律的晶體管微縮是提升性能的主要路徑。但在2005-2007年間，隨著晶體管進入納米尺度，量子隧穿效應引發的漏電流呈指數增長，閾值電壓難以繼續降低，最終導致該定律失效。此時，工藝微縮帶來的性能增益已無法抵消功耗的快速增長，著名的"功耗墻"問題開始顯現。

單純依靠縮小晶體管尺寸來提升性能的方法不再可行，部分工程師開始轉向專用硬件，即專門為了某種或某幾種計算設計的計算硬件，例如Google的TPU（Tensor Processing Unit，張量處理器），就是一款專為加速機器學習任務而設計的專用硬件。然而，專用計算硬件只能聚焦于某一類或者某幾類特定的計算任務，在處理其他任務時則可能力不從心。

而GPU則是向通用性演進的典型代表。雖然其最初設計目標是為圖形渲染加速，但高度并行的SIMT（單指令多線程）架構意外契合了通用計算的演進需求，無論是基于CUDA的深度學習訓練，還是通過OpenCL加速的流體仿真，都能通過高度并行獲得遠超CPU的計算性能。

2. GPU的“快”

(1) 為什么快？

高計算并發：與CPU相比，GPU將更大比例的芯片面積分配給流處理器（如NVIDIA的CUDA核心），相應地減少控制邏輯（control logic）所占的面積，從而在高并行負載下獲得更高的單位面積性能。

低內存延遲：內存訪問導致的延遲也是影響性能的一大因素，GPU通過在其每個核心上運行大量線程的方式，來應對并掩蓋因全局內存訪問導致的延遲。這種設計使得GPU即使在面臨較慢的內存訪問時，也能維持高效的計算性能。具體來說，每個SIMT核心同時管理多組線程（多個warp，一個warp 32個線程），當某個warp因為等待內存數據而暫停時，GPU可以迅速切換到另一個warp繼續執行。這種快速切換使得GPU能夠在等待內存數據返回的同時，保持高利用率，從而有效地“隱藏”了內存訪問延遲。

特化內存與計算架構：GPU通常配備高帶寬的顯存（如GDDR6或HBM），能夠快速讀取和寫入數據。如NVIDIA A100使用HBM2e顯存最高可達到1.6TB/s帶寬，是普通DDR5內存（51.2GB/s）的31倍。計算架構方面，GPU集成專用計算單元實現硬件級加速，例如，NVIDIA的Tensor核心針對結構化稀疏計算做專門設計，在低精度損失的情況下，可以極大得提升計算性能。

(2) 有多快？

理論算力計算：GPU算力常以FLOPS（Floating-Point Operations Per Second，每秒浮點運算次數）來表示，通常數量級為T(萬億)，也即是大家聽到的TFLOPS。最常見的計算方式為CUDA核心計算法。

# CUDA核心計算法
算力（FLOPS）= CUDA核心數 × 加速頻率 × 每核心單個周期浮點計算系數
# 以A100為例
A100的算力（FP32單精度）= 6912（6912個CUDA核心） × 1.41（1.41GHz頻率） × 2（單周期2個浮點計算） = 19491.84 GFLOPS ≈ 19.5 TFLOPS

實測性能評估：通過計算只能得到紙面上的理論算力，如果同學們手上真的有GPU，那么實測性能評估則可以直接讓你獲取你的GPU的性能。此處為大家提供幾種最常見的實測方式和思路。

首先推薦一個非常實用的工具 GPU-Z，它是一款免費工具，可提供計算機中顯卡的詳細參數信息，支持實時監控 GPU 負載、溫度、顯存使用情況等關鍵數據，是排查顯卡性能問題或計算故障的實用診斷工具。GPU-Z是監控工具，而3DMark則是最流行的性能測試工具，通過模擬高負載游戲場景評估電腦圖形處理能力（在steam平臺即可購買，電腦上有GPU的同學不妨買來跑個分試試）。

最后再介紹一下GEMM（General Matrix Multiplication，通用矩陣乘法），這是一種經典的并行計算領域的計算密集型應用，與跑分工具這樣的封裝好的峰值性能測試工具相比，GEMM的重點反而不是進行性能測試，而是不斷調整優化逼近理論峰值的過程。GEMM通過執行時間 T 和總操作數（M×K與K×N的兩矩陣相乘）計算實測算力：

算力 = 總操作數 / 執行時間 = A(M, K) × B(K, N)/ T =  2 × M × N × K / T

如果實測算力低于GPU理論峰值算力，則表明可能存在低效內存訪問、計算資源利用率低、未充分利用硬件加速單元等問題，這些問題均可通過逐步優化來解決，以逼近理論峰值，當然也有溫度/功耗問題和顯存帶寬瓶頸等硬問題，但影響較小。對實際操作進行GPU編程有興趣的同學可以選擇深入了解GEMM，學習實現的比較好的GEMM庫是如何優化以逼近理論峰值的，在這個過程中深入理解GPU計算和編程。

3. GPU架構概述

在這里作者要做一個簡單的說明，現代的GPU架構，先不論不同廠家，僅NVIDIA一家就有數十年的架構迭代史，其中涉及的各種優化改進，限于篇幅，本文不可能一一介紹。但是，要想完整了解整個GPU架構的發展，作者認為可以分兩步走：以NVIDIA為例，就是“從0到Fermi“，和”從Fermi到Blackwell“。Fermi架構是現代通用GPU架構的基石，其中許多核心設計思想傳承至今，而此后直到作者撰文的2025年最新的Blackwell架構，都可以看做在基礎上的一路迭代。本文介紹的重點為兩步走里的第一步，即講解現代通用GPU中的基石級的通用技術與設計，讀者邁好第一步，就可以以此為基礎廣泛探索。

第一張圖為Fermi架構圖（來自Fermi架構白皮書），完整的Fermi架構GPU由4個GPC組成（黃色框），每個GPC有4個流式多處理器SM (Streaming Multiprocessor, 紅色框），每個SM又有32個CUDA Core，此外還有L1、L2 Cache、共享內存、顯存等組件。而每個SM、每個CUDA Core的結構則可見第二張圖。這樣看還是過于復雜，為了更清晰的從原理上了解通用GPU機構，本文將根據以下的簡化通用GPU架構圖講解，介紹GPU架構使用的術語也將傾向于學術界常見的通用術語：

SIMT核心（SIMT Core）是GPU的核心計算單元，類似于CPU的多核集群，負責協調和管理大量線程的并行執行，對應NVIDIA 架構中的SM。SIMT（Single Instruction, Multiple Threads，單指令多線程），是GPU的核心執行模型，其本質是通過統一指令指揮多個線程并行處理不同數據。后文將做單獨講解。多個SIMT核心組成SIMT Core Cluster，對應NVIDIA的GPC，每個Cluster/GPC可以看做是一個可完整運作的mini GPU，而實際的GPU由多個GPC組成，也就是大家常說的“多核”。

在同一個SIMT核心內運行的線程可以通過共享內存（Shared Memory）來進行彼此通信同步，SIMT核心內還包含一級指令和數據緩存，用來減少與低級內存的交互次數從而提高性能。而SIMT Core Cluster之間通過Interconnection Network通信。

除SIMT核心外，另一重要部分是內存和內存管理，在圖中即簡化為Memory Partition和GDDR部分。Memory Partition部分管理顯存的訪問，跨SM的L2全局一致性緩存也位于此處。GDDR，即為大家常常提到的顯存，其是位于GPU芯片外部的專用內存，用于存儲圖形數據等，相比于CPU的普通內存通常針對訪問延遲和帶寬進行優化。

二、GPU編程

本章將介紹如何編寫程序使用GPU完成非圖形類的計算，介紹重點在于揭示GPU的通用編程模式，以及程序執行的流程，并非專門的GPU編程教學。

1. 程序如何執行？以SAXPY為例

SAXPY，即將向量X的元素乘以A，再加上向量Y。以下是用C語言實現的CPU計算SAXPY的代碼：

// SAXPY函數實現
void saxpy(int n, float a, float *x, float *y) {
    for (int i = 0; i < n; i++) {
        y[i] = a * x[i] + y[i];
    }
}

int main() {
    float a = 2.0;
    int n; // 向量長度
    float *x; // 向量x
    float *y; // 向量y
    // 此處省略內存分配、元素賦值、長度指定
    // ...
    // 調用SAXPY函數
    saxpy(n, a, x, y);

    return 0;
}

針對上述CPU計算代碼，將代碼改寫為使用CUDA編寫的在GPU上運行SAXPY：

__global__ void saxpy(int n, float a, float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

int main() {
    float a = 2.0;
    int n; // 向量長度
    float *hx; // host向量x
    float *hy; // host向量y
    // 此處省略內存分配、元素賦值、長度指定
       
    // GPU內存分配
    int vector_size = n * sizeof(float); // 向量數據大小
    float *dx; // device向量x
    float *dy; // device向量y
    cudaMalloc(&dx, vector_size);
    cudaMalloc(&dy, vector_size);
    
    // 將host向量內容拷貝到device向量
    cudaMemcpy(dx, hx, vector_size, cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, vector_size, cudaMemcpyHostToDevice);
    
    // 執行saxpy
    int t = 256; // 每個thread block的線程數
    int blocks_num = (n + t - 1) / t; // thread block數量
    saxpy<<<blocks_num, t>>>(n, a, dx, dy);
    
    // 將device向量y內容（計算結果）拷貝到host向量y
    cudaMemcpy(hy, dy, vector_size, cudaMemcpyDeviceToHost);
    
    // ... (剩余邏輯)
    
    return 0;
}

(1) 設備側與主機側

GPU編程的思維是將GPU當作CPU的協同外設使用，通常GPU自身無法獨立運行，需要CPU指定任務，分配數據，驅動運行。第一行的__global__關鍵字，表示這段函數是內核函數（kernel，注意與Linux內核無關），是交給GPU執行的，而main函數則無此標識，由CPU執行。通常，將交GPU執行的代碼部分稱為設備(device)代碼，而交給CPU執行的代碼部分稱為主機(host)代碼。host與device是CUDA編程慣用的風格，CPU稱為host側，而GPU稱為device側。

main函數中的cudaMalloc、cudaMemcpy，是CPU操作GPU內存的操作，在分離式GPU架構（也就是獨顯）中，CPU分配內存用于GPU計算，再將數據傳輸到分配的內存空間，然后在GPU上啟動內核函數。GPU執行的內核函數只能從分配的GPU內存空間讀取數據。代碼中的host向量對應CPU內存的數據，而device向量則代表GPU內存的數據。

值得一提的是，近年來統一內存（unified memory）在GPU的應用中逐漸流行，統一內存是指一種允許CPU和GPU共享同一段地址空間的內存架構，這種架構下可以實現CPU和GPU之間數據交換的自動化，開發者不需要手動管理數據在CPU到GPU之間的傳輸。

(2) 線程組織

完成內存分配和數據拷貝后，CPU觸發GPU執行saxpy內核函數。觸發時同時指定了執行內核函數的線程的組織形式。在CUDA編程中，線程以thread，thread block，grid的層級結構進行組織，如上圖所示：

• 線程（thread，綠色部分）：最基本的執行單元。線程包含獨立寄存器狀態和獨立程序計數器。
• 線程塊（thread block，黃色部分）：由多個線程組成的集合，支持一維、二維或三維結構。線程塊內的線程可以通過共享內存進行通信，線程塊之間無法通過共享內存通信，但可通過全局內存進行數據交互。
• Warp（藍色線框）：硬件底層概念，GPU實際運行時將32個線程組成一個warp，同一warp內的線程同步執行相同的指令。
• 線程塊與warp的關系：warp是底層概念，NVIDIA的warp固定包含32個線程，warp是線程硬件調度的最小粒度。線程塊是軟件概念，線程塊有多少個線程組成由代碼指定。在運行時，硬件會將線程塊中的線程32個為一組打包成多個warp進行調度，因此，線程塊里的線程數最好為32的整數倍，以避免為拼湊完整warp而自動分配無效線程造成資源浪費
• 網格（grid，總體）：網格是所有線程塊的集合，支持一維、二維或三維結構，覆蓋整個計算任務的運行范圍。

thread，thread block，grid，warp是NVIDIA的術語，而對于AMD，四者又有其獨特的稱呼，因為本文使用的例子為CUDA編程，GPU編程部分的講解也將使用NVIDIA的術語體系，下表為術語對照表：

區別于NVIDIA，AMD的一個wavefront由64個work item組成。線程塊有時也被稱為CTA（Co-operative Thread Array)。

代碼執行saxpy部分：

// 執行saxpy
int t = 256; // 每個thread block的線程數
int blocks_num = (n + t - 1) / t; // thread block數量
saxpy<<<blocks_num, t>>>(n, a, dx, dy);

此處指定線程塊為一維的，一個每個線程塊（thread block）有256個線程（thread）。又計算得到了線程塊的數量block_num，指定網格（grid）也為一維，一個網格中有block_num個線程塊。最后，用<<< >>>三個尖括號包含網格的線程塊數、線程塊的線程數，指定一個grid有block_num個線程塊，一個線程塊有256個線程。

(3) 線程塊數量的計算

一個線程塊由多少個線程組成可以指定，與此不同的是，線程塊本身的數量則是由計算規模決定的，這段代碼根據向量的長度計算了線程塊的數量：

int blocks_num = (n + t - 1) / t; // thread block數量

這樣計算的目的是保證線程數量足夠，即每一個計算單元都有一個線程負責計算。

例如，如果向量長度n=250，則block_num = (250 + 256 - 1) / 256 = 1，每個線程塊有256個線程，那么要保證每個向量元素有一個線程負責計算，1個線程塊就夠了。又例如，如果向量長度n=257，則block_num = (257 + 256 - 1) / 256 = 2，需要兩個線程塊才能提供足夠的線程，當然，本例子中的兩個線程塊足以提供512個線程，有很多線程實際上是閑置了。

總結上述計算方式，可以得到計算線程塊數量時最常見的向上取整編程范式：

// B：線程塊數，N：問題規模，T：線程塊內線程數
B = (N + T - 1) / T

(4) 指定線程執行內核函數指令

最后，我們來關注saxpy內核函數本身，main函數中分配的每個線程都會并發地執行這段代碼：

__global__ void saxpy(int n, float a, float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

此處為每個線程分配了一個其所屬的向量元素，然后驅動線程分別完成計算。

首先計算i，i為線程的編號，blockIdx是block在grid上的坐標，blockDim則是block本身的尺寸，threadIdx為thread在block上的坐標。此前提到我們的grid、block都是一維的，因此只需要取其X維度，因此block的編號就直接取blockIdx.x，而一個block有blockDim.x個線程，線程編號為threadIdx.x。

假設當前線程是第二個線程塊上的第10個線程，即第266個線程，則其index應為265：

i = blockIdx.x * blockDim.x + threadIdx.x = 1 * 256 + 9 = 265

得到線程編號i后，第3行判斷i是否落在[0,n]區間內，n為線程總數。如果為否，則該線程就是前面提到的多分配的閑置線程，不調度。而對于需要調度的線程，則根據自己的線程編號，讀取源向量不同位置的元素，執行計算，并將結果寫入結果向量的不同位置。這樣，我們就為不同線程安排了獨立的工作，讓他們并發地完成工作。

2. 多維線程組織結構

截止到這里我們提到的grid、thread_block都是一維的，實際可以支持一維、二維、三維，這里再舉一個三維的例子：

// 主機端調用代碼
void launch_kernel_3d() {
    // 三維數據尺寸
    int dimX = 64    int dimY = 32    int dimZ = 16;
    // 定義三維線程塊（Block）和網格（Grid）
    dim3 blockSize(8, 4, 4);  // 指定每個塊包含8x4x4=128個線程
    dim3 gridSize(
        (dimX + blockSize.x - 1) / blockSize.x, // X方向塊數
        (dimY + blockSize.y - 1) / blockSize.y, // Y方向塊數
        (dimZ + blockSize.z - 1) / blockSize.z  // Z方向塊數
    );
    // 啟動內核函數
    kernel_3d<<<gridSize, blockSize>>>(d_data, dimX, dimY, dimZ);
}

使用dim3（CUDA數據結構）來承載三維grid、thread block的尺寸。grid為三維，因此要計算X、Y、Z三個維度上thread_block的數量，仍套用前文提到的向上取整計算方法。而如果是三維的grid、block，其計算線程編號時就需要取X、Y、Z三個維度：

// 核函數定義（處理三維數據）
__global__ void kernel_3d(float* data, int dimX, int dimY, int dimZ) {
    // 計算三維索引
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;

    if (x < dimX && y < dimY && z < dimZ) {
        // 處理三維數據（例如：三維矩陣元素操作）
        int idx = x + y * dimX + z * dimX * dimY; // 線程編號
        data[idx] *= 2.0f; // 示例：每個元素翻倍
    }
}

3. SIMT

前文提到：SIMT（Single Instruction, Multiple Threads，單指令多線程），由NVIDIA提出，是現代通用GPU的核心執行模型，甚至可以說正是SIMT的出現，使得GPU從一種處理圖形計算的專用硬件，進化為處理各類計算的通用處理器。SIMT的本質是通過統一指令指揮多個線程并行處理不同數據，結合上述例子，此處展開講解。

SIMT本質上是一種并行計算的范式，要徹底理解SIMT，以及SIMT存在的意義，就必須從另一種更基礎的并行計算的范式——SIMD講起。因為SIMT是對SIMD進行“線程級抽象”得到的，或者說，SIMT是“基于Warp的SIMD”。

SIMD（Single Instruction Multiple Data，單指令多數據），即：在同一時刻向多個數據元素執行同樣的一條指令。SIMD范式常見的一種實現是CPU的向量化運算，將N份數據存儲在向量寄存器里，執行一條指令，同時作用于向量寄存器里的每個數據。可見SIMD，特別是向量化運算，是一種偏硬件底層的并行計算優化，而SIMT范式則是通過線程編程模型隱藏了底層SIMD的執行細節。

在向量化運算實現的SIMD，有N個這樣流程并發執行：“指令+操作數→結果”，而SIMT的設計思想，則將“指令+操作數”抽象成了“線程”，線程可以看做是打包了指令和操作數的一個執行單元：線程包含獨立寄存器狀態（操作數）和程序計數器（指令）。在軟件編程時，程序以線程為單位進行調度，編程者只需要關注安排多少線程執行哪些指令，而無需過多考慮底層細節。這使得編程模型更接近多線程CPU，降低開發者適配難度。

SAXPY例子中的內核函數，就是以SIMT模型進行編程的，安排所有線程執行相同的指令，但每個線程執行指令時的指令操作數均不同，這便是SIMT：

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
    // 每個線程都執行這條指令，每個線程讀取不同元素執行相同計算
    y[i] = a * x[i] + y[i];
}

而在各線程實際運行時，硬件層面便會回歸SIMD范式。繼續以SAXPY為例，實際執行時GPU硬件會將其組織為warp，warp中的每個線程基于唯一索引i，訪問不同的內存位置，以不同的數據執行相同的指令，這便是SIMD：

// 一個Warp中每個線程的執行流程（線程0-31）
//【指令 + 操作數 = 結果】的SIMD范式
i = 0 → y[0] = a*x[0] + y[0]  
i = 1 → y[1] = a*x[1] + y[1]
...
i = 31 → y[31]= a*x[31]+ y[31]

傳統的SIMD關注的是一條條指令本身的執行方式，而SIMT則將SIMD“包了一層”，底層實現SIMD，表面上提供線程級編程模型，讓編程者很大程度上可以從串行的角度思考，而屏蔽了很多并行角度的執行細節。

這種編程便利最好的體現就是在出現分支（如if-else）時：Warp執行每個Branch Path，執行某個path時，不在那個path上的線程閑置不執行，線程活躍狀態通過一個32位的bitmask標記，分支收斂時再對齊匯總到下一段指令等等。后文將對這一過程作詳細講解，而在這里讀者只需要理解到，如果只有底層SIMD，那么這一切復雜流程都要編程者自己思考+編排，而在SIMT編程模型下編程者只需要編寫分支代碼，把這些編排交給硬件底層即可。

4. 指令集與編譯

剛才我們講解了CUDA C語言編寫的SAXPY，到這里，只是到了高級語言層面，眾所周知，高級語言需要轉換為機器碼才能被機器執行，本節將簡單介紹CUDA C/C++的程序的編譯流程，以及CUDA的PTX、SASS指令集。

(1) 指令集：SASS、PTX

SASS（Streaming Assembly）是GPU的機器指令集，是實際在GPU上執行的指令。SASS指令集直接對應GPU架構（Maxwell、Pascal等），雖然不是嚴格的一一對應，但通常每個GPU架構有專屬的SASS指令集，因此需要針對特定架構進行編譯。

PTX（Parallel Thread Execution）是一種中間表示形式，位于高級GPU編程語言（如CUDA C/C++）和低級機器指令集（SASS）之間。PTX與GPU架構基本無耦合關系，它本質上是從SASS上抽象出來的一種更上層的軟件編程模型，PTX的存在保證了代碼的可移植性（同一份PTX分發到不同架構上轉為對應SASS）與向后兼容性（可將PTX代碼轉為最新GPU架構對應的SASS）。

PTX是開發者可編程的最底層級，而SASS層則是完全閉源的，這也是NVIDIA的“護城河”之一。

(2) 編譯流程

CUDA程序的編譯由NVCC（NVIDIA CUDA Compiler）完成。

首先，NVCC完成預處理；隨后分類代碼為設備代碼和主機代碼，NVCC驅動傳統的C/C++編譯器主機代碼的編譯和匯編；對于設備代碼，NVCC將其編譯針對某架構的SASS，編譯過程中涉及C --> PTX --> SASS的轉化，但通常不顯式表現出來，生成的PTX/SASS碼也會被直接嵌入最終的可執行文件。

運行期，GPU會優先查找可執行文件中是否有適合當前架構的SASS，如有則直接執行。若無，則GPU驅動（driver）會使用JIT（Just-In-Time）編譯手段，將PTX碼編譯為當前架構對應的SASS再執行（前提是可執行文件必須包含PTX）。

三、SIMT核心架構

前面兩章，我們主要從總體概述和軟件編程的角度了解了GPU。相信不少同學在校園課程中，曾學習過CPU的核心架構，我一直以為，在了解了底層硬件是如何運作之后，我們看待處理器/硬件的視角才會有本質上的轉變，從一個用戶（這是執行我代碼的黑盒）轉變為一個專業技術人員（這是中央處理器）。因此，我們將更進一步，從更偏硬件的視角進一步了解GPU架構。

1. 軟硬分界線

前文提到SIMT核心也就是NVIDIA的SM，也給出了來自Fermi白皮書的SM結構圖。但是，線程以Warp為單位在SM上執行，具體如何執行，執行的流程是什么，每個組件發揮什么作用，單單從結構體是看不出來的，因此我們需要引入SM的指令流水線結構圖來進行講解：

如圖所示，SIMT核心流水線從運行的處理階段可以分為SIMT前端和SIMD后端兩個部分：

• SIMT前端：主要負責指令的獲取、譯碼和發射、分支預測、以及線程的管理和調度。這部分設計的組件對應SM結構圖中的藍色、橙色部分（Warp Scheduler、Register File）。
• SIMD后端：主要負責完成計算。這部分設計的組件對應SM結構圖中的綠色部分（Core）。

SIMT前端與SIMD后端的劃分本質上是控制流與數據流的解耦，SIMT前端關注指令流/控制流，而SIMD后端關注單個指令執行/數據流。

SIMT前端在硬件運行時“落實”了程序對線程的調度：SIMT前端以warp為單位調度線程，其包含的指令緩存（I-Cache）、解碼器和程序計數器PC組件集中管理線程的指令流，并使用SIMT堆棧等技術實現線程間的條件分支獨立控制流。SIMD后端主要負責執行實際的計算任務。在SIMT前端確定了warp要執行的指令后，指令發射，SIMD后端負責高效地完成一條條指令。具體的數據計算單元ALU，以及存取計算數據的寄存器訪問（Operand Collector）、寄存器文件（Register File）、內存讀寫（Memory）位于此處。

說到這里，這么多組件、組件之間有各種配合，不少同學估計已經要繞暈了。下面本文如果平鋪直敘地直接深入一個個組件的細節，就會變得難以理解。因此，下面本文將采取一種“三步走”的講解策略，先構建一個能執行計算任務的“最小系統”流水線，然后逐步向其中添加優化與功能，最終經過三步，構建出上圖中完整的流水線架構。

2. 第一步：最小可用系統

如上圖，我們將SIMT內核的架構做了最大可能的簡化，構成了一個“最簡GPU”。這個最小可用系統由6部分構成，此6個組件相互配合，使得我們的最簡GPU可以做到最簡的指令執行功能：即順序執行每一條指令，一條指令執行完再執行下一條：

• Fetch：取指令
• Decode：指令解碼
• SIMT Stack：SIMT堆棧，管理線程束的分支執行狀態，下文講解
• Issue：指令發射
• ALU：算數邏輯單元，代表執行計算的組件
• MEM：存儲器訪問單元，代表對L1 Cache、共享內存等各層級內存訪問的管理。

其中1、2、4、5、6部分是在CPU上久而有之的“老面孔”了，本文不多做解釋。本節將重點介紹GPU獨有的“新面孔”：SIMT堆棧。

(1) 分支發散：哪些線程執行哪條指令？

在GPU并行計算的發展歷程中，SIMT堆棧是早期架構解決線程分支管理問題的核心機制。

現實中的計算任務常包含大量條件分支（if-else、循環等）。在遇到條件分支發散（Branch Divergence）當線程束內線程選擇不同執行路徑時，會產生線程發散（Thread Divergence）：

如上圖，起初有5個線程執行相同的指令，直到分支發散處，根據SIMT的特性：多線程執行相同指令，但每個線程有自己獨立的數據，假設此處是一個if-else，有不同數據的線程將得到不同的條件判斷結果，2個線程進入if分支，3個線程進入else分支，進入不同分支的線程執行的指令流自然不同。

此處便出現了線程發散，即同一warp內的線程要執行不同指令，單由于線程以warp為最小單位調度，同一時鐘周期內同一warp內的線程必須執行相同的指令，那么不同執行分支的線程就需要分開調度，例如一個時鐘周期調度該warp執行if分支（if分支的線程活躍），下個時鐘周期再調度該warp執行else分支的線程（else分支的線程活躍）。也就是說，以warp為單位調度不代表每次調度warp，其中全部32個線程都活躍，也可以只有部分線程活躍，其余線程閑置。

分支發散帶來的復雜性不僅是線程指令流的發散，還有調度順序。如上圖，if-else分支發散后，分支聚合，5個線程執行紅色部分，但依賴if和else分支線程的運行結果，那么就要求藍色部分和黃色部分先執行完，再執行紅色部分。

為解決分支發散時的線程調度，NVIDIA于2008年在Tesla架構中首次引入SIMT堆棧，并作為2010年Fermi架構的核心技術，其核心思想是：

• 路徑跟蹤：當線程束遇到分支時，通過堆棧記錄所有可能執行路徑的上下文（如程序計數器PC、活躍線程掩碼）。
• 串行化執行：依次調度warp中每個分支路徑上的線程，其他線程暫時閑置。
• 重新收斂：在所有路徑執行完畢后，恢復完整warp的并行執行。

(2) SIMT堆棧

為了介紹SIMT堆棧的工作原理，我們引入一個稍復雜一點的分支發散例子，如下圖中的左圖，是一個程序的分支流，其中有兩層嵌套的if-else。而下方右圖則用表格的形式展示了左圖程序執行過程中SIMT堆棧的情況：表格最下行為棧頂，三行分別為聚合點PC、下條指令PC和活躍掩碼（Active Mask）。

聚合點PC，即分支聚合點的指令指針，例如，對于B、F這一分支發散，其聚合點PC就是G。

下條指令PC，顧名思義，就是當前指令的下一條指令的PC，如A的下條指令PC為B。

活躍掩碼（Active Mask），代表了哪些線程執行這條指令，本例子中假設有4個線程，而活躍掩碼就有4位，每一位分別對應一個線程，這一位為0，則線程不執行這條指令，為1則執行，例如，指令B的活躍掩碼為1110，代表前三個線程執行B，而第四個線程執行else分支的F（因此F的活躍掩碼為0001）。

觀察執行A、B、C、D時的SIMT堆棧，可以得到SIMT堆棧的運行方式：在遇到分支發散時，先將分支聚合點壓入堆棧，隨后壓入各分支的指令，各分支指令執行完畢后，回到聚合點，執行聚合點的指令。

我們跟著例子走一遍：

• 執行指令A，發現有分支發散。此時先將分支聚合點G壓棧，再將兩分支F、B先后壓棧。
• 執行棧頂的B，發現又有分支發散。此時先將聚合點E壓棧，再將兩分支D、C壓棧。
• 執行棧頂的C、D，回到聚合點E。后續按彈棧順序，再執行F、G，完成執行。

通過以上調度策略，保證了存在依賴時的正確性，例如，如果執行E依賴執行A、B、C、D的執行結果，SIMT棧剛好保證了E在ABCD后執行。

至于壓入各分支時的壓棧順序，如壓入C、D時的順序，因為C、D二者之間不存在依賴關系，從正確性角度而言，CD或者DC順序都可以，此時通常從性能角度出發，優先壓入有更少線程執行的指令（線程少離棧頂遠，線程多離棧頂近），從而保證有更多線程執行的指令先彈棧執行，這樣做有助于盡量減少棧的層數，提高性能。

(3) SIMT堆棧的問題

盡管SIMT堆棧在早期GPU架構中實現了分支管理能力，但其設計本質上面臨多重硬件與效率瓶頸，難以適應現代計算任務（光線追蹤、AI訓練推理等）對復雜控制流的需求：

• 傳統方案依賴固定深度的硬件堆棧，每個線程束需獨立維護堆棧，導致寄存器占用率攀升。
• 堆棧通常只有4-8級最大深度，這就意味著如果程序控制流過于復雜，例如，在訓練Transformer模型時，自注意力機制可能觸發數十/上百層條件判斷，遠超堆棧容量。
• 每次分支發散時，硬件需執行壓棧，并在路徑切換時彈棧。例如，一個包含5層嵌套if-else的著色器，需至少10次堆棧操作（進入和退出各一次）。隨著程序變得復雜，此類操作越來越多，會造成顯著的流水線延遲。
• 最后，由于堆棧的嚴格后進先出（LIFO）特性要求分支路徑必須按嵌套順序執行，很容易造成負載失衡甚至死鎖。例如，在光線追蹤中，部分線程可能因等待材質紋理讀取而停滯，而其他線程已完成計算，但受限于堆棧順序無法提前推進。

(4) 獨立線程調度

在Volta之前的架構（如Pascal、Fermi）中，在分支線程調度上，由SIMT完成調度，而Warp作為基本調度單元，所有線程共享統一的PC和活動掩碼，當Warp內線程執行不同分支路徑時，需按路徑順序串行執行。例如：線程0-3執行分支A的指令，線程4-31執行分支B的指令，則必須排隊執行，一部分線程先執行分支A的指令，另一部分線程必須等待。

而從Volta架構開始，引入了獨立線程調度（Independent Thread Scheduling）。每個線程擁有獨立的程序計數器(PC)和執行狀態寄存器，允許同一Warp內的線程在不同分支路徑上并行執行指令流。但硬件層面仍以Warp為基本調度單元。

圖出處：NVIDIA Volta架構白皮書

(5) 無堆棧分支收斂

同時，也是從Volta架構開始，隨著獨立線程調度的引入，傳統SIMT堆棧被棄用，分支收斂機制也升級到了無堆棧分支重新收斂（Stackless Branch Reconvergence）機制，通過收斂屏障（Convergence Barriers）技術來低成本解決分支代碼執行調度問題，獨立線程調度為無堆棧分支重新收斂提供了硬件支持。

無堆棧收斂屏障機制的核心手段之一是屏障參與掩碼（Barrier Participation Mask）與線程狀態協同管理，其核心思想可以通過ADD和WAIT操作來展示：

• ADD（屏障初始化）：當Warp執行到分支發散處前，通過專用ADD指令，活躍線程將其標識位注冊到指定收斂屏障的32位掩碼中，標記參與該屏障的線程組。
• WAIT（屏障同步）：在預設的收斂點（如分支匯合處），硬件插入WAIT指令。到達此處的子線程組將線程狀態標記為“阻塞”，并更新屏障狀態寄存器。當所有參與線程均抵達屏障后，調度器才重新激活完整線程束。

為了便于理解，下面用一個圖表示一個簡單的的ADD，WAIT的例子：

另外，通過新增的syncwarp()函數，開發者也可手動指定分支后的同步點，強制線程在特定位置重新收斂。

相比于SIMT堆棧，收斂屏障只需要使用僅需位掩碼和狀態寄存器，對于一個Warp（32個線程），一個屏障只需要32bit（每個bit對應一個線程），操作成本和硬件資源占用均極低，且不會再有堆棧深度限制，可以支持任意深度的條件分支嵌套。這一設計使得現代GPU（如NVIDIA Volta+架構）在復雜控制流場景下仍能保持高吞吐量，成為實時光追、AI推理等應用的關鍵支撐。

3. 第二步：動態指令調度以提高并發

在第一步構建的最小可用系統中，采用的是“一條指令執行完再執行下一條”的最簡執行策略。前文提到過，GPU為了隱藏內存訪問的延遲，需要在內存訪問指令為執行完前，先分配warp去執行其他指令。這里的策略其實就是動態指令調度，根據指令依賴關系和執行單元可用性，動態決定指令發射順序。

但此處有一個重要條件，就是先分配執行的這個其他指令，不能依賴于未完成指令的結果，否則無法執行。因此，需要先判斷指令之間是否存在依賴關系，才能選擇出不依賴未完成指令的指令進行執行。為了分析指令之間的依賴關系，以支持亂序執行，第二步為我們的系統增加了I-Cache、I-Buffer和ScoreBoard三個組件，并且ALU和MEM又多了一個指向ScoreBoard的“回寫”操作。

I-Cache（指令緩存）、I-Buffer（指令緩沖區）：緩存從內存中讀取的指令，和解碼后的指令。此二者將一系列指令存放在一起，用于進行依賴分析，并在分析結束后快速讀取指令進行亂序執行。I-Cache和I-Buffer為指令依賴分析提供了數據，ScoreBoard（計分牌）則是實際執行依賴分析操作的組件。

GPU計分板的核心目標是檢測指令間的數據依賴關系（如RAW、WAR、WAW），并控制指令發射順序以避免沖突。數據依賴關系反映到硬件層面體現為對寄存器的讀寫依賴關系，因此，GPU的計分板被設計為一個bitmap，其記錄了每一條未完成指令的目標寄存器，即如果這條指令要寫寄存器R1，則將R1對應的bit置為1。在指令完成后，再將R1對應的bit寫回0。該流程如下圖所示：

由于寄存器是線程私有的，需要為每個線程分配足夠的寄存器，因此SIMT核心中的寄存器數量是很大的，即便做到一個寄存器只需要一個bit表示狀態，ScoreBoard也會變得過大。因此，實際設計中，每個warp維護一個自己的ScoreBoard，由于每個warp同一時間只能執行同一條指令，一條指令能訪問的寄存器也是有限的，因此每個warp的ScoreBoard有3-4bit即可，每一個bit稱為一個表項（entry）。

在判斷一條指令是否能執行時，將該指令的源/目標寄存器與其所屬warp的計分板表項做比較（計算AND），生成依賴位向量（Dependency Bit Vector）。如果依賴位向量有任何一位為1，則說明存在數據沖突（依賴），該指令不能執行，反之如果全部為0，則可以發射執行。

4. 第三步：提高并發指令的數據供給效率

(1) 并發指令數據訪問

寄存器是處理器內部的高速存儲單元，用于臨時存放指令執行過程中所需的操作數、中間結果和地址信息。在GPU中，每個SIMT核心都擁有獨立的寄存器文件（寄存器的集合體，本質上是一組寄存器組成的存儲陣列）。

第二步引入的計分板（ScoreBoard）機制，解決了時序維度上的數據依賴問題，從而支持發射無依賴指令進行延遲隱藏，除了時序上的復雜性，指令并行還會帶來空間上的復雜性，即大量并發指令同時嘗試訪問寄存器文件獲取指令數據，寄存器文件必須支持多warp并發訪問。

(2) 簡單粗暴：多端口寄存器文件

端口（port），是讀寫存儲單元的接口。每多一個端口，存儲單元就可以多支持一個并發讀寫操作，單端口的情況下，同時只能支持一個讀或一個寫，若一個讀操作與一個寫操作并發，則只能串行執行，而增加一個端口，稱為雙端口，則此時的一讀一寫就可以并發完成。

因此，為了支持大量warp并發訪問寄存器數據，一個簡單粗暴的做法是，為寄存器文件設計足夠多的端口，來容納所有并發讀寫操作。

盡管多端口設計在理論上可行，但其硬件代價呈指數級增長，包括導致芯片面積暴增，同時，動態功耗會隨端口數平方增長、高訪問延遲等。因此，簡單的硬件堆料是低效且不可取的。

(3) 單端口內存庫

寄存器文件與共享內存的并發訪問沖突，本質上源于一個根本矛盾：存儲單元的物理端口數量與程序所需的并發訪問量之間的不匹配。若將多端口設計比作“拓寬車道”，則單端口內存庫（Single-Ported Memory Bank）更像是“優化交通規則”——通過精細化調度，在有限硬件資源下挖掘最大效率。

在計算機存儲體系結構中，banking是一種將存儲體分成多個獨立的部分（bank），每個bank可以獨立訪問，從而提高并行訪問能力的技術。如圖所示，單端口內存庫將寄存器文件分成多個bank后，每個bank可以獨立進行讀寫操作，每個bank只有一個端口，如果同一時刻只有同一個線程訪問，則可以成功，但如果有一個以上線程并發訪問，則將產生訪問沖突。

如此以來，分為多個bank的寄存器文件，一定程度上模擬了多端口寄存器文件的行為，即支持了跨bank的并發讀寫操作。單端口內存庫也是GPU核心架構中最常見的片上存儲單元微架構，除寄存器外，其同樣應用于共享內存。

同時，為了進一步優化，有很多通過寄存器布局優化減少Bank沖突的機制。其中常見的有：

• 交錯寄存器布局（Interleaved Register Allocation）。讓不同的warp的同編號寄存器分配到不同bank上。如warp0的R0分配到bank0，而warp1的R0則分配到bank1。這種布局方式在warp均勻調度發射指令（常見調度模式，大量warp輪流執行）時可以有效地防止沖突。
• 動態Bank分配（Dynamic Bank Allocation）。根據指令的寄存器訪問模式，動態調整邏輯寄存器到物理Bank的映射關系，避免靜態固定映射導致的沖突。
• 編譯器驅動的寄存器分配優化。編譯器在代碼生成階段，通過智能分配寄存器，減少Bank沖突。
• 在以上基礎上，發展出了混合Bank設計（Hybrid Banking）。將寄存器文件劃分為不同特性的Bank子集，采取不同的布局分配機制，針對不同訪問模式優化。

(4) 還有沖突：Operand Collector

接前文的例子，不論是“拓寬車道”還是“優化交通規則”，總會有車道爭搶的問題，那么也就總是需要“路口紅綠燈”來居中協調。

針對單bank的并發操作還是會引發數據沖突，這時就需要引入Operand Collector（操作數收集器）進行指令的統一調度。Operand Collector是 GPU 流水線中的一個關鍵硬件模塊，負責在指令執行前收集所有必需的操作數（即寄存器或內存中的數據）。它的核心目標是解決寄存器文件（Register File）的 Bank 沖突問題，并通過動態調度最大化寄存器訪問的并行性，從而提升指令吞吐量。

當指令進入寄存器讀取階段（Register Read Stage）時，系統為其分配一個收集單元（Collector Unit），每個收集單元為一條指令服務，負責緩存該指令所需的所有源操作數（如 ADD R1, R2, R3 中的 R2 和 R3）。收集單元向寄存器文件發送讀請求，獲取源操作數。例如，指令 ADD R1, R2, R3 需要讀取 R2 和 R3。

當不同指令出現數據沖突時，Operand Collector將動態調度這些沖突的請求，將沖突請求分配到不同周期排隊執行。若進入排隊狀態，收集單元暫存已就緒的操作數，直到所有操作數準備完畢，指令拿到操作數發射執行。

四、總結

講到這里，已歷上萬字，我們從引言中“Dennard Scaling”的失效開始，引入GPU出現的背景，又介紹了GPU的通用性，以及高并發、低延遲保證的高計算速度。隨后，我們以最常見的CUDA為例，介紹了GPU編程的基礎，SIMT與SIMD，編譯鏈接的過程。最后，我們深入硬件層面，分為三步走，先用最簡系統“run起來”，然后分別解決了指令依賴問題，以及并發執行中的數據訪問沖突問題，構建并了解了一個通用GPU核心的架構。

本文介紹的GPU知識，只是對各廠商、各架構設計做“求同存異”后，得到的主干性的、通用性的基礎知識，而GPU作為當代最為炙手可熱的科技產品之一，其發展是日新月異的。筆者希望這些基礎知識可以作為有興趣的讀者的“指路牌”，指引讀者在本文建立起來的基礎視野上，進一步探索。