📚 课程目录第2章 GPU架构与CUDA编程模型基础

GPU架构与CUDA编程模型基础：SM、线程层次、内存层次、warp调度

在进入CUDA kernel与TensorRT推理瓶颈定位之前，必须先建立对GPU硬件架构与CUDA软件模型的清晰认知。很多性能分析之所以无从下手，根源在于开发者将GPU视为一个“黑盒”，不清楚kernel中的线程如何在硬件上被调度、数据如何在不同层级的内存间流动、以及SM（Streaming Multiprocessor，流式多处理器）的资源分配如何成为瓶颈。本章将从SM的物理结构出发，串联线程层次、内存层次与warp调度机制，为后续的kernel剖析和TensorRT层级调优打下坚实的硬件与软件映射基础。

2.1 SM（流式多处理器）：GPU计算的核心单元

核心原理

SM是GPU硬件架构中最关键的计算单元。以NVIDIA Ampere架构的A100为例，一颗GA100芯片集成了108个SM，每个SM包含64个FP32 CUDA核心、32个FP64核心、4个第三代Tensor Core，以及4个纹理单元。SM本质上是一个独立的计算“小处理器”，拥有自己的指令调度器、寄存器堆、共享内存和L1缓存。GPU执行计算任务时，并不是整个芯片统一调度，而是以SM为单位并行运行多个线程块（Block）。

理解SM需要抓住三个关键点：第一，SM是资源分配的基本单位，一个Block只能被调度到一个SM上执行，不能跨SM拆分；第二，SM内部的CUDA核心数量决定了其单周期指令吞吐能力；第三，SM的寄存器堆和共享内存容量决定了能同时驻留多少线程——这直接关系到kernel的occupancy（占用率）。

判断方法：如何评估SM资源瓶颈

在定位kernel性能瓶颈时，可以通过NVIDIA Nsight Compute或CUDA Runtime API获取SM的关键指标：

SM占用率（Achieved Occupancy）：实际活跃warp数量与理论最大warp数量的比值。低于50%通常意味着资源受限。
寄存器溢出（Register Spilling）：当每个线程需要的寄存器数量超过SM寄存器堆可分配量时，数据会溢出到本地内存（Local Memory），导致性能断崖式下降。
共享内存Bank冲突（Bank Conflict）：多个线程同时访问共享内存同一Bank时产生串行化，降低有效带宽。

// 查询设备SM属性
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("SM数量: %d\n", prop.multiProcessorCount);
printf("每个SM最大线程数: %d\n", prop.maxThreadsPerMultiProcessor);
printf("共享内存/SM: %zu bytes\n", prop.sharedMemPerMultiprocessor);

案例拆解：一个典型的SM资源受限场景

假设我们在A100上运行一个自定义的卷积kernel，配置为每个Block 256线程，每个线程使用40个寄存器。A100每个SM有65536个寄存器，那么每个SM最多可驻留的线程数为 65536 / 40 = 1638，即约6个Block（1638/256≈6.4）。此时SM的线程占用率为 1638 / 2048 = 80%，看起来尚可。但如果将每个线程的寄存器数提升到64个，则每个SM最多驻留1024个线程，占用率降至50%。此时kernel的访存延迟无法被有效隐藏，SM的调度器经常处于空闲等待状态，GPU利用率大幅下降。

经验提醒：在TensorRT中，层融合（Layer Fusion）有时会增加寄存器压力。如果发现融合后kernel的occupancy明显下降，需要权衡计算密度提升与占用率下降之间的得失。通常建议保持occupancy在75%以上。

2.2 线程层次：Grid-Block-Warp-Thread的映射关系

核心原理

CUDA的线程层次是软件抽象，与硬件SM结构存在明确的映射关系。从大到小依次为：Grid（网格） → Block（线程块） → Warp（线程束） → Thread（线程）。一个Grid包含多个Block，一个Block包含多个线程（通常为32的倍数），而GPU硬件实际调度的最小单位是Warp——32个线程组成一个Warp，它们在同一时刻执行同一条指令（SIMT模型，单指令多线程）。

这种层次设计的意义在于：Block是软件层面的资源分配单位，一个Block内的线程可以通过共享内存和同步屏障（__syncthreads()）协作；Warp是硬件层面的执行调度单位，SM的warp调度器每个周期从就绪的warp中选择一个发射指令。理解这一映射关系，是优化kernel的关键——Block大小设置不当会导致warp数量不足或资源浪费。

操作流程：合理配置线程层次

确定Block大小：通常设置为128~512线程，且为32的倍数。过小（如32）会导致Block数量过多，调度开销增大；过大（如1024）会限制每个SM可驻留的Block数量。
计算Grid大小：根据数据总量和Block大小计算，确保覆盖所有数据元素。
检查warp数量：每个SM的warp数量 = Block内warp数 × 驻留Block数。A100每个SM最大64个warp（2048线程）。
验证资源限制：使用cudaOccupancyMaxActiveBlocksPerMultiprocessor()函数计算理论最大驻留Block数。

// 计算最优Block大小
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
kernel<<<numBlocks, blockSize>>>(d_data, N);

常见误区

风险提示：很多开发者误以为Block越大越好。实际上，当Block超过512线程时，每个SM能驻留的Block数量减少，可能导致warp调度灵活性下降。例如，一个1024线程的Block在A100上只能驻留2个Block（共64个warp），而4个256线程的Block同样占满64个warp，但后者在Block边界上的同步开销更小，且能更灵活地填充SM资源。此外，Block内线程数不是32的倍数时，最后一个warp会有空闲线程，浪费计算资源。

2.3 内存层次：从寄存器到全局内存的延迟与带宽

核心原理

GPU内存层次的设计目标是：用高速小容量缓存掩盖低速大容量显存的访问延迟。从快到慢依次为：寄存器（Register） → 共享内存（Shared Memory） → L1缓存 → L2缓存 → 全局内存（Global Memory）。各层级的关键差异如下表所示：

内存类型	作用域	延迟	容量/SM	生命周期
寄存器	线程私有	<1ns	65536×32bit	线程执行期间
共享内存	Block内共享	~20ns	最多228KB	Block执行期间
L1缓存	SM内	~30ns	256KB	自动管理
L2缓存	芯片级	~200ns	40MB (A100)	自动管理
全局内存	所有线程	~400-800ns	40GB (HBM2e)	程序运行期间

在TensorRT推理中，全局内存带宽往往是首要瓶颈。例如，一个FP16精度的BERT推理，模型权重约400MB，如果全部从全局内存读取，在A100的2TB/s带宽下，仅权重加载就需要200μs。因此，数据复用是优化的核心——通过共享内存和寄存器缓存反复使用的数据，减少全局内存访问次数。

判断方法：识别内存瓶颈

使用Nsight Compute的Memory Workload Analysis可以快速判断：

Memory Throughput：如果超过80%，说明kernel受内存带宽限制。
L1/TEX Hit Rate：低于50%时，考虑优化数据局部性。
Shared Memory Bank Conflicts：冲突率超过10%时，需要调整数据布局或访问模式。

案例拆解：共享内存优化矩阵转置

以经典的矩阵转置为例（2048×2048 float矩阵），朴素实现直接读写全局内存，每个线程处理一个元素。由于转置的访存模式导致非合并访问（读行写列），全局内存带宽利用率极低。优化方案是使用共享内存作为中转：

__global__ void transpose_shared(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE];
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    
    // 合并读入共享内存
    if (x < n && y < n) tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    __syncthreads();
    
    // 转置写出
    int x2 = blockIdx.y * TILE + threadIdx.x;
    int y2 = blockIdx.x * TILE + threadIdx.y;
    if (x2 < n && y2 < n) out[y2 * n + x2] = tile[threadIdx.x][threadIdx.y];
}

该优化将全局内存访问次数从2次（读+写）降为2次但均为合并访问，实测带宽利用率从35%提升到85%以上。在TensorRT中，类似的tiling策略被广泛用于卷积和矩阵乘的kernel实现中。

2.4 Warp调度与SIMT执行模型

核心原理

Warp是GPU执行指令的最小硬件单位，由32个线程组成。SM中的warp调度器负责在每个时钟周期从多个就绪的warp中选择一个发射指令。由于所有线程共享程序计数器（PC），它们必须执行同一条指令——这就是SIMT（单指令多线程）模型。当warp内线程出现分支分歧（如if-else语句）时，调度器会串行执行每个分支路径，导致部分线程空闲，称为warp divergence。

warp调度的核心价值在于延迟隐藏：当某个warp因访问全局内存而等待时，调度器立即切换到另一个就绪的warp执行计算，从而保持SM的CUDA核心始终忙碌。因此，SM中驻留的warp数量越多，延迟隐藏能力越强。

操作流程：分析warp调度效率

检查warp占用率：通过Nsight Compute的Scheduler Stats查看每个周期发射的指令数。
识别分支分歧：使用Branch Efficiency指标，低于100%说明存在分歧。
分析活跃warp数：Active Warps per Scheduler，A100每个SM有4个调度器，每个调度器管理16个warp。
优化分支结构：将数据按条件预排序，或使用算术方式替代分支。

常见误区与复盘

  关键认知：warp divergence并不总是坏事。如果分支的两个路径计算量都很大，且数据分布均匀，串行执行两个路径的总时间可能接近并行执行的时间。真正的性能杀手是极度不平衡的分支——例如99%的线程走if路径，1%走else路径，else路径的warp几乎完全空闲。在TensorRT中，动态形状（Dynamic Shape）可能导致不同warp处理不同长度的序列，加剧divergence。解决方案是使用固定长度padding或对数据按长度排序后分桶处理。

另一个常见误区是忽略warp同步的代价。__syncthreads()会阻塞Block内所有warp直到全部到达屏障点。如果Block内warp数量过多（如32个warp），同步开销会显著增加。在TensorRT的plugin开发中，应尽量减少同步点，或使用更细粒度的同步原语（如cooperative groups）。

2.5 综合案例：从硬件视角定位TensorRT推理瓶颈

假设我们使用TensorRT部署一个ResNet-50模型，在A100上测得推理延迟为1.2ms，但理论计算时间仅需0.3ms。通过Nsight Systems分析发现，GPU利用率仅为45%，且kernel之间存在明显的空闲间隙。进一步使用Nsight Compute分析最耗时的卷积kernel，发现其occupancy仅为35%，原因是每个线程使用了128个寄存器，导致SM驻留线程数严重不足。

针对该瓶颈，我们采取以下步骤：第一，在TensorRT中设置BuilderFlag::kFP16启用FP16精度，减少寄存器压力；第二，使用IAutotuner搜索更优的tile大小和block配置；第三，将多个小卷积层融合为一个大kernel，减少kernel启动开销和全局内存往返。优化后，occupancy提升至72%，推理延迟降至0.6ms，GPU利用率达到78%。

这个案例说明：硬件架构知识是性能分析的基石。如果不理解SM资源限制和warp调度机制，很难定位到寄存器溢出这一根因，更无法制定有效的优化策略。

总结提醒：本章的核心知识点可归纳为三点：SM是资源分配与执行的基本单位，Block映射到SM、Warp映射到调度器；内存层次决定了数据访问的代价，优化目标是最大化寄存器/共享内存复用，最小化全局内存访问；warp调度是延迟隐藏的机制，保持高occupancy是kernel高效执行的前提。在后续章节中，我们将基于这些硬件知识，深入剖析CUDA kernel的性能分析方法与TensorRT的层融合、精度校准等优化技术。