课程结构（来自 syllabus）

课件 new-pdc00-syllabus 把课程组织为五个模块（本完整版覆盖前四个模块的已讲内容 Lect01–10）：

各讲速览

课程地图

建议按「概念 → 单卡编程 → 算子优化 → 编译自动化 → 多卡分布式 → 通信」的顺序理解，每一层都在解决上一层撞到的「墙」。

第 1 讲 · 引论：为什么需要并行

本讲是《并行与分布式计算（AI 方向）（Intro-to-PDC, AI-centric version）》的开篇（Lecture 01: Introduction，2026 春季学期，授课教师罗国杰 Guojie Luo）。主线是：以现代生成式 AI（generative AI）对算力、存储与通信的巨大需求为背景，说明通用 CPU 为何在 AI 上「失败」、GPU 为何崛起，进而回顾当代计算架构的演进谱系（GPU / TPU / NPU / 空间可重构 / 存算一体 / 晶圆级 / 超节点），最后用模型参数与显存容量的鸿沟（内存墙）引出核心问题——「并行到底能不能帮上忙」，并用 Amdahl 定律与 Gustafson 定律给出「不能」与「能」两个答案。

1.1 「吃 Token」的生成式 AI 应用（Token-Hungry Generative AI）

课件用一组当代应用说明 AI 对算力/token 的消耗正在爆发式增长：

• LLM 驱动的自主智能体（LLM-powered autonomous agents）：需要反复调用大模型进行推理与规划，token 消耗巨大。
• AI 手机与 AI PC（AI mobile phone and AI PC）：端侧也在集成大模型能力。
• Token 消耗趋势（Token consumption trends）：整体呈快速上升。
• 视频生成（Video generation）：生成式内容进一步推高算力需求。

图片来源（课件标注）：lilianweng.github.io、openrouter.ai、36kr.com（含 doubao phone）、github.com/baaivision/URSA。

1.2 AI 硬件谱系：按部署场景分类

课件给出一张「AI 硬件谱系」示意，按算力/功耗从大到小分为四类场景，并各举例子：

1.3 AI/ML 的硬件需求：三大维度

课件从三个维度刻画 AI/ML 对硬件的需求：

• 模型参数（Model Parameters）：呈指数增长（exponential growth）。
• 算力需求（Computational Demand）：其增速超过摩尔定律（outpacing Moore's Law）；课件还引用了 OpenAI 关于「计算成本 vs 美国 GDP」的对比。
• 通信规模（Communication Scale）：成为新的关键瓶颈（the new critical bottleneck）。

课件引用来源：medium.com 的 AI 计算成本分析、ourworldindata.org 的 AI 参数量数据。

1.4 通用 CPU 为何「失败」于 AI

1.4.1 CPU 的历史地位与两大危机

课件先肯定 CPU 的历史地位：数十年来，中央处理器（Central Processing Unit, CPU）一直是通用计算的普适引擎，其设计面向顺序（sequential）、通用（general-purpose）任务。但面对 AI，它遭遇「两大危机（The Two Crises）」：

课件引用：A. Gholami 等，"AI and Memory Wall"，IEEE Micro, vol. 44, no. 3, pp. 33-39, 2024。

1.4.2 架构失配与结论

• 失配（The Mismatch）：CPU 的顺序设计对 AI 负载（如矩阵乘 matrix multiplication）中固有的大规模并行从根本上是低效的。
• 利用率极低：CPU 在神经网络负载上只有 5–10% 的利用率，大约只交付 $100$ GFLOPS，却要消耗数百瓦（hundreds of watts）功耗。
• 结论（Conclusion）：AI 的指数级需求暴露了通用计算模型的极限，迫使转向专用硬件（specialized hardware）。

1.5 GPU 的崛起（The Rise of the GPU）

• 意外的英雄（An Unlikely Hero）：图形处理器（Graphics Processing Unit, GPU）诞生于 1990 年代末，本是为电子游戏图形而生，天然拥有数千个用于并行任务的简单核心。
• 架构上的共鸣（Architectural Resonance）：图形运算的核心数学（矩阵与向量运算）与深度学习的数学完全一致，使 GPU 成为一个「完美但纯属意外」的契合体。
• 软件催化剂 CUDA（The Software Catalyst）：NVIDIA 的 CUDA 平台（2007 年）把 GPU 解放出来用于通用计算，让不具备深厚图形学背景的 AI 研究者也能用上它的算力。
• 结果（The Result）：AI 模型的训练时间从「周」骤降到「小时」，GPU 由此确立为 AI 训练的事实标准（de facto standard）。

课件引用：J. D. Owens 等，"GPU Computing"，Proceedings of the IEEE, vol. 96, no. 5, pp. 879-899, 2008。

1.6 GPU 架构实例：Blackwell（GB202）

1.6.1 整芯片层级（GPC / TPC / SM）

• 完整的 GB202 GPU 包含 12 个图形处理簇（Graphics Processing Clusters, GPC）。
• 共 $12\times 8 = 96$ 个纹理处理簇（Texture Processing Clusters, TPC）。
• 共 $96\times 2 = 192$ 个流式多处理器（Streaming Multiprocessors, SM）。
• 配备 512-bit 内存接口，含 $16$ 个 32-bit 内存控制器。

1.6.2 单个 SM 的内部构成

• 每个 SM 含 128 个 CUDA Core、1 个第 4 代 RT Core、4 个第 5 代 Tensor Core、4 个纹理单元（Texture Units）。
• 片上存储：256 KB 寄存器文件（Register File） + 128 KB L1/共享内存（L1/Shared Memory）。
• 此外每个 SM 还含 2 个 FP64 Core（图中未画出）；其 FP64 的 TFLOP 速率仅为 FP32 的 $1/64$。

课件引用：NVIDIA RTX Blackwell GPU Architecture 白皮书。

1.7 第 5 代 Tensor Core 的 MMA

• 第 5 代 Tensor Core 的矩阵乘累加（MMA，PTX 中为 tcgen05.mma）完全不再用寄存器保存矩阵；操作数改为驻留在共享内存（shared memory）与张量内存（Tensor Memory）中。
• 指令由单线程（single-thread）发射，而非以前的线程束组（warp-group）。

1.8 Tensor Core 在 NVIDIA GPU 中的扩展趋势

编者补充 这一「算力涨、访存不涨」的剪刀差，正是后文「内存墙」与「通信是新瓶颈」的微观体现。

1.9 计算架构的演进谱系（Compute Architectures Evolve）

课件把当代计算架构归纳为若干「系列（series）」，说明系统在并行度与异构性两方向同时扩展：

1.10 华为达芬奇架构（Huawei DaVinci）

• Cube（立方计算单元）：$4096$（即 $16^3$）个 FP16 MAC + $8192$ 个 INT8 MAC。
• Vector（向量单元）：2048-bit 的 INT8/FP16/FP32 向量运算，含特殊函数（激活函数、非极大值抑制 NMS、ROI、SORT 等）。
• 显式内存层级设计（Explicit memory hierarchy）：由 MTE 管理。

课件引用：H. Liao 等，"DaVinci: A Scalable Architecture for Neural Network Computing"，IEEE Hot Chips 31 (HCS), 2019。

1.11 Intel 高级矩阵扩展（AMX）

• 策略（Strategy）：为 CPU 增强 AI 加速能力。
• 架构（Architecture）：它是 x86 指令集（ISA）的扩展，不是独立芯片。
• 核心组件：Tiles——二维寄存器文件（2D register files）；TMUL——专用矩阵乘指令。
• 采用显式的加载/存储/配置模型；指令包括 tilecfg、tileload、tmul、tilestore。
• 初始数据类型支持：INT8、BF16（后续扩展 FP16）。

1.12 AMD 的 SoC：Zen + RDNA + XDNA

课件以 AMD 的一体化 SoC 为例，展示 CPU（Zen）、GPU（RDNA）与 NPU（XDNA）在同一芯片上的异构集成。编者补充 这与 1.11 的 Intel AMX 同属「让 CPU 平台获得 AI 加速」的思路，只是 AMD 通过独立的 XDNA NPU 而非纯 ISA 扩展来实现。

1.13 存算一体架构（Compute-in-Memory，例：d-Matrix）

• 算——数字存算一体（DIMC）：把计算单元物理集成在 SRAM 内部以避免数据移动；全数字精度，支持 MXINT8/MXINT4 块浮点格式；高能效达 $38$ TOPS/W，而 H100 约为 $\sim 2.8$ TOPS/W。
• 存——依赖片上 SRAM 与创新内存层级：Stash（权重存储）每切片 $6$ MB、聚合带宽 $150$ TB/s；Global Memory（激活值存储）每切片 $4$ MB；LPDDR5（$256$ GB 容量）用于处理超大模型或 KV Cache。
• 扩展性与软件栈：Scale-Out 互连技术 + 透明网卡，绕过复杂的 RDMA；Aviator 软件栈中的 Host-Free 运行时由设备端 RISC-V 核心管理计算图执行，避免 Host CPU 瓶颈。

1.14 晶圆级集成架构（Wafer-Scale，例：Cerebras）

• WSE-3 规格（对比 NVIDIA H100）：面积 $46{,}225\ \text{mm}^2$（约 $57\times$）；核心数 $900{,}000$ 个 AI core（约 $52\times$）；片上内存 $44$ GB SRAM（约 $880\times$）；内存带宽 $21$ PB/s（约 $7000\times$）。
• WSE-3 Core：计算与内存紧密排列；支持高性能 AI 计算（对 16/8b 数据分别为 8/16-way SIMD）；高带宽 SRAM 与 Cache。
• 推理方案：片上流水线并行（Pipeline Parallelism）——将模型各层物理映射到晶圆的不同区域，通信延迟近乎为零。

1.15 超节点架构（Supernode，例：NVL72 与 CloudMatrix 384）

课件图片来源：zhuanlan.zhihu.com/p/1919101924445254162。

1.16 模型参数 vs GPU 显存容量（内存墙）

课件用「NLP 摩尔定律 vs GPU 显存增长」的对照数据，揭示内存墙（memory wall）：

1.17 并行到底能不能帮上忙？两条定律

课件用两条定律给出「不能」与「能」两个方向的回答。记号（Notations）统一为：$s$ 为加速比（speedup），$p$ 为可并行比例（parallel fraction），$n$ 为处理器数（number of processors）。

Amdahl 定律在实践中（in practice）会更糟：现实里往往并行度不足（not enough parallelism），再叠加数据依赖（data dependencies）与 I/O 瓶颈（I/O bottlenecks），进一步压低可达加速比。

1.18 本讲小结（Summary）

• 需求：当代 AI 带来指数级的算力、存储与通信需求。
• 架构：出现并行且异构（Parallel and Heterogeneous）的架构谱系。
• 并行能帮忙吗？「不能」——Amdahl 定律（规模固定）；「能」——Gustafson 定律（规模可扩展）。
• 系统如何「扩展」：现代计算系统在并行度（计算单元数量）与异构性（计算单元类型）两个维度同时扩展——这要归因于 Dennard 缩放的终结。
• 如何编程这个「怪物」？理想上应有全自动的编译器与运行时（fully-automatic compilers and runtimes），但目前尚未达到——因此仍需要能「hack 系统」的领域专家（domain experts）。

第 2 讲 · 异构编程：CUDA 核心

本讲聚焦异构计算（heterogeneous computing）的编程模型，讲师为罗国杰（Guojie Luo）。课件三条主线分别是：主机、设备与互连（Host, Device, and Interconnect）、为 CUDA 核心编程（Programming CUDA Cores）、以及为张量核心编程（Programming Tensor Cores，留待后续讲次）。围绕 CPU（host）与 GPU（device）如何协同、如何通过 PCIe / NVLink 互连、CUDA 编程模型（thread / block / grid / warp）、SIMT 执行、内存模型、占用率（occupancy）、控制分歧与谓词执行、合并访存与 bank conflict 等展开。本讲主要覆盖本讲的前两条主线；张量核心部分明确留到后续讲次。

2.1 AI 硬件谱系与 AI 芯片的两种形态

课件先给出 AI 硬件的谱系（以英伟达等为例），从云端到边缘、再到极小设备：

AI 芯片以英伟达为例有两种形态：

• 独立 AI 芯片（standalone）：如 DGX A100。
• 集成 AI 芯片（integrated）：如 Thor。

2.2 CPU 与 XPU 的异构集成

课件把 CPU 与 XPU（各类加速器）的异构集成（Heterogeneous Integration）分为两种层级：

2.3 PCIe 基础知识

• 市场采用情况（截至 $2025$ 年 $8$ 月）：主流为 PCIe 4.0；最新设备用 PCIe 5.0；PCIe 6.0 / 7.0 尚未量产。
• PCIe 通道（Lanes）：「x」= 通道对（TX/RX）的数量。例如 PCIe 5.0，x1 $=$ 4GB/s，x16 $=$ 32GB/s。
• 典型配置：x4 多用于大多数 NVMe SSD、NIC；x16 多用于大多数 GPU、FPGA。

常见 PCIe 设备（课件示例）

2.4 PCIe 在 AI 中的角色（1）：GPU–Host 连接

截至 $2025$ 年 $8$ 月，GPU 大多通过 PCIe 与 host 相连：

• Host $\to$ GPU：命令（CUDA kernels、OpenCL 二进制）、训练数据、模型权重等。
• GPU $\to$ Host：GPU 状态（kernel 完成、CUDA events）、计算结果（梯度、推理输出）等；GPU 经 PCIe 连到主板。

数据传输路径的演进

2.5 PCIe 在 AI 中的角色（2）：GPU–GPU 连接

在 $2016$ 年 NVLink 发布之前，GPU 之间也通过 PCIe 互连，用于数据 / 流水线 / 张量 / 专家并行（Data / Pipeline / Tensor / Expert Parallelism）。

2.6 NVLink：GPU 互连

• NVLink 1.0 于 $2016$ 年随 P100 GPU 首次引入：$4$ 条 NVLink，每条 $8$ 个通道（channel），双向带宽 $40$ GB/s，合计 $160$ GB/s。
• 已快于 PCIe 3.0 x16（$32$ GB/s）。
• 基于高速差分信号的串行通信（serial communication based on high-speed differential signaling）；NVLink 1.0 每通道 $5$ GB/s。
• 最新一代为 $2024$ 年的 NVLink 5.0。

2.7 异构 Host–Device 交互流程

课件对比了「不使用统一内存」与「使用统一内存（Unified Memory）」两种主机侧编排流程：

2.8 从 GPU 到通用 GPU（GPGPU）

课件回顾了 GPU 走向通用计算（General-Purpose GPU, GPGPU）的演进：

• 传统 GPU 架构把核心按图形流水线分工：vertex cores（顶点）、geometry cores（几何）、pixel cores（像素）。不同图像输入会造成各类核心的利用率不足（underutilization）。
• 解决方案：统一着色器核心（Unified Shader Cores）。代表产品：$2004$ 年用于 XBox 360 的 ATI Xenos；$2006$ 年的 NVIDIA GeForce 8800GTX。

CUDA 之前用 GPU 做通用编程

• Larsen & McAllister（SC 2001）：用图形硬件做大规模矩阵乘法，通过「可视化」简单并行算法的计算过程得到结果；受限于当时图形硬件的精度，属前瞻性研究。
• Thompson、Hahn & Oskin（MICRO-35, 2002）：提出编程框架并应用于矩阵乘法、3-SAT 等问题，相比 CPU 实现有惊人的性能提升，同时分析了瓶颈。
• 该 MICRO'02 论文第 6 节「Analysis」对未来图形架构提出了改进建议：更快的内存接口、从显卡到主存的 DMA 硬件、操作系统管理显存、更好的算术精度控制、逻辑布尔运算、跨顶点程序调用保留状态的能力、以及一个编译器。

2.9 CUDA：Compute Unified Device Architecture

课件列出的 CUDA 奠基性文献：

• J. Nickolls, “GPU parallel computing architecture and CUDA programming model,” IEEE Hot Chips 19 (HCS), Aug. 2007.
• J. Nickolls, I. Buck, M. Garland, K. Skadron, “Scalable Parallel Programming with CUDA,” Queue, 6(2):40–53, Mar. 2008.
• E. Lindholm, J. Nickolls, S. Oberman, J. Montrym, “NVIDIA Tesla: A Unified Graphics and Computing Architecture,” IEEE Micro, 28(2):39–55, Mar. 2008.

2.10 CUDA 编程的六个核心问题与概念对照

课件把 CUDA（以及推广的异构）编程归纳为一张问题—概念对照表：

此外还有一些杂项主题（Misc.）：调试工具（debugging tools）、有用的库（useful libraries）。

2.11 CPU 与 GPU 如何交互

一则关于 host 与 device 的故事：

• host ≈ CPU；device ≈ GPU。在 OpenCL / ROCm 中，CPU 也可以充当 device。
• CPU–GPU 接口带宽（课件数据）：x86–NV5090 @ PCIe 5.0 为 $64$ GB/s；Grace-Hopper NVLink C2C 为 $900$ GB/s；Grace-Blackwell NVLink C2C 为 $1.8$ TB/s。
• kernel ≈ 面向 host 的一个 device 函数接口。
• 数据管理靠 xxxMalloc()、xxxMemcpy()（CUDA 中即 cudaMalloc() / cudaMemcpy()）。

Host 程序

• 作为普通 C/C++ 应用的一部分运行在 CPU 上。
• 启动（launch）一个由 kernel 线程块组成的 grid。
• 当所有线程都终止后，调用才返回。

Kernel 程序

• __global__ 标记一个 kernel 函数。
• 线程根据自己在块内的位置（threadIdx）以及所属块在 grid 中的位置（blockIdx）计算其线程 id。

2.12 如何表达并行：线程层次（Thread Hierarchy）

Grid：一维线程块布局

grid 中的线程可访问以下内置变量（以课件一维示例，某 Thread 处于第 $4$ 个块的第 $0$ 号线程，每块 $4$ 线程，共 $8$ 个块）：

• 所属块编号：blockIdx.x（示例中 $= 4$）。
• 块内线程编号：threadIdx.x（示例中 $= 0$）。
• 块的维度（每块线程数）：blockDim.x（示例中 $= 4$）。
• grid 中块的数量：gridDim.x（示例中 $= 8$）。

Grid：二维线程块布局

• 每种颜色是一个线程块，每个小方格是一个线程；概念与一维相同。
• 二维下每个块与线程都有二维索引：块 ID 为 blockIdx.x、blockIdx.y；块内线程 ID 为 threadIdx.x、threadIdx.y。

2.13 如何表达同步

• __syncthreads()：块内线程同步屏障。
• xxxDeviceSynchronize()：设备级同步（如 cudaDeviceSynchronize()）。
• events（事件）。

2.14 如何管理与访问数据：CUDA 内存模型

• 在不使用统一内存时，host 与 device 拥有各自独立的地址空间（distinct host and device address spaces）。
• 通过 cudaMalloc() 在 GPU 上分配内存。
• 通过 cudaMemcpy() 在两个地址空间之间搬运数据。

kernel 可见的三类地址空间

2.15 kernel 代码如何在 GPU 上执行：SIMD + 流水线

• SIMD 执行可与流水线（pipelining）相结合。
• 所有 ALU 执行同一条指令。
• 流水线把指令拆成多个阶段（phases）。
• 当第一条指令完成时（课件例中为 $4$ 个周期），下一条指令已就绪可执行。

2.16 SIMD、SPMD 与 SIMT 辨析

SIMT：如何表达一个线程的工作（单指令流）

一个「尴尬并行（embarrassingly parallel）」循环可翻译为 GPU kernel：

// 串行循环：
for (int i = 0; i < N; i++)
    h_a[i] *= 2.0;

// 翻译成 GPU kernel：
__global__ void myKernel(int N, double *d_a) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) d_a[i] *= 2.0;
}

• 将从 host 启动的 device 函数称为 kernel，用 __global__ 属性声明。
• kernel 应声明为 void。
• 传入 kernel 的所有指针必须指向 device 上的内存。
• 所有线程「同时」执行 kernel 主体；每个线程用其唯一的 thread 与 block ID 计算全局 ID。
• grid 中的线程数可能多于 $N$（故需 if (i < N) 边界判断）。

SIMT：如何表达线程结构（从 host 启动 grid）

dim3 threads(256, 1, 1);            // 一个线程块的 3 维尺寸
dim3 blocks((N + 256 - 1) / 256, 1, 1);  // grid 的 3 维尺寸
myKernel<<<blocks, threads, 0, 0>>>(N, a);

其中 <<<blocks, threads, 0, 0>>> 为 kernel 启动的执行配置语法（依次为 grid 维度、block 维度、动态共享内存字节数、stream）。

2.17 优化（Optimizations）

• 资源感知优化（resource-aware）：关注的资源包括 SM、寄存器（registers）、共享内存（shared memory）、内存带宽（memory bandwidth）。
• 手段（不完全列表）：
  • 选择「最优」块大小，以最大化活跃 warp 数，从而提升内存带宽利用率；
  • 在划分后把（任务, 数据）映射到合适的 thread id，以：避免控制分歧（control divergence）、合并全局访存（coalesce global memory access）、避免共享内存 bank conflict。
• 需借助性能分析工具（performance analysis tools）。

2.18 占用率（Occupancy）：寄存器

寄存器可用量是较大 kernel 的主要限制因素之一。

2.19 占用率：共享内存

• 每个 SM 上的共享内存有限（近期 GPU 为几十至一百多 KB）。
• 共享内存限制每 SM 的活跃块数。
• 每块的数据量可能与线程数无关（如直方图 histograms 固定），也可能随线程数变化（如矩阵乘法、卷积）。

2.20 占用率：线程

• GPU 对每块最大线程数有硬件限制：AMD GPU 每块 $256$ 线程；NVIDIA GPU 每块 $512$ 线程。
• NVIDIA GPU 有每 SM 的限制（随型号而定），例如每 SM $768$ 或 $1024$ 线程、每 SM $8$ 或 $16$ 个 warp。
• AMD GPU 有全 GPU 范围的 wavefront 限制，例如 5870 GPU 上 $496$ 个 wavefront（约 $25$ 个 wavefront / 约 $1600$ 线程每计算单元）。

2.21 占用率：限制因素与权衡

• 三大因素（寄存器、共享内存、线程/块上限）的最小值决定一个计算单元的活跃线程数（即占用率）。
• 因素间相互作用复杂：限制因素可能是线程粒度或 wavefront 粒度；改变工作组大小会影响寄存器或共享内存用量；略微减少任一因素（如寄存器用量）可能让另一个工作组也能活跃。
• NVIDIA 的 Nsight occupancy calculator 可把这些因素可视化，帮助权衡。

占用率的权衡（Tradeoffs）

• 共享内存与寄存器在其它 warp 执行时仍持久驻留于 SM，因此上下文切换开销低。
• 每 SM 可支持的活跃 warp 数有限，由每块共享内存与每线程寄存器用量决定。
• 可用一个称为利用率（utilization）的指标表达每 SM 可能的活跃 warp 数。
• 活跃 warp 越多，NV 与 AMD 硬件上的延迟隐藏（latency hiding）越好。

2.22 分歧控制流（Divergent Control Flow）

• 在 AMD 与 NVIDIA 上，一个 wavefront / warp 内的指令都以锁步（lockstep）方式发射。
• 但每个 work item 可以走与同 wavefront 其它线程不同的路径。
• 若 wavefront 内的 work item 走上分歧的控制流路径，硬件会用掩码（mask）屏蔽无效路径。
• 应把分支限制在 wavefront 粒度，以防发射被浪费的指令。

2.23 谓词执行（Predication）与控制流

当同一条指令发给 warp 中所有线程、而它们要走不同执行路径时如何处理？

• 谓词执行（predication）是缓解条件分支代价的方法，对短代码段的分支尤其有益。
• 其基础在于：执行一条指令再丢弃其结果，可能与执行一个条件判断一样高效。
• 编译器可能用分支谓词化来替换 switch 或 if-then-else 语句。

GPU 上的谓词执行

• Predicate（谓词）是根据条件设为 true / false 的条件码。
• 条件流的两个分支都被调度执行：谓词为 true 的指令提交结果；谓词为 false 的指令不写结果、不读操作数。
• 仅对非常短的条件分支能带来性能收益。

__kernel void test() {
    int tid = get_local_id(0);
    if (tid % 2 == 0)
        Do_Some_Work();
    else
        Do_Other_Work();
}
// Predicate = True  给线程 0,2,4,...
// Predicate = False 给线程 1,3,5,...
// else 分支时谓词取反

分歧 vs 无分歧：两个案例

谓词执行对性能的影响

设 if 分支耗时 $t_1$、else 分支耗时 $t_2$。谓词执行会先执行 if 分支（$t_1$），随后丢弃无效结果、反转掩码，再执行 else 分支（$t_2$）。因此总时间为：

即两条路径被串行执行，这也是控制分歧代价的直观来源。

2.24 Warp Voting 示例（Histogram256）

• 每指令的隐式同步使得 warp voting 之类的技术成为可能，对不支持共享内存原子操作的设备有用。课件以 $256$ 桶（bin）直方图为例，假设 $64$ 线程的块。
• per-thread 子直方图所需共享内存：$256 \text{ bins} \times 4 \text{ B} \times 64 \text{ threads} = 64$ KB；而 G80 GPU 只有 $16$ KB 共享内存，放不下。
• per-warp 子直方图所需共享内存：$256 \text{ bins} \times 4 \text{ B} \times 2 \text{ warps} = 2$ KB，可行且能让多个工作组同时活跃。

用 warp voting 处理 warp 内冲突写

• 用warp 内线程 ID 作 tag标记对 per-warp 子直方图的写；下一轮循环时线程可检查自己上一次的写是否成功。
• 课件示例把 $32$ 位无符号整数拆成 $5$ 位 tag $+$ $27$ 位计数：读当前值、加 tag 与自增后的计数写回、再检查是否成功提交，未成功则回到 while 循环重试。
• 最坏情况：当 $32$ 个线程都写同一个 bin 时，需要 $32$ 次迭代。

void addData256(
    volatile __local uint *l_WarpHist,
    uint data, uint workitemTag) {
    unsigned int count;
    do {
        count = l_WarpHist[data] & 0x07FFFFFFU;   // 读当前值
        count = workitemTag | (count + 1);        // 加 tag 与自增计数
        l_WarpHist[data] = count;
    } while (l_WarpHist[data] != count);          // 未成功则重试
}

2.25 使用 Wavefront/Warp 的陷阱

• OpenCL 规范未涉及 warp / wavefront，也未提供跨平台查询其大小的手段。AMD GPU（5870）每 wavefront $64$ 线程，NVIDIA 每 warp $32$ 线程；NVIDIA 的 OpenCL 扩展只对 NVIDIA 硬件返回 warp size。
• 跨设备维持性能与正确性更难：把代码硬编码为 $32$ 线程/warp，在 AMD（$64$）上会浪费执行资源；硬编码为 $64$ 在 NVIDIA 上可能导致竞争并影响本地内存预算。
• 可移植性做法：在 JIT 时确定 warp size——检测运行在 AMD 还是 NVIDIA，并向构建命令追加 -DWARP_SIZE。

线程小结（Threads: Summary）

• 为性能，工作组内的分歧应限制在 wavefront / warp 粒度。
• 在「避免分歧的方案」与「可快速谓词化的简单代码」之间存在权衡——分支通常高度偏置且局部化，从而形成短的谓词化块。
• 应最大化任意时刻活跃的 wavefront 数以隐藏延迟；活跃数由寄存器、本地内存等资源需求决定。
• 针对特定 wavefront 的实现能带来更优化的实现，但因 AMD 与 NVIDIA 的 wavefront 大小不同，维护性能与正确性可能困难。

2.26 合并访存（Coalescing Memory Accesses）

示例与总线寻址

• 数组 X 是指向整型数组（每元素 $4$ 字节）的指针，位于地址 0x00001232。某线程访问 int tmp = X[0];。
• 假设 GDDR 内存总线为 $32$ 字节（$256$ 位）宽；字节可寻址的总线访问必须按总线宽度对齐，故低 $5$ 位被屏蔽。
• 期望地址 0x00001232，总线掩码 0xFFFFFFE0，实际总线访问 0x00001220。0x00001220~0x0000123F 范围内的任意访问都产生地址 0x00001220。
• 该范围全部数据都被返回：本例中 $4$ 字节有用、$28$ 字节被浪费。

合并的作用

• 为充分利用总线，GPU 在可能时把多个线程的访问合并（coalescing）成更少的请求。
• 考虑 int tmp = X[threadIdx.x + blockDim.x * blockIdx.x];：前 $16$ 个线程访问 0x00001232~0x00001272。若每次单独发送，则 $16$ 次访问共 $64$ 字节有用、$448$ 字节浪费；注意同一 $32$ 字节范围内的访问返回完全相同的数据。
• 当线程访问落在同一 $32$ 字节范围时被合并，每个范围只访问一次。本例只需 $3$ 次访问；若数组起始是 $256$ 位对齐，则只需 $2$ 次。

2.27 线程映射（Thread Mapping）

• work item（GPU 线程）$\Leftrightarrow$ 数据项：恰当的映射能与硬件对齐、带来巨大性能收益；不当的映射对性能是灾难性的。
• 以串行矩阵乘法为例：适合输出数据分解（output data decomposition）。创建 $N \times M$ 个线程，等效地消去外两层循环；每个线程做 $P$ 次计算，内层循环留在 kernel 中。
• 问题：索引空间应是 $M \times N$ 还是 $N \times M$？

两种映射在功能上等价，性能差异源于全局内存总线上的访存模式：

• 假设行主序（row-major）数据，一行中的数据（相邻列的元素）在内存中连续存储。
• 为保证合并访问，同 wavefront 的连续线程应映射到矩阵的列（第二维）——这对矩阵 B、C 给出合并访问；对矩阵 A，访问模式由迭代变量 i3 决定，线程映射不影响它。
• 映射 1 中，连续线程 tx 映射到 C 的不同行、非连续线程 ty 映射到 B 的列 $\to$ 低效访存。
• 映射 2 中，连续线程 tx 映射到 B、C 中的连续元素 $\to$ 两矩阵访问均合并（合并程度取决于工作组与数据尺寸）。

2.28 矩阵转置（Matrix Transpose）

• 转置定义直白：$\text{Out}(x, y) = \text{In}(y, x)$。
• 无论选哪种线程映射，读 / 写中总有一个是合并访问、另一个是非合并访问；数据须先读入临时位置（如寄存器）再写到新位置。
• 若用本地（共享）内存缓冲读写之间的数据，可重排线程映射，使读、写两个方向都合并——但要求工作组为方形（square）。此时 global mem index 与 local mem index 的对应关系被重排（如 $0,1,2,3 \to 0,4,8,12$）。
• 课件在 AMD 5870 上对比朴素（Naive）与优化（Optimized，用本地内存 + 线程重映射）两个转置 kernel，随矩阵阶数（$1024$/$2048$/$3072$/$4096$）增大，优化版耗时明显更低。

2.29 Bank Conflict（本地/共享内存）

共享内存被划分为多个 bank，访存延迟取决于同一 bank 上的冲突数：

2.30 参考资料、CUDA 与架构简史

HIP & CUDA 参考资料（课件列出）

• AMD ROCm 文档、ROCm examples；HIP CPU Runtime（header-only 库，让 CPU 执行未经修改的 HIP 代码）；Omniperf（系统性能分析工具）。
• CUDA Toolkit 文档；NVIDIA Technologies – Architectures。

CUDA「不完整」简史

NVIDIA GPU 架构简史

2.31 本讲小结

• GPGPU 的起源（Origins of GPGPU）。
• CUDA 中的一些非常基础的概念（host/device、kernel、thread/block/grid、warp、内存模型、occupancy、divergence、coalescing、bank conflict）。
• NVIDIA GPU 架构。
• 调试工具预览。
• CUDA / 架构历史回顾。

第 3 讲 · CUDA 实例：归约与分块卷积

本讲通过两个经典 CUDA 实例，把第 2 讲的编程模型（thread / block / warp、shared memory、__syncthreads()、constant memory 等）落到实处：一是并行归约（parallel reduction）——把 $N$ 个元素求和，看它如何从最朴素的 atomicAdd 一路优化到高效的树形归约（tree-based reduction），逐层消除串行化、bank conflict 与 warp divergence；二是分块卷积（tiled convolution）——在内存带宽成为瓶颈时，如何用 shared memory 分块（tiling）复用数据、处理 halo / ghost 元素，并显著提升算术访存比（arithmetic-to-global-memory-access ratio）。课件配套代码参考 NVIDIA cuda-samples 中的 reduction 示例，及 UW CSE P590（Dr. Hari Sadasivan）的讲义。

3.1 归约问题与两条基本思路

归约（reduction）指用一个可结合的二元运算（这里以加法为例）把 $N$ 个输入元素汇聚为单个结果。课件对比了两条基本思路：

block 内 / block 间同步的层次

• block 内同步（intra-block synchronization）：用 __syncthreads() 作为 block 内的 barrier，保证一轮归约的 shared memory 写对块内所有线程可见后，才进入下一轮读。
• block 间同步（inter-block synchronization）：CUDA 没有跨 block 的全局 barrier，课件给出三条实现全局同步的途径：
  • kernel 结束（kernel termination）——kernel 返回是天然的全局同步点，随后可选择：
    • 在 CPU 上做最终归约（final reduction on CPU），或
    • 在 GPU 上再启动一个归约 kernel（launch new reduction kernel on GPU）；
  • 全局内存上的原子操作（atomic operations in global memory）。

3.2 树形归约的 CUDA 实现要点

课件把一个 block 内的归约 kernel 拆成四个阶段，并给出对应的关键机制：

关键 CUDA 机制

// 动态 shared memory：block 内共享
extern __shared__ int sdata[];

// 每轮归约前的 barrier
__syncthreads();

// 启动语法：G 个 block、每 block B 个线程、
// 第三个参数是动态 shared memory 字节数
reduce<<<G, B, B*sizeof(int)>>>(in, out, n);

① 加载阶段（Load phase）

• sdata[tid] = (i < n) ? g_idata[i] : 0;——每个线程恰好写一个元素，线程之间无竞争（no contention）。
• 越界线程（out-of-bound）存入 0，因为 0 是加法的单位元（identity），不影响求和结果。

② barrier：__syncthreads()

• 每次 shared memory 写阶段之后必须调用，保证块内全部 $B$ 个（如 256 个）写对块内所有线程全局可见。

③ 归约循环（Reduction loop）

• 步长（stride）$s$ 从 blockDim.x/2 开始，每轮减半（$s \to s/2$）。
• 线程 tid 读 sdata[tid + s]——读到的总是另一个线程写入的位置。
• 每轮之后再次 __syncthreads()，确保本轮写完成后才进入下一轮读。

④ 最终写回（Final write）

• if (tid == 0) g_odata[blockIdx.x] = sdata[0];——每个 block 仅一次无原子（atomic-free）的全局写，代价很低。

__global__ void reduce(int *g_idata, int *g_odata, int n) {
    extern __shared__ int sdata[];
    int tid = threadIdx.x;
    int i   = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = (i < n) ? g_idata[i] : 0;   // ① 加载，越界填 0
    __syncthreads();                          // ②

    for (int s = blockDim.x/2; s > 0; s >>= 1) {   // ③ 归约循环
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid == 0) g_odata[blockIdx.x] = sdata[0];  // ④ 每 block 一次写
}

3.3 静态 vs. 动态 shared memory 与占用率权衡

静态分配 vs. 动态分配

内存访问模式（Memory access pattern）

• sdata[tid] += sdata[tid + s]; 每次是两次 shared 读 + 一次写。
• 所有访问都落在 block 的 48 KB 片上（on-chip）区域内。
• 归约循环期间没有全局总线流量（no global bus traffic）——这正是把数据先搬进 shared memory 的收益所在。

占用率权衡（Occupancy trade-off）

3.4 优化一：避免 bank conflict（连续寻址）

shared memory 被划分为若干 bank；若同一 warp 内多个线程在同一轮访问同一个 bank（且不是广播同一地址），硬件会把这些访问串行化，造成 bank conflict。

冲突版本：交错寻址（interleaved addressing）

• 索引 idx = 2*s*tid——随着 $s$ 增大，步长会变成 32 的倍数。
• 当 $s = 16$ 时：线程 $0..15$ 映射到 idx = 0, 32, 64, …，全部落在 bank 0 → 16 路（16-way）冲突。
• 硬件把这 16 次访问串行化 → 该轮延迟放大 16 倍。

无冲突版本：连续寻址（sequential addressing）

if (tid < s) sdata[tid] += sdata[tid + s];

• 当 $s = $ blockDim/2 时：tid=0→bank 0，tid=1→bank 1，…，tid=31→bank 31 ✓。
• 当 $s = 1$ 时：tid=0 触及 bank 0 和 1，tid=1 触及 bank 2 和 3，… ✓。
• 任何一轮都不会有两个活跃线程共享同一个 bank——冲突被彻底消除。

二维 shared 数组的 padding 技巧

• __shared__ float s[32][33];——多出的一列使每一行相对上一行偏移 1 个 bank。
• 第 $i$ 行起始于 bank $i \bmod 32$，从而消除按列访问（column-access）时的 bank conflict。
• 可用 Nsight Compute 剖析冲突：指标 l1tex__data_bank_conflicts_pipe_lsu_mem_shared。

3.5 优化二：避免 warp divergence（warp 感知寻址）

同一 warp 内的 32 个线程以 SIMT 方式执行；若一条分支使 warp 内部分线程走一条路、部分走另一条路，硬件会把两条分支串行执行，即 warp divergence（warp 分歧）。

版本 A —— 发散版（divergent）

• 判定条件 tid % (2*s) == 0。
• $s = 1$ 时：warp 0 中 {0,2,4,…,30} 活跃，{1,3,5,…,31} 空闲——活跃与空闲混在同一 warp。
• 混合的 warp 会串行执行两个分支；每一轮都发生分歧，且随 $s$ 增大越来越严重。

版本 B —— warp 感知版（warp-aware）

• 判定条件 tid < s。
• $s = $ blockDim/2 = 128 时：线程 $0$–$127$ 活跃（warp 0–3），$128$–$255$ 空闲（warp 4–7）。
• 每个 warp 要么整体活跃、要么整体空闲——不存在混合 warp。
• 空闲 warp 直接不被调度，没有串行化开销。

附加优化：最后一个 warp 的循环展开（last-warp unrolling）

• 当 $s \le 32$ 时只剩 1 个 warp 活跃，此时 __syncthreads() 已非必需（warp 内本就同步执行）。
• 可用 volatile 修饰的 sdata 指针手动展开（unroll）最后 6 次加法（对应 $s = 32,16,8,4,2,1$），省去同步开销。
• 剖析指标（Nsight）：sm__sass_average_branch_targets_threads_uniform，用于衡量分支是否统一。

3.6 分块卷积（tiled convolution）：动机与基本概念

卷积（convolution）把一个 filter（滤波器 / mask）滑过输入数组，对每个输出位置做加权求和。朴素实现里每个输出元素都要重复从全局内存读入相邻的输入元素，导致内存带宽成为瓶颈（memory bandwidth bottleneck）。课件给出的解法是分块（tiled）卷积算法，同时适用于一维（1D）与二维（2D）卷积。

输入 tile 与输出 tile（Input / Output tiles）

• 输出 tile（output tile）：由每个 block 负责计算的那一组输出元素构成。
• 输入 tile（input tile）：为计算某个输出 tile 中的所有 $P$ 个元素，所需要的那组输入 $N$ 元素。
• 由于卷积需要邻域，输入 tile 比输出 tile 更大：

  $\text{输入 tile 维度} = \text{输出 tile 维度} + 2 \times \text{filter radius}$

• 滤波器（filter）放在 constant memory（常量内存）中（读多写少、带广播缓存，适合被所有线程共享读取）。

halo 元素与 ghost 元素

• halo cells（晕圈元素）：输出 tile 边缘的输出元素在卷积时会用到相邻 tile 的输入元素，这些"借用"来的边缘输入就是 halo。课件指出：一个 block 输入 tile 的 halo 元素，同时也是相邻 tile 的内部元素（internal elements）——这正是数据可复用的根源。
• 加载输入 tile 时按三段处理：加载左侧 halo（left halo）→ 加载内部元素（internal elements）→ 加载右侧 halo（right halo）。
• ghost cells（幽灵元素）：处于输入数组边界之外的位置。计算边界处的输出时会访问这些越界元素；课件的处理是给它们填 0（见 3.9 边界条件）。

3.7 shared memory 数据复用与 1D 分块卷积核

• 观察（Observation）：同一 block 内的线程会加载一些相同的输入元素——存在大量重复的全局读取。
• 优化（Optimization）：让每个线程只加载一个输入元素到 shared memory，其余线程再从 shared memory 读取该元素，从而实现数据复用（data reuse）。

课件给出的 1D 分块卷积核有两种：一种是基础的 1D Convolution Tiled Kernel；另一种是带缓存的 1D Convolution Tiled Kernel with Caching——后者利用「halo 元素本就是相邻 tile 的内部元素」的事实，把这些 halo 交由（L2 / 常规）缓存来提供，从而只把内部元素显式装入 shared memory。

3.8 算术访存比（Arithmetic to Global Memory Access Ratio）

课件用一个 $M \times M$ 的滤波器量化 tiling 的收益。定义算术访存比为「每字节全局访问所支撑的运算量（OP/B）」。

无分块（Without tiling）

有分块（With tiling）

设输入 tile 维度为 $T$，则输出 tile 维度为 $T - M + 1$。

3.9 输入/输出 tile 尺寸不匹配与坐标映射

尺寸不匹配的挑战（Difference in Tile Sizes）

• 挑战（Challenge）：输入 tile 与输出 tile 维度不同，二者关系为 $\text{输入 tile 维度} = \text{输出 tile 维度} + 2 \times \text{filter radius}$。
• 解法（Solution）：启动足够多的线程以加载整个输入 tile 到 shared memory，然后只用其中一个子集（subset）的线程去计算并写出输出 tile。

输入分块的高层策略（Input Tiling）

• 把一整块 $N$ 的 tile 载入 LDS / shared memory：所有线程都参与加载；随后一个子集的线程在 shared memory 中使用这些 $N$ 元素做计算。
• 核心权衡：输入 tile 必须比输出 tile 更大，我们采用「把输入 tile 载入 shared memory」的策略。

处理不匹配 / 从输出坐标映射到输入坐标

• 让 thread block 的大小匹配输入 tile：每个线程加载输入 tile 的一个元素。
• 因此部分线程不参与输出计算——会出现 if 语句与控制分歧（control divergence）。
• 课件给出以 filter radius $=2$（对应 $5\times 5$ filter）为例的坐标换算：把输出坐标 $(row\_o, col\_o)$ 平移得到输入坐标 $(row\_i, col\_i)$。

int tx = threadIdx.x;
int ty = threadIdx.y;

// 输出坐标
int row_o = blockIdx.y * TILE_SIZE + ty;
int col_o = blockIdx.x * TILE_SIZE + tx;

// 输入坐标（相对输出平移一个 filter radius = 2）
int row_i = row_o - 2;
int col_i = col_o - 2;

3.10 边界条件与输出写回

加载 halo 与边界条件（Boundary Conditions）

• 加载 halo 时，若线程要加载的输入位置落在数组 $N$ 之外（即 ghost 元素），应返回 / 存入 0.0。
• 计算边界输出的线程会访问越界（out of bounds）的输入元素，这些即 ghost 元素。
• 解法：对越界的输入元素，向 shared memory tile 中存 0。

// 加载输入 tile 到 shared memory，越界填 0
if ((row_i >= 0) && (row_i < N.height) &&
    (col_i >= 0) && (col_i < N.width)) {
    Ns[ty][tx] = N.elements[row_i * N.width + col_i];
} else {
    Ns[ty][tx] = 0.0f;   // ghost 元素填 0
}

计算输出（Calculating Output）

• 如前所述，部分线程不参与输出计算：仅当 ty < TILE_SIZE && tx < TILE_SIZE 的线程子集执行卷积累加。
• 累加使用常量内存中的 mask Mc 与 shared memory 中的 Ns。

if (ty < TILE_SIZE && tx < TILE_SIZE) {
    float output = 0.0f;
    for (int i = 0; i < 5; i++) {
        for (int j = 0; j < 5; j++) {
            output += Mc[i][j] * Ns[i + ty][j + tx];
        }
    }
    // 见下方写回
}

写回输出（Writing Output）

• 部分线程不写输出：只有落在有效输出范围内的线程才写回全局内存。

if (row_o < P.height && col_o < P.width) {
    P.elements[row_o * P.width + col_o] = output;
}

2D 卷积的完整形态

• 课件最终给出两个 2D 分块卷积核：2D Convolution with Tiling and Constant Memory（滤波器置于常量内存 + shared memory 分块），以及在其上进一步利用缓存的 2D Convolution with Tiling & Constant Memory with Caches（把 halo 交给缓存、只显式装入内部元素，减少 shared memory 装载与控制分歧）。

第 4 讲 · Attention 与 FlashAttention

本讲围绕注意力机制（Attention）展开，从 RNN 到 Transformer 的演进讲起，梳理自注意力（self-attention）与 Q/K/V 三元组的信息检索视角、多头注意力（Multi-Head Attention）与位置编码（Positional Encoding），再进入本课程真正关心的并行与系统层面：标准注意力的 $O(n^2)$ 计算/显存复杂度、GPU 存储层次带来的「内存墙（Memory Wall）」问题，以及 FlashAttention 如何用分块（Tiling）、重计算（Recomputation）、核融合（Kernel Fusion）与 Online Softmax 把 HBM 访存量从 $O(N^2)$ 降到 $O(N)$。

4.1 从 RNN 到 Transformer（From RNN to Transformer）

课件先给出序列建模的发展脉络，注意力思想正是在这条脉络中逐步登场并最终「一统天下」：

4.2 什么是注意力（What is Attention?）

课件从心理学/神经科学出发引入「注意力」概念：在最一般的意义上，注意力可被描述为一种整体的警觉水平或与环境交互的能力（an overall level of alertness or ability to engage with surroundings，Attention in Psychology, Neuroscience, and Machine Learning, 2020）。在视觉神经科学中，早期神经元只对简单视觉属性（强度对比、颜色对抗、朝向、运动方向与速度、立体视差等）敏感，随着从低级到高级视觉区推进，神经元调谐越来越专门化，高级区域的神经元只对角点/交界、明暗形状线索或特定真实物体视图响应（Computational Modelling of Visual Attention, 2001）。

注意力机制在深度学习中的三个关键节点：

• Bahdanau Attention（2014）—— 注意力的兴起：为神经机器翻译引入「联合学习对齐与翻译（Jointly Learning to Align and Translate）」。
• Luong Attention（2015）—— 提升效率：提出「全局注意力（Global Attention）」，简化并改进了 Bahdanau 的方法。
• Transformer（2017）—— Attention Is All You Need：完全抛弃循环（RNN）与卷积（CNN），仅依赖注意力（relying solely on Attention）。

由此衍生出多种注意力形式：自注意力（Self-Attention）、多头注意力（Multi-Head Attention）、稀疏注意力（Sparse Attention）等。

4.3 Transformer 架构背景（Background: Transformer）

课件从高层视角（a high-level look）概览 Transformer 的编码器（Encoder）与解码器（Decoder）结构：

• 编码流（Encoding Flow）：自注意力层中各路径相互依赖（dependent paths）；前馈层中各路径相互独立（interdependent → 可并行）。编者补充：前馈层逐位置独立，因此天然可并行，这是 Transformer 相比 RNN 更易并行的关键。
• 高层直觉（Self-Attention at a High Level）：以句子「The animal didn't cross the street because it was too tired」为例，自注意力能让「it」正确地关注到「the animal」，从而消解指代。

4.4 Q/K/V 矩阵：把自注意力看作信息检索（Q/K/V as Information Retrieval）

自注意力可以理解为一次信息检索（Information Retrieval）。每个 token 被变换成三个向量，构成三元组（The Triplets）：

机制（The Mechanism）分三步：

• 相似度（Similarity）：计算当前 token 的 Query 与所有其他 token 的 Key 之间的兼容性/相似度。
• 权重（Weights）：把得分经过 Softmax 得到注意力权重。
• 输出（Output）：最终表示是所有 Value 的加权和（weighted sum of the Values）。

4.4.1 注意力的计算流水线（Computation Pipeline of Attention(Q,K,V)）

整个计算可归纳为：先由输入分别算出 Query、Key、Value 矩阵，再做打分与加权（Score and Weight）。用矩阵形式写出标准注意力：

4.5 多头注意力（Multi-Head Attention）

• 多头注意力扩展了模型关注不同位置的能力（expands the model's ability to focus on different positions）。
• 它为注意力层提供了多个「表示子空间」（representation subspaces）：每个头在不同子空间中独立做注意力。
• 各头的输出最终被拼接并压缩回单个矩阵（condense these matrices down into a single matrix）。
• 直觉示例仍是「it → the animal」：不同的头可以捕捉不同的关注关系。

4.6 位置编码与残差（Positional Encoding & Residuals）

• 位置编码（Positional Encoding）：由于注意力本身不含顺序信息，需加入位置编码向量，其取值遵循特定模式（follow a specific pattern）。课件给出维度为 $4$ 的玩具示例，以及 Tensor2Tensor 的真实实现——对 $20$ 个词（行）、嵌入维度 $512$（列）的位置编码可视化。
• 残差与层归一化（Residual Connections & Layer Normalization）：每个子层外包残差连接 + LayerNorm；课件示例给出由 $2$ 层堆叠编码器与解码器构成的 Transformer。

4.7 解码器与最终输出层（Decoder, Final Linear & Softmax）

• 解码起点：以起始符号 [BOS]（Begin of Sequence）开始，重复解码过程直到产生特殊结束符号。
• 最终线性层 + Softmax 层（The Final Linear and Softmax Layer）：把解码器输出映射到词表上的概率分布。

4.7.1 仅解码器结构（LLM — Decoder-only）

大语言模型广泛采用仅解码器（Decoder-Only）结构，专为自回归（Autoregression）而生：根据「此前出现的一切」预测下一个 token。其优势：

• 简单高效（Simplicity and Efficiency）：单一训练目标——预测下一个 token；无需昂贵稀缺的平行语料（parallel datasets）；可高效扩展到数万亿 token。
• 与生成任务完美契合（Perfect Match for Generative Tasks）。
• 灵活性与涌现能力（Flexibility and Emergent Abilities）。
• 可扩展性与产业验证（Scalability and Industry Validation）。

课件同时对比了 Decoder-Only vs Encoder-Decoder 两种范式。

4.8 注意力的复杂度（Complexity of Attention）

在 Transformer 自注意力中，我们要计算所有 token 两两之间的注意力得分：

• Query $Q$ 是 $(n,d)$ 矩阵（$n$ 个 token，每个维度 $d$）。
• Key $K$ 是 $(n,d)$ 矩阵。
• 由此形成一个 $n\times n$ 的注意力得分矩阵 $QK^\top$。
• 再乘以 $(n,d)$ 的 Value $V$。

4.9 内存墙问题（The Memory Wall Problem）

课件的一句话总结："Attention is all you need until you need memory."（在你需要内存之前，Attention 就是你需要的一切。）自注意力的灵活性代价高昂——其复杂度随输入序列长度二次增长。Transformer 在两处「撞墙」：

• KV Cache 爆炸（KV Cache Explosion）：推理时缓存的 Key/Value 随序列长度累积。
• 带宽瓶颈（The Bandwidth Bottleneck）：数据在存储层次间反复搬运，带宽成为瓶颈。

4.9.1 标准注意力的数据流（Data Flow of Standard Attention）

课件以 $N=4096,\ d=128,\ \text{FP16}$ 为例，把标准注意力拆成三步，逐步分析其 HBM（高带宽显存）访存：

• Step 1 计算得分：$S=QK^\top/\sqrt{d}$ —— 从 HBM 读 $Q,K$，把 $N\times N$ 的 $S$ 写回 HBM。
• Step 2 应用 Softmax：$P=\text{softmax}(S)$ —— 读回 $S$，写出 $N\times N$ 的 $P$。
• Step 3 计算输出：$O=PV$ —— 读回 $P$ 与 $V$，写出 $O$。

可见标准实现把庞大的 $N\times N$ 中间矩阵反复写回又读回 HBM，累积出巨大的 Total HBM Traffic（总 HBM 访存量）。

4.9.2 「内存墙」的本质（The "Memory Wall" of Standard Attention）

• 问题：标准注意力复杂度 $O(n^2)$。
• 痛点：显存消耗随序列长度 $n$ 平方增长，长上下文处理时频繁 OOM（Out-of-Memory）。
• 关键之问：我们是否真的需要显式生成并存储那个巨大的 $N\times N$ 注意力矩阵？（Do we truly need to materialize and store that massive $N\times N$ attention matrix?）

4.10 GPU 存储层次（GPU Memory Hierarchy）

理解内存墙必须理解 GPU 存储层次的速度落差（约 10 倍）。以 A100 为例：

SRAM（L1/Shared）比 HBM 快约 10$\times$，但容量极小（几百 KB 量级），因此整张 $N\times N$ 矩阵（如 $N=4096$）根本放不进 SRAM，必须分块处理。

4.10.1 注意力的根本问题（Root Problem of Attention）

4.11 FlashAttention 核心思想：更少的搬运（Core Idea: Less Movement）

FlashAttention 的设计哲学是「优化数据搬运，而不仅仅优化计算」（Optimize Data Movement, Not Just Computation）——更少搬运，更快速度（Less Movement, More Speed）。它依赖两大技术支柱：

4.11.1 分块概念（Tiling Concept）

• 为什么要分块？ GPU SRAM（如 A100 的 L1/Shared Memory）极快但容量极小（通常几百 KB）；传统 $N\times N$ 矩阵（如 $N=4096$）远超 SRAM 容量，必须切成能放进缓存的「tiles」。
• 如何工作？ 把输入 $Q,K,V$ 划分为块（tiles），用嵌套循环把这些块逐片从 HBM 载入 SRAM 做局部计算。
• 性能收益：部分点积（partial dot-products）与 Softmax 累加完全在 SRAM 内完成，从而免去把 $N\times N$ 中间结果写回慢速 HBM。

4.11.2 核融合（Kernel Fusion）

4.12 SRAM 调度与前向实现（SRAM Scheduling & Forward Pass）

调度策略（Tiling）：$Q$ 按行切分；$K$ 与 $V$ 按列切分。采用双重循环（Dual-Loop）执行：

• 外层循环（Outer Loop）：把一个 $K$、$V$ 块载入 SRAM。
• 内层循环（Inner Loop）：把一个 $Q$ 块载入，做点积，并更新输出。
• 结果：中间注意力矩阵在 SRAM 内被即算即用（computed and consumed entirely within SRAM），从不触及 HBM。

4.13 Online Softmax（在线 Softmax）

4.13.1 分块的挑战：Softmax 的全局依赖（Global Dependency of Softmax）

为什么不能直接在各个块上独立计算 Softmax？

• 分块困境（Tiling Dilemma）：Softmax 需要该行的全局最大值 $m$（用于数值稳定）与全局求和 $L$（用于归一化）。
• 冲突（Conflict）：处理第一个块时，我们并不知道剩余块的最大值或总和；按标准逻辑必须先扫描整行。
• 过渡之问：如何在只看到部分数据（partial chunk）的情况下，计算或「预测」出正确的 Softmax 输出？

4.13.2 安全 Softmax 回顾（Review of Safe Softmax）

• 必要性：直接计算 $e^{x_i}$ 在 $x_i$ 很大时（如 $x_i=1000$）会导致浮点数上溢（Floating Point Overflow）。
• 解法：对向量中每个元素先减去最大值 $m=\max(x_i)$，再做指数化。

4.13.3 Online Softmax 算法（Online Softmax Algorithm）

Online Softmax 的思想是：边遍历分块、边维护「运行最大值」与「运行归一化和」；当新块带来更大的最大值时，对此前已累加的结果按比例缩放校正，从而只看到部分数据也能得到正确的 Softmax 输出。课件以 $S=[x_0,x_1,x_2,x_3]$、$V=[v_0,v_1,v_2,v_3]$ 为例，配合公式推导（Formula Derivation）展示这一过程。

// Online Softmax（伪代码，编者补充）
m = -inf; l = 0; o = 0
for x_k, v_k in stream(S, V):
    m_new = max(m, x_k)
    l = l * exp(m - m_new) + exp(x_k - m_new)
    o = o * exp(m - m_new) + exp(x_k - m_new) * v_k
    m = m_new
return o / l

4.14 本讲小结（Conclusion）

• Attention：以 Q/K/V 信息检索视角理解自注意力，$\text{Attention}(Q,K,V)=\text{softmax}(QK^\top/\sqrt{d_k})V$。
• Transformer 架构：编码器/解码器、多头注意力、位置编码、残差与 LayerNorm、Decoder-Only 的 LLM 范式。
• 内存墙（Memory Wall）：标准注意力的二次复杂度 $O(n^2)$ 带来访存受限与显存爆炸。
• FlashAttention：通过 Tiling（分块）与 Recomputation（重计算）（并配合 Kernel Fusion 与 Online Softmax），把 HBM 访存流量从 $O(N^2)$ 降到 $O(N)$。
• 核心概念：Self-Attention、Multi-Head Attention、Positional Encoding。

第 5 讲 · 卷积、非线性与算子融合

本讲（Lecture 05: Convolution, Nonlinear, and Operator Fusion，罗国杰老师主讲）以卷积算子（convolution operator）为主线，围绕「如何在 GPU 上把卷积算得又对又快」展开：先剖析朴素 7 层循环卷积的低效，引入 Roofline 模型与算子的 compute-bound / memory-bound 分类；再系统对比三种卷积实现——直接卷积（Direct Convolution）、Img2Col + GEMM、Winograd 卷积；随后转入 Element-wise（逐元素）与 Reduction（归约）两类基本算子的并行特性、原子操作陷阱与树形归约；最后落到归一化与激活函数（Normalization & Activation）以及全讲的高潮——算子融合（Operator Fusion）：动机、类型、代价、反例，并以 Fused Softmax 作为完整实战范例。

5.1 2D 卷积算子的基本原理

• 输入（Inputs）：输入特征图（input feature maps）、权重矩阵（weight matrix）、偏置（bias，可选）。
• 输出（Outputs）：输出特征图（output feature maps）。
• 核心运算（Core operation）：乘累加（Multiply-Accumulation, MAC）。
• 朴素实现（Native Implementation）：7 层嵌套循环（7-level loop）——遍历输出通道、输入通道、输出高、输出宽、卷积核高、卷积核宽（外加 batch）逐一做 MAC。

5.2 朴素 7 层循环的低效

课件以常规卷积算子（conventional convolution operator）为例指出：朴素 7 层循环在数学上完全正确，但在结构上是灾难性的（Mathematically correct. Structurally disastrous.）——它无法利用现代硬件的数据复用与访存局部性，算力利用率极低。其浮点运算量为：

5.3 性能分析工具：Roofline 模型

课件先给出两个 GPU 核心指标（GPU core metrics）：

• 算力 / 计算能力（Compute Power, FLOPS）：GPU 每秒能执行的浮点运算次数。得益于更多的计算核心与更高的时钟频率，算力增长非常迅速。
• 内存带宽（Memory Bandwidth）：GPU 每秒能从全局显存（global memory）读出或写入的数据总量。
• 内存墙（Memory Wall）：内存带宽的增速远远落后于算力的增速，导致 GPU 计算核心常常在空闲等待慢速内存供数。

5.4 算子类型：Compute-bound vs Memory-bound

编者补充 这一分类是本讲的「诊断表」：卷积属 compute-bound → 用 GEMM/Winograd 提升算力利用；而激活、归一化、逐元素运算属 memory-bound → 用算子融合减少访存。

5.5 Img2Col 变换：把几何转成算术

Img2Col（image-to-column）的核心思想是「Converting geometry to arithmetic」——把空间卷积（spatial convolution）问题转化为通用矩阵乘法（General Matrix Multiply, GEMM）。

• 如何构造矩阵？把每个卷积滑动窗口内的输入元素展开（unfold）成一列，所有窗口拼成「输入特征矩阵」；卷积核展平成「权重矩阵」。二者做一次 GEMM 即得到输出。
• 为何这样做？为了借用厂商高度优化的 GEMM / BLAS 库：Intel MKL、NVIDIA cuBLAS、Ascend catlass 等。
• GEMM 为何快？厂商 GEMM 通常做了三类深度优化：① 微内核优化（micro-kernel optimization）；② 内存打包与布局排序（memory packing & layout ordering）；③ 缓存分块（cache blocking）。

5.6 Winograd 卷积

Winograd 属于快速滤波算法（fast filtering algorithm），通过增加加法、减少乘法来降低卷积的乘法次数（对乘法代价远高于加法的硬件尤其划算）。

5.6.1 一维例子 $F(2,3)$

• 常规算法：$2\times 3 = 6$ 次乘法。
• Winograd 用中间量 $m_1\sim m_4$ 改写结果，只需 4 次乘法 + 11 次加法 + 2 次算术右移（right arithmetic shift）。
• 其中 $g_1+g_2$ 可复用（reused），进一步减少运算。

5.6.2 矩阵形式与二维推广

5.6.3 tile 越大越好吗？$F(4\times4,3\times3)$

• $F(2\times2,\ 3\times3)$ 高效；$F(4\times4,\ 3\times3)$ 理论上可将乘法减少约 6.00$\times$。
• 但代价急剧上升：① 变换矩阵（transform matrix）复杂得多；② 变换本身的复杂度会吞掉减少乘法带来的收益（overwhelm any savings）；③ 变换矩阵元素的数值幅度（magnitude）随 tile 增大而增大 → 数值稳定性变差。

5.7 三种卷积实现对比（本讲小结）

5.8 逐元素运算（Element-wise）

• 核心逻辑：对集合中的每个元素独立施加相同变换，元素之间互不干扰（no interference between elements）。
• 并行性（Parallelism）：各计算分支 $f(x)$ 之间无依赖，可分布到成千上万个线程并发执行。
• GPU 友好（GPU Friendly）：与 SIMD 执行模型完美匹配。

5.9 归约运算（Reduction）

• 核心逻辑：通过一个二元运算（binary operation，如 Max）把序列中的多个元素合并为更少的结果。
• 数据流（Data flow）：朴素形式下，每个元素的计算依赖前一个元素的累积结果——存在数据依赖。

5.9.1 Element-wise vs Reduction 核心差异

5.10 复习：GPU 内存层次结构

• 线程层级：线程（Threads）被组织成块（Blocks），块被分派到 SM（Streaming Multiprocessor）上执行。
• 内存可见性：全局内存（Global Memory）对所有线程可见；共享内存（Shared Memory）仅对同一 Block 内的线程可见。

编者补充 越靠近计算核（寄存器、Shared Memory）带宽越高、容量越小；算子融合的本质就是让数据尽量停留在上层（寄存器 / SRAM），避免下探到 HBM。

5.11 归约中的数据竞争（Data Competition）

考虑一个基本归约：全局内存中的值为 $3$，两个线程都想把它加 $1$。

• 场景 1（正确串行）：线程 1 读 $3$、加 $1$、写回 $4$；线程 2 读 $4$、写回 $5$。结果正确。
• 场景 2（竞争 / race）：线程 1 读 $3$，线程 2 也读 $3$；两者都在 $3$ 上操作，各自写回 $4$。结果错误（丢失一次更新）。

5.11.1 解决方案：原子操作（Atomic Operations）

• 核心逻辑：保证执行期间该内存位置不被其他操作访问（锁定 read-modify-write 全过程）。
• 局限：atomicAdd 虽然保证结果正确，但性能通常很差，只适合稀疏访问（sparse access）。

5.11.2 原子操作的性能瓶颈

• 串行化（Serialization）：GPU 原本强大的并行处理能力退化为单核串行性能。
• 指令开销更高：原子操作需要在底层锁定内存位置，以防 read-modify-write 过程被其他 Worker 打断。

5.12 归约的并行化：树形结构

• 若运算满足结合律（associative property），例如 $a+(b+c)=(a+b)+c$，即可用树形结构（tree structure）并行化。
• 先让不同 Worker 计算局部归约结果（local reduction），再由 Combiner 函数合并。
• 串行朴素归约为 $O(N)$；树形并行归约深度降为 $O(\log N)$。

5.12.1 归约的优化目标

• 目标：达到 GPU 峰值性能（optimal peak GPU performance）。
• 瓶颈：访存瓶颈（Memory Access Bottleneck）——算术强度极低，每加载一个元素只做 1 次运算（1 FLOP per Element）。
• 结论：归约的优化目标是最大化带宽利用率（maximize bandwidth utilization）。

5.13 Triton 示例：向量加法与求和

• 向量加法（Vector Addition）：最简单的并行模式——每个线程处理一个或多个独立数据点，线程间无需通信。
• 求和（Sum）：归约相对复杂，需要树形并行结构；但 Triton 把底层树形交换逻辑封装进 tl.sum，可直接调用。

@triton.jit
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offs < n
    x = tl.load(x_ptr + offs, mask=mask)
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(out_ptr + offs, x + y, mask=mask)   # element-wise，线程完全独立

# 求和：Triton 自动实现树形归约
s = tl.sum(x, axis=0)

5.14 海量数据的归约与 Kernel 分解

问题：不同线程分布在不同 SM 上，如何共享中间结果？理想做法是「全局同步（global synchronization）」——若 CUDA 支持 Global Barrier，则所有 block 在第一轮归约后自动同步，再递归归约到单个值。

5.14.1 解决方案：Kernel 分解（Kernel Decomposition）

• 既然无法在 Kernel 内同步，就用结束 Kernel 来同步——Launch 的结束天然是一个全局同步点。
• 多级归约（Multi-level reduction）：
  • Level 0：启动大量 Block，各自归约一部分数据，把结果写回全局内存。
  • Level 1：启动少量 Block，读取 Level 0 结果做二次归约。

5.15 激活函数（Activation Function）

• 定义：为神经网络引入非线性能力（nonlinear capabilities），使模型能拟合复杂的函数映射。
• 以 GELU 为例：因包含超越函数（transcendental functions，如 erf），其计算强度较高。
• 关键洞察：GELU 独立执行时受内存带宽约束（memory-bound）；但当它与 GEMM 融合时，复杂的 erf 计算能显著降低 kernel 延迟——因为可复用已在片上的数据。

5.16 LayerNorm

• 地位：Transformer 等现代神经网络的核心组件，核心思想是对单个样本的所有特征做归一化。
• 计算依赖：计算均值与方差需要遍历整行数据，本质上是一次 Reduction 操作。

5.17 从计算图到执行：算子的生命周期

课件把一个算子从「计算图」走到「硬件执行」的过程分为三层：

5.17.1 Kernel Launch 优化：CUDA Graph

• 把一系列 GPU 操作（Kernels、Memcpys）录制成静态拓扑图（static topology）。
• 捕获阶段（capture）：记录所有 kernel 的执行顺序、参数、显存地址、调度路径。
• 重放阶段（replay）：CPU 只需下发一条超轻量指令即可触发整张图在 GPU 上自动流转——把「通用 GPU kernel」的多次启动开销压到「最小 GPU kernel」。

5.18 算子间的数据传递

• 默认模式（Default）：Conv 的结果必须先写回全局显存，Scale 再从显存读出——产生一次多余的全局显存往返。
• 算子融合（Operator Fusion）：逐元素运算可与其他算子无缝集成，显著减少访存与 kernel 启动开销。

5.19 计算图全局优化：算子融合

深度学习中的典型融合场景（typical scenarios）：

5.19.1 算子融合的优点

5.19.2 算子融合的问题

5.19.3 反例一：XLA 的贪心合并

• 场景：输入的强制布局为 $\{3,2,1,0\}$，而算子（Custom-call）的强制布局为 $\{1,0,2,3\}$。
• 计算 Add 时，XLA 发现两个输入布局不同，为对齐必须插入一个 Copy 操作。
• 为减少 kernel 启动开销，XLA 强行合并 Add 与 Copy；但不一致的 I/O 布局破坏了内存合并访问（memory coalescing），使访存变得随机紊乱，最终导致 9 倍性能退化。

5.19.4 反例二：Fused Adam

• PyTorch 自动启用的 Fused AdamW 在 A40 GPU 上出现剧烈性能下降。可能原因：
  • Fuse 模式使用的寄存器数量约为标准模式的两倍。
  • Fused AdamW 按张量（per-tensor）触发，导致 Kernel 启动次数过多。

5.20 完整范例：Fused Softmax

性能瓶颈（未融合时）：对 $M\times N$ 的输入，共读取 $5MN+2M$ 个元素、写入 $3MN+2M$ 个元素。频繁的全局内存读写造成严重的带宽浪费。

5.20.1 融合的基本思路

• Load：所有线程协作并行地从 Global Memory 读入数据，写入 Shared Memory。
• Sync & Reduce：在 Shared Memory 中做多轮迭代，全程在片上高速执行，不消耗显存带宽。
  • CUDA：每一步用 __syncthreads() 同步；
  • Triton：自动同步。
• Write：只把最终结果写回 Global Memory。

5.20.2 Softmax Kernel 实现要点

• SRAM 驻留（SRAM Residency）：整行数据加载后常驻片上，直到计算完成才写回。
• 算子融合：把 Max、Sub、Exp、Sum 合并进一个 Kernel；内部自动实现并行树结构（parallel tree structure）。

@triton.jit
def softmax_kernel(out_ptr, in_ptr, n_cols, BLOCK_SIZE: tl.constexpr):
    row = tl.program_id(0)
    offs = tl.arange(0, BLOCK_SIZE)                 # 连续索引，自动生成
    x = tl.load(in_ptr + row * n_cols + offs, mask=offs < n_cols, other=-float('inf'))
    x = x - tl.max(x, axis=0)                       # Max + Sub（融合，片上）
    num = tl.exp(x)                                 # Exp
    den = tl.sum(num, axis=0)                       # Sum（树形归约）
    tl.store(out_ptr + row * n_cols + offs, num / den, mask=offs < n_cols)

5.20.3 Softmax Host 端的三项优化

5.20.4 幕后优化一：连续寻址（Continuous Addressing）

• 交错寻址（Interleaved addressing）可能引发 bank conflict（存储体冲突）。
• 连续寻址（相邻线程访问连续数据）能最大化内存带宽。
• 相比 CUDA，Triton 通过显式调用 tl.arange()自动生成连续索引，更难写出错误代码。

5.20.5 幕后优化二：算法级联（Algorithm Cascading）

• 让单个线程处理多组数据，增大每线程工作量，从而隐藏延迟并减少 Block 启动开销。
• 若 BLOCK_SIZE 远大于物理线程数，Triton 编译器会在每个物理线程内自动执行循环，处理 $\text{BLOCK\_SIZE} / \text{线程数}$ 份数据。

第 7 讲 · Tensor Cores 的演进

本讲主题为 NVIDIA Tensor Core 的演进（The Evolution of NVIDIA Tensor Cores），由罗国杰老师主讲（Intro-to-PDC，Spring 2026）。本讲从性能第一性原理（Performance First Principles）出发，引出矩阵乘加（MMA，matrix multiply-accumulate）这一核心运算，然后沿 Volta → Turing → Ampere → Hopper → Blackwell 五代架构，逐代梳理 Tensor Core 在 MMA 尺寸、支持精度、访存机制与编程接口上的演进。全程贯穿一个核心观点：计算便宜、搬数据昂贵，几乎所有 Tensor Core 的优化都与访存改进（memory access improvements）相伴。

7.1 性能第一性原理（Performance First Principles）

• Amdahl 定律（Amdahl's Law）：增加算力资源只能加速并行部分（scaling compute resources only speeds up the parallel portion）。
• 算术强度（Arithmetic intensity）：用 roofline 分析（roofline analysis）衡量，即计算量与访存量之比。
• 数据移动是原罪（Data movement is the cardinal sin）：
  • 计算便宜、搬数据昂贵（Computation is cheap and data movement is expensive）。
  • 几乎所有 Tensor Core 优化都与访存改进相伴（almost all tensor core optimizations are in conjunction with memory access improvements）。

7.2 PTX 编程模型（PTX Programming Model）

一个 PTX 程序描述由大量 GPU 线程执行的内核函数（kernel function）。经典的 SIMT 组件如下：

SM 将每个线程映射到一个标量处理核心（scalar processing core），即 CUDA core。

7.3 CUDA 存储层次（CUDA Memory Hierarchy）

线程层次与存储层次一一对应：

7.4 MMA 概述（MMA Overview）

给定矩阵，乘加指令（multiply and accumulate, MMA）计算：

• $A$ 是 $M \times K$ 矩阵；
• $B$ 是 $K \times N$ 矩阵；
• $C$ 与 $D$ 是 $M \times N$ 矩阵。

矩阵形状记作 mMnNkK 或 $M \times N \times K$。数据流动路径为 SMEM $\leftrightarrow$ RF $\leftrightarrow$ Tensor Core。

7.5 Tensor Core 演进总表（Tensor Core Evolution）

编者补充

[自注] 上表严格照抄课件第 6 页。课件此表未列出 Blackwell 一行，其内容见后续 7.11–7.15。课件标题中提到 Blackwell 支持 FP8/FP4/微缩放（MX）等更低精度，但演进总表本身只覆盖到 Hopper。

7.6 Volta（2017）：Tensor Core 的起源

为什么 NVIDIA 要加入 Tensor Core（Why NVIDIA Added Tensor Cores）：

• 受 Google TPU 启发（Motivated by Google TPU）。
• 对传统基于 SIMT 的 MMA，指令开销远高于运算本身（instruction overhead is much higher than operations themselves）。
• NVIDIA 设计了半精度矩阵乘加指令 HMMA（half-precision matrix multiply and accumulate）。
• Volta 的 Tensor Core 在 Volta 架构开发非常晚期才加入，仅在流片（tape out）前几个月——这印证了 NVIDIA 调整架构的速度之快。

第 1 代 Tensor Core：Warp 作用域 MMA（Warp-scoped MMA）。即由一个 warp 协作完成 MMA。课件此处（第 9–11 页）另附了 Tensor Core 架构的推测图（Deduced，来源见参考链接）。

7.7 Turing（2018）：低精度与 ldmatrix

• 支持更低精度 MMA（INT8 / INT4 / Binary）。
• 引入 ldmatrix 访存指令。相较传统 LDS（load shared）逐线程加载的诸多低效：
  • LDS 逐线程执行，只写入线程本地寄存器（thread-local registers）。
  • SMEM 采用 $32$ 位访问粒度；LDS.b16 每线程浪费 $16$ 位。
  • 矩阵转置加载需两条 LDS.b16 指令，进一步增加数据浪费。
  • 在 warp 层面，大量并行 LDS 指令增加指令开销（instruction overhead）。
  • Bank 冲突（Bank conflicts）进一步降低有效加载效率。

ldmatrix 的 PTX 语法：

ldmatrix.sync.aligned.shape.num{.trans}{.ss}.type r, [p];

.shape = {.m8n8};
.num   = {.x1, .x2, .x4};
.ss    = {.shared};
.type  = {.b16};

7.8 Ampere（2020, A100）：异步数据拷贝（Async Data Copy）

• 可异步地在共享内存与全局内存间拷贝数据（asynchronously copy data from shared memory to global memory）。
• 对比 Volta：在 Volta 中，线程必须先把数据从全局内存（GMEM）加载到寄存器，再存入 SMEM。
• Ampere 直接从 GMEM 取数并存入 SMEM（可选经过 L1），从而释放更多寄存器给 MMA 指令（freeing up more registers for MMA instructions）。
• PTX 指令：cp.async。

7.9 Hopper（2022）：Tensor Memory Accelerator（TMA）

• TMA 是一个专用硬件单元，加速 GMEM 与 SMEM 之间大批量数据的异步传输（bulk asynchronous copies）。
• CTA 内单个线程即可发起一次 TMA 拷贝，而非 warp 级发起（not warp-wide）。
  • TMA 让线程腾出来执行其他独立工作，例如地址生成（address generation）。
  • PTX 指令：cp.async.bulk。
• 但对小请求，TMA 加载因地址生成开销（address generation overhead）而延迟高于常规异步拷贝。
  • NVIDIA 建议对大数据拷贝使用 TMA，以摊薄开销（amortize the overhead）。
  • 例如 LLM 推理中，TMA 不适合按小块加载 KV cache 的场景，但当每块是 $16$ 字节的倍数时效果良好。

TMA 优势（TMA Advantages）——通过拷贝描述符生成地址（Address Generation via Copy Descriptor）：

• 简化编程模型：TMA 接管步长（stride）、偏移（offset）与边界（boundary）的计算。
• 无缝扩展到 scale-up 域（NVLink）：只要交换了内存句柄（memory handler），即可无差别地对远端 rank 进行 LD/ST（load/store to/from remote ranks agnostically）。

7.10 Hopper：Warpgroup 级异步 MMA（wgmma）与线程块簇

Warpgroup 级异步 MMA（Warpgroup-level Asynchronous MMA, wgmma）：

• 一个 warpgroup（$4$ 个 warp）协同执行一次 MMA 操作。
• 支持更宽的形状（wider range of shapes），例如 m64nNk16，其中 $N$ 可取 $8$ 到 $256$ 之间 $8$ 的倍数。
• 矩阵加载/存储缓冲（Matrix load and store buffers）：
  • 操作数矩阵 $A$ 可位于 RF 或 SMEM；
  • 操作数矩阵 $B$ 只能经 SMEM 访问；
  • warpgroup 中所有线程共同在寄存器中持有输出矩阵。

线程块簇 / 分布式共享内存（Thread Block Cluster / Distributed Shared Memory, DSMEM）：

• 新增线程层次级别线程块簇（Thread Block Cluster），映射到同一图形处理簇（Graphics Processing Cluster, GPC）内物理相邻的一组 SM。
• 在 CTA（映射到 SM）与 grid（映射到整个 GPU）之间提供更细粒度的控制。
• 分布式共享内存（DSMEM）：可用 load、store、atomic 操作直接访问簇内其他 SM 的共享内存。
• DSMEM 将线程块间的数据交换加速约 $7\times$，有利于频繁交换数据的访存受限（memory-bound）场景（课件举例：LLM 推理？）。

TMA Multicast（TMA 多播）：

• TMA 可将数据从 GMEM 加载到线程块簇中多个 SM 的 SMEM，由多播掩码（multicast mask）指定目标。
• 减少 L2 缓存流量，进而减少 HBM 流量（reduces L2 cache traffic and subsequently reduces HBM traffic）。

异步事务屏障（Asynchronous Transaction Barrier）：

• 不再只统计线程到达（thread arrivals），还统计 SMEM 存储的事务数（transactions）。
• 实现带隐式同步的数据交换（data exchange with implied synchronization）。

7.11 Blackwell：专用 Tensor Memory（TMEM）

• 引入专用张量内存（Dedicated Tensor Memory, TMEM），以缓解 TC 运算中极端的寄存器压力（extreme register pressure）。
• 每个 SM 上 TMEM 有 $128$ 行（lanes）$\times$ $512$ 列、每格 $4$ 字节，共 $256$ KB，恰好等于一个 SM 上寄存器文件的大小。
• 受限的访存模式（Restricted memory access pattern）：需要一个 warpgroup 才能访问整个 TMEM，warpgroup 中每个 warp 只能访问特定的一组 lane。
  • 优点：硬件设计者可减少访问端口数量，节省芯片面积（saving chip space）。
  • 缺点：尾处理（epilogue）操作需要一整个 warpgroup 来完成。

7.12 Blackwell：第 5 代 MMA（tcgen05.mma）

• 第 5 代 MMA（PTX 中的 tcgen05.mma）彻底不再用寄存器持有矩阵，操作数改为驻留在共享内存与 Tensor Memory（SMEM 与 TMEM）。
• 由单线程发起，而非 warpgroup 发起（issued by single-thread instead of warp-group）。

7.13 Blackwell：CTA Pair 与 MMA.2SM

• 线程块簇中的两个 CTA，若其在簇内的 CTA rank 仅最后一位不同（如 $0$ 和 $1$、$4$ 和 $5$），则组成一个 CTA 对（CTA Pair）。
  • 一个 CTA 对映射到一个纹理处理簇（Texture Processing Cluster, TPC），TPC 由两个 SM 组成，多个 TPC 再组成一个 GPC。
• MMA.2SM 用 $2$ 个 SM 协同完成一次 MMA 操作：
  • 在 CTA 对层面，由 CTA 对中leader CTA 的单个线程发起 MMA.2SM；
  • 通过（更快的 TPC?）DSMEM 通信；
  • 这种共享降低了 SMEM 的容量与带宽需求（reduces both SMEM capacity and bandwidth requirements）。

7.14 Blackwell：原生细粒度量化（Native Fine-grained Quantization）

• 量化粒度从逐张量（per-tensor）→ 逐组（per-group），缩放（scaling）与 MMA 融合（fused），并支持微缩放格式（microscaling, MX formats）。
• Hopper 用 $22$ 位累加 FP8，限制了动态范围（limiting dynamic range）。DeepSeek 采用 CUDA core promotion 做 FP32 累加，并与 MMA 计算重叠（overlapping with MMA compute）。

编者补充

[自注] 课件标题写作「Amphere」，规范拼写应为 Ampere；此处已按规范术语呈现。上文各代所支持的数据类型（FP16/BF16/TF32/FP64/INT8/INT4/Binary/FP8/FP4/MX）均取自课件演进总表与 Blackwell 各页原文，未额外编造未出现的算力数字。

7.15 回顾：吞吐翻倍，延迟未降（Looking Back）

这正好呼应了本讲开头「数据移动是原罪」的主线：Tensor Core 越快，越需要访存机制（ldmatrix、cp.async、TMA、DSMEM、更大的 staging buffer）持续跟进，否则计算单元会因搬数据而挨饿。

7.16 一些延伸思考（Random Thoughts）

• 基于 tile 的编程与微架构演进（Tile-based programming）：cuTile 支持从 Blackwell 起步；Rubin 微架构是否会包含 cuTile 专属方案（cuTile-exclusive recipes）？
• 计算与通信，谁是性能杀手？
  • 在 Hopper 上训练 DSV3（EP64），$380$ TFLOPS，计算 : EP 通信 $= 1:1$；
  • 在 GB200 上，$1000$ TFLOPS，计算 : EP 通信 $= 4:1$，MFU 与 Hopper 相近。
• 如何快速把 kernel 评测并接入端到端流水线？始终以 PyTorch 作为张量接口，或提供 Python 绑定，以便 (1) 原生管理内存/流/同步/D2H·H2D/分布式编排；(2) 作为插件无缝接入既有框架（如 DeepEP、DeepGEMM、TE、cuDNN、nvshmem、Triton 等）。

7.17 补充：CUDA 性能剖析工具（CUDA Profiling Tools）

7.18 参考资料（References）

• SemiAnalysis：NVIDIA Tensor Core Evolution — from Volta to Blackwell。
• 微信公众号技术解读文章两篇（Tensor Core 架构推测与演进）。
• NVIDIA CUDA 官方文档、Hopper 白皮书（whitepaper）与 GTC 教程。

第 8 讲 · AI 编译器

本讲（Lecture 08: AI Compilers，主讲罗国杰 Guojie Luo）系统介绍深度学习编译（Deep Learning Compilation）的整体流程、中间表示（Intermediate Representations, IR）与优化方法（Optimization Methods）。课件大纲（Outline）分为四大块：①深度学习编译流程（含 torch.compile）；②中间表示（MLIR、LLVM、TVM Relay IR、Triton…）；③优化技术；④硬件映射与电子设计自动化（Hardware mapping / EDA，如把数据流图映射到 CGRA 数据流架构，以及 FlagTree 框架与实践）。本讲内容较多，围绕"从高层框架到硬件指令"的编译主线逐层展开。

8.1 通用机器学习编译流程（General ML Compilation Flow）

课件给出的通用编译流程（General Machine Learning Compilation Flow）由若干"抽象表示阶段（Abstract Representations）"串联而成，每个阶段各有关注点，主要处理张量（Tensor）与张量函数（Tensor function），并把自身表示部署（deploy）给后续阶段。整体流水线如下：

图片来源：blog.rijuyuezhu.top / mlc.ai/summer22-zh（课件标注）

8.2 计算图（Computation Graph）

8.2.1 为什么需要计算图（Why Computation Graph?）

• 是反向传播（Back-Propagation, BP）链式法则（Chain Rule）的表示。
• 便于自动优化与 lowering（Easy for Automatic Optimization and Lowering）。
• 支持多层表示（Multi-layer Representation）。

图中元素（Elements in Graph）：

• 节点（Node）= 张量（Tensor）。
• 边（Edge）= 函数（Function），可能是若干原始算子（primitive operators）的组合。

8.2.2 静态计算图（Static Graph, TensorFlow）

特点：在计算之前完整构建整张图（Construct the Graph Completely before Computation）。流程为：定义变量 → 定义函数 → 执行前向计算 → 得到反向计算结果。

8.2.3 动态计算图（Dynamic Graph, PyTorch）

特点：边算边建（on-the-fly）——在当前计算之后再构建前驱图（Construct Precedent Graph after Current Computation）。流程为：定义变量 → 执行前向计算 → 边算边构建计算图 → 执行反向传播。

示例来源：geeksforgeeks.org（Dynamic vs Static Computational Graphs）

8.2.4 PyTorch 中计算图如何构建

• 为需要梯度的张量（Tensor Requiring Gradient）创建 AutogradMeta 对象：为存储梯度分配内存空间；分配累加器（accumulator）把后续链上的所有梯度相加。
• 边算边追踪（Trace On-the-fly）计算过程：为每个运算创建 grad_fn 对象；每个运算持有指向其前驱（precedent）的指针。

视频来源：pytorch.org/blog/computational-graphs-constructed-in-pytorch

8.2.5 优化后的计算图（Optimized Computation Graph）

• 它仍是一张用函数操纵张量的计算图，只是用更高效的算子表示函数，例如 kernel 融合（Kernel fusion）、等价重排序（Equivalent reordering）、平衡（Balancing）。
• 是机器学习编译中的关键阶段（Critical Stage）：处于早期、优化空间大（large optimization space），直接影响内存、效率等；可利用各种图操作与优化技术。

8.3 torch.compile：现代计算图方式

torch.compile 体现了动静态权衡（Dynamic-Static Tradeoff）：PyTorch 与 TensorFlow 都扩展了各自的计算图（torch.fx 与 TensorFlow 2.0），既易于开发又能为执行高效优化。其主思想是边基准测试边优化（Benchmark and Optimization On-the-fly）——只要当前计算不改变，就优化它；从而做到首次运行后更快的编译与执行（faster compilation and execution after first run）。

torch.compile 三大关键组件（Key Components）

图片来源：pytorch.ac.cn/blog/accelerated-pytorch-inference

8.4 编译器实例（一）：Glow / ELL

8.4.1 Glow（Graph Lowering compiler）

• 开源社区项目，2018 年由 Facebook 提出，是面向硬件加速器（hardware accelerators）的机器学习编译器与执行引擎。
• 采用两阶段中间表示（Two-phase IR），为多种嵌入式与服务器级硬件生成专门调优的机器码。
• 采用提前编译（Ahead-of-time, AOT）：最小化运行时开销，以节省磁盘空间、内存、启动时间等。
• 设计为高层框架的后端（backend）：输入是计算图，输出是支持多种硬件类型的原生机器码（可编译为含 native code 的目标文件 object files）。
• 两级 IR 分工：
  • 高层 IR（High-level IR）：让优化器做领域特定优化（domain-specific optimizations）。
  • 低层、基于指令的仅地址 IR（instruction-based address-only IR）：做内存相关优化——指令调度（Instruction scheduling）、静态内存分配（Static memory allocation）、拷贝消除（Copy elimination）。

8.4.2 ELL（Embedded Learning Library）

• 微软（Microsoft）的嵌入式学习库，用于把软件部署到资源受限平台（resource constrained platforms），如小型单板计算机。
• 是面向嵌入式智能的交叉编译器（cross-compiler）：编译器本身运行在你的笔记本/台式机上，生成的机器码运行在单板计算机上。

8.5 编译器实例（二）：TVM

8.5.1 TVM 内部流程（Internal Flow）

TVM 从 PyTorch / TensorFlow / ONNX 接入，内部流程为：

• Relay：函数式（Functional）、静态类型（statically typed）的中间表示（IR）。

8.5.2 目标与自动调度（Goal & Auto-scheduling）

目标：自动把张量运算（如 matmul、conv2d）转成高效的代码实现。

def fuse_add_exp(a: ty.handle, c: ty.handle) -> None:
    A = tir.match_buffer(a, (64,))
    C = tir.match_buffer(c, (64,))
    B = tir.alloc_buffer((64,))
    with tir.block([64], "B") as [vi]:
        B[vi] = A[vi] + 1
    with tir.block([64], "C") as [vi]:
        C[vi] = exp(B[vi])

8.5.3 TVM 的图能力示例

• Relay 算子融合：把多个算子融合成单个 fused operation，再针对目标专门优化。
• 多设备划分（Partition）：把网络切分到多个设备上运行。
• 为不同数据布局生成高效代码（data layouts）：
  • NHWC：Number（样本数）、Height、Width、Channels，TensorFlow 常用。
  • NCHW：Number、Channels、Height、Width，PyTorch 常用。编者补充：txt 原文此处写作"NHCW"，按标准布局命名应为 NCHW
  布局显著影响内存访问（memory access）。

图片来源：Sameer Farooqui, How to use Apache TVM to optimize your ML models

8.6 中间表示（Intermediate Representations, IR）

IR 为变换（transformations）提供更易操作的表示。IR 聚焦图优化，对之前的输入格式无感（ignorant to previous input format），从而：减少人工优化实现细节的工作量；利用社区成果（如 MLIR、Triton…）；易于部署到新兴硬件。

8.6.1 不同 IR 层次（一）：基于 DAG

• 表示程序的计算流（computation flow）：高层——直接对接计算图；可以是符号式（Symbolic）或命令式（Imperative）。
  • 符号式：关注用户定义表达式的语义（semantics）。
  • 命令式：关注实际的操作序列（sequence of operations）。
• PyTorch.fx：符号式，把优化留给其他模块。
• MXNet：命令式，也支持符号式。

8.6.2 符号式 vs 命令式（Symbolic vs Imperative）

8.6.3 不同 IR 层次（二）：基于 Let-binding

• 把一个变量与某张量或表达式关联（associate）；确定计算的作用域（scope）：确定所有张量的产出，并汇集这些变量的映射（map）。
• 可实现常量折叠（constant folding）与公共子表达式消除（common-subexpression elimination, CSE）等平凡优化。
• TVM Relay 承担了这类角色的一部分。

8.6.4 不同 IR 层次（三）：领域特定与混合型

• DS-IR（领域特定 IR）：目标代码 lowering 优化，是最后一步，利用硬件特定抽象。
• 混合型（Hybrid，如 TVM Relay）：结合多个阶段，兼具 DAG 型与 let-binding 型 IR 的优点。
• Tensor Comprehensions 可视为 DS-IR：对张量运算采用领域特定抽象与数学表达式（多面体模型 polyhedra model），为 GPU 等加速器实现优化与高效代码生成。

8.7 MLIR

8.7.1 MLIR 简介：可复用的编译流水线

• MLIR 是一条可复用的编译流水线（Reusable Compilation Pipeline）：可复用的 IR 与 Pass；序列化/反序列化系统；多个方言（dialect）可共存。
• 方言翻译图（Dialect Translation Graph）中的分类：
  • Payload：实际运算（$+\ -\ \times\ /$）。
  • Structure：控制流（Control Flow）、函数（Function）。
  • Tensor：高层张量。
  • Buffer：低层指针。

8.7.2 洞见：方言混合（Dialect Mixing）

例：PyTorch 生成神经网络，要表示如下计算，会混用多个方言：

Lowering 过程：先降控制流（linalg $\to$ affine $\to$ scf $\to$ cf），再降实际算子（替换 math、arith $\to$ LLVM IR）。方言与 lowering 方法模块化可复用：loop tiling 时无需关心具体计算，计算优化时无需关心函数是什么。

8.7.3 洞见：多级优化（Multi-level Optimization）

• IR 优化可在不同阶段进行：linalg 能检测出某矩阵被转置两次；affine 能做loop tiling；arith 能复用常量表达式。
• 多级优化的价值：在合适的层次而非最低层的 LLVM-IR 上优化——linalg 容易发现"矩阵转置了两次"，而 LLVM-IR 很难发现。

8.7.4 MLIR 基本结构（Tree-Based Structure）

• Operation（操作）：单个操作，接受操作数（Operands）并产生结果（results）。
• Block（块）：基本块，一串操作的列表；块有块参数（block parameters），跳转到块时可传入不同参数。
• Region（区域）：一般是一个函数或一个循环；Op 可以包含 Region，如 if 或 for。

8.7.5 MLIR 使用要点（输入/输出、codegen、Op 结构）

• 输入输出：项目配置用 CMake；为 MLIRContext 导入所需方言；用 parseSourceFile 读入；用 op->print 输出（调试用）。
• 代码生成（Code Generation）：创建 OpBuilder；所有 Op 只能由 builder 创建（MLIR 管理内部类型）；用 builder.getXXXType 或 XXXType::get 构造类型；在 builder 中设置 InsertionPoint 以填充 Block。
• Op 结构：包含 Operand（输入）、Attribute（编译器已知的常量）、Result（结果）、Region（内部区域）。Attribute 是编译期已知的值，Operand 只在运行时已知；Operand 有 Type，Type 也可作为 Attribute。
• Op 的类型变换：所有 Op 有统一表示 Operation（含 Operand/Result/Attribute）；每个 Op 都是指向 Operation 的指针，但各自定义数据解释方式（如 Add 把首个操作数当 Lhs，Constant 无操作数）。
• Downcast / Eq / Hash：LLVM 提供 cast、dyn_cast；Operation* 的相等是指针相等而非值相等，可作为 Map 的 key；Type 与 Attr 同样按指针 casting 处理（专用指针/通用指针/Value）。

8.7.6 MLIR 图结构（Graph Structure）

• Op 之间的连接：Value 来自 Result 或 Argument；Op 有自己的 Result；Operands 是指向其他 Result 的指针；OpOperand 是存储 operand 的结构。
• Use Chain（使用链）：Value 会串起所有使用链——Use = OpOperand，User = Op；Value 的主要操作是 replaceAllUse（不自行实现 replace，细节见 tutorial）。
• 迭代与操作：可递归迭代子节点、迭代基本块、插入/删除 Op、为 Op 查找父节点。

8.7.7 用 TableGen 定义 IR

• MLIR 用 TableGen 定义新 IR：自动生成每个参数的 Getter/Setter（Generate API automatically）。
• 每个 Operand 是带不同约束（constraints）的 Value，产生不同类型。
• 约束（Constraints）：参数可含 attribute；约束可引用 mlir/IR/CommonAttrConstraints.td 与 mlir/IR/CommonTypeConstraint.td（TypeConstraint / AttrConstraint）。
• 其它特性（见 tutorial）：verifier、emitError（把 error 绑定到 Op，emitWarning 类似）、assemblyFormat（更优雅的输出）、Variadic、customBuilder、extraFunction。

8.7.8 IR Trait 与 Pattern Rewrite

• IR Trait：有效减少 IR 的工作量。
  • SideEffectInterfaces（内存副作用）：Pure 可自动 CSE 与 DCE；MemRead / MemWrite / MemAlloc / MemFree 表示其它内存副作用。
  • InferTypeOpAdaptor（自动类型推断）：用 builder 时可跳过返回类型；SameOperandsAndResultType 表示输入输出同类型。
  • FunctionOpInterface：函数接口（见 tutorial）。
• SideEffect Trait：在定义 Op 时加上 trait，即可用于 CSE、DCE Pass；canonicalize（标准化）、cse（去除公共子表达式），如 ./toy-opt -cse、./toy-opt -canonicalize。
• Pattern Rewrite（模式重写）：Pattern 匹配 IR 的一个子图，Rewrite 用新子图替换；许多 IR 操作都可看作模式重写，如算术优化 $x\times 2 \to x + x$、表达式 lowering（用低层 Op 重写高层 Op）。只需定义 Pattern，MLIR 负责调度——用贪心算法（greedy）自动匹配替换。

8.7.9 Pattern 定义与使用

• matchAndRewrite：返回 success（匹配）/ failure（无法匹配）。
• PatternRewriter：具体修改器，可记录并撤销替换——在部分替换后 undo，或当 Pattern 收益不足时放弃。例：把自定义 AddOp 替换为 arith 的 AddOp。
• ConversionTarget：指定哪些 Op 合法/非法；RewritePatternSet：指定一组 Pattern。三种接受替换的设置：
  • partialConversion：只要替换合法就保留。
  • fullConversion：不断改写非法输入直到全部合法。
  • greedyPatternRewrite：无需 target，贪心地尽可能多地尝试替换。

8.7.10 方言转换（Dialect Conversion）与 Materialization

• 不仅转换 Op，也转换 Type。TypeConverter 是 MLIR 提供的方言转换工具：addConversion（添加转换规则，起关键作用）、addTargetMaterialization（把 SourceType 转成 TargetType 的代码块）、addSourceMaterialization（反向转回）。
• MLIR 用Conversion Pattern 做类型转换（含一个 adapter）：自动尝试转换 Op 的 operand 类型；转换失败则保持不变并把结果存入 adapter；一般可用 adaptor 中的值完成转换。
• 类型转换依赖特殊 Op unrealized_conversion_cast，表示尚未实现的类型转换：输入类型错时插入 TargetMaterialization，输出类型错时插入 SourceMaterialization；冗余转换（如 A→B→A）会被移除或报错。
• 自定义 Materialization 场景：如自定义 float32 需转为 float8，希望尽量少调用函数；若 float32 作为参数且他人要用该函数则转换不可避免。MLIR 有专门 Pass/Pattern 移除 unrealized conversion。
• 多层 Pattern 混合：可直接从高层 IR 降到最低层（原 →Arith→LLVM，现直接到 LLVM）；可直接把 pattern 加入 pattern set，MLIR 总会尝试匹配更多 pattern 并校验结果是否满足 target。

8.8 Triton

8.8.1 Triton 简介与编译流程

• 用类 NumPy 风格写 GPU 代码（Write GPU code in a Numpy-like style），相比 CUDA 更简洁。
• 编译流程：Dialect $\to$ Triton IR $\to$ LLVM IR $\to$ PTX。
• 通过对 block 的操作暴露实例内并行（intra-instance parallelism）：每个 kernel 实例在单线程上执行 block 操作；Triton 按 block 变量的生存期（live range）分配 shared memory；自动按实例与 SM 并行化操作（分析每个操作的迭代空间再划分 partition）。
• Triton JIT 把输入当作指针而非张量（用指针算术而非索引运算符）：对内存访问的低层控制很重要（尤其块稀疏 block-sparse）；同时抽象掉所有并发相关问题（concurrency）。

@jit
def add(X, Y, Z, N):
    pid = program_id(0)
    idx = pid * 512 + arange(512)   # block of indices
    mask = idx < N
    x = load(X + idx, mask=mask)
    y = load(Y + idx, mask=mask)
    store(Z + idx, x + y, mask=mask)

# 启动（无 thread-block 概念）
grid = (ceil_div(N, BLOCK),)
add[grid](x, y, z, x.shape[0])

• 每个 "program" 运行在一个 CUDA block 上，操作大小为 2 的幂的定长张量；程序员无需考虑 threads/warps。注意：RBLOCK 是张量大小，不是 block 中的线程数。

@triton.jit
def sum_kernel(in_ptr, out_ptr, xnumel, rnumel, RBLOCK: tl.constexpr):
    xindex = tl.program_id(axis=0)
    rindex = tl.arange(0, RBLOCK)
    offset = xindex * rnumel + rindex
    data = tl.load(in_ptr + offset, mask=rindex < rnumel)
    res = tl.sum(data, axis=0)
    tl.store(out_ptr + xindex, res)

8.8.2 为何尽量用 Triton 而非 CUDA

• 朴素（naive）CUDA 代码会错过优化，例如向量化加载（vectorized loads）。
• 线程通信对程序员和编译器都更难推理。
• CUDA 模板里每个线程的工作被写死（set in stone），束缚了编译器的手脚。

8.8.3 Triton 编译器三段

• 大多数 Triton IR 算子继承自内建 MLIR 方言：Arith（addf/addi/constant/cmpf/cmpi…）、Math（exp/sin/cos/log/absi/absf/sqrt）、StructuredControlFlow（for/if/while/yield/condition），从而减少到 LLVM 的工作量。

8.8.4 Autotuner

• 一个小组件，根据执行时间找最佳配置（best config）。
• Triton只关心 tile 大小（normal / micro / nano 三个层级），比 TVM 简单得多。
• 一行代码即可开启优化（Optimize with a single line of code）。

来源：openai.com/index/triton；Inside Triton Compiler, Peter Bell

8.9 TileLang：面向 AI Kernel 的分块编程模型

8.9.1 三层编程接口与编译流程

编译流程：Tile Program（计算的高层规约）→ 带 Tile 库的 Tile Program → 带线程原语的 Tile Program → IRModule → 源代码生成（C / CUDA / HIP / LLVM…）。

8.9.2 基于 Tile 的编程模型

• Tile 声明——Tile 作为一等对象（First-Class Objects）：一个 tile 表示一块有形状的数据，可被 warp、线程块或等价并行单元拥有和操纵；A、B 缓冲区按分块（tiled chunks）读取。执行上下文包含线程块索引（bx、by）与线程数，便于 TileLang 自动优化。
• 显式硬件内存分配（Explicit Hardware Memory Allocation）：暴露面向用户的 intrinsics 把 tile 缓冲区放到硬件内存层级：
  • T.alloc_shared：分配片上快速存储（on-chip），对应 NVIDIA GPU 的 shared memory。
  • T.alloc_fragment：在 fragment memory 中分配累加器，对应 寄存器文件（register files）。
  • T.clear（初始化硬件缓冲区）、T.fill（几乎等同 clear）、T.copy（global 与硬件专属内存间的数据传输）、T.Parallel（数据赋值也可并行执行）。

8.9.3 以数据流为中心的 Tile 算子

• 抽象出一组 Tile 算子供用户使用或自定义；每个 Tile 算子须实现两个关键接口：
  • Lower：定义高层 Tile 算子如何 lowering 为低层 IR。
  • InferLayout：确定与该算子关联的内存与循环布局。
• 预定义算子：gemm（高度优化的通用矩阵乘）、reduce（跨维度聚合的灵活高效归约）、atomic（并行环境下对 shared/global 内存安全更新的原子操作）。

8.9.4 调度注解与原语（Schedule Annotations & Primitives）

8.9.5 后端组成与自动调优

• 前端层：JIT——用 TileLang DSL 写 kernel，用 @tilelang.jit 装饰器编译。
• 编译流水线：从高层 IR 到优化 IR 的两阶段变换。
• 代码生成：TVM 后端——借 TVM 生成目标专属设备代码（CUDA / HIP / Metal）。
• 运行时系统：编译好的 kernel 被缓存，经 adapter 层执行。
• 自动调优系统（Autotuning）：探索配置空间寻找最优参数。三步：用带保留优化参数的 Tile Language 实现目标程序 → 通过手工搜索或 Carver 自动生成候选配置 → 并行编译并基准测试候选配置以选出最佳。（可调参数如 block_M/block_N/block_K、num_stages、thread_num、enable_rasterization。）
• 优化效果：课件展示在 H100 上的 MLA 解码与 Flash Attention 性能。

来源：arxiv.org/pdf/2504.17577v2；tilelang.com；github.com/tile-ai/tilelang

8.10 计算图优化（Graph Optimizations）

优化发生在不同层次的 IR 上。按 openMLSys 与课件划分为节点级、块级、数据流级、量化相关与硬件相关等层次。

8.10.1 节点级优化（Node Level）

8.10.2 块级优化（Block Level）

• 代数化简（Algebraic Simplification）：用更简形式替换复杂表达式。
• 算子下沉（Operator Sinking）：把计算移到更靠近其消费者的位置，改善内存访问模式、减少中间存储。
• 算子融合（Operator Fusion）：把多个张量运算合并为单个 fused operation，直接减少中间存储。

8.10.3 算子融合专题（Operator Fusion）

是硬件（GPU）优化中最有效的方法之一：减少中间内存存储、最小化内存传输、改善缓存局部性、减少 kernel 启动开销。

扩展—聚合—重构流程（Expansion-Aggregation-Reconstruction）用于获得更大的融合探索空间：

• Expansion（扩展）：把运算展开成基本（原始）算子。
• Aggregation（聚合）：把所有算子纳入更大的算子组合。
• Reconstruction（重构）：在原始与扩展算子上探索融合机会。

8.10.4 数据流级优化（Data flow-level）

• 公共子表达式消除（CSE）：识别并消除冗余计算——若表达式 $E$ 之前已算且未变，则在图中用其结果替换 $E$ 的所有实例。
• 死代码消除（DCE）：识别并移除结果/效果未被使用的代码段，通常在其它图优化之后进行。
• 布局变换（Layout Transformations）：优化图中张量的数据布局，增强内存访问与缓存利用率。

8.10.5 量化相关优化（Quantization-related）

• Rescale 节点折叠与合并：Rescale 算子调整数据值的 scale/range（ONNX 的 QuantizedLinear、Glow 的 RescaleQuantized）；可折叠进前一个数值计算节点，也可与其它操作数结合以提高效率——如与 MAX 操作结合，把 max 两侧归一到统一 scale，用更少比特实现更高硬件效率。
• 常量量化：把 Quantize(Constant) 替换为一个更新后的简单常量。
• Profile 引导量化（Profile-Guided Quantization）：预测网络各阶段可能的数值范围。

8.10.6 硬件相关优化（Hardware Specific）

• Kernel 自动调优（Auto-Tuning）：为特定计算 kernel 搜索并选择最佳配置。
• 内存规划（Memory Planning）（对优化效果起关键作用）：内存复用（reuse）、内存对齐（alignment）；张量压缩/解压缩（标量量化 Scalar / 向量量化 Vector）；利用内存层级（调 batchsize 匹配 cache、利用高带宽内存 HBM）；内存池化（Memory pooling）——用 CPU/其它设备内存扩容，并发数据交换（swapping）而不影响前台计算。

8.11 计算算子调度（Computation Operator Scheduling）

两个关键问题（Two Key Questions）：①函数中每个值何时、在何处计算？②数据存在哪里？主要手段：

// 沿 m、n 各按因子 32 tiling；沿 k 按因子 4 split 并 reorder
for (m.outer: 0..32)
 for (n.outer: 0..32) {
  for (m.inner.init: 0..32)
   for (n.inner.init: 0..32)
     C[...] = 0f32
  for (k.outer: 0..256)
   for (k.inner: 0..4)
    for (m.inner: 0..32)
     for (n.inner: 0..32) {
       cse_var_3 = n.outer*32
       cse_var_2 = m.outer*32768 + m.inner*1024
       cse_var_1 = cse_var_2 + cse_var_3 + n.inner
       C[cse_var_1] += A[...] * B[...]
     }
 }

8.11.1 基于随机变换的搜索

• 在随机变换上搜索：随机生成多种调度并选合适者；跨多进程并行基准测试；用代价模型（cost models）避免每次都实测；在 trace 上做进化搜索（evolutionary search）而非每次随机采样。
• 用优化后的调度/IR替换原始的，并可与融合等其它优化组合。

8.11.2 多面体模型（Polyhedral Models）

• 从调度搜索空间中自动选择算子调度；把每个循环抽象成多维空间：计算实例是点（points），实例间依赖是空间中的线（lines）。
• 通过空间变换调整每个实例的执行顺序，从而消除循环依赖、增强并行可能性。

// 原始（存在循环依赖）
for (i = 0; i < N; i++)
 for (j = 1; j < N; j++)
  a[i+1][j] = a[i][j+1] - a[i][j] + a[i][j-1];

// 变换后（沿绿色分块/虚线切分，去除依赖）
for (i_new = 0; i_new < N; i_new++)
 for (j_new = i+1; j_new < i+N; j_new++)
  a[i_new+1][j_new-i_new] = a[i_new][j_new-i_new+1]
                          - a[i_new][j_new-i_new]
                          + a[i_new][j_new-i_new-1];

8.11.3 Halide：适合多面体模型编程

• 核心思想：把算法与其调度分离（Separate the algorithm from its schedule）——程序员可改变 schedule，为同一算法表达多种可能的组织方式。
• 支持 split、reorder…；最终 lowering 到 LLVM 完成代码生成。

8.12 硬件映射与阅读材料

课件大纲还包含硬件映射与电子设计自动化（Hardware mapping / EDA）：把数据流图映射到数据流架构（CGRA），以及 FlagTree 框架与实践（详见课程后续/实践部分）。

第 9 讲 · 分布式并行：数据 / 张量 / 流水线 / 混合

当模型规模跨过单卡显存与算力上限，训练必须横向扩展（scale horizontally）：把显存与计算切分到多张 GPU、多台机器。本讲从训练显存危机出发，系统讲解四类并行范式——数据并行（Data Parallelism, DP）、张量并行（Tensor Parallelism, TP）、流水线并行（Pipeline Parallelism, PP），以及把三者组合起来的混合并行（Hybrid Parallelism）。课件顺序为 4 个子课件：09a 数据并行、09b 张量并行、09c 流水线并行、09d 混合并行。

并行的三条正交轴（总览）

来源：new-pdc09a~d（data / tensor / pipeline / hybrid parallelism）。此表整合自 09a 第 2–3 页与 09d 第 3–4 页。

9.1 数据并行（Data Parallelism, DP）

9.1.1 训练的显存危机（The Memory Crisis）

从推理过渡到训练，显存构成发生质变：

• 推理显存主要由模型权重与瞬态 KV cache 主导。
• 训练显存暴涨：需保留整张计算图（computational graph）以做反向传播，并维护高精度优化器状态（optimizer states）。单个加速器不够用，计算与显存都必须横向扩展。

训练显存来自两方面：其一是存储模型权重（通常 FP16），其二是存储优化器状态（Adam 的动量 momentum 与方差 variance，为精度通常存 FP32）。

完整显存构成（09a 第 22 页）

9.1.2 数据并行基础与 DDP

• DP 是最基础的分布式训练策略，工作在 SPMD（Single-Program Multiple-Data）范式下。
• 训练数据集被划分为互不重叠（non-overlapping）的 mini-batch。
• 每张 GPU 持有一份完整、相同的模型权重与优化器状态副本（replica）。
• 每个加速器在自己的数据分片上独立做前向与反向。
• PyTorch 中 DP 的标准模块是 DistributedDataParallel（DDP）。

梯度分桶（Gradient Bucketing in DDP）

• 为每个参数张量单独发起 All-Reduce 会造成网络拥塞（congestion）。
• DDP 用梯度分桶缓解：参数梯度被组织进可配置的内存桶（buckets）；只有当一个桶在反向过程中被填满时才触发该桶的 All-Reduce。
• 这大幅减少网络调用次数，优化带宽利用。

计算-通信重叠（Compute-Communication Overlap）

通信（红块）不在反向计算的关键路径（蓝块）上，因此可以很好地重叠。依据两条依赖关系：

1. grad_calc(layer i, gpu k) 只依赖 grad_calc(layer i-1, gpu k)；
2. all-reduce(layer i, gpu k) 依赖 grad_calc(layer i, all gpus)。

于是某层梯度算完即可 All-Reduce，同时继续算下一层梯度，通信被计算掩盖。

9.1.3 DDP 的显存冗余问题

• 标准 DDP 存在严重的显存冗余（memory redundancy）。
• 在 8 张 GPU 上训练时，集群被迫维护 8 份完整的模型权重、梯度和庞大的 FP32 优化器状态副本。
• 冗余复制消耗高带宽显存（HBM）却不带来额外算力。
• 这严格把最大可训练模型限制在单卡容量内，成为 LLM 无法逾越的壁垒。

9.1.4 ZeRO：消除 DP 的显存冗余

ZeRO（Zero Redundancy Optimizer，零冗余优化器）的核心思想：把昂贵的状态切分开，并利用 reduce-scatter 的等价性来完成分片式的同步。

ZeRO Stage 1：优化器状态分片（$P_{os}$）

• 把优化器状态（一阶 + 二阶矩，即动量与方差）跨 GPU 切分。
• 每个 worker 仍持有全部参数，在自己的数据子集上算梯度。
• 每个 worker 只负责更新一个参数分片（shard）。

执行流程：

1. 每个 worker 在自己数据子集上算出完整梯度；
2. ReduceScatter 梯度：每个 worker 拿到与自己参数分片对应的完整梯度（通信量 = #params）；
3. 每个 worker 用「梯度 + 状态」更新自己的参数分片；
4. 做一次 AllGather 同步参数（通信量 = #params）。

$K$ 为每参数优化器状态占用的字节数（FP32 动量典型为 $12$）。式中 $4$ 对应 FP16 参数（$2$）+ FP16 梯度（$2$）。

ZeRO Stage 2：优化器状态 + 梯度分片（$P_{os+g}$）

• 在 Stage 1 基础上，把梯度也跨机器分片，沿用相同思路。
• 难点：因为是数据并行，每个 worker 仍必须算出完整梯度，但绝不会同时实例化整个梯度向量。

执行流程：

1. Step 1：大家沿计算图增量式反向。
   • 1a：某层梯度一算完，立即 reduce 发送给对应的 worker；
   • 1b：一旦某梯度在反向图中不再需要，立即释放（free）。
2. Step 2：每台机器用「梯度 + 状态」更新自己的参数；
3. Step 3：AllGather 参数。

ZeRO Stage 3（即 FSDP）：分片一切

• 把一切都分片，包括参数本身！
• 沿用「增量式通信/计算」思路：在遍历计算图时按需（on demand）发送和请求参数。
• 参数/梯度被请求/发送后立即释放。
• 通信与计算重叠：AllGather 在前向进行时一次性发生，从而掩盖通信开销。
• FSDP（Fully Sharded Data Parallel）即 ZeRO-3 在 PyTorch 中的实现。

9.1.5 通信开销小结（09a 第 21 页）

9.2 张量并行（Tensor Parallelism, TP）

9.2.1 DP 的两大局限（引出 TP）

• 计算扩展受限：数据并行下 #machines $<$ batch size（且接近上限时通信开销很高），而增大 batch size 又有收益递减（diminishing returns）。
• 模型放不下：ZeRO-1 和 ZeRO-2 无法扩展显存；ZeRO-3 原理上很好，但可能很慢，且不能减少激活显存。

9.2.2 张量并行的基本概念

• LLM 绝大多数参数位于线性投影（Linear Projections）中：
  • Attention：Q、K、V 与输出（Output）投影；
  • MLP：上投影（Up）与下投影（Down）。
• 这些投影都是通用矩阵乘法（GEMM，General Matrix Multiplication），形如 $Y = XW$。
• TP 有两种把 GEMM 切到不同 GPU 的方式：按行切（row-wise）与按列切（column-wise）。

9.2.3 拆解 MLP 块（Megatron-LM 切法）

一个 MLP 块由「线性升维 $\to$ 非线性激活（GeLU）$\to$ 线性投影回 hidden 维 $\to$ dropout」组成，写作 $Y = \text{Dropout}\big(\text{GeLU}(XA)\,B\big)$ 形式。第一个权重记为 $A$，第二个记为 $B$。

• 把 $A$ 按行切（row-wise split）：需要在过 GeLU 前同步（因为 GeLU 非线性，$\text{GeLU}(X_1 A_1 + X_2 A_2) \neq \text{GeLU}(X_1 A_1) + \text{GeLU}(X_2 A_2)$）。
• 把 $A$ 按列切（column-wise split）：$A = [A_1, A_2]$，各列输出 $XA_i$ 可独立喂进 GeLU，无需同步。
• $A$ 按列切 $\to$ 输出 $Y$ 也按列切；再把 $B$ 按行切（$B = [B_1; B_2]$），做分块矩阵乘法（block-wise matmul），中途无需同步。

对 Attention 的切分同理采用交错切分（interleaving split）以最小化通信：多头注意力天然可按头分组做列切（Q/K/V 投影），输出投影按行切，前反向同样各一次 all-reduce。

9.2.4 何时用 TP、TP 的局限

• 局限：TP 只切 GEMM；LayerNorm 与 Dropout 仍被复制（replicated）。
• 因此每个 TP 分区上仍有完整的激活显存开销 $[\text{batch\_size}\times\text{seq\_len},\ \text{hidden\_dim}]$。

为解决这一冗余，Megatron 提出序列并行（Sequence Parallel）[MLSys'23]：把 LayerNorm/Dropout 沿序列维（seq_len）切分，与 TP 组合，进一步降低复制的激活显存。编者补充 09b 第 16 页仅列出「Sequence Parallel [MLSys'23]」标题，未展开正文；此处按其定位补充一句说明。

9.3 流水线并行（Pipeline Parallelism, PP）

9.3.1 跨节点扩展的动机

• TP 计算高效，但数学上必须同步 All-Reduce，因此不适合跨节点（inter-node）扩展——以太网/InfiniBand 延迟会灾难性拖垮 GPU 利用率。
• 要把深层 LLM 跨多节点扩展，系统必须纵向切分（slice longitudinally）模型。
• PP 把连续的层块分配到不同的物理节点上（K 个 stage）。激活值与部分梯度在相邻 stage 间来回传递。

通信模式为点对点（P2P）Send/Recv，只依赖激活值。

9.3.2 朴素流水线的气泡问题

朴素 PP 采用 Forward-then-Backward（F-then-B）时间线：每个 GPU 大部分时间都在空转（idling），等待反向传播回传——这段空闲即气泡（bubble）。

解决思路：Batch split（拆 batch）让 GPU 忙起来。

GPipe（Google, 2019）

• 把全局 mini-batch 划分为 $M$ 个更小的 micro-batch（微批）。
• 为最大化 TFLOPS，增大全局 batch size 以塞进更多 micro-batch 填满流水线。
• micro-batch 越多，气泡占比越低。

气泡率公式 编者补充

9.3.3 为什么用 PP

• 省显存（相比 DDP）；
• 通信性质好（相比 FSDP 和 TP）：只依赖激活、且是点对点通信；
• 因此通常在较慢的网络链路（即跨节点）上用 PP，以获得更好的显存维度扩展。

9.3.4 调度优化：F-then-B vs 1F1B

GPipe（F-then-B）的显存危机

• GPipe 严格同步：所有 $M$ 个前向必须全部完成，才能开始任何反向。
• 反向绝对需要中间激活，因此每个 stage 必须同时缓存全部 $M$ 个 micro-batch 的激活。
• 随 $M$ 增大以缩小气泡，激活显存线性增长 $O(M)$，迅速触发 OOM。

1F1B（PipeDream）

• PipeDream 提出 1F1B（One-Forward-One-Backward）调度：预热（warm-up）阶段结束后，GPU 交替做「一次前向、立即一次反向」。
• 优先执行反向，改变缓存数据的生命周期：激活被消费成梯度后立即丢弃。
• 任一 GPU 需存的最大在途（in-flight）激活严格由流水线深度 $K$ 界定，而非 $M$。
• 这把峰值激活显存与 micro-batch 数解耦，允许（在显存允许下）无限扩大 $M$。

9.3.5 1F1B 的局限与 Interleaved 1F1B

• 要让气泡率可忽略（$<5\%$），$M$ 必须极高。
• 但全局 batch size 受优化器收敛上限数学约束。
• 若全局 batch 固定，增大 $M$ 会把单个 micro-batch 缩到接近 $1$。
• 过小的 micro-batch 无法喂饱现代 GPU 的 Tensor Core，破坏算术强度（arithmetic intensity）。
• 因此需要一种不增大 $M$ 也能减小气泡的新调度。

9.3.6 进阶 PP

• Zero-bubble Pipeline（ICLR'24）：核心洞见是分离阻塞与非阻塞计算，以实现更灵活的调度。小规模用整数线性规划（ILP）同时优化显存与端到端 makespan；大规模用启发式（heuristics）。相比 Interleaved 1F1B，PP 通信次数更少，气泡相近或更少。
• DeepSeek-V3 DualPipe（2025）：采用特定的分片（sharding）策略进一步降低气泡（DeepSeek-V3, arXiv:2412.19437v1）。

9.4 混合并行（Hybrid Parallelism）

9.4.1 三轴正交与乘性缩放

• 正交轴（Orthogonal Axes）：DP（batch）、TP（层内/hidden 维）、PP（层间/深度）相互正交，可自由组合。
• 缩放墙（The Scaling Wall）：孤立范式都会失效——DP 撞显存墙、TP 撞网络瓶颈、PP 需海量 micro-batch 导致显存爆炸。
• 资源乘性缩放（Multiplicative Scaling）：总用量随各维度相乘组合（$\text{Total} = T \times P \times D$，即 TP $\times$ PP $\times$ DP 度数相乘得到总 GPU 数）。
• 协同（Synergy）：TP 减小单层显存 $\to$ 可塞下更大 micro-batch $\to$ 降低 PP 气泡率。

9.4.2 3D 并行系统架构

硬件共生（Hardware Symbiosis）：把 DP、PP、TP 嵌套（nest）起来以匹配数据中心拓扑。

• 内层循环（TP）：限制在单节点内，走 NVLink 承载繁重的 MatMul 通信；
• 中层循环（PP）：跨节点扩展，走 RDMA/InfiniBand，因 stage 边界通信量低、可容忍带宽；
• 外层循环（DP）：在逻辑上包裹 TP+PP，全局扩展以获得海量 batch 吞吐。

案例：Megatron-Turing NLG 530B

有效算力与通信代价建模（09d 第 6 页）

9.4.3 扩展到 4D / 5D：CP 与 EP

• 催化剂：超长序列（massive sequence lengths）与混合专家（MoE, Mixture of Experts）打破了 3D 并行。
• 上下文并行（Context Parallelism, CP）：把 token 跨 GPU 切分，通过 ring 协议交换 KV 状态，防止序列引发的 OOM。
• 专家并行（Expert Parallelism, EP）：用 All-to-All 路由把 token 动态派发到离散的 MoE 专家上。
• 5D 公式：$$\text{TotalGPUs} = \text{TP} \times \text{CP} \times \text{EP} \times \text{PP} \times \text{DP}$$
• 结果：击破每一个瓶颈——稠密权重、batch 吞吐、无限上下文、稀疏计算。

ZeRO 与 PP 集成的瓶颈

• 冲突：ZeRO-3 的动态 All-Gather 取参与 PP 高频的 micro-batch 切换相冲突。
• 问题：造成重复参数通信的网络风暴（network storm）。
• ZeroPP 解法：引入分块任务交错（blockwise task-interleaved）的 PP 调度：在丢弃权重前把每层的微计算批量化，大幅削减网络调用同时保住显存收益。

9.4.4 自动并行（Automatic Model Parallelism, AMP）

• 手工脆弱性：硬编码原语（如 All-Reduce）在切换新架构或非对称硬件时会失效。
• AMP：把分布逻辑抽象为一个编译器优化问题——将神经网络表示为计算 DAG、把硬件表示为拓扑图，自动搜索最优 sharding。
• 影响：研究者只写单设备代码，编译器合成 $n$ 维执行计划。

FlexFlow：SOAP 搜索空间

• SOAP 四维：Sample $=$ DP（batch 划分）、Operation $=$ PP（图节点分布）、Attribute $=$ CP（序列/图像特征划分）、Parameter $=$ TP（权重矩阵切分）。
• 搜索引擎：用 MCMC（马尔可夫链蒙特卡洛）探索庞大配置空间。
• Oracle 模拟器：毫秒级预测策略性能（而非数小时），发现比手工调优快达 3.3$\times$ 的非直觉方案。

Alpa：分层空间搜索

• 扁平 MCMC 在 100B+ 参数处失效（指数复杂度）。
• Alpa 按数据中心带宽层级重构搜索：
  • Level 1（Intra-op，节点内）：ILP 求解高带宽「device mesh」内的 TP/DP 分配；
  • Level 2（Inter-op，节点间）：动态规划（DP）求解跨远端 mesh 的 PP 分级。
• 结果：在复杂 MoE 上比 DeepSpeed 等专用手工系统快达 9.7$\times$。

其他自动并行系统

• GSPMD（编译器式 SPMD）：开发者写单设备 JAX/TF 代码 + 稀疏 NamedSharding 提示，XLA 补全 pass 前后向传播标注、自动推导所有未标注中间张量的 sharding，在 $2048$ TPU 核上达 62% 利用率。
• Unity（代数与并行联合优化）：标准编译器先融合算子再并行，可能锁死高效网络策略；Unity 通过图代换联合优化代数融合与并行，甚至会反直觉地拆开（unfuse）算子以启用更快通信，吞吐提升达 3.6$\times$。
• Galvatron（系统级混合策略搜索）：面向生产的大 Transformer，显式集成激活检查点（Activation Checkpointing）与 ZeRO offloading。两步流程：决策树按硬件显存预算剪掉不可能策略 $\to$ 动态规划逐层分配最优并行；成本模型独特地考虑计算/通信重叠带来的 GPU 性能下降。

9.5 本讲速查与自测

第 10 讲 · 通信原语、MPI 与集合通信库

本讲把大模型（LLM）训练/推理的通信基础设施自底向上讲透：先建立"点对点（Point-to-point）"与"集合通信（Collective communication）"的概念框架，逐一给出 Broadcast、Reduce、Gather、Scatter、AllReduce、AllGather、ReduceScatter、AllToAll 八大原语的语义；再引入 α-β 通信代价模型并推导各原语（尤其是 Broadcast、AllReduce）的通信量与轮数下界；然后落到编程接口层——MPI（Message Passing Interface）的编程模型与 API；最后落到工程实现层——集合通信库（Collective Communication Library, CCL），重点是 NVIDIA 的 NCCL、Baidu 的 Ring-AllReduce 与 NCCL 的双二叉树（Double Binary Tree），并说明它们如何在 PyTorch DDP/FSDP、Horovod 等深度学习框架中支撑数据并行。

10.1 通信原语（Communication Primitives）

10.1.1 AI 系统中的通信数据通路（Communication Datapath）

课件把 AI 系统的通信通路分为两大层次：

• Scale-up（节点内 intra-node）：NVLink、华为 HCCS、Unified Bus。节点内 GPU 之间通过 NVLink 或 PCIe 通信。
• Scale-out（跨节点 inter-node）：InfiniBand（IB）、RoCE。跨节点通过 RDMA NIC（如 IB 网卡）通信。

10.1.2 集合通信 vs. 点对点通信（Collective vs. Point-to-point）

• 集合通信（Collective）：通信发生在一组进程（a group of processes）之间，而不仅仅是成对进程之间。例如 AllReduce、AllGather、Broadcast。
• 点对点（Point-to-point）：
  • 双边（Two-side）：Send / Recv。
  • 单边（One-side）：Put / Get。

课件（10b）进一步给出正式定义：集合通信是涉及一组处理单元（节点）、并在这些节点全部或部分之间完成数据传输的通信；数据传输过程中可能施加归约算子（reduction operator）或其他数据变换。它是消息传递范式（message-passing paradigm）的自然延伸，通常通过库接口或语言构造暴露，MPI 就是提供丰富集合通信操作的典型接口。

10.1.3 八大集合通信原语语义

课件（10a）逐页给出 NCCL 文档中各原语的精确语义（root rank 举例为 rank 2，共 rank 0–3）：

10.1.4 常用集合通信分类（10b / 10c）

数据再分布操作又分为：有根（Rooted）——假定所有节点都知道 root 节点的身份（如 Broadcast/Scatter/Gather）；无根（Non-rooted）——无指定 root（如 AllGather/All-to-all）。

10.1.5 应用示例：矩阵-向量乘法（GEMV）

课件（10b / 10c）用矩阵-向量乘 $A\times b$ 说明集合通信如何嵌入并行算法，给出三种并行程序：行块条带（Rowwise block striped）、列块条带（Columnwise block striped）、棋盘块（Checkerboard block）。

• 行划分（Row-Partitioned）算法阶段：每个进程持有 $A$ 的第 $i$ 行与向量 $b$ → 做内积计算（Inner product computation）得到分量 $c_i$ → 用 All-gather 通信把各分量拼成完整结果向量。
• 列划分（Col.-Partitioned）算法阶段：每个进程持有 $A$ 的第 $i$ 列与 $b_i$ → 做乘法（Multiplications） → 用 All-to-all exchange 交换部分结果 → 各进程累加部分结果（Sum partial results）。
• 棋盘块在 $p$ 为完全平方数（square number）与非平方数时映射方式不同，需要仔细（carefully）安排进程网格。

10.1.6 FSDP / ZeRO：数据并行中的集合通信

课件（10a）以 PyTorch 的 DP / DDP / FSDP（Data Parallelism / Distributed DP / Fully Sharded DP）为例说明 AllReduce 的实际用途：

• DP：多个模型副本放在不同设备上；归约（reduce）在 node 0 上执行 → 负载不均衡（imbalance）。推理很简单。
• DDP（推荐）：使用优化过的 ring-based all-reduce + 分布式支持，具备鲁棒的多进程支持。
• 训练显存分析：模型参数只占训练总显存的一部分，还需考虑优化器状态（optimizer states）与梯度（gradients）。核心思想是去除冗余数据（removing redundant data），即 ZeRO（Zero Redundancy Optimizer）。

10.2 MPI（Message Passing Interface）

10.2.1 分布式内存与进程间通信

课件（10b）先给出并行化的场景：混合同构（Hybrid & Homogeneous）与混合异构（Hybrid & Heterogeneous）的分布式内存系统。进程间通信（IPC）分两类：

• 共享内存系统内部（within shared-memory）：POSIX 共享内存、POSIX 消息队列等。
• 跨共享内存系统（across shared-memory）：POSIX socket、消息传递（message passing）、RPC（配合 protobuf 等）、REST API 等。

此外还有分布式共享内存（Distributed Shared Memory, DSM）（K. Li & P. Hudak, 1989 的共享虚拟内存一致性工作）以及 PGAS 数据访问模型下的 NVSHMEM / OpenSHMEM。

10.2.2 MPI 的历史与标准

• 1980 年代末：各厂商各有专用库。
• 1989：Oak Ridge 国家实验室开发 PVM（Parallel Virtual Machine）。
• MPI 标准演进：1992 开始制定 → 1994 v1.0 → 1997/2008/2009 v2.0/2.1/2.2 → 2012/2015 v3.0/3.1 → 2021/2023 v4.0/4.1 → 进行中 v5.0。
• 今天：MPI 是主导的消息传递库标准。实现有 MPICH2、Open MPI 等。

10.2.3 编程模型与最小 API 集合

MPI 采用消息传递模型（Message-Passing Model）。最小例程集合如下：

Hello World 示例（10b）：

#include <mpi.h>
int main(int argc, char *argv[])
{
    int npes, myrank;
    MPI_Init(&argc, &argv);
    MPI_Comm_size(MPI_COMM_WORLD, &npes);
    MPI_Comm_rank(MPI_COMM_WORLD, &myrank);
    printf("From process %d out of %d, Hello World!\n",
        myrank, npes);
    MPI_Finalize();
    return 0;
}

10.2.4 编译与运行

• 编译：mpicc -o foo foo.c、mpic++ -o bar bar.cpp。mpicc/mpic++ 只是 wrapper——用 mpicc --show 可看到它其实展开成 gcc 加上一堆 include / lib 路径与 -lmpi -lopen-rte -lopen-pal 等链接选项。
• 运行：mpirun -np 4 foo 或 mpiexec；如 mpirun -np 8 ./a.out。
• 指定主机：使用 --hostfile 与 --host 选项。

10.2.5 SPMD 与 MPMD

• SPMD（Single Program Multiple Data）：所有进程运行同一份程序，按 rank 分支执行不同逻辑。mpirun -np 8 ./hello。经典写法：rank 0 用 MPI_Send 向 rank 1..size−1 发送消息，其余 rank 用 MPI_Recv 接收（消息 tag=11、类型 MPI_CHAR）。
• MPMD（Multiple Program Multiple Data）：不同进程运行不同程序。mpirun -np 1 ./master : -np 7 ./slave（master 发送、slave 接收）。

Hostfile 示例（slots 决定各主机分到多少进程）：

host_ada       slots=2 max_slots=8
host_barbara   slots=2 max_slots=8

• mpirun -hostfile <file> -np 3 ./hello → host_ada 上 2 个、host_barbara 上 1 个进程。
• -np 4 → 各 2 个；-np 5 → host_ada 3 个、host_barbara 2 个。
• -np 17 → 报错：系统中没有足够的 slots（not enough slots available）。

10.2.6 MPI vs. Socket API；进程启动与阻塞语义

课件（10b）对比裸 socket 与 MPI：用 socket 收发要自己写 while(count>0){ n=write(fd,buf,count); count-=n; buf+=n; } 这样的循环并做特殊错误处理；MPI 则把这些封装起来并额外提供：

• 进程启动与关闭（startup / shutdown）：一个 manager 进程用 ssh 启动其他进程，或联系守护进程（daemons）启动（更可扩展 more scalable）。
• 阻塞 / 非阻塞收发：阻塞 MPI_Send/MPI_Recv；非阻塞 MPI_Isend/MPI_Irecv 配合 MPI_Test*/MPI_Wait*。
• 集合通信（collective communications）。

10.2.7 组与通信子（Groups & Communicators）

MPI 用 communicator 界定通信范围（默认 MPI_COMM_WORLD），rank 是进程在 communicator 内的编号。构造子通信子的流程：

• 用 MPI_Group_incl 从全局 group 里选出子集形成新 group；
• 用 MPI_Comm_create 为新 group 创建新 communicator；
• 用 MPI_Comm_rank 确定在新 communicator 中的 rank；
• 用任意 MPI 消息传递例程在其中通信；
• 结束时用 MPI_Comm_free、MPI_Group_free 释放（可选）。

10.2.8 α-β 通信代价模型（Cost Model）

建模假设（本讲分析专用）：所有节点通过通信网络连通；节点每次只做单端口（single port）通信（同一时刻至多参与一次通信操作——可以是单向 send/recv、双向电话式（telephone-like）收发、或同时向两个不同节点收发的全双向 sendreceive）；通信介质同构且全互连（homogeneous and fully connected）。

10.2.9 Broadcast 的下界与算法

• MST（Minimum Spanning Tree）算法：把节点集二分成两个大致相等的子集；从 root 向不含 root 的子集里某节点发 $x$（该节点成为局部 root）；在两个子集内递归广播。总代价：$$\lceil \log_2 p \rceil\,(\alpha + n\beta)$$它在 $\alpha$ 项上达到下界，但 $\beta$ 项差一个对数因子。MST 构造的是二项生成树（binomial spanning tree）。（名称是 misnomer：同构介质下所有广播树都是生成树。）
• 流水线（Pipelining）：$n$ 很大时把消息切成 $k$ 块、每块 $n/k$。在"路径（path）"广播树上：首块经 $p-1$ 轮到达节点 $p-1$，其余 $k-1$ 块再 $k-1$ 轮，总代价 $(p+k-2)(\alpha + n\beta/k)$。可用于固定度数的树（二叉树、Fibonacci 树），但对非常数度数树（二项树、星形图）无益。
• 第三个下界（轮数）：$k-1+\lceil\log_2 p\rceil$——首块 $\lceil\log_2 p\rceil$ 轮加上后续 $k-1$ 轮。
• 同时生成树（Simultaneous / ESBT）：在 $p=2^d$ 的 $d$ 维超立方体上用 $d$ 棵边不相交的生成二项树（edge-disjoint spanning binomial trees, ESBT），均以节点 0 为根，第 $j$ 棵树广播第 $j, j+d, j+2d,\dots$ 块，实现最优 $k-1+\log_2 p$ 轮流水线（S. L. Johnsson & C. T. Ho, 1989）。

10.2.10 物理网络拓扑（Physical Network Topologies）

课件列出常见拓扑：Ring（环）、Hypercube（超立方体）、Torus（环面，含 2D Torus）、Fat Tree（胖树）、Dragonfly(+) 等。其中 Leiserson (1985) 证明 胖树是"通用网络（Universal Networks）"：Fat-tree $U(A)$ 能以面积 $A_U\in O(A\log^2 A)$、路由时间 $O(\lambda\log^2 A)$（$\lambda$ 为负载因子）模拟任意路由器 $A$。

10.3 集合通信库（CCL）：算法复杂度、NCCL 与实践

10.3.1 各原语的通信复杂度

课件（10c）在 α-β 模型下给出若干原语的经典复杂度公式（源：Grama et al., Introduction to Parallel Computing）：

10.3.2 AllReduce 算法（重点）

AllReduce 在深度学习中的意义（10c）：损失 $L(w)=\sum_{i\in S}\ell(x_i,y_i;w)$，整体梯度 $\nabla_w L(w)=\sum_{i\in S}\nabla_w\ell(x_i,y_i;w)$。数据并行把数据集切分 $S = S_1\cup S_2\cup\cdots\cup S_p$，第 $k$ 节点算局部梯度 $\nabla_w L(w)^{(k)}=\sum_{i\in S_k}\nabla_w\ell$，整体梯度 $\nabla_w L(w)=\sum_{k=1}^{p}\nabla^{(k)}$ —— 这个"各卡梯度求和并广播回所有卡"正是一次 AllReduce。

递归倍增 / 减半-倍增（recursive doubling / halving-doubling）：编者补充 由 10.3.1 的公式可见，用递归倍增实现的 Reduce-Scatter + AllGather 组合，$\alpha$ 项为 $2\alpha\log_2 p$、$\beta$ 项为 $2n\beta(p-1)/p$ 量级，其延迟随 $p$ 对数增长——这正是"halving-doubling AllReduce"相对 Ring 在小消息下更优的原因，与课件"双二叉树对数延迟"的结论一致。

10.3.3 MPI 与 NCCL 的算子命名对照

10.3.4 NCCL 与 CCL 生态、拓扑感知

• NCCL（NVIDIA Collective Communications Library）：现代集合通信库，提供 ncclAllReduce、ncclReduce、ncclBroadcast、ncclAllGather、ncclReduceScatter 等原语。
• 拓扑感知（topology-aware）：集合通信的实现高度依赖硬件拓扑——NVIDIA GPU 走 CLOS 拓扑，华为 NPU 走全互连（Full-mesh）拓扑，节点内经 NVLink/PCIe、跨节点经 IB/RoCE，库据此选择 Ring / 树 / 递归倍增等不同算法与路径。
• 与 MPI 的关系：编者补充 NCCL 只做 GPU 间集合通信，不提供进程启动/rank 管理；实践中常与 MPI（负责进程启动、rank 分配）或 PyTorch 分布式后端配合使用。MPI 是通用消息传递标准，NCCL/gloo 是面向深度学习、拓扑感知的集合通信实现。
• 在深度学习框架中的使用：PyTorch DDP 默认使用 ring-based all-reduce（GPU 上即 NCCL 后端）；FSDP（ZeRO-3）用 Reduce-Scatter + All-Gather；Uber Horovod 采用 Baidu 的 Ring-AllReduce。

互动自测

以下题目覆盖第 1–10 讲全部考点，支持单选与多选，自动判分。题库见 app.js 中的 QUIZ_BANK。

考前速查

把整门课最该记住的公式与结论压成几张表，配合正文各讲食用。