2. GPU计算基础

在讨论大规模 GPU 集群推理服务性能优化之前，必须先建立对 GPU 计算底层的准确认知。推理服务的延迟、吞吐与显存占用，本质上都是由 GPU 的硬件架构、存储层次和并行执行模型共同决定的。本章不讨论具体框架或算子库，而是聚焦于硬件与执行模型本身，为后续的算子优化、显存规划与多卡调度建立共同语言。

2.1 GPU 架构的核心组织：从 SM 到线程

现代 NVIDIA GPU（Volta 及以后架构）的基本计算单元是 流式多处理器（Streaming Multiprocessor, SM）。一个 GPU 由数十到上百个 SM 组成，每个 SM 内部包含若干 CUDA 核心（即 FP32 算术单元）、Tensor Core、特殊函数单元（SFU）、共享内存（Shared Memory）、寄存器文件（Register File）以及调度器（Scheduler）。

理解 GPU 架构的关键在于：它不是“很多个 CPU 核心的简单堆叠”，而是一个以“线程束（Warp）”为调度单位的 SIMT（单指令多线程）执行机器。每个 Warp 包含 32 个线程，这 32 个线程在同一时刻执行同一条指令，但操作各自不同的数据。如果一个 Warp 内出现分支发散（即部分线程走 if，部分走 else），则两个分支会被串行执行，造成性能损失。

以 A100 GPU 为例，其关键规格如下表所示：

参数	A100 (GA100)	H100 (GH100)	对推理的影响
SM 数量	108	132	并行度上限
每 SM 最大线程数	2048	2048	决定可驻留线程总量
每 SM 共享内存（可配置）	164 KB	228 KB	决定算子内数据复用能力
寄存器文件 / SM	64K × 32-bit	64K × 32-bit	限制每线程寄存器用量
Tensor Core 能力	第三代	第四代（支持 FP8）	矩阵乘法的吞吐关键

在推理场景中，SM 的数量决定了 GPU 能同时容纳多少计算任务。如果某个算子的线程块（Block）数量远小于 SM 数量，则 GPU 利用率会明显不足。反之，如果线程块过多但每个块占用资源过大，则可能因寄存器或共享内存不足导致驻留率下降。

2.2 显存层次结构与数据搬运成本

GPU 的存储体系分为多个层级，从寄存器、共享内存、L1 缓存、L2 缓存，到显存（HBM）乃至主机端内存。每一级的容量和带宽差异巨大，而推理性能优化的核心之一，就是让数据尽可能停留在靠近计算单元的层级。

以下是一个典型的数据访问延迟与带宽对比（以 A100 为参考）：

• 寄存器（Register）：容量极小（每线程最多 255 个），访问延迟约 1 个时钟周期，带宽最高。
• 共享内存（Shared Memory）：每 SM 最多 164 KB，延迟约 20-30 周期，带宽约为显存的 10 倍以上。需要显式管理，用于线程块内数据交换。
• L1 缓存：与共享内存共用物理存储，延迟约 30-40 周期，对全局内存访问透明缓存。
• L2 缓存：全 GPU 共享，容量 40 MB（A100），延迟约 200 周期，用于跨 SM 的数据共享。
• 显存（HBM2e/HBM3）：容量 40-80 GB，带宽约 2-3 TB/s，但延迟高达 400-800 周期。这是绝大多数推理模型参数和中间激活值的存放位置。
• 主机内存（PCIe/NVLink）：带宽仅 32 GB/s（PCIe Gen4）或 600 GB/s（NVLink），延迟在微秒级，通常只用于数据初始化或超大模型无法完全驻留显存时的交换。

在推理服务中，显存带宽往往是比算力更先触达的瓶颈。例如，一个 7B 参数的 FP16 模型权重占用约 14 GB 显存。在自回归解码阶段，每个 token 的生成都需要读取全部权重参数。假设批量大小为 1，生成一个 token 需要读取 14 GB 数据，以 A100 的 2 TB/s 带宽计算，理论最短时间为 7 毫秒——这远高于实际计算所需的时间。因此，减少权重读取量（如量化、批处理）或提高数据复用（如连续批处理）是推理优化的核心方向。

2.3 并行计算模型：Grid-Block-Thread 与执行流程

CUDA 的并行模型采用三层结构：Grid（网格）→ Block（线程块）→ Thread（线程）。一个 Kernel 启动时，会创建一个 Grid，Grid 由多个 Block 组成，每个 Block 包含多个 Thread（通常为 128 到 1024 个）。Block 内的线程可以通过共享内存和同步屏障（__syncthreads()）协作，而不同 Block 之间无法直接通信，只能通过全局内存。

硬件调度时，Block 被分配到 SM 上执行。一个 SM 可以同时驻留多个 Block，但受限于线程数、寄存器和共享内存总量。当某个 Block 因等待显存访问而阻塞时，SM 会切换到其他驻留的 Warp 执行，以此隐藏内存延迟。这种机制称为 延迟隐藏（Latency Hiding），是 GPU 高吞吐的关键。

一个典型的 Kernel 执行流程如下：

// 以向量加法为例，展示 Grid-Block-Thread 映射
__global__ void vectorAdd(const float* A, const float* B, float* C, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        C[idx] = A[idx] + B[idx];
    }
}

// 启动配置：256 线程/块，共 n/256 个块
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(A, B, C, n);

在推理框架（如 TensorRT、vLLM）中，算子库（如 cuBLAS、cuDNN）已经封装了底层的线程调度。但理解这一模型有助于解释为什么某些算子（如 LayerNorm）在 batch 较小时性能不佳——因为其计算量不足以填满 SM 的并行度，导致延迟隐藏效果差。

2.4 典型推理场景中的并行特征分析

以 GPT 类模型的推理为例，其计算过程可分为两个阶段：预填充（Prefill）和解码（Decode）。这两个阶段对 GPU 资源的利用模式截然不同。

• Prefill 阶段：处理输入序列（如 512 个 token），计算量密集，属于典型的计算密集型（Compute-bound）任务。此时矩阵乘法（GEMM）占据主导，Tensor Core 利用率高，SM 的算术单元被充分使用。
• Decode 阶段：逐 token 生成输出，每次只处理 1 个 token 的向量。此时权重矩阵的读取量远大于计算量，属于访存密集型（Memory-bound）任务。GPU 的算力大量空闲，而显存带宽成为瓶颈。

这种差异直接影响了优化策略：Prefill 阶段应关注 Tensor Core 的利用率和矩阵分块策略；Decode 阶段则应关注权重数据的复用（如增大 batch 以摊销权重读取）和减少访存次数（如 KV Cache 的紧凑存储）。

2.5 常见误区与性能陷阱

在 GPU 推理优化实践中，以下误区经常导致性能不达预期：

2.6 性能分析工具与关键指标

在实际优化中，应使用 NVIDIA 提供的性能分析工具定位瓶颈。以下是最常用的工具和指标：

• ncu（Nsight Compute）：用于 Kernel 级别的详细分析，可查看 SM 吞吐、访存吞吐、Warp 状态、寄存器溢出等。
• nsys（Nsight Systems）：用于应用级时间线分析，查看 Kernel 执行顺序、GPU 空闲时间、数据拷贝耗时。
• DCGM（Data Center GPU Manager）：用于集群级监控，采集 GPU 利用率、显存使用、温度、功耗等。

关键性能指标包括：

指标	含义	参考阈值
SM 吞吐（Compute Throughput）	SM 内算术单元的平均利用率	Prefill 阶段应 > 60%
显存带宽利用率	实际访存带宽 / 理论峰值带宽	Decode 阶段通常 > 80%
Warp 驻留率（Occupancy）	活跃 Warp 数 / 最大可驻留 Warp 数	建议 > 50%，但需结合访存比判断
寄存器溢出（Spill）	寄存器不足时数据被挤入本地内存	应为 0，否则性能骤降

在分析推理服务时，建议先使用 nsys 查看整体时间分布，确认瓶颈是 Kernel 计算、数据拷贝还是 CPU 端调度；再使用 ncu 对耗时 Top 3 的 Kernel 做深入分析，检查其是计算受限还是访存受限。这一流程可以避免盲目优化。