↓ 跳过正文
111qqz的小窝
111qqz的小窝
  1. Posts/

GPU Execution Model：为什么大量 Warps 能隐藏内存延迟？

·3238 字·7 分钟
Deep-Learning Gpu-Performance CUDA CMU15-442 Ml-Infra

之前写了 GPU 存储层级、Register Tiling 和 Cache Tiling、FlashAttention 的分块策略——这些文章都在讨论"怎么减少数据搬运"、“怎么让数据在片上多用几次”。但我发现自己一直跳过了一个更前面的问题：GPU 到底怎么执行这些 kernel？

我知道 GPU 有很多核心，知道 CUDA 把线程组织成 Block 和 Grid，知道 Warp 是 32 个线程。但如果有人问：一次 HBM 访问要等几百个 cycle，GPU 的计算单元这段时间在干什么？ 我之前的回答大概只能含糊地说"切换到别的线程"——具体怎么切换、为什么能切换、切换的代价是什么，并没有真正想清楚。

这篇从 CMU 15-442 的 Lecture 06 出发，把这个问题理一遍。

GPU 要的是总吞吐，不是单线程速度
#

CPU 和 GPU 面对的问题不同。

CPU 要跑通用程序——分支多、数据依赖复杂、串行逻辑长。所以 CPU 投入大量晶体管做乱序执行、分支预测、多级 Cache，目标是让单个线程跑得尽可能快。

GPU 面对的典型工作负载是矩阵乘法、卷积、向量运算——成千上万个元素做几乎相同的计算，天然可以并行。GPU 的设计选择是：不追求单线程低延迟，而是把晶体管堆成大量简单的计算单元，同时维持大量执行上下文，追求整体吞吐。

这个选择带来一个直接后果：GPU 没有 CPU 那种深度乱序执行流水线来掩盖单次内存访问的延迟。一次 HBM 访问可能要等几百个 cycle。如果只有一个执行上下文，这几百个 cycle 计算单元就是空转。

所以问题变成了：GPU 靠什么避免这种空转？

从一个 Vector Add 开始
#

先看最简单的 CUDA kernel：

 1__global__ void vector_add(
 2    const float* a, const float* b,
 3    float* c, int n
 4) {
 5    int i = blockIdx.x * blockDim.x + threadIdx.x;
 6    if (i < n) c[i] = a[i] + b[i];
 7}
 8
 9// Host 端启动：
10vector_add<<<4, 256>>>(a, b, c, n);

这里的 <<<4, 256>>> 创建了一个 Grid，包含 4 个 Block，每个 Block 有 256 个 Thread。总共 1024 个 Thread。

NVIDIA GPU 把每 32 个 Thread 组织成一个 Warp——这是硬件调度的基本单位。所以每个 Block 有 256 / 32 = 8 个 Warp，整个 Grid 有 32 个 Warp。

这些概念有一个逻辑层到物理层的映射关系：

  • Grid / Block / Thread / Warp 是逻辑上的任务组织。
  • SM（Streaming Multiprocessor）是物理上的执行单元。
  • 一个 Block 在执行期间被分配到某个 SM。
  • 一个 SM 可以同时驻留多个 Block。
  • SM 内部的 Warp Scheduler 负责选择就绪的 Warp 发射指令。
GPU 执行模型：逻辑任务到物理硬件

这里有一个类比帮助我理解。CPU 的 Process 拥有独立的资源（地址空间、文件描述符），OS Thread 是 Process 内的调度单位。类似地，CUDA Block 拥有一组共享资源（Shared Memory），Warp 是 SM 内的调度单位。

但这个类比有明确的局限：普通 CUDA Block 之间共享 Global Memory 地址空间，不像 CPU Process 之间地址空间隔离。Warp 内的 32 个 Thread 以 SIMT（Single Instruction, Multiple Threads）模式执行——同一条指令应用到不同的数据上——这和 CPU OS Thread 各自独立的执行流也不同。

还有一点需要注意：1024 Thread / Block 是某些 GPU 架构下的最大允许数量，不是每个 Block 的固定大小。<<<4, 256>>> 里的 256 是程序员选择的。

Resident 不等于 Executing
#

上面说的"一个 SM 可以同时驻留多个 Block"——这里"驻留"是什么意思？

一个 Warp 被分配到 SM 之后，它的执行上下文（寄存器、程序计数器）就一直保留在 SM 上，直到 Block 执行完毕。这些 Warp 叫 Resident Warp。

但 Resident 不代表正在执行。SM 内部的 Warp 可以处于不同状态：

  • Resident：已驻留 SM，拥有执行上下文和分配的资源。
  • Eligible：下一条指令的操作数已经就绪，可以被 Scheduler 选择。
  • Issued：被 Warp Scheduler 选中，当前 cycle 正在发射指令。
SM 内部：Warp 状态与 Scheduler 选择

一个 SM 可能驻留几十个 Warp，但每个 cycle 只有少数 Eligible Warp，Scheduler 从中选择一个或几个发射。

GPU 并不是让所有 Resident Warp 在同一个 cycle 同时执行指令。 它的做法是在就绪的 Warp 之间快速交替——而这种交替几乎没有代价，因为每个 Warp 的执行上下文始终驻留在硬件上，不需要像 CPU 上下文切换那样保存和恢复寄存器。

如果 Warp 内部遇到分支，不同 Thread 走了不同的路径（Branch Divergence），硬件会串行执行各个路径。这会浪费一些执行槽位，但不是本篇的重点。

用其他工作填补等待时间
#

现在可以回答开头的问题了。

假设 Warp 0 发射了一条 Global Memory Load 指令。后续计算依赖这个 Load 的结果，数据还没回来——Warp 0 没法继续，变成 Not Eligible。

但 SM 上还驻留着 Warp 1、2、3…… 如果它们的下一条指令已经准备好了（操作数就绪），Warp Scheduler 可以立刻选择它们发射指令。

Latency Hiding：多 Warp 交错执行隐藏内存延迟

Warp 0 等待的那几百个 cycle，被其他 Warp 的有用工作填满了。等数据回来，Warp 0 重新变成 Eligible，Scheduler 会在合适的时机再选择它。

这就是 Latency Hiding。一个很重要的认知：

Latency Hiding 没有降低单次内存访问的延迟——HBM 该等多久还是多久。它的作用是让计算单元在等待期间不闲着，去执行其他已经准备好的工作。

除了多个 Warp 之间的交错，单个 Warp 内部如果存在多条相互独立的指令，也能帮助隐藏延迟。看两段代码：

1// 串行依赖：第二次 Load 地址取决于第一次 Load 的结果
2float x = a[i];
3float y = b[(int)x];
1// 两次独立的 Load
2float x = a[i];
3float y = b[i];

第二种写法的两次 Load 没有数据依赖，硬件可以同时发射（或在指令流水线中重叠），不需要等第一个 Load 完成才能发射第二个。这种能力叫 Memory-Level Parallelism——即使只有一个 Warp，多个在途的 Memory Request 也能部分隐藏延迟。

实际 kernel 的延迟隐藏同时依赖这两个层面：Warp 之间的交错（Thread-Level Parallelism）和 Warp 内的独立指令（Instruction-Level Parallelism）。

更多 Warps 不一定更快
#

既然更多 Warp 意味着更多可以交替的执行上下文，那是不是应该把 SM 塞满？

这就涉及 Occupancy 的概念。Occupancy 定义为：

$$ \text{Occupancy} = \frac{\text{Resident Warps}}{\text{Maximum Resident Warps}} $$

拿 A100 举例：108 个 SM，每个 SM 最多驻留 64 个 Warp（2048 个 Thread）。如果某个 kernel 的资源配置导致每个 SM 只能驻留 32 个 Warp，Occupancy 就是 50%。

直觉上 Occupancy 越高越好——更多 Resident Warp 意味着更大的调度池，更容易找到 Eligible Warp 来填补等待。

但事实不总是这样。

考虑两个极端情况。Kernel A：Occupancy 100%，但每个 Warp 的指令序列是一条很长的依赖链，每条指令都等前一条的结果。大量 Warp 驻留在 SM 上，但它们大部分时间都不是 Eligible。Kernel B：Occupancy 50%，但每个 Warp 有大量相互独立的指令。尽管 Warp 数量少一半，Scheduler 总能找到可以发射的指令。

Kernel B 完全可能比 Kernel A 更快。

另一个角度：当 HBM 带宽已经被打满（所有 SM 发出的 Memory Request 总量已经达到了 HBM 可持续的吞吐上限），继续增加 Resident Warp 不会带来更多有效带宽。这时候瓶颈在带宽，不在调度池的大小。之前那篇讨论过 memory bound 和 compute bound 的区分，这里是同一个道理。

还有一个实际约束：Occupancy 受寄存器数量和 Shared Memory 容量的限制。每个 Thread 使用的寄存器越多，SM 能同时驻留的 Warp 就越少。Register Tiling 那篇提过 tile 太大导致 register spilling 的问题——在 GPU 上还多了一层：寄存器用多了会直接压低 Occupancy，减少可调度的 Warp 数量。

Occupancy 是隐藏延迟的手段之一，不是优化的最终目标。 最终目标是计算单元和内存带宽的有效利用率——这取决于指令的独立性、内存访问模式、带宽是否饱和等多个因素，而不单单是"SM 上塞了多少个 Warp"。

写在最后
#

串起来看，GPU 执行模型的因果链其实很清晰：

GPU 追求高吞吐 → Kernel Launch 暴露大量独立工作（Grid / Block） → Block 被调度到 SM → 每个 SM 驻留多个 Warp → Warp Scheduler 在就绪的 Warp 之间交替发射 → 某个 Warp 等待内存时，其他 Warp 填补空隙 → 计算单元保持忙碌。

但这条链有上限。当带宽饱和、或所有 Warp 都在等同一种资源时，更多的 Warp 也无法进一步提高吞吐。

GPU 不是靠降低单次访存延迟来提高性能——它靠的是在等待期间做别的事情。理解这一点，是理解后续所有 GPU 性能优化（Tiling、Occupancy 调优、Kernel Fusion）的前提。

参考资料
#

相关文章

从 Register Tiling 到 Cache Tiling：理解高性能计算中的 Data Reuse

·4020 字·9 分钟
Deep-Learning GEMM Ml-Infra Gpu-Performance CMU15-442
最近在看 CMU 15-442 的 Hardware Acceleration。课件翻到 Register Tiling，一个 2×3 output tile 的例子。我盯着那页 slide 看了一会儿——Register Tiling 这个词不陌生，但要我说清楚它和 Cache Tiling 的区别，以及它们到底在优化什么，我发现自己的理解一直停在"把矩阵切成小块"这个表面。

从 Dense FFN 到 MoE：Router、Token Dispatch 与 Grouped GEMM

·5127 字·11 分钟
Deep-Learning MoE CUDA Ml-Infra Transformer
上一篇 MoE 笔记把架构家族理了一遍——稀疏、稠密、Soft MoE，gate 在前还是在后，决定了算力形态和适用场景。 但那篇有一段伪代码一直让我不太舒服： 1for prob, idx in zip(topk_probs.unbind(1), topk_idx.unbind(1)): 2 out += prob.unsqueeze(1) * self.experts[idx](x) 逻辑上没问题。但想象一下真实场景：一个 batch 有几千个 token，每个 token 选 2 个 expert，就是上万个 token-expert 计算任务。用 Python for loop 逐个调 expert 的 forward，GPU 利用率会低得离谱。

GPU 的数据存在哪？——从 HBM 到片上存储的存储层级

·2379 字·5 分钟
Deep-Learning Gpu-Performance HBM FlashAttention 6.S191 Ml-Infra
最近在看 MIT 6.S191 Lecture 9 的时候，反复听到一句话：Data Movement Matters。 之前写过 Arithmetic Intensity 和 Compute/Memory Bound，也刚写完 FlashAttention 的分块和 Online Softmax。这些文章里"HBM"和"片上存储"出现了很多次，但我发现自己其实一直没有认真想过一个问题：HBM 到底是什么？它全称叫 High Bandwidth Memory——高带宽内存——既然带宽已经很高了，为什么还要费这么大劲减少对它的访问？