GPU架构基础：GPU与CPU的区别、CUDA核心、流式多处理器（SM）、显存层次
核心原理：CPU与GPU的“分工哲学”
要理解GPU架构，首先要回答一个根本问题：为什么我们不能直接用CPU完成大规模并行计算？答案藏在两种芯片的设计目标差异中。CPU是为低延迟优化的“单核强者”，它把大量晶体管用于分支预测、乱序执行、大容量缓存，目的是让单条指令流尽可能快地跑完。而GPU是为高吞吐量优化的“众核劳模”，它牺牲单线程速度，换取数千个线程同时执行的能力。
英伟达将这种差异概括为“Latency-Oriented vs Throughput-Oriented”。CPU的典型规格是8~16个高性能核心，每个核心拥有独立的L1/L2缓存和复杂的控制单元；而GPU（如A100）拥有6912个CUDA核心，但控制单元极简，缓存也以共享内存和纹理缓存为主。下图展示了二者的核心结构差异：
从图中可以直观看到，CPU的每个核心都配备复杂的控制单元和私有缓存，适合执行分支密集、依赖强的串行任务；而GPU的SM（流式多处理器）内部包含大量轻量级CUDA核心，它们共享控制逻辑和共享内存，适合执行数据并行任务——比如矩阵乘法、图像滤波、神经网络卷积等。
CUDA核心与流式多处理器（SM）
CUDA核心是英伟达GPU中最基本的计算单元，本质上是一个单精度浮点/整数运算单元（FP32/INT32）。它没有CPU核心那样的复杂分支预测和乱序执行能力，但胜在数量庞大。以Ampere架构的A100为例，每个SM包含64个FP32 CUDA核心、32个INT32核心和32个FP64核心。当执行一个kernel时，线程被组织为线程束（warp），每个warp包含32个线程，SM以warp为单位调度执行——这意味着一个warp中的32个线程在同一时刻执行同一条指令（SIMT模式）。
SM是GPU的“心脏”，它负责线程的创建、调度和执行。一个典型的SM包含以下关键组件：
- CUDA核心阵列：分为4个处理块（Processing Block），每个块包含16个FP32核心和8个INT32核心。
- 共享内存（Shared Memory）：可编程的片上高速缓存，容量通常为64KB~228KB（取决于架构），用于同一线程块内的数据交换。
- 寄存器文件（Register File）：每个SM拥有64K个32位寄存器，分配给线程使用，寄存器溢出会导致性能骤降。
- 线程束调度器（Warp Scheduler）：每个SM有4个调度器，每个调度器每时钟周期可发射一条指令。
- 特殊功能单元（SFU）：用于执行超越函数（如sin、cos、log等）。
理解SM的关键在于线程层次结构。CUDA的线程组织为三层：grid（网格）→ block（线程块）→ thread（线程）。一个grid包含多个block，一个block包含多个线程（最多1024个）。block被调度到SM上执行，block内的线程通过共享内存通信，而不同block之间无法直接通信，只能通过全局内存。这种设计使得GPU可以透明地扩展：无论硬件上有多少个SM，同一个kernel都能运行，只是block被分配到不同SM上并行执行。
显存层次：从寄存器到全局内存
GPU的存储体系是性能优化的核心战场。下表总结了从最快到最慢的各级存储及其关键特征：
| 存储层级 | 容量（典型值） | 延迟 | 作用域 | 用途 |
|---|---|---|---|---|
| 寄存器文件 | 每个SM 256KB | ~1周期 | 单个线程 | 局部变量、中间计算 |
| 共享内存 | 每个SM 48~228KB | ~20周期 | 线程块内 | 数据复用、块内通信 |
| L1缓存 | 每个SM 128KB~256KB | ~30周期 | 线程块内 | 缓存全局内存访问 |
| L2缓存 | 40MB（A100） | ~200周期 | 整个GPU | 跨SM数据共享 |
| 全局内存（显存） | 40GB~80GB（HBM2e） | ~400-800周期 | 所有线程 | 主数据存储 |
| 主机内存（PCIe） | 系统内存 | 数微秒 | CPU+GPU | 数据交换 |
一个常见的性能误区是：把所有数据都放在全局内存中，然后抱怨GPU“慢”。实际上，全局内存的带宽虽然很高（A100超过2TB/s），但延迟也极高。正确的做法是：
- 将数据从全局内存加载到共享内存，供线程块内反复使用。
- 利用寄存器缓存每个线程私有的中间结果。
- 通过合并访问（coalesced access）让相邻线程访问相邻内存地址，最大化内存带宽利用率。
下面是一个典型的矩阵乘法优化示例，展示了如何利用共享内存和寄存器减少全局内存访问：
// 矩阵乘法 C = A * B，每个线程计算一个输出元素
__global__ void matmul_kernel(float *A, float *B, float *C, int N) {
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
float sum = 0.0f;
for (int tile = 0; tile < N / BLOCK_SIZE; ++tile) {
// 协作加载数据到共享内存
As[threadIdx.y][threadIdx.x] = A[row * N + tile * BLOCK_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(tile * BLOCK_SIZE + threadIdx.y) * N + col];
__syncthreads(); // 确保块内所有线程加载完成
for (int k = 0; k < BLOCK_SIZE; ++k) {
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads(); // 防止覆盖共享内存
}
C[row * N + col] = sum;
}这段代码中，每个线程块负责计算一个BLOCK_SIZE × BLOCK_SIZE的输出子矩阵。通过将A和B的分块加载到共享内存，每个数据元素被重复使用了BLOCK_SIZE次，从而将全局内存访问量减少了BLOCK_SIZE倍。这是GPU优化的经典“分块（tiling）”技巧。
案例与场景：从理论到实践
让我们通过一个实际案例来串联上述概念。假设我们要在GPU上执行一个图像模糊操作（3×3卷积核）。输入图像为1920×1080像素，每个像素为RGB三通道浮点数。
第一步：线程映射。我们让每个线程处理一个输出像素。图像被划分为16×16像素的线程块，每个块包含256个线程。这样，整个图像需要(1920/16) × (1080/16) = 120 × 68 = 8160个线程块。
第二步：数据复用。每个输出像素需要读取周围3×3=9个输入像素。如果直接从全局内存读取，每个像素会被相邻的9个线程重复读取，产生大量冗余访问。优化方案是：每个线程块先将对应的18×18像素区域（包含边界）加载到共享内存，然后所有线程从共享内存读取数据。
第三步：边界处理。图像边缘的像素缺少邻居，通常采用“钳位”（clamp）策略，即超出边界的坐标返回边界像素值。这可以在加载共享内存时判断并处理。
通过这种优化，全局内存访问量从1920×1080×9次降低到1920×1080次（每个像素只被加载一次），性能提升可达5~8倍。这个案例充分体现了共享内存作为“用户管理的缓存”的价值——它比硬件自动管理的L1缓存更可控，也更高效。
误区与复盘：常见性能陷阱
误区一：线程数越多越好。实际上，当线程数超过GPU可同时驻留的上限时，多余的线程只是排队等待，并不会增加并行度。每个SM能驻留的最大线程数受限于寄存器数量和共享内存容量。例如，如果每个线程使用64个寄存器，一个SM的64K寄存器只能支持1024个线程，而SM的硬件上限是2048个线程。此时应减少每线程寄存器使用量，或调整block大小。
误区二：忽略warp发散。当warp内的32个线程执行不同的分支路径时（例如if (threadIdx.x % 2 == 0)），GPU会串行执行每个分支，导致性能下降。应尽量让warp内的线程保持统一执行路径，或将分支条件改为算术运算（如a = cond ? x : y）。
复盘建议：使用NVIDIA Nsight Compute工具分析kernel时，重点关注以下指标：Achieved Occupancy（实际占用率）、Shared Memory Bank Conflicts（共享内存冲突）、Global Load/Store Efficiency（全局内存访问效率）。如果占用率低于50%，优先检查寄存器或共享内存是否成为瓶颈；如果全局内存效率低于80%，检查是否存在非合并访问。
最后需要强调的是，GPU架构的演进始终围绕“提高吞吐量”这一核心目标。从Volta的Tensor Core到Hopper的Transformer Engine，英伟达不断在SM中增加专用计算单元，但CUDA核心和SM的基本编程模型保持稳定。理解本节的基础概念，是后续学习Tensor Core、CUDA编程优化、多GPU通信等高级主题的基石。