推理服务框架与调度
大模型推理从单次 forward 到高吞吐线上服务,需要一整套调度和内存管理机制。本文覆盖 vLLM/SGLang 体系中最高频的面试考点。
Q:FlashAttention 为什么能加速,计算过程是什么
来源:京东 / AI Infra 一面 · 快手 / AI Infra 一面 · 快手 / AI Infra 二面 · 阿里国际 / AI Infra 实习 · 混元 / AI Infra · 面经总结
普通回答:FlashAttention 通过 tiling 减少 HBM 访问来加速。
更好的回答:
标准 Attention 的问题:
Q (N×d) × K^T (d×N) → S (N×N) → softmax → P (N×N) × V (N×d) → O (N×d)
- 中间矩阵 S 和 P 大小 O(N²),必须写回 HBM
- N=4096, heads=32 → S 和 P 各占 ~2 GB,反复读写 HBM 是瓶颈
FlashAttention 核心思想——IO-Aware Tiling:
不存储完整的 N×N attention matrix,改为分块计算 + online softmax:
将 Q 分成 Br 大小的块,K/V 分成 Bc 大小的块
for each Q_block (外层循环):
初始化 O_block = 0, m = -inf, l = 0 (running max 和 running sum)
for each K_block, V_block (内层循环):
1. 从 HBM 加载 Q_block, K_block, V_block 到 SRAM (Shared Memory)
2. 计算局部 S_block = Q_block × K_block^T (在 SRAM 中)
3. Online Softmax 更新:
- m_new = max(m_old, rowmax(S_block))
- 修正旧累加: O_block *= exp(m_old - m_new) * (l_old / l_new)
- 累加新块: O_block += exp(S_block - m_new) × V_block
- 更新 l_new
4. 不写回 S_block 到 HBM!
写回 O_block 到 HBM
为什么能加速:
- HBM 访问量从 O(N²) 降到 O(N²d/M):M 是 SRAM 大小,分块后每块在 SRAM 内完成
- 不存储 N×N 矩阵:空间从 O(N²) 降到 O(N)
- 实际收益:2-4× wall-clock 加速(Attention 部分),长序列收益更大
V1 → V2 → V3 的改进:
| 版本 | 关键改进 |
|---|---|
| V1 | 基础 tiling + online softmax |
| V2 | 外层循环改为 Q(非 K/V)→ 减少非矩阵乘 FLOPs,更好利用 Tensor Core;前向减少一次 rescale |
| V3 | 利用 Hopper 的 TMA + warpgroup,producer-consumer 异步流水线,WGMMA 指令 |
FlashDecoding(decode 阶段的 FlashAttention):
- Decode 时 Q 只有 1 个 token,但 KV 很长
- 在 KV 的 seq_len 维度并行(多个 block 各算一段),最后 combine kernel 合并
- Combine kernel 占总耗时很小(<5%),因为只是对 block 数做 reduction
考察点:核心是 online softmax + tiling 避免 N×N 矩阵写回 HBM,追问通常到 V2 为何外层改 Q。
Q:PagedAttention 具体是什么 / 为什么效果好
来源:京东 / AI Infra 一面 · 字节 / AI Infra 后端 · 快手 / AI Infra 实习 · 美团北斗 / AI Infra · 小厂 / AI Infra
普通回答:PagedAttention 像操作系统的虚拟内存一样管理 KV Cache。
更好的回答:
问题背景:
- KV Cache 是推理最大的显存消耗(7B 模型 batch=32 可占 40+ GB)
- 传统方式为每个请求预分配最大序列长度的 KV 空间 → 巨大浪费(平均利用率仅 20-40%)
- 不同请求长度不同,预分配导致内存碎片,batch size 受限
PagedAttention 设计:
借鉴 OS 虚拟内存的分页机制:
- 将 KV Cache 分成固定大小的 block(如 16 tokens 一块)
- 维护一个 block table(类似页表)记录每个请求的逻辑 block → 物理 block 映射
- 物理 block 从全局 block pool 按需分配,用完归还
逻辑视图(请求A): [block0] → [block1] → [block2] → ...
↓ ↓ ↓
物理显存: [slot 5] [slot 12] [slot 3] ← 不连续!
为什么效果好:
- 消除碎片:物理 block 不需要连续,任意空闲 slot 都能用 → 显存利用率接近 100%
- 按需分配:生成几个 token 就分配几个 block,不预分配最大长度
- Prefix Cache / Copy-on-Write:相同 prefix 的请求可共享 block(如 system prompt),写时复制
- 结果:同等显存下 batch size 提升 2-4×,吞吐翻倍
Attention 计算变化:
- 标准 Attention 要求 KV 在内存中连续
- PagedAttention kernel 需要根据 block table 做间接寻址——按 block 读取 KV,计算局部 attention 后合并
- 计算量不变,多了一次查表开销(可忽略)
vLLM 如何结合:
- Scheduler 管理 block 分配/回收/抢占
- 抢占(preemption):显存不够时 swap out 某个请求的 KV block 到 CPU,后续 swap in 恢复
- Prefix Caching:相同 prompt 前缀的 block 可以复用(hash 匹配)
考察点:理解分页思想如何解决 KV Cache 碎片和预分配浪费,以及对 attention kernel 的影响。
Q:Chunk Prefill 是什么 / 为了解决什么问题
来源:京东 / AI Infra 一面 · 鹅厂 / AI Infra 实习
普通回答:把 prefill 分成多个 chunk 执行。
更好的回答:
问题:
- 长 prompt 的 prefill 一次计算耗时很长(如 32K tokens 可能几秒)
- 在此期间,正在 decode 的请求被阻塞 → 它们的 TPOT(Time Per Output Token)剧增
- 用户体验:已经在对话中的请求突然”卡住”
Chunk Prefill 解决方案:
将一个长 prefill 请求切分为多个 chunk,与 decode 请求交替执行:
时间线(无 Chunk Prefill):
|---- Prefill A (32K tokens, 3s) ----| Decode B | Decode B |
↑ B 被阻塞 3 秒
时间线(有 Chunk Prefill):
|Prefill A chunk1|Decode B|Prefill A chunk2|Decode B|Prefill A chunk3|Decode B|
↑ B 每隔一个 chunk 就能执行一次
关键设计点:
- Chunk size 选择:太小 → prefill 效率低(每块 kernel launch overhead);太大 → decode 延迟高
- 通常 chunk_size = 512~2048 tokens
- 不是 A 的 prefill 和 A 的 decode 交替——是 A 的 prefill chunk 和 其他请求 B 的 decode 交替
- 一个 iteration 内:scheduler 将一个 prefill chunk + 多个 decode 请求组成一个 batch
与 PD 分离的关系:
- PD 分离:prefill 和 decode 分到不同 GPU 节点
- Chunk Prefill:在同一节点上通过时间片共享解决延迟问题
- 两者可以结合:prefill 节点内部也可以用 chunk prefill 来服务多个请求
SGLang 的实现:
- RadixAttention + Chunk Prefill 结合
- Scheduler 每个 step 决定分配多少 budget 给 prefill vs decode
考察点:理解 prefill 阻塞 decode 的问题本质,以及 chunk granularity 的 trade-off。
Q:Continuous Batching / vLLM 的调度机制
来源:小厂 / AI Infra RL · 快手 / AI Infra 实习 · 阿里国际 / AI Infra 实习
普通回答:Continuous Batching 是动态加入和移除请求。
更好的回答:
Static Batching 的问题:
- 传统方式:凑齐一个 batch 的请求,等所有请求都生成完(最长的那个)才开始下一个 batch
- 短请求早完成但要等长请求 → GPU 利用率低,延迟高
Continuous Batching(iteration-level scheduling):
- 每个 decode step 都是一个调度点
- 一个请求完成(生成 EOS)→ 立即移出 batch,新请求立即加入
- batch 大小动态变化,GPU 始终满载
vLLM Scheduler 流程:
每个 step:
1. 检查 waiting 队列有无新请求需要 prefill
2. 检查 running 队列中有无请求完成(EOS / max_tokens)
3. 检查显存是否够分配新 block
4. 如果不够 → preempt(抢占优先级最低的请求,swap out 其 KV Cache)
5. 组装本步 batch: prefill tokens + decode tokens
6. 执行 forward → 输出 token
7. 更新状态
抢占(Preemption)策略:
- 显存不够时,选择一个请求暂停
- Swap:将其 KV Cache 从 GPU 搬到 CPU(后续 swap in 恢复)
- Recompute:直接丢弃 KV Cache,恢复时重新 prefill
- 策略:FCFS(先来先服务)或 priority-based
SGLang 的 RadixAttention:
- 用 Radix Tree 管理所有请求的 KV Cache prefix
- 自动识别共享 prefix 并复用(比 vLLM 的 hash-based prefix cache 更灵活)
- 对多轮对话、few-shot 场景效果好
考察点:理解 iteration-level 调度的核心优势,以及显存不够时的抢占策略。
Q:Speculative Decoding(投机采样)原理
来源:面经总结 · 小厂 / AI Infra
普通回答:小模型先生成 draft,大模型验证。
更好的回答:
动机:
- Decode 阶段是 memory-bound,每步都要读全部权重但只算 1 token
- 如果能一次验证多个 token → 有效利用 compute → 加速
流程:
1. Draft Phase: 小模型(如 68M 参数)快速自回归生成 K 个 draft tokens
- 小模型速度极快(参数少 → 读取少)
- 生成 draft tokens: t1, t2, ..., tk 以及对应概率 q(t_i)
2. Verify Phase: 大模型一次性 forward 处理 K+1 个 position
- 输入: [context, t1, t2, ..., tk]
- 输出: 每个位置的概率分布 p(t_i)
3. Accept/Reject (保证分布不变):
for i = 1 to K:
if p(t_i) >= q(t_i): accept t_i
else: 以概率 p(t_i)/q(t_i) accept, 否则 reject 并 resample from (p - q)+
第一个 reject 之后的所有 token 丢弃
4. 期望加速: 如果 acceptance rate = α, 期望每次生成 1/(1-α) 个 token
关键性质:
- 数学上保证输出分布与大模型独立解码完全一致(无质量损失)
- 加速倍数取决于 draft model 与 target model 的分布接近程度
变体:
- Medusa:不用独立 draft model,在大模型上加多个 head 并行预测未来 token
- EAGLE:用大模型的 hidden states 训练轻量 draft head
- Self-speculative:大模型跳层做 draft
- MTP(Multi-Token Prediction):DeepSeek V3 的方案,训练时就学多步预测
系统挑战:
- Draft model 和 target model 需要共享 KV Cache 管理
- Verify 的 batch 维度不规则(每个请求 accept 数量不同)
- vLLM 和 SGLang 的实现差异在于 KV Cache 管理和 batch 组织方式
考察点:理解 accept/reject 保证分布不变的数学原理,以及与 MTP/Medusa 的对比。
Q:CUDAGraph 是什么 / 为什么更适合 decode 阶段
来源:小马智行 / AI Infra 实习
普通回答:CUDAGraph 把一系列 kernel 录制下来重放,减少 launch overhead。
更好的回答:
问题:
- 每次 kernel launch 有 CPU 端开销(~5-10μs/kernel)
- 大模型一次 forward 有数百个小 kernel
- Decode 阶段计算量很小(每步只处理 1 token),但 kernel 数不变
- 当计算时间接近 launch overhead 时,CPU 成为瓶颈
CUDAGraph 原理:
// 录制阶段(只执行一次)
cudaStreamBeginCapture(stream);
kernel_A<<<...>>>();
kernel_B<<<...>>>();
kernel_C<<<...>>>();
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&graphExec, graph);
// 重放阶段(每次 decode step)
cudaGraphLaunch(graphExec, stream); // 一次 launch 执行整个图
- 录制时不真正执行,只记录 kernel 序列和参数
- 重放时整个图作为一个单元提交给 GPU → CPU 只有一次 launch 开销
为什么更适合 decode:
- Decode 每步 input shape 固定(batch_size × 1)→ 图可以复用
- Prefill 的 seq_len 每次不同 → 每次 shape 变化需要重新录制 → 不适合
- Decode 的单步计算快,launch overhead 占比大 → CUDAGraph 收益明显
为什么占更多显存:
- CUDAGraph 录制时绑定了所有 kernel 的输入输出地址
- 这些 tensor 在图的生命周期内不能被释放或移动
- 等于把 decode 一步需要的所有中间 tensor “钉住”了
- 对比正常执行:PyTorch 可以在 kernel 间复用显存
vLLM 中的实践:
- 为不同 batch_size 预录制多个 CUDAGraph
- Decode 时根据当前 batch size 选择对应的 graph 重放
- Batch size 变化时切换 graph(有少量 overhead)
考察点:理解 CUDAGraph 的 trade-off:减少 launch overhead vs 增加显存占用。
Q:Prefix Cache 机制
来源:字节 / AI Infra 后端
普通回答:缓存相同 prefix 的 KV Cache 避免重复计算。
更好的回答:
应用场景:
- System prompt:同一个应用的所有请求共享相同的 system prompt(可能几千 token)
- Few-shot examples:相同的 few-shot 示例
- 多轮对话:前几轮的上下文
实现方式:
- Hash-based(vLLM):
- 将 token sequence hash 为 key
- 全局维护
hash → KV Cache blocks的映射 - 新请求来时,检查 prefix 的 hash 是否命中 → 命中则直接复用 block
- Radix Tree(SGLang):
- 用 Radix Tree(前缀树)组织所有活跃/缓存的 KV Cache
- 自动找到最长公共前缀
- 支持更细粒度的 prefix 共享(不限于完整 prompt 匹配)
收益:
- 4K token 的 system prompt × 100 并发 → 节省 99% 的重复 prefill 计算
- 多轮对话第 N 轮只需 prefill 新消息部分
LRU 淘汰:
- 显存有限,缓存的 prefix 需要 LRU 淘汰
- 热门 prefix(高命中率)常驻,冷 prefix 被回收
考察点:理解 prefix cache 的收益场景和 hash vs radix tree 的实现区别。
Q:AF 分离是什么 / 为什么有了 PD 分离还要 AF 分离
来源:快手 / AI Infra 一面
普通回答:AF 分离是把 Attention 和 FFN 分开。
更好的回答:
PD 分离的局限:
- PD 分离解决了 prefill(compute-bound)和 decode(memory-bound)的资源争抢
- 但 decode 节点内部,每个 Transformer 层仍包含 Attention 和 FFN,它们特征不同:
- Attention:需要读 KV Cache(随 seq_len 线性增长),是 decode 的显存瓶颈
- FFN:固定大小的权重矩阵乘,计算模式固定
AF 分离(Attention-FFN Disaggregation):
- 将 Attention 计算和 FFN 计算分配到不同的硬件
- Attention 节点:配大显存(存 KV Cache),带宽优先
- FFN 节点:配高算力,对显存需求相对固定
为什么有价值:
- MoE 模型中 FFN 参数量巨大(DeepSeek V3 有 256 个 Expert),但每 token 只激活 8 个
- Attention 参数少但 KV Cache 大(长序列场景)
- 分离后:FFN 节点可以用 EP 做专家并行,Attention 节点专注 KV Cache 管理
Mooncake 系统的 KV-Cache Centric 设计:
- KV Cache 作为独立的分布式存储层
- Prefill 节点生成 KV Cache → 存入分布式 KV Store
- Decode 节点从 KV Store 拉取需要的 KV Cache
- 好处:KV Cache 可以跨请求共享、按需加载、独立扩缩容
考察点:理解 PD 分离只是第一步,AF 分离进一步细分了 decode 内部的异构计算需求。
Q:如何设计高吞吐低延迟的推理服务架构
来源:网易 / AI Infra 校招 · OPPO / AI Infra 实习 · 阿里云 / AI Infra
普通回答:用 vLLM 部署,加量化和 Tensor Parallel。
更好的回答:
系统设计需要考虑的层面:
1. 请求调度层:
- Continuous Batching(iteration-level scheduling)
- 优先级调度:延迟敏感请求优先
- Admission control:过载时拒绝新请求而非所有请求变慢
2. 显存管理层:
- PagedAttention:消除 KV Cache 碎片
- Prefix Cache:复用公共前缀的 KV Cache
- 动态 batch size:根据当前 KV Cache 占用调整并发数
3. 计算优化层:
- FlashAttention:减少 Attention 的 HBM 访问
- 量化(INT8/FP8):减少带宽需求 + 利用低精度 Tensor Core
- CUDAGraph:decode 阶段减少 launch overhead
- Speculative Decoding:提升单请求 decode 速度
4. 分布式层:
- Tensor Parallel(节点内 NVLink)
- Pipeline Parallel(跨节点)
- PD 分离 + Chunk Prefill
5. 可观测性:
- 指标:TTFT(首 token 延迟)、TPOT(每 token 延迟)、Throughput(tokens/s)
- Profiling:哪个阶段是瓶颈
- SLO 监控:延迟 P99 是否达标
延迟 vs 吞吐 Trade-off:
- 大 batch → 高吞吐但每请求延迟增加(排队 + 共享 compute)
- 小 batch → 低延迟但 GPU 利用率低
- 解法:根据 SLO 动态调整 batch 上限
考察点:系统性地覆盖从调度到计算到分布式的全链路,而不是单点技术。