如果你现在随手打开任何一份 GPU 计算相关的招聘要求几乎都能看到“熟悉 CUDA 执行模型”这一条。但把时间拨回 2010 年情况完全不是这样。那时候的 GPU 虽然有很强的数学吞吐可要说“通用计算架构”总感觉差点意思。直到英伟达发布 Fermi也就是第一代真正意义上的完整 GPU 计算架构这个局面才被扭转。Fermi 对应的产品是 GTX 480、GTX 470以及面向计算市场的 Tesla 20 系列。核心芯片叫 GF100采用 40nm 工艺晶体管数量达到 30 亿。很多人只知道它是一块很热的老显卡却不理解它为什么被称为“架构里程碑”。这篇文章我想用从业者的视角把 Fermi 从头拆一遍它诞生前 GPU 计算的困境、GF100 的内部设计逻辑、真实产品的代价以及它对后来十年 GPU 计算的深远影响。这篇笔记适合几类人刚接触 GPU 计算的在校学生、在做计算卡选型时犹豫的工程师、单纯好奇“为什么一张 2010 年的显卡能成为里程碑”的硬件爱好者。我不打算堆参数表而是想讲清楚每一处设计选择背后的“为什么”。1. Fermi 诞生之前为什么 GPU 计算一直“缺一块完整的骨架”1.1 从 G80 的“凑合”到 GT200 的“算力有余、架构欠账”英伟达在 2006 年发布 G80 时做了一件革命性的事把 GPU 里各自独立的顶点处理器和像素处理器统一成了一种可编程的流处理器。同一个硬件既可以算顶点也可以算像素并且第一次允许开发者用类 C 语言写通用代码。G80 就是 CUDA 的起点。但起点归起点G80 骨子里还是一款图形卡。它的计算部分更像是“在图形流水线上加了个通用后门”寄存器文件不够大共享内存容量极其有限L2 缓存被图形任务拆分程序写起来有一堆潜规则。到了 2008 年的 GT200英伟达把流处理器数量从 128 个提高到 240 个单精度浮点峰值大幅提升但芯片底层的调度模型、内存层次、并发能力基本没有本质变化。结果是算力翻倍利用率却未必翻倍。很多算法换成 GT200 后确实更快但瓶颈还是那个瓶颈——数据喂不进去或者喂进去之后取不出来。这好比一辆赛车G80 时代给你装了个宽敞的副驾驶位但整台车的底盘、悬挂、刹车还是按赛道用途调校的。GT200 换了个更大排量的发动机可底盘没动开着还是别扭。真正的问题在于GPU 原本是一台“为图形渲染设计的机器”想要把它变成“为通用计算设计的机器”需要连底盘一起重画。这个底盘就是 Fermi 做的事。1.2 计算卡用户的三个痛点数据局部性、可靠性与并发能力在 Fermi 之前不少做科学计算和数据分析的团队已经拿 GPU 测试过一些算法结论往往是“快但是难用”。我总结下来痛点集中在三个地方。第一个是数据局部性。共享内存是 GPU 上程序员能显式控制的缓存G80/GT200 时代每个 SM 只有 16KB而且不可配置。做分块矩阵乘法、规约、直方图这类高复用算法时要么小心翼翼反复把数据切小要么干脆放弃共享内存直接打全局内存。全局内存的延迟又高程序性能忽上忽下。第二个是可靠性。长时间运行的科学计算最怕内存翻转。游戏卡运行几小时后如果数据出错画面也就一帧的瑕疵重开就是了但一个跑了两天的分子动力学模拟如果某次访存被宇宙射线干扰得出的整个结果作废。Fermi 之前 GPU 几乎没有任何内存纠错保护这在专业计算领域是不可接受的。第三个是并发能力。老架构同一时刻往往只能让一个内核跑满整个 GPU。当你只想往 GPU 里塞一个小任务时其他硬件都在空转浪费严重。这三个痛点单独看都是“可以忍”的但放一起GPU 就永远只是“加速器玩具”而不是“正规计算设备”。Fermi 把这三大块一起补上了所以它才配叫“第一个完整 GPU 计算架构”。2. GF100 的内部兵法SM、GPC 与缓存架构到底改了什么2.1 SM 单元的重构32 个 CUDA 核心与真正的双精度GF100 整颗芯片有 512 个 CUDA 核心分成 16 个 SM每个 SM 包含 32 个 CUDA 核心。别小看这个“32”它是第一个能在一个 SM 内完整容纳一条线程束warp32 个线程的设计。此前每个 SM 只有 8 个流处理器一条 warp 要跨多个 SM 执行线程间通信要靠全局内存中转慢得离谱。Fermi 让 warp 的概念真正落到一个物理单元内部调度效率完全不在一个量级。每个 SM 内部有 4 个 Warp 调度器每个调度器可以独立从就绪线程束中挑选指令送入执行单元。配合 32 个 CUDA 核心理论上一个周期能发射多条线程束指令。这意味着 SM 的流水线可以获得更好的填充率小的分支、延时、访存等待都能被其他线程束掩盖。对程序员来说最直观的感受就是“同样的算法怎么换个架构就快了这么多”。双精度方面Fermi 每个 SM 内置 4 个双精度单元支持 FMA 指令——乘法和加法一条指令完成精度和速度都更好。整体双精度吞吐大约是单精度的 1/8。这个数字放到今天看一般但在当时是跨越式的此前 GPU 的双精度性能弱到几乎不能用于严肃计算Fermi 第一次让 GPU 可以做大规模高精度数值计算。2.2 共享内存加 L1 的“可裁剪缓存”以及那颗巨大的 768KB L2Fermi 在每个 SM 上放了 64KB 的 SRAM它可以灵活划分给共享内存和 L1 缓存。你既可以切成 16KB L1 48KB 共享内存让显式复用占大头也可以切成 48KB L1 16KB 共享内存让自动缓存处理更多访存模式。这种“可裁剪”设计非常重要因为不同算法对缓存的需求完全不同。分块矩阵乘法想要尽量大的显式暂存区而稀疏数据访问可能更需要自动缓存来吸收随机性。老架构是“给什么吃什么”Fermi 是“你要什么我给你配什么”。真正的杀手锏是那颗 768KB 的统一 L2 缓存。作为对比GT200 的 L2 只有 256KB而且不少路径是为图形渲染设计的。Fermi 的 L2 是被所有 SM 共享的统一缓存一个 SM 写入的数据另一个 SM 可以快速读到。GPU 终于有了和 CPU 三级缓存类似的“最后一公里”数据交换通道。同时64 位和 128 位等不同访存模式都能在 L2 层面更好地交织。这个设计让很多以前必须反复访问显存的算法现在可以命中 L2能耗和延迟都大幅下降。我后来在调内核时最常做的一件事就是调整数据布局去“喂饱”这颗 L2效果往往立竿见影。2.3 核内并发、全局原子操作与内存纠错通用计算的“三大件”如果说 SM 和缓存是硬件底座那么 Fermi 在功能特性上做的三件事才是它彻底从“图形卡”跨入“计算卡”的关键。第一是并发内核执行。Fermi 首次允许同一个上下文中的多个内核同时运行。你可以创建多个 CUDA 流stream把不同规模的内核放进不同的流里小内核会自动填补大内核留下的执行空隙。以前 GPU 一次只能跑一个内核像一条单车道公路Fermi 直接修成了多车道小任务、大任务可以并行交错。第二是更完善的原子操作。全局内存和共享内存的原子操作在 Fermi 上得到统一支持且性能明显改善这使得哈希表、计数器、排序等复杂并行算法终于可以高效落地。第三就是前面提到的 ECC 内存纠错。Fermi 在显存、寄存器和缓存路径上加入了纠错支持可以纠正单比特错误、检测双比特错误。这三样让 GPU 第一次具备了“可靠地大规模执行通用计算”的完整能力。我把这套组合称为“三大件”调度/并发、原子操作、纠错。调度决定利用率原子操作决定表达能力纠错决定可信度。三者齐备GPU 才敢真正承接核心业务负载。3. 从纸面参数到实测红利CUDA 时代“第一个完整计算栈”的意义3.1 编译器和编程模型的质变你可以用新思维写内核Fermi 不只是硬件变化它配套的 CUDA 工具链也产生了质变。更聪明的编译器可以更好地完成寄存器分配开发者可以更自由地使用共享内存、L2 缓存和流并发这些显式能力。我记得第一次在 Fermi 卡上写流式并发代码时感觉生产力完全不一样了cudaStream_t stream1, stream2; cudaStreamCreate(stream1); cudaStreamCreate(stream2); kernelAgrid, block, sharedMemBytes, stream1(dataA); kernelBgrid, block, sharedMemBytes, stream2(dataB); cudaStreamSynchronize(stream1); cudaStreamSynchronize(stream2);这段代码在 Fermi 之前也能写但内核实际上很难真正并行执行。到了 Fermi我亲手测过两个规模悬殊的内核放在不同流里的情况小内核会在大内核的空隙之间穿插执行整卡吞吐提升肉眼可见。这种“新思维”还体现在共享内存的动态配置上。以前改一个块大小就可能撞上 16KB 的限制逼着你去拆算法在 Fermi 上我只需在每次启动内核前按需配置共享内存大小算法结构不用大改。配套工具也在同期完善。事件探查器profiler可以对每个 SM 的微指令事件计数调试器可以查看每个线程的寄存器和共享内存状态。一个“完整计算栈”真正成型了语言、编译器、运行时、库、调试工具全部到位。3.2 实测带来的红利从流体求解器到并发流式处理纸面参数再好也要看真实负载。我在一轮实测里印象最深的是流体求解器。这个程序的核心是一个三维格点迭代每个格点计算时需要读取邻近六个格点的数据。老架构上我把数据按块切小才能塞进共享内存总会引入大量边界同步。换成 Fermi 后共享内存上限从 16KB 提到 48KB一个块能容纳的数据范围大了一截边界交换次数明显减少再加上 768KB L2 兜底块与块之间的热数据复用也更容易命中缓存。最终整个内核的性能提升不只来自核心数量变多更来自数据结构可以保持更自然的形态。另一个典型场景是流式处理。某一次处理一批大小不一的异步任务时我把不同耗时的调用拆分到不同流部分任务在 PCIe 传输的同时其他内核已经在 GPU 上执行。借助 Fermi 的并发内核执行和异步拷贝整条流水线几乎一直处于满载状态。这种优化模式在今天的 CUDA 编程里非常普遍但在当时是“终于可以放心用了”。如果只盯着峰值浮点性能Fermi 的优势会被低估。它在访存密集型、短任务并发、混合流水线这类“非理想条件”下的表现才是架构完整的真正证明。4. 真实产品与真实代价GTX 480 的发热、良率与 Tesla 20 系列的取舍4.1 消费卡 GTX 480完整架构的阉割与功耗代价GF100 完整设计是 512 个 CUDA 核心但第一代消费卡 GTX 480 实际对外发布的是 480 个核心——英伟达屏蔽了一个 SM原因很简单功耗和良率。40nm 工艺在 2010 年刚刚成熟制造一颗 30 亿晶体管的芯片难度极高。即便屏蔽了核心GTX 480 的 TDP 也大约在 250W 上下满载时散热器风扇转速高到像机箱里多了台吹风机。我记得自己第一次跑满负载测试盯着温度监控曲线确实有点怀疑人生。但 GTX 480 又是当时单精度性能最强的消费级 GPU。很多兜里预算有限的开发者都会买它来跑 CUDA。它的价值在于让普通开发者第一次用相对低廉的价格接触到完整 Fermi 架构的核心能力可配置共享内存、统一 L2、原子操作、并发内核。至于双精度和 ECC游戏卡能跑出个影子但真正能发挥全力的场景还是在专业卡上。所以如果你想用 Fermi 做严肃计算老老实实选 Tesla 20 系列才是正路。4.2 专业卡 Tesla 20 系列ECC 计算的正确打开方式Tesla 20 系列有多个型号显存从 3GB 到 6GB 不等面向的是工作站和服务器市场。它与游戏卡最大的区别不在于核心数而在于稳定性设计。ECC 支持的加入意味着显存错误可以被检测和纠正这对科学计算是刚需。Tesla 20 系列还经过更严格的工作负载验证适合长时间高负载运行。有人会问同样都是 GF100游戏卡和计算卡差距有那么大吗我的回答是物理芯片可能同源但验证标准、驱动支持、显存配置和可靠性设计决定了两者的应用边界。游戏卡跑三天三夜出一次静默错误可能谁都没察觉可一旦计算程序在关键迭代步里得到错误数据整套结果报废。专业计算选型本质上买的是“确定性”。HPC 用户也更看重双精度。Tesla 20 系列虽然双精度同样是单精度的 1/8 左右但表现稳定配合 ECC 与完整驱动栈让大量涉及线性代数、有限差分、分子动力学的程序第一次真正迁移到了 GPU 上。这里我特别想提醒做选型的工程师如果只是跑深度学习推理单精度游戏卡够了但凡是长时间高精度仿真务必把 ECC 当成硬指标。4.3 半代产品的意义计算市场的试错与定型Fermi 的激进也带来了代价。300mm²级别的芯片面积、40nm 初期的良率、高功耗散热压力这些都让英伟达在第一代产品上做了不少妥协。GTX 480 后期其实还经历过一次步进改良到了后续基于 GF110 的产品才把功耗温度拉回合理区间。但“第一个完整架构”本来就不是为了第一代产品的一步到位。Fermi 的意义在于它用一次大的架构重构换来了后续十年 GPU 计算的稳步演进。虽然很多消费者记住的是“GTX 480 很热”但工程师记住的是第一次能在一个可编程、可纠错、可并发调度的 GPU 上构建严肃计算系统。这种定型和试错是任何一次真正的架构创新都绕不开的路。5. Fermi 的遗产它如何悄悄决定了后来十年 GPU 计算的模样5.1 Kepler、Maxwell 乃至后继者都躲不开的“Fermi 坐标系”2012 年的 Kepler 在 Fermi 基础上大幅提升能效比每个 SM 能容纳的 CUDA 核心数涨到 192 个2014 年的 Maxwell 进一步优化每瓦性能还把 L2 缓存做得更大。表面上看这些架构比 Fermi 先进很多但底层的那套“坐标系”——线程束调度模型、块与网格的编程层次、共享内存与 L1 统一配置、统一 L2 缓存、流并发、ECC 选项——都是从 Fermi 定下来的。你至今在 CUDA 里写的grid, block本质上还是在对着 Fermi 设计的硬件抽象做映射。一个网格由若干线程块组成线程块被调度到某个 SM 上执行块内的线程通过共享内存和同步协作线程束是硬件调度的基本单位。这套模型在今天仍然适用就是最好的遗产证明。我甚至觉得Fermi 最重要的遗产不是某个具体硬件参数而是让“GPU 计算架构”这个概念本身变得清晰了调度器、内存层次、并发粒度、可靠性机制这些设计维度第一次系统地出现在 GPU 上。后来的每一代架构创新都是在 Fermi 搭建好的骨架上做肌肉强化。5.2 从实验室玩具到正式产品管线计算生态的成型Fermi 出现之后GPU 计算才真正从论文里的“加速实验”走向了工业软件的产品化。大型流体力学、结构分析、分子模拟、金融风险计算等软件开始认真提供 CUDA 后端不再只是用 GPU 跑几个样例。这背后离不开配套数学库的成熟线性代数库、快速傅里叶变换库等都被重新优化针对 Fermi 的内存层次做了深度调优。后来的深度学习浪潮恰恰是建立在这套已经成型的 GPU 计算基础设施上的。AI 训练需要的高带宽、大容量存储、可靠的长期运行、流并发能力以及完善的调试和性能分析工具都在 Fermi 时代被打通了脉络。没有这套地基后来“用 GPU 训练大模型”就不会来得那么顺。5.3 对今天选型者和开发者的几句实在话聊到最后我想给今天的开发者和选型者一些实际的建议。第一挑选计算设备时不要只看核心数量和单精度峰值。缓存结构、共享内存配额、双精度吞吐、ECC 支持、多流并发能力这些参数对真实负载的影响往往比核心数更大。第二如果你的工作负载需要长时间大规模计算优先考虑专业计算卡而不是游戏卡。游戏卡削减双精度、屏蔽部分可靠性功能是行业常态别用“参数错觉”给自己埋坑。第三写内核时优先思考数据局部性再考虑并行度。Fermi 用一颗 768KB L2 扭转了 GPU 的访存体验这个“缓存优先”的思路在今天依然是性能优化的第一抓手。第四要善用每代架构的特性而不是写一套代码想跑遍所有卡。我踩过不少坑之后才意识到真正高效的做法是翻开陌生架构的白皮书时先看共享内存、L2 和 warp 调度再看浮点峰值。我自己后来的习惯是每一次拿到新卡都先写一个简单的访存带宽测试摸清 L2 的脾气然后再把老内核逐步搬上去。这个习惯就是当年对着 Fermi 反复调优时养成的。回过头看那种“架构完整”带来的踏实感到今天都很少有其他硬件能给我同样的体验。
