本页速览 GPU 是 LLM 推理的主力硬件——SM、warp、Tensor Core、HBM、L2 cache 的体系结构决定了一切优化方向。本文讲透 GPU 编程模型、各代 GPU 算力与带宽参数、性能调优 checklist（occupancy、coalesced access、bank conflict、async memcpy、pipelining）、cudnn/cublas/cutlass 库。

GPU 体系结构与优化

概念定义：GPU 不是"大 CPU"

GPU 的设计哲学与 CPU 完全不同——CPU 优化低延迟（单线程跑超快），GPU 优化高吞吐（成千上万线程并行）。理解 GPU 体系结构是理解所有推理优化的基础——显存层次、Roofline、算子融合、权重量化都建立在这套体系之上。

理解 GPU 的三个关键认知：

GPU 是众核架构——单线程比 CPU 慢，但成千上万个线程并行后总吞吐远超 CPU；
GPU 用带宽换延迟——单次访存延迟 200-400 cycle（CPU 50 cycle），但带宽 3-8 TB/s；
GPU 的算力来源是 Tensor Core——不是普通 ALU，而是矩阵计算专用单元（每个时钟周期算一个 8×8×8 矩阵乘加）。

LLM 推理工程师必须懂 GPU——否则没法解释"为什么量化能加速 4 倍"或"为什么 batch 加大吞吐能 30 倍"。

一、GPU 体系结构

从晶圆到 SM 到 Warp 到 Thread

text

┌─────────────────────────────────────────────────────────┐
│  GPU（如 H100）                                          │
│  ┌──────────────────────────────────────────────┐       │
│  │  SM 0   SM 1   SM 2  ... SM 131（132 个 SM）│       │
│  │  ┌──────────────┐                            │       │
│  │  │ SM 内部:     │                            │       │
│  │  │ - 4 个 Block │                            │       │
│  │  │ - 64 个 Warp │                            │       │
│  │  │ - 2048 thread│                            │       │
│  │  │ - Tensor Core│                            │       │
│  │  │ - 228KB SRAM │                            │       │
│  │  └──────────────┘                            │       │
│  └──────────────────────────────────────────────┘       │
│  ┌──────────┐  ┌──────────┐                            │
│  │ L2 Cache │  │ HBM      │                            │
│  │ 50 MB    │  │ 80 GB    │                            │
│  └──────────┘  └──────────┘                            │
└─────────────────────────────────────────────────────────┘

层	数量（H100）	资源	角色
GPU	1	全 SM + L2 + HBM	整体调度
SM	132	4 个 warp scheduler + Tensor Core + SRAM	流多处理器，单 SM 跑 1 个或多个 block
Block	软件定义	多个 warp + shared memory	block 内可同步、共享 SRAM
Warp	32 thread	单指令多数据（SIMT）	GPU 的最小执行单元，32 thread 同步执行
Thread	单执行单元	寄存器	算单元素

Warp 是最小执行单元

GPU 不能调度"单个 thread"——一个 warp 的 32 个 thread 同步执行相同指令。如果有分支（if/else）让 32 个 thread 走不同路径，所有 thread 都执行所有分支，最后按 mask 选择结果——这就是 warp divergence（分支发散），是性能杀手。写 kernel 要尽量减少 warp 内部 branch。

SM 的关键资源

每个 SM（H100）有：

64 个 warp scheduler（每时钟周期可发射 4 条指令）；
4 个 Tensor Core 单元（每个时钟周期算一个 16×8×16 matmul）；
228 KB shared memory（可配置成不同比例的 SRAM/L1）；
256 KB L1 cache（包括上面 228 KB）；
65536 个 32-bit 寄存器（per SM）。

二、Tensor Core：算力的来源

Tensor Core（Volta 起引入）是 GPU 的矩阵计算专用单元——比普通 CUDA Core 快几十倍。

计算能力

一个 Tensor Core 在一个时钟周期内算 D = A × B + C（A 是 m×k、B 是 k×n、C/D 是 m×n 的矩阵）。不同精度下：

精度	Tensor Core shape	单 op FLOPS
FP64	8×8×4	256 FLOPS
FP32 (TF32)	8×8×8	512 FLOPS
FP16/BF16	8×8×16	1024 FLOPS
INT8	8×8×16	2048 OPS
FP8 (H100+)	8×8×32	4096 FLOPS
FP4 (B200+)	32×32×32	32768 FLOPS

为什么 Tensor Core 比普通 CUDA Core 快几十倍

普通 CUDA Core 是标量 ALU（单时钟算 1-2 个 FLOPS）。Tensor Core 是矩阵专用单元（单时钟算 1024 个 FLOPS）——差距来自数据复用：传统 CUDA Core 算 C[i,j] = ΣA[i,k]·B[k,j]，每个元素 A 和 B 要读多次；Tensor Core 算一个 8×8×16 矩阵，A/B 各读一次就够——算术强度从 ~1 提升到 ~50，正好跑满 Roofline 算力上限。

算力计算

H100 SXM5：

132 SM × 4 Tensor Core/SM = 528 Tensor Core；
每时钟周期算 528 × 1024 = 540,672 FLOPS（BF16）；
主频 1.83 GHz → 989 TFLOPS（540,672 × 1.83 GHz ≈ 989 TF）。

三、各代 GPU 关键参数

GPU	架构	SM 数	主频	BF16 算力	INT8 算力	HBM 容量	HBM 带宽	FP8 算力
A100 80GB	Ampere	108	1.41 GHz	312 TF	624 TOPS	80 GB	2.0 TB/s	—
H100 SXM5	Hopper	132	1.83 GHz	989 TF	1979 TOPS	80 GB	3.35 TB/s	1979 TF
H200 SXM	Hopper	132	1.83 GHz	989 TF	1979 TOPS	141 GB	4.8 TB/s	1979 TF
H20	Hopper	78	1.5 GHz	148 TF	296 TOPS	96 GB	4.0 TB/s	296 TF
L40S	Ada	142	1.11 GHz	362 TF	733 TOPS	48 GB	0.866 TB/s	733 TF
B200 SXM	Blackwell	168	2.0 GHz	2.25 PF (FP4)	—	192 GB	8.0 TB/s	4.5 PF
B100	Blackwell	128	1.95 GHz	1.8 PF (FP4)	—	192 GB	8.0 TB/s	3.6 PF

LLM 推理选 GPU 的实战速查

大模型在线服务（70B+）：H200（带宽 4.8TB/s + 141GB 显存）
中等模型在线服务（7B-30B）：H100 SXM5 或 H20
离线批处理：A100 80GB（性价比仍不错）
边缘 / 中小模型：L40S（48GB，性价比）
极致算力 / FP4：B200（Blackwell 旗舰）
A100 升级路径：A100 → H200 是 LLM 推理最优升级路径（带宽 + 显存同时跳）

四、CUDA 编程模型回顾

grid / block / thread 三级

cuda

// CUDA kernel 调用
my_kernel<<<numBlocks, threadsPerBlock>>>(arg1, arg2);
//                ↑               ↑
//              grid            block
//              (3D)            (3D)

grid：所有 block 的集合，跨 SM 调度；
block：单 SM 内执行，可同步、共享 shared memory；
thread：单执行单元，有寄存器。

CUDA C++ kernel 模板

cuda

__global__ void vector_add(const float* a, const float* b, float* c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) c[i] = a[i] + b[i];
}
// 调用：
vector_add<<<ceil(N/256.0), 256>>>(d_a, d_b, d_c, N);

Shared Memory：block 内 SRAM

cuda

__shared__ float tile[16][16];   // block 内共享
// 多个 thread 协作把数据从 HBM 加载到 shared memory
tile[ty][tx] = a[row * K + tx];
__syncthreads();                   // 等 block 内所有 thread 都完成加载
// 后续计算反复用 tile，不再读 HBM

Shared memory tiling 的本质

是手动 cache 管理——程序员决定哪些数据进 SRAM、何时使用、何时丢。CPU 的 cache 是硬件自动管理，GPU 的 shared memory 是软件管理——程序员的自由度大、责任也大（管理不好反而慢）。

Warp shuffle：32 thread 间超快通信

cuda

// 32 个 thread 间传数据，不走 shared memory
float my_val = ...;
float other_val = __shfl_xor_sync(0xffffffff, my_val, 1);  // 与相邻 thread 交换

warp shuffle 比 shared memory 快 2-3 倍，常用于 reduction（求和、求最大值）。

五、性能调优 Checklist

写完 kernel 还要调优——这是 CUDA 工程师的主战场。NVIDIA 推荐的 checklist：

1. Occupancy（占用率）

occupancy = 活跃 warp 数 / SM 最大 warp 数。低 occupancy → SM 算力闲置。提高方法：

减少 per-thread 资源使用（寄存器、shared memory）；
增加 grid 数量（让 SM 跑多个 block）；
调 launch bounds（__launch_bounds__ 限制寄存器数量）。

但 occupancy 不是越高越好——80% 以上时再提高收益递减，因为已经能隐藏延迟。

2. Coalesced Access（合并访存）

warp 内 32 个 thread 访问连续内存 → 1 次 HBM 读；散乱访问 → 32 次 HBM 读。

cuda

// 好（coalesced）：
a[threadIdx.x + blockIdx.x * blockDim.x]   // 32 个 thread 访问连续地址

// 坏（uncoalesced）：
a[threadIdx.y * N + threadIdx.x]            // 32 个 thread 访问非连续地址

coalesced 是 memory-bound 算子最重要的优化——非 coalesced 比 coalesced 慢 5-10 倍。

3. Bank Conflict（共享内存 bank 冲突）

shared memory 分 32 个 bank，每个 bank 串行服务。如果同一时钟周期 32 个 thread 访问同一 bank → 32 次串行。

cuda

// 坏（bank conflict）：
tile[threadIdx.x][0]   // 32 个 thread 都访问 bank 0

// 好（避免 conflict）：
tile[threadIdx.x][threadIdx.x % 32]  // 32 个 thread 访问不同 bank

4. Async Memcpy（异步传输）

H100 起的 cp.async 指令、TMA 单元让 HBM→SRAM 传输异步化：

cuda

// 老式：同步加载
tile[ty][tx] = a[row * K + col];   // 等 HBM 读
__syncthreads();

// 新式：异步加载 + 流水线
cp_async(&tile[ty][tx], &a[row * K + col]);
// 在等待期间算其他数据（如上一个 tile 的 matmul）

异步加载 + 双缓冲（double buffering）能把 HBM 延迟隐藏 100%。

5. Pipelining（流水线）

把数据加载、计算、写回三个阶段流水化：

text

朴素:   load1 → compute1 → write1 → load2 → compute2 → write2 → ...
流水:   load1 ─────────────→
                  compute1 ─────────────→
                                write1 ─────────────→
                              load2 ─────────────→
                                          compute2 ──→

流水化让 SM 在等待 HBM 时算别的数据——LLM 推理 kernel（如 Marlin）都用 2-4 级流水。

调优优先级

先修算法（找对的算法/并行策略）—— 收益 5-10×；
再修 coalesced/bank conflict（访问模式）—— 收益 3-5×；
再修 occupancy/async/pipeline（隐藏延迟）—— 收益 1.5-2×；
最后用 Tensor Core/FP8（硬件指令）—— 收益 2-4×。 颠倒顺序会浪费——在错算法上做调优是徒劳。

六、GPU 库生态

不要什么都自己写——能用成熟库就用：

库	用途	优势
cuBLAS / cuBLASLt	矩阵乘（GEMM）	NVIDIA 最优 GEMM 实现
cuDNN	CNN + attention + RNN	标准 DNN 算子库
CUTLASS	矩阵乘模板（C++）	可定制 epilogue 融合
CUB / Thrust	reduction/scan/sort	标准并行算法库
NCCL	多 GPU 通信	all-reduce/all-gather 最优实现
TensorRT / TRT-LLM	整图优化 + 部署	NVIDIA 闭源旗舰
cuQuantum	量子模拟	量子计算专用
nvJPEG / nvTIFF	图像编解码	多媒体专用

不要重造轮子

99% 的 LLM 推理不需要手写 kernel——vLLM/TensorRT-LLM 已经把 matmul/attention/layernorm 优化到极致。手写 kernel 是最后 10% 极致优化才需要——且成本高（人月级别）。

七、CUDA Profiling：用 Nsight 找瓶颈

性能优化离不开 profiling：

Nsight Systems：系统级 trace，找 CPU-GPU 同步、kernel launch 开销、内存传输瓶颈；
Nsight Compute：单 kernel 级 profiling，给 Roofline 上限、occupancy、bank conflict、coalesced 详细数据；
PyTorch Profiler：Python 层级，找哪些 op 慢、哪些 op 占内存。

典型 workflow：

text

Nsight Systems → 找最慢 kernel → Nsight Compute 单独分析 → 找瓶颈（coalesced？occupancy？）→ 改 kernel → 重测

详见benchmarking 实践。

八、分布式推理与多 GPU

单 GPU 装不下大模型（如 Llama-405B 需 800GB）时，必须用多 GPU：

并行策略	切分维度	通信量	适用
Tensor Parallel (TP)	把每层权重切成 N 份	all-reduce per layer	单机多 GPU，延迟敏感
Pipeline Parallel (PP)	把层切成 N 段	pipeline bubble	跨机，吞吐密集
Expert Parallel (EP)	按 expert 切分	all-to-all	MoE 模型
Data Parallel (DP)	数据并行	all-reduce grad	训练为主，推理少用

LLM 推理 TP 最常用——单机 8 卡把 70B 模型切成 8 份，每卡装 10GB。详见distributed-inference 案例研究。

九、权衡与取舍

GPU 型号选择：在线服务选 H200（带宽大）；训练选 B200（算力大）；边缘选 L40S（性价比）；
Tensor Core vs 普通 CUDA Core：算 matmul 用 Tensor Core（mandatory），算 elementwise 用 CUDA Core；
精度选择：训练 FP16/BF16；推理 INT8/FP8/INT4/FP4；H100+ 优先 FP8（免校准精度好）；
库 vs 手写：能用 cuBLAS/cuDNN 就别手写；需要融合时用 CUTLASS 或 Triton；极致才用 CUDA C++；
编译 vs 手调：torch.compile 自动优化 80%；最后 20% 手调（Nsight 找瓶颈）。

参考资料

NVIDIA H100 White Paper —— Hopper 架构与 Tensor Core 细节
NVIDIA Blackwell Architecture White Paper —— B200/B100 架构与 FP4
NVIDIA CUDA C++ Programming Guide —— CUDA 编程官方文档
NVIDIA CUTLASS Documentation —— GEMM 模板库
NVIDIA Nsight Compute Documentation —— kernel 级 profiling
Mark Harris. CUDA C++ Best Practices Guide —— CUDA 调优 checklist 权威指南
Kirk & Hwu. Programming Massively Parallel Processors（4th Ed., 2022） —— CUDA 教学权威教材
Patterson & Hennessy. Computer Organization and Design: RISC-V Edition —— 体系结构基础

GPU 体系结构与优化

GPU 体系结构与优化 ​

概念定义：GPU 不是"大 CPU" ​

一、GPU 体系结构 ​

从晶圆到 SM 到 Warp 到 Thread ​

SM 的关键资源 ​

二、Tensor Core：算力的来源 ​

计算能力 ​

算力计算 ​

三、各代 GPU 关键参数 ​

四、CUDA 编程模型回顾 ​

grid / block / thread 三级 ​

CUDA C++ kernel 模板 ​

Shared Memory：block 内 SRAM ​

Warp shuffle：32 thread 间超快通信 ​

五、性能调优 Checklist ​

1. Occupancy（占用率） ​

2. Coalesced Access（合并访存） ​

3. Bank Conflict（共享内存 bank 冲突） ​

4. Async Memcpy（异步传输） ​

5. Pipelining（流水线） ​

六、GPU 库生态 ​

七、CUDA Profiling：用 Nsight 找瓶颈 ​

八、分布式推理与多 GPU ​

九、权衡与取舍 ​

延伸阅读 ​

参考资料 ​

GPU 体系结构与优化

概念定义：GPU 不是"大 CPU"

一、GPU 体系结构

从晶圆到 SM 到 Warp 到 Thread

SM 的关键资源

二、Tensor Core：算力的来源

计算能力

算力计算

三、各代 GPU 关键参数

四、CUDA 编程模型回顾

grid / block / thread 三级

CUDA C++ kernel 模板

Shared Memory：block 内 SRAM

Warp shuffle：32 thread 间超快通信

五、性能调优 Checklist

1. Occupancy（占用率）

2. Coalesced Access（合并访存）

3. Bank Conflict（共享内存 bank 冲突）

4. Async Memcpy（异步传输）

5. Pipelining（流水线）

六、GPU 库生态

七、CUDA Profiling：用 Nsight 找瓶颈

八、分布式推理与多 GPU

九、权衡与取舍

延伸阅读

参考资料