作者 | binnnliu

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

一、圖形渲染到GPGPU

1. 為圖形而生

GPU最初的使命是加速圖形渲染。而渲染一幀圖像，本質(zhì)上就是對(duì)數(shù)百萬個(gè)像素點(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ù)界的探索

一批敏銳的研究人員意識(shí)到，GPU的本質(zhì)就是一個(gè)擁有數(shù)百甚至數(shù)千個(gè)核心的大規(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é)摿?，開始不再局限于加速圖形渲染，主動(dòng)擁抱GPGPU。

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

2007年，NVIDIA正式推出了CUDA平臺(tái)。CUDA的革命性在于，它提供了一套簡(jiǎn)單的編程模型，讓開發(fā)者能用近似C語言的方式，輕松地駕馭GPU內(nèi)部成百上千個(gè)并行核心。 開發(fā)者無需再關(guān)心復(fù)雜的圖形接口，可以直接編寫在數(shù)千個(gè)線程上并發(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是整個(gè)系統(tǒng)的核心，是總指揮，GPU的任務(wù)指令是由CPU分配的。

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

• MMIO（內(nèi)存映射I/O）由CPU直接控制數(shù)據(jù)讀寫，操作系統(tǒng)會(huì)把設(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通常會(huì)支持多內(nèi)存通道，即多IMC和Memory Channel的組合，以滿足日益增長的數(shù)據(jù)處理需求。

三、一個(gè)簡(jiǎn)單的應(yīng)用

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

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

實(shí)現(xiàn)兩個(gè)長度為 23? (約10億) 的浮點(diǎn)數(shù)數(shù)組的相加。其中，一個(gè)數(shù)組 (x) 的所有元素初始化為 1.0，另一個(gè)數(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í): 整個(gè)程序從啟動(dòng)到結(jié)束（real time）耗時(shí) 21.4秒。這額外的時(shí)間主要消耗在分配8GB內(nèi)存（new float[N]）以及初始化數(shù)組上。

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

這里的代碼后面會(huì)詳細(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上啟動(dòng)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í)的絕對(duì)優(yōu)勢(shì)。CPU需要串行執(zhí)行10億次加法(此處只考慮單核場(chǎng)景)，而GPU則將任務(wù)分配給成千上萬個(gè)線程同時(shí)處理。

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

四、編譯-Fat Binary

nvcc  add.cu -o add_cu -gencode arch=compute_75,code=sm_75 上面的例子中，我們看到這個(gè)編譯指令。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) 代號(hào) 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)一個(gè)可執(zhí)行文件中，形成所謂的胖二進(jìn)制 (Fat Binary)。它“胖”在包含了一份主機(jī)代碼和多份針對(duì)不同GPU架構(gòu)的設(shè)備代碼。

1. 程序加載 - cubin loading

(1) 程序啟動(dòng)

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

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

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

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

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

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

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

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

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

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

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

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

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

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

④ 上下文就緒

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

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

(1) 檢測(cè)硬件

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

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

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

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

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

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

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

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

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

一個(gè)常見的誤解是CPU會(huì)直接、實(shí)時(shí)地控制GPU。實(shí)際上，考慮到CPU和GPU是兩個(gè)獨(dú)立的處理器，并且通過PCIe總線連接，直接的、同步的控制會(huì)帶來巨大的延遲和性能開銷。因此，現(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可以理解為相對(duì)于 Ring Buffer 基地址 (Base Address) 的偏移量)

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

2. CPU 執(zhí)行到cudaMalloc

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

• CPU 線程調(diào)用 cudaMalloc()。CUDA 運(yùn)行時(shí)庫將此請(qǐng)求轉(zhuǎn)發(fā)給 NVIDIA 驅(qū)動(dòng)程序
• 驅(qū)動(dòng)程序向物理VRAM Allocator請(qǐng)求物理內(nèi)存，向 UVA Manager 請(qǐng)求虛擬地址，更新UVA映射表；（物理VRAM Allocator是跨進(jìn)程的，維護(hù)整個(gè)GPU 物理顯存的使用情況）
• 更新 GPU page table[Command Buffer + Doorbell方式,特定的、高優(yōu)先級(jí)的通道，非默認(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ū)動(dòng)分配物理內(nèi)存；而malloc本身是C庫函數(shù)(用戶態(tài)), 向操作系統(tǒng)“批發(fā)”大塊內(nèi)存，然后在用戶程序請(qǐng)求時(shí)“零售”出去。避免內(nèi)存分配時(shí)昂貴的系統(tǒng)調(diào)用和缺頁異常開銷

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

(3) 釋放策略

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

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

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

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

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

• 驅(qū)動(dòng)將Kernel函數(shù)啟動(dòng)所需信息打包成一個(gè)命令命令包括：?jiǎn)?dòng)Kernel函數(shù)，Kernel函數(shù)對(duì)應(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ì)列中取出命令包。
• 它開始解碼這個(gè)命令包，識(shí)別出這是一個(gè)“Kernel函數(shù)啟動(dòng)”任務(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 的維度，將整個(gè)計(jì)算任務(wù)分發(fā)成一個(gè)個(gè)獨(dú)立的Thread Blocks。
• GPU的全局調(diào)度器（GigaThread Engine）,將Thread Blocks分配給有空閑資源的 SM。一個(gè)線程塊從生到死都只會(huì)在一個(gè) SM 上執(zhí)行，不會(huì)遷移。

④ 線程塊調(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];
  }
}

• 每個(gè) SM 接收到一個(gè)或多個(gè)線程塊，SM 內(nèi)部的硬件調(diào)度器 (Scheduler)進(jìn)一步將每個(gè)線程塊內(nèi)部的線程，按照threadIdx的順序，每 32 個(gè)線程劃分成一個(gè) Warp。比如，一個(gè)有 256 個(gè)線程的線程塊，會(huì)被劃分為 8 個(gè) 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)一個(gè)線程塊完成了所有計(jì)算，它所占用的 SM 資源（如寄存器、共享內(nèi)存）會(huì)被釋放，SM 可以接收新的線程塊。
• 當(dāng)整個(gè) Grid 的所有線程塊都執(zhí)行完畢，這個(gè)Kernel函數(shù)啟動(dòng)任務(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,  一個(gè)GPU包含多個(gè)GPC, 一個(gè)GPC包含多個(gè)TPC

(2) TPC

Texture Processing Cluster, 一個(gè)TPC包含多個(gè)SM

(3) SM

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

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

一個(gè)硬件單元，專門處理**FMA(Fused Multiply-Add)**操作，能在一個(gè)時(shí)鐘周期內(nèi)完成一個(gè)小的矩陣乘加運(yùn)算（一個(gè)4x4的FP16矩陣相乘后累加到另一個(gè)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)鍵組件

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

2. 接口

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

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

其中HBM和L2 Cache是整個(gè)GPU共享的;

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

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

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

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

1. 編程模型

將一個(gè)待批量并發(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ù)局部性需求，比如，我要處理一個(gè)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: 每個(gè) Block 在 x 維度上有多少個(gè) Thread。（在我們例子中是256）。
• blockIdx.x: 當(dāng)前 Block 在 Grid 中的 x 坐標(biāo)。范圍是 0 到 gridDim.x - 1。
• gridDim.x: Grid 在 x 維度上有多少個(gè) Block。

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

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

而當(dāng)我們要處理一個(gè)二維矩陣或圖像時(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. 硬件層面

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

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

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

所以整個(gè)編程模型大概就是：

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

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

3. 隱藏延遲 - hide latency

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

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

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

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

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

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

(1) Registers

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

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

(2) Shared Memory

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

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

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

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

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

八、SIMD vs SIMT

前面CUDA Demo中我們已經(jīng)知道Kernel函數(shù)add會(huì)被啟動(dòng)成茫茫多的線程執(zhí)行,每個(gè)線程通過計(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的一條指令一次只能操作單個(gè)數(shù)據(jù)。例如，一次浮點(diǎn)加法就是double + double;

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

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

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

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

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

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

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

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

3. Warp Divergence

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

(1) Pre-Volta

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

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

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

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

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

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

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

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

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

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

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

(3) 同步機(jī)制

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

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

九、總結(jié)

至此，我們大體了解了AI Infra場(chǎng)景下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場(chǎng)景下的工作流程，然而AI Infra背景下，單GPU往往是夠用的，另外這里Cuda Streams、Unified Memory、MPS都沒提，留給后續(xù)填坑了。下一篇將詳細(xì)講解GPU的性能優(yōu)化相關(guān)知識(shí)。