本文將從 GPU 最核心的 SIMT 執(zhí)行模型出發(fā)，層層拆解 Grid/Block 配置、Warp 調(diào)度與 SM 占用率的深層關(guān)系，最終落地到內(nèi)存合并訪問的實(shí)戰(zhàn)優(yōu)化，
參考資料：《Programming MassivelyParallel Processors A Hands-on Approach》 Fourth Editing

一、SIMT vs SIMD：

要理解 GPU 的執(zhí)行邏輯，首先必須搞懂 GPU 的 SIMT（單指令多線程） 與我們熟知的 CPU SIMD（單指令多數(shù)據(jù)） 的核心區(qū)別。兩者都源于“單指令驅(qū)動(dòng)多份并行計(jì)算”的思想，但在抽象層次、硬件實(shí)現(xiàn)和編程體驗(yàn)上有著本質(zhì)差異。

1. 核心差異與編程體驗(yàn)

• SIMD（如 x86 AVX2 向量化指令）：一條指令直接操作一個(gè)“寬寄存器”中的多份連續(xù)數(shù)據(jù)。編程呈現(xiàn)為“寬寄存器”思維，程序員必須顯式操作批量數(shù)據(jù)、處理打包與解包。

• SIMT（GPU 執(zhí)行模型）：一條指令驅(qū)動(dòng)多個(gè)獨(dú)立的線程并行執(zhí)行，每個(gè)線程有自己的寄存器上下文、程序計(jì)數(shù)器。編程呈現(xiàn)為“多線程”思維，代碼邏輯幾乎與串行 C++ 一致，底層向量化完全由硬件自動(dòng)映射。

我們通過代碼直觀對(duì)比這種差異：

SIMD 代碼示例（需顯式操作寬寄存器）

// SIMD思維：顯式操作寬寄存器，一次處理8個(gè)float元素 (假設(shè)AVX2 256位)
void vector_add_simd(float* a, float* b, float* c, int n) {
    // 步長固定為8，對(duì)應(yīng)256位AVX寄存器的容量
    for (int i = 0; i < n; i += 8) {
        // 顯式加載256位寬數(shù)據(jù)到寄存器
        __m256 va = _mm256_load_ps(&a[i]);
        __m256 vb = _mm256_load_ps(&b[i]);
        // 單條指令完成8組數(shù)據(jù)的加法
        __m256 vc = _mm256_add_ps(va, vb);
        // 顯式將結(jié)果寫回內(nèi)存
        _mm256_store_ps(&c[i], vc);
    }
}

SIMT 代碼示例（符合直覺的多線程抽象）

// SIMT思維：多線程抽象，代碼邏輯與串行完全一致
__global__ void vector_add_simt(float* a, float* b, float* c, int n) {
    // 每個(gè)線程獲取自己的獨(dú)立全局索引
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    // 每個(gè)線程只處理自己對(duì)應(yīng)的1個(gè)元素
    if (gid < n) {
        c[gid] = a[gid] + b[gid];
    }
}
// 啟動(dòng)時(shí)無需手動(dòng)寫步長循環(huán)，只需啟動(dòng)足夠的線程：vector_add_simt<<<Grid, Block>>>(...);

2. 分支處理：SIMT 靈活性的核心體現(xiàn)

對(duì)于 SIMD 來說，遇到 if-else 分支時(shí)只能依賴掩碼屏蔽，兩個(gè)分支都得串行走一遍，毫無靈活性。 而在 SIMT 架構(gòu)中，同一個(gè) Warp（線程束，固定 32 個(gè)線程） 內(nèi)的線程可以獨(dú)立走不同分支。硬件能夠自動(dòng)標(biāo)記非活躍線程，僅執(zhí)行對(duì)應(yīng)分支的活躍線程。

// SIMT分支處理：代碼極度自然，硬件自動(dòng)處理活躍掩碼
__global__ void branch_demo(float* data, int n) {
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid >= n) return;
    
    // 哪怕在同一個(gè)Warp內(nèi)，這32個(gè)下屬線程也可以走完全不同的分支
    if (data[gid] > 0) {
        data[gid] = logf(data[gid]);
    } else {
        data[gid] = fabsf(data[gid]) + 1e-6;
    }
}

這種分支獨(dú)立能力，是 GPU 能從單純的“圖形渲染器”走向復(fù)雜通用并行計(jì)算（GPGPU）的核心基石。

二、Grid/Block 配置與 Warp 調(diào)度的映射邏輯

明白了 SIMT 是如何讓單個(gè)線程跑起來的，接下來必須厘清開發(fā)者配置好的 <<<Grid, Block>>> 是如何映射到 GPU 物理硬件上的。

1. 硬件執(zhí)行的層級(jí)關(guān)系

• Grid（全局線程網(wǎng)格）：決定了總?cè)蝿?wù)量，會(huì)被分配到整個(gè) GPU 上。
• Block（線程塊）：是分配給 SM（流式多處理器，GPU 核心計(jì)算單元） 的最小單位。同一個(gè) Block 內(nèi)的所有線程一定會(huì)被全部分配到同一個(gè) SM 上執(zhí)行，絕不會(huì)跨 SM 拆分。
• Warp（線程束，固定 32 個(gè)線程）：是 SM 調(diào)度執(zhí)行的最小物理單位。Block 被塞進(jìn) SM 后，會(huì)被硬件自動(dòng)切分為若干個(gè) Warp交由調(diào)度器管理。

2. Block 大小必須是 32 的整數(shù)倍

既然硬件只認(rèn) 32 人為一個(gè)物理調(diào)度組（Warp），那么 Block 的大小規(guī)劃就顯得至關(guān)重要。 如果開發(fā)者隨意將 Block 設(shè)為 40 線程，它會(huì)被強(qiáng)行拆分為 2 個(gè) Warp。第一個(gè) Warp 32 人滿載，第二個(gè) Warp 僅 8 人干活，其余 24 個(gè)物理線程坑位完全閑置，白白浪費(fèi) 75% 的算力資源。

三、延遲隱藏與多 Warp 調(diào)度

既然 Block 已經(jīng)扔給 SM 執(zhí)行了，為什么我們還需要極其龐大的總線程數(shù)量？這引出了 GPU 擁有超高吞吐量的核心秘密——延遲隱藏機(jī)制。

用切換代替等待

GPU 中不同存儲(chǔ)的延遲差異極大：寄存器訪問只需 1 個(gè)時(shí)鐘周期，而全局內(nèi)存（顯存）訪問高達(dá) 200~400 個(gè)時(shí)鐘周期。 如果只給 SM 分配等同于運(yùn)算核心數(shù)量的線程，一旦這些線程觸發(fā)了讀取全局內(nèi)存的指令，所有核心將陷入數(shù)百個(gè)周期的死等，利用率直接歸零。

為此，SM 采用了后廚訂單式的“超額分配”邏輯：給 SM 分配遠(yuǎn)多于其計(jì)算核心數(shù)的 Warp。 當(dāng)某個(gè) Warp 觸發(fā)長延遲訪存進(jìn)入等待狀態(tài)時(shí)，SM 會(huì)以零開銷瞬間切換到另一個(gè)就緒的 Warp 去執(zhí)行算術(shù)指令。用其他線程的“計(jì)算”來填補(bǔ)當(dāng)前線程的“等待時(shí)間”，從而讓核心時(shí)刻保持滿載。

四、衡量指標(biāo)：SM 占用率（Occupancy）與底層博弈

既然需要“足夠多”的 Warp 才能隱藏延遲，我們將「實(shí)際分配給 SM 的 Warp 數(shù)量」與「硬件支持的最大 Warp 數(shù)量」比值，定義為 SM 占用率 (Occupancy)。它是發(fā)揮 GPU 性能的必要非充分條件。

決定占用率能否拉滿的，是 SM 內(nèi)部四大物理資源的分配“木桶效應(yīng)”：

1. 線程槽與塊槽：以 A100 GPU 為例，單 SM 上限為 2048 個(gè)線程與 32 個(gè) Block 槽位。若開發(fā)者將 Block 極小化設(shè)為 32 線程，受限于 32 個(gè)塊槽上限，SM 最多只能容納 32×32=1024 個(gè)線程，由于塊槽先滿，剩余 50% 線程槽位直接作廢。

2. 寄存器 (Registers )的影響：A100 單 SM 有 65536 個(gè)寄存器。若要 2048 線程實(shí)現(xiàn) 100% 占用，單線程可用極限為 32 個(gè)寄存器（65536 ÷ 2048）。若你的代碼邏輯稍稍復(fù)雜，多加了 2 個(gè)局部變量導(dǎo)致單線程需 33 個(gè)寄存器，總線程數(shù)就會(huì)被硬生生降檔（不得不踢出部分 Block），占用率瞬間暴跌至 75%。

   ?? 溢出陷阱：當(dāng)使用編譯指令（如 -maxrregcount）強(qiáng)行壓低寄存器上限去換取占用率時(shí)，超出的變量會(huì)“溢出”到極慢的本地內(nèi)存，導(dǎo)致核函數(shù)整體耗時(shí)翻倍，這是典型的飲鴆止渴。

3. 共享內(nèi)存 (Shared Mem)：A100 單 SM 配置最高可達(dá) 164 KB 共享內(nèi)存。單塊申請(qǐng)的共享內(nèi)存過大，同樣會(huì)成為鎖死 SM 吞吐的短板。

五、榨干帶寬：內(nèi)存合并訪問（Memory Coalescing）

GPU 的全局內(nèi)存是以 128 字節(jié)的緩存行 為單位進(jìn)行交易的。合并訪問是指：同一個(gè) Warp 內(nèi)的 32 個(gè)線程訪問的物理地址是連續(xù)對(duì)齊的，這樣硬件僅需發(fā)起 1 次內(nèi)存事務(wù)即可讀滿 128 字節(jié)，帶寬利用率 100%。若地址嚴(yán)重分散，硬件將被迫發(fā)起多達(dá) 32 次獨(dú)立的內(nèi)存請(qǐng)求，帶寬效率不足 5%，并引發(fā)極其嚴(yán)重的排隊(duì)擁塞。

以下是實(shí)現(xiàn)合并訪問的 4 個(gè)核心代碼策略及正反面對(duì)比：

策略 1：保持線程 ID 與數(shù)據(jù)索引的線性映射

這是最基礎(chǔ)的原則：讓相鄰的線程，訪問相鄰的內(nèi)存元素。

? 正確寫法：線程ID直接對(duì)應(yīng)數(shù)據(jù)索引

__global__ void linear_access_kernel(float* data, int n) {
    // 線程gid直接訪問data[gid]，相鄰線程訪問相鄰的4字節(jié)地址
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid < n) {
        data[gid] = compute(data[gid]);
    }
}

? 錯(cuò)誤寫法：大步長跨越訪問（完全撕裂帶寬）

__global__ void stride_access_kernel(float* data, int n) {
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    // 相鄰線程地址間隔 400 字節(jié)，硬件被迫發(fā)起 32 次獨(dú)立內(nèi)存事務(wù)！
    int idx = gid * 100;
    if (idx < n) {
        data[idx] = compute(data[idx]);
    }
}

策略 2：優(yōu)先基于塊內(nèi)線程的線性處理

當(dāng)按塊分解任務(wù)時(shí)，優(yōu)先針對(duì) threadIdx.x（而非 blockIdx.x）實(shí)現(xiàn)連續(xù)索引，確保落在同一個(gè) Warp 里的 32 個(gè)線程地址連續(xù)。

? 正確寫法：塊內(nèi)連續(xù)映射

__global__ void block_linear_kernel(float* global_data, int n) {
    int block_start = blockIdx.x * blockDim.x;
    // 線程在塊內(nèi)的局部偏移 threadIdx.x 保證了連續(xù)性
    int gid = block_start + threadIdx.x;
    
    if (gid < n) {
        global_data[gid] = process(global_data[gid]);
    }
}

策略 3：使用寬類型結(jié)構(gòu) (float4) 合并提取

在處理多通道向量或圖像像素時(shí)，直接使用 CUDA 內(nèi)置類型 float2/float4 可以成倍提升每條取指指令的吞吐量，且完美保留了內(nèi)存合并的特性。

? 正確寫法：寬類型合并訪問

__global__ void float4_access_kernel(float4* data, int n) {
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid < n) {
        // 相鄰線程各訪問緊挨著的 16 字節(jié)，硬件自動(dòng)優(yōu)化為數(shù)次完整的 128 字節(jié)請(qǐng)求
        float4 val = data[gid]; 
        val.x = process(val.x);
        val.y = process(val.y);
        val.z = process(val.z);
        val.w = process(val.w);
        data[gid] = val; // 連續(xù)寫回
    }
}

策略 4：使用共享內(nèi)存梳理亂序訪問（Scatter 重組）

面對(duì)業(yè)務(wù)剛需的不可控隨機(jī)位置更新，怎么辦？答案是增加極低延遲的共享內(nèi)存作為“中轉(zhuǎn)緩沖”。

? 正確寫法：用 Shared Memory 吸收不規(guī)則訪問代價(jià)

__global__ void shared_mem_reorder_kernel(float* global_src, int* indices, float* global_dst, int n) {
    __shared__ float shared_buf[256];
    
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Step 1：100% 連續(xù)合并讀取全局內(nèi)存，將其存入共享緩沖
    if (gid < n) {
        shared_buf[tid] = global_src[gid];
    }
    __syncthreads(); // 等待同塊內(nèi)所有緩沖就緒
    
    // Step 2：在共享內(nèi)存內(nèi)做雜亂的隨機(jī)訪問（成本比讀全局內(nèi)存低 20 倍）
    if (gid < n) {
        int dst_idx = indices[gid]; 
        global_dst[dst_idx] = shared_buf[tid]; 
    }
}

?? 注意：務(wù)必在日常開發(fā)中自查并規(guī)避三大反面模式：大于 1 的步長訪問、基于隨機(jī)數(shù)組的不可見索引訪問、未對(duì)齊 128 字節(jié)邊界的錯(cuò)位偏移訪問。