作者 | binnnliu

AI 流行的當(dāng)下，你有沒有想過：大模型推理服務(wù)到底怎么跑起來的？大模型推理服務(wù)的運(yùn)行過程中，CPU和GPU分別負(fù)責(zé)哪些工作？用GPU一定比CPU跑的快么？哪些場景需要用GPU?

一、圖形渲染到GPGPU

1. 為圖形而生

GPU最初的使命是加速圖形渲染。而渲染一幀圖像，本質(zhì)上就是對數(shù)百萬個像素點(diǎn)進(jìn)行相似的計(jì)算，這天然就是一種大規(guī)模并行任務(wù)。

2. 可編程性的開啟 (2001)

NVIDIA發(fā)布GeForce 3，首次引入可編程著色器 (Programmable Shaders)。實(shí)質(zhì)上允許開發(fā)者為 GPU 編寫軟件，讓GPU的眾多并行處理單元去同時(shí)執(zhí)行，以精確控制光照和顏色如何加載到顯示器上。這是朝著加速計(jì)算方向邁出的重要一步，因?yàn)樗试S開發(fā)者直接為 GPU 編寫軟件。

3. 學(xué)術(shù)界的探索

一批敏銳的研究人員意識到，GPU的本質(zhì)就是一個擁有數(shù)百甚至數(shù)千個核心的大規(guī)模并行架構(gòu)，其浮點(diǎn)運(yùn)算吞吐量遠(yuǎn)超當(dāng)時(shí)的CPU。他們的核心想法是：能不能用GPU進(jìn)行科學(xué)計(jì)算?  開始探索利用GPU計(jì)算科學(xué)計(jì)算問題，從而利用GPU的算力。這便是GPGPU（通用計(jì)算GPU）的萌芽。但是門檻非常高， 需要開發(fā)者同時(shí)精通圖形學(xué)和科學(xué)計(jì)算。

4. NVIDIA的抉擇

NVIDIA敏銳地捕捉到了GPGPU的發(fā)展?jié)摿Γ_始不再局限于加速圖形渲染，主動擁抱GPGPU。

2006年，發(fā)布了第一款為通用計(jì)算設(shè)計(jì)的統(tǒng)一架構(gòu)GPU -  GeForce 8800 GTX 顯卡（G80架構(gòu)）。它將GPU內(nèi)部的計(jì)算單元統(tǒng)一起來，形成了一個龐大的、靈活的并行核心陣列，為通用計(jì)算鋪平了硬件道路。

2007年，NVIDIA正式推出了CUDA平臺。CUDA的革命性在于，它提供了一套簡單的編程模型，讓開發(fā)者能用近似C語言的方式，輕松地駕馭GPU內(nèi)部成百上千個并行核心。 開發(fā)者無需再關(guān)心復(fù)雜的圖形接口，可以直接編寫在數(shù)千個線程上并發(fā)執(zhí)行的程序。至此終結(jié)了GPGPU編程的蠻荒時(shí)代，讓GPU計(jì)算真正走下神壇，成為開發(fā)者觸手可及的強(qiáng)大工具。

隨著深度學(xué)習(xí)的發(fā)展與流行，CUDA生態(tài)系統(tǒng)目前也成為NVIDIA最深、最寬的護(hù)城河。

參考鏈接：nvidia-past-present-and-future

二、CPU/GPU異構(gòu)計(jì)算架構(gòu)

CPU是整個系統(tǒng)的核心，是總指揮，GPU的任務(wù)指令是由CPU分配的。

CPU通過PCIe總線給GPU發(fā)送指令和數(shù)據(jù)交互。而PCIe支持DMA和MMIO兩種通訊模式:

• MMIO（內(nèi)存映射I/O）由CPU直接控制數(shù)據(jù)讀寫，操作系統(tǒng)會把設(shè)備地址映射到CPU的虛擬空間中，適合小數(shù)據(jù)量的指令交互
• DMA（直接內(nèi)存訪問）則允許設(shè)備繞過CPU直接訪問系統(tǒng)內(nèi)存，專為大數(shù)據(jù)塊的高效傳輸設(shè)計(jì)。

CPU通過IMC和Memory Channel訪問內(nèi)存，為了提升數(shù)據(jù)傳輸帶寬，高端CPU通常會支持多內(nèi)存通道，即多IMC和Memory Channel的組合，以滿足日益增長的數(shù)據(jù)處理需求。

三、一個簡單的應(yīng)用

講道理，對于開發(fā)來說，再通俗易懂的語言描述都不如一個簡單Demo來的實(shí)在。

Demo代碼來自even-easier-introduction-cuda,可在collab測試運(yùn)行下述代碼。

實(shí)現(xiàn)兩個長度為 23? (約10億) 的浮點(diǎn)數(shù)數(shù)組的相加。其中，一個數(shù)組 (x) 的所有元素初始化為 1.0，另一個數(shù)組 (y) 的所有元素初始化為 2.0，我們計(jì)算 y[i] = x[i] + y[i]。

1. CPU的實(shí)現(xiàn)

#include <iostream>
#include <math.h>
#include <chrono>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
    for (int i = 0; i < n; i++)
        y[i] = x[i] + y[i];
}

int main(void)
{
    int N = 1<<30;

    float *x = new float[N];
    float *y = new float[N];

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    auto start = std::chrono::high_resolution_clock::now();

    // Run kernel on 1M elements on the CPU
    add(N, x, y);

    auto stop = std::chrono::high_resolution_clock::now();
    auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(stop - start);
    std::cout << "CPU 'add' function execution time: " << duration.count() << " ms" << std::endl;

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(y[i]-3.0f));

    std::cout << "Max error: " << maxError << std::endl;

    delete [] x;
    delete [] y;

    return 0;
}

性能表現(xiàn)：

g++ add.cpp -o add
time ./add

CPU 'add' function execution time: 3740 ms
Max error: 0

real 0m21.418s
user 0m15.798s
sys 0m4.400s

• 計(jì)算耗時(shí): 核心的add函數(shù)耗時(shí) 3740毫秒。
• 總耗時(shí): 整個程序從啟動到結(jié)束（real time）耗時(shí) 21.4秒。這額外的時(shí)間主要消耗在分配8GB內(nèi)存（new float[N]）以及初始化數(shù)組上。

2. GPU的實(shí)現(xiàn)

這里的代碼后面會詳細(xì)解讀，此處看懂含義即可。

• 分配內(nèi)存: 分別在CPU（Host）和GPU（Device, cudaMalloc）上分配內(nèi)存。
• 數(shù)據(jù)傳輸 (H2D): 將CPU上的輸入數(shù)據(jù) (h_x, h_y) 拷貝到GPU顯存 (d_x, d_y)。
• 執(zhí)行Kernel函數(shù): 在GPU上啟動addKernel函數(shù)，利用其大規(guī)模并行能力進(jìn)行計(jì)算。
• 數(shù)據(jù)傳回 (D2H): 將GPU計(jì)算完成的結(jié)果 (d_y) 拷貝回CPU內(nèi)存 (h_y) 以便后續(xù)使用或驗(yàn)證。

#include <iostream>
#include <math.h>

#define CUDA_CHECK(call) \
do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error in %s at line %d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while (0)


// __global__ 關(guān)鍵字聲明的函數(shù)被稱為Kernel函數(shù)
__global__
void add(int n, float *x, float *y)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index < n) {
        y[index] = x[index] + y[index];
    }
}

int main(void)
{
    int N = 1 << 30;
    size_t bytes = N * sizeof(float);


    float *h_x, *h_y;
    h_x = new float[N];
    h_y = new float[N];

    float *d_x, *d_y;
    CUDA_CHECK(cudaMalloc(&d_x, bytes));
    CUDA_CHECK(cudaMalloc(&d_y, bytes));

    for (int i = 0; i < N; i++) {
        h_x[i] = 1.0f;
        h_y[i] = 2.0f;
    }


    CUDA_CHECK(cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_y, h_y, bytes, cudaMemcpyHostToDevice));

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    CUDA_CHECK(cudaEventRecord(start));

    int blockSize = 256;
    int numBlocks = (N + blockSize - 1) / blockSize;
    add<<<numBlocks, blockSize>>>(N, d_x, d_y);

    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));

    float milliseconds = 0;
    CUDA_CHECK(cudaEventElapsedTime(&milliseconds, start, stop));
    std::cout << "GPU Kernel 'add' execution time: " << milliseconds << "         ms" << std::endl;

    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    CUDA_CHECK(cudaMemcpy(h_y, d_y, bytes, cudaMemcpyDeviceToHost));

    float maxError = 0.0f;
    for (int i = 0; i < N; i++) {
        maxError = fmax(maxError, fabs(h_y[i] - 3.0f));
    }
    std::cout << "Max error: " << maxError << std::endl;

    delete[] h_x;
    delete[] h_y;

    CUDA_CHECK(cudaFree(d_x));
    CUDA_CHECK(cudaFree(d_y));

    return 0;
  }

(1) 性能表現(xiàn)

nvcc  add.cu -o add_cu -gencode arch=compute_75,code=sm_75
time ./add_cu

GPU Kernel 'add' execution time: 48.6738 ms
Max error: 0

real 0m19.413s
user 0m15.308s
sys 0m4.014s

• 計(jì)算耗時(shí): GPUKernel函數(shù)的執(zhí)行耗時(shí)僅為 48.7毫秒。
• 總耗時(shí): 程序總耗時(shí)為 19.4秒。

(2) 性能分析

單看核心計(jì)算任務(wù)，GPU (48.7ms) 的速度是CPU (3740ms) 的 約75倍。這完美體現(xiàn)了GPU在處理數(shù)據(jù)并行任務(wù)時(shí)的絕對優(yōu)勢。CPU需要串行執(zhí)行10億次加法(此處只考慮單核場景)，而GPU則將任務(wù)分配給成千上萬個線程同時(shí)處理。

但是雖然GPU計(jì)算本身極快，但程序的總耗時(shí) (19.4s) 卻和CPU版本 (21.4s) 相差無幾。這是為什么呢？主要是CPU和GPU通訊的開銷。這里下一篇文章會詳細(xì)介紹。

四、編譯-Fat Binary

nvcc  add.cu -o add_cu -gencode arch=compute_75,code=sm_75 上面的例子中，我們看到這個編譯指令。add.cu被編譯為二進(jìn)制文件add_cu。它具體是怎么做的呢？

(1) 主機(jī)代碼編譯: 將C/C++代碼（在CPU上運(yùn)行的部分）交由系統(tǒng)的主機(jī)編譯器（如GCC、MSVC）編譯成標(biāo)準(zhǔn)的CPU目標(biāo)代碼。

(2) 設(shè)備代碼編譯: 將在__global__函數(shù)（如add）中定義的GPU代碼，編譯成兩種主要格式：

• SASS (Streaming Assembler): 這是特定GPU架構(gòu)的原生機(jī)器碼。例如，為NVIDIA T4 GPU (Turing 架構(gòu) 代號 sm_75 ) 編譯的SASS，只能在該架構(gòu)上最高效地運(yùn)行。NVCC可以為多種指定的架構(gòu)預(yù)編譯多份SASS代碼。
• PTX (Parallel Thread eXecution): arch=compute_75 指示編譯器生成一份 PTX 代碼，確保程序能在任何不低于 Turing 架構(gòu)的新 GPU 上通過 JIT 編譯運(yùn)行（向前兼容性）。

這兩種設(shè)備代碼連同主機(jī)代碼一起，被打包進(jìn)一個可執(zhí)行文件中，形成所謂的胖二進(jìn)制 (Fat Binary)。它“胖”在包含了一份主機(jī)代碼和多份針對不同GPU架構(gòu)的設(shè)備代碼。

1. 程序加載 - cubin loading

(1) 程序啟動

操作系統(tǒng)加載可執(zhí)行文件，CPU 開始執(zhí)行主機(jī)代碼。

(2) 首次 CUDA 調(diào)用

當(dāng)代碼第一次調(diào)用任何 CUDA API 函數(shù)時(shí)（比如 cudaSetDevice, cudaMalloc，或者第一個Kernel函數(shù)啟動），CUDA 運(yùn)行時(shí)庫 (CUDA Runtime Library) 會被初始化。

此處就是所謂的GPU上下文初始化/CUDA上下文初始化，主要步驟：

① 硬件準(zhǔn)備與喚醒

從低功耗的待機(jī)模式喚醒，進(jìn)入高性能的計(jì)算模式；

加載驅(qū)動模塊（如NVIDIA CUDA Driver或AMD ROCm），并檢測可用GPU設(shè)備及其屬性（如顯存大小、計(jì)算能力、NVLink連接）。

② CUDA上下文數(shù)據(jù)結(jié)構(gòu)創(chuàng)建

CPU側(cè)創(chuàng)建上下文信息的數(shù)據(jù)結(jié)構(gòu)：創(chuàng)建一個統(tǒng)一虛擬地址空間（UVA)，這個空間可以將所有的系統(tǒng)內(nèi)存和所有GPU的內(nèi)存都映射進(jìn)來，共享一個單一的虛擬地址空間。(每次cudaMalloc都會增加一條記錄)

③ 特定GPU上創(chuàng)建上下文

a. 在顯存中為當(dāng)前進(jìn)程分配并建立頁表結(jié)構(gòu)：

• NVIDIA驅(qū)動程序（在CPU上）查詢其內(nèi)部維護(hù)的、用于管理GPU物理顯存的數(shù)據(jù)結(jié)構(gòu)（即VRAM Allocator，跨進(jìn)程維護(hù)），以找到一個空閑的物理地址。CPU本地軟件操作，不涉及與GPU的硬件通信。
• CPU在自己的內(nèi)存（RAM）里，準(zhǔn)備好了要寫入的數(shù)據(jù)內(nèi)容；
• NVIDIA驅(qū)動程序（在CPU上）命令DMA引擎將對應(yīng)數(shù)據(jù)復(fù)制到顯存；

b. 分配Pinned Memory命令緩沖區(qū)

c. 通過MMIO配置GPU的MMU硬件(PMMU 控制寄存器)，告訴它頁表的起始位置

④ 上下文就緒

上下文完全建立，后續(xù)的Kernel函數(shù)啟動、內(nèi)存拷貝等命令可以通過流 (Stream) 機(jī)制提交到其命令緩沖區(qū)，由GPU異步執(zhí)行。

2. 首次調(diào)用add<<<...>>>()時(shí)，進(jìn)行Kernel函數(shù)加載

(1) 檢測硬件

它會查詢當(dāng)前的 GPU，識別出具體架構(gòu)。

(2) 尋找最佳匹配 (SASS)

然后，它會在 Fat Binary 的設(shè)備代碼段中進(jìn)行搜索，尋找有沒有預(yù)編譯好的、針對 sm_75 的 SASS 代碼。

(3) 沒有找到完全匹配的 SASS 代碼

如果沒有找到完全匹配的 SASS 代碼運(yùn)行時(shí)會找到 PTX 中間代碼，并調(diào)用集成在 GPU 驅(qū)動中的 JIT (Just-In-Time) 編譯器將其即時(shí)編譯(JIT)為目標(biāo)GPU的SASS代碼; (cpu上完成)；

為了避免每次運(yùn)行程序都重新進(jìn)行 JIT 編譯，NVIDIA 驅(qū)動通常會緩存 JIT 編譯的結(jié)果。NVIDIA驅(qū)動會在用戶的home目錄下創(chuàng)建一個計(jì)算緩存，通常是 ~/.nv/ComputeCache。

(4) cubin loading (cubin 是 CUDA binary 的縮寫)

• 將準(zhǔn)備好的 SASS 代碼（無論是來自 Fat Binary 還是 JIT 編譯的結(jié)果）申請顯存空間；通過DMA復(fù)制到顯存；
• 驅(qū)動程序在其內(nèi)部的表格中，將Kernel函數(shù) add 與其在 VRAM 中的地址關(guān)聯(lián)起來。后續(xù)調(diào)用 add<<<...>>>() 時(shí)，運(yùn)行時(shí)會將一個包含該 VRAM 地址的啟動命令提交到流中，由 GPU 異步執(zhí)行。

五、程序執(zhí)行 - Kernel Launch

一個常見的誤解是CPU會直接、實(shí)時(shí)地控制GPU。實(shí)際上，考慮到CPU和GPU是兩個獨(dú)立的處理器，并且通過PCIe總線連接，直接的、同步的控制會帶來巨大的延遲和性能開銷。因此，現(xiàn)代GPU采用了一種高效的異步通信模型，其核心就是 命令緩沖區(qū)（Command Buffer）與門鈴（Doorbell）機(jī)制。這也是CUDA Streaming的底層通訊機(jī)制。

1. Command Buffer + Doorbell 機(jī)制

• cpu先把需要執(zhí)行的命令寫到ring buffer命令緩沖區(qū)(Pinned Memory,位于主機(jī)內(nèi)存); 更新w_ptr
• 在適當(dāng)?shù)臅r(shí)候通過MMIO設(shè)置Doorbell Register，告訴GPU有新任務(wù)需要處理
• GPU上的DMA引擎將ring buffer命令緩沖區(qū)[r_ptr, w_ptr)復(fù)制到顯存中，然后開始執(zhí)行；(其中w_ptr和r_ptr可以理解為相對于 Ring Buffer 基地址 (Base Address) 的偏移量)

下面對于部分由代表型的API的執(zhí)行邏輯進(jìn)行單獨(dú)闡述。

2. CPU 執(zhí)行到cudaMalloc

cudaMalloc 是一個同步阻塞調(diào)用，它不使用上述的流式命令緩沖區(qū)機(jī)制。（CUDA 11.2+支持cudaMallocAsync可實(shí)現(xiàn)異步分配）

• CPU 線程調(diào)用 cudaMalloc()。CUDA 運(yùn)行時(shí)庫將此請求轉(zhuǎn)發(fā)給 NVIDIA 驅(qū)動程序
• 驅(qū)動程序向物理VRAM Allocator請求物理內(nèi)存，向 UVA Manager 請求虛擬地址，更新UVA映射表；（物理VRAM Allocator是跨進(jìn)程的，維護(hù)整個GPU 物理顯存的使用情況）
• 更新 GPU page table[Command Buffer + Doorbell方式,特定的、高優(yōu)先級的通道，非默認(rèn)的Stream]，刷新TLB
• 返回虛擬內(nèi)存指針

與malloc的不同之處：

(1) Lazy Allocation vs. Eager Allocation

malloc支持overcommit，實(shí)際物理內(nèi)存的分配發(fā)生在訪問時(shí)(Lazy Allocation)，通過缺頁中斷（Page Fault）按需映射到物理內(nèi)存；而cudaMalloc是同步分配連續(xù)的物理顯存(Eager Allocation)，保證了后續(xù)使用的確定性，但初始開銷更高。

(2) system call overhead

cudaMalloc直接陷入內(nèi)核，調(diào)用GPU驅(qū)動分配物理內(nèi)存；而malloc本身是C庫函數(shù)(用戶態(tài)), 向操作系統(tǒng)“批發(fā)”大塊內(nèi)存，然后在用戶程序請求時(shí)“零售”出去。避免內(nèi)存分配時(shí)昂貴的系統(tǒng)調(diào)用和缺頁異常開銷

• 申請<128KB內(nèi)存時(shí)，會優(yōu)先在freelist中查找是否有合適的空閑 Chunk，沒有找到，才會通過brk系統(tǒng)調(diào)用向操作系統(tǒng)申請內(nèi)存
• 申請>=128KB內(nèi)存時(shí)，會直接通過mmap系統(tǒng)調(diào)用向操作系統(tǒng)申請內(nèi)存，free時(shí)也會直接釋放

(3) 釋放策略

cudaFree會直接釋放，而free對于brk/sbrk分配的內(nèi)存不會直接釋放(物理內(nèi)存和虛擬內(nèi)存都不釋放，為了避免Page Fault引入的性能開銷就沒有釋放物理內(nèi)存)，用戶態(tài)維護(hù)freelist，同時(shí)會合并連續(xù)空閑的虛擬地址空間，有效減少內(nèi)存碎片(coalescing)。

3. CPU 執(zhí)行到 cudaMemcpy、cudaMemset

通過Command Buffer + Doorbell 機(jī)制提交命令到GPU; 然后同步或者異步等待。

4. CPU 執(zhí)行到Kernel函數(shù)add<<<...>>>()

(1) CPU側(cè)：命令打包與提交

• 驅(qū)動將Kernel函數(shù)啟動所需信息打包成一個命令命令包括：啟動Kernel函數(shù)，Kernel函數(shù)對應(yīng)的add SASS 代碼的入口地址，執(zhí)行配置(Grid 維度、Block 維度、共享內(nèi)存大小等)、參數(shù)指針(GPU虛擬地址)
• 將命令寫入主機(jī)端的 Pinned Memory Ring Buffer
• 通過 MMIO 寫 Doorbell 寄存器，通知 GPU

(2) GPU側(cè): 命令獲取與運(yùn)行

① 通過 DMA 從 Pinned Memory 讀取Ring buffer部分內(nèi)容

② 命令解碼

• GPU 的命令處理器 (Front-End) 從其內(nèi)部隊(duì)列中取出命令包。
• 它開始解碼這個命令包，識別出這是一個“Kernel函數(shù)啟動”任務(wù)，并解析出所有的執(zhí)行參數(shù)（Grid/Block 維度、Kernel函數(shù)地址等）。

③ 工作分發(fā)

int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, d_x, d_y);

• 命令處理器根據(jù) Grid 的維度，將整個計(jì)算任務(wù)分發(fā)成一個個獨(dú)立的Thread Blocks。
• GPU的全局調(diào)度器（GigaThread Engine）,將Thread Blocks分配給有空閑資源的 SM。一個線程塊從生到死都只會在一個 SM 上執(zhí)行，不會遷移。

④ 線程塊調(diào)度與執(zhí)行

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

• 每個 SM 接收到一個或多個線程塊，SM 內(nèi)部的硬件調(diào)度器 (Scheduler)進(jìn)一步將每個線程塊內(nèi)部的線程，按照threadIdx的順序，每 32 個線程劃分成一個 Warp。比如，一個有 256 個線程的線程塊，會被劃分為 8 個 Warps (Warp 0: 線程 0-31, Warp 1: 線程 32-63, ...)。
• SM 內(nèi)部的硬件調(diào)度器 (Scheduler) Warps分配給 SM 內(nèi)的CUDA Cores 和其他執(zhí)行單元（如 Tensor Cores）去執(zhí)行。
• CUDA 核心開始執(zhí)行位于指定 SASS 地址的機(jī)器指令，進(jìn)行實(shí)際的計(jì)算。

⑤ 完成與資源回收

• 當(dāng)一個線程塊完成了所有計(jì)算，它所占用的 SM 資源（如寄存器、共享內(nèi)存）會被釋放，SM 可以接收新的線程塊。
• 當(dāng)整個 Grid 的所有線程塊都執(zhí)行完畢，這個Kernel函數(shù)啟動任務(wù)就算完成了。

Grid、Thread Block、Warp、Thread、SM這些概念到底是干啥的。下面結(jié)合GPU的硬件架構(gòu)詳細(xì)介紹。

六、GPU的硬件架構(gòu)

如上是NVIDIA GA100 GPU的架構(gòu)圖：

A100 GPU 架構(gòu)圖

1. 計(jì)算單元

(1) GPC

Graphics Processing Cluster,  一個GPU包含多個GPC, 一個GPC包含多個TPC

(2) TPC

Texture Processing Cluster, 一個TPC包含多個SM

(3) SM

Streaming Multiprocessor, SM是GPU執(zhí)行計(jì)算任務(wù)的核心單元，它是：

• CUDA Cores (執(zhí)行FP32/INT32等通用計(jì)算的ALUs/FPUs)
• Tensor Cores

一個硬件單元，專門處理**FMA(Fused Multiply-Add)**操作，能在一個時(shí)鐘周期內(nèi)完成一個小的矩陣乘加運(yùn)算（一個4x4的FP16矩陣相乘后累加到另一個4x4矩陣上）深度學(xué)習(xí)絕大部分的計(jì)算都是FMA操作，NVidia工程師為此專門設(shè)計(jì)專用計(jì)算單元。

• 寄存器 (Register File)、共享內(nèi)存 (Shared Memory)
• L1數(shù)據(jù)緩存/指令緩存 (L1 Data Cache / Instruction Cache)
• Warp調(diào)度器 (Warp Scheduler) 等關(guān)鍵組件

單個SM的架構(gòu)圖如下：

2. 接口

• PCIe 負(fù)責(zé)CPU與GPU的通訊，DMA模式
• NVLINK 負(fù)責(zé)GPU間的通訊

3. 內(nèi)存與緩存

其中HBM和L2 Cache是整個GPU共享的;

而L1 Cache/Shared Memory則是SM維度獨(dú)享的;

Shared Memory是每個SM內(nèi)部的一塊高速、可編程的片上緩存。同一線程塊（Block）內(nèi)的所有線程都可以訪問它，速度遠(yuǎn)快于訪問全局顯存（HBM）。它是實(shí)現(xiàn)Block內(nèi)線程高效協(xié)作和數(shù)據(jù)交換的核心，對于矩陣乘法等需要數(shù)據(jù)復(fù)用的算法至關(guān)重要。

速度由快到慢依次為 寄存器 -> L1 Cache -> L2 Cache -> HBM -> DRAM(主機(jī)內(nèi)存)

七、編程模型 vs 硬件執(zhí)行模型

1. 編程模型

將一個待批量并發(fā)的數(shù)據(jù)組織成Grid、Thread Block、Thread的結(jié)構(gòu)。

Grid和Thread Block可以是1維的也可以是2維或者3維的。這里這么設(shè)計(jì)，感覺主要是為了讓程序員可以根據(jù)實(shí)際處理的結(jié)構(gòu)能夠更自然的思考,同時(shí)可以覆蓋數(shù)據(jù)局部性需求，比如，我要處理一個1維數(shù)據(jù)，自然的我們就可以把Grid和Thread Block定義為1維的。比如上面例子中計(jì)算1維數(shù)組的加法，就可以用1維的Grid和Thread Block。

int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, d_x, d_y);

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

Grid視圖：

這行代碼是CUDA編程的基石(SIMT)，它將軟件層面的線程坐標(biāo)映射到數(shù)據(jù)上的全局索引。

• threadIdx.x: 當(dāng)前 Thread 在其 Block 內(nèi)的 x 坐標(biāo)。范圍是 0 到 blockDim.x - 1。
• blockDim.x: 每個 Block 在 x 維度上有多少個 Thread。（在我們例子中是256）。
• blockIdx.x: 當(dāng)前 Block 在 Grid 中的 x 坐標(biāo)。范圍是 0 到 gridDim.x - 1。
• gridDim.x: Grid 在 x 維度上有多少個 Block。

blockIdx.x * blockDim.x計(jì)算出了當(dāng)前線程塊之前所有線程塊包含的線程總數(shù)（偏移量），再加上threadIdx.x，就得到了當(dāng)前線程在整個Grid中的全局唯一ID。這保證了10億個元素，每個都能被一個特定的線程處理到。

這里解釋下上面提到的數(shù)據(jù)局部性:  y[index] = x[index] + y[index]; 可以合并訪存 (Coalesced Memory Access)。即一個Warp中的32個線程訪問連續(xù)的32個內(nèi)存地址，GPU硬件可以將其合并成一次或少數(shù)幾次寬內(nèi)存事務(wù)，極大提升訪存效率。

而當(dāng)我們要處理一個二維矩陣或圖像時(shí)，最自然的思考方式就是二維的。這時(shí)候我們可以用2維的Grid和Thread Block。

dim3 blockSize(16, 16); // 16x16 = 256 線程/塊
dim3 gridSize((N + blockSize.x - 1) / blockSize.x, (N + blockSize.y - 1) / blockSize.y);

__global__ void matrixMulGPU(const float* A, const float* B, float* C, int N) {

    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}

Grid視圖：

2. 硬件層面

將整個GPU的運(yùn)算單元分為 GPU、SM、Warp和Core。

軟件層面將grid切分成多個Thread Block是為了對硬件的抽象，這樣程序員就不必關(guān)心GPU具體有多少個物理核心、多少個SM。

Thread Block是最小的“資源分配與調(diào)度”單位，Warp是最小的硬件調(diào)度單位。

所以整個編程模型大概就是：

一個任務(wù)軟件層面上被分為Grid和Thread Block，Thread Block被分配給硬件的SM，SM又將Thread Block按照32個Thread為一組分成Warp，分配給Warp scheduler執(zhí)行。

最終的視圖大概是這樣的：

3. 隱藏延遲 - hide latency

前面已經(jīng)看到一個計(jì)算任務(wù)對應(yīng)一個Grid，一個Grid又由多個Thread Block組成，GPU的全局調(diào)度器（GigaThread Engine）將Thread Blocks分配給有空閑資源的 SM。(多個Thread Blocks可以被分配給一個SM,取決于共享內(nèi)存、寄存器使用的使用情況)

一個Thread Block被分解成多個Warp（例如，一個1024線程的Block被分解成32個Warp)。SM內(nèi)部的調(diào)度硬件，會將這32個Warp分配給它內(nèi)部的4個Warp Scheduler。通常會盡量均勻分配，比如每個Warp Scheduler分到8個Warp。

而一個Warp Scheduler同一時(shí)刻只能運(yùn)行一個Warp, 當(dāng)某個正在執(zhí)行的Warp因?yàn)榈却齼?nèi)存而暫停時(shí)，它可以立刻從剩下的Warp中挑選一個就緒的來執(zhí)行。這就是所謂的隱藏延遲 (hide latency)。而如何充分利用這個特性呢？給每個Warp Scheduler足夠多的可切換的Warp。

每個SM都包含一個巨大、單一的物理寄存器文件，為實(shí)現(xiàn)零開銷Warp上下文切換的提供了硬件基礎(chǔ)。這是與CPU昂貴的上下文切換（需要保存和恢復(fù)大量狀態(tài)）的根本區(qū)別。

要讓每個 Warp Scheduler (Warp 調(diào)度器) 有足夠的可切換 Warp，其本質(zhì)是提高 GPU 的占用率。占用率指的是一個 SM  上實(shí)際活躍的 Warp 數(shù)量與該 SM 理論上能支持的最大 Warp 數(shù)量的比例。

一個 SM 能同時(shí)運(yùn)行多少 Warp（**一個 SM 在同一時(shí)刻只能為一個 Kernel 服務(wù)，但可以同時(shí)運(yùn)行該Kernel的多個線程塊(只要資源允許)**），取決于以下三個主要資源的限制：

(1) Registers

每個線程都需要使用寄存器來存儲其局部變量。一個 SM 上的寄存器總數(shù)是固定的

假設(shè)一個 SM 有 65536 個寄存器，最大支持 2048 個線程 (64 Warps)。 每個Kernel需要 64 個寄存器，那么一個 Block (假設(shè) 256 線程) 就需要 256 * 64 = 16384 個寄存器。這個 SM 最多可以容納 65536 / 16384 = 4 個這樣的 Block，也就是 1024 個線程 (32 Warps)，占用率為 50%。如果 Kernel 每個線程需要 128 個寄存器，那么這個 SM 只能容納 2 個這樣的 Block，占用率就更低了。

(2) Shared Memory

共享內(nèi)存是分配給每個線程塊 (Block) 的、速度很快的片上內(nèi)存。一個 SM 上的共享內(nèi)存總量是固定的。

假設(shè)一個 SM 有 96KB 共享內(nèi)存，最大支持 16 個 Block。如果Kernel 每個 Block 需要 32KB 共享內(nèi)存，那么這個 SM 最多只能同時(shí)運(yùn)行 96KB / 32KB = 3 個 Block。在這個場景下，共享內(nèi)存成為了主要的限制因素。這就將 SM 上并發(fā)的 Block 數(shù)量上限從硬件支持的 16 個銳減到了 3 個，從而嚴(yán)重限制了 SM 上的總并發(fā) Warp 數(shù)量，降低了占用率。

(3) 線程塊/線程數(shù)限制

每個 SM 架構(gòu)本身就有硬件限制，比如一個 SM 最多能同時(shí)調(diào)度多少個 Block（例如 16 或 32），以及最多能同時(shí)管理多少個線程（例如 2048）。這個是硬性上限，無法通過代碼改變。

不過提高 GPU 的占用率來隱藏延遲也不是萬能的，隱藏延遲的有效性，本質(zhì)上取決于 Warp調(diào)度器是否有“就緒態(tài)”的Warp可供切換。比如：如果一個Kernel非常簡單，每個線程只使用極少的寄存器，并且不使用共享內(nèi)存，那么一個SM上可能會駐留大量的Warp。但如果這個Kernel的計(jì)算是訪存密集型且延遲很高的，同時(shí)計(jì)算/訪存指令比例很低，那么即使占用率達(dá)到100%，Warp調(diào)度器可能依然會“無Warp可調(diào)”，因?yàn)樗蠾arp都在等待數(shù)據(jù)返回。這時(shí)候我們就不得不提另外一個概念，訪存比(Ratio = Total Bytes / Total FLOPs)或者計(jì)算強(qiáng)度(Roofline,I = Total FLOPs / Total Bytes), 說白了，就是看一個程序是計(jì)算密集型(Compute-bound)還是IO(內(nèi)存訪問)密集型(Memory-bound)。可以使用NVIDIA Nsight Compute分析Kernel函數(shù)的占用率和計(jì)算強(qiáng)度。 不過這里不做延伸了，放到下篇性能優(yōu)化中講。

八、SIMD vs SIMT

前面CUDA Demo中我們已經(jīng)知道Kernel函數(shù)add會被啟動成茫茫多的線程執(zhí)行,每個線程通過計(jì)算 blockIdx 和 threadIdx 來處理不同的數(shù)據(jù)。

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

從程序員的角度看，我們似乎是在編寫多線程（Multiple Threads）程序。但從硬件的角度看，它是如何讓這么多線程同時(shí)執(zhí)行同一條指令（Single Instruction）的呢？

這種 "單指令，多線程"（Single Instruction, Multiple Threads, SIMT）的編程模型，正是CUDA的魅力所在。SIMT通過線程編程模型巧妙的隱藏了底層SIMD的執(zhí)行細(xì)節(jié)。而要理解SIMT，就不得不提在CPU中廣泛使用的SIMD技術(shù)。

1. SIMD(Single instruction, multiple data)

在傳統(tǒng)的標(biāo)量計(jì)算模型中，CPU的一條指令一次只能操作單個數(shù)據(jù)。例如，一次浮點(diǎn)加法就是double + double;

當(dāng)處理如圖形、音頻或科學(xué)計(jì)算中常見的大規(guī)模數(shù)據(jù)集時(shí)，這種“一次一個”的模式效率極低，因?yàn)槲覀冃枰獙Ａ繑?shù)據(jù)重復(fù)執(zhí)行完全相同的操作，這暴露了標(biāo)量處理的瓶頸。

為了打破這個瓶頸，現(xiàn)代CPU集成了SIMD（單指令，多數(shù)據(jù)）架構(gòu)。CPU增加了能容納多個數(shù)據(jù)元素的寬向量寄存器（如256位的YMM寄存器），以及能夠并行處理這些數(shù)據(jù)的執(zhí)行單元。

比如_mm256_add_pd cpu可以同時(shí)進(jìn)行4對double的加法運(yùn)算(256位的寄存器, 256/64=4)

為了加速多媒體和科學(xué)計(jì)算，Intel不斷引入更強(qiáng)大的SIMD指令集，從MMX的64位 -> SSE的128位 -> AVX的256位 -> AVX-512的512位。

但是SIMD偏硬件底層，編程不友好：

• 手動打包解包向量
• 手動處理if else邏輯

2. SIMT(Single instruction, multiple thread)

為了解決編程不友好的問題，NVIDIA提出SIMT（Single Instruction, Multiple Threads）。SIMT是CUDA編程的基石，是GPU從一種處理圖形計(jì)算的專用硬件，進(jìn)化為GPGPU的基礎(chǔ)。

具體實(shí)現(xiàn)簡單來說就是：同一時(shí)刻，Warp調(diào)度器只發(fā)布一條指令，后端仍然以SIMD的模式執(zhí)行，而具體哪些線程執(zhí)行依賴活動掩碼控制。（ps: 下圖為Pre-Volta的一個示意圖，Volta以及之后的架構(gòu)由于線程獨(dú)立PC和Stack的出現(xiàn)，SIMT Stack已被淘汰）

SIMT巧妙的隱藏了SIMD的復(fù)雜性，程序員只需要思考單個線程的邏輯，大大降低了心智負(fù)擔(dān)。比如，如下代碼每個thread都執(zhí)行相同的代碼，但是由于每個thread都會計(jì)算出特有的index，所有其實(shí)都在處理不同的數(shù)據(jù)。

int i = blockIdx.x * blockDim.x + threadIdx.x;
C[i] = A[i] + B[i];

3. Warp Divergence

每個Warp中的32個線程必須同步的執(zhí)行相同的指令序列(SIMT是基于Warp的SIMD)，這就導(dǎo)致在處理if-else時(shí)，GPU需要串行執(zhí)行每個分支，導(dǎo)致算力浪費(fèi)。

(1) Pre-Volta

在Pre-Volta架構(gòu)中，一個Warp（32個線程）共享同一個程序計(jì)數(shù)器（PC）。這意味著它們在代碼中的位置必須時(shí)刻保持一致。

如下圖所示：由于硬件需要串行執(zhí)行不同的代碼分支，導(dǎo)致一部分線程在另一部分執(zhí)行時(shí)只能空閑（Stall），造成了嚴(yán)重的并行效率損失。

Warp具體是怎么處理分支邏輯的呢？ 利用SIMT Stack記錄所有可能執(zhí)行路徑的上下文，遇到分支時(shí)，通過活動掩碼標(biāo)記需要執(zhí)行的活躍線程。當(dāng)前分支執(zhí)行完時(shí),硬件會去檢查SIMT Stack是否還有其他可執(zhí)行分支。最終所有分支執(zhí)行完成后，在匯合點(diǎn)（Reconvergence Point）恢復(fù)Warp中所有線程的執(zhí)行。

這里有個問題，如上圖，如果執(zhí)行B的時(shí)候因?yàn)榈却齼?nèi)存而暫停時(shí)，有沒有可能切到另外一個分支執(zhí)行X；Thread層面的隱藏延遲？

在Pre-Volta架構(gòu)中，答案是不能。因?yàn)檎麄€Warp共享一個程序計(jì)數(shù)器和狀態(tài)，需要為每個線程配備獨(dú)立的程序計(jì)數(shù)器（PC）和棧（Stack）。

2. Post-Volta Volta及后續(xù)架構(gòu)

Volta及后續(xù)架構(gòu)為每個線程配備獨(dú)立的程序計(jì)數(shù)器（PC）和棧（Stack）。

但是在任何時(shí)刻，Warp調(diào)度器還是只發(fā)布一條指令,即指令緩存（I-Cache）、指令獲取單元（Fetch）、指令解碼單元（Decode）都是Warp級別共享的。這意味著，盡管線程擁有獨(dú)立的PC，但一個Warp內(nèi)的線程不能在同一時(shí)鐘周期執(zhí)行不同的指令。

為什么不能讓一個Warp中的32個線程在同一時(shí)刻執(zhí)行32條不同的指令？ MIMD,multiple instruction, multiple thread， 恭喜你發(fā)明了多核cpu架構(gòu)。GPU的定位就是并行計(jì)算，沒必要搞MIMD；另外這樣搞導(dǎo)致硬件成本和功耗成本都大幅提升。算是硬件效率與執(zhí)行靈活性的一個trade-off。

這樣Volta及后續(xù)架構(gòu)，在Warp調(diào)度器同一時(shí)刻只發(fā)布一條指令的情況下，利用獨(dú)立程序計(jì)數(shù)器（PC）和活動掩碼（Active Mask）就可以實(shí)現(xiàn)智能調(diào)度。硬件通過在不同周期、用不同的“活動掩碼”來執(zhí)行不同的指令，巧妙地"編織"出了多線程獨(dú)立執(zhí)行的假象。說白了，就是當(dāng)一個Warp中的某些線程因?yàn)榈却齼?nèi)存操作而暫停時(shí)，調(diào)度器可以切換執(zhí)行同一個Warp下的其他線程，從而實(shí)現(xiàn)所謂的“線程級延遲隱藏”。實(shí)際上，這樣也難以避免Warp Divergence導(dǎo)致的算力浪費(fèi)，只是通過thread層面的隱藏延遲減少了部分因等待內(nèi)存而導(dǎo)致算力浪費(fèi)。

這里值得一提的是，獨(dú)立PC和Stack的引入同時(shí)也解決Pre-Volta架構(gòu)可能會死鎖的問題。(Pre-Volta架構(gòu)由于其剛性的SIMT執(zhí)行模型，在處理Warp內(nèi)部分線程依賴另一部分線程的場景時(shí)，易產(chǎn)生死鎖。)

(3) 同步機(jī)制

前面提到了Warp層面和thread層面的延遲隱藏，那當(dāng)我們Warp間或者同一個Warp中的不同thread間需要同步時(shí)，怎么辦呢？

• __syncthreads() 它保證一個Block內(nèi)的所有線程都執(zhí)行到這個Barriers后，才能一起繼續(xù)往下執(zhí)行。
• __syncwarp() 它保證一個Warp內(nèi)的32個線程都執(zhí)行到這個Barriers后，才能繼續(xù)往下執(zhí)行。

九、總結(jié)

至此，我們大體了解了AI Infra場景下GPU的工作流程與編程模式：

• 從圖形專用到GPGPU演進(jìn)
• CPU和GPU的協(xié)作通訊 命令緩沖區(qū)（Command Buffer）+ 門鈴（Doorbell）
• CUDA程序的生命周期
• CUDA的編程模型（Grid -> Block -> Thread）與GPU的硬件架構(gòu)（GPU -> SM -> Warp -> Core）
• SIMT通過線程編程模型隱藏了底層SIMD的執(zhí)行細(xì)節(jié)
• Warp層面和thread層面的延遲隱藏,以及各自層面的同步函數(shù)(__syncthreads() 和 __syncwarp())

本文旨在了解單GPU場景下的工作流程，然而AI Infra背景下，單GPU往往是夠用的，另外這里Cuda Streams、Unified Memory、MPS都沒提，留給后續(xù)填坑了。下一篇將詳細(xì)講解GPU的性能優(yōu)化相關(guān)知識。