如果你已经被“GPU程序跑得慢”、Nsight Profile 里一堆看不懂的指标，或者“明明照着教程写的，怎么还是快不起来”折磨过，我猜你缺的可能不是一个更好的 API 写法，而是一张关于 GPU 硬件本体的地图。网上讲 CUDA/ROCm 的教程一抓一大把，但大多数只讲接口怎么调，不讲芯片内部那套运行逻辑。恰恰是这套逻辑，决定了你的代码能跑多快、显存会不会爆、为什么同样一段代码在别人的卡上飞快、在你的卡上却慢得离谱。
这篇文章是 GPU 硬件原理架构系列的第一篇。我不打算上来就堆术语，而是想顺一条线走清楚：GPU 为什么被设计成如今这个样子、芯片里到底有哪些关键部件、几万个线程是怎么被组织起来“同时”执行的、数据在这些部件之间又是怎么流动的，以及硬件事实如何反过来决定了 CUDA 和 ROCm 的标准写法。看完这篇，你再去看那些“GPU 编程最佳实践”，会发现很多规则可以自己推导出来，而不是死记硬背。
1. CPU和GPU的底层分歧：延迟优先还是吞吐优先
1.1 同样是芯片，为什么设计思路差了这么远
CPU 的设计目标非常明确：让单条指令流执行得尽可能快。为了这个目标，一颗现代 x86 处理器的面积里塞了大量“不直接算数”的电路——分支预测器、乱序执行引擎、重排序缓冲区、超大规模的多级缓存，以及跨核心的一致性协议。这些电路存在的意义只有一个：猜测程序下一步要干什么，提前把数据准备好，让那几个运算单元别闲着。单核里真正的 ALU/FPU 占比通常也就是 20% 到 30%。
GPU 把目标换了：不再追求单条指令多快，而是追求单位时间里能执行多少条指令。所以 GPU 上那些“伺候单线程”的复杂控制逻辑全被砍掉，换成大量简单的计算单元。以一块 NVIDIA 的 SM 为例，里面有上百个 FP32 计算单元，但几乎没有分支预测，也没有乱序执行。整张卡积少成多，可以同时跑几万甚至几十万条线程。
行业里习惯用一句话概括：CPU 是延迟导向，GPU 是吞吐导向。CPU 擅长把单个任务尽快做完，GPU 擅长把海量同类任务成批做完。如果你手头的问题是单任务线性逻辑，GPU 不仅没有优势，反而会因为更高的访存延迟和更低的单核频率跑得更慢。很多入门者误以为“GPU 核多就一定快”，实际上只有任务具备足够的并行度，GPU 的优势才会体现出来。
1.2 晶体管预算、时钟和带宽，全都在配合同一个目标
任何芯片设计都绕不开“晶体管的钱花在哪”这个问题。一颗现代高端 CPU 的 L3 缓存动辄几十 MB，加上分支预测器和乱序执行那套机构，晶体管数量非常可观；而一张旗舰 GPU，比如 H100，800 亿晶体管的绝大多数投给了计算单元和与之配套的数据通路。GPU 的 L2 缓存相对它的算力来说小得多，这不是省成本，而是设计哲学使然：它默认你不会只靠一两个线程去跑，而是有海量线程在同时运行。
时钟频率上也能看出差异。消费级 CPU 睿频能干到 5GHz 以上，GPU 核心频率通常只有 1.5GHz 到 2.2GHz。更低的主频换来了更高的流处理器数量和更可控的功耗，因为 GPU 本来就是靠“宽度”而不是“高度”取胜。一块 GPU 上可能同时驻留几万个线程，靠整体规模碾压单核的时钟优势。
内存系统更是天差地别。CPU 的内存带宽依赖内存通道数量，顶尖服务器 CPU 也就几百 GB/s；GPU 走的是超宽显存总线，HBM 这类堆叠显存能提供 3TB/s 以上的带宽。这一条对 GPU 太重要了——神经网络训练、渲染、科学计算全都是访存密集型任务，GPU 能成为 AI 计算主力，超高带宽功不可没。所以判断一张 GPU 值不值，不能只看核心数和频率，更要看它“喂得饱自己”的能力：显存带宽、L2 带宽、寄存器文件总量，以及片上互联。这也是后面几节要展开的内容。
2. 宏观拆解GPU：从GPC到SM，芯片是分块管理的
2.1 先看全局：GPC/TPC/SM这套层级是怎么来的
现代 NVIDIA GPU 的组织方式是分区块的：几个 SM 组成一个 TPC，几个 TPC 组成一个 GPC，整颗芯片由多个 GPC 构成。以 Ampere 架构 GA102 为例，7 个 GPC 分布在芯片四周，中间是跨芯片共享的 L2 缓存和访存接口。AMD 这边结构类似，Shader Engine 下挂若干 Compute Unit（CU）。叫法不同，但分块管理的思想完全一致。
分层不是软件工程师拍脑袋定的，主要是三个原因。第一，物理设计上芯片面积太大，必须分成多个区域独立管理时钟和功耗，否则局部发热和时钟偏差会失控。第二，内存和缓存系统需要就近布局，数据通路越短，能效越好。第三，调度需要合适的粒度，把任务分配到 SM 之后，SM 内部的资源管理才能做得精细。
这个层级直接对应到 CUDA 的编程模型：一个 kernel 启动后以 grid 形式存在，grid 里的 thread block 会被调度到各个 SM 上，一个 block 只能待在一个 SM 上，不能拆分。所以 block 的大小和资源占用，直接影响一个 SM 能同时容纳多少个 block，也就影响整体占用率。调优时反复折腾 block 尺寸，本质上就是在调整这张“任务分配表”。
2.2 SM内部解剖：一个车间里的关键工种
把镜头拉近到 SM，一个现代 SM 就像一个小车间，里面有几种互相配合的“工种”。以 Ampere 架构为例，一个 SM 包含 128 个 FP32 计算单元（也就是大家常说的 CUDA core）、若干 SFU（Special Function Unit，负责 sin/cos 等超越函数）、若干 LSU（Load/Store Unit，负责访存指令），还有 4 个 Tensor Core 和 4 个 Warp Scheduler。
运算单元负责算术，Tensor Core 专门做矩阵乘加，从 Volta 开始引入，是如今深度学习训练和推理的算力主力。Warp Scheduler 则负责不停地从就绪 warp 里选一条指令派发到执行单元，SM 里有多个调度器，可以并行派发。寄存器文件是 SM 里非常庞大的一块资源，Ampere 一个 SM 有 65536 个 32 位寄存器，合计 256KB，比普通 CPU 的一个核心多得多。Shared Memory 和 L1 缓存共用一片片上 SRAM，靠配置决定多少给缓存、多少给显式共享。
这里的核心关系是：SM 的并发能力由寄存器文件、Shared Memory、线程数上限共同决定，三种资源任何一个先被占满，都会限制驻留线程总量。这直接导致一个常见现象：你以为限制性能的是“线程数”，很多时候其实是寄存器或 Shared Memory 先满了。
2.3 用手算推一遍：为什么线程数和寄存器会打架
这里给一个实际的计算例子。假设某个 SM 寄存器总量 65536 个，你的 kernel 每个线程用 64 个寄存器，那么一个 256 线程的 block 需要 256×64=16384 个寄存器，这个 SM 最多能驻留 65536/16384=4 个 block，也就是 1024 个线程，对应 SM 最大并发 2048 线程的 50% 占用率。
如果每个线程用到 128 个寄存器，一个 block 就要 32768 个寄存器，SM 只能驻留 2 个 block，也就是 512 线程，占用率降到 25%。如果编译器发现寄存器不够，会把部分局部变量溢出（spill）到本地内存——实际上就是显存——性能立刻崩。
所以调整 block 大小、launch_bounds这些参数，本质上就是在“线程并发数”和“每个线程可用寄存器数”之间找平衡。我的经验是：不要迷信最大占用率，先用 profiler 看 kernel 是访存瓶颈还是计算瓶颈，再决定往哪个方向调。一个计算密集的 kernel，占用率略低但每个线程手里寄存器充裕，往往比强行拉高占用率导致溢出更快。
3. Warp执行模型：GPU性能最核心的“隐藏开关”
3.1 32个线程绑成一个warp，是指令派发的最小单元
在 NVIDIA 硬件里，32 个线程组成一个 warp，SM 以 warp 为单位派发指令。一个 warp 里的 32 个线程执行同一条指令，指令的操作数来自各自不同的寄存器，这就是 SIMT（Single Instruction, Multiple Threads）。AMD 里的叫法不同，GCN/CDNA 时代叫 wavefront，64 线程一组；RDNA 后改成 32 线程的 wave32。名字不同，原理一致。
为什么是 32？一个常见的解释是：warp 太大，分支分歧的代价会变高；warp 太小，取指与调度的开销占比不划算。同时 32 个线程读 32 个 4 字节的 float，正好是 128 字节，和缓存行粒度吻合，非常适合做合并访存。这个数字从 Fermi 时代延续至今，已经成为 CUDA 生态的基础假设，大量优化技巧都是围着它转的。
一个 256 线程的 block 会被切成 8 个 warp。warp 内的线程用 lane id 标识，范围是 0 到 31。写 kernel 时经常看到 threadIdx 和 lane id 的换算关系，但要注意 block 维度与 warp 切分方式不是简单的“行优先一眼看穿”，真要吃透，最好的办法是用 Nsight 或者写个小 kernel 打印 threadIdx 和 warp 的关系。
3.2 延迟隐藏：GPU不靠大缓存靠“人海战术”
CPU 降低访存延迟靠缓存和乱序执行；GPU 没有那么多缓存，也不做乱序，靠的是并发掩盖延迟。当某个 warp 执行的指令需要从显存取数，它要等几百个周期才能拿到数据，此时这个 warp 暂时派不上用场。不过 SM 里还有其他 warp，调度器立刻切入一个可以执行的 warp，让计算单元一直有事干。
这个过程叫延迟隐藏（latency hiding）。要隐藏得彻底，SM 里活动 warp 的数量得足够多。粗略估算一下：如果显存访问延迟大约 600 个周期，SM 每周期派发一条指令，那么至少需要几百条“在飞行”的指令，也就是要几十个 warp 同时在驻留，才能把访存等待完全掩盖掉。这就是 occupancy（占用率）的意义：活动 warp 数除以 SM 最大 warp 数。
占用率太低，延迟盖不住，算力闲置；占用率太高，寄存器或 Shared Memory 不够，反而触发溢出。实际调优时，我通常先跑 Nsight，看 Warp Stall 和 Occupancy 两个指标，对比不同 block 大小下的表现，而不是一上来就抄网上推荐的“256 线程”或“512 线程”。不同 kernel 的最优配置差别很大，必须实测。
3.3 分支分歧：同一个warp，同一时间只能走一条路
因为 warp 是同步执行一条指令，所以如果 warp 内线程出现分支分叉，硬件就必须让它们依次执行每个分支。比如 if (tid % 2 == 0) 走 A，else 走 B，执行顺序就是先跑 A（一半线程干活，另一半被掩蔽，也就是空转），再跑 B（另一半干活），总耗时可能直接翻倍。
这不是死锁，也不是错误，但它是一种不可见的浪费。对于性能敏感的 kernel，尽量让 warp 内所有线程走同一条路径。比如处理循环边界条件时，可以把边界处理的代码单独用一个 warp 或一个 block 处理，而不是让每个 warp 都带一个 if。还有一个常用技巧：很多条件判断可以用位运算、乘掩码等方式改写成算术操作，从根源上避开分支分歧。理解这一点之后，再看网上那些“用 min/max 代替 if”的写法，就知道背后的原因了。
4. 存储层级：GPU的“血液循环系统”
4.1 五级存储，每一级负责什么
GPU 的存储大致分五级，访问速度从快到慢，容量从大到小。理解这个结构，是分析一切性能问题的起点。
| 层级 | 位置 | 典型容量 | 延迟量级 | 说明 |
|---|---|---|---|---|
| 寄存器文件 | SM 片上 | 256KB/SM | 1~2 周期 | 线程私有 |
| Shared Memory / L1 | SM 片上 | 数十KB~百KB | 20~30 周期 | block 内共享，可编程 |
| L2 缓存 | 全芯片 | 20~50MB | 200 周期左右 | 跨 SM 共享 |
| 显存 HBM/GDDR | 片外 | 数十GB | 500~1000 周期 | 主存 |
| 主机内存 / PCIe | 系统 | 系统内存 | 微秒级 | 数据交换 |
寄存器文件在每个 SM 里都很夸张，相当于给每个线程预分配了多个私人储物柜，只要不用太多，热点数据全都可以放里面。Shared Memory 和 L1 在物理上是同一块 SRAM，靠配置决定多少给自动缓存、多少给显式共享。L2 是整颗芯片统一调度的最后一级片上缓存，跨 SM 的数据共享和原子操作都要经过它。最后落到显存，容量最大，但速度掉了一个量级。
4.2 内存合并：GPU访存的第一法则
GPU 访存的效率高度依赖 warp 内 32 个线程访问的地址是否连续。如果线程 tid 访问的地址是 base + tid * 4（每个线程读一个 4 字节的 float），那么 32 个线程访问的是 128 字节连续区域，硬件可以合并成一两次显存事务完成；如果访问 base + tid * 128，每个线程间隔很大，就会变成几十次独立事务，有效带宽暴跌。
这就是常说的内存合并（memory coalescing），它在任何 GPU 架构上都成立，因为片外访存的代价实在太贵。优化的常见思路是调整数据结构布局：把“按对象组织”改成“按字段组织”，也就是 SoA（Structure of Arrays）而不是 AoS（Array of Structures）。举个例子，一个粒子系统需要位置、速度、质量三个属性，AoS 是每个粒子一个结构体，SoA 是三个独立数组。GPU 线程要连续访问所有粒子的速度时，SoA 天然合并，AoS 则会出现跨步访问。这项改动在很多场景里能带来好几倍的访存加速，而且改造成本通常不高。
4.3 算力再高也怕带宽饿死：Roofline视角
GPU 算力增长一直快于显存带宽增长。拿 H100 举例，FP32 算力大约 60 多 TFLOPS，HBM3 显存带宽约 3.3TB/s。如果一条指令需要从显存读两个 FP32 数做一次加法，理论刚需带宽等于算力乘以 8 字节，远超显存能提供的量。所以绝大多数真实 kernel 都是访存受限的，也就是常说的 memory-bound。
判断一个 kernel 是 compute-bound 还是 memory-bound，最通用的工具是 Roofline 模型：纵轴是算力，横轴是算术强度（FLOP/Byte），斜线部分是带宽限制区，水平部分是算力限制区。先估算自己 kernel 的算术强度，再和 GPU 的“FLOPS/带宽”比值对比，就能知道优化方向。这个比值在 H100 上大约是 20 FLOP/Byte，在消费级卡上通常更低。实际优化时，如果程序跑在斜线区，做再多的指令优化都没用，得想办法减少访存量：数据复用、缓存、合并访问、降低精度，才是正路。
5. 架构怎么倒逼出编程模型：反推几条实战准则
5.1 CUDA线程模型就是硬件模型的倒影
CUDA 的 grid/block/thread 三层模型，不是抽象出来的摆设，它就是硬件执行层的直接映射：一个 thread 对应一个 SIMT lane；32 个 thread 组成一个 warp，对应硬件派发单位；一个 block 被调度到一个 SM 上，内部再切成 warp；整个 grid 横跨整张 GPU。
理解这个对应关系，很多“规定”就变得合理了。比如为什么 block 内可以用 shared memory 和 __syncthreads() 同步，而 block 之间不行？因为 shared memory 就在 SM 里，block 里的线程都在同一个“车间”，同步是片上操作；跨 block 只能经过 L2 和全局内存，代价高，所以 CUDA 连原生 grid 级同步都很少见。再比如，为什么性能建议里常说“相邻线程访问相邻地址”？因为相邻线程就是同一个 warp 的相邻 lane，它们的访存在硬件里天然可以被合并。这不是风格偏好，是硬件事实。
5.2 三条能直接上手的架构思维
把前面的内容浓缩成三条实战判断，遇到性能问题可以按这个顺序排查。
第一，先判断方向。一个 kernel 慢，先看是卡在访存还是卡在计算。Nsight Compute 里看 SOL（Speed of Light）相关的两个占比，哪个接近 100% 就说明瓶颈在哪，再决定优化动作。方向判断错了，后面全是无用功。
第二，算资源账。每个线程用多少寄存器、多少 shared memory、block 开多大，先用资源总量算一遍能驻留多少个 block、多少个 warp。很多时候性能差的根源就是寄存器 spill，或者 shared memory 不够导致占用率太低。算完这笔账，很多“玄学”调优立刻变成确定性工程。
第三，改数据布局优先于改指令。因为 GPU 普遍 memory-bound，优化访存模式通常比优化算术指令收益大得多。把 AoS 改成 SoA，加restrict帮助编译器做向量化，用 float4 这类向量类型一次取 16 字节，都是立竿见影的手段。这些写法的共同目的，就是把“每次访存能搬回来的有用字节数”尽量拉高。
我最初学 GPU 编程时也走了一段弯路：拿着 CPU 的思维去优化，在分支和循环上死磕，结果收益甚微。后来把架构文档反复读了几遍，对照 Nsight 面板一点点核对，才意识到 GPU 调优的钥匙几乎全在硬件上。这个系列我计划按这个顺序写下去：本篇是整体架构，下一篇拆 warp 调度与指令流水线，再往后是 Tensor Core 的矩阵乘原理、显存子系统与多卡互联。如果你也在折腾 GPU 相关的东西，先把这一层的底子打牢，后面讨论算子开发、深度学习训练的性能调优，都会轻松很多。
