企业官网建设流程全解析

如果你已经照着前两章的内容把第一个 CUDA 程序跑通了，接下来大概率会遭遇一个很常见的尴尬：同样的 kernel，在朋友那台新显卡上飞快，换到自己机器上却像老牛拉车。我第一次碰到这个问题时，以为是代码里埋了什么逻辑坑，反复查了一下午，最后才发现——两张显卡的硬件差异，比代码差异大得多。做 GPU 并行编程，如果对英伟达 GPU 架构只有模糊印象，性能优化的天花板基本就被锁死在想象力上了。

这篇第三章，我打算把英伟达 GPU 的架构与代系关系完整地理一遍：SM 到底怎么工作、Warp 为什么永远贴着 32 条线程、Compute Capability 和sm_XX这些开发时天天见到的名词背后对应着哪一代硬件，以及从 Fermi 一路走到 Blackwell，代系更迭到底在改什么。它适合刚掌握 CUDA 基础语法、正准备深入性能调优的读者，也适合那些想搞清楚自己到底该买什么显卡的人。看懂这一章，后续再聊优化就有共同的坐标系了。

1. 先放下“核数”思维：GPU 的真实工作方式

1.1 GPU 不是“核特别多的 CPU”

很多人第一次看 GPU 参数，第一反应是“哇，几千个核心”，然后顺理成章地以为并行编程就是把 CPU 的多线程代码复制过来，线程多开几倍就行。这个误解会带偏后面所有优化方向。

CPU 的设计哲学是少数几个大核、高频率、复杂的乱序执行和分支预测、大容量缓存。它擅长处理依赖强、逻辑复杂的串行任务，那种“一个步骤卡住，后面全部等它”的问题它能硬算扛过去。GPU 则完全是另一套思路：塞进去上千个轻量计算单元，每个单元本身很朴素，但数量大、集体行动。单个 CUDA Core 并不能独立执行复杂的控制流，它必须依赖一个共同的“指挥单元”，以小组为单位整齐推进。

打个比方：CPU 像一位全能工匠，从测量、切割到打磨全流程都能独立完成；GPU 像一条几千人的流水线，每个人只重复做一个简单动作，但组合起来产量惊人。理解了这个区别，你才会明白为什么 GPU 编程里“线程数量”的设计远比“线程优先级”重要。

1.2 隐藏延迟：用线程数量换时间

为什么 GPU 需要维持那么多线程？核心原因在于“访存延迟隐藏”。

假设一次全局内存访问需要几百个周期。CPU 的做法是砸大量缓存、乱序执行、分支预测，尽量让算力等内存的时间缩短；GPU 的做法更直接——如果一条线程在等内存返回，调度器立刻把其他线程切进来执行。因为线程切换只是换一套寄存器上下文，开销近乎为零，所以只要线程数量足够多，计算单元几乎总能保持忙碌。

这带来一个反直觉的结论：在 GPU 上，线程太少比线程太多更危险。不少人写 kernel 时把 block 调小、线程调少，觉得“反正数据就这么多”，结果算力利用率极低。实际测下来，往往是线程数量上去了，延迟被盖住了，吞吐量才会质变。这是硬件设计决定的，不是玄学。

1.3 CUDA 的三层线程结构，是对硬件的一种“告白”

CUDA 没有把底层全部藏起来，它反而把硬件调度单位显式暴露给程序员。Grid、Block、Thread这三档就是软件对硬件的映射：

• Thread是最小执行单元，一条线程跑一个独立的数据路径。
• 多条Thread组成Block，一个 Block 会被完整分配到同一个 SM 上，Block 内的线程可以通过共享内存协作。
• 多个Block组成Grid，Grid 里的 Block 由 GPU 的调度器分发给所有 SM。

这样设计的原因很实际：同一个 Block 里的线程可以通过共享内存同步通信，而不同 Block 之间不能被假定为同时运行，通信只能走全局内存，还得靠多次 kernel 调用或者 atomics 这类笨办法。

所以“Block 会被完整放在一个 SM 上”这句话是理解很多限制的钥匙：为什么共享内存容量有限制？因为一个 SM 的共享内存就那么大；为什么 Block 不能开得无限大？因为一方面 SM 最多只能同时容纳那么多线程，另一方面硬件本身对 Block 尺寸有上限。我见过有人把 Block 设成(256, 256)然后编译失败，就是没想清楚这一层。

2. 英伟达架构的“族谱”：从 Fermi 到 Blackwell

2.1 主流动脉一览表

英伟达 GPU 的架构更迭，从开发者角度可以看作一条清晰的线索。我把主流代系的代表参数整理成了一个表，信息按“开发者最关心的口径”压缩过。

注意看 Compute Capability 这一列，你会发现它并不跟架构代号“等价跳号”。Maxwell 直接跳过了 4.x，Ada Lovelace 又落在了 8.9。这是因为计算能力版本表的是“硬件指令代次”，只要指令集没有断代式变化，英伟达就不愿意让代号断得太碎。

2.2 架构代号不是“换马甲”，SM 内部重构才是关键

很多人以为新一代显卡只是频率高一点、流处理器多一点、显存大一点。真正拉开性能差距的，往往是 SM 内部结构的重构。

Fermi 时代一个 SM 里有 32 个 CUDA Core；到了 Kepler，SMX 一个 SM 塞下 192 个 CUDA Core，调度器的分工也变了；Maxwell 又主动“精简”回 128 个配合低功耗设计；到了 Ampere 架构消费级 GPU，英伟达把 FP32 和 INT32 做成两条独立管道，很多型号的普通浮点单元数量直接翻了一倍。同样叫 SM，不同代系里面能并行推进的数学指令数量完全不同。

Tensor Core 的出现更是一个分水岭。Volta 之前，矩阵乘要靠一堆 FFMA 指令一步步算；Volta 开始，硬件里多了一个专门做矩阵乘加的张量单元，一条指令算一个子矩阵。本质上，它把“软件循环”变成了“硬件指令”，这对 AI 训练的加速是核弹级别的。

所以看代系关系时，别只盯着工艺制程数字。7nm 换 5nm 固然重要，但 SM 内部有没有结构性变化，才是决定代码上限的东西。

2.3 同一个代系，也可能有两种“脾气”

这里要特别提一个很容易踩的坑：同一个架构代号里，不同型号的 Compute Capability 未必相同。比如 Ampere 就有 8.0 和 8.6 两拨；数据中心旗舰 A100 是 8.0，普通消费级 RTX 30 系是 8.6。它们都叫 Ampere，但指令集细节有差异。

这意味着判断硬件能力时，不能只看销售文案里写的“Ampere”就默认一样，必须在代码里读取实际的 major/minor 值。很多人拿着网上抄来的-arch=sm_80编译参数，放到自己的 RTX 3060 上跑，直接报 “no kernel image”，就是因为没搞清楚自己的卡其实是sm_86。记住这句话：架构代号是市场语言，Compute Capability 才是开发者语言。

3. 开发中真正打交道的是 Compute Capability 与 sm_XX

3.1 major.minor 到底代表什么

Compute Capability 是一串major.minor编号，例如7.5、8.6、9.0。它描述的是这个 GPU 支持的硬件指令代次，直接影响：

• 支持哪些 PTX/SASS 指令（有没有 Tensor Core、有没有特定原子操作）。
• Thread Block 的最大尺寸、单个 SM 的最大线程数。
• 共享内存容量上限、寄存器配额。
• 运行时库会据此决定启用哪些优化路径。

驱动程序和 CUDA runtime 按这套编号对外提供功能。也就是说，你在写代码时查询prop.major和prop.minor，本质上就是在问“眼前这个硬件到底属于哪一代、懂哪些指令”。

3.2 nvcc 的编译目标：arch、code、gencode 三件套

我把这三件套理解递进着看：

• -arch=compute_XX指定生成 PTX 的目标虚拟架构。PTX 是 CUDA 的中间表示，类似 Java 的 bytecode，它可以跨代兼容，只要驱动够新，PTX 能被 JIT 实时编译成新硬件的机器码。
• -code=sm_XX指定生成 SASS 的目标实际架构。SASS 是真正绑定硬件的机器码，不能跨代运行。
• 如果不写-code，-arch的参数会同时用于生成 SASS，所以这时候你的可执行文件就只认那一个确切代系。

实际项目里，我建议用-gencode一次性打包多个目标：

nvcc -gencode arch=compute_80,code=sm_80 \ -gencode arch=compute_86,code=sm_86 \ -gencode arch=compute_90,code=sm_90 \ -o app main.cu

这样编出的 fatbin 里包含多套 SASS，驱动会在运行时挑最匹配的那套加载。代价是编译时间和二进制体积变大，但换来的是不同机器上都能跑，省事很多。

还有一种写法是同时带上 PTX 作为兜底：

nvcc -gencode arch=compute_90,code=compute_90 ...

code=compute_90表示把 PTX 也塞进 fatbin，这样未来新一代硬件出现时，如果驱动够新，可以通过 JIT 把 PTX 编成新机器的机器码运行。这是老程序能“莫名其妙”在新显卡上跑起来的原因之一。

3.3 拿到新机器，先跑几行代码“验明正身”

我现在的习惯是，任何一台新机器到手，先写一个十几行的小程序把硬件底细打印出来。核心就是cudaGetDeviceProperties：

#include <cstdio> #include <cuda_runtime.h> int main() { cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); printf("GPU: %s\n", prop.name); printf("Compute Capability: %d.%d\n", prop.major, prop.minor); printf("SM count: %d\n", prop.multiProcessorCount); printf("Max threads per block: %d\n", prop.maxThreadsPerBlock); printf("Shared memory per SM: %zu bytes\n", prop.sharedMemPerMultiprocessor); printf("Global memory: %.1f GB\n", prop.totalGlobalMem / 1024.0 / 1024.0 / 1024.0); return 0; }

编译之后会输出类似“Compute Capability: 8.6，SM count: 28”这样的信息。看到输出，再决定后面编译参数和优化策略往哪个方向走。这一步花不了一分钟，却能省下后面好几个小时的瞎试。

4. SM 与内存层次：代系变化背后的“暗线”

4.1 SM 内部零件一直在变

每个 SM 的基本组成有 CUDA Core、调度器、寄存器堆、共享内存、L1 缓存，还有后来的 Tensor Core 和 RT Core。代际更迭时，这些部件的数量配额和协同方式都在变。

Fermi 的 SM 很简约，32 个 CUDA Core 带一个共享内存单元；Kepler 把“大 SM”路线做到一个高峰，SMX 里塞了 192 个 CUDA Core；Maxwell 考虑到能效，又回到小 SM 多 SM 的路子；Volta 引入了独立线程调度，同一 warp 内的线程不必再像以前那样必须同进同退。对写代码的人来说，这些变化最直接的体感是：以前必须小心翼翼避免 warp 内分支发散，Volta 之后发散没以前那么致命，但想拿满性能，最优写法仍然要照顾到 warp 的一致性。

Ampere 的消费级 SM 是又一个转折点。FP32 和 INT32 拆成独立管道后，很多类型代码的普通浮点吞吐翻倍。如果你只盯着“流处理器总数”看，会觉得提升不大，但跑起 FP32 密集的算子，差距特别明显，因为它能同时执行浮点和整数指令。

4.2 缓存、显存与带宽，谁才是主角

GPU 性能的另一个发动机是显存带宽。旗舰级 GPU 的带宽常年是几百 GB/s 起步，如今能过 TB/s。为了喂饱这些带宽，L2/L1 缓存的容量、命中率策略也在不断进化。

Maxwell 时代就发现，与其做一堆复杂的能耗换一点缓存命中率，不如简化设计、腾出空间给并行单元，于是它成了能效大跃进的一代。Pascal 开始把统一内存虚拟寻址做得更实用，CPU 和 GPU 可以共享一套地址空间，拷贝数据的繁琐度大幅下降。Hopper 则把 NVLink 的带宽抬到很高，多卡训练时的数据交换速度成为核心卖点。

这种演进的启示是：你在做优化时，先问一句“这个瓶颈是算力还是访存”。如果是访存密集，那换一张带宽更高的卡或者优化访存模式，比优化计算指令更有效。很多新手盯着 ALU 指令使劲调，最终发现带宽早就跑满了，白忙一场。

4.3 用一次 occupancy 计算，看清代际差异在代码里的体现

Occupancy 指的是一个 SM 上活跃 warp 数与最大 warp 数的比值，它直接反映隐藏延迟的能力。我拿一个假设的 kernel 来演示代际差异如何影响结果。

假设每条线程用 40 个寄存器，每个 block 设 128 条线程，每 block 需要 100 字节共享内存。

• 老架构某卡，SM 最大线程数 1536，寄存器堆算下来能容纳约 1638 条线程，共享内存 48KB，则每个 SM 最多放 12 个 block，线程数 1536。此时占用率 = 1536 / 1536 = 100%。
• 新架构某卡，SM 最大线程数 2048，共享内存更大，但寄存器仍是每线程 40 个的上限，那么寄存器限制会让实际线程数卡在 1638，占用率约 80%。看起来“掉”了 20%，但新卡的每 SM 总吞吐能力更高，实际性能反而可能翻倍。

这个例子说明，占用率这个百分比不能脱离硬件代系孤立比较。同样一个 kernel，在不同代系里最优 block 尺寸可能完全不同。我踩过的坑是拿一份老项目参数直接套到新卡上，结果它在老架构上占用率很高，新架构上反而因为共享内存配额调整而变差。所以调优的时候，一定要对着cudaGetDeviceProperties的实测值重新算一遍。

5. 落地经验：同一段代码在不同代系上表现分化的原因与对策

5.1 共享内存 bank conflict 和 warp 调度

共享内存为了高吞吐，被分成了 32 个 bank。一个 warp 的 32 条线程同时访问共享内存时，理想情况是每条线程命中不同 bank，一拍完成；如果有两条线程访问同一个 bank 的不同地址，就会产生 bank conflict，需要多一拍去串行处理。这个规则 Fermi 时代就有，一直延续到今天。

不同代系的冲突惩罚不完全一样，Maxwell、Pascal 之后有些情况能自动合并优化，但代码本身写得“歪”了，在什么代系都慢。经典解法是填充 padding：

__shared__ float tile[32][33]; // 33 而不是 32，分散 bank 映射

把第 32 列改成 33 列，看似浪费了一点点内存，却能让后续每一行的访问自然错开 bank，很多二维 stencil 算子的速度立刻往上走。这种优化跟代系没太大关系，属于基本功，但能不能坚持写进代码里，才是老手和新手的区别。

5.2 Tensor Core 不是“免费的午餐”

Volta 之后，Tensor Core 成了 AI 算力的主引擎。但要注意，它只对特定数据布局和操作类型生效。普通 CUDA 代码如果不显式调用相关 API，或者不走 cuBLAS 的 TF32 路径，Tensor Core 是不会被用到的。

我在实际项目里发现，很多人买了一张 Ampere 卡，以为跑普通 FP32 运算也会自动翻倍提速，结果完全不是这么回事。普通浮点性能靠的是 SM 里的 FP32 单元数量，Tensor Core 只服务矩阵乘加类的规则密集计算。另外，如果在代码里用了 Tensor Core 特殊指令，老架构编译会直接失败。这时候一般用条件编译和运行时判断：

if (prop.major >= 8) { // 走 Tensor Core 路径 } else { // 走普通 CUDA Core 路径 }

这个分支不复杂，但能让同一份代码在不同代系之间安全落地。

5.3 我的一些实测习惯和建议

说了这么多，我把目前自己执行的一套流程总结在这里，基本可以复制使用。第一，新机器先跑 deviceQuery，把 Compute Capability、SM 数量、共享内存上限记录下来，别凭猜。第二，编译时用-gencode覆盖目标机型，不要图省事只写一个-arch。第三，调优前先问自己瓶颈类型，访存密集还是计算密集，再去动参数。第四，参考老教程里的固定数值时要小心，共享内存上限、最大 block 尺寸这些在不同代系间变量很大，一定要用实测值重新核算。

至于买显卡，如果主要做 AI 训练，优先看 Tensor Core 能力、显存容量和带宽；如果主力是通用并行计算、图像处理这类非张量任务，SM 数量和频率反而更关键。把这句话记在心里，能少花不少冤枉钱。

最后再分享一个小经验：不要迷信网上的 benchmark 平均值。同一张卡，驱动版本、CUDA 版本、编译参数不同，跑同一个 kernel 都可能差出 20% 到 30%。架构代系是硬件基础，但软硬件配合才是最终性能的放大器。弄清楚代系这件事，说到底就是让你在做每一个决定时，都知道自己站在哪一层硬件之上。