刚开始接触 GPU 和 CUDA 时，很容易把 Stream、Grid、Block、Warp、SM 和 CUDA Core 混在一起。它们其实分属不同层次：有些是程序员用来描述计算的抽象，有些是 GPU 内部真实存在的硬件，还有一些负责把软件定义的任务映射到硬件上执行。

本文沿一次 CUDA 任务从 CPU 提交到 GPU 完成的过程展开。先以 GA100 为例认识 GPU 的主要硬件，再看 Stream 如何承载工作、Kernel 如何组织 Grid、Block 与 Thread，以及这些 Thread 如何组成 Warp 并在 SM 上执行；随后继续梳理数据经过的内存层级与不同范围的同步机制。最后的小 Batch Decode 只是一个综合例子，用来观察这些概念如何在真实计算中共同发挥作用。

GPU 基础架构

GPU 的整体结构：以 GA100 为例

本文以 NVIDIA Ampere 架构的 GA100 作为具体例子。它的主要硬件层级如下：

GA100
├── GPC
│   └── TPC
│       └── SM
│           ├── SM Processing Block（本文简称 SMSP）
│           │   ├── Warp Scheduler
│           │   ├── Register File
│           │   └── Execution Units
│           └── Shared Memory / L1 Data Cache
├── L2 Cache
└── HBM2

从计算角度看，GPC 和 TPC 主要用于组织大量 SM；CUDA 程序不会直接指定工作运行在哪个 GPC 或 TPC。真正承载 CUDA Thread Block、执行指令的核心硬件是 SM。GPC 之外的 L2 Cache 由整个 GPU 共享，HBM2 则是承载 Device Memory 的片外内存。

GA100 完整 GPU 框图，展示 GPC、TPC、SM、L2 Cache、HBM2 Memory Controller、NVLink 和 PCI Express Host Interface 的位置关系

完整 GA100 GPU 框图。图片来源：NVIDIA, “Ampere Architecture In-Depth”, Figure 4。

Streaming Multiprocessor

SM 是 Thread Block 的驻留与资源分配边界。一个 Block 会在同一个 SM 上完成执行，并使用该 SM 提供的寄存器、Shared Memory 和执行资源。以 GA100 为例，SM 内部可以概括为：

  • SM Processing Block（本文简称 SMSP）：SM 内的执行分区；一个 SM 包含 4 个 SMSP，每个 SMSP 管理一组常驻 Warp。
  • Warp Scheduler / Dispatch Unit：选择就绪的 Warp，并将指令发往执行单元。
  • Register File：保存各个 Thread 的寄存器数据。
  • Execution Units：执行具体指令；FP32、INT32 和 FP64 单元负责常规算术，Tensor Core 负责矩阵运算，Load/Store Unit 负责访存，SFU 负责特殊数学函数。
  • Shared Memory / L1 Data Cache：在 GA100 上共享 SM 的片上存储资源。Shared Memory 由 Thread Block 显式使用，L1 Data Cache 则由硬件自动缓存访存数据。

GA100 Streaming Multiprocessor 由四个 SM Processing Block 组成，每个分区包含 Warp Scheduler、Dispatch Unit、Register File、FP32、INT32、FP64、Tensor Core 和 Load/Store Unit

GA100 Streaming Multiprocessor 的内部结构。图片来源：NVIDIA, “Ampere Architecture In-Depth”, Figure 5。

到这里先只需要记住两个边界：Block 驻留在 SM，Warp 则由 SM 内的 SMSP 管理和执行。后面的 CUDA 编程模型会解释 Block 与 Warp 从何而来。

CPU 如何向 GPU 提交工作

Stream：从提交任务到取得结果

从 CPU 看，GPU 是一个异步执行设备。CPU 不会像调用普通函数那样，启动 Kernel 后原地等待返回值；它会把数据传输、Kernel Launch 和 Event 等操作依次提交到 CUDA Stream，然后继续执行自己的代码。

下面是一条最简单的处理链：输入先从 Host Memory 复制到 Device Memory，Kernel 在 GPU 上产生输出，输出再被复制回 Host Memory。

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream);
vectorAdd<<<grid, block, 0, stream>>>(d_a, d_b, d_c, n);
cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream);
cudaEventRecord(done, stream);

// CPU 可以在这里继续处理其他工作。

cudaEventSynchronize(done);
// 到这里，CPU 才能安全使用 h_c 中的结果。

同一 Stream 中的操作按提交顺序执行，因此 D2H Copy 不会早于 Kernel，Event 也不会早于 D2H Copy 完成。异步 API 的返回只表示操作已经提交，并不表示 GPU 已经完成。CPU 只有在等待 Event、Stream 或整个 Device 之后，才能确认相应结果已经可用。

CPU 发出异步 CUDA 调用后继续执行其他工作，GPU 则按照 Stream 顺序执行 H2D Copy、Kernel、Event 和 D2H Copy

CUDA Stream 中的异步提交与执行顺序。图片来源：NVIDIA CUDA Programming Guide, “Asynchronous Execution”, Figure 20。

如果结果仍要交给下一个 Kernel，就没有必要先传回 CPU：继续把下一个 Kernel 提交到同一 Stream，它会在前一个 Kernel 完成后直接读取 Device Memory 中的结果。不同 Stream 中没有依赖的操作则可能并发执行；如果一个 Stream 需要使用另一个 Stream 产生的数据，可以通过 Event 建立 GPU 侧依赖，而不必让 CPU 同步等待。

Stream 描述了任务的提交顺序和结果何时可见，但没有解释一个 Kernel 进入 GPU 后如何执行。接下来从执行流、数据流与同步等待三个角度继续下钻。

GPU 如何执行这些工作

执行流：从 Kernel Launch 到执行单元

一次 Kernel Launch 会创建许多执行同一段 Kernel 代码的 Thread。CUDA 将这些 Thread 组织成 Grid 和 Block，让程序可以用固定规模的硬件处理任意规模的数据。

CUDA 预先定义了三层并行结构：

  • Grid 是一次 Kernel Launch 创建的全部 Block。
  • Block 是 Grid 中可以独立执行的 Thread 分组，也是 Shared Memory 和块内同步的作用范围。不同 Block 的执行顺序不受保证。
  • Thread 是执行 Kernel 代码的最小逻辑实例，通过自身在 Grid 和 Block 中的坐标确定要处理的数据。

下面的一维向量加法展示了最基本的索引方式：

__global__ void vectorAdd(
    const float* a,
    const float* b,
    float* c,
    int n
) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize, 0, stream>>>(d_a, d_b, d_c, n);

<<<gridSize, blockSize, 0, stream>>> 指定这个 Grid 包含 gridSize 个 Block，每个 Block 包含 blockSize 个 Thread。Kernel 中的每个 Thread 再通过 blockIdx、blockDim 和 threadIdx 找到自己负责的元素。Grid 可以远大于 GPU 当前能够同时执行的规模，尚未运行的 Block 会等待硬件资源空闲。

以上是 CUDA 暴露给程序员的逻辑结构。落实到 NVIDIA GPU 上，Block 是分配到 SM 的基本单位：一个 Block 的所有 Thread 都在同一个 SM 上执行，一个 SM 则可以同时驻留多个 Block。整体映射关系如下：

CUDA 编程模型中的 Thread、Thread Block 和 Kernel Grid，以及对应的 CUDA Core、Streaming Multiprocessor 和 GPU 硬件层级

CUDA 编程模型与 GPU 硬件层级的对应关系。图片来源：Modal GPU Glossary, “What is a CUDA Thread Block?”。

上图只表示 CUDA 编程层级与 GPU 硬件层级的宏观对应关系，Thread 与 CUDA Core 并不是一对一映射。Block 驻留到 SM 后，其中连续的 32 个 Thread 会组成一个 Warp；SMSP 的 Warp Scheduler 以 Warp 为单位选择并发射指令，指令随后被送往相应的执行单元，并作用于 Warp 中当前处于活跃状态的 Thread。

一个 Thread Block 按照连续的 Thread ID 被划分成若干组，每组 32 个 Thread 形成一个 Warp，并由 SM 调度执行

Thread Block 的 Warp 划分方式。图片来源：NVIDIA CUDA Programming Guide, “Hardware Multithreading”, Figure 22。

问题也就来了：一个 Kernel 动辄启动成千上万个 Block，每个 Block 又包含数百个 Thread，远超 GPU 在同一时刻能够真正执行的规模。面对这么多 Block 和 Warp，有限的 SM 和执行单元是怎么把它们高效地“跑起来”的？

答案是将空间并行与时分复用结合起来：不同 SM 可以同时执行不同 Block；同一个 SM 也可以驻留多个 Block，因此可供调度的 Warp 往往来自不同 Block。某个 Warp 因访存或指令依赖而停顿时，Warp Scheduler 会继续发射其他已经就绪的 Warp，让执行单元保持忙碌。不同 Block 的 Warp 由此在执行单元上交错运行，使多个 Block 能够并发推进，并通过以 Warp 为单位的时分复用隐藏延迟。

Warp Scheduler 从多个 Thread Block 的常驻 Warp 中选择已经就绪的 Warp 发射指令，并在某个 Warp 停顿时穿插执行其他 Warp

Warp 调度与延迟隐藏示意图。图片来源：《CUDA 如何调度 kernel 到指定的 SM？——Zihao Zhao》。

数据流：数据如何参与计算

在 A100 系统中，数据由 CPU 经 PCIe 写入 HBM。Kernel 发起 Global Memory 访问后，请求可能由 SM 内的 L1 Data Cache 或 GPU 共享的 L2 Cache 命中，操作数最终进入 Register File，供执行单元使用。Shared Memory 则是一条由 Kernel 显式控制的数据复用路径，并非自动缓存。

硬件结构位置可见范围常见用途
CPU DRAMHostCPU保存输入数据和 CPU 需要读取的结果
HBMGPU 片外整个 GPU保存权重、激活值、KV Cache 和 Kernel 输出
L2 CacheSM 之外整个 GPU缓存所有 SM 的 Global Memory 访问
L1 Data CacheSM 内当前 SM缓存当前 SM 的访存数据
Shared MemorySM 内同一 Block复用数据和在线程间通信
Register FileSM / SMSP 内每个 Thread 逻辑私有保存指令操作数和中间结果
  • vectorAdd 的路径是 HBM → Cache → Register File → FP32 Unit → HBM。数据几乎不复用，通常受 HBM 带宽限制。
  • 矩阵乘法的路径是 HBM / L2 → Shared Memory → Register File → Tensor Core。数据在片上反复复用，以减少 HBM 访问；GA100 的异步数据拷贝指令还可以跳过 Register File，直接把数据从 Global Memory 搬入 Shared Memory。

优化数据流的核心，就是让数据尽量少跨越 PCIe 和 HBM，并在片上多次复用。

Kernel 输出通常保留在 HBM 中供后续 Kernel 使用；只有 CPU 确实需要结果时才进行 D2H Copy。

同步与等待：什么时候可以继续执行

从 Stream 的视角看，Kernel 通常像普通程序一样顺序执行：同一 Stream 中，后一个 Kernel 会等待前一个 Kernel 完成。Kernel 边界因此也是最常见的全 Grid 同步点。

phase1<<<grid, block, 0, stream>>>(data);
phase2<<<grid, block, 0, stream>>>(data);

但一个 Kernel 内部并没有这样的全局顺序。不同 Block、Warp 和 Thread 会并发或交错推进；与其说它们“乱序执行”，更准确的说法是 CUDA 不保证它们之间的先后关系。只要每个 Thread 处理独立数据，这通常不是问题；当多个 Thread 需要交换数据时，才必须显式同步。

例如，一个 Block 分批把数据搬入 Shared Memory 并反复计算时，需要两个 Barrier：

__shared__ float tile[256];

for (int offset = 0; offset < n; offset += blockDim.x) {
    tile[threadIdx.x] = input[offset + threadIdx.x];
    __syncthreads();  // 等整个 Block 写完这一批数据

    consume(tile);
    __syncthreads();  // 等整个 Block 用完，再覆盖 tile
}

__syncthreads() 会等待整个 Block，并保证 Barrier 之前的 Shared Memory 写入对块内 Thread 可见。如果协作只发生在一个 Warp 内，则可以缩小同步范围：

__shared__ float warpData[32];

if (threadIdx.x < warpSize) {
    warpData[threadIdx.x] = input[threadIdx.x];
    __syncwarp();

    if (threadIdx.x == 0)
        output[0] = warpData[1];
}

这里 __syncwarp() 只等待同一个 Warp 中参与的 Thread。无论使用哪种 Barrier，参与同步的 Thread 都必须到达对应调用；否则可能造成错误或停滞。两者都不能同步不同 Block，跨 Block 的阶段同步通常仍由前后两个 Kernel 提供。

不同 Stream 之间默认没有执行顺序。如果一个 Stream 的 Kernel 要读取另一个 Stream 产生的数据，可以让生产者记录 Event，再让消费者等待这个 Event：

producer<<<grid, block, 0, producerStream>>>(data);
cudaEventRecord(ready, producerStream);

cudaStreamWaitEvent(consumerStream, ready, 0);
consumer<<<grid, block, 0, consumerStream>>>(data);

这个依赖完全由 GPU 维护，不会阻塞 CPU。

CPU 只有在需要读取结果时才必须等待 GPU。cudaStreamSynchronize() 会阻塞到指定 Stream 完成，cudaDeviceSynchronize() 则等待整个 Device；更细粒度的做法是在 Stream 中记录一个 Event，把它当作完成标记：

cudaEventRecord(done, stream);

cudaError_t status = cudaEventQuery(done);
if (status == cudaSuccess) {
    use_result();
} else if (status == cudaErrorNotReady) {
    do_other_cpu_work();
}

Event 只有在它之前提交到该 Stream 的操作全部完成后才会变为完成状态。CPU 可以用 cudaEventQuery() 非阻塞地检查，也可以在真正需要结果时调用 cudaEventSynchronize() 阻塞等待。所谓异步，并不是 CPU 不再关心结果，而是先提交工作和完成标记，继续处理其他任务，之后再查询或等待这个标记。

为什么 Decode 需要 Continuous Batching

前面的知识可以用一个经典问题串起来：为什么 LLM Decode 在 Batch 很小时很难用满 GPU，而 Continuous Batching 又能显著提高吞吐？

以一个隐藏维度为 4096 的线性层为例。单个请求每次只生成一个 Token，因此核心计算接近一次矩阵向量乘法：y=xWy=xW，其中 xx 的形状是 1×40961\times4096，WW 的形状是 4096×40964096\times4096。如果权重使用 FP16，忽略输入和输出后，这次计算大致需要：

读取权重：4096 × 4096 × 2 Byte ≈ 32 MiB
完成计算：4096 × 4096 × 2 FLOP ≈ 33.6 MFLOP
算术强度：约 1 FLOP / Byte

作为量级对比，A100 40GB 的 HBM2 峰值带宽是 1555 GB/s，读取 32 MiB 至少需要约 22 微秒；它的 FP16 Tensor Core 峰值计算吞吐是 312 TFLOPS，完成 33.6 MFLOP 理论上只需要约 0.1 微秒。虽然真实执行还会受到 Cache、指令和调度开销影响，但这个超过两个数量级的下限已经说明：Batch 为 1 时，GPU 大部分时间是在搬权重，而不是做乘加。

这时即使 Grid 中有大量 Block、每个 SMSP 也驻留了多个 Warp，Warp Scheduler 最多只能在访存等待期间切换到其他 Warp；当所有 Warp 都在等待 HBM，增加并行 Thread 也无法突破带宽上限。问题不在于“没有足够多的 Thread”，而在于每从 HBM 读入一份权重，只做了很少的计算。

如果一次同时处理 BB 个 Token，xx 就从一个向量变成 B×4096B\times4096 的矩阵。同一块权重从 HBM 搬入 L1 / Shared Memory 后，可以被多个 Token 复用，再由 Warp 将数据送入 Register File 和 Tensor Core。计算量随 BB 增长，权重读取量却不需要同比增长，矩阵向量乘法也逐渐变成更适合 GPU 的矩阵乘法。

Continuous Batching 做的正是持续维持这个 BB：推理引擎在每个 Decode 迭代结束时移除已经完成的请求，再把新请求加入下一批，而不是等待整批请求全部结束。CPU 只需要在本轮 Sample 和结果回传完成后更新 Batch，Stream 与 Event 则让它在 GPU 执行期间准备其他工作。Batch 越大，权重复用和吞吐通常越好，但单个请求的排队时间与 KV Cache 占用也会增加。现代推理调度器的核心工作，就是在这组吞吐、延迟与显存约束之间寻找平衡。

它没有改变 GPU 的执行方式，只是让每轮 Decode 有更多工作可调度，也让从 HBM 读入的权重服务更多 Token。这正是前文执行流与数据流在实际工作负载中的交点。

参考资料

  1. NVIDIA CUDA Programming Guide: Programming Model
  2. NVIDIA CUDA Programming Guide: Writing SIMT Kernels
  3. NVIDIA CUDA Programming Guide: Asynchronous Execution
  4. NVIDIA CUDA Programming Guide: Hardware Multithreading
  5. NVIDIA CUDA Programming Guide: Synchronization Primitives
  6. NVIDIA CUDA C++ Best Practices Guide
  7. NVIDIA, “Ampere Architecture In-Depth”
  8. Orca: A Distributed Serving System for Transformer-Based Generative Models
  9. 《CUDA 如何调度 kernel 到指定的 SM？——Zihao Zhao》