2. GPU硬件执行模型

要理解批处理吞吐与延迟的平衡，必须先看清GPU的硬件执行模型。CPU擅长用少量复杂核心处理串行逻辑，而GPU的设计哲学是“用海量简单核心掩盖内存与指令延迟”。这个哲学落到硬件上，就是一组层层嵌套的结构：线程（thread）组成线程块（block），线程块被调度到流多处理器（SM）上执行，而SM内部又以warp为基本调度单位。与此同时，寄存器、共享内存、全局内存构成一条从快到慢、从小到大、从私有到共享的存储层次。批处理的吞吐与延迟，正是这些结构相互作用的结果。

一、SM：吞吐的物理载体

SM（Streaming Multiprocessor）是GPU的核心执行单元，一块A100有108个SM，一块消费级RTX 4090有128个SM。每个SM内部包含若干CUDA Core（整数/浮点运算单元）、Tensor Core、加载/存储单元（LSU）、特殊功能单元（SFU），以及寄存器文件、共享内存/L1缓存、warp调度器。SM的关键特征是“多发射、多驻留”：它同时驻留多个线程块，每个时钟周期从就绪的warp中挑选指令发射，从而在某个warp等待内存时切换到另一个warp继续计算。

因此，SM的吞吐不取决于单个线程跑得多快，而取决于“同一时刻有多少warp处于就绪状态”。如果驻留的warp太少，或者它们都在等内存，SM的运算单元就会空转，吞吐下降。这就是为什么“占用率（occupancy）”是GPU优化的核心指标之一。

二、warp：调度的最小单位

GPU以32个线程为一个warp进行调度。同一个warp内的32个线程执行同一条指令，但操作各自的数据——这就是SIMT（单指令多线程）模型。warp是延迟隐藏的基本单位：当warp A发出全局内存加载指令后需要等待数百个时钟周期，调度器会立刻切换到warp B、warp C继续执行，直到warp A的数据返回。只要就绪warp足够多，内存延迟就能被“藏”在计算背后。

但warp也带来一个代价：分支发散。如果warp内32个线程走不同的if分支，硬件只能串行执行各分支路径，吞吐按路径数成比例下降。例如一个warp中16个线程走if、16个走else，实际执行时间接近两倍。因此，让同一warp内的线程尽量走相同路径，是提升吞吐的重要原则。

三、线程块与资源分配

线程块是编程和资源分配的单位。一个block内的线程可以通过共享内存和__syncthreads()协作，但不同block之间默认无法直接通信。block被整体调度到一个SM上，其占用的寄存器、共享内存、线程槽位在生命周期内一直保留。这决定了SM能同时驻留多少个block：

每个SM可驻留block数 = min(
    线程槽位上限 / block线程数,
    寄存器文件大小 / (block线程数 × 每线程寄存器数),
    共享内存大小 / 每block共享内存,
    block数硬上限
)

举例：某SM有64K个32位寄存器、100KB共享内存、最多2048线程。若一个block有256线程、每线程用64个寄存器、每block用16KB共享内存，则寄存器限制为 65536/(256×64)=4个block，共享内存限制为 100/16≈6个block，线程限制为 2048/256=8个block，最终只能驻留4个block，即1024线程，占用率50%。若把每线程寄存器降到32个，寄存器限制变为8个block，占用率可达100%。这说明寄存器用量直接决定延迟隐藏能力。

四、内存层次：延迟与带宽的阶梯

GPU存储层次从近到远依次是：寄存器、共享内存/L1、L2、全局内存（DRAM）。寄存器访问约1个周期，共享内存约20~30周期，L2约200周期，全局内存约400~800周期。容量则相反：寄存器每SM仅256KB，共享内存约100~228KB，L2数十MB，全局内存数十GB。

这个阶梯决定了优化策略：把频繁访问的数据放进寄存器或共享内存，把大块数据留在全局内存但用合并访问（coalescing）提高有效带宽。所谓合并访问，是指同一warp的32个线程访问连续地址，硬件可将其合并为一次128字节事务；若地址分散，则退化为多次事务，有效带宽成倍下降。

五、结构如何共同决定吞吐与延迟

把上述结构串起来看，批处理的吞吐与延迟平衡可以这样理解：

• 延迟来自内存层次：全局内存访问需要数百周期，这是无法消除的物理事实。批处理的价值在于用大量独立任务填满这些等待窗口。
• 吞吐来自warp级并行：SM靠驻留多个warp，在某个warp等待时切换执行，把延迟“藏”起来。驻留warp越多，隐藏能力越强。
• 驻留数受资源约束：寄存器、共享内存、线程槽位三者取最小值。寄存器用量过高会直接压低占用率，进而削弱延迟隐藏。
• 批大小影响并行度：批太小，总线程数不足以填满SM；批太大，单次处理时间变长，端到端延迟上升。平衡点在于让SM保持足够驻留warp，同时不让单个请求等待过久。

六、一个可操作的参数检查示例

假设你在调优一个矩阵乘kernel，block设为256线程，每线程使用80个寄存器，每block共享内存24KB。目标SM有64K寄存器、100KB共享内存、2048线程上限。计算驻留数：

寄存器限制：65536 / (256 × 80) = 3.2 → 3 个block
共享内存限制：100 / 24 = 4.16 → 4 个block
线程限制：2048 / 256 = 8 个block
实际驻留：min(3, 4, 8) = 3 个block = 768 线程
占用率：768 / 2048 = 37.5%

此时瓶颈在寄存器。若通过减少临时变量把每线程寄存器降到64个，寄存器限制变为 65536/(256×64)=4 个block，占用率升到50%；若进一步降到48个，限制变为5个block，占用率62.5%。每降低一档寄存器用量，就多一个block的延迟隐藏能力。这就是“寄存器—占用率—吞吐”的直接传导链。

反过来，如果共享内存用量从24KB增到40KB，共享内存限制变为2个block，即使寄存器降到48个，实际驻留也只有2个block，占用率反而下降。这说明资源之间是相互制约的，优化时必须同时观察三项限制，而不是只盯一个。

七、小结

SM是吞吐的物理载体，warp是延迟隐藏的调度单位，线程块是资源分配的单位，而寄存器、共享内存、全局内存构成延迟与容量的阶梯。批处理吞吐与延迟的平衡，本质上是在这些约束下寻找一个驻留warp足够多、单请求等待又不过长的配置点。理解这套模型，后续的占用率调优、内存访问优化和批大小选择才有依据。