1. GPU架构基础：从CPU到GPU的思维转变

做GPU优化，首先得理解GPU到底是个什么东西。我见过不少从CPU转过来的同学，上来就按CPU的思路写CUDA代码，结果性能惨不忍睹。说白了，GPU和CPU的设计哲学完全不同——CPU是「全能选手」，GPU是「流水线工人」。

1.1 GPU vs CPU：设计目标的根本差异

CPU的设计目标是「低延迟」——它要快速响应各种不同类型的任务。所以CPU有庞大的缓存、复杂的控制单元、强大的分支预测。而GPU的设计目标是「高吞吐」——它不在乎单个任务慢一点，但同一时间要处理成千上万个任务。

对比维度 CPU GPU
核心数量 4-16个高性能核心 数千个轻量级核心
控制单元 复杂（分支预测、乱序执行） 简单（顺序执行为主）
缓存大小 大（MB级别） 小（KB级别）
内存带宽 50-100 GB/s 900-2000 GB/s
适用场景 串行、复杂逻辑 并行、数据密集型

核心观点：CPU擅长「做一件事很快」，GPU擅长「同时做很多事」。如果你的算法是串行的，老老实实用CPU；如果是数据并行的，GPU才是你的菜。

1.2 SM与CUDA Core：GPU的「车间」和「工人」

GPU内部的基本计算单元叫SM（Streaming Multiprocessor，流式多处理器）。你可以把SM想象成一个「车间」，每个车间里有多个CUDA Core——也就是「工人」。

以NVIDIA的Ampere架构为例，每个SM包含128个CUDA Core。一个GPU可能有几十到上百个SM。所以总CUDA Core数量轻松上万。

我的经验：刚开始做CUDA时，我总以为CUDA Core越多越好。后来发现，真正限制性能的往往是「怎么把数据喂给这些Core」。你想想看，一万个工人等着干活，但原材料（数据）送不过来，那工人再多也没用。

SM内部还有几个关键部件：

  • Warp Scheduler：负责调度线程。32个线程组成一个warp，这是GPU执行的基本单位。
  • Register File：寄存器文件，每个SM有数万个寄存器。
  • Shared Memory：共享内存，SM内部的高速存储。
  • Tensor Core：专门做矩阵乘法的加速单元（从Volta架构开始引入）。

1.3 显存层次结构：速度与容量的博弈

GPU的存储层次和CPU类似，但带宽和延迟的差异更大。我习惯把显存层次比作「仓库-工作台-手边工具」的关系。

GPU显存层次结构 Global Memory（全局显存） 容量：8-80 GB | 带宽：900-2000 GB/s | 延迟：400-800 cycles L2 Cache 容量：6-40 MB | 延迟：200-400 cycles L1 / Shared Memory 容量：48-164 KB/SM | 延迟：~30 cycles Register ~1 cycle 越往上越快、越小 越往下越慢、越大 💡 优化核心： 让数据尽量留在上层

1.3.1 Global Memory（全局显存）

这是GPU最大的存储空间，所有线程都能访问。但它的延迟高达400-800个时钟周期。我刚开始做优化时，经常因为频繁访问Global Memory导致性能瓶颈。后来学乖了——能少访问就少访问。

避坑指南：我曾经写过一个矩阵乘法的kernel，每次计算都从Global Memory读数据，结果性能还不如CPU。后来改成用Shared Memory做分块，速度直接提升了20倍。记住：Global Memory是「万恶之源」，能不用就不用。

1.3.2 Shared Memory（共享内存）

Shared Memory位于SM内部，延迟只有30个时钟周期左右。同一个Block内的线程可以共享这块内存。它就像车间里的「公共工作台」——大家把原材料搬上来，一起加工。

使用Shared Memory的典型模式：

// 从Global Memory加载数据到Shared Memory
__shared__ float s_data[BLOCK_SIZE];
s_data[threadIdx.x] = g_data[blockIdx.x * blockDim.x + threadIdx.x];
__syncthreads();  // 等待所有线程加载完成

// 在Shared Memory上计算
float result = s_data[threadIdx.x] * 2.0f;

1.3.3 Register（寄存器）

寄存器是GPU里最快的存储，访问只需要1个时钟周期。每个SM有数万个寄存器，但每个线程能用的数量有限（通常是255个）。

我个人的习惯是：尽量把频繁使用的变量声明为局部变量，编译器会自动分配到寄存器。但要注意——寄存器用太多会导致「寄存器溢出」，数据被挤到Local Memory（实际在Global Memory里），性能反而下降。

小技巧：用 --ptxas-options=-v 编译选项可以查看每个kernel用了多少寄存器。如果超过32个，就要考虑优化了。

1.4 内存带宽与延迟：性能的「水龙头」和「水管」

带宽决定了单位时间内能传输多少数据，延迟决定了每次访问需要等多久。GPU的带宽很高（HBM2e可达2 TB/s），但延迟也很高（几百个时钟周期）。

为什么会这样？因为GPU的设计思路是「用并行掩盖延迟」——当一个warp在等数据时，立刻切换到另一个warp执行。只要你有足够多的warp，延迟就被「隐藏」了。

存储层级 典型带宽 典型延迟 优化策略
Register ~20 TB/s 1 cycle 多用局部变量
Shared Memory ~10 TB/s ~30 cycles 数据复用、分块
L2 Cache ~2 TB/s 200-400 cycles 合并访问、缓存友好
Global Memory ~1.5 TB/s 400-800 cycles 合并访问、减少访问次数

关键认知：GPU优化的本质就是「管理数据流」——让数据尽量留在寄存器或Shared Memory里，减少对Global Memory的访问。你想想看，如果每次计算都要去Global Memory读数据，那再多的CUDA Core也只能干等着。

嗯，这一章的内容就到这里。记住这些基础概念，后面我们会一步步深入——从内存合并到Bank Conflict，从Occupancy到异步拷贝。打好基础，后面的优化技巧才能用得上。

← 上一章 📖 返回目录 下一章 →