CUDA并行编程必知：英伟达GPU架构与代系关系解析

1. 为什么并行编程要先认 GPU 代系

1.1 架构名就是 GPU 的“DNA”

很多人在接触 CUDA 的时候，第一个下意识去看的是显存多大、核心数多少、频率多高。这些数字当然有用，但如果你写并行程序是为了跑性能，真正决定上限的其实是 GPU 架构本身。英伟达每一代 GPU 的架构代号，从 Kepler、Maxwell、Pascal、Volta、Turing，到 Ampere、Hopper、Ada Lovelace，再到现在的 Blackwell，从来都不是随便起的名字。每一代都对应着 SM（Streaming Multiprocessor）内部结构、缓存层级、调度器、特殊计算单元、访存通路甚至虚拟内存模型的改变。

打个比方，CPU 时代你看一颗处理器，除了主频和核心数，还得问一句“第几代”，因为代际决定了指令集扩展、单核 IPC、内存控制器效率。GPU 更是这样，而且比 CPU 更极端：两代架构之间，同样是“几千个 CUDA 核心”，实际跑一次矩阵乘可能差出 3 倍以上。原因很简单，核心数量只是表面，每个时钟周期里这些核心能不能吃饱数据、能不能同时读写共享内存、能不能利用专用的矩阵单元，才决定真实速度。

所以这一章的主题很明确：把英伟达 GPU 架构与代系关系搞清楚。这是“掌握 GPU 并行编程”系列的第三章，前两章讲线程模型和内存模型，第三章就是回到硬件底座。没有这一层，你会经常遇到同样的 kernel 在不同的卡上表现天差地别，却完全不知道去哪找原因。

1.2 算力版本不是跑分，是兼容性接口

在看 CUDA 文档的时候，你会频繁遇到一个组合：compute_80、sm_80、compute_86、sm_86之类。这个数字其实是“算力版本”（Compute Capability），简写为 CC。格式是主版本号.次版本号，比如 8.0、8.6、9.0。

主版本号代表大架构更新，次版本号代表同一代架构里的小幅优化或功能裁剪。比如 Ampere 这一代，数据中心级的 A100 对应 CC 8.0，消费级的一些 Ampere 卡则对应 CC 8.6。它们都属于 Ampere 大代系，但 SM 内部的指令调度、L2 大小、Tensor Core 代数并不完全一样。

很多新手容易把 CC 当成性能评分，看到 8.6 比 8.0 高，就以为 8.6 更“强”。其实 8.6 和 8.0 只是功能集不同，不是同一条线的数字大小关系。8.0 的核心和 8.6 的核心往往面向不同市场，一个侧重高算力调度，一个可能更注重图形和混合精度。真正决定代码写法的，是你编译时指定的目标算力版本，以及运行时驱动的 JIT 能力。

1.3 代系决定了你能用哪些“新玩具”

为什么必须在写代码前先确认 GPU 代系？因为这些年并行计算的玩法变化太大了。早年的 CUDA 程序，基本就是启动线程、读写显存、手工用__syncthreads()同步。到了 Pascal 引入了统一内存和 NVLink，Volta 引入了 Tensor Core 和独立线程调度，Ampere 引入了异步拷贝cp.async，Hopper 又加入了线程块簇和 TMA（Tensor Memory Accelerator）。

也就是说，如果你手里是一张老架构卡，很多新写法根本编译不进去；如果你手里是新卡，却还用十年前的写代码方式，性能也一定不理想。代系关系不是一个表格里躺着的东西，它直接影响你每一步优化手段的选择。

2. 从 Kepler 到 Ampere：架构演进主线

2.1 Kepler 到 Pascal：从堆核心走向能效和互联

英伟达 GPU 架构的早期阶段，思路比较直接：通过增加 SM 数量和 CUDA 核心数来提升吞吐。Kepler 这一代把 SM 结构改成了 SMX，目的在于提高单 SM 的计算密度，同时加入了动态并行能力，让 GPU 内部可以自己启动 kernel，避免频繁回 CPU。动态并行听起来很高级，但实际开发里限制不少，很多场景下反而引入调度复杂度，所以在后来的架构里用得并不算多。

到了 Maxwell，设计目标明显转向能效比。它把 SM 重新组织成更灵活的分区结构，不用像以前那样为了榨干核心数而死磕指令发射带宽。结果就是同样功耗下，性能反而提升明显，而且编程模型没有发生剧烈变化。很多早期的并行程序在 Maxwell 上跑，往往只需要重新编译就能获得不错的加速。

Pascal 是一个重要的分水岭。它第一次在专业卡上大规模使用 HBM 显存，带宽一下子拉高到 700GB/s 以上。同时也引入了 NVLink，让多卡之间的通信不再是 PCIe 那点可怜带宽。对于并行编程来说，Pascal 的出现意味着“多 GPU 程序”和“统一寻址”开始成为可落地的事情。cudaMallocManaged统一内存也是从这个时期开始被更多人拿来做数据迁移，虽然性能要小心调，但开发体验确实好多了。

从 CPU 视角看，这三代其实是在摸索一条主线：先把计算能力做上去，再把数据通路做宽。在 Pascal 时代，你写一个完全靠全局内存访存的 kernel，只要做好合并访问，基本可以跑满带宽。但如果你没有做访存合并，再宽的带宽也救不了你。

2.2 Volta 到 Turing：Tensor Core 正式登场

Volta 这一代在通用计算领域影响力极大，原因有两个。

第一个原因是独立线程调度。早期 CUDA 执行模型里，一个 warp 的线程基本是锁步的，遇到分支时会做掩码操作，某个线程走到 A 分支、其他线程被掩蔽，整体还是同步推进。Volta 改成每个线程有自己的程序计数器，调度粒度更细。这带来的最大好处是，某些分支复杂的代码不再那么怕发散，编译器可以更大胆地做指令级优化。当然，__syncthreads()语义也变得更复杂了，一个 warp 里的线程可能不在同一个执行点，所以千万不能假设线程天然同步。

第二个原因是 Tensor Core。这个名字现在无人不知，但第一次出现时，很多写通用并行计算的人没太当回事。Tensor Core 不是用来跑通用逻辑的，它是专门为矩阵乘法和卷积设计的专用硬件，一个时钟周期能干很多次浮点乘加。对 AI、深度学习、科学计算里的线性代数运算，性能直接起飞。

到了 Turing，英伟达又把 Tensor Core 改进了一轮，同时加入了 RT Core。RT Core 是光线追踪专用硬件，对通用并行编程来说没什么直接意义，但 Turing 的 Tensor Core 开始支持 INT8、INT16 等更低精度运算，这让很多推理任务可以在不那么高的精度下获得巨大吞吐。并行编程的思维也开始变化：同一份数据，用 FP32 算、用 TF32 算、用 FP16 算，得到的性能可能差一个数量级，而你是否决定用混合精度，完全取决于算法误差能否接受。

2.3 Ampere 之后：异步编程成了通关钥匙

Ampere 这一代最值得并行程序员关注的，不是又多了多少核心，而是cp.async。在它出现之前，你要把全局内存里的数据搬到共享内存，一般得写__ldg或者普通的load到寄存器，再store到共享内存，期间线程会被卡在访存等待上。cp.async不一样，它是一条异步拷贝指令，发出去之后线程不用等数据到齐，可以继续做其他事，后续再用commit_group和wait_group把数据同步回来。这个机制配合 double buffering，能让访存延迟被掩盖得很漂亮。

Ampere 还引入了细粒度稀疏性支持，Tensor Core 在做矩阵运算时可以跳过数值为零的元素，理论吞吐能再翻倍。当然，实际达到多少要看数据的稀疏率和存储格式是否配合，但至少说明硬件对“结构化的数据跳过”开始有原生能力了。

Hopper 则把异步和集群概念再往前推了一步，加入了线程块簇（Thread Block Clusters），让多个线程块可以共享分布式共享内存，并且 TMA 可以直接在全局内存和共享内存之间做 bulk 传输。这意味着很多以前需要手工循环搬运的数据，现在可以一口气交给硬件去搬。你在写 kernel 的时候，不再只是选 block 大小，而是要考虑 cluster 怎么组织、TMA 的 descriptor 怎么配。

我个人的感觉是，Ampere 之后的并行编程已经进入了一个新阶段：谁能把内存搬移和计算重叠得好，谁就能赢。单纯比较 CUDA 核心数量已经没有多少意义了，异步性、内存通路、张量单元利用率才是真正的分水岭。

3. 软件栈和代系之间，到底怎么对上暗号

3.1 一张表认清架构与算力版本

写 CUDA 程序时，最基础的动作是确认目标算力版本。我把 CPU 视角里常见的几个架构整理成了表格，方便你做快速对照。这里的“主要芯片代号”是硬件设计文档里的说法，不一定对应最终市场上的产品名。

注意，架构代号和算力版本并不总是一一对应。同一个大架构里，不同芯片可能使用不同的次版本号，甚至同一个芯片在不同市场下还会被屏蔽部分 SM。所以最靠谱的做法，不是背表格，而是用deviceQuery去拿真实 GPU 的major.minor。

3.2 CUDA 编译选项和 PTX 兼容链

CUDA 的编译过程是分层的。源代码先被编译器转成 PTX（Parallel Thread Execution），这是一种虚拟指令集，跟具体硬件无关。然后 PTX 再被 ptxas 编译成 SASS，这才是 GPU 真正执行的机器码。

你可以通过编译选项控制最终产物：

nvcc -arch=compute_80 -code=sm_80,compute_80 -o axpy axpy.cu

这里compute_80表示编译器生成 CC 8.0 的 PTX，sm_80表示生成 Ampere 架构的 SASS。最终文件里会同时塞进 SASS 和 PTX 两套代码。运行的时候，如果驱动发现当前 GPU 的算力版本正好适配sm_80，直接用 SASS；如果当前 GPU 是更新的架构，没有现成 SASS，就会尝试用 PTX 做 JIT 编译，生成一份适配当前 GPU 的代码。

这个机制特别重要。很多人只写-arch=sm_90，导致在旧卡上直接报invalid device function。反过来只写compute_80，会让每次运行都触发 JIT，首帧性能很差。稳妥的做法是多个目标组合编译：

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

这样既有针对 Hopper 的原生代码，又有 Ampere 的 SASS 和 PTX 兜底，老卡也能跑，只是老卡可能走 JIT 或原生 SASS，性能不会太差。

3.3 在代码里给不同代系写“专属路径”

除了编译选项，你还可以在源代码里用__CUDA_ARCH__这个宏判断当前编译目标。它在设备侧编译时会被自动定义成算力版本数字，比如编译sm_80时等于 800，编译sm_90时等于 900。

典型用法是这样的：

__global__ void kernel(const float* in, float* out) { #if __CUDA_ARCH__ >= 800 // Ampere 及以上，可以用 cp.async 的思路 #else // 老架构走普通访存路径 #endif }

但要注意，__CUDA_ARCH__是编译期常量，不是运行期变量。同一时刻只能反映“当前这段代码正在为哪个架构编译”，不会理解为你手头 GPU 的实际型号。所以如果要根据运行时真实 GPU 做分支，还是得靠cudaGetDeviceProperties拿major和minor。

4. 实操：确认当前 GPU 属于哪个代系

4.1 一条命令和小程序拿到全部信息

拿到一台新机器，我习惯先做三件事：看系统里的 GPU 名称、确认算力版本、确认 SM 数量。nvidia-smi能显示名称和显存，但算力版本不一定直接列出来，最可靠的方式是用 CUDA 自带的例子。

如果你装了 CUDA Toolkit，可以找到 samples 里的deviceQuery：

cd /usr/local/cuda/samples/1_Utilities/deviceQuery sudo make ./deviceQuery

它会打印出 CUDA 驱动版本、GPU 名称、算力版本、SM 数量、共享内存大小、全局显存大小、内存时钟这些信息。对于并行程序设计，这是真正的起点。

如果手边只有 Python，也可以用 PyTorch 快速看：

import torch if torch.cuda.is_available(): print(torch.cuda.get_device_name(0)) print(torch.cuda.get_device_capability(0))

输出类似(8, 6)，表示 CC 8.6，也就是 Ampere 架构的消费级芯片。PyCUDA 也可以：

import pycuda.driver as drv drv.init() dev = drv.Device(0) print(dev.name(), dev.compute_capability())

如果都不想装，写一个几十行的 CUDA C 程序也可以：

#include <cstdio> #include <cuda_runtime.h> int main() { int dev = 0; cudaDeviceProp prop; cudaGetDevice(&dev); cudaGetDeviceProperties(&prop, dev); printf("%s\n", prop.name); printf("CC: %d.%d\n", prop.major, prop.minor); printf("SMs: %d\n", prop.multiProcessorCount); printf("Global mem: %.2f GB\n", (double)prop.totalGlobalMem / 1024.0 / 1024.0 / 1024.0); return 0; }

编译时不需要指定太高的架构，随便选一个能覆盖目标卡的就行：

nvcc -arch=sm_80 device_query.cu -o device_query ./device_query

你会看到类似：

某型号 Ampere GPU CC: 8.6 SMs: 82 Global mem: 23.99 GB

这些数据是后续所有优化判断的基准。

4.2 运行时判断加编译期分支的组合打法

真正的生产代码里，编译期分支和运行期分支往往要一起用。编译期__CUDA_ARCH__保证代码能在不同架构下生成不同的设备指令；运行期cudaGetDeviceProperties保证主机侧知道当前用户到底插了什么卡。

一个常见的模式是：在主机侧准备配置结构体，根据prop.major决定 kernel 启动的 grid 和 block 大小，或者决定选哪个 kernel 函数。比如在 Ampere 及以上用支持cp.async的 kernel，在老架构上退回到普通版本。

cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); if (prop.major >= 8) { kernel_async<<<grid, block, shared_bytes>>>(...); } else { kernel_legacy<<<grid, block, shared_bytes>>>(...); }

这个判断不是可有可无。同一个算法，在 Ampere 上不用异步拷贝，性能可能差距在 20% 到 40% 之间；在老架构上你强制使用新特性，直接编译失败或者运行报错。所以“检测代系 + 分支”不是炫技，是兼容性的基础。

4.3 一个可复现的调优对比案例

为了说明代系对代码的影响，我在一次内部性能调试里整理过一个模拟数据：同一份矩阵乘法 kernel，分别在 Pascal 和 Ampere 上跑。

第一版是朴素写法，线程按行分块读取全局内存，不做显式的 shared memory 重用。结果在 Pascal 上跑得还算中规中矩，但在 Ampere 上也没有快很多，因为瓶颈根本不在计算，而在访存。第二版写了一个 32x32 的 tile 版本，用 shared memory 缓存子矩阵，并且每个线程负责 4 个输出元素。这个版本在 Pascal 上明显提升了，Ampere 上也提升，但线程块大了之后同步开销开始变大。

第三版换成 Ampere 专属路径：使用cp.async把下一个 tile 提前搬进 shared memory，同时当前 tile 正在计算，形成两段流水线。这个版本的数值表现和性能如下：

这个例子说明，架构代系直接决定了你能选择的优化手段。同一份代码，如果只用通用方式，Ampere 相对 Pascal 的提升不会特别夸张；一旦针对 Ampere 的异步拷贝特性重写，差距立刻拉到 2 倍以上。

5. 常见问题与避坑实录

5.1 编译时指定了过高的算力版本

最典型的问题是：代码在别人的新卡上编译通过了，拿回自己的老卡上运行，结果一启动就报cudaErrorInvalidDeviceFunction。原因是编译时只生成了新架构的 SASS，老卡没有对应的机器码。

排查思路很简单。第一，用cudaGetDeviceProperties确认目标卡 CC；第二，看 NVCC 编译命令里-gencode是否包含compute_XX的 PTX 兜底；第三，不要只用-arch=sm_90，至少要加上一个compute_XX。如果你希望程序在多个代系上都跑得不错，最省事的就是编译多套 fatbin。

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

这个命令产出体积大一点，但换来的兼容性很值。尤其当你做的是给别人用的库，别想着偷懒只针对最高配编译。

5.2 把“CUDA 核心数”当成唯一指标

很多产品参数特别喜欢强调“几千个 CUDA 核心”。但如果你真的写并行程序，会发现核心数多并不等于你快。架构不同，单个核心的 FMA 单元数量、寄存器文件大小、调度器效率、访存延迟完全不同。

典型的例子是，Pascal 和 Ampere 同样叫“CUDA 核心”，但 Ampere 的 SM 支持cp.async、异步屏障、更宽的张量指令。一个矩阵乘 kernel 在 Ampere 上如果只走通用 FP32 路径，可能只有 Moore 架构下原生 Tensor Core 路径的三分之一性能。所以看一个 GPU 能不能跑好你的程序，先看架构再看核心数，顺序不能反。

5.3 新特性直接写进老卡分支，不判断可用性

有人写了__CUDA_ARCH__ >= 800的分支，结果在 Pascal 上运行时发现，编译期没有报错，运行期却出了问题。原因在于__CUDA_ARCH__判断的是“当前正在编译的目标架构”，不是“运行时的实际 GPU”。

解决办法我再说一遍：设备侧分支用__CUDA_ARCH__，主机侧分支用cudaDeviceProp.major/minor。两者不能互相替代。遇到诡异行为时，先用cudaGetLastError()拿到返回码，看是不是从cudaLaunchKernel或cudaMemcpyAsync返回了设备不支持的指令。

我见过不少坑，最后总结成一张排查表：

个人心得

这套架构代系的东西，我最早也是死记硬背，后来发现真正帮我记住的，是每次换卡后重新编译、重新调优的过程。拿到新卡，先别急着看频率，先看 CC 和 SM 数量，再想算法里哪个环节能吃到新架构的红利，这样优化才有方向。

另外一个很实用的小技巧：写共享代码时，尽量把“纯通用路径”和“架构专属路径”分开，不要为了炫技把所有代码都绑死在最新特性上。通用路径保证所有卡都能正确运行，专属路径通过分支选择，在新卡上把性能拉满。这样你调试的成本最低，别人用起来也最省心。

并行编程里，硬件和软件从来不是隔开的。理解了英伟达 GPU 架构与代系关系，很多原本要花几个小时定位的问题，其实一眼就能看出原因。