作者 | leowwlwang

“你買(mǎi)的4090多少錢(qián)？”、“H100性能真厲害！” ，GPU的價(jià)格性能一直是大家樂(lè)于談?wù)摰脑掝}，作者也經(jīng)常可以在茶余飯后聽(tīng)到這樣的討論。在熱火朝天地談?wù)撔阅苤笜?biāo)、價(jià)格以外，本著”知其然也要知其所以然“的道理，作者學(xué)習(xí)整理了GPU本身的工作原理，編程模型，架構(gòu)設(shè)計(jì)，在這里將我的學(xué)習(xí)筆記整理成文與大家分享，希望與大家一起 “知其然也要知其所以然”。

一、引言

1. Why GPU?

為什么要使用GPU？很多同學(xué)的第一反應(yīng)就是“快”，這當(dāng)然沒(méi)錯(cuò)。而一個(gè)更嚴(yán)謹(jǐn)?shù)恼f(shuō)法是，GPU兼顧了“通用性”與“高效性”，才使得其一步步成為高性能計(jì)算的首選。

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

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

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

推導(dǎo)到芯片設(shè)計(jì)領(lǐng)域：晶體管縮小，芯片能塞入的晶體管更多，同時(shí)保持整體能耗穩(wěn)定，推動(dòng)計(jì)算機(jī)性能持續(xù)提升。

在計(jì)算機(jī)發(fā)展的前四十年間，基于Dennard定律的晶體管微縮是提升性能的主要路徑。但在2005-2007年間，隨著晶體管進(jìn)入納米尺度，量子隧穿效應(yīng)引發(fā)的漏電流呈指數(shù)增長(zhǎng)，閾值電壓難以繼續(xù)降低，最終導(dǎo)致該定律失效。此時(shí)，工藝微縮帶來(lái)的性能增益已無(wú)法抵消功耗的快速增長(zhǎng)，著名的"功耗墻"問(wèn)題開(kāi)始顯現(xiàn)。

單純依靠縮小晶體管尺寸來(lái)提升性能的方法不再可行，部分工程師開(kāi)始轉(zhuǎn)向?qū)Ｓ糜布?，即?zhuān)門(mén)為了某種或某幾種計(jì)算設(shè)計(jì)的計(jì)算硬件，例如Google的TPU（Tensor Processing Unit，張量處理器），就是一款專(zhuān)為加速機(jī)器學(xué)習(xí)任務(wù)而設(shè)計(jì)的專(zhuān)用硬件。然而，專(zhuān)用計(jì)算硬件只能聚焦于某一類(lèi)或者某幾類(lèi)特定的計(jì)算任務(wù)，在處理其他任務(wù)時(shí)則可能力不從心。

而GPU則是向通用性演進(jìn)的典型代表。雖然其最初設(shè)計(jì)目標(biāo)是為圖形渲染加速，但高度并行的SIMT（單指令多線程）架構(gòu)意外契合了通用計(jì)算的演進(jìn)需求，無(wú)論是基于CUDA的深度學(xué)習(xí)訓(xùn)練，還是通過(guò)OpenCL加速的流體仿真，都能通過(guò)高度并行獲得遠(yuǎn)超CPU的計(jì)算性能。

2. GPU的“快”

(1) 為什么快？

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

低內(nèi)存延遲：內(nèi)存訪問(wèn)導(dǎo)致的延遲也是影響性能的一大因素，GPU通過(guò)在其每個(gè)核心上運(yùn)行大量線程的方式，來(lái)應(yīng)對(duì)并掩蓋因全局內(nèi)存訪問(wèn)導(dǎo)致的延遲。這種設(shè)計(jì)使得GPU即使在面臨較慢的內(nèi)存訪問(wèn)時(shí)，也能維持高效的計(jì)算性能。具體來(lái)說(shuō)，每個(gè)SIMT核心同時(shí)管理多組線程（多個(gè)warp，一個(gè)warp 32個(gè)線程），當(dāng)某個(gè)warp因?yàn)榈却齼?nèi)存數(shù)據(jù)而暫停時(shí)，GPU可以迅速切換到另一個(gè)warp繼續(xù)執(zhí)行。這種快速切換使得GPU能夠在等待內(nèi)存數(shù)據(jù)返回的同時(shí)，保持高利用率，從而有效地“隱藏”了內(nèi)存訪問(wèn)延遲。

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

(2) 有多快？

理論算力計(jì)算：GPU算力常以FLOPS（Floating-Point Operations Per Second，每秒浮點(diǎn)運(yùn)算次數(shù)）來(lái)表示，通常數(shù)量級(jí)為T(mén)(萬(wàn)億)，也即是大家聽(tīng)到的TFLOPS。最常見(jiàn)的計(jì)算方式為CUDA核心計(jì)算法。

# CUDA核心計(jì)算法
算力（FLOPS）= CUDA核心數(shù) × 加速頻率 × 每核心單個(gè)周期浮點(diǎn)計(jì)算系數(shù)
# 以A100為例
A100的算力（FP32單精度）= 6912（6912個(gè)CUDA核心） × 1.41（1.41GHz頻率） × 2（單周期2個(gè)浮點(diǎn)計(jì)算） = 19491.84 GFLOPS ≈ 19.5 TFLOPS

實(shí)測(cè)性能評(píng)估：通過(guò)計(jì)算只能得到紙面上的理論算力，如果同學(xué)們手上真的有GPU，那么實(shí)測(cè)性能評(píng)估則可以直接讓你獲取你的GPU的性能。此處為大家提供幾種最常見(jiàn)的實(shí)測(cè)方式和思路。

首先推薦一個(gè)非常實(shí)用的工具 GPU-Z，它是一款免費(fèi)工具，可提供計(jì)算機(jī)中顯卡的詳細(xì)參數(shù)信息，支持實(shí)時(shí)監(jiān)控 GPU 負(fù)載、溫度、顯存使用情況等關(guān)鍵數(shù)據(jù)，是排查顯卡性能問(wèn)題或計(jì)算故障的實(shí)用診斷工具。GPU-Z是監(jiān)控工具，而3DMark則是最流行的性能測(cè)試工具，通過(guò)模擬高負(fù)載游戲場(chǎng)景評(píng)估電腦圖形處理能力（在steam平臺(tái)即可購(gòu)買(mǎi)，電腦上有GPU的同學(xué)不妨買(mǎi)來(lái)跑個(gè)分試試）。

最后再介紹一下GEMM（General Matrix Multiplication，通用矩陣乘法），這是一種經(jīng)典的并行計(jì)算領(lǐng)域的計(jì)算密集型應(yīng)用，與跑分工具這樣的封裝好的峰值性能測(cè)試工具相比，GEMM的重點(diǎn)反而不是進(jìn)行性能測(cè)試，而是不斷調(diào)整優(yōu)化逼近理論峰值的過(guò)程。GEMM通過(guò)執(zhí)行時(shí)間 T 和總操作數(shù)（M×K與K×N的兩矩陣相乘）計(jì)算實(shí)測(cè)算力：

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

如果實(shí)測(cè)算力低于GPU理論峰值算力，則表明可能存在低效內(nèi)存訪問(wèn)、計(jì)算資源利用率低、未充分利用硬件加速單元等問(wèn)題，這些問(wèn)題均可通過(guò)逐步優(yōu)化來(lái)解決，以逼近理論峰值，當(dāng)然也有溫度/功耗問(wèn)題和顯存帶寬瓶頸等硬問(wèn)題，但影響較小。對(duì)實(shí)際操作進(jìn)行GPU編程有興趣的同學(xué)可以選擇深入了解GEMM，學(xué)習(xí)實(shí)現(xiàn)的比較好的GEMM庫(kù)是如何優(yōu)化以逼近理論峰值的，在這個(gè)過(guò)程中深入理解GPU計(jì)算和編程。

3. GPU架構(gòu)概述

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

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

SIMT核心（SIMT Core）是GPU的核心計(jì)算單元，類(lèi)似于CPU的多核集群，負(fù)責(zé)協(xié)調(diào)和管理大量線程的并行執(zhí)行，對(duì)應(yīng)NVIDIA 架構(gòu)中的SM。SIMT（Single Instruction, Multiple Threads，單指令多線程），是GPU的核心執(zhí)行模型，其本質(zhì)是通過(guò)統(tǒng)一指令指揮多個(gè)線程并行處理不同數(shù)據(jù)。后文將做單獨(dú)講解。多個(gè)SIMT核心組成SIMT Core Cluster，對(duì)應(yīng)NVIDIA的GPC，每個(gè)Cluster/GPC可以看做是一個(gè)可完整運(yùn)作的mini GPU，而實(shí)際的GPU由多個(gè)GPC組成，也就是大家常說(shuō)的“多核”。

在同一個(gè)SIMT核心內(nèi)運(yùn)行的線程可以通過(guò)共享內(nèi)存（Shared Memory）來(lái)進(jìn)行彼此通信同步，SIMT核心內(nèi)還包含一級(jí)指令和數(shù)據(jù)緩存，用來(lái)減少與低級(jí)內(nèi)存的交互次數(shù)從而提高性能。而SIMT Core Cluster之間通過(guò)Interconnection Network通信。

除SIMT核心外，另一重要部分是內(nèi)存和內(nèi)存管理，在圖中即簡(jiǎn)化為Memory Partition和GDDR部分。Memory Partition部分管理顯存的訪問(wèn)，跨SM的L2全局一致性緩存也位于此處。GDDR，即為大家常常提到的顯存，其是位于GPU芯片外部的專(zhuān)用內(nèi)存，用于存儲(chǔ)圖形數(shù)據(jù)等，相比于CPU的普通內(nèi)存通常針對(duì)訪問(wèn)延遲和帶寬進(jìn)行優(yōu)化。

二、GPU編程

本章將介紹如何編寫(xiě)程序使用GPU完成非圖形類(lèi)的計(jì)算，介紹重點(diǎn)在于揭示GPU的通用編程模式，以及程序執(zhí)行的流程，并非專(zhuān)門(mén)的GPU編程教學(xué)。

1. 程序如何執(zhí)行？以SAXPY為例

SAXPY，即將向量X的元素乘以A，再加上向量Y。以下是用C語(yǔ)言實(shí)現(xiàn)的CPU計(jì)算SAXPY的代碼：

// SAXPY函數(shù)實(shí)現(xiàn)
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; // 向量長(zhǎng)度
    float *x; // 向量x
    float *y; // 向量y
    // 此處省略?xún)?nèi)存分配、元素賦值、長(zhǎng)度指定
    // ...
    // 調(diào)用SAXPY函數(shù)
    saxpy(n, a, x, y);

    return 0;
}

針對(duì)上述CPU計(jì)算代碼，將代碼改寫(xiě)為使用CUDA編寫(xiě)的在GPU上運(yùn)行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; // 向量長(zhǎng)度
    float *hx; // host向量x
    float *hy; // host向量y
    // 此處省略?xún)?nèi)存分配、元素賦值、長(zhǎng)度指定
       
    // GPU內(nèi)存分配
    int vector_size = n * sizeof(float); // 向量數(shù)據(jù)大小
    float *dx; // device向量x
    float *dy; // device向量y
    cudaMalloc(&dx, vector_size);
    cudaMalloc(&dy, vector_size);
    
    // 將host向量?jī)?nèi)容拷貝到device向量
    cudaMemcpy(dx, hx, vector_size, cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, vector_size, cudaMemcpyHostToDevice);
    
    // 執(zhí)行saxpy
    int t = 256; // 每個(gè)thread block的線程數(shù)
    int blocks_num = (n + t - 1) / t; // thread block數(shù)量
    saxpy<<<blocks_num, t>>>(n, a, dx, dy);
    
    // 將device向量y內(nèi)容（計(jì)算結(jié)果）拷貝到host向量y
    cudaMemcpy(hy, dy, vector_size, cudaMemcpyDeviceToHost);
    
    // ... (剩余邏輯)
    
    return 0;
}

(1) 設(shè)備側(cè)與主機(jī)側(cè)

GPU編程的思維是將GPU當(dāng)作CPU的協(xié)同外設(shè)使用，通常GPU自身無(wú)法獨(dú)立運(yùn)行，需要CPU指定任務(wù)，分配數(shù)據(jù)，驅(qū)動(dòng)運(yùn)行。第一行的__global__關(guān)鍵字，表示這段函數(shù)是內(nèi)核函數(shù)（kernel，注意與Linux內(nèi)核無(wú)關(guān)），是交給GPU執(zhí)行的，而main函數(shù)則無(wú)此標(biāo)識(shí)，由CPU執(zhí)行。通常，將交GPU執(zhí)行的代碼部分稱(chēng)為設(shè)備(device)代碼，而交給CPU執(zhí)行的代碼部分稱(chēng)為主機(jī)(host)代碼。host與device是CUDA編程慣用的風(fēng)格，CPU稱(chēng)為host側(cè)，而GPU稱(chēng)為device側(cè)。

main函數(shù)中的cudaMalloc、cudaMemcpy，是CPU操作GPU內(nèi)存的操作，在分離式GPU架構(gòu)（也就是獨(dú)顯）中，CPU分配內(nèi)存用于GPU計(jì)算，再將數(shù)據(jù)傳輸?shù)椒峙涞膬?nèi)存空間，然后在GPU上啟動(dòng)內(nèi)核函數(shù)。GPU執(zhí)行的內(nèi)核函數(shù)只能從分配的GPU內(nèi)存空間讀取數(shù)據(jù)。代碼中的host向量對(duì)應(yīng)CPU內(nèi)存的數(shù)據(jù)，而device向量則代表GPU內(nèi)存的數(shù)據(jù)。

值得一提的是，近年來(lái)統(tǒng)一內(nèi)存（unified memory）在GPU的應(yīng)用中逐漸流行，統(tǒng)一內(nèi)存是指一種允許CPU和GPU共享同一段地址空間的內(nèi)存架構(gòu)，這種架構(gòu)下可以實(shí)現(xiàn)CPU和GPU之間數(shù)據(jù)交換的自動(dòng)化，開(kāi)發(fā)者不需要手動(dòng)管理數(shù)據(jù)在CPU到GPU之間的傳輸。

(2) 線程組織

完成內(nèi)存分配和數(shù)據(jù)拷貝后，CPU觸發(fā)GPU執(zhí)行saxpy內(nèi)核函數(shù)。觸發(fā)時(shí)同時(shí)指定了執(zhí)行內(nèi)核函數(shù)的線程的組織形式。在CUDA編程中，線程以thread，thread block，grid的層級(jí)結(jié)構(gòu)進(jìn)行組織，如上圖所示：

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

thread，thread block，grid，warp是NVIDIA的術(shù)語(yǔ)，而對(duì)于AMD，四者又有其獨(dú)特的稱(chēng)呼，因?yàn)楸疚氖褂玫睦訛镃UDA編程，GPU編程部分的講解也將使用NVIDIA的術(shù)語(yǔ)體系，下表為術(shù)語(yǔ)對(duì)照表：

區(qū)別于NVIDIA，AMD的一個(gè)wavefront由64個(gè)work item組成。線程塊有時(shí)也被稱(chēng)為CTA（Co-operative Thread Array)。

代碼執(zhí)行saxpy部分：

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

此處指定線程塊為一維的，一個(gè)每個(gè)線程塊（thread block）有256個(gè)線程（thread）。又計(jì)算得到了線程塊的數(shù)量block_num，指定網(wǎng)格（grid）也為一維，一個(gè)網(wǎng)格中有block_num個(gè)線程塊。最后，用<<< >>>三個(gè)尖括號(hào)包含網(wǎng)格的線程塊數(shù)、線程塊的線程數(shù)，指定一個(gè)grid有block_num個(gè)線程塊，一個(gè)線程塊有256個(gè)線程。

(3) 線程塊數(shù)量的計(jì)算

一個(gè)線程塊由多少個(gè)線程組成可以指定，與此不同的是，線程塊本身的數(shù)量則是由計(jì)算規(guī)模決定的，這段代碼根據(jù)向量的長(zhǎng)度計(jì)算了線程塊的數(shù)量：

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

這樣計(jì)算的目的是保證線程數(shù)量足夠，即每一個(gè)計(jì)算單元都有一個(gè)線程負(fù)責(zé)計(jì)算。

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

總結(jié)上述計(jì)算方式，可以得到計(jì)算線程塊數(shù)量時(shí)最常見(jiàn)的向上取整編程范式：

// B：線程塊數(shù)，N：?jiǎn)栴}規(guī)模，T：線程塊內(nèi)線程數(shù)
B = (N + T - 1) / T

(4) 指定線程執(zhí)行內(nèi)核函數(shù)指令

最后，我們來(lái)關(guān)注saxpy內(nèi)核函數(shù)本身，main函數(shù)中分配的每個(gè)線程都會(huì)并發(fā)地執(zhí)行這段代碼：

__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];
    }
}

此處為每個(gè)線程分配了一個(gè)其所屬的向量元素，然后驅(qū)動(dòng)線程分別完成計(jì)算。

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

假設(shè)當(dāng)前線程是第二個(gè)線程塊上的第10個(gè)線程，即第266個(gè)線程，則其index應(yīng)為265：

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

得到線程編號(hào)i后，第3行判斷i是否落在[0,n]區(qū)間內(nèi)，n為線程總數(shù)。如果為否，則該線程就是前面提到的多分配的閑置線程，不調(diào)度。而對(duì)于需要調(diào)度的線程，則根據(jù)自己的線程編號(hào)，讀取源向量不同位置的元素，執(zhí)行計(jì)算，并將結(jié)果寫(xiě)入結(jié)果向量的不同位置。這樣，我們就為不同線程安排了獨(dú)立的工作，讓他們并發(fā)地完成工作。

2. 多維線程組織結(jié)構(gòu)

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

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

使用dim3（CUDA數(shù)據(jù)結(jié)構(gòu)）來(lái)承載三維grid、thread block的尺寸。grid為三維，因此要計(jì)算X、Y、Z三個(gè)維度上thread_block的數(shù)量，仍套用前文提到的向上取整計(jì)算方法。而如果是三維的grid、block，其計(jì)算線程編號(hào)時(shí)就需要取X、Y、Z三個(gè)維度：

// 核函數(shù)定義（處理三維數(shù)據(jù)）
__global__ void kernel_3d(float* data, int dimX, int dimY, int dimZ) {
    // 計(jì)算三維索引
    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) {
        // 處理三維數(shù)據(jù)（例如：三維矩陣元素操作）
        int idx = x + y * dimX + z * dimX * dimY; // 線程編號(hào)
        data[idx] *= 2.0f; // 示例：每個(gè)元素翻倍
    }
}

3. SIMT

前文提到：SIMT（Single Instruction, Multiple Threads，單指令多線程），由NVIDIA提出，是現(xiàn)代通用GPU的核心執(zhí)行模型，甚至可以說(shuō)正是SIMT的出現(xiàn)，使得GPU從一種處理圖形計(jì)算的專(zhuān)用硬件，進(jìn)化為處理各類(lèi)計(jì)算的通用處理器。SIMT的本質(zhì)是通過(guò)統(tǒng)一指令指揮多個(gè)線程并行處理不同數(shù)據(jù)，結(jié)合上述例子，此處展開(kāi)講解。

SIMT本質(zhì)上是一種并行計(jì)算的范式，要徹底理解SIMT，以及SIMT存在的意義，就必須從另一種更基礎(chǔ)的并行計(jì)算的范式——SIMD講起。因?yàn)镾IMT是對(duì)SIMD進(jìn)行“線程級(jí)抽象”得到的，或者說(shuō)，SIMT是“基于Warp的SIMD”。

SIMD（Single Instruction Multiple Data，單指令多數(shù)據(jù)），即：在同一時(shí)刻向多個(gè)數(shù)據(jù)元素執(zhí)行同樣的一條指令。SIMD范式常見(jiàn)的一種實(shí)現(xiàn)是CPU的向量化運(yùn)算，將N份數(shù)據(jù)存儲(chǔ)在向量寄存器里，執(zhí)行一條指令，同時(shí)作用于向量寄存器里的每個(gè)數(shù)據(jù)。可見(jiàn)SIMD，特別是向量化運(yùn)算，是一種偏硬件底層的并行計(jì)算優(yōu)化，而SIMT范式則是通過(guò)線程編程模型隱藏了底層SIMD的執(zhí)行細(xì)節(jié)。

在向量化運(yùn)算實(shí)現(xiàn)的SIMD，有N個(gè)這樣流程并發(fā)執(zhí)行：“指令+操作數(shù)→結(jié)果”，而SIMT的設(shè)計(jì)思想，則將“指令+操作數(shù)”抽象成了“線程”，線程可以看做是打包了指令和操作數(shù)的一個(gè)執(zhí)行單元：線程包含獨(dú)立寄存器狀態(tài)（操作數(shù)）和程序計(jì)數(shù)器（指令）。在軟件編程時(shí)，程序以線程為單位進(jìn)行調(diào)度，編程者只需要關(guān)注安排多少線程執(zhí)行哪些指令，而無(wú)需過(guò)多考慮底層細(xì)節(jié)。這使得編程模型更接近多線程CPU，降低開(kāi)發(fā)者適配難度。

SAXPY例子中的內(nèi)核函數(shù)，就是以SIMT模型進(jìn)行編程的，安排所有線程執(zhí)行相同的指令，但每個(gè)線程執(zhí)行指令時(shí)的指令操作數(shù)均不同，這便是SIMT：

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

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

// 一個(gè)Warp中每個(gè)線程的執(zhí)行流程（線程0-31）
//【指令 + 操作數(shù) = 結(jié)果】的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]

傳統(tǒng)的SIMD關(guān)注的是一條條指令本身的執(zhí)行方式，而SIMT則將SIMD“包了一層”，底層實(shí)現(xiàn)SIMD，表面上提供線程級(jí)編程模型，讓編程者很大程度上可以從串行的角度思考，而屏蔽了很多并行角度的執(zhí)行細(xì)節(jié)。

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

4. 指令集與編譯

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

(1) 指令集：SASS、PTX

SASS（Streaming Assembly）是GPU的機(jī)器指令集，是實(shí)際在GPU上執(zhí)行的指令。SASS指令集直接對(duì)應(yīng)GPU架構(gòu)（Maxwell、Pascal等），雖然不是嚴(yán)格的一一對(duì)應(yīng)，但通常每個(gè)GPU架構(gòu)有專(zhuān)屬的SASS指令集，因此需要針對(duì)特定架構(gòu)進(jìn)行編譯。

PTX（Parallel Thread Execution）是一種中間表示形式，位于高級(jí)GPU編程語(yǔ)言（如CUDA C/C++）和低級(jí)機(jī)器指令集（SASS）之間。PTX與GPU架構(gòu)基本無(wú)耦合關(guān)系，它本質(zhì)上是從SASS上抽象出來(lái)的一種更上層的軟件編程模型，PTX的存在保證了代碼的可移植性（同一份PTX分發(fā)到不同架構(gòu)上轉(zhuǎn)為對(duì)應(yīng)SASS）與向后兼容性（可將PTX代碼轉(zhuǎn)為最新GPU架構(gòu)對(duì)應(yīng)的SASS）。

PTX是開(kāi)發(fā)者可編程的最底層級(jí)，而SASS層則是完全閉源的，這也是NVIDIA的“護(hù)城河”之一。

(2) 編譯流程

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

首先，NVCC完成預(yù)處理；隨后分類(lèi)代碼為設(shè)備代碼和主機(jī)代碼，NVCC驅(qū)動(dòng)傳統(tǒng)的C/C++編譯器主機(jī)代碼的編譯和匯編；對(duì)于設(shè)備代碼，NVCC將其編譯針對(duì)某架構(gòu)的SASS，編譯過(guò)程中涉及C --> PTX --> SASS的轉(zhuǎn)化，但通常不顯式表現(xiàn)出來(lái)，生成的PTX/SASS碼也會(huì)被直接嵌入最終的可執(zhí)行文件。

運(yùn)行期，GPU會(huì)優(yōu)先查找可執(zhí)行文件中是否有適合當(dāng)前架構(gòu)的SASS，如有則直接執(zhí)行。若無(wú)，則GPU驅(qū)動(dòng)（driver）會(huì)使用JIT（Just-In-Time）編譯手段，將PTX碼編譯為當(dāng)前架構(gòu)對(duì)應(yīng)的SASS再執(zhí)行（前提是可執(zhí)行文件必須包含PTX）。

三、SIMT核心架構(gòu)

前面兩章，我們主要從總體概述和軟件編程的角度了解了GPU。相信不少同學(xué)在校園課程中，曾學(xué)習(xí)過(guò)CPU的核心架構(gòu)，我一直以為，在了解了底層硬件是如何運(yùn)作之后，我們看待處理器/硬件的視角才會(huì)有本質(zhì)上的轉(zhuǎn)變，從一個(gè)用戶(hù)（這是執(zhí)行我代碼的黑盒）轉(zhuǎn)變?yōu)橐粋€(gè)專(zhuān)業(yè)技術(shù)人員（這是中央處理器）。因此，我們將更進(jìn)一步，從更偏硬件的視角進(jìn)一步了解GPU架構(gòu)。

1. 軟硬分界線

前文提到SIMT核心也就是NVIDIA的SM，也給出了來(lái)自Fermi白皮書(shū)的SM結(jié)構(gòu)圖。但是，線程以Warp為單位在SM上執(zhí)行，具體如何執(zhí)行，執(zhí)行的流程是什么，每個(gè)組件發(fā)揮什么作用，單單從結(jié)構(gòu)體是看不出來(lái)的，因此我們需要引入SM的指令流水線結(jié)構(gòu)圖來(lái)進(jìn)行講解：

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

• SIMT前端：主要負(fù)責(zé)指令的獲取、譯碼和發(fā)射、分支預(yù)測(cè)、以及線程的管理和調(diào)度。這部分設(shè)計(jì)的組件對(duì)應(yīng)SM結(jié)構(gòu)圖中的藍(lán)色、橙色部分（Warp Scheduler、Register File）。
• SIMD后端：主要負(fù)責(zé)完成計(jì)算。這部分設(shè)計(jì)的組件對(duì)應(yīng)SM結(jié)構(gòu)圖中的綠色部分（Core）。

SIMT前端與SIMD后端的劃分本質(zhì)上是控制流與數(shù)據(jù)流的解耦，SIMT前端關(guān)注指令流/控制流，而SIMD后端關(guān)注單個(gè)指令執(zhí)行/數(shù)據(jù)流。

SIMT前端在硬件運(yùn)行時(shí)“落實(shí)”了程序?qū)€程的調(diào)度：SIMT前端以warp為單位調(diào)度線程，其包含的指令緩存（I-Cache）、解碼器和程序計(jì)數(shù)器PC組件集中管理線程的指令流，并使用SIMT堆棧等技術(shù)實(shí)現(xiàn)線程間的條件分支獨(dú)立控制流。SIMD后端主要負(fù)責(zé)執(zhí)行實(shí)際的計(jì)算任務(wù)。在SIMT前端確定了warp要執(zhí)行的指令后，指令發(fā)射，SIMD后端負(fù)責(zé)高效地完成一條條指令。具體的數(shù)據(jù)計(jì)算單元ALU，以及存取計(jì)算數(shù)據(jù)的寄存器訪問(wèn)（Operand Collector）、寄存器文件（Register File）、內(nèi)存讀寫(xiě)（Memory）位于此處。

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

2. 第一步：最小可用系統(tǒng)

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

• Fetch：取指令
• Decode：指令解碼
• SIMT Stack：SIMT堆棧，管理線程束的分支執(zhí)行狀態(tài)，下文講解
• Issue：指令發(fā)射
• ALU：算數(shù)邏輯單元，代表執(zhí)行計(jì)算的組件
• MEM：存儲(chǔ)器訪問(wèn)單元，代表對(duì)L1 Cache、共享內(nèi)存等各層級(jí)內(nèi)存訪問(wèn)的管理。

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

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

在GPU并行計(jì)算的發(fā)展歷程中，SIMT堆棧是早期架構(gòu)解決線程分支管理問(wèn)題的核心機(jī)制。

現(xiàn)實(shí)中的計(jì)算任務(wù)常包含大量條件分支（if-else、循環(huán)等）。在遇到條件分支發(fā)散（Branch Divergence）當(dāng)線程束內(nèi)線程選擇不同執(zhí)行路徑時(shí)，會(huì)產(chǎn)生線程發(fā)散（Thread Divergence）：

如上圖，起初有5個(gè)線程執(zhí)行相同的指令，直到分支發(fā)散處，根據(jù)SIMT的特性：多線程執(zhí)行相同指令，但每個(gè)線程有自己獨(dú)立的數(shù)據(jù)，假設(shè)此處是一個(gè)if-else，有不同數(shù)據(jù)的線程將得到不同的條件判斷結(jié)果，2個(gè)線程進(jìn)入if分支，3個(gè)線程進(jìn)入else分支，進(jìn)入不同分支的線程執(zhí)行的指令流自然不同。

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

分支發(fā)散帶來(lái)的復(fù)雜性不僅是線程指令流的發(fā)散，還有調(diào)度順序。如上圖，if-else分支發(fā)散后，分支聚合，5個(gè)線程執(zhí)行紅色部分，但依賴(lài)if和else分支線程的運(yùn)行結(jié)果，那么就要求藍(lán)色部分和黃色部分先執(zhí)行完，再執(zhí)行紅色部分。

為解決分支發(fā)散時(shí)的線程調(diào)度，NVIDIA于2008年在Tesla架構(gòu)中首次引入SIMT堆棧，并作為2010年Fermi架構(gòu)的核心技術(shù)，其核心思想是：

• 路徑跟蹤：當(dāng)線程束遇到分支時(shí)，通過(guò)堆棧記錄所有可能執(zhí)行路徑的上下文（如程序計(jì)數(shù)器PC、活躍線程掩碼）。
• 串行化執(zhí)行：依次調(diào)度warp中每個(gè)分支路徑上的線程，其他線程暫時(shí)閑置。
• 重新收斂：在所有路徑執(zhí)行完畢后，恢復(fù)完整warp的并行執(zhí)行。

(2) SIMT堆棧

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

聚合點(diǎn)PC，即分支聚合點(diǎn)的指令指針，例如，對(duì)于B、F這一分支發(fā)散，其聚合點(diǎn)PC就是G。

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

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

觀察執(zhí)行A、B、C、D時(shí)的SIMT堆棧，可以得到SIMT堆棧的運(yùn)行方式：在遇到分支發(fā)散時(shí)，先將分支聚合點(diǎn)壓入堆棧，隨后壓入各分支的指令，各分支指令執(zhí)行完畢后，回到聚合點(diǎn)，執(zhí)行聚合點(diǎn)的指令。

我們跟著例子走一遍：

• 執(zhí)行指令A(yù)，發(fā)現(xiàn)有分支發(fā)散。此時(shí)先將分支聚合點(diǎn)G壓棧，再將兩分支F、B先后壓棧。
• 執(zhí)行棧頂?shù)腂，發(fā)現(xiàn)又有分支發(fā)散。此時(shí)先將聚合點(diǎn)E壓棧，再將兩分支D、C壓棧。
• 執(zhí)行棧頂?shù)腃、D，回到聚合點(diǎn)E。后續(xù)按彈棧順序，再執(zhí)行F、G，完成執(zhí)行。

通過(guò)以上調(diào)度策略，保證了存在依賴(lài)時(shí)的正確性，例如，如果執(zhí)行E依賴(lài)執(zhí)行A、B、C、D的執(zhí)行結(jié)果，SIMT棧剛好保證了E在ABCD后執(zhí)行。

至于壓入各分支時(shí)的壓棧順序，如壓入C、D時(shí)的順序，因?yàn)镃、D二者之間不存在依賴(lài)關(guān)系，從正確性角度而言，CD或者DC順序都可以，此時(shí)通常從性能角度出發(fā)，優(yōu)先壓入有更少線程執(zhí)行的指令（線程少離棧頂遠(yuǎn)，線程多離棧頂近），從而保證有更多線程執(zhí)行的指令先彈棧執(zhí)行，這樣做有助于盡量減少棧的層數(shù)，提高性能。

(3) SIMT堆棧的問(wèn)題

盡管SIMT堆棧在早期GPU架構(gòu)中實(shí)現(xiàn)了分支管理能力，但其設(shè)計(jì)本質(zhì)上面臨多重硬件與效率瓶頸，難以適應(yīng)現(xiàn)代計(jì)算任務(wù)（光線追蹤、AI訓(xùn)練推理等）對(duì)復(fù)雜控制流的需求：

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

(4) 獨(dú)立線程調(diào)度

在Volta之前的架構(gòu)（如Pascal、Fermi）中，在分支線程調(diào)度上，由SIMT完成調(diào)度，而Warp作為基本調(diào)度單元，所有線程共享統(tǒng)一的PC和活動(dòng)掩碼，當(dāng)Warp內(nèi)線程執(zhí)行不同分支路徑時(shí)，需按路徑順序串行執(zhí)行。例如：線程0-3執(zhí)行分支A的指令，線程4-31執(zhí)行分支B的指令，則必須排隊(duì)執(zhí)行，一部分線程先執(zhí)行分支A的指令，另一部分線程必須等待。

而從Volta架構(gòu)開(kāi)始，引入了獨(dú)立線程調(diào)度（Independent Thread Scheduling）。每個(gè)線程擁有獨(dú)立的程序計(jì)數(shù)器(PC)和執(zhí)行狀態(tài)寄存器，允許同一Warp內(nèi)的線程在不同分支路徑上并行執(zhí)行指令流。但硬件層面仍以Warp為基本調(diào)度單元。

圖出處：NVIDIA Volta架構(gòu)白皮書(shū)

(5) 無(wú)堆棧分支收斂

同時(shí)，也是從Volta架構(gòu)開(kāi)始，隨著獨(dú)立線程調(diào)度的引入，傳統(tǒng)SIMT堆棧被棄用，分支收斂機(jī)制也升級(jí)到了無(wú)堆棧分支重新收斂（Stackless Branch Reconvergence）機(jī)制，通過(guò)收斂屏障（Convergence Barriers）技術(shù)來(lái)低成本解決分支代碼執(zhí)行調(diào)度問(wèn)題，獨(dú)立線程調(diào)度為無(wú)堆棧分支重新收斂提供了硬件支持。

無(wú)堆棧收斂屏障機(jī)制的核心手段之一是屏障參與掩碼（Barrier Participation Mask）與線程狀態(tài)協(xié)同管理，其核心思想可以通過(guò)ADD和WAIT操作來(lái)展示：

• ADD（屏障初始化）：當(dāng)Warp執(zhí)行到分支發(fā)散處前，通過(guò)專(zhuān)用ADD指令，活躍線程將其標(biāo)識(shí)位注冊(cè)到指定收斂屏障的32位掩碼中，標(biāo)記參與該屏障的線程組。
• WAIT（屏障同步）：在預(yù)設(shè)的收斂點(diǎn)（如分支匯合處），硬件插入WAIT指令。到達(dá)此處的子線程組將線程狀態(tài)標(biāo)記為“阻塞”，并更新屏障狀態(tài)寄存器。當(dāng)所有參與線程均抵達(dá)屏障后，調(diào)度器才重新激活完整線程束。

為了便于理解，下面用一個(gè)圖表示一個(gè)簡(jiǎn)單的的ADD，WAIT的例子：

另外，通過(guò)新增的syncwarp()函數(shù)，開(kāi)發(fā)者也可手動(dòng)指定分支后的同步點(diǎn)，強(qiáng)制線程在特定位置重新收斂。

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

3. 第二步：動(dòng)態(tài)指令調(diào)度以提高并發(fā)

在第一步構(gòu)建的最小可用系統(tǒng)中，采用的是“一條指令執(zhí)行完再執(zhí)行下一條”的最簡(jiǎn)執(zhí)行策略。前文提到過(guò)，GPU為了隱藏內(nèi)存訪問(wèn)的延遲，需要在內(nèi)存訪問(wèn)指令為執(zhí)行完前，先分配warp去執(zhí)行其他指令。這里的策略其實(shí)就是動(dòng)態(tài)指令調(diào)度，根據(jù)指令依賴(lài)關(guān)系和執(zhí)行單元可用性，動(dòng)態(tài)決定指令發(fā)射順序。

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

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

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

由于寄存器是線程私有的，需要為每個(gè)線程分配足夠的寄存器，因此SIMT核心中的寄存器數(shù)量是很大的，即便做到一個(gè)寄存器只需要一個(gè)bit表示狀態(tài)，ScoreBoard也會(huì)變得過(guò)大。因此，實(shí)際設(shè)計(jì)中，每個(gè)warp維護(hù)一個(gè)自己的ScoreBoard，由于每個(gè)warp同一時(shí)間只能執(zhí)行同一條指令，一條指令能訪問(wèn)的寄存器也是有限的，因此每個(gè)warp的ScoreBoard有3-4bit即可，每一個(gè)bit稱(chēng)為一個(gè)表項(xiàng)（entry）。

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

4. 第三步：提高并發(fā)指令的數(shù)據(jù)供給效率

(1) 并發(fā)指令數(shù)據(jù)訪問(wèn)

寄存器是處理器內(nèi)部的高速存儲(chǔ)單元，用于臨時(shí)存放指令執(zhí)行過(guò)程中所需的操作數(shù)、中間結(jié)果和地址信息。在GPU中，每個(gè)SIMT核心都擁有獨(dú)立的寄存器文件（寄存器的集合體，本質(zhì)上是一組寄存器組成的存儲(chǔ)陣列）。

第二步引入的計(jì)分板（ScoreBoard）機(jī)制，解決了時(shí)序維度上的數(shù)據(jù)依賴(lài)問(wèn)題，從而支持發(fā)射無(wú)依賴(lài)指令進(jìn)行延遲隱藏，除了時(shí)序上的復(fù)雜性，指令并行還會(huì)帶來(lái)空間上的復(fù)雜性，即大量并發(fā)指令同時(shí)嘗試訪問(wèn)寄存器文件獲取指令數(shù)據(jù)，寄存器文件必須支持多warp并發(fā)訪問(wèn)。

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

端口（port），是讀寫(xiě)存儲(chǔ)單元的接口。每多一個(gè)端口，存儲(chǔ)單元就可以多支持一個(gè)并發(fā)讀寫(xiě)操作，單端口的情況下，同時(shí)只能支持一個(gè)讀或一個(gè)寫(xiě)，若一個(gè)讀操作與一個(gè)寫(xiě)操作并發(fā)，則只能串行執(zhí)行，而增加一個(gè)端口，稱(chēng)為雙端口，則此時(shí)的一讀一寫(xiě)就可以并發(fā)完成。

因此，為了支持大量warp并發(fā)訪問(wèn)寄存器數(shù)據(jù)，一個(gè)簡(jiǎn)單粗暴的做法是，為寄存器文件設(shè)計(jì)足夠多的端口，來(lái)容納所有并發(fā)讀寫(xiě)操作。

盡管多端口設(shè)計(jì)在理論上可行，但其硬件代價(jià)呈指數(shù)級(jí)增長(zhǎng)，包括導(dǎo)致芯片面積暴增，同時(shí)，動(dòng)態(tài)功耗會(huì)隨端口數(shù)平方增長(zhǎng)、高訪問(wèn)延遲等。因此，簡(jiǎn)單的硬件堆料是低效且不可取的。

(3) 單端口內(nèi)存庫(kù)

寄存器文件與共享內(nèi)存的并發(fā)訪問(wèn)沖突，本質(zhì)上源于一個(gè)根本矛盾：存儲(chǔ)單元的物理端口數(shù)量與程序所需的并發(fā)訪問(wèn)量之間的不匹配。若將多端口設(shè)計(jì)比作“拓寬車(chē)道”，則單端口內(nèi)存庫(kù)（Single-Ported Memory Bank）更像是“優(yōu)化交通規(guī)則”——通過(guò)精細(xì)化調(diào)度，在有限硬件資源下挖掘最大效率。

在計(jì)算機(jī)存儲(chǔ)體系結(jié)構(gòu)中，banking是一種將存儲(chǔ)體分成多個(gè)獨(dú)立的部分（bank），每個(gè)bank可以獨(dú)立訪問(wèn)，從而提高并行訪問(wèn)能力的技術(shù)。如圖所示，單端口內(nèi)存庫(kù)將寄存器文件分成多個(gè)bank后，每個(gè)bank可以獨(dú)立進(jìn)行讀寫(xiě)操作，每個(gè)bank只有一個(gè)端口，如果同一時(shí)刻只有同一個(gè)線程訪問(wèn)，則可以成功，但如果有一個(gè)以上線程并發(fā)訪問(wèn)，則將產(chǎn)生訪問(wèn)沖突。

如此以來(lái)，分為多個(gè)bank的寄存器文件，一定程度上模擬了多端口寄存器文件的行為，即支持了跨bank的并發(fā)讀寫(xiě)操作。單端口內(nèi)存庫(kù)也是GPU核心架構(gòu)中最常見(jiàn)的片上存儲(chǔ)單元微架構(gòu)，除寄存器外，其同樣應(yīng)用于共享內(nèi)存。

同時(shí)，為了進(jìn)一步優(yōu)化，有很多通過(guò)寄存器布局優(yōu)化減少Bank沖突的機(jī)制。其中常見(jiàn)的有：

• 交錯(cuò)寄存器布局（Interleaved Register Allocation）。讓不同的warp的同編號(hào)寄存器分配到不同bank上。如warp0的R0分配到bank0，而warp1的R0則分配到bank1。這種布局方式在warp均勻調(diào)度發(fā)射指令（常見(jiàn)調(diào)度模式，大量warp輪流執(zhí)行）時(shí)可以有效地防止沖突。
• 動(dòng)態(tài)Bank分配（Dynamic Bank Allocation）。根據(jù)指令的寄存器訪問(wèn)模式，動(dòng)態(tài)調(diào)整邏輯寄存器到物理Bank的映射關(guān)系，避免靜態(tài)固定映射導(dǎo)致的沖突。
• 編譯器驅(qū)動(dòng)的寄存器分配優(yōu)化。編譯器在代碼生成階段，通過(guò)智能分配寄存器，減少Bank沖突。
• 在以上基礎(chǔ)上，發(fā)展出了混合Bank設(shè)計(jì)（Hybrid Banking）。將寄存器文件劃分為不同特性的Bank子集，采取不同的布局分配機(jī)制，針對(duì)不同訪問(wèn)模式優(yōu)化。

(4) 還有沖突：Operand Collector

接前文的例子，不論是“拓寬車(chē)道”還是“優(yōu)化交通規(guī)則”，總會(huì)有車(chē)道爭(zhēng)搶的問(wèn)題，那么也就總是需要“路口紅綠燈”來(lái)居中協(xié)調(diào)。

針對(duì)單bank的并發(fā)操作還是會(huì)引發(fā)數(shù)據(jù)沖突，這時(shí)就需要引入Operand Collector（操作數(shù)收集器）進(jìn)行指令的統(tǒng)一調(diào)度。Operand Collector是 GPU 流水線中的一個(gè)關(guān)鍵硬件模塊，負(fù)責(zé)在指令執(zhí)行前收集所有必需的操作數(shù)（即寄存器或內(nèi)存中的數(shù)據(jù)）。它的核心目標(biāo)是解決寄存器文件（Register File）的 Bank 沖突問(wèn)題，并通過(guò)動(dòng)態(tài)調(diào)度最大化寄存器訪問(wèn)的并行性，從而提升指令吞吐量。

當(dāng)指令進(jìn)入寄存器讀取階段（Register Read Stage）時(shí)，系統(tǒng)為其分配一個(gè)收集單元（Collector Unit），每個(gè)收集單元為一條指令服務(wù)，負(fù)責(zé)緩存該指令所需的所有源操作數(shù)（如 ADD R1, R2, R3 中的 R2 和 R3）。收集單元向寄存器文件發(fā)送讀請(qǐng)求，獲取源操作數(shù)。例如，指令 ADD R1, R2, R3 需要讀取 R2 和 R3。

當(dāng)不同指令出現(xiàn)數(shù)據(jù)沖突時(shí)，Operand Collector將動(dòng)態(tài)調(diào)度這些沖突的請(qǐng)求，將沖突請(qǐng)求分配到不同周期排隊(duì)執(zhí)行。若進(jìn)入排隊(duì)狀態(tài)，收集單元暫存已就緒的操作數(shù)，直到所有操作數(shù)準(zhǔn)備完畢，指令拿到操作數(shù)發(fā)射執(zhí)行。

四、總結(jié)

講到這里，已歷上萬(wàn)字，我們從引言中“Dennard Scaling”的失效開(kāi)始，引入GPU出現(xiàn)的背景，又介紹了GPU的通用性，以及高并發(fā)、低延遲保證的高計(jì)算速度。隨后，我們以最常見(jiàn)的CUDA為例，介紹了GPU編程的基礎(chǔ)，SIMT與SIMD，編譯鏈接的過(guò)程。最后，我們深入硬件層面，分為三步走，先用最簡(jiǎn)系統(tǒng)“run起來(lái)”，然后分別解決了指令依賴(lài)問(wèn)題，以及并發(fā)執(zhí)行中的數(shù)據(jù)訪問(wèn)沖突問(wèn)題，構(gòu)建并了解了一個(gè)通用GPU核心的架構(gòu)。

本文介紹的GPU知識(shí)，只是對(duì)各廠商、各架構(gòu)設(shè)計(jì)做“求同存異”后，得到的主干性的、通用性的基礎(chǔ)知識(shí)，而GPU作為當(dāng)代最為炙手可熱的科技產(chǎn)品之一，其發(fā)展是日新月異的。筆者希望這些基礎(chǔ)知識(shí)可以作為有興趣的讀者的“指路牌”，指引讀者在本文建立起來(lái)的基礎(chǔ)視野上，進(jìn)一步探索。