📚 课程目录第4章 GPU指令集与执行模型

GPU指令集与执行模型：SIMT执行模型、warp调度、指令级并行（ILP）与线程级并行（TLP）

核心原理：SIMT——单指令流下的线程束伪装

GPU并非像CPU那样依赖复杂的乱序执行硬件来挖掘指令级并行，而是采用SIMT（单指令多线程）模型，将“并行计算”伪装成“顺序指令流”。在SIMT中，warp（线程束）是硬件调度的基本单位，通常由32个线程组成。这32个线程在同一时钟周期内执行同一条指令，但各自操作不同的数据。从程序员视角看，每个线程拥有独立的程序计数器（PC）和寄存器状态；但从硬件视角看，一个warp内的线程共享同一个取指与译码单元。

这种设计的核心优势在于：控制逻辑的摊薄。CPU为每个线程维护复杂的分支预测、乱序执行窗口，而GPU将32个线程的控制开销合并为一份，将晶体管预算几乎全部投入到ALU（算术逻辑单元）和寄存器文件中。代价则是：当warp内线程发生分支发散（即不同线程走不同分支）时，硬件必须串行执行每个分支路径，并屏蔽不活跃的线程，导致执行效率下降。

Warp调度机制与延迟隐藏

现代GPU的流处理器（SM）中通常包含4个或更多warp调度器，每个调度器负责管理一组warp。调度器在每个时钟周期内从就绪warp中选择一个，将其指令发射到执行单元。所谓“就绪”，是指该warp当前没有等待内存访问、同步障碍或固定延迟指令。这种机制的核心价值在于延迟隐藏：当某个warp因访问全局内存而阻塞时（典型延迟400~800个周期），调度器立即切换到其他warp执行计算，从而让ALU始终处于忙碌状态。

要达到理想的延迟隐藏效果，每个SM需要足够的活跃warp数量。以NVIDIA Ampere架构为例，一个SM最多可驻留64个warp（2048个线程）。如果每个线程仅使用少量寄存器，则能容纳更多warp；反之，寄存器压力过大会限制warp驻留数量，导致延迟隐藏能力下降。因此，寄存器分配与占用率（occupancy）之间存在直接权衡。

ILP与TLP：两种并行维度的协同

指令级并行（ILP）指在单个线程内，多条无依赖关系的指令可以并行执行。GPU的ILP主要来源于编译器优化——通过循环展开和指令重排，将独立的浮点运算或整数运算排列在一起，使流水线能够连续发射。例如，对以下循环进行4次展开：

// 原始循环
for (int i = 0; i < N; i++) {
    c[i] = a[i] * b[i] + c[i];
}

// 4次展开后（编译器自动生成）
for (int i = 0; i < N; i += 4) {
    c[i]   = a[i]   * b[i]   + c[i];
    c[i+1] = a[i+1] * b[i+1] + c[i+1];
    c[i+2] = a[i+2] * b[i+2] + c[i+2];
    c[i+3] = a[i+3] * b[i+3] + c[i+3];
}

展开后，4次乘加运算之间没有数据依赖，编译器可以将它们调度到不同的执行流水线中，实现每个周期发射多条指令。然而，GPU的ILP能力远不如CPU——CPU的乱序执行窗口可达数百条指令，而GPU主要依赖编译器静态调度，ILP窗口通常只有几条到十几条指令。

线程级并行（TLP）则是GPU的真正强项。通过创建成千上万个线程，GPU利用warp调度器在不同warp之间切换来隐藏延迟。TLP与ILP的关系可以概括为：TLP是GPU吞吐量的基石，ILP是补充手段。当TLP不足时（例如线程数太少或占用率低），提升ILP可以部分弥补延迟隐藏的缺口；当TLP充足时，过度追求ILP反而可能因寄存器占用增加而降低占用率。

案例：矩阵乘法的并行策略选择

以经典的SGEMM（单精度矩阵乘法）为例，假设矩阵尺寸为4096×4096。我们对比三种实现策略：

策略	线程组织	ILP程度	预期性能
朴素实现	每个线程计算1个输出元素	低（循环内串行乘加）	约20%峰值
ILP优化	每个线程计算4个输出元素（沿K维展开）	中（4路独立乘加）	约45%峰值
TLP+分块	16×16线程块，每线程计算8×8输出子块	高（8路乘加+寄存器分块）	约85%峰值

第三种策略之所以最优，是因为它同时利用了共享内存分块（减少全局内存访问）和寄存器分块（每个线程计算多个输出，提升数据复用率）。每个线程内8路独立的乘加指令提供了足够的ILP，而大量线程块保证了TLP。实际测试中，该策略在A100上可达到约90%的浮点峰值。

误区与复盘：常见性能陷阱

误区1：盲目追求100%占用率。高占用率并不总是等于高性能。当计算密度高且数据复用好时，较低的占用率（如50%）配合更高的ILP和寄存器分块，反而能减少缓存压力，获得更好性能。

误区2：忽略分支发散。在warp内使用if (threadIdx.x % 2 == 0)这类条件判断，会导致warp串行执行两个分支，性能减半。应尽量让分支条件基于warp粒度（如warp_id % 2），或使用算术方式替代分支。

误区3：过度循环展开。展开因子过大会导致寄存器溢出（spill），将局部变量存储到本地内存，反而引入数百周期的内存访问延迟。一般建议展开因子控制在2~8之间，具体需结合寄存器压力测试。

复盘建议：使用NVIDIA Nsight Compute分析kernel时，重点关注Warp Stall Sampling和Occupancy指标。如果发现“等待固定延迟指令”占比过高，说明ILP不足；如果“等待内存”占比高，则需增加TLP或优化数据复用。通过调整__launch_bounds__限制最大线程数，可以迫使编译器更激进地使用寄存器，从而提升ILP。

理解SIMT模型的关键在于：warp是硬件执行的真实单位，而线程只是编程抽象。性能优化的本质，就是让每个warp尽可能高效地执行，同时提供足够的warp来隐藏各种延迟。ILP和TLP并非非此即彼的选择，而是需要根据具体kernel的计算特性、内存访问模式和硬件资源限制，找到最优的平衡点。