1. 为什么并行编程要先认 GPU 代系1.1 架构名就是 GPU 的“DNA”很多人在接触 CUDA 的时候第一个下意识去看的是显存多大、核心数多少、频率多高。这些数字当然有用但如果你写并行程序是为了跑性能真正决定上限的其实是 GPU 架构本身。英伟达每一代 GPU 的架构代号从 Kepler、Maxwell、Pascal、Volta、Turing到 Ampere、Hopper、Ada Lovelace再到现在的 Blackwell从来都不是随便起的名字。每一代都对应着 SMStreaming 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 引入了统一内存和 NVLinkVolta 引入了 Tensor Core 和独立线程调度Ampere 引入了异步拷贝cp.asyncHopper 又加入了线程块簇和 TMATensor 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 到 TuringTensor 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 视角里常见的几个架构整理成了表格方便你做快速对照。这里的“主要芯片代号”是硬件设计文档里的说法不一定对应最终市场上的产品名。架构代号主要芯片代号算力版本这代值得关注的并行特性KeplerGK104 / GK1103.0 / 3.5 / 3.7动态并行、SMX 结构MaxwellGM204 / GM1075.0 / 5.2 / 5.3能效优化、独立共享内存分配PascalGP100 / GP1026.0 / 6.1 / 6.2统一内存、NVLink、HBM2VoltaGV1007.0Tensor Core、独立线程调度TuringTU102 / TU1167.5RT Core、INT8/INT16 Tensor CoreAmpereGA100 / GA1028.0 / 8.6第三代 Tensor Core、cp.async、稀疏性HopperGH1009.0线程块簇、TMA、DPX 指令注意架构代号和算力版本并不总是一一对应。同一个大架构里不同芯片可能使用不同的次版本号甚至同一个芯片在不同市场下还会被屏蔽部分 SM。所以最靠谱的做法不是背表格而是用deviceQuery去拿真实 GPU 的major.minor。3.2 CUDA 编译选项和 PTX 兼容链CUDA 的编译过程是分层的。源代码先被编译器转成 PTXParallel Thread Execution这是一种虚拟指令集跟具体硬件无关。然后 PTX 再被 ptxas 编译成 SASS这才是 GPU 真正执行的机器码。你可以通过编译选项控制最终产物nvcc -archcompute_80 -codesm_80,compute_80 -o axpy axpy.cu这里compute_80表示编译器生成 CC 8.0 的 PTXsm_80表示生成 Ampere 架构的 SASS。最终文件里会同时塞进 SASS 和 PTX 两套代码。运行的时候如果驱动发现当前 GPU 的算力版本正好适配sm_80直接用 SASS如果当前 GPU 是更新的架构没有现成 SASS就会尝试用 PTX 做 JIT 编译生成一份适配当前 GPU 的代码。这个机制特别重要。很多人只写-archsm_90导致在旧卡上直接报invalid device function。反过来只写compute_80会让每次运行都触发 JIT首帧性能很差。稳妥的做法是多个目标组合编译nvcc -gencode archcompute_80,codesm_80 \ -gencode archcompute_80,codecompute_80 \ -gencode archcompute_90,codesm_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 里的deviceQuerycd /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 -archsm_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_asyncgrid, block, shared_bytes(...); } else { kernel_legacygrid, 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 正在计算形成两段流水线。这个版本的数值表现和性能如下优化阶段Pascal 上耗时模拟值Ampere 上耗时模拟值主要瓶颈朴素全局访存4.01 ms3.42 ms访存延迟32x32 tile shared2.85 ms2.10 ms同步和寄存器溢出tile cp.async 双缓冲2.88 ms1.21 msPascal 无法使用 cp.async这个例子说明架构代系直接决定了你能选择的优化手段。同一份代码如果只用通用方式Ampere 相对 Pascal 的提升不会特别夸张一旦针对 Ampere 的异步拷贝特性重写差距立刻拉到 2 倍以上。5. 常见问题与避坑实录5.1 编译时指定了过高的算力版本最典型的问题是代码在别人的新卡上编译通过了拿回自己的老卡上运行结果一启动就报cudaErrorInvalidDeviceFunction。原因是编译时只生成了新架构的 SASS老卡没有对应的机器码。排查思路很简单。第一用cudaGetDeviceProperties确认目标卡 CC第二看 NVCC 编译命令里-gencode是否包含compute_XX的 PTX 兜底第三不要只用-archsm_90至少要加上一个compute_XX。如果你希望程序在多个代系上都跑得不错最省事的就是编译多套 fatbin。nvcc -gencode archcompute_80,codesm_80 \ -gencode archcompute_86,codesm_86 \ -gencode archcompute_90,codesm_90 \ -gencode archcompute_90,codecompute_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返回了设备不支持的指令。我见过不少坑最后总结成一张排查表现象可能原因解决思路编译成功运行时报 invalid device functionfatbin 里没有当前 GPU 可用的 SASS也没有可 JIT 的 PTX增加compute_XXPTX 或对应架构的 SASS同一个 kernel 在新老卡上性能差距巨大没有针对架构差异做访存和异步优化先看内存工作量再考虑cp.async和 Tensor Core老卡上跑出了错误结果__CUDA_ARCH__判断被误用走了不支持的设备指令路径用主机侧cudaGetDeviceProperties做运行时分支显存带宽跑不满访存没有合并或 tile 大小跟 L2 不匹配检查访问 pattern尝试 128 字节对齐、增大 L2 命中率个人心得这套架构代系的东西我最早也是死记硬背后来发现真正帮我记住的是每次换卡后重新编译、重新调优的过程。拿到新卡先别急着看频率先看 CC 和 SM 数量再想算法里哪个环节能吃到新架构的红利这样优化才有方向。另外一个很实用的小技巧写共享代码时尽量把“纯通用路径”和“架构专属路径”分开不要为了炫技把所有代码都绑死在最新特性上。通用路径保证所有卡都能正确运行专属路径通过分支选择在新卡上把性能拉满。这样你调试的成本最低别人用起来也最省心。并行编程里硬件和软件从来不是隔开的。理解了英伟达 GPU 架构与代系关系很多原本要花几个小时定位的问题其实一眼就能看出原因。