导言：一个被低估的转折点

2012 年秋天，多伦多大学一间学生宿舍里，两块价值几百美元的 GeForce GTX 580 游戏显卡连续运转了好几天。运行在它们上面的程序叫 AlexNet，是一个有 6000 万参数的卷积神经网络。几周后，它在 ImageNet 图像识别竞赛上把第二名的错误率甩开了将近 11 个百分点——15.3% 对 26.2%——这个差距大到让整个学术界意识到，深度学习不是噱头，而是真的可行30。

这件事后来被反复讲述，因为它常被当作「现代 AI 的起点」。但很少有人追问一个更基础的问题：为什么训练它用的是显卡，而不是当时数据中心里成千上万的 CPU？答案藏在一个缩写里——CUDA。如果没有 NVIDIA 在 2006 到 2007 年间押下的那一注，AlexNet 的作者们就得用图形 API 把神经网络「伪装」成画三角形的任务，那条路几乎走不通。

CUDA 全称 Compute Unified Device Architecture（统一计算设备架构）。它表面上是一套让程序员用类 C 语言给显卡写程序的工具，本质上是一个赌注：赌「图形处理器」这种为了画游戏画面而生的芯片，可以被改造成通用的、可编程的大规模并行计算引擎；并且赌世界上有足够多的计算任务，天然适合这种「宽而不快」的算力形态。

这个赌注花了将近二十年才完全兑现。今天，从训练 GPT 系列大模型，到天气预报、分子动力学、自动驾驶、电影特效，背后跑的几乎都是 CUDA。NVIDIA 也因此从一家显卡公司变成了全球市值最高的公司之一。

这份材料想说清楚的，不是「CUDA 有多厉害」，而是更本质的几件事：它从哪来、每一代变了什么、它能做什么不能做什么、它背后的根本假设是什么、怎么给它写程序、以及硬件到底是怎么把这套编程模型变成现实的。

第一部分 · 广度全景

一、史前史：在 CUDA 之前，人们怎么「骗」显卡干活

要理解 CUDA 为什么是个转折点，得先看清它出现之前的世界有多别扭。

1.1 GPU 最初只会做一件事：把三角形画到屏幕上

上世纪九十年代末到本世纪初，显卡（GPU，Graphics Processing Unit，这个词由 NVIDIA 在 1999 年随 GeForce 256 推广开来）是一种高度专用的芯片。它的全部使命是图形渲染：接收一堆三维空间里的三角形顶点，经过坐标变换、光照计算、光栅化，最后给屏幕上每个像素算出颜色。

早期 GPU 的内部是一条固定的「流水线」（fixed-function pipeline），各个阶段写死在硬件里，程序员只能调参数，不能改逻辑。转折出现在 GPU 引入「可编程着色器」（programmable shader）之后：NVIDIA 在 2001 年的 GeForce 3 上首次让程序员能写「顶点着色器」，在 2003 年的 GeForce FX 一代支持了浮点运算5。着色器本质上是一小段程序，对每个顶点或每个像素并行地跑一遍。

关键洞察就藏在这里：一段程序，同时作用在成千上万个数据元素上——这不正是科学计算里最常见的形态吗？矩阵乘法、向量运算、物理模拟，本质上都是「对一大堆数据做相同的操作」。

1.2 GPGPU：把计算问题「翻译」成画图问题

2003 年成了一个标志性年份。两个研究组各自独立发现，可以用 GPU 解线性代数问题，而且比当时的 CPU 还快5。这就是 GPGPU（General-Purpose computing on GPU，图形处理器上的通用计算）的开端。同年 SIGGRAPH 上，Mark Harris、David Luebke、Ian Buck 等人做了影响深远的 GPGPU 课程报告，把这件事推向主流视野5。

但「翻译」这个词道出了当时的痛苦。GPU 那时只认图形 API（OpenGL、Direct3D），程序员要做通用计算，必须把问题硬塞进图形的概念框架里：

  • 把输入数据伪装成「纹理」（texture，本来是贴在物体表面的图片）；
  • 把计算逻辑写成「片元着色器」（fragment shader，本来是给像素上色的）；
  • 触发计算的办法是「画一个铺满屏幕的矩形」，逼着 GPU 对每个像素跑一遍着色器；
  • 计算结果是一张「渲染出来的图」，再想办法读回来。

这套做法被戏称为「用着色器骗 GPU」。它有几个致命限制：着色器之间不能随意读写内存（只能各自管各自的像素）、没有真正的整数运算、调试极其困难、而且每换一个算法都要重新折腾一遍图形管线5。一句话：能用，但反人类。

1.3 Brook：把「流」变成一等公民

破局的人是 Ian Buck，斯坦福大学的一名博士生。他的博士课题就是「流式计算」（stream computing），成果是一门叫 Brook 的编程语言3。

Brook 的核心思想是引入一个叫「流」（stream）的抽象：流就是一组需要做相同处理的数据元素的集合；而「核函数」（kernel）就是作用在流上的运算。程序员只管描述「对这些数据做什么」，至于底层是怎么映射到 GPU 的纹理、着色器上的，由编译器去操心3。

这个设计的精神后来几乎原封不动地进了 CUDA：让 C 程序员看到熟悉的概念，把图形管线的脏细节藏起来。Brook 的论文标题就叫《Brook for GPUs: Stream Computing on Graphics Hardware》3。

2004 年，NVIDIA 把 Buck 招了进去，让他和公司里负责 GPU 计算架构的 John Nickolls 搭档。两人一起把 Brook 演化成 CUDA，融合了 Brook 的流式思想和 NVIDIA 自家图形着色语言 Cg 的优点，再吸收 Brook 用户的反馈4。这是一段值得记住的细节：CUDA 不是凭空设计出来的，它是学术界的探索（Brook）被工业界收编、再和硬件团队深度协同的产物。

证据等级说明：Buck 2004 年入职 NVIDIA、与 Nickolls 合作演化 CUDA 这一线索，来自 NVIDIA 官方技术博客与 ACM Communications 的回顾文章，属 A/B 级，可信度较高24。

二、2006-2007：G80 与 CUDA 1.0，为什么是革命

2.1 硬件先行：G80 的「统一着色器」

2006 年 11 月，NVIDIA 发布了代号 G80 的芯片，搭载在 GeForce 8800 GTX 显卡上，对应的微架构代号叫 Tesla（注意：这里的 Tesla 是架构代号，和后来的数据中心产品品牌 Tesla、以及汽车公司都没关系）6。

G80 做了一件在图形史上也算革命的事：统一着色器架构（Unified Shader Architecture）。在它之前，GPU 内部的顶点处理单元和像素处理单元是分开的、写死的两套硬件。如果一个场景顶点多像素少，像素单元就闲着；反之亦然。G80 把它们合并成一池子通用的浮点处理器（NVIDIA 叫它们 stream processor，流处理器），按需动态分配去做顶点活儿还是像素活儿，实现了负载均衡6。

这个改动对图形是优化，对通用计算却是质变。一旦芯片里有了一大堆「什么都能算」的通用浮点单元，它就不再只是图形加速器，而成了一个可编程的并行计算平台67。G80 内部有 128 个这样的核心，组织成「流多处理器」（Streaming Multiprocessor，简称 SM）这个直到今天都还是 NVIDIA GPU 心脏的结构单元8。

NVIDIA 的架构师 Jonah Alben 后来回忆：「我们基本上把 NV30/NV40 的整个着色器架构扔掉了，从头做了一个全新的通用处理器架构，也就是 SIMT。」8 他还说，G80 引入的基本结构「直到今天（指 Pascal 时代）都还在」8。这句话点出了一个重要事实：CUDA 的硬件地基，在 2006 年就基本定型了。

2.2 软件随后：CUDA 1.0

硬件有了通用计算能力，还需要一套不用「骗 GPU」就能编程的软件。这就是 CUDA。NVIDIA 在 2007 年发布了第一版 CUDA 工具包，提供一套基于 C 语言（C99 超集）的编程模型，让程序员第一次能直接给 GPU 写通用程序，而不必碰 OpenGL 的纹理和着色器1。

关于发布日期的一处分歧：CUDA 1.0 的确切公开日期存在不同说法。英文维基百科记作 2007 年 2 月 16 日，而大量二手资料（包括一些架构综述）记作 2007 年 6 月1。还有资料把「2006 年 11 月」当作 CUDA 发布时间，但那其实是 G80 硬件的发布月份5。比较稳妥的表述是：硬件（G80/Tesla）在 2006 年底问世，CUDA 软件平台在 2007 年随之推出。这里区分硬件和软件两个时间点很重要。

为什么说这是革命？因为它一次性解决了 GPGPU 时代所有的别扭：程序员可以用熟悉的 C 写 kernel，可以让线程自由地读写显存任意位置（gather/scatter），可以用真正的整数和指针，可以在芯片上开辟一块由程序员手动管理的高速「共享内存」让线程协作。计算不再需要伪装成画图1。

三、架构演进编年史：从 Tesla 到 Blackwell

从 2006 年至今，NVIDIA 大约每两年推出一代新架构，每一代都用一位科学家的名字命名，并对应一个或几个 CUDA 版本、一个「计算能力」（compute capability，简称 cc）编号。计算能力是一个像 7.0、8.6 这样的版本号，用来标识某块 GPU 支持哪些硬件特性，是 CUDA 编程里非常核心的概念1。

下面这张表是理解 CUDA 历史的骨架（年份指架构首发大致年份）：

架构代号计算能力大致年份标志性新能力
Tesla1.0–1.32006–2009统一着色器、SIMT、共享内存、首次支持双精度（1.3）
Fermi2.0–2.12010–2011真正的 L1/L2 缓存层级、ECC、完整 IEEE-754 双精度、C++ 支持
Kepler3.0–3.72012–2015动态并行（3.5）、Hyper-Q、统一内存雏形、能效大幅提升
Maxwell5.0–5.32014–2016能效优先重设计、专用共享内存（与 L1 分离）
Pascal6.0–6.22016–2017真·统一内存（缺页迁移）、NVLink、HBM2、FP16
Volta7.0–7.22017–2018首代 Tensor Core、独立线程调度、HBM2、NVLink 2
Turing7.52018–2020RT Core（光线追踪）、第二代 Tensor Core、INT8/INT4
Ampere8.0–8.72020–2021第三代 Tensor Core、异步拷贝、结构化稀疏、TF32、MIG
Ada Lovelace8.92022–2023第四代 Tensor Core、FP8、第三代 RT Core
Hopper9.02022–2023线程块簇、分布式共享内存、TMA、Transformer Engine、FP8
Blackwell10.0–10.32024+第二代 Transformer Engine、FP4、双 die 封装、NVLink 5

（来源汇总自 CUDA 维基百科版本表及各代架构白皮书191213。）

下面逐代说清楚每一代真正改变了什么。

Tesla（2006–2009）：奠基

第一代不仅奠定了 SM、SIMT、共享内存这些至今未变的概念，还在末期的 GT200 芯片（计算能力 1.3）上加入了双精度浮点单元，让 GPU 第一次能严肃地做科学计算6。早期 Tesla 没有缓存、没有 ECC（纠错内存），双精度性能也很弱——它更像一个原型，证明了路线可行。

Fermi（2010）：从原型到生产工具

Fermi 是第一代「认真对待计算」的架构。它引入了真正的两级缓存层级（每个 SM 一个 L1、全芯片共享一个 L2），支持 ECC 纠错内存（数据中心刚需），把双精度性能提上来，完整实现 IEEE-754-2008 浮点标准（含融合乘加 FMA），并第一次真正支持 C++ 特性（虚函数、new/delete 等）8。如果说 Tesla 证明了路能走，Fermi 让这条路能跑生产负载。

Kepler（2012）：动态并行与能效

Kepler 把 SM 重新设计为「SMX」，每个 SMX 塞进 192 个 CUDA 核心，并把指令调度从硬件挪到编译器（软件调度），换取更好的能效8。它最重要的编程能力是动态并行（Dynamic Parallelism，计算能力 3.5 引入）：GPU 上正在跑的 kernel 可以自己再启动新的 kernel，不必每次都回到 CPU 发指令1。这对递归、自适应网格、图算法这类「计算量事先不确定」的问题很关键。Hyper-Q 则允许 CPU 多个进程的任务同时塞进 GPU 的多条硬件队列，提高利用率。

Maxwell（2014）：把每一瓦电用到极致

Maxwell 处在台积电 28 纳米工艺停滞期，没法靠制程进步，于是把重心全放在能效上。它把 SM 重新切分（SMM），让共享内存彻底独立于 L1 缓存（之前两者在 Fermi/Kepler 上共用一块物理存储、靠配置切分），简化了调度逻辑8。Maxwell 的意义更多在工程层面：证明了架构精细化能在工艺不变的情况下大幅提升每瓦性能。

Pascal（2016）：统一内存与互联

Pascal 借助台积电 16 纳米 FinFET 工艺大幅提升密度，但它对 CUDA 编程模型最大的贡献是真正的统一内存（Unified Memory）。它带来 49 位虚拟地址空间和一个「页面迁移引擎」（Page Migration Engine），让 CPU 和 GPU 第一次能共享同一个虚拟地址空间，数据在谁需要时按缺页（page fault）自动迁移到谁那里，而不必程序员手动拷贝2122。这条后面会细讲。Pascal 还引入了 NVLink 高速互联（替代 PCIe 做 GPU 间通信）、HBM2 高带宽显存、以及为深度学习准备的半精度 FP16 运算1。

Volta（2017）：两件改变历史的事

Volta 的 V100 是一个分水岭，它干了两件事，每件都深刻影响了之后所有 GPU。

第一件是 Tensor Core。这是一种全新的硬件单元，专门做一件事：小矩阵的乘加（一次完成一个 4×4 矩阵乘法累加）。V100 上每个 SM 有 8 个 Tensor Core，全卡 640 个9。它专为深度学习里铺天盖地的矩阵乘法设计，用混合精度（FP16 输入、FP32 累加）换取数倍乃至十几倍的吞吐。可以说，今天所有大模型的算力，根子上都是 Tensor Core 在扛。

第二件是独立线程调度（Independent Thread Scheduling，ITS）。在 Volta 之前，一个 warp（32 个线程的捆绑执行单元）共享一个程序计数器，遇到分支必须整体串行处理，而且程序员不能假设 warp 内线程会自己「错开」执行——这导致很多隐含依赖的代码有死锁风险。Volta 给每个线程独立的程序计数器和调用栈，让 warp 内线程可以真正独立地推进、暂停、重新汇合9。这让更复杂的细粒度同步成为可能，但也要求程序员改掉一些旧习惯（后面深度部分会讲）。

Turing（2018）：光追与推理

Turing 主要面向游戏和图形，引入了 RT Core（专做光线追踪的硬件）和第二代 Tensor Core，并新增 INT8、INT4 整数运算支持，瞄准的是深度学习推理（inference）场景——推理对精度要求比训练低，用低精度整数能省电提速8。

Ampere（2020）：异步与稀疏

A100 这一代把规模和灵活性都往前推了一大步。它引入第三代 Tensor Core，支持一种叫 TF32 的新格式（用 Tensor Core 跑普通 FP32 训练而几乎不改代码）、支持结构化稀疏（让权重矩阵按 2:4 模式稀疏化，吞吐翻倍）；硬件层面引入异步拷贝（async copy），让数据从显存搬到共享内存的过程可以和计算重叠，不再占用寄存器中转1。它还提供 MIG（Multi-Instance GPU），把一块大 GPU 在硬件层面切成多个互相隔离的小 GPU，方便云上多租户共享。

Hopper（2022）：为 Transformer 而生

H100 几乎是为大语言模型量身定做的。它的关键创新有几项1012：

  • 线程块簇（Thread Block Cluster）：在原有的「网格→线程块→线程」三层结构上方再加一层，让一组线程块可以协同，并直接访问彼此的共享内存（叫分布式共享内存，Distributed Shared Memory），不必绕道全局显存10。
  • TMA（Tensor Memory Accelerator，张量内存加速器）：一个专用硬件单元，负责在显存和共享内存之间异步搬运大块（最多五维）张量数据，把搬数据这件繁琐又关键的事从线程手里接管过去10。
  • Transformer Engine：软硬结合，能在每一层自动在 FP8 和 16 位精度之间智能切换、自动处理缩放，号称把大模型训练提速最多 9 倍、推理最多 30 倍（相对 A100）1131。

Blackwell（2024+）：FP4 与「机柜即一块 GPU」

最新的 Blackwell（B200/GB200）把激进推到新高度1314：

  • 双 die 封装：单个 GPU 由两块芯片通过一条 10 TB/s 的 NV-HBI 接口连成一体，对软件呈现为单一 GPU；
  • 第二代 Transformer Engine 与 FP4：用「微张量缩放」技术做到 4 比特浮点推理，B200 的密集 FP4 算力达 9 PFLOPS14；
  • GB200 NVL72：把 72 块 Blackwell GPU 用 NVLink 5 连成一个域，对外像一块巨型 GPU，NVLink 单卡带宽 1.8 TB/s，是 H100 的两倍14。

可以看到一条清晰的主线：从单个芯片，到机柜级别的一整套系统，被当作单一的计算单元来编程。CUDA 这二十年，是计算单元的「尺度」不断放大的历史。

四、生态：CUDA 真正的护城河

如果 CUDA 只是一门给 GPU 写程序的语言，它不会有今天的地位。让它难以被取代的，是层层叠叠的库和工具生态。绝大多数人用 CUDA，其实并不直接写 kernel，而是调用这些库27。

  • cuBLAS：GPU 上的基础线性代数库（矩阵乘、向量运算），是几乎一切数值计算的地基。
  • cuDNN：深度学习专用原语库（卷积、池化、归一化等），2014 年发布。它的出现意义重大——把 cuDNN 接进 Caffe 这样的框架，标准模型能直接快 36% 还省内存28。今天 PyTorch、TensorFlow 底下跑的都是它。
  • Thrust：类似 C++ STL 的并行算法库（排序、归约、扫描），让人不写 kernel 也能用 GPU。
  • NCCL：多 GPU、多机之间的集合通信库（all-reduce 等），分布式训练大模型离不开它。
  • CUTLASS：开源的 CUDA C++ 模板库，把高性能矩阵乘法（GEMM）的实现拆解开放出来，让人能自定义、能把激活函数等操作和矩阵乘融合在一起，性能逼近 cuBLAS29。
  • CUDA Graphs：把一连串 kernel 的依赖关系预先「录制」成一张图，之后整体一次性提交执行，省掉每个 kernel 单独启动的 CPU 开销（在深度学习里这种开销常达每次 20–200 微秒）27。
  • MPS（Multi-Process Service）：让多个进程的任务在一块 GPU 上真正并发执行，提升利用率32。

这里有一条值得点出的暗线：深度学习浪潮反过来塑造了 CUDA。2012 年 AlexNet 用 GTX 580 训练成功后30，NVIDIA 敏锐地意识到 AI 是 GPU 的巨大新市场，于是从 Volta 开始把 Tensor Core、混合精度、Transformer Engine 这些专为神经网络设计的硬件一代代加进来。是 AI 的需求拉着 CUDA 往「矩阵乘法机器」的方向进化，而 CUDA 提供的算力又反过来让 AI 模型越做越大。这是一个相互加强的循环，也是过去十年最重要的技术故事之一。

五、CUDA 到底能做什么、不能做什么

横向铺开后，有必要冷静界定一下它的能力边界。

CUDA 擅长的，是数据并行（data-parallel）、吞吐量导向（throughput-oriented）的任务——也就是「对海量数据做大体相同的操作，且这些操作之间相对独立」的问题。典型应用包括：

  • 深度学习训练与推理（本质是海量矩阵乘法）；
  • 科学计算与模拟：流体力学、分子动力学、气候模拟、有限元；
  • 线性代数与信号处理：大矩阵分解、FFT；
  • 图像/视频处理、计算机视觉、医学影像重建；
  • 金融领域的蒙特卡洛模拟、风险计算；
  • 密码学、基因测序比对等。

CUDA 不擅长甚至不适合的，是延迟敏感、强串行、分支密集、数据依赖复杂的任务：

  • 需要尽快算完一件小事的低延迟任务（GPU 靠「同时算很多」取胜，而非「一件事算得快」）；
  • 控制流高度发散、每个数据走的路径都不一样的任务（会触发 warp 分支发散，效率暴跌，后面细讲）；
  • 数据量太小，搬到 GPU 的开销超过计算收益的任务；
  • 强顺序依赖、无法并行的算法（比如某些递推）。

理解这条边界，需要进入第二部分——它背后是 GPU 与 CPU 在设计哲学上的根本分歧。

第二部分 · 深度本质

第一部分回答了「是什么、有什么、从哪来」。这一部分回答更难也更本质的问题：为什么 GPU 适合这件事？这套编程模型的底层逻辑是什么？硬件又是怎么把它变成现实的？

六、核心技术假设：GPU 为什么能这样设计

6.1 CPU 与 GPU 的根本分歧：延迟 vs 吞吐

理解 CUDA 的第一把钥匙，是搞清楚 CPU 和 GPU 在优化目标上的根本对立。

CPU 是延迟优化（latency-optimized）的机器。它假设你给它的任务大多是串行的、有先后依赖的，所以它拼命让单个任务尽快完成：高主频、深流水线、复杂的乱序执行、分支预测、以及占据芯片大半面积的多级缓存——缓存的作用是把数据放在离核心很近的地方，让那一个正在执行的线程几乎不用等待15。CPU 是「少数几个非常聪明、跑得飞快的工人」。

GPU 是吞吐优化（throughput-optimized）的机器。它假设你给它的任务有海量个、彼此独立，所以它不在乎单个任务多快完成，只在乎单位时间内能完成的任务总量。它把芯片面积主要花在算术运算单元上，而不是缓存和控制逻辑。GPU 是「成千上万个不那么聪明、单个也不快、但人多势众的工人」15。

这两条路线没有谁对谁错，是面对不同问题的不同最优解。CUDA 的全部设计，都建立在「我面对的是海量独立任务」这个假设之上。

6.2 最关键的一招：用并行隐藏延迟，而不是用缓存消除延迟

这是整个 GPU 计算最核心、也最反直觉的思想，值得反复琢磨。

访问显存（global memory）非常慢，一次访问要等几百个时钟周期。CPU 的对策是用大缓存尽量避免去访问主存。GPU 走了完全相反的路：它不试图消除这个等待，而是用「换人干活」来掩盖它。

具体来说：当一个 warp（一组线程）发出一条访存指令、开始漫长等待时，GPU 的调度器立刻切换到另一个已经准备好的 warp 去执行；那个 warp 也卡住了，就再换下一个。只要 SM 上「待命」的 warp 足够多，总有人能干活，执行单元就一直不空闲，那几百周期的访存延迟就被「藏」在了别人的计算里17。

这就是延迟隐藏（latency hiding）。它要求 GPU 同时持有远超执行单元数量的线程——这叫超额订阅（oversubscription）。CPU 切换线程很贵（要保存/恢复寄存器、刷新流水线），所以养不起太多线程；GPU 则把每个线程的上下文（主要是寄存器）常驻在芯片上，切换几乎零成本，于是可以养成千上万个线程随时待命17。

一句话总结这个本质假设：CPU 靠缓存让一个线程不必等待；GPU 靠海量线程让等待不要紧。 这就是为什么 GPU 需要那么多并行任务才能跑满——如果你只给它几个线程，访存一卡，整块芯片就闲着，性能崩塌。这也直接解释了第五节里 CUDA「不适合数据量太小、强串行任务」的边界从何而来。

6.3 异构计算模型：主机与设备的分工

CUDA 程序天生是「两个世界」的协作，这叫异构计算（heterogeneous computing）：

  • 主机（host）：指 CPU 及其内存。负责程序的主流程、串行逻辑、I/O、以及调度 GPU 干活。
  • 设备（device）：指 GPU 及其显存。负责被「外包」出去的大规模并行计算部分。

典型流程是：CPU 准备数据 → 把数据拷到 GPU 显存 → CPU 启动一个 kernel 让 GPU 并行计算 → 计算完把结果拷回 → CPU 继续。其中那个被丢给 GPU 并行执行的函数，就叫 kernel（核函数）16。

这个模型的本质，是承认「没有一种处理器擅长所有事」，于是让 CPU 干它擅长的串行控制，让 GPU 干它擅长的并行计算，各取所长。CUDA 后来的很多演进（统一内存、动态并行）都是在削减这两个世界之间的摩擦——但摩擦从未完全消失，数据在 host 和 device 间搬运的开销，至今仍是性能优化的关键战场。

七、SIMT：CUDA 执行模型的灵魂

要真正理解 GPU 怎么执行代码，必须搞懂 SIMT，以及它和两个近亲（SIMD、SMT）的区别。这是很多人含混不清、却最能体现 CUDA 本质的地方15。

7.1 三个缩写的本质区别

SIMD（Single Instruction, Multiple Data，单指令多数据）：一条指令同时操作一个数据向量的多个元素，所有「通道」（lane）严格锁步、共用一个程序计数器、一个指令流。CPU 里的 AVX、SSE 指令就是这种。它高效，但僵硬——程序员必须显式地把数据打包成向量，而且所有通道只能走同一条路，遇到「有的元素要这样、有的要那样」就很难处理15。

SMT（Simultaneous Multithreading，同时多线程）：让一个核心同时持有多个独立线程的状态，靠在它们之间快速切换来填满执行单元、掩盖停顿。Intel 的超线程（Hyper-Threading）就是它。它灵活（每个线程完全独立），但要的线程不多，扩展不到成千上万15。

SIMT（Single Instruction, Multiple Threads，单指令多线程）是 NVIDIA 的发明，可以理解为 SIMD 和 SMT 的杂交，取两者之长15：

  • 像 SIMD 一样，硬件把 32 个线程捆成一个 warp，让它们尽量执行同一条指令，从而省下取指、译码的开销（一条指令喂给 32 个线程）；
  • 但又像 SMT 一样，每个线程在程序员看来是独立的——它有自己的寄存器、可以访问不同的内存地址、甚至可以走不同的分支15。

关键在于：SIMD 把「并行」暴露给程序员（你得自己写向量代码），SIMT 把「并行」藏起来（你只管写单个线程的逻辑，硬件负责把 32 个凑成一组跑）。这就是为什么 CUDA 代码看起来像在写一个普通线程的串行程序，而实际上有成千上万个它在同时跑——编程心智负担被大幅降低了，这是 CUDA 易用性的根源。

7.2 warp 分支发散：SIMT 的甜蜜与代价

SIMT 的灵活是有代价的，最典型的就是分支发散（branch divergence）。

一个 warp 的 32 个线程共享取指逻辑，理想情况下它们走同一条路。但如果代码里有 if (条件)，而 warp 内有的线程条件为真、有的为假，怎么办？硬件的做法是：先执行 if 分支，此时只让条件为真的线程「激活」，其余线程被遮蔽（masked off）、空转；再执行 else 分支，反过来。也就是说，两条分支被串行地各跑一遍，每次只有一部分线程真正在干活18。

后果是：如果一个 warp 内线程走的路径五花八门，效率会急剧下降——极端情况下 32 条不同路径，性能可能掉到 1/32。这就是为什么前面说 CUDA「不适合控制流高度发散的任务」。优秀的 CUDA 程序员会刻意安排数据和线程的对应关系，让同一个 warp 里的线程尽量走相同的分支。分支跑完后，线程会重新汇合（reconverge）继续齐步走18。

7.3 Volta 的独立线程调度：一次模型升级

前面历史部分提到 Volta 引入了独立线程调度（ITS）。现在可以讲清楚它到底改了什么。

Volta 之前，一个 warp 只有一个程序计数器，32 个线程在硬件层面无法真正「各走各的」。这带来一个隐患：很多老代码隐含假设了「warp 内线程是锁步的」，比如用这个假设来做线程间通信，省掉显式同步。这种「warp 同步编程」（warp-synchronous programming）在简单情况下能用，但很脆弱，复杂同步下还会死锁9。

Volta 给每个线程独立的程序计数器和调用栈，硬件可以让 warp 内线程真正独立地推进、在任意指令处暂停、再重新汇合9。好处是能写更复杂、更细粒度的线程协作（比如同一 warp 内不同线程持有不同的锁）。代价是：那些依赖「隐式锁步」的老代码可能出错。所以从 CUDA 9 开始，NVIDIA 强制要求 warp 内通信必须用带 _sync 后缀的显式原语（如 __shfl_sync、__ballot_sync），通过一个掩码明确指定哪些线程参与，不能再偷懒依赖隐式同步20。这是一个「模型变严格、但更安全」的典型演进。

八、编程模型：层次结构与内存模型

现在进入「怎么给 CUDA 写程序」。核心是两套层次结构：线程的层次、内存的层次。

8.1 线程的三/四层结构

当你启动一个 kernel，会同时产生海量线程，它们被组织成严格的层次16：

  • 线程（thread）：最小执行单位，每个线程跑一份 kernel 代码，有自己的寄存器和局部变量，通过内置变量知道「我是第几个」从而处理不同的数据。
  • 线程块（block）：若干线程的组（比如 256 个）。同一个 block 内的线程有两项特权：能访问同一块共享内存（shared memory），能通过 __syncthreads() 做块内同步（所有线程在此等齐再继续）。一个 block 会被整体分配到某一个 SM 上执行，不会跨 SM16。
  • 网格（grid）：一次 kernel 启动产生的所有 block 的集合。
  • 线程块簇（cluster）：Hopper 新增的中间层，让若干个 block 组成一簇，能协同并互相访问共享内存（分布式共享内存）10。

而在硬件执行时，block 里的线程又会被自动切成一个个 warp（32 个线程一组）——warp 是调度和执行的真正单位，但它在编程模型里是「半隐藏」的：你写代码时面对的是 thread 和 block，但要写出高性能代码，又必须时刻想着 warp（因为分支发散、访存合并都发生在 warp 粒度）。这种「编程概念」与「硬件现实」之间的微妙错位，是 CUDA 进阶的关键门槛。

这个层次结构为什么这么设计？因为它对应了硬件的物理层次：thread 对应 CUDA 核心、warp 对应 SIMT 执行宽度、block 对应一个 SM、grid 对应整块 GPU。编程模型是硬件结构的镜像。

8.2 内存的多层结构

CUDA 暴露给程序员的内存种类之多，常让初学者困惑。但它们各有用途，理解它们就是理解性能优化的一大半19：

内存类型速度作用域谁管理典型用途
寄存器 register最快单个线程编译器线程的局部变量
共享内存 shared很快一个 block程序员手动block 内线程协作、缓存复用数据
L1/L2 缓存快L1 单 SM / L2 全卡硬件自动自动缓存全局访问
全局内存 global慢（几百周期）所有线程程序员主数据存储，即显存
局部内存 local慢（其实在显存里）单个线程编译器寄存器装不下时溢出到这里
常量内存 constant快（有专用缓存）所有线程只读程序员所有线程都读的小量只读数据
纹理内存 texture快（有专用缓存）所有线程只读程序员有空间局部性的只读数据，自带插值

寄存器最快但数量有限，是稀缺资源（这点后面讲 occupancy 时很关键）。共享内存是 CUDA 性能优化的核心武器——它是一块由程序员手动管理的、SM 上的高速便签纸（scratchpad），速度接近寄存器。注意「局部内存」这个名字有误导性：它名为 local，物理上却在慢速显存里，是寄存器不够用时的溢出区，应尽量避免19。

8.3 两个性能命门：访存合并与 bank 冲突

有了内存层次，就有两个最常见、也最致命的性能陷阱。

访存合并（memory coalescing）针对全局内存。当一个 warp 的 32 个线程同时访问全局内存时，如果它们访问的是连续的、对齐的地址（线程 0 访问第 0 个元素、线程 1 访问第 1 个……），硬件能把这些访问合并成一两次大的内存事务，效率最高。反之，如果访问地址东一个西一个（比如跨步、随机），硬件就得发起很多次小事务，有效带宽暴跌19。所以「让相邻线程访问相邻数据」是 CUDA 第一性能准则。这也解释了为什么矩阵运算里行优先/列优先、数据布局如此讲究。

bank 冲突（bank conflict）针对共享内存。共享内存被硬件切成 32 个「bank」（存储体），目的是让 warp 的 32 个线程能同时各访问一个 bank、并行无阻。但如果多个线程同时访问落在同一个 bank 的不同地址，就发生 bank 冲突，硬件只能把这些访问串行化，速度按冲突的线程数成倍下降23。常见的解决办法是给数组加一列「填充」（padding）改变地址映射，或者用「swizzling」（地址重排）技巧错开访问23。

这两个陷阱有一个共同的本质：GPU 的内存系统是为「整个 warp 一起、规整地访问」而设计的，任何打破规整性的访问模式都会被惩罚。 这是数据并行硬件的内在逻辑。

8.4 线程协作：同步、原子、warp 原语

线程之间要协作，CUDA 提供了几层工具：

  • __syncthreads()：block 内屏障同步，所有线程到此等齐。常用于「先一起把数据搬进共享内存，等齐，再一起用」这种模式20。
  • 原子操作（atomic）：保证「读-改-写」不被打断，多个线程安全地更新同一个变量（如全局计数、直方图累加）。代价是会串行化对该地址的访问。
  • warp 级原语：在一个 warp 的 32 个线程之间直接交换数据，不经过共享内存，极快。最重要的是 shuffle（__shfl_sync，让线程直接读另一个线程的寄存器）和 vote（__ballot_sync/__any/__all，对 32 个线程的布尔值做快速归约，比如「有没有任何线程满足条件」）20。这些是写高性能归约、扫描算法的利器。
  • Cooperative Groups（协作组，CUDA 9 引入）：一套更灵活、更安全的同步抽象，让程序员能显式地定义任意粒度的线程组（一个 warp、一个 block、甚至整个 grid）并对其同步，取代以前各种隐式、脆弱的做法20。

九、性能优化的思维框架

写出能跑的 CUDA 代码不难，写出快的 CUDA 代码是门手艺。其底层是几个相互关联的概念。

9.1 占用率：不是越高越好

占用率（occupancy）指一个 SM 上「实际活跃的 warp 数」与「硬件支持的最大 warp 数」之比17。它直接关系到延迟隐藏：活跃 warp 越多，访存停顿时越有人能顶上，越能藏住延迟。

但占用率受资源约束：每个 SM 的寄存器和共享内存是固定的，被所有常驻线程瓜分。如果你的 kernel 每个线程用了很多寄存器，或每个 block 用了很多共享内存，能同时容纳的 warp 就少，占用率就低17。

一个常见误区是「占用率越高越好」。事实是：占用率只要高到足以隐藏延迟就够了，再往上提升收益递减，甚至因为每线程资源被压缩而变慢17。Vasily Volkov 的经典研究甚至展示了低占用率下靠指令级并行也能跑满性能的反例17。所以占用率是手段不是目的，目的永远是让执行单元别闲着。

9.2 Roofline：判断你被什么卡住了

Roofline 模型是判断「该往哪优化」的思维工具。它把一个 kernel 的性能上限画成一条折线：横轴是计算强度（每搬运一字节数据，做了多少次浮点运算），纵轴是可达性能。折线由两段构成：一段斜线（被内存带宽限制，memory-bound）、一段水平线（被算力峰值限制，compute-bound）17。

它的实践价值在于：先搞清楚你的 kernel 是被带宽卡住还是被算力卡住，再对症下药。如果是 memory-bound，优化访存（合并访问、用共享内存复用数据、减少数据搬运）才有用，拼命优化计算指令是白费;如果是 compute-bound，反过来。深度学习里的大矩阵乘法通常 compute-bound（这正是 Tensor Core 的用武之地），而很多逐元素操作、归一化则是 memory-bound。

9.3 流与并发：让 GPU 别等 CPU

流（stream）是 CUDA 里表达并发的机制。一个 stream 是一串按顺序执行的操作（拷贝、kernel、再拷贝……）；不同 stream 之间则可以并发22。最经典的用法是把数据分块，用多个 stream 流水线化：当 stream A 在算第一块时，stream B 同时在把第二块数据拷进来，让「拷贝」和「计算」重叠，藏住数据传输的开销。

更进一步，Kepler 的 Hyper-Q、以及 MPS 服务，让多个 CPU 进程的任务也能并发塞进同一块 GPU32；CUDA Graphs 则把整串操作打包成图、一次提交，省掉反复启动的 CPU 开销27。这些都指向同一个目标：让昂贵的 GPU 时刻有活干，别卡在等 CPU 发指令上。

9.4 常见陷阱清单

把前面散落的坑集中起来，方便对照：

  • 访存不合并：相邻线程访问不相邻数据 → 带宽浪费。
  • 共享内存 bank 冲突 → 访问被串行化。
  • 分支发散：同一 warp 内线程走不同路 → 串行执行各分支。
  • 寄存器溢出：单线程用太多寄存器 → 溢出到慢速 local memory，且压低占用率。
  • 数据搬运过度：在 host 和 device 间反复拷贝小数据 → 开销盖过收益。
  • 占用率过低：资源用太满导致活跃 warp 太少 → 藏不住延迟。
  • 依赖隐式 warp 同步：Volta 之后不带 _sync 的旧式 warp 通信 → 可能出错。
  • 用 GPU 跑串行/小规模任务 → 杀鸡用牛刀，反而更慢。

十、底层实现：硬件如何兑现这套模型

最后一层，把编程模型和硬件结构对应起来。这是「本质」的最深一层。

10.1 SM 的内部结构

流多处理器（SM）是 GPU 的基本计算积木，一块 GPU 由几十到上百个 SM 组成。以 Volta V100 的 SM 为例，它内部分成 4 个处理块（processing block），每块包含9：

  • 16 个 FP32（单精度）核心、16 个 INT32（整数）核心、8 个 FP64（双精度）核心；
  • 2 个 Tensor Core；
  • 一个 warp 调度器（warp scheduler）和一个分发单元（dispatch unit）；
  • 一块 L0 指令缓存；
  • 一个 64 KB 的寄存器文件（register file）。

这个结构直接对应编程模型：寄存器文件存放所有常驻线程的寄存器（这就是为什么寄存器是稀缺资源、用多了就限制占用率）；warp 调度器就是 SIMT 的硬件实现核心——它每个周期挑选一个「准备好」的 warp，把它的下一条指令分发给执行单元9。

10.2 warp 调度器：延迟隐藏的发动机

warp 调度器是把第六节那个抽象思想（用并行隐藏延迟）落到硅片上的关键部件。它维护着一批常驻 warp 的状态，每个时钟周期做一件事：从所有「就绪」（操作数齐备、没在等访存）的 warp 里挑一个，发射它的指令9。

当一个 warp 因访存而停顿，它就从「就绪」名单里暂时移除，调度器毫不停留地去发射别的就绪 warp。因为切换只是换一个指针指向不同的寄存器区域（上下文本来就常驻片上），所以零开销17。这就是为什么 GPU 能用「人海战术」藏住延迟——整个机制就压缩在 warp 调度器每周期一次的选择里。一个 SM 上能同时常驻几十个 warp（如最多 64 个），它们轮番上阵，让算术单元几乎不空闲。

从 Tesla 到现在，SM 的演进主线就是：调度器从 1 个增到多个、core 数量和种类不断丰富、加入 Tensor Core / RT Core 等专用单元，但「warp 调度器轮换就绪 warp 来隐藏延迟」这个内核，二十年没变8。

10.3 内存层级的硬件实现及其演变

硬件内存层级和编程模型里的内存种类大致对应，但有一段有趣的演变史8：

  • 寄存器文件：每个 SM 一块，被该 SM 上所有线程瓜分。
  • 共享内存与 L1 的纠缠：早期 Tesla 没有 L1。Fermi 引入 L1，并且让 L1 和共享内存共用同一块物理 SRAM，靠配置切分比例（比如 48KB 共享 + 16KB L1，或反过来）。Maxwell 又把两者分开成独立物理块。到了 Volta/Turing，再次把它们合并成一块统一的、可配置的存储（NVIDIA 称之为「combined L1/shared memory」），让它能根据负载灵活当缓存或当便签纸用9。这段「分了又合、合了又分」的历史，反映了在「程序员手动管理（共享内存）」与「硬件自动管理（缓存）」之间反复权衡。
  • L2 缓存：全芯片共享，是所有 SM 访问显存的统一中转。它的容量一路猛增，到 Hopper 已达 50 MB 级别（早期 Fermi 仅几百 KB）。
  • 显存（DRAM）：从早期 GDDR，到 Pascal 引入 HBM2（高带宽内存，3D 堆叠、超宽位宽），再到 HBM3。带宽是 GPU 的命脉，HBM 把它推到每秒数 TB10。

10.4 编译工具链：PTX、SASS 与前向兼容的精妙设计

CUDA 怎么从源码变成 GPU 上跑的机器码？这套工具链里藏着一个非常聪明的设计，专门解决「今天编译的程序，怎么在未来还没造出来的 GPU 上跑」的问题2425。

编译分两个阶段：

  1. nvcc（NVIDIA 的编译器驱动）先把你的 CUDA C++ 设备代码编译成 PTX。
  2. 再由 ptxas 工具把 PTX 编译成 SASS，也就是某个具体 GPU 架构的真实机器码（这一两阶段流程及 -arch/-code 等编译选项的细节可参见相关整理26）。

关键在 PTX 和 SASS 的分工24：

  • PTX（Parallel Thread Execution，并行线程执行）是一种虚拟指令集，是一台「抽象 GPU」的汇编语言。它不对应任何真实硬件，而是表达所有 NVIDIA GPU 的共同特征。PTX 从 CUDA 诞生第一天起就存在，且保持相对稳定、向前兼容24。
  • SASS（Streaming ASSembler）是真实硬件的原生机器码，每一代架构都不一样，且不对外承诺稳定24。

这套两层结构带来一个关键能力：前向兼容（forward compatibility）。一个编译好的 CUDA 程序，可以把 PTX 和针对若干已知架构的 SASS 一起打包进一个文件，叫 fatbin（胖二进制）25。运行时：

  • 如果 fatbin 里有匹配当前 GPU 的现成 SASS，直接用；
  • 如果没有（比如这块 GPU 是程序编译之后才出的新架构），NVIDIA 驱动会临时充当 JIT 编译器（即时编译），把打包进去的 PTX 现场编译成这块新 GPU 的 SASS25。

正因如此，2015 年编译的、内嵌了 PTX 的老程序，今天放到 Blackwell 上也能跑——驱动会把 PTX 即时翻译成 Blackwell 的机器码。这就是 PTX 作为「虚拟 ISA」的精妙之处：它把「程序」和「具体硬件」解耦，让 NVIDIA 能放手大改底层 SASS、每代 GPU 重新设计硬件指令集，而不破坏既有软件生态。这种「稳定的虚拟层 + 自由演进的物理层」思路，和 Java 字节码、LLVM IR 异曲同工，是 CUDA 生态能持续二十年向前兼容的制度性保障24。

10.5 上下文、调度与并发的硬件支撑

最后补上执行层面的几块拼图：

  • 上下文（context）：每个使用 GPU 的进程有自己的 CUDA 上下文，封装了它的地址空间、内存分配、流等状态，相当于该进程在 GPU 上的「会话」。
  • 块到 SM 的调度：当 kernel 启动，GPU 的全局调度器（GigaThread engine）把一个个 block 分派到有空闲资源的 SM 上。一个 block 一旦上了某个 SM 就待到执行完，期间它的 warp 由该 SM 的 warp 调度器轮换。block 数量通常远多于 SM 能同时容纳的，于是它们排队、滚动地被处理——这天然实现了「同一份代码自动适配不同规模 GPU」（小卡 SM 少就慢慢跑，大卡 SM 多就并行更多 block）16。
  • 并发与抢占：早期 GPU 一次只能跑一个 kernel，且不能被打断。Kepler 的 Hyper-Q 让多队列并发；Pascal 引入指令级抢占（preemption），让长时间运行的 kernel 能被中断，改善了多任务和调试体验。这些硬件机制，是上层流、MPS、MIG 等并发能力的根基。

结语处的几个开放问题

把 CUDA 二十年的脉络摊开后，有几件事仍悬而未决，留给读者继续琢磨。

CUDA 的护城河到底有多深？它的统治力一半来自硬件，一半来自生态——cuDNN、cuBLAS 这些库背后是无数人年的优化，以及 PTX 这套向前兼容机制锁定的软件资产。竞争者（AMD 的 ROCm、各种编译器中间层、以及 PyTorch 等框架试图抽象掉底层）都在试图绕过它，但「换一套生态」的迁移成本至今高得惊人。这层护城河是会被时间侵蚀，还是会随 AI 需求继续加固？

专用化会走多远？从通用浮点核心（G80），到 Tensor Core（Volta），到 Transformer Engine 和 FP4（Hopper/Blackwell），硬件越来越为「矩阵乘法 + 神经网络」这一种负载量身定做。这提升了 AI 性能，但也让 GPU 离「通用计算」的初心越来越远。当一块芯片大半面积都在为 Transformer 服务时，它还算「通用」吗？如果下一代主流 AI 架构不再是 Transformer，这些硬件会不会成为包袱？

编程模型的复杂度谁来消化？从最初「写个 kernel」的简洁，到如今要同时操心 warp、bank、occupancy、stream、cluster、TMA、混合精度——CUDA 的心智负担在持续增加。大量复杂度被 CUTLASS、Triton、编译器和上层框架吸收掉了，普通从业者越来越不需要直接碰这些。那么未来「会写 CUDA」还会是一种稀缺技能，还是会像「会写汇编」一样退回到极少数底层专家手里？

这些问题没有标准答案，但理解了前面这十节，至少能在它们面前站得住脚——知道每个判断背后的硬件约束和历史脉络是什么。这，大概就是真正「懂 CUDA」与「会用 CUDA」之间的区别。

参考文献

  1. CUDA — Wikipedia（版本史、计算能力与架构映射表、各计算能力特性）。https://en.wikipedia.org/wiki/CUDA 访问日期 2026-06。

  2. The Origins of GPU Computing — Communications of the ACM（GPU 计算起源、Buck 与 Nickolls 协作历史）。https://cacm.acm.org/federal-funding-of-academic-research/the-origins-of-gpu-computing/

  3. Ian Buck, Tim Foley et al., “Brook for GPUs: Stream Computing on Graphics Hardware”（斯坦福，CUDA 前身 Brook 论文）。https://graphics.stanford.edu/papers/brookgpu/brookgpu.pdf

  4. Inside the Programming Evolution of GPU Computing — NVIDIA Technical Blog / The Next Platform（Brook 到 CUDA 的演化，融合 Cg）。https://developer.nvidia.com/blog/inside-the-programming-evolution-of-gpu-computing/

  5. General-purpose computing on graphics processing units — Wikipedia（GPGPU 史前史、2001/2003 关键节点、着色器通用计算）。https://en.wikipedia.org/wiki/General-purpose_computing_on_graphics_processing_units

  6. Tesla (microarchitecture) — Wikipedia（G80 统一着色器架构、SM、流处理器、CUDA 引入）。https://en.wikipedia.org/wiki/Tesla_(microarchitecture)

  7. E. Lindholm et al., “NVIDIA Tesla: A Unified Graphics and Computing Architecture”, IEEE Micro 2008（Tesla 架构一手论文，SIMT 与统一架构原理）。

  8. Fabien Sanglard, “A history of NVidia Stream Multiprocessor”（SM 逐代演进、core 数、warp 调度器、Alben 引述）。https://fabiensanglard.net/cuda/

  9. NVIDIA Tesla V100 GPU Architecture Whitepaper, WP-08608-001, 2017（Volta SM 结构、Tensor Core、独立线程调度、L1/共享内存合并）。https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf

  10. NVIDIA Hopper Architecture In-Depth — NVIDIA Technical Blog（线程块簇、分布式共享内存、TMA、HBM3、L2）。https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/

  11. NVIDIA H100 Tensor Core GPU Architecture Whitepaper, GTC 2022（Hopper 一手架构白皮书）。

  12. Hopper (microarchitecture) — Wikipedia（线程块簇、TMA、分布式共享内存定义）。https://en.wikipedia.org/wiki/Hopper_(microarchitecture)

  13. Blackwell (microarchitecture) — Wikipedia（双 die、NV-HBI、208 亿晶体管、4NP 工艺）。https://en.wikipedia.org/wiki/Blackwell_(microarchitecture)

  14. NVIDIA Blackwell Architecture / GB200 NVL72 — NVIDIA（第二代 Transformer Engine、FP4、微张量缩放、NVLink 5、NVL72）。https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/

  15. “SIMD < SIMT < SMT: parallelism in NVIDIA GPUs” — yosefk.com（SIMT 与 SIMD/SMT 本质区别、CPU/GPU 设计哲学）。https://www.yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html

  16. CUDA Programming Guide — Programming Model — NVIDIA（kernel、grid/block/thread 层次、block 到 SM 调度、异构模型）。https://docs.nvidia.com/cuda/cuda-programming-guide/

  17. Vasily Volkov, “Understanding Latency Hiding on GPUs”, UC Berkeley EECS-2016-143（延迟隐藏、占用率、超额订阅的深入分析）。https://www2.eecs.berkeley.edu/Pubs/TechRpts/2016/EECS-2016-143.pdf

  18. “Benchmarking the cost of thread divergence in CUDA” — arXiv:1504.01650（warp 分支发散的代价与机制）。

  19. Using Shared Memory in CUDA C/C++ — NVIDIA Technical Blog（共享内存、访存合并、各类内存用途）。https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/

  20. Using CUDA Warp-Level Primitives — NVIDIA Technical Blog（shuffle/vote/ballot、_sync 原语、Cooperative Groups、__syncthreads）。https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/

  21. Unified Memory for CUDA Beginners — NVIDIA Technical Blog（统一内存编程模型）。https://developer.nvidia.com/blog/unified-memory-cuda-beginners/

  22. Beyond GPU Memory Limits with Unified Memory on Pascal — NVIDIA Technical Blog（Pascal 49 位虚拟地址、页面迁移引擎、缺页迁移、流与预取）。https://devblogs.nvidia.com/beyond-gpu-memory-limits-unified-memory-pascal/

  23. CUDA Shared Memory Bank / Swizzling — Lei Mao’s Log Book（bank 冲突机制、padding 与 swizzling 解法）。https://leimao.github.io/blog/CUDA-Shared-Memory-Bank/

  24. Understanding PTX, the Assembly Language of CUDA — NVIDIA Technical Blog（PTX 作为虚拟 ISA、SASS、JIT、前向兼容）。https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/

  25. CUDA Compiler Driver NVCC — NVIDIA Documentation（两阶段编译、cubin、fatbin、二进制兼容规则）。https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/

  26. CUDA Compilation — Lei Mao’s Log Book（nvcc/ptxas 流程、-arch/-code/-gencode）。https://leimao.github.io/blog/CUDA-Compilation/

  27. CUDA Refresher: The GPU Computing Ecosystem — NVIDIA Technical Blog（cuBLAS/cuFFT/Thrust/NCCL/CUDA Graphs 生态全景）。https://developer.nvidia.com/blog/cuda-refresher-the-gpu-computing-ecosystem/

  28. Sharan Chetlur et al., “cuDNN: Efficient Primitives for Deep Learning”, arXiv:1410.0759（cuDNN 设计与对 Caffe 的提速数据）。https://arxiv.org/pdf/1410.0759

  29. CUTLASS: Fast Linear Algebra in CUDA C++ — NVIDIA Technical Blog（CUTLASS 模板库、GEMM 分层、算子融合）。https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/

  30. AlexNet — Wikipedia（2012 ImageNet、GTX 580、cuda-convnet、15.3% vs 26.2% 错误率）。https://en.wikipedia.org/wiki/AlexNet

  31. H100 Transformer Engine Supercharges AI Training — NVIDIA Blog（Transformer Engine、FP8 与 16 位自动切换、提速倍数）。https://blogs.nvidia.com/blog/h100-transformer-engine/

  32. Multi-Process Service (MPS) — NVIDIA Documentation（MPS、Hyper-Q、多进程并发执行）。https://docs.nvidia.com/deploy/mps/