MLSys 中文整理

Appearance

原文：第 3 章 说明：忠实翻译原网页内容，并补入与经典文献、业界系统的对照。术语首次出现给出英文锚点。

第 3 章 GPU 架构与机器学习系统 ​

本章在体系中的位置

这一章紧接第 2 章：第 2 章把 CPU 拆到极限——流水线、分支预测、乱序执行、缓存、功耗墙——然后停在"单条线程的延迟哲学到头了"。第 3 章回答那个悬着的问题：如果延迟哲学到头了，换一套哲学能走多远？ GPU 就是那套新哲学的产物。

在 16 周课程的坐标里，本章是第 2 章「CPU 基础」的延续，也是第 4 章 CUDA 编程的架构地基。第 4 章你会反复回头引用这里的 SM、warp、shared memory、occupancy——这些概念先在这章找到物理出处，再在第 4 章变成代码。

在 MLSys / AI Infra 体系的坐标里，这一章给出"硬件坐标轴"：后面讲编译器的分块与融合、讲分布式训练的数据搬移、讲推理系统的算子优化，最后都要落回本章的吞吐哲学、内存层级和精度-架构连接来解释。

第一次听到"GPU 架构"这个词，大多数人脑子里冒出的画面是：一个算术单元更多的 CPU。GPU 无非是更快一点的 CPU。这个图景不算全错——两者都是处理器，都在对数据执行指令。但作为解释，它浅得危险。它漏掉了两个架构分道扬镳的根本原因，也藏住了机器学习系统学生能从硬件学到的最重要的一课：

架构不是凭空设计的。架构是围绕工作负载设计的。

工作负载变了，架构迟早要变；工作负载没变，架构的演化就始终被它的结构卡住。这是本章（也是这场讲座）反复强调的核心主张：工作负载为王（workload is the king）。

GPU 不是"更多的算力"，而是对另一个优化问题给出的不同答案。CPU 的优化目标是相对复杂的串行线程的低延迟执行；GPU 的优化目标是大量可并行执行的简单操作的高吞吐执行。一句话之差，后果巨大。它解释了为什么 CPU 把大片芯片面积花在控制逻辑、推测、缓存和分支预测上，而 GPU 把面积几乎全给了算术资源和"让很多线程同时活着"的机制；它解释了为什么 GPU 先在图形里兴旺；为什么早期 GPGPU 编程那么痛苦；为什么 Brook 和后来的 CUDA 那么重要；为什么深度学习称王之后张量核出现了。它也解释了为什么直到今天，机器学习系统的大部分工作，本质上都是算法、编译器、运行时与硬件之间的协同设计。

本章的目标不是让你背一串历史事实，而是帮你建立一张因果图景。读完你应该不仅知道现代 GPU 里有什么，还知道它为什么有这些组件、这些组件为什么在那个时间点出现、是什么工作负载逼出了它们，以及为什么机器学习恰好和 GPU 那么契合、别的负载却不。

3.1 为什么这一章重要 ​

先做一个小思想实验。给你两个盒子：一个标着 CPU，一个标着 GPU，里面各有一块芯片。现在有一个程序，只有一条极难伺候的、串行依赖极深的计算链，比如一段反复依赖前一步结果的迭代。哪个盒子跑得快？CPU。再换一个程序：对 个互不相关的数，每个都做一遍同样的乘加。哪个盒子跑得快？GPU。两种答案都不神秘——但如果你只会说"GPU 快因为它核多"，你就错过了真正的问题：为什么这两个盒子被造得如此不同？

答案是它们被优化的问题不同。CPU 回答"怎么把一条线程跑得最快"；GPU 回答"怎么让一整套独立操作以最大吞吐流过芯片"。这两个问题是设计一切的起点，也是第 2 章 CPU 故事的终点和第 3 章 GPU 故事的起点。

教学目的。 讲座强调，这一章不是要你记住一串架构部件的名字。它想让你理解一条因果链：工作负载结构 → 架构选择 → 硬件组件 → 编程模型。这条链上每一环都拴着下一环。等你把链看通了，现代 GPU 里那些看似杂乱的部件——着色器、SM、张量核、TMA、精度格式——就都不是孤立的名词，而是同一套工作负载哲学的多次变奏。

关键约束。 记住一个反直觉的事实：现代芯片的硬约束是功耗，不是晶体管数。这个约束在 3.2 节会反复出现，它也是"为什么 GPU 长成那样"的物理底色。

3.2 从单线程思维到吞吐思维 ​

讲座用 CPU 与 GPU 的对照开场。这不是营销分类，而是理解现代加速器设计哲学最短的路径。为了让对照精确，先界定"线程"。本课程说线程时，用的不是操作系统的完整定义，而是一个计算定义：线程是一条打算按顺序执行的指令流。你写 C/C++/Python 时，默认一条指令跟在另一条后面；代码里微小的并行机会不会显式标出，发现它们的责任落在编译器、运行时和处理器硬件身上。

直觉。 现代 CPU 极其精致，恰恰因为它想方设法让一条线程跑得快：取指、译码、分支预测、寄存器重命名、依赖跟踪、乱序执行、深缓存层级压延迟……这些组件大多不直接对数据做算术，它们存在的唯一目的是让一两条复杂线程以最小延迟执行。但就算卷了几十年，从一条线程里能榨出的指令级并行（instruction-level parallelism，ILP）通常只有 4 到 8 条——而拿到这点并行度已经需要软硬件的极大复杂。你写的循环越依赖前一步，能并行的指令越少，CPU 那套精密机械就越派不上用场。

GPU 换了一个观察起点。假设工作负载不是一条必须极快的难线程，而是一大堆可以用同样方式处理、彼此几乎没有依赖的数据元素。这时候再把芯片面积投给"从一条线程里榨取有限 ILP 的精密机械"就没有意义了——自然的做法是简化控制侧、复制计算侧。这就是从 CPU 思维到 GPU 思维的概念一跃。

讲座把这个结构命名为流（stream）。流不只是"数据流向某处"的直觉；流是一种工作负载结构：一大批可以独立（或足够独立）处理的数据元素。元素之间依赖弱，系统就能同时跑很多线程，而不必为协调付代价。于是 GPU 几乎不关心单线程延迟，它只关心总吞吐。这是一个根本的架构选择，不是实现细节。

CPU 视角（线程中心）：优化一条线程的延迟
GPU 视角（流中心）：优化许多同类操作的总吞吐

把两者的芯片面积直觉摆出来：CPU 芯片上直接算数据的算术单元只占一小块——许多 CPU 的算术逻辑单元（ALU）占芯片面积不足 10%，其余全在给控制和访存当后勤。GPU 则反过来：算术单元占大头，控制逻辑薄薄一层。把 CPU "变"成 GPU 的粗劣心法：先拆掉大部分单线程优化逻辑，留下算术部分，把它复制很多份，让多条数据共用同一条指令流（取指与译码只做一次，算术执行做很多遍），再布置存储与调度，让大量操作同时"在空中"。这不是 GPU 历史的晶体管级还原，却是解释"吞吐处理器为什么长得和延迟处理器这么不一样"的好模型。

关键约束。 这里必须点破一件事：现代高性能处理器的硬约束是功耗，不是晶体管数或硅面积。学生总以为瓶颈在面积，其实现代芯片是功耗受限的机器。激进提频的功耗代价极高——动态功耗随频率线性增长、随电压平方增长，而漏电又限制电压下探。这正是 GPU 走向"许多中频并行算术单元"、而不是"少数极高频率的推测核心"的原因。GPU 接受较低的单线程速度，用规模补偿。这个交易只在工作负载允许时才成立——如果负载天生串行，规模补偿就失效。

小结：CPU 用复杂换单线程低延迟；GPU 用简单+复制换总吞吐。选择哪种哲学，由工作负载说了算。

3.3 GPU 三代演化 ​

站在工作负载中心看 GPU 历史，脉络立刻清晰。讲座把演化切成三代。

第一代（约 1990 年代末以前）：固定功能机器。 GPU 主要是一台固定功能（fixed-function）加速器：加速渲染，但流水线的操作大部分由硬件预先定死。干本行很快，却不是通用计算平台。你没法让"画三角形"的硬件去解方程——至少不是以程序员的身份。

第二代（2000 年代初起）：可编程图形处理器。 晶体管预算大到"有限可编程"成为可能，"GPU"一词从"图形加速"变成"可编程图形处理器"。图形流水线的关键阶段可以由软件改写，顶点与像素着色器出现。这时研究者开始问：图形芯片能不能干图形以外的事？

第三代（约 2010 年起）：通用并行计算平台。 非图形用例足够清楚、足够重要，GPU 开始明明白白演化成通用并行计算平台。早期由科学计算、线性代数、数据挖掘、图像处理驱动；后期由深度学习和大型语言模型主导。这里讲座有个关键观点：转换之后，GPU 架构不再只是硬件故事，而变成软硬件协同设计（hardware-software co-design）的故事——被 CUDA 和 CUDA 激活的工作负载深刻塑造。

这个三代框架之所以值得记，是因为它阻止你把今天的 GPU 想象成"一次性出现的"。现代 GPU 是分层的历史物件：图形逻辑、通用计算逻辑、专用张量硬件、多层存储、软件抽象，因为回应不同时期的不同压力而共存。不懂分层，现代 GPU 看起来就是一坨随便堆的模块；懂了，它就是一连串对不断变化的工作负载的应答。

小结：固定功能 → 可编程 → 通用并行计算。每一代都不是升级，而是换了一个工作负载答案。

3.4 图形流水线：GPU 计算的起源 ​

要明白 GPU 为什么先对非图形应用有用，得先懂图形流水线本身。讲座仔细重建了这条流水线，这里也值得用叙述再走一遍。

直觉。 想象一个用网格（mesh）表示的三维物体，里面全是顶点。流水线的第一项任务是顶点变换（vertex transformation）：把顶点从模型空间变到世界空间，再变到相机空间，最后变到屏幕空间。平移、旋转、缩放这些几何操作天然就是矩阵运算——你在这里已经能看到日后 GPGPU 思想的种子：图形不只是画画，它充满了结构化的数值计算。

顶点变换之后是形状/基元组装（primitive assembly）：把相邻顶点连成几何基元，通常是三角形。三角形之所以关键，是因为它定义表面——表面能与光交互、能在空间定向、能被硬件高效逼近。但此刻三角形还是几何对象，还不是离散的像素。这一步交给光栅化（rasterization）：把几何基元变成片段（fragment）/像素，确定"屏幕上哪些离散位置属于这个三角形"。到这儿，流水线从几何计算进入图像空间计算。

接下来是像素着色器（pixel shader，也叫片段着色器 fragment shader），决定每个片段的最终值：颜色计算、光照、纹理查找、纹理值混合、反射、阴影、法线贴图……一串效果都在这里算。流水线的收尾是光栅操作（raster operations，ROP）：深度测试、模板测试、混合、抗锯齿，把结果写进内存。

要点。 图形时代的 GPU，有些阶段可编程、有些固定。到 GeForce 7800 那一代，顶点和像素阶段可编程，ROP 基本固定。这个不对称很重要：GPU 里已经有大量可编程浮点算力，但它仍嵌在一条面向图形的流水线里。顶点着色器（vertex shader）处理"每个顶点"的几何变换，像素着色器处理"每个片段"的颜色计算——两个阶段一前一后，都是"同一条程序、海量数据元素"的执行模式，只是数据规模和访存模式不同。这种对称结构让研究者特别容易想到：把这两级"元素级程序"抽出来通用化，不就是一台并行计算机吗？

幻灯片里那句"GPU 是一台巨大的着色器（the GPU is a massive shader）"极具启发性：早期可编程 GPU 的计算心脏全在着色阶段。这正是研究者开始问"这些阶段能不能改做科学计算"的原因——如果像素着色器每周期能做大量算术，而每个像素又基本独立计算，那么结构相同的非图形负载，是不是也能在这上面跑？

3.5 可编程，不等于程序员友好 ​

讲座在这里做了一个微妙但必须的区分：一台机器可以是"可编程"的，同时极其不便编程。先学了 CUDA 的学生常常不知道早期有多痛，所以这个区分值得停下来说清楚。

GeForce 7800 上，可编程部分确实能算：浮点算术、点积这类向量运算、纹理访问、特殊函数、一些控制流、寄存器与输出间的数据移动——硬件是强的。但从科学家、金融研究员、ML 从业者眼里看，它极度不友好。

为什么不友好？ 因为即使你的问题不是图形，你也得用图形的语言思考。片段着色器有足够的算力做非图形的事，但周围系统还期待纹理、光栅化通道、渲染目标和帧缓冲。你想算 FFT、矩阵乘、分割，就得把问题映射到图形基元上。这不是语法别扭，而是强迫你错误地思考问题：你不是想"我有一个输入数组和一个要施加的函数"，而是想"我怎么把数据存成像素、怎么让图形系统渲染出一张像素值就是答案的图"。

帧缓冲在这里是个很好的镜子。普通图形里流水线把结果写进帧缓冲，因为目的是显示；可科学计算里，一个阶段的输出通常是下一阶段的输入，帧缓冲是错误的目的地。研究者学会把输出重定向进纹理内存，让一个渲染通道的结果成为下个通道的输入。这个 hack 居然能用，但它不得不存在这件事本身，就说明了这个编程模型有多陌生。

小结：可编程性解决的是"能不能表达"，不是"好不好表达"。早期 GPU 只解决了前者。

3.6 早期 GPGPU：图形之上的 hack ​

早期 GPGPU 可以概括成一句话：系统性地劫持图形流水线。讲座把这段逻辑拆得很清楚。

存储放哪、算术谁来？ 任何通用计算至少需要两样东西：存放数据的地方，和变换数据的算术单元。图形时代的 GPU 里，纹理内存提供存储，片段着色器提供算术。

为什么数据能存成纹理？因为大量数值工作负载本来就有结构：一维数组 → 一维纹理，矩阵 → 二维纹理，体积物理场 → 三维纹理。科学数据结构的每个元素都可以当成像素/纹素看待——这不是硬件设计者最初的设想，但表示足够接近，让这个把戏可行。

为什么计算能塞进片段着色器？因为片段着色器本来就是"把算术和过滤操作独立施加到大量图像空间元素上"。如果非图形负载也是由大量独立数据元素组成，片段着色器就能扮演计算内核（kernel）的角色。这才是图形与后来 GPU 计算之间真正的概念桥：两者共享一个结构性质——大量输出可以由对应输入并行算出。

但映射仍然痛苦。GPU 不支持科学程序员习惯的数组循环，人们得把光栅化本身当迭代机制：渲染一个和纹理同尺寸的屏幕对齐矩形，强制着色器对每个纹素执行一次；负载超出可渲染域，就把数据切成分块（tile）分多趟处理。这些解法很聪明，却暴露了底层错配：一条渲染流水线被掰成计算引擎的形状。

关键约束。 片段着色器能做的：ALU 运算、向量操作、特殊函数单元、纹理访问、控制流指令、数据移动、输出写回。但内存语义高度受限：位置 的输出通常必须写回对应的位置 ，没有散写（scatter）、没有共享内存、没有原子操作，读-改-写模式受限。图形能忍受，因为图形问题不要任意写模式；通用计算就太憋屈了。

于是早期 GPGPU 程序员要手管渲染通道、手把数组编码成纹理、躲开"非法"访存模式、甚至靠可视化调试。讲座把这个荒谬概括得极好：要用 GPU 跑领域应用，你几乎得想象怎么画出一张像素即答案的图。这就是早期 GPGPU 又兴奋又悲惨的原因——性能承诺是真的，抽象是错的。

3.7 Brook 与正确抽象的寻找 ​

从那个别扭世界走向可用世界，绕不开 Ian Buck 的 Brook 项目。讲座把 Brook 不是当作一个小语言提案，而是当作计算机史上的转折时刻——这个定性是站得住的。

动机。 2000 年代初有一条性能趋势：GPU 的浮点吞吐增长远超 CPU，而且两条曲线在发散。关键不是某个基准上 GPU 碰巧更快，而是差距看起来会越拉越大。若真如此，糟糕的编程模型就会成为进步的日益严重的路障——硬件机会大到不能无视。

于是 Buck 问了三个问题：

  1. GPU 编程能不能简化？
  2. 基于 GPU 的计算，正确的抽象是什么？
  3. GPU 上能高效实现的问题范围到底有多大？

这三个问题紧密耦合：坏抽象让简单的事变难、让性能不可预测；好抽象能暴露硬件真正适合的那类工作负载。

Brook 的答案。 优雅之处在于它匹配了硬件与负载的真实结构。Brook 引入两个概念：流（stream）——需要相似计算的一组记录；内核（kernel）——独立施加到流元素上的函数。这一手漂亮，因为它精确表达了让 GPU 值得用的两个条件：流暴露数据并行；内核承载实际计算，且可以按算术强度来分析。

编程模型刻意保持简单，不是天外语言，而是带适度扩展的 C。这是个关键设计选择——发明全新语言，采纳就难；贴着熟悉的语法，程序员就能写"看起来像一个逐元素函数"的代码，把映射交给编译器和运行时：

c
kernel void foo (float a<>, float b<>, out float result<>) {
    result = a + b;
}

这个模型的妙处在于：程序员每次只思考一个数据元素，实现却在整条大流上复制这份计算。普通 C 里你要写个全数组循环；Brook 里你只写一次局部计算，全局复制交给系统。这已经极其接近现代 GPU 编程的心态。

Brook 还配齐了让抽象落地的工具链：编译器 brcc 把 Brook 翻译成 Cg/HLSL 等着色器形式，运行时 brt 管理流纹理与内核执行。这提醒我们：编程模型从来不只是语法，它永远包含编译、资源管理与执行支持。

接轨：Brook 的完整论述见 Ian Buck 等, BrookGPU: Stream Computing on Graphics Hardware, SIGGRAPH 2004。它是"流编程"这条线最早成名的代表，日后 CUDA 的 stream 与 kernel 概念在这里已有雏形。

3.8 为什么有些应用在 GPU 上赢，有些输 ​

Buck 工作最有力的地方，也是讲座最有教学力的时刻之一：它不把 GPU 加速当魔法，而是问：GPU 到底什么时候真的胜过 CPU？答案系在两个反复出现的性质上，每个机器学习系统学生都该记住：数据并行（data parallelism）与算术强度（arithmetic intensity）。

讲座把它放进一条更大的系统原则里：算力在变便宜，存储相对慢，通信被物理封顶。所以现代性能论证不能停在数算术单元：一台机器标称算力再高，如果数据喂不进去，真实性能就坍缩到通信边界。算术强度不是边角术语，它是对"每搬一个字的数据，负载里有没有足够的有用工作来撑起一台吞吐机器"的紧凑提问。这个概念后来被 Roofline 模型正式化为性能分析的核心轴——你会在本书后面再见到它。

两个性质。 数据并行 = 很多数据元素能独立处理；算术强度 = 每单位传输数据上执行的计算量高到通信成本不占主导。负载光有"大量元素"不够——如果每个元素只做一次平凡运算，搬数据往返的开销足以吞掉任何算术优势。用一句话记：

公式直觉。 矩阵乘法是典范： 矩阵乘 矩阵，算术量 ，存储量 ，比值 。矩阵越大，作为 GPU 负载越划算——这就是大模型机器学习天然贴 GPU 的原因之一：它恰恰被这类操作主导。

讲座用 SAXPY、FFT、分割、SGEMV 等早期 Brook 评测说明这点。逐一看：SAXPY（标量乘加：）是教科书级的"每个元素一点运算、互不相干"——元素级并行拉满，数据复用为零，GPU 在这种负载上几乎必赢。FFT 有分段依赖（蝶形结构要逐级推进），还特别吃缓存局部性，CPU 的深缓存反而是优势。分割（segmentation） 是图像处理里的常见操作，元素独立、天然并行，是 GPU 的好菜。SGEMV（稀疏矩阵-向量乘）算术强度远低于矩阵乘——每个非零元搬进来只做一次乘加，通信成本压不住。这些例子的意义不是考古而是概念：有的负载数据复用少、有效并行高，于是加速巨大；有的要更多通信、更多同步，或当时硬件啃不动的结构复用，结果平平。它们训练你结构性而非情绪化地推理：GPU 抽象地不好也不坏，它的收益只在负载具备正确性质时出现。

还有传输开销。早期 CPU 与 GPU 内存空间分离，应用住在 CPU 侧，数据要过互连送进 GPU、算完再送回来。讲座强调这成本不能无视：基线应该对全套端到端路径（含启动与传输开销）比，而不只是对裸 GPU 算术吞吐。这个心态至今仍要紧——现代加速器很强，但搬数据依然贵。

80% 的启示。 Brook 证明抽象未必毁性能：某些情形下它的实现能达到手写 GPU 版本约 80% 的性能。这个结果极其重要——它意味着高层抽象可以足够接近硬件效率，在实践中有意义，值得为之设计。

3.9 从 Brook 到 CUDA：软硬件协同设计被挑明 ​

讲座在这里用了一个极好的教学问题：Brook 之后还缺什么？ 如果 Brook 已经给了更好的编程模型，故事为什么不就此结束？答案：Brook 再出色，仍是建在为图形设计的硬件之上。软件能虚拟化、能重新解释硬件，却无法从根本上移除那些本不为通用负载准备的硬件约束。

Ian Buck 加入 NVIDIA 后，事情变了。在斯坦福他只能搭软件侧；在 NVIDIA 他能影响硬件侧。从那时起，GPU 演化被 CUDA 的需求和 CUDA 激活的工作负载明确驱动——这是现代计算里软硬件协同设计最重要的案例之一：同一批人同时设计语言与它要跑的硬件。

需要的改动是实质性的。 图形时代约束要移除或削弱：指令集要变丰富——在浮点着色操作之外补上整数、位操作、更灵活的控制流、更通用的访存指令；通用 load/store 取代面向图形的访问模式；散写/聚集（scatter/gather）要支持；指针寻址要变得有意义；线程标识符要支持一维/二维/三维，让向量、矩阵、物理网格的程序员能自然表达问题。

最重要的是内存层级要变。早期图形流水线没有通用并行程序做线程协作所需要的、显式管理的片上共享内存（shared memory）。CUDA 时代的 GPU 引入共享内存，让局部线程组内部快速通信，不必每次打回片外 DRAM。这是意义深远的架构一步：生产者-消费者结构、分块线性代数从此能高效利用局部复用。

讲座把这个十年的演化概括得极好：2000 到 2010，GPU 硬件从"面向图形的可编程"走向"真正的通用可编程"，CUDA 是这次变化的软件面孔。表面上看 CUDA 可以叫"跑在 GPU 上的 C"，但这句话别让你低估它。CUDA 真正做的，是在程序员与一台吞吐导向的处理器阵列之间建立了一份新契约：

c
__global__ void KernelFunc(...);
__shared__ int SharedVar;
KernelFunc<<<500, 128>>>(...);

注意表层语法变化有多小：一个内核限定符、一个共享内存限定符、启动语法指定执行配置，加上 cudaMalloc / cudaFree / cudaMemcpy 等显式内存管理。CUDA 的天才部分在于：它几乎不让你牺牲熟悉感，同时把机器的本质结构暴露出来。

3.10 CUDA 生态的崛起 ​

讲座正确地坚持：工业平台不只是语言 + 编译器 + 运行时。学术界，原型能靠"证明可行"成功；工业界，平台成功要靠周围长出生态。所以 CUDA 的成长不只涉及语言与硬件设计，还包括库、工具、文档、教学推广、给专家的底层入口。

这解释了讲座为什么要提领域库和分层软件栈：有人要高层接口，有人要在 PTX 附近干活，还有人要常见例程的优化库。于是 CUDA 积攒出 cuBLAS（稠密线性代数）、cuFFT（傅里叶变换）、cuSPARSE（稀疏运算）、cuRAND（随机数）、图像视频处理库等一整套生态。教学含义在于：高性能计算平台之所以有用，不只是暴露硬件，而是把常见模式变成可复用、高度优化的组件。

讲座还把这条生态线接到今天的 ML 实践：PTX 层的底层优化工作，以及 Triton 这类现代抽象。这是一座重要的桥——它提醒你抽象与性能之争从未真正解决：新层会不断出现，因为我们总在可编程性与控制力之间寻找甜点位。今天你写 Triton kernel、写 torch.compile 的 inductor 后端，做的还是同一件事：在抽象与性能之间再找一个新平衡。

3.11 统一着色器架构与流式多处理器（SM）的登场 ​

GPU 历史上一个里程碑是 GeForce 8800 一代引入、与 Tesla 品牌绑定的统一着色器架构（unified shader architecture）。此前可编程资源与特定图形阶段绑得很死——顶点单元只能干顶点、像素单元只能干像素；此后可编程资源变得灵活、通用。讲座强调：从这时起，我们今天讨论 GPU 的架构词汇开始成形。

最重要的词是流式多处理器（streaming multiprocessor，SM）——NVIDIA GPU 的基本执行簇。一个 SM 里有多个 CUDA core、load/store 单元、特殊函数单元、大寄存器文件、共享内存、调度逻辑与指令派发机制。SM 教学上重要，是因为它是抽象编程模型与物理机器相遇的地方：以后学 block、共享内存、warp、occupancy，只要意识到"SM 就是这些现象发生的局部街区"，全都不难。

Fermi 时代的 SM 长这样（示意）：多个 CUDA 处理器、load/store 单元、特殊函数单元、可配置的共享内存与 L1 缓存、大寄存器文件、warp 调度器。架构信息很明确：SM 不是 CPU 意义上的单核，而是一个吞吐导向的簇，同时支持大量线程，靠本地存储与调度让一部分线程等内存时计算不停。每个 SM 能同时驻留的线程数受硬件资源（寄存器文件大小、共享内存容量、线程槽位）封顶——这个"能塞进多少线程"的指标就是日后的 occupancy（占用率），第 4 章你会在调优时反复和它打交道。

要点。 CUDA core 是处理浮点、整数、位运算的通用算术执行单元，但它是标量的。这一点极其重要——学生常把 CUDA core 想象成矩阵处理器，它不是。一个 CUDA core 每周期每线程只处理一个标量操作，大规模矩阵乘是无数标量操作协调出来的（直到专用张量硬件出现）。

接轨：SM 的组成与线程层级定义见 NVIDIA, CUDA C++ Programming Guide 的 "Hardware Model" 一节；Fermi 架构细节见 NVIDIA Fermi Compute Architecture Whitepaper（2010）。GeForce 8800 那代统一着色器架构的公开描述见 NVIDIA 当年的 GeForce 8800 架构资料。

3.12 为什么 CUDA 核不足以满足现代 AI ​

到 2010 年前后，GPU 在"能跑各种并行负载"的意义上已真正通用。但可编程性不再是主要瓶颈后，性能重新成为主问题。讲座把这次转换说得很尖锐：如果现代 AI 负载由矩阵乘法主导，而 CUDA core 本质是标量单元，那只用 CUDA core 做矩阵乘，就不是花硅面积和能耗预算的最优方式。

工作负载视角再次关键。主导 AI 负载不是任意标量算术，而是高度结构化的稠密线性代数。矩阵乘在 GPU 视角下性质很美：大规模并行、算术强度高。但光有这些不保证硬件效率最高——硬件原语是标量、应用原语是矩阵分块乘加，这中间仍有一层错配。硬件层面，你等于用一千次标量运算去模拟一个本可直接做的矩阵分块运算，每次标量运算都要过取指、派发、寄存器读写——全是浪费。

张量核（tensor core） 就是回答这个错配的。与其让机器用天文数字个标量操作合成大矩阵积，张量核直接加速矩阵分块操作。看矩阵乘法的机制：乘积由反复的行列点积构成，脑内旋转其中一个操作数视角，就能理解成大量乘加（multiply-accumulate）结构在对应元素上并行工作。张量核把这个结构内化进硬件。

这改变了原语的粒度。CUDA core 最好理解为线程级标量引擎；张量核最好理解为分块级矩阵引擎。这个转变大幅提升单位面积的浮点吞吐和能效——因为每个有用算术操作摊薄的控制与通信开销变少了。

CUDA core 心智模型：标量/向量算术
张量核心智模型：矩阵分块乘加

讲座强调差距的量级：张量核进场后，矩阵密集负载比只靠 CUDA core 可以有数量级的性能提升。现代 GPU 上，硅预算的很大一块给张量式计算、另一大块给 CUDA 式通用计算、剩下的给喂它们的内存与支撑结构——具体百分比代代不同，但总信息是对的：AI 一旦称王，架构就重新分配面积来回应。张量核不是"顺手加的协处理器"，它是负载结构逼出来的：矩阵乘占主导，就把芯片面积押给"一次做一块矩阵乘"的专用电路，而不是通用标量流水线。这也是"架构围绕工作负载设计"这一章母题最直白的例证——你在 GPU 芯片上看到的每一个功能块，都能追问出"它是被哪一类工作负载逼出来的"。

接轨：张量核首次出现在 NVIDIA Volta（Tesla V100，2017），Hopper（H100）与 Blackwell 世代逐代强化并扩充精度档位。

3.13 精度格式：数值与架构的深层连接 ​

张量核不只是换计算原语，还和换数值表示绑在一起。讲座从张量核自然滑向精度格式，是有道理的：现代 AI 硬件的吞吐增益，不只来自更多晶体管或更好的组织，更来自高效处理降精度表示、同时保住模型质量。

历史路径。 科学计算曾长期宠 FP64（双精度）——天文、物理模拟离了高精度会算错；GPU 后来把 FP32（单精度）做成图形与通用计算的中心——渲染的精度需求刚好够用；深度学习把 FP16 和 BF16 推到前台——FP16 尾数少但动态范围小，BF16 牺牲尾数、保留和 FP32 相同的指数范围，专门为训练时的梯度缩放而生；最近的 AI 硬件拥抱 FP8，有些场景甚至 FP4。讲座对"数字表示"的重要性措辞很强，完全合理——AI 硬件长期性能增益里有很大一块，来自减少每个值的比特数，从而降低存储成本、降低搬移成本、让每周期能做更多运算。这在架构上的连锁反应很直接：少一半比特，同一带宽下能喂给计算单元的数据翻倍，张量核同一硬件还能吃下两倍的操作数。

对机器学习系统学生这是特别重要的一课：硬件性能不能脱离算法稳定性单独分析。从 FP32 到 FP16 不只是硬件事件，它要求训练程序、归一化行为、缩放方法、数值策略都容忍更低精度；再推到 FP8/FP4，挑战再来一遍。讲座暗示剩下的"低垂果实"可能在从 4 比特往 1 比特推进——但这种进步不可能只来自硬件，模型和优化方法必须配套重新设计。换句话说，精度是算法与架构共谋的产物，不是硬件单方面决定的开关。

接轨：张量核的精度档位与格式演进见 NVIDIA H100（Hopper） 与 Blackwell 架构白皮书；训练侧的低精度策略（loss scaling、混合精度）对照 PyTorch 的 AMP 实践。

3.14 新瓶颈：数据搬移 ​

张量计算变快之后，一个新瓶颈必然上位：把数据及时送到计算单元。这是现代 AI 硬件最深的架构转换之一。机器每周期能做的算术远超前代，但数据仍从受延迟与带宽约束的片外内存层级到达，系统就越来越多地变成搬移受限（movement-bound）而不是计算受限（compute-bound）。

物理本质。 就算用高带宽内存（high bandwidth memory，HBM）这类更高带宽的技术，片外访问的延迟也不会消失。数据从片外内存出发，穿过缓存、共享内存、寄存器，然后进入计算单元。从 HBM 到张量核的路很长，每一级都算数。讲座因此强调：通信被物理封顶——你可以绕开瓶颈、隐藏它、重叠它、减少空转，但没法一念消除。

GPU 对内存延迟的经典应答是线程级延迟隐藏（thread-level latency hiding）：一组线程等数据时，另一组就绪线程顶上算。这也是 GPU 让这么多线程常驻的原因之一。但算术单元越专化、越强，光靠"线程多"就不总够了——线程切换本身也耗指令槽与寄存器；机器还需要专门机制，用大块、结构化的方式搬数据。

小结：算得快之后，喂数据成了主战场。架构的重心从"怎么算"移向"怎么搬、何时搬、少搬多少"。

3.15 TMA 与异步数据搬移 ​

张量内存加速器（Tensor Memory Accelerator，TMA） 就是讲座介绍的这类机制之一。理解 TMA 最快的办法是类比 DMA（direct memory access）：传统 DMA 场景里 CPU 只指定一次传输，专用硬件搬整块数据，而不是让 CPU 逐元素搬。TMA 对面向张量的 GPU 负载做类似的事：用硬件支持在全局内存与共享内存之间结构化搬移张量分块，而不是让一堆普通线程手动搬。

好处有三个。

降指令开销。 普通线程逐个加载会吞指令槽和寄存器资源；一次 TMA 传输顶掉成百上千条加载指令。

降同步痛苦。 大量线程各自请求内存，请求与返回的时机各不相同，容易在流水线里戳出气泡；硬件管理的大块传输时序更规整。

更有效的计算-通信重叠。 传输可以异步进行，别的工作继续跑——这是让"搬"和"算"并行推进的关键机制之一，也是现代推理/训练引擎里 overlap 流水线的底层依赖。

讲座也小心地记录了代价：硬件接管越多，运行时行为对程序员越不透明。显式的软件管理搬移费力但可预测；硬件管理的异步搬移整体可能更快，但更难精确推理。这是一个重要的系统权衡——性能与可预测性不总是一起上升。

处理 TMA 最有价值的地方在于教会你超越静态框图思考：硬件特性不是"存在"而已，它们改变优化的性质。TMA 一进来，一些旧的软件瓶颈消失，但冒出一些新的推理难题。现代系统工作经常就是这个样子：解决一个问题，把困难挪到别处。

3.16 特殊指令与"更多 FLOPS"到底意味着什么 ​

讲座还有一块很有价值的讲法，关于特殊指令：FMA、半精度点积、张量核矩阵指令、整数矩阵指令。学生常被 teraFLOPS 的营销话术唬住，以为这些数字纯粹来自更多核或更高主频。实际上被支持的指令的语义极其重要。

融合乘加（FMA）是现成的例子。机器学习和线性代数里 这种表达式不断出现——它就是一个 MAC（multiply-accumulate）结构，矩阵乘、卷积、点积全由它堆出来。若乘与加是两条指令，中间结果得写出来再读回去，白白费通信和延迟；融成一条，中间流量就没了，还只做一次舍入（rounding），数值上更稳。于是指令本身就承载了对一个常见计算母题的优化。讲座调侃说厂商把乘和加都算进宣传的浮点总量"几乎像作弊"——但深层点不在营销：硬件吞吐很大程度上取决于硬件直接支持哪些复合操作。同样的主频和核数，指令语义不同，可达吞吐可以差出几倍——厂商宣传的 "teraFLOPS" 通常就按 FMA 计数，一次 FMA 算作两个 FLOP，于是标称峰值天然是"指令级吞吐的两倍"。

半精度点积、张量矩阵指令延续同一逻辑：与其让软件从标量操作拼出每个模式，硬件越来越多地暴露与负载循环代数模式匹配的原语。在矩阵主导的 AI 世界，这完全合理。工作负载再次决定哪些指令配得上硬件支持。

讲座对长期性能增益的粗略分解也很好用。性能增长不来自单一来源，而是多个：数字表示改进、指令变丰富、制程进步、稀疏性利用、芯片面积变大。引人注意的是"表示"这一项贡献了非常大的份额——这正是低位机器学习不是小众话题的原因，它站在性能故事的中心。

3.17 在更大系统里的 GPU：加速器 ​

讲完这些，讲座提醒：GPU 通常仍是挂在 CPU 上的加速器（accelerator）。这是一个重要的系统视角。GPU 不是孤岛：人类开发者通常从 CPU 侧编辑代码、跑开发工具、启动作业、调试、管理整个应用流程，然后调用 GPU 干计算密集的部分。

主机-设备关系塑造软件模型。 数据通常在主机内存准备好，拷到设备内存，被一个或多个内核处理，再拷回来。尽管越来越多的集成系统在模糊主机与设备内存的界线（统一内存、PCIe 之上的内存一致性模型），这个心智模型依然重要。别把 GPU 编程想成"写并行代码"；把它想成把应用切分成主机驻留与设备驻留两部分，并仔细管理它们之间的通信：

c
// 概念性的主机-设备流程
prepare_host_data();
cudaMemcpy(device_input, host_input, ...);
Kernel<<<grid, block>>>(device_input, device_output);
cudaMemcpy(host_output, device_output, ...);

讲座强调，最后那步拷贝回来是真实的，不能挥手抹掉。忽略数据搬移，你的性能直觉就是错的。加速器很强，但边界算数。 第 3.8 节说的"传输开销"在这里以最朴素的形式重现：一次往返的 cudaMemcpy 可能比你精心优化的内核还贵——把 GB 级张量搬过 PCIe（几十 GB/s 量级带宽），光传输时间就够你把整块数据在 GPU 上算好几遍。这也是为什么现代系统拼命减少主机-设备往返：算子融合、持久化 kernel（persistent kernel，内核常驻 GPU、数据流式喂入）、统一虚拟内存与零拷贝映射，全都是"别把数据来回搬"这一原则的工程化。

3.18 内核思维：一份局部计算，无数全局实例 ​

内核（kernel）是 GPU 计算的入口。但概念重点不是"它跑在设备上"，而是：它表达一份会被实例化很多次的局部计算模式。写内核时，你通常不会把问题的整个全局控制流写在一个地方；你写的是"一个线程对一份数据/一个局部任务该做什么"，然后机器启动海量这样的线程。

这就是 GPU 模型既优雅又一开始令人困惑的原因。优雅在于局部规则只写一次；困惑在于你必须记住：全局行为是无数局部实例同时执行涌现出来的。讲座讲 Brook、CUDA、线程映射时反复回到这个想法。这个局部-全局之分值得再强调一次，因为它是几乎所有并行编程的核心——从 CUDA kernel 到 MapReduce，到分布式数据并行里的"每卡一份模型副本"（一份局部逻辑，多个全局实例），都是同一个模式。

3.19 为什么网格、块、线程要分层 ​

讲座里一个最有洞察力的课堂提问是：为什么 CUDA 执行模型要用网格（grid）、块（block）、线程（thread）三层？ 如果我要处理一百万个数据元素，为什么不干脆铺开一百万个线程、让运行时全自动决定一切？

答案：因为硬件内存层级不是平的。

数据可能从全局内存出发，但工作一分配下去，局部性（locality）就开始算数。同处一个 SM 的线程可以共享共享内存和邻近的寄存器状态这些快的本地资源；分到不同 SM 的线程要通信就贵得多。于是软件模型把线程分组为块（block），期望一块块落在一个 SM 上执行、吃本地存储与同步的好处。

这是学生最该内化的架构对应之一：

层级含义
网格（Grid）一个内核对应的全部工作
块（Block）打算共享一个 SM 本地资源的线程局部组
线程（Thread）那份局部计算的一个逻辑实例

关键推理。 层次结构存在，是因为内存层级存在。如果内存系统完全平坦，平的线程空间也许就够；但因为通信成本取决于线程住哪，执行模型必须暴露一部分局部结构。这个逻辑在分布式系统里以更大的尺度重复：数据并行时为什么要有"本地组"？因为跨节点同步比跨线程同步贵几个数量级——层级跟着通信成本走，这条规律从 SM 一直管到数据中心。

接轨：grid/block/thread 的完整定义、尺寸上限与可嵌套性见 NVIDIA CUDA C++ Programming Guide 的 "Execution Model" 一节。

3.20 线程索引：执行空间到数据空间的映射 ​

层级有了，每个线程还得知道自己该处理哪个数据元素。一维情形下 CUDA 用一个简单的线性化公式：设 blockIdx.x 是块索引、blockDim.x 是每块线程数、threadIdx.x 是块内线程索引，则全局线程索引为

讲座用每块 4 线程的简单例子推过一遍：第一个块覆盖索引 0 到 3，第二块覆盖 4 到 7，依此类推。把公式拆开看：blockDim.x 是"块有多长"，blockIdx.x 是"这是第几块"，blockIdx.x × blockDim.x 给你本块在全局中的起点偏移，再加块内的 threadIdx.x 就落到具体元素。这个一行公式其实是整个执行空间→数据空间映射的样板：任何把多维数据压成一维来并行处理的系统，最终都要算这么一笔账。

一维很简单，但讲座也点破为什么学生觉得 CUDA 编程心理上难：真实数据常常是多维的，而硬件索引与内存布局偏偏常要线性化或混合维推理。矩阵、图像、张量天然活在二维及以上，指针算术和底层存储却想要一维偏移。这中间的脑内换算极易出错、极其耗神——一个 2D 块索引和一个 3D 数据坐标，谁跟谁对齐，错一位就是错一整片。CUDA 为此提供了 threadIdx.y、threadIdx.z 等多维索引，以及 blockDim 的多维版本，把"块在网格里是 2D 的、线程在块里是 3D 的"直接暴露给程序员——但你仍然要自己维护"索引 ↔ 数据布局"的对应关系，因为内存本身是线性的一串地址。

对系统教育来说这是很有价值的观察。它解释了为什么高阶编译器与张量语言在 CUDA 可用之后依然重要：它们不只是便利，它们卸掉了把高维计算映射到低级执行结构的真实认知负担。你写一个 torch.matmul，编译器替你做完了这整套"高维语义 ↔ 一维索引"的换算；而这套换算本身，正是手写 CUDA 时最贵的脑力支出。

3.21 选网格与块大小的实际后果 ​

讲座收尾在一个非常实际的话题：CUDA 里网格与块大小该怎么选？ 讲座许诺后面回 CUDA 编程再细讲，但架构逻辑现在就能记下。

因为一个 warp 有 32 个线程，块大小通常选 32 的倍数，让硬件通道尽量用满。实践中 128、256 这类块大小很常见，因为它们同时平衡好几个目标：大到能在块内暴露有用并行，又不到总是榨干每个 SM 的共享内存或寄存器。

关键约束。 别把这些数字当魔法常数。它们是架构妥协。好块大小必须：尊重 warp 大小；在有益时给同一个 SM 上驻留多个块/多个 warp 留空间；不过度消耗共享内存或寄存器。每块共享内存用太多，会压减一个 SM 上能同时住的块数；寄存器压力过大，会压低占用率（occupancy）；块形状不好还可能和数据的自然布局错位，伤到合并访存（coalescing）和局部性。

这也是矩阵负载常用 16×16 这种二维块形状的原因——不是审美，是它天然贴矩阵分块的几何、共享内存分段和 warp 友好的执行。同理，网格要大到覆盖整个数据集，而且通常远大于 SM 数量，让调度器总有额外工作可拿、负载不均能被隐藏。这些不是零散的"CUDA 技巧"，它们是我们已经研究过的内存层级与执行层级直接推出来的结论。记住这条推导路径，比记住任何一组"推荐值"都值钱——因为硬件每代在变，推导路径不变。

3.22 SIMT、warp 与分叉 ​

讲座接着转向 NVIDIA GPU 的定义性执行概念：SIMT——single instruction, multiple threads，单指令多线程。SIMT 与 SIMD 相关但不等同：SIMD 里程序员通常要显式推理向量通道；SIMT 里程序员被允许想"许多标量线程"，但硬件仍然以共享同一条指令流的组为单位执行。

那个组就叫 warp（NVIDIA 术语，直接保留原文；早期机译常误作"扭曲"），在 NVIDIA GPU 上通常包含 32 个线程。硬件以锁步（lockstep）单位调度、执行 warp——一次取指、一次派发，32 个线程的算术单元并行执行同一条指令。这让编程模型比裸 SIMD 友好：每个线程仍有自己的线程标识符，也可以有自己的控制流。但分叉（divergence）的硬件后果不会消失：warp 内的线程走了不同的分支，warp 就无法真正同时执行所有路径，机器必须串行化这些路径，效率下降。

c
if (threadIdx.x % 2 == 0) {
    // 偶数线程走这里
} else {
    // 奇数线程走这里——同一个 warp 得分成两趟执行
}

所以分叉最好理解成吞吐问题，而不是正确性问题：分叉的程序照样给出正确答案，只是硬件利用率不好——那个 warp 里 32 个线程，跑每个分支时只有一半在干活，另一半空转。想象一下：32 个工人排成一排，命令是"偶数号先做动作 A，奇数号后做动作 B"。硬件唯一的选择是把队伍分成两趟，一趟做 A 时另一半干站着。分支嵌套越深、分歧越交错，空转的比例越大。所以 GPU 程序员一条不成文的铁律是：把分叉写进不同的 warp，让 warp 内部的 32 个线程尽量走同一条路——数据排布对齐、把条件判断从热循环里搬出来，都是为这个目的服务的。

讲座还简要提到动态 warp 成形（dynamic warp formation）——一种把程序计数器（PC）相同的线程重新聚组、缓解分叉的硬件技术。记不记得住这个术语无所谓，要记住的是底层原则：机器偏爱连贯执行，控制不规则就降吞吐。这也是为什么现代编译器极力把分歧压进不同 warp、用整块（unrolled）代码和数据对齐来讨好执行模型。

接轨：SIMT 与 warp 的权威定义见 NVIDIA, CUDA C++ Programming Guide；Fermi 白皮书对 SIMT 硬件行为有更细致的描述。

3.23 为什么机器学习那么贴 GPU ​

讲座的最终综合明确回到机器学习。现代 ML——尤其大模型训练与推理——里我们想优化的一切，都被同时具备两个惊人性质的运算主导：丰富的（abundant）数据并行，和高算术强度（尤其在大矩阵/张量运算里）。这两个性质恰好是 GPU 几十年演化一直在喂的东西：从图形时代的像素级并行，到 GPGPU 时代的流并行，再到深度学习时代的矩阵分块并行，硬件始终在奖励"并行 + 高复用"的负载。

这个观察也解释了为什么某些更早的神经架构不那么 GPU 友好。RNN、LSTM 这类循环网络有很强的顺序依赖——一个时间步的隐状态是下一步的输入，时间轴并行不起来；Transformer 核心运算暴露的并行多得多——矩阵乘与注意力点积全是可分块的，虽然它也带来自己的系统瓶颈（内存搬移、注意力的数据流挑战）。讲座用这组对照讲一个更大的观点：与并行硬件对齐的算法形式，在实践中可扩展性好得多。架构史一再复现这个规律：能塞进吞吐机器的算法结构，赢得部署。

所以机器学习落在 GPU 上不是偶然。它落在这里，是因为其主导计算的结构性质，正好匹配 GPU 从图形起源一路积累的结构长处。图形本身，说到底，早就满是矩阵运算、滤波、插值、大规模并行片段处理。回头看，从图形加速到 AI 加速的过渡顺理成章——但只有理解了让这次过渡成为可能的整条抽象与再设计链，它才显得顺理成章。

3.24 更深的系统一课：硬件成熟时，压力向上移 ​

讲座最耐人回味的反思在结尾附近，话题从硬件演化转向未来研究。过去二十年，GPU 栈的硬件、架构、软件系统层已经榨出了海量优化；平台一旦成熟到那个程度，靠同一路工程再榨就越来越难，剩下的机会越来越要求"意识到系统约束的算法创新"。

这是对本科生绝佳的一课，它教你系统前沿怎么移动：技术早年，基础平台建设是主挑战——先把能用的硬件造出来，把驱动、编译器、运行时凑齐；后来协同设计称王——语言设计者与芯片设计者坐在一起，让软件与硬件互相迁就，Brook→CUDA 是这条线的原型；再后来算法结构与编译器智能成为关键杠杆——硬件能榨的都榨完了，下一步改进必须改算法形态（如更稀疏的注意力、更低精度的表示）或让编译器自动搜索更优的映射（如 auto-tuning、TVM/LLVM 生态、Triton 编译器）。讲座讲低位训练、未来模型架构变化、自动化调参/编译器优化日益重要的角色时，指的就是这次转移。想一想：当张量核、TMA、精度格式这些硬件杠杆都被一代代压榨干净之后，下一轮数量级的提升从哪里来？几乎只能从"让算法本身更贴合硬件"与"让软件自动找到更好的映射"里来——低位训练（FP8/FP4 甚至更低）、对内存层级友好的模型结构（状态空间模型替代全注意力）、编译器里越来越强的自动调度与成本模型。每一行都要求你同时懂算法与硬件，这正是机器学习系统这门课的训练目标。

换句话说，GPU 架构的故事不会止于张量核或 TMA。它延续到这样的问题：什么模型在更低精度下仍然数值稳定？什么样的注意力或状态空间机制能更好利用内存层级？编译器能在张量调度、分块、融合上自动化到什么程度？优化空间的哪些部分该人搜、哪些该学习系统搜？这些不再是纯粹的硬件问题，但没有硬件理解，一个都答不了。成熟硬件把压力向上推给了算法与编译器，而回答这些问题仍然需要知道硬件在哪。

3.25 全章小结 ​

现在可以把整章的心智弧线连成一条线。GPU 一开始是固定功能图形加速器；晶体管预算上来后，它在选定的图形阶段变得可编程；这些可编程阶段暴露出惊人大的浮点吞吐。研究者意识到：如果非图形数据能用图形数据结构表示、非图形计算能用图形通道表达，GPU 能加速的就远不止渲染。这个洞见催生早期 GPGPU——但编程模型不自然，因为程序员不得不用图形的语言思考。

Brook 提供了正确的早期抽象：把 GPU 重新定义成"带内核的流处理器"，而不是"纹理、三角形、帧缓冲的机器"。CUDA 随后把这个转变工业化、普遍化，NVIDIA 同时重新设计硬件，让 GPU 不再需要靠图形把戏假装通用计算机。统一着色器架构、流式多处理器、更丰富的指令、通用访存、共享内存，把 GPU 变成真正的并行计算平台。之后深度学习称王，张量核、低精度格式、TMA 这类先进数据搬移机制，把机器推得和矩阵/张量密集负载更紧地咬合。

这一切对机器学习系统为什么重要，现在清楚了：现代 AI 负载恰恰是那种奖励"围绕并行与算术强度建造的高吞吐架构"的负载。但用好这份契合需要理解整条栈：负载结构、内存层级、执行模型、主机-设备交互、编译器行为、硬件专化。这就是为什么 GPU 架构不只是这门课的一章——它是组织性的章节之一，它解释了为什么栈的其余部分长成那样。

所以本章之后自然的下一步，不是为硬件事实而硬件事实，而是学着给这台机器编程。一旦你明白 GPU 为什么这样组织，内核、内存层级、网格、块、线程、warp、分叉、编译器优化这些概念就不再是孤立的规则，而是一个围绕吞吐、局部性、并行结构设计的处理器的软件可见的影子。第 4 章，CUDA，就从这里开始。

延伸阅读 ​