课程结构（来自 syllabus）

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

复习路线图

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

考点权重（编者估计）

L01 · 引论：为什么需要并行

来源：new-pdc01-intro.pdf

1.1 AI 的硬件需求与三大瓶颈

• 模型参数呈指数增长（BERT 0.34B → GPT-3 175B → DeepSeek-V3 671B(MoE) → GPT-4.5 ≈12.8T(MoE)）。
• 算力需求增速超过摩尔定律。
• 通信规模成为新的关键瓶颈。

CPU 为什么「失败」于 AI

GPU 的崛起是「意外的英雄」：它本为图形（矩阵/向量运算）而生，恰与深度学习的数学一致；2007 年 CUDA 让 GPU 可用于通用计算，训练时间从「周」降到「小时」。

1.2 「内存墙」：参数 vs 显存

1.3 并行能帮上忙吗？两条定律

课件例子：10 个处理器、80% 并行、20% 串行，Amdahl 给出 $S=\dfrac{1}{0.2+0.8/10}=3.57$，远达不到 10×。现实中还有数据依赖、I/O 瓶颈等进一步削弱并行度。

L02 · 异构编程：CUDA 核心

来源：new-pdc02-hetero.pdf

2.1 主机、设备与互连

• host ≈ CPU；device ≈ GPU（在 OpenCL/ROCm 中 CPU 也可作设备）。
• host–device 互连：PCIe、NVLink-C2C。课件数据：x86–5090 走 PCIe 5.0 约 64 GB/s；Grace-Hopper NVLink C2C 900 GB/s；Grace-Blackwell 1.8 TB/s。
• device–device 互连：PCIe、NVLink、RDMA。NVLink 1.0（2016，P100）双向 160 GB/s，已快于 PCIe 3.0 x16（32 GB/s）。
• GPUDirect：Storage（NVMe→GPU 一次拷贝）、P2P（GPU 间直接 Load/Store）。

2.2 线程层次：表达并行

线程通过内置变量定位：threadIdx、blockIdx、blockDim、gridDim。全局 id 常写作 i = threadIdx.x + blockIdx.x * blockDim.x。

__global__ void myKernel(int N, double *d_a) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) d_a[i] *= 2.0;
}
// 启动：
dim3 threads(256,1,1);
dim3 blocks((N+255)/256,1,1);
myKernel<<<blocks,threads>>>(N, a);

2.3 SIMD / SPMD / SIMT 三者辨析

2.4 占用率（Occupancy）：三大资源限制

每 SM 的活跃线程数受以下三者最小值限制：

2.5 控制分歧、谓词执行

warp 内线程锁步执行。若走不同分支（divergence），硬件用掩码屏蔽无效路径，各路径被串行执行。短分支可用谓词执行（predication）：执行后丢弃无效结果，避免分支跳转。建议把分支控制在 warp 粒度（整 warp 走同一路径则不分歧）。

2.6 访存优化：合并访问与 bank conflict

• 合并访存（coalescing）：总线按对齐宽度（如 32 字节）访问；同一 warp 的线程若访问落入同一段，多次访问被合并成更少请求，提升带宽利用率。
• 线程映射（thread mapping）：行主序数据下，应让 warp 内连续线程映射到列（第二维），实现合并访问；矩阵乘「映射 2」远优于「映射 1」即源于此。
• 矩阵转置：无论怎么映射，读/写总有一个不合并；用 shared memory 缓冲并重排线程映射，可让读写都合并（工作组需为方形）。
• Bank conflict：shared memory 分 bank，多个线程命中同一 bank 的不同地址 → 串行化；命中不同 bank（无冲突）或全部命中同一地址（广播）→ 无延迟。

L03 · CUDA 实例：归约与分块卷积

来源：new-pdc03-cuda-examples.pdf

3.1 并行归约（reduction）

从 atomicAdd 到树形归约

// 关键机制
extern __shared__ int sdata[];   // block 内共享
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];

两个关键优化

为什么不在设备端做全局同步？

3.2 分块卷积（tiled convolution）

• 瓶颈：内存带宽 → 用分块算法 + shared memory 复用。
• 输入 tile vs 输出 tile：输入 tile 维度 = 输出 tile 维度 + $2\times$ filter radius（边界需邻接 tile 的 halo 元素；越界 ghost 元素填 0）。
• 滤波器放在 constant memory；输入 tile 放 shared memory，多个线程复用。
• 尺寸不匹配的处理：启动足够线程加载整个输入 tile，再用子集线程计算并写出输出（会有 if 与控制分歧）。

__global__ void reduceOptimized(float* input, float* output, int N) {
    extern __shared__ float sdata[];
    int tid = threadIdx.x;
    int i   = blockIdx.x * blockDim.x + threadIdx.x;
    sdata[tid] = (i < N) ? input[i] : 0.0f;   // 1) 先搬到 shared
    __syncthreads();
    for (int s = blockDim.x/2; s > 0; s >>= 1) { // 2) 连续寻址 + warp 感知
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }
    if (tid == 0) output[blockIdx.x] = sdata[0];
}

L04 · Attention 与 FlashAttention

来源：new-pdc04-attention.pdf

4.1 自注意力与 Q/K/V

把自注意力看作信息检索：每个 token 生成三个向量——Query「我在找什么」、Key「我包含什么」、Value「我提供的内容」。

4.2 复杂度与「内存墙」

4.3 GPU 存储层次（A100 例）

SRAM 比 HBM 快约 10×，但容量很小 → 不能整张 $N\times N$ 矩阵放进去。

4.4 FlashAttention：两大支柱

4.5 Online Softmax

分块的困境：softmax 需要全局最大值（数值稳定）与全局求和（归一化），但处理第一块时看不到后续块。

L05 · 卷积、非线性与算子融合

来源：new-pdc05-conv-nonlinear-fusion.pdf

5.1 Roofline 与算子分类

Roofline 用「算力 (FLOPS)」与「内存带宽」两条天花板，结合算子的算术强度判断瓶颈。课件强调：算力增速远超带宽增速 → 「内存墙」使计算核常等数据。

5.2 卷积三种实现

• Img2Col：把空间卷积转成 GEMM → 调用 cuBLAS/MKL 等高度优化库；代价是数据冗余。
• Winograd $F(2,3)$：常规 6 次乘法 → 4 次乘法 + 11 加 + 2 移位。tile 越大乘法省得越多，但变换复杂度与数值误差也越大（$F(4×4,3×3)$ 理论省 6×，但变换会吃掉收益）。

5.3 Element-wise vs Reduction

原子操作的代价：atomicAdd 虽正确，但需锁定内存位置、read-modify-write 串行化，GPU 并行退化为单核性能，只适合稀疏访问。满足结合律的归约应用树形结构（$O(N)$ 工作量、$O(\log N)$ 深度）。

5.4 算子融合（Operator Fusion）

融合 Softmax 例：Load（协作读入 shared）→ Sync & Reduce（片上多轮，CUDA 用 __syncthreads()，Triton 自动）→ Write（仅最终结果写回）。把 Max、Sub、Exp、Sum 融成一个 kernel；连续寻址避免 bank conflict；Brent 定理：用 $O(N/\log N)$ 线程每线程串行处理 $O(\log N)$ 元素，总代价 $O(N)$。

L07 · Tensor Cores 的演进

来源：new-pdc07-tensor-cores.pdf

7.1 性能第一性原理

• Amdahl 定律：增加算力只能加速并行部分。
• 算术强度（roofline 分析）。
• 数据移动是「原罪」：计算便宜、搬数据昂贵 → 几乎所有 Tensor Core 优化都伴随访存改进。

7.2 MMA 与存储层次

MMA（multiply-accumulate）计算 $D = A\times B + C$，$A$ 为 $M\times K$、$B$ 为 $K\times N$、$C/D$ 为 $M\times N$，形状记作 $M\times N\times K$。数据在 SMEM ↔ RF ↔ Tensor Core 间流动。

7.3 演进总趋势

7.4 TMA（Hopper）要点

• 专用硬件，加速 GMEM↔SMEM 的大块异步拷贝（cp.async.bulk）。
• CTA 内单线程即可发起，自动处理 stride/offset/边界，释放线程做别的事。
• 对小请求反而延迟高（地址生成开销）→ 适合大块拷贝以摊薄开销（如 LLM 推理中按 16B 倍数分块）。

L08 · AI 编译器

来源：new-pdc08-ai-compilers.pdf

8.1 通用编译流程

总目标：最小化内存占用、提升执行效率、可扩展到多异构节点。

8.2 计算图：静态 vs 动态

计算图：节点=张量，边=函数；用于表示反向传播的链式法则，便于自动优化与 lowering。

8.3 torch.compile 三大组件

8.4 IR 层次与 MLIR / TVM / Triton

• MLIR：可复用的多层 IR/Pass，多 dialect 共存（tensor/arith/math/linalg…），逐层 lowering + pattern rewrite；可在合适层次做优化（如 linalg 检测矩阵两次转置，affine 做 loop tiling）。
• TVM：Relay（函数式、静态类型 IR）→ TE → TIR；AutoTVM（模板搜索）→ Auto-scheduler/Ansor（无需手写模板）。
• Triton：以 block/tile 为粒度、类 NumPy 风格写 GPU 代码，自动处理线程/warp、向量化、shared memory 分配；autotuner 只调 tile 大小，比 TVM 简单。编译流程：Triton IR → TritonGPU IR → LLVM IR → PTX。
• TileLang：三层接口（硬件无关/带 tile 库/带线程原语），tile 作为一等对象。

8.5 图优化分层

算子调度：Tiling（分块利用局部性）、Reorder（重排、利用结合律 $A@(B@v) < (A@B)@v$）、Split。

L09 · 训练/推理中的并行：DP / TP / PP / 混合

来源：new-pdc09a~d（data / tensor / pipeline / hybrid parallelism）

9.1 三条正交的并行轴

9.2 训练显存从哪来

9.3 数据并行与 ZeRO

• DDP：每 GPU 完整模型副本，处理不重叠 mini-batch，反向后 All-Reduce 同步梯度。梯度分桶减少网络调用、与反向计算重叠。
• ZeRO-1：分片优化器状态；ZeRO-2：再分片梯度；ZeRO-3（FSDP）：连参数一起分片，按需 all-gather 取参、用完即弃。

K 为每参数优化器状态字节数（FP32 动量典型 12B）。

9.4 张量并行（Megatron）

LLM 参数大多在线性投影（Attention 的 Q/K/V/O，MLP 的上/下投影），即 GEMM $Y=XW$。MLP 块切法：

• 第一个矩阵 $A$ 按列切 → 各列输出可独立过 GeLU（无需同步）。
• 第二个矩阵 $B$ 按行切 → 块乘积无需中途同步。
• 前向 $f$ 为恒等、$g$ 为 all-reduce；反向相反。

9.5 流水线并行

为什么用 PP：省显存（vs DDP），通信只依赖激活、点对点（vs FSDP/TP），适合慢速链路（跨节点）。局限：降气泡需大 M，但全局 batch 受收敛限制，M 大则 micro-batch 太小、浪费 Tensor Core。

9.6 混合并行（3D / 5D）与自动并行

• 3D：内层 TP（节点内 NVLink）× 中层 PP（跨节点 IB）× 外层 DP（全局）。案例 MT-NLG 530B：8-way TP × 35-way PP × 16-way DP = 4480 A100。
• 5D：新增 CP（上下文并行，ring 交换 KV 防长序列 OOM）与 EP（专家并行，All-to-All 路由 MoE）。TotalGPUs = TP×CP×EP×PP×DP。
• 自动并行：FlexFlow（SOAP：Sample=DP/Operation=PP/Attribute=CP/Parameter=TP + MCMC）、Alpa（层级搜索：intra-op ILP + inter-op DP）、GSPMD、Unity、Galvatron。

L10 · 通信原语、MPI 与集合通信库

来源：new-pdc10a~c（communication primitives / MPI & CCL）

10.1 通信数据通路

• Scale-up（节点内）：NVLink、HCCS、Unified Bus。
• Scale-out（跨节点）：InfiniBand (IB)、RoCE（RDMA NIC）。
• 集合通信涉及一组进程（如 AllReduce、AllGather、Broadcast）；点对点：Send/Recv（双边）、Put/Get（单边）。

10.2 七大集合通信原语

10.3 α-β 通信模型与 Broadcast

10.4 Ring All-Reduce 与双二叉树

10.5 MPI 基础

编译：mpicc -o foo foo.c（mpicc 只是 wrapper）；运行：mpirun -np 8 ./foo。SPMD：-np 8 ./hello；MPMD：-np 1 ./master : -np 7 ./slave。MPI 名称：MPI_Bcast / MPI_Reduce / MPI_Allgather / MPI_Reduce_scatter / MPI_Allreduce / MPI_Alltoall。

作业精解索引

各次作业的代表性大题已分散在对应讲次，点击展开解答。下面是速查表（题面来自作业 PDF，解答中标 [自注] 者为编者补充推导）。

互动自测

考前速查

核心公式

易混对照

记忆口诀 编者补充