GPU Kernel、异步 Launch、Stream 与融合 Kernel 全解

本文目标：把"GPU 在跑 kernel"这件事从原理讲到工程——kernel 是什么、CPU/GPU 怎么协作、为什么 launch 是异步、synchronize 和 stream 怎么用、为什么要融合 kernel、怎么写/怎么调。读完你能回答面试官"你说的 kernel 是什么"“为什么 PyTorch op 测时间要 sync”“stream 是干什么的”“为什么要融合算子”。

和前面几篇呼应：Megatron Timer 反复强调 cuda.synchronize()，本篇讲清它为什么是命门。

一、先澄清：此 kernel 非彼 kernel

“kernel” 有两个常见含义，别混：

操作系统内核 GPU kernel
是什么 OS 的核心，常驻 ring0 管硬件 一段在 GPU 上并行执行的函数
谁跑 CPU GPU 上成千上万个线程
例子 Linux kernel、Windows kernel vector_add、matmul、flash_attention

本文讲的是 GPU kernel。一句话定义：

kernel 是你写的一段被编译成 GPU 指令、由 GPU 上成千上万个线程并行执行的函数。CPU 调用它时只是"派单"，真正干活的是 GPU。

最小例子，矩阵加法 C = A + B，CPU 写一行 C = A + B，底层翻译成 CUDA kernel：

1
2
3
4
5
// 这就是一个 kernel: 被 GPU 上几万个线程同时执行的函数
__global__ void vector_add(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 我是第几个线程
    if (i < N) C[i] = A[i] + B[i];                   // 每个线程算一个元素
}

CPU 调用 vector_add<<<blocks, threads>>>(...)：

  • 立刻返回（异步），只是把这个 kernel 丢进 GPU 的执行队列。
  • GPU 的 SM（流多处理器）拿到 kernel，启动成千上万线程，每个算一个 C[i]，全部并行。

二、CPU/GPU 协作模型：派单与施工

把 CPU 当项目经理、GPU 当工地，最能记住：

1
2
3
4
5
6
7
8
9
10
CPU(项目经理)                    GPU(工地)
   │
   │ 1. 准备数据在显存(或交给 DMA 搬)
   │ 2. 填派工单: kernel名 + 线程数 + 参数  ──► [GPU 执行队列] ──► SM 开始跑
   │ 3. 立刻继续往下写下一张派工单          (异步, CPU 不等)
   │    ...又一个 kernel 丢进队列...
   │
   │ 需要结果时: cuda.synchronize() ─────► 阻塞 CPU, 等队列全跑完
   ▼                                          │
   拿到结果 ◄─────────────────────────────────┘

关键事实：

  • CPU 提交 kernel（launch）极快（微秒级），它只填一个"派工单"丢进队列。
  • GPU 跑 kernel 慢（毫秒~秒级），但成千上万线程并行。
  • 因为 launch 快、GPU 慢，CPU 会一口气把一串 kernel 全丢进队列，GPU 按顺序慢慢跑。这就是"异步 pipeline"——CPU 和 GPU 重叠工作，GPU 不空等、CPU 不干等。

为什么必须异步（不异步就慢）

如果每次 launch 都同步等 GPU 跑完再返回，CPU 要傻等——GPU 算 1ms，CPU 也要等 1ms 才能发下一个 kernel。异步则 CPU 可以在 GPU 跑第 1 个 kernel 的同时把第 2、3、4 个都丢进队列，CPU 的 launch 时间和 GPU 的计算时间重叠。这是 GPU 编程高性能的基础：让 CPU 永远在 launch、GPU 永远在算，两者谁都不闲着。

三、异步 Launch 与 synchronize（Timer 的命门）

3.1 异步带来的"测不准"

回到 Megatron Timer 为什么必须 sync：

1
2
3
4
start_time = time.time()       # CPU 此刻
C = A + B                      # launch 一个 matmul kernel, CPU 立刻继续
end_time = time.time()         # CPU 几微秒后就到了这里
print(end_time - start_time)   # 几微秒 ← 完全错! GPU 可能还没开始跑

不 sync 时，time.time() 测的是 “CPU 提交 kernel 的时间”，而 GPU 上这个 kernel 可能还在队列里排队、甚至还没开始。测出来的值可能是真实 GPU 计算时间的零头——完全不可信。

3.2 torch.cuda.synchronize() 的作用

1
torch.cuda.synchronize()      # 阻塞 CPU, 等 GPU 把队列里所有 kernel 全跑完

调用后 CPU 卡住，直到 GPU 把当前 stream 上所有已 launch 的 kernel 全部执行完。这时再取 time.time()，才反映 GPU 真实完成时间。

所以 Megatron Timer 的 start/stop 都先 synchronize()：

  • start 前 sync：确保起点是"GPU 之前全干完了"的时刻，不被前面残留 kernel 污染。
  • stop 前 sync：确保终点是"这段 kernel 真跑完"的时刻。

3.3 sync 的代价

synchronize() 会让 CPU 停下来等 GPU，打破异步重叠——CPU 不能在 GPU 跑时继续 launch 后续 kernel，于是 CPU 和 GPU 串行化，性能下降。所以：

  • 生产训练不要乱插 sync（除非必要）。
  • 测时间时才 sync（且要能用 log_level 关掉，见 Megatron Timer 的 DummyTimer）。
  • 需要拿 GPU 结果到 CPU 用时才 sync（如 loss.item() 打印、判断是否 NaN、控制流分支）。

面试金句：“CUDA op 异步 launch，CPU 侧时间戳测的是提交时间不是计算时间，必须 synchronize() 等队列排空才可信；但 sync 会打破 CPU/GPU 重叠拖慢训练，所以只在测时间、取结果到 CPU、做控制流时用，生产里要尽量少。”

3.4 何时该 sync（经验）

场景 要不要 sync
测一段 op 耗时 要（前后各一次）
loss.item() 打印 / 检查 NaN 要（隐式 sync：.item()/.cpu() 会触发）
控制流：if 某条件 launch 不同 kernel 要（拿条件值到 CPU）
正常前向/反向一串 op 不要，保持异步
print 一个 GPU tensor 隐式 sync

注意 .item()、.cpu()、.tolist()、对 GPU tensor 做 print、if 判断 GPU tensor 值——这些操作会隐式触发 sync，是训练慢的常见暗坑。推理/训练代码里循环内 loss.item() 每步都 sync 一次，会显著拖慢。

四、CUDA Stream 与 Event：谁先谁后、能不能并行

4.1 Stream（流）是什么

stream 是 GPU kernel 的执行队列。同一 stream 里的 kernel 严格按 launch 顺序串行执行；不同 stream 的 kernel 可以并行执行（如果 GPU 资源够）。

1
2
stream A:  [k1] → [k2] → [k3]      (串行)
stream B:  [m1] → [m2]             (和 A 里的可以重叠并行)
  • 默认所有 op 走 default stream（stream 0），全串行。
  • 想让两段独立计算重叠（如"算梯度"和"搬数据"重叠），把它们放不同 stream：
    1
    2
    3
    4
    5
    6
    7
    8
    9
    s1 = torch.cuda.Stream()
    s2 = torch.cuda.Stream()
    with torch.cuda.stream(s1):
        # 这些 kernel 进 s1 队列
        grad = model(x)
    with torch.cuda.stream(s2):
        # 这些 kernel 进 s2 队列, 可和 s1 并行
        data.copy_(host_data)
    torch.cuda.synchronize()   # 等两条都跑完

4.2 典型用途：计算/通信重叠、H2D/计算重叠

这是 stream 最有价值的场景：

  • 计算与通信重叠：一边 GPU 算梯度（stream A），一边 NCCL all-reduce 上一段梯度（stream B，通信 kernel 也在 GPU 上但走不同 stream），两者并行——这就是 Megatron DDP “bucket ready 触发 async all-reduce 与反向 overlap” 的底层。
  • Host-to-Device 搬数据和计算重叠：下一 batch 的数据 H2D 拷贝（stream B）和当前 batch 前向（stream A）并行，CPU/PCIe 传输藏进算力里。
  • prefill/decode 分流：vLLM、推理引擎里有时把 prefill 和 decode 放不同 stream 重叠。

4.3 Event（事件）：跨 stream 同步

不同 stream 要同步"等 A 这步完成 B 才开始"，用 event：

1
2
3
4
5
6
7
e = torch.cuda.Event()
with torch.cuda.stream(s1):
    work()
    e.record()        # 在 s1 这里插个标记
with torch.cuda.stream(s2):
    e.wait()          # s2 等到 s1 的标记点再继续
    depend_on_work()

Event 还可用来测 GPU 时间（比 time.time()+sync 更精细）：

1
2
3
4
5
6
7
start = torch.cuda.Event(enable_timing=True)
end   = torch.cuda.Event(enable_timing=True)
start.record()
...  # 一串 op
end.record()
torch.cuda.synchronize()
print(start.elapsed_time(end))   # 毫秒, 只测这段 GPU kernel 时间

Event 记录的是 GPU 时间线上的点，不受 CPU launch 异步干扰，是测 GPU kernel 耗时的正确姿势。Megatron Timer 用 time.time()+sync 是更简单粗粒度的版本，精细 profile 用 Event 或 nsys/ncu。

面试金句：“stream 是 kernel 执行队列，同 stream 串行、跨 stream 可并行；计算/通信重叠、H2D/计算重叠靠多 stream。跨 stream 同步用 Event，Event 也能精确测 GPU kernel 时间（不受 launch 异步干扰）。”

五、Grid / Block / Thread：kernel 的并行结构

一个 kernel 启动时，要告诉 GPU 派多少线程、怎么组织。CUDA 用三层结构：

1
2
3
4
5
6
7
8
9
Grid  ──► 一堆 Block ──► 每个 Block 一堆 Thread

       Grid (整个 kernel 的线程网格)
       ┌──────────┐
       │ Block(0,0)│ Block(1,0) │ ...     ← GridDim, 通常按数据形状分
       │ Block(0,1)│ Block(1,1) │ ...
       └──────────┘
       每个 Block 里:
       Thread(0) Thread(1) ... Thread(255)  ← BlockDim, 通常 128/256
  • Thread：最小执行单元，每个 thread 跑一份 kernel 代码、算一个/几个元素。
  • Block：一组 thread，同一 block 内 thread 可共享 shared memory、可 __syncthreads() 同步。一个 block 跑在一个 SM 上。
  • Grid：所有 block 的集合，跨 SM 调度。

典型配法：每个 thread 算一个输出元素，blockDim=256，gridDim=ceil(N/256)。

关键：算力靠"喂饱 SM"

GPU 有几十上百个 SM，每个 SM 能同时跑成千上万线程。kernel 要派足够多的 thread 才能占满 SM——这叫** occupancy（占用率）。如果 kernel 只派几百个 thread（比如 batch=1 推理），SM 大量空闲，算力浪费，这就是 decode 阶段算力受限不足、访存受限为主**的根源。vLLM 用 continuous batching 把多个请求的 decode token 攒一起喂 GPU、用 CUDA Graph 跳 launch，都是为了解决这个。

SM 内的存储层次（影响 kernel 性能）

1
2
3
4
寄存器 (per-thread, 最快, 几 KB)       ← thread 私有, ~1 cycle
Shared Memory (per-block, 快, ~100KB)  ← block 内共享, 用 __syncthreads 同步
L1/L2 Cache (SM 内/全局)               ← 硬件管
Global Memory (显存 HBM, 慢)           ← 所有人可访问, 带宽高但延迟大

写高性能 kernel 的核心是：把热点数据放 shared memory/寄存器、少访问 global memory（HBM）。这就是 FlashAttention 的精髓——把 attention 的中间 s2s^2 矩阵不写回 HBM、用 shared memory/寄存器在线算 softmax，把 O(s2)O(s^2) HBM 访存降到 O(s)O(s)。

显存篇提过"训练访存受限、推理 decode 访存受限"——kernel 跑多快，不是看你算多快，而是看你从 HBM 读数据快不快。访存受限下，优化方向是"少读 HBM"（融合 kernel、tile+shared memory、减少中间结果落地），不是"算更快"。

六、融合 Kernel：为什么要把多个算子合成一个

6.1 朴素 op 链的问题

一段 Y = activation(BiasAdd(MatMul(X, W)))，朴素实现是三个独立 kernel：

1
2
3
launch matmul      → GPU 算 → 写结果到 HBM
launch bias_add    → 从 HBM 读结果 → 加 bias → 写回 HBM
launch activation  → 从 HBM 读 → 激活 → 写回 HBM

三个问题：

  1. 多次 HBM 读写：中间结果 matmul 输出 被写 HBM 又被下一 kernel 读 HBM，带宽浪费、延迟叠加。访存受限下这是性能杀手。
  2. launch 开销：每个 kernel launch 有固定 CPU 开销（几微秒）+ GPU 启动开销。算子小、launch 多时（decode 阶段每步几十个小算子），launch 开销占比极高。
  3. 不能跨 kernel 优化：每个 kernel 独立编译，没法把"读-算-写"合并优化。

6.2 融合的解法

融合 kernel（fused kernel）：把多个算子合成一个 kernel，中间结果留在寄存器/shared memory，不落地 HBM，一次 launch 跑完。

1
launch fused_kernel  → 线程读 X, W → 算 matmul → 直接在寄存器里 add bias → 激活 → 写最终 Y

收益：

  • 少 HBM 读写（最大头，访存受限下决定性提速）。
  • 少 launch 开销。
  • 更好的数据局部性（在 shared memory 里 tile 复用）。

6.3 训练/推理里的融合 kernel 实例

  • FlashAttention：融合 QK^T → mask → softmax → ×V，把 O(s2)O(s^2) 中间矩阵不落地 HBM。Megatron/vLLM 默认用。
  • FusedMLP：融合 gate × up → activation → × down（SwiGLU 全融合），Megatron fusions/ 下。
  • FusedLayerNorm + bias + residual：Megatron 把 residual = x + dropout(layernorm(x + bias)) 融成一步。
  • Fused Adam（optimizer）：把一阶/二阶矩更新 + 权重更新融成一个 kernel，Megatron/VLLM 用。
  • Fused softmax / rope / cross_entropy：Megatron csrc/、vLLM v1/sample/ops。

6.4 融合的代价与边界

  • 实现复杂：手写 CUDA/Triton 工作量大、可维护性差。
  • 通用性差：融合的特定 op 组合换形状/换激活可能要重写。
  • 不一定都更快：compute-bound 的大 kernel（大 GEMM）本身 HBM 访存占比低，融合收益小；访存-bound 的小算子链融合收益大。

面试金句：“融合 kernel 解决两个问题：访存受限下减少中间结果落地 HBM（性能大头）、launch 受限下减少 launch 次数。FlashAttention/FusedMLP/FusedAdam 都是典型。算 compute-bound 的大 GEMM 融合收益小，访存-bound 的小算子链融合收益大——这就是为什么 decode 阶段（小算子多）特别需要融合和 CUDA Graph。”

七、CUDA Graph：把一串 launch 录成一次 replay

7.1 问题：launch 开销在 decode 里很重

推理 decode 阶段每步只产一两个 token，但每步要 launch 几十个 kernel（mlp/attention/norm/rope/sample…）。每个 launch 几微秒 CPU 开销 × 几十个 × 几百步 = CPU 被 launch 占满、GPU 等待 launch。这就是 launch-bound。

7.2 CUDA Graph 的解法

CUDA Graph：在"录制"阶段把一串 kernel launch 记录成一个图结构（依赖关系），"重放"阶段一次 cudaGraphLaunch 把整图丢给 GPU，CPU 只需一次 launch 调用，GPU 自己按图依赖调度。

1
2
录制(capture): 跑一遍 op 链 → 记录成 graph(节点=kernel, 边=依赖)
重放(replay):  一次 cudaGraphLaunch → GPU 按图跑完全部 kernel

收益：

  • CPU launch 开销从几十次降到 1 次——decode 延迟大幅下降。
  • GPU 自己管依赖，减少 CPU↔GPU 同步。

代价：

  • 图的输入/输出地址固定，shape 要固定（变 shape 要重新抓图或抓多个图）。
  • 要占额外显存放图。
  • 动态控制流（if/while 依赖运行时数据）难进图。

7.3 工程实例

  • vLLM：v1/cudagraph_dispatcher.py + worker/gpu/cudagraph_utils.py——启动时为若干个固定 num_tokens 组合抓图，运行时 dispatch_cg_and_sync_dp 按实际 token 数选匹配图 replay。decode 受益最大，prefill（shape 多变）收益小。
  • 静态形状训练 / inference：很多推理引擎对固定 shape 模型用 CUDA Graph 加速。

和融合 kernel 的区别：融合是"把多个算子合成一个 kernel"（减 HBM 访存+减 launch），CUDA Graph 是"把多个 kernel launch 合成一次提交"（只减 launch，不改 kernel）。两者互补——融合解决访存+launch，Graph 解决 launch。decode 慢两者都要用。

八、怎么写 kernel：Triton 与 CUDA C

8.1 两条路

  • CUDA C/C++：原生 NVIDIA，性能最高、控制最细，但难写难调。Megatron csrc/、vLLM csrc/ 里大量手写 CUDA kernel。
  • Triton：OpenAI 出的 Python DSL（@triton.jit），编译成 GPU 代码，比 CUDA 易写、性能接近，FlashAttention 早期版、很多 fused kernel 现在用 Triton 写。生态火。
1
2
3
4
5
6
7
8
9
10
11
12
# Triton 写个 vector_add, 比 CUDA 简洁
import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, z_ptr, N, BLOCK: tl.constexpr):
    pid = tl.program_id(0)                       # 我是第几个 block
    offs = pid * BLOCK + tl.arange(0, BLOCK)     # 这个 block 管哪些元素
    mask = offs < N
    x = tl.load(x_ptr + offs, mask=mask)
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(z_ptr + offs, x + y, mask=mask)

8.2 什么时候要自己写

  • 框架没现成融合 kernel、且是性能热点（attention、mlp、optimizer、采样）。
  • 特殊算子（MLA、MoE dispatch、量化）。
  • 极致性能优化（cudnn/cutlass 调不动的特殊 shape）。

但优先用现成的（FlashAttention、cudnn、cutlass、Triton 模板），手写 kernel 是最后手段——调到最优极费时。

九、怎么调 kernel 性能：profiling

写完/优化 kernel 要测、要找瓶颈：

  • torch.cuda.Event：粗粒度测段 GPU 时间（前面讲过）。
  • torch.profiler：导出 Chrome trace，看每个 kernel 的 wall-clock、占空比、launch 间隔。
  • nvidia-smi：看 GPU 利用率、显存、温度（粗粒度，util 高不代表没瓶颈）。
  • nsys（Nsight Systems）：系统级 timeline，看 CPU launch 和 GPU kernel 的时间线、stream 并行情况、有没有气泡、sync 在哪卡住。最该会的工具。
  • ncu（Nsight Compute）：kernel 级细粒度，看每个 kernel 的算力/访存/占用率/瓶颈（compute bound 还是 memory bound）、给出优化建议。调单个 kernel 用它。

典型调优流程：

  1. nsys 看哪段有气泡、哪个 kernel 占时最长。
  2. 对最耗时的 kernel 用 ncu 看是 compute bound 还是 memory bound。
  3. memory bound → 减少 HBM 访存（融合、tile、shared memory、向量化读取）；compute bound → 用 tensor core、减冗余计算。
  4. launch bound → CUDA Graph 或融合 kernel。

面试金句：“调 GPU 性能先 nsys 看 timeline 找最长/有气泡的 kernel，再 ncu 看单个 kernel 是 compute bound 还是 memory bound，对症下药——memory bound 减 HBM 访存（融合/tile），compute bound 上 tensor core，launch bound 用 CUDA Graph。nvidia-smi util 高不等于没瓶颈，可能正卡在访存。”

十、串起训练/推理：kernel 视角看性能

把这套知识用到你熟悉的训练/推理场景：

训练（Megatron）

  • 前向/反向一串 kernel 异步 launch，CPU 不停发，GPU 不停算，两者重叠——这就是高性能的基础。
  • all-reduce/all-gather 通信也是 kernel（NCCL kernel，在 GPU 上跑），放单独 stream 和计算 kernel 重叠（DP overlap）。
  • Timer 要 synchronize() 才能测准各段 kernel 时间，但 sync 拖慢训练所以用 log_level 控制。
  • 融合 kernel（FusedMLP/FlashAttention/FusedAdam）减访存+减 launch，是访存受限训练的核心优化。

推理（vLLM）

  • prefill 算力受限：大 GEMM kernel，算得满，launch 开销占比小，CUDA Graph 收益小。
  • decode 访存受限+launch 受限：每步几十个小 kernel、batch 小 SM 空闲——continuous batching 攒 token 喂饱 SM + 融合 kernel + CUDA Graph 跳 launch 三件套。
  • 多 stream：prefill/decode 分流、KV 搬运和计算重叠。
  • 异步 scheduler：CPU 调度下一步时 GPU 还在算上一步，靠的就是 launch 异步——CPU 在 launch 间隙做调度，不阻塞 GPU。

三种瓶颈（务必分清）

瓶颈类型 含义 优化方向
compute bound 算力打满，访存不是瓶颈 上 tensor core、减冗余计算、更高精度融合
memory bound 访存打满，算力空闲（最常见） 减 HBM 读写、融合 kernel、tile+shared memory
launch bound launch 开销占主导（小算子多） CUDA Graph、融合 kernel

大模型训练/推理绝大多数时间是 memory bound（显存篇讲过访存受限），所以"减访存"是优化主线——融合、tile、量化（少读权重）、KV cache 复用（少重算）全是减访存。

十一、面试速答清单

Q1：GPU kernel 是什么？和 OS kernel 区别？

GPU kernel 是一段在 GPU 上由成千上万线程并行执行的函数；CPU 调用它（launch）只是把派工单丢进队列，真正执行的是 GPU。OS kernel 是常驻 ring0 管硬件的操作系统核心，跑在 CPU。两者同名但完全不同。

Q2：为什么 PyTorch 测 op 时间必须 cuda.synchronize()？sync 的代价？

CUDA op 异步 launch，time.time() 测的是 CPU 提交 kernel 的时间，GPU 可能还没开始跑，数据不可信。synchronize() 阻塞 CPU 等 GPU 队列排空才反映真实计算时间。代价是打破 CPU/GPU 重叠、串行化拖慢，所以只在测时间、取结果到 CPU（.item()/.cpu()）、做控制流时用，生产尽量少。

Q3：stream 是什么？怎么用它做计算/通信重叠？

stream 是 kernel 执行队列，同 stream 串行、跨 stream 可并行。把独立的两段计算放不同 stream（如算梯度放 stream A、NCCL all-reduce 放 stream B）就能让它们并行——这是 DDP 反向 overlap 的底层。跨 stream 同步用 Event，Event 也能精确测 GPU kernel 时间。

Q4：为什么要融合 kernel？收益和代价？

朴素 op 链每个 kernel 把中间结果写 HBM、下个 kernel 又读 HBM，访存浪费+延迟叠加；launch 多次开销大。融合把多算子合成一个 kernel，中间结果留寄存器/shared memory 不落地 HBM、一次 launch 跑完。收益是减访存（memory bound 下决定性）+减 launch；代价是难写难维护、通用性差。FlashAttention/FusedMLP/FusedAdam 是典型。

Q5：CUDA Graph 解决什么？和融合 kernel 区别？

解决 launch bound——把一串 kernel launch 录成图，运行时一次 cudaGraphLaunch 提交，CPU 只调一次 launch。区别：融合是"多算子合成一个 kernel"（减访存+减 launch），Graph 是"多 kernel launch 合成一次提交"（只减 launch，不改 kernel）。decode 阶段 launch 受限两者互补都要用。vLLM 为固定 num_tokens 抓图、按 token 数 dispatch 选图。

Q6：memory bound / compute bound / launch bound 怎么区分、怎么优化？

memory bound 是访存打满算力空闲（大模型最常见），优化靠减 HBM 读写（融合/tile/shared memory/量化）；compute bound 是算力打满，优化靠 tensor core/减冗余计算；launch bound 是小算子多 launch 开销主导，靠 CUDA Graph/融合。先 nsys 看 timeline 找最长/有气泡的 kernel，再 ncu 看单 kernel 是 compute 还是 memory bound 对症下药。nvidia-smi util 高不代表没瓶颈，可能正卡访存。

Q7：为什么 decode 阶段 GPU 利用率低、特别需要优化？

decode 每步只产一两个 token，kernel 规模小、batch 小，SM 喂不饱（occupancy 低）、算力空闲，同时几十个小算子 launch 开销占比高——既 memory bound 又 launch bound。所以 vLLM 用 continuous batching 攒 token 喂饱 SM、融合 kernel 减访存、CUDA Graph 跳 launch 三件套。prefill 是大 GEMM 算力受限，这些优化收益小。

Q8：FlashAttention 融合了什么、为什么快？

融合 QK^T → mask → softmax → ×V，把 attention 中间 O(s2)O(s^2) 矩阵不写回 HBM、用 shared memory/寄存器在线分块算 softmax，把 HBM 访存从 O(s2)O(s^2) 降到 O(s)O(s)。memory bound 下减访存是决定性收益，还顺带省了 O(s2)O(s^2) 显存。这是融合 kernel 减访存最经典的例子。

十二、一张图收口

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
CPU(派单)  ─launch(kernel)─►  [GPU stream 执行队列]  ─► SM 上成千线程并行跑 kernel
   │                          同 stream 串行, 跨 stream 可并行(Event 同步)
   │
   │ 要结果/测时间: cuda.synchronize() 阻塞 CPU 等队列排空
   ▼
GPU 结果

kernel 结构:  Grid > Block > Thread   (thread 私有寄存器, block 共享 shared memory, 全局 HBM)
访存层次:     寄存器(~1cyc) > shared mem(~100cyc) > HBM(慢,带宽高延迟大)
              ↑ 高性能 kernel 把热点数据放寄存器/shared, 少读 HBM

三大瓶颈:
  memory bound(最常见, 访存打满) → 融合/tile/shared mem/量化(减 HBM)
  compute bound(算力打满)       → tensor core/减冗余
  launch bound(launch 主导)     → CUDA Graph/融合

调优:  nsys(timeline 找最长/气泡) → ncu(单 kernel compute 还是 memory bound) → 对症

主线一句话：GPU kernel 是被成千线程并行执行的函数、CPU 异步 launch 派单进 stream 队列、要测准或拿结果才 sync（sync 拖慢所以要少用）；高性能靠"喂饱 SM + 少读 HBM + 少 launch"——融合 kernel 解决访存+launch、CUDA Graph 解决 launch、多 stream 做计算通信重叠；大模型绝大多数时间 memory bound，所以"减访存"是优化主线。 把 launch 异步、sync、stream、融合、三种瓶颈讲顺，GPU 性能这关就稳了。

参考资料

  • NVIDIA CUDA C++ Programming Guide（grid/block/thread、stream、event、memory hierarchy）
  • NVIDIA Nsight Systems / Nsight Compute 文档
  • FlashAttention: Fast and Memory-Efficient Exact Attention Dao et al., 2022
  • Triton: https://triton-lang.org
  • 与本文 Megatron Timer、显存计算法则、vLLM 源码精读、RDMA 篇交叉对照