近年来，主流大语言模型架构正经历从标准多头注意力（MHA）向多查询注意力（MQA）、分组查询注意力（GQA）及多头潜在注意力（MLA）的范式转移。这一演进的核心驱动力在于解决自回归解码阶段 KV Cache 带来的显存容量与带宽瓶颈（即“内存墙”问题），旨在通过降低访存开销来显著提升推理吞吐量与长文本处理能力。

当前 transformer 的组件中，尤其是 attention 部分从硬件适配和自身结构上有多方面的优化。如 Qwen3-Next、Llama 的 GQA，DeepSeek 的 MLA，minimax-M2 的 MHA 等。那这些优化到底是出于何种目的？

首先简单赘述一下 LLM 在 inference 阶段的两个核心步骤：prefill和decode。

prefill阶段，模型接收完整的输入序列，并行计算输入词元的

decode阶段，模型接收部分输入序列，通过自回归解码生成输出序列。在这个过程中，模型需要对每个位置的 token 进行注意力计算，以理解输入序列的上下文。

Read More

Context、RAG、Memory 不是互斥，而是互补。上下文工程用于会话即时优化，RAG 用于把权威文档注入生成，长期记忆用于跨会话个性化。

Context、RAG、Memory 对比

简言之，context engineering 是通过控制 prompt 本身的格式和指令来让用户“说对话”，而 RAG 是通过引入外部证据来抑制幻觉，也就是“说对事”。Memory 则是通过持久化状态来构建个体认知，让用户“记得谁在说话”。

常见实现方式

通用的步骤： 1. 先用上下文工程 —— 0成本提效，先跑通闭环 2. 再上RAG —— 只对关键文档启用，避免全库检索 3. 最后建记忆 —— 仅存“用户身份+关键状态+偏好”，非全部对话 4. 缓存+压缩+遗忘 —— 成本杀手三件套 5. 隐私即设计 —— GDPR（欧盟标准真是严格）不是补丁，是架构底座

其中涉及的工具： - 上下文：LangChain（模板）、PromptLayer - RAG：LlamaIndex + Chroma/FAISS + GPT-4-turbo - 记忆：Redis（键值）、Pinecone（向量）、Zep（带策略）

Context 和 Prompt

其中 context 和 prompt 是互补的，context 是会话即时优化，属于 数据，prompt 是会话长期优化，属于 指令。我们给模型的不是 “所有信息”，而是 “怎么用关键信息”。

使用示例：会议议程生成工作流

用户问：“生成下周会议议程，基于上次记录和政策文档，标注来源。”

Context（原材料池，非直接输入）

• 记忆_2024-05-10：上次会议“预算审批延迟”
• RAG_003：《差旅报销政策 v3》第 2 条：“超 5k 需 CTO 审批”
• API_日程：张三 6/3 14:00–15:00 不可用
• 冗余原文：2000字会议纪要全文 （不进 prompt）

Prompt（最终输入模型的文本）

[SYSTEM]
你是一个专业会议助理。输出格式严格为：
1. 每项议程：【标题】（来源：记忆_X / 文档_Y）
2. 总时长≤30分钟
3. 冲突时间自动跳过
4. 无来源内容，禁止生成

[USER]
请基于以下摘要生成议程：
---
记忆_2024-05-10：预算审批因数据缺失延迟，需补充财务报表。
文档_003：超5k预算需CTO审批。
API_日程：张三 6/3 14:00–15:00 不可用。
---
请生成30分钟议程，按优先级排序，每项标注来源。

在后面的内容中，简略使用 Context 来替代 “context+prompt” 的组合，实际指输入给 AI/Agent 的内容。

Context、RAG、Memory 在有限资源下的取舍

原则

• Context/RAG/Memory → 对应 可控性 | 可验证性 | （记忆）连续性
• 可控性 | 可验证性 | 连续性 → 先选其二，三者全满必“费钱&费事&费人”
• 限制：上下文窗口、语义冲突、延迟成本、合规摩擦
• 三者全部进步，需要技术底层全面革新

每对组合 = 一种系统架构

其中记忆是“主观历史”，RAG 是“客观事实”，可控是“规则约束”——三者在语义层天然对抗

折中的办法 - 是“逼近”，不是“解决”

1. 分层召回：先记忆，再RAG，择优注入（不全塞）
2. 动态摘要：记忆压缩为语义标签，减少 token 消耗，过多反而质量下降（也是一种通胀）
3. 可信度引擎：给每条证据打分（来源权威性 + 时间新鲜度 + 用户确认）
4. 异步预载：后台缓存高频记忆/文档，不卡用户响应
5. 后处理控制：模型输出后，用规则层过滤/重写，保留可控性

所有方案共性：不追求“同时全量”，而追求“智能选择”

产品在一定条件下，实现“可接受解”的路线图

永远保留一层策略审计层（极大可能性是类Excel的产品呈现形式） —— 即使选了“可验证+记忆”，也别让模型自由发挥。

Context、RAG、Memory 三元结构

S = (C, R, M)—— 会话上下文、外部证据、长期记忆 ——三者并行检索 → 融合裁剪 → 有痕写回

伪代码表示核心流程和接口

用户提问 → 并行查C/R/M → 按分数和token预算裁剪 → 拼成带来源的prompt → 生成 → 选中结果触发写M（需确认）

retrieve_c()           # 会话内最近对话/API结果
retrieve_r(query, k)   # 外部知识库：检索
retrieve_m(user, query, k)  # 个人记忆：语义/时序双模式
assemble_prompt(context, rag, mem, budget)  # 按分数+预算拼prompt（带来源标签）
write_mem(user, item, confirm=False)  # 默认需用户确认，否则不写

示例说明 Context/RAG/Memory

客服+产品手册+个人偏好场景

[SYSTEM]
你是一个智能客服助手。回答必须：
- 仅使用下方标注来源的信息
- 无法确认时，回答“根据现有信息无法确定”
- 每个结论后标注来源（如：记忆_3 / 文档_P12）
- 避免重复、避免假设、避免情感化语言

[USER]
用户问题：{user_query}

[CONTEXT]
--- 记忆_1（用户偏好） ---
{user_preference_summary}
--- 文档_P12（产品手册） ---
{doc_snippet_1}
--- 文档_P15（FAQ） ---
{doc_snippet_2}
--- 记忆_5（历史工单） ---
{ticket_summary}
--- （最多再加 1 条 Top-k 检索结果） ---

请按上述信息作答，不要扩展。

回填优先级策略：

1. 用户偏好（影响语气/定制）
2. 当前问题直接相关文档
3. 历史相似工单
4. RAG 检索 Top-k（仅当匹配度 > 阈值）

Read More

在当今的 LLM 应用中，Agent 是一个至关重要的概念。它能帮助 LLM 完成代码生成、问题解答、多轮对话等复杂任务。然而，在众多 LLM 应用中，最成功的那些往往不依赖于复杂的架构或特殊的计算库，而是采用简单、通用的 Agent 模式。

什么是Agent

“Agent”可以有多种定义方式。一些客户将 Agent 定义为能够长时间独立运行、使用各种工具来完成复杂任务的完全自主的系统。另一些客户则用这个术语来描述遵循预定义 workflow 的、更具规定性的实现。在 Anthropic，我们将所有这些变体归类为 agentic systems，但在 workflows 和 agents 之间做出了一个重要的架构区分：

• Workflows 是指大语言模型（LLM）和工具通过预定义的代码路径被编排起来的系统。
• 另一方面，Agents 则是指大语言模型（LLM）动态地指导其自身流程和工具使用，并保持对任务完成方式的控制权的系统。

什么时候使用或者不使用 Agent

在使用 LLM 构建应用程序时，我们建议寻找尽可能简单的解决方案，仅在需要时才增加复杂性。这可能意味着根本不构建 agentic systems。Agentic systems 通常以延迟和成本换取更好的任务性能，你应该考虑这种权衡在何时是合理的。

当需要更高的复杂性时，对于定义明确的任务，workflows 提供了可预测性（predictability）和一致性（consistency），而当需要大规模的灵活性和模型驱动的决策时，agents 则是更好的选择。然而，对于许多应用程序来说，通过检索和上下文示例来优化单次 LLM 调用通常就足够了。

何时以及如何使用 Frameworks

当前有许多 frameworks 可以用来帮助 agentic systems 更简单地被用到，包括：

• The Claude Agent SDK;
• Amazon Bedrock’s AI Agent framework;
• Rivet, a drag and drop GUI LLM workflow builder; and
• Vellum, another GUI tool for building and testing complex workflows.

这些框架通过简化标准的底层任务（如调用 LLMs、定义和解析 tools、以及将调用链接在一起）来让起步变得容易。然而，它们常常会创造出额外的抽象层，这可能会掩盖底层的提示（prompts）和响应，从而使调试变得更加困难。它们也可能诱使你在一个更简单的设置就足够的情况下增加不必要的复杂性。

我们建议开发者从直接使用 LLM APIs 开始：许多模式只需几行代码即可实现。如果你确实使用了框架，请确保你理解其底层代码。对底层机制的错误假设是客户出错的一个常见来源。

搭建 blocks，workflows，以及agents

在本节中，我们将探讨我们在生产环境中看到的 agentic systems 的常见模式。我们将从我们的基础构建 blocks ——增强型 LLM——开始，并逐步增加复杂性，从简单的组合式 workflows 到自主的 agents。

搭建 blocks：The augmented LLM

Agentic systems 的基础构建块是一个通过增强功能（如 retrieval、tools 和 memory）得到增强的 LLM。我们目前的模型能够主动使用这些能力——生成自己的搜索查询、选择合适的工具，并决定要保留哪些信息。

我们建议在实现时关注两个关键方面：根据你的具体用例定制这些能力，并确保它们为你的 LLM 提供一个简单、文档齐全的接口。虽然实现这些增强功能的方法有很多，但其中一种方法是通过我们最近发布的 Model Context Protocol，它允许开发者通过一个简单的客户端实现，与一个不断增长的第三方工具生态系统集成。

在本文的其余部分，我们将假设每次 LLM 调用都能访问这些增强功能。

Workflow: Prompt chaining

提示链 (Prompt chaining) 将一个任务分解为一系列步骤，其中每个 LLM 调用处理前一个调用的输出。你可以在任何中间步骤上添加程序化的检查（见下图中的“gate”），以确保流程仍在正轨上。

何时使用此 workflow：此 workflow 非常适用于任务可以轻松、清晰地分解为固定子任务的情况。其主要目标是通过使每个 LLM 调用成为一个更简单的任务，来用延迟换取更高的准确性。

适用于提示链 workflow 场景的例子：

• 生成营销文案，然后将其翻译成另一种语言。
• 编写文档大纲，检查大纲是否符合某些标准，然后根据大纲编写文档。

Workflow：Routing

路由 (Routing) 对输入进行分类，并将其导向一个专门的后续任务。这个 workflow 允许关注点分离，并构建更专门化的 prompts。如果没有这个 workflow，针对一种输入的优化可能会损害在其他输入上的性能。

何时使用此 workflow：当存在可以更好地分开处理的不同类别，并且分类可以被 LLM 或更传统的分类模型/算法准确处理时，路由 (Routing) 在复杂任务中表现良好。

适用于路由 workflow 场景的例子：

• 将不同类型的客户服务查询（一般问题、退款请求、技术支持）引导到不同的下游流程、prompts 和 tools 中。
• 将简单/常见的问题路由到更小、更具成本效益的模型（如 Claude Haiku），并将困难/不寻常的问题路由到功能更强大的模型（如 Claude Sonnet），以优化最佳性能。

Workflow: Parallelization

LLM 有时可以同时处理一个任务，并以编程方式聚合它们的输出。这种 workflow，即并行化 (parallelization)，主要有两种变体：

• 分片 (Sectioning)：将一个任务分解为并行运行的独立子任务。
• 投票 (Voting)：多次运行同一个任务以获得多样化的输出。

何时使用此 workflow：当划分的子任务可以为了速度而并行化时，或者当需要多个视角或尝试以获得更高置信度的结果时，并行化 (Parallelization) 是有效的。对于具有多种考虑因素的复杂任务，LLM 通常在每个考虑因素由单独的 LLM 调用处理时表现更好，从而可以集中关注每个特定方面。

适用于并行化 workflow 场景的例子：

分片 (Sectioning)：

• 实现 guardrails，其中一个模型实例处理用户查询，而另一个实例筛选不当内容或请求。这通常比让同一个 LLM 调用同时处理 guardrails 和核心响应表现得更好。
• 自动化评估 (evals) 以评估 LLM 性能，其中每个 LLM 调用评估模型在给定 prompt 上性能的不同方面。

投票 (Voting)：

• 审查一段代码的漏洞，其中几个不同的 prompts 会审查代码，如果发现问题就进行标记。
• 评估一段给定的内容是否不当，通过多个 prompts 评估不同方面或要求不同的投票阈值来平衡假阳性和假阴性。

Workflow: Orchestrator-workers

在 orchestrator-workers (协调器-工作者) workflow 中，一个中心的 LLM 会动态地分解任务，将它们委托给 worker LLMs，并综合它们的结果。

何时使用此 workflow：此 workflow 非常适用于你无法预测所需子任务的复杂任务（例如，在编码中，需要更改的文件数量以及每个文件中更改的性质可能取决于任务）。虽然它在结构上与并行化相似，但关键区别在于其灵活性——子任务不是预先定义的，而是由 orchestrator (协调器) 根据具体输入动态决定的。

适用于 Orchestrator-workers workflow 场景的例子：

• 每次都对多个文件进行复杂更改的编码产品。
• 涉及从多个来源收集和分析信息以获取可能相关信息的搜索任务。

Workflow: Evaluator-optimizer

在 evaluator-optimizer (评估器-优化器) workflow 中，一个 LLM 调用生成响应，而另一个 LLM 则在循环中提供评估和反馈。

何时使用此 workflow：当我们有明确的评估标准，并且迭代优化能提供可衡量的价值时，此 workflow 特别有效。两个合适的迹象是：首先，当人类明确表达他们的反馈时，LLM 的响应可以得到明显改善；其次，LLM 能够提供此类反馈。这类似于人类作者在撰写一篇精炼文档时可能经历的迭代写作过程。

适用于 Evaluator-optimizer workflow 场景的例子：

• 文学翻译，其中存在翻译器 LLM 最初可能无法捕捉到的细微差别，但评估器 LLM 可以提供有用的评论。
• 复杂的搜索任务，需要多轮搜索和分析以收集全面的信息，其中评估器决定是否需要进一步搜索。

Agents

随着 LLMs 在关键能力——理解复杂输入、进行推理和规划、可靠地使用 tools、以及从错误中恢复——方面日趋成熟，Agents 正在生产环境中崭露头角。Agents 的工作始于人类用户的命令或互动式讨论。一旦任务明确，agents 就会独立规划和操作，并可能返回给人类以获取更多信息或判断。在执行过程中，至关重要的是 agents 在每一步都从环境中获取“ground truth”（例如 tool 调用结果或代码执行情况）来评估其进展。然后，Agents 可以在 checkpoints 或遇到 blockers 时暂停以获取人类反馈。任务通常在完成后终止，但包含停止条件（例如最大迭代次数）以保持控制也很常见。

Agents 可以处理复杂的任务，但它们的实现通常很简单。它们通常只是在循环中基于环境反馈使用 tools 的 LLMs。因此，清晰周到地设计 toolsets 及其文档至关重要。我们将在附录2 (“Prompt Engineering your Tools”) 中详细阐述工具开发的最佳实践。

何时使用 agents：Agents 可用于开放式问题，这些问题难以或不可能预测所需的步骤数，并且你无法硬编码固定的路径。LLM 可能会运行多轮，你必须对其决策有一定程度的信任。Agents 的自主性使其成为在受信任环境中扩展任务的理想选择。

agents 的自主性意味着更高的成本和复合错误的潜力。我们建议在沙盒环境中进行广泛测试，并配备适当的 guardrails。

适用于 agents 场景的例子：

以下示例来自我们自己的实现：

• 一个用于解决 SWE-bench 任务的编码 Agent，其中涉及根据任务描述对多个文件进行编辑；
• 我们的“计算机使用”参考实现，其中 Claude 使用计算机来完成任务。

Combining and customizing these patterns

这些构建块并非规定性的。它们是开发者可以根据不同用例塑造和组合的常见模式。成功的关键，与任何 LLM 功能一样，在于衡量性能和迭代实现。重申一遍：只有在能够证明增加复杂性可以显著改善结果时，你才应该考虑这样做。

Summary

在 LLM 领域取得成功，关键不在于构建最复杂的系统，而在于构建满足你需求的正确系统。从简单的 prompts 开始，通过全面的评估来优化它们，只有在更简单的解决方案无法满足需求时，才添加多步骤的 agentic systems。

在实现 agents 时，我们尝试遵循三个核心原则：

• 保持 agent 设计的简洁性。
• 通过明确展示 agent 的规划步骤来优先考虑透明度。
• 通过详尽的 tool 文档和测试，精心打造你的 agent-computer interface (ACI)。

Frameworks 可以帮助你快速入门，但在转向生产环境时，不要犹豫，减少抽象层，用基础组件来构建。通过遵循这些原则，你可以创建出不仅功能强大，而且可靠、可维护、并受用户信赖的 agents。

原文：Building effective agents by Anthropic

Read More

自从出行游玩要带上小朋友之后，我发现我和我太太能选择的旅行目的地范围越来越小了，总得围绕着路程近、交通方便、环境好、食物适应这几个论点展开，但在这之上，有一个点更是王炸——就是有海滩。对于学龄前的小朋友来说，海滩真是一个绝佳的地方，他们可以在沙滩上瞎玩一整天，而丰盛的海鲜又很满足我们这种生长在沿海地区人的口味。甚至是像我儿子那样的婴儿，都可以在海滩晒晒太阳吹吹海风。我和我太太上次来泰国应该已经是新冠疫情肆虐之前的2019年底了，甚至对我太太来说，这已经是她第四次来泰国了。但是这个国家似乎有着很奇妙的魅力，让人每次来都有新的体验。

清迈仿佛一封没拆封的信笺

我们原本的行程只是想带着小朋友找个有海有沙滩的地方趟几天。因为带着两个学龄前的小朋友，基本上就可以告别人文历史氛围比较重的景点了。我们那时候备选的有冲绳、富国岛、长滩，最后还是选择了比较保险的泰国普吉岛。但是杭州直飞普吉的航班又太贵，于是我提议要么去清迈转机，顺便逛一逛这座历史悠久群山环绕，有着“泰北玫瑰”之称的城市。

清迈有着700多年的建城历史，曾是兰纳王国的首都。这座城市被群山环抱，古城呈正方形布局，至今仍保留着部分古城墙和护城河。与曼谷的喧嚣不同，清迈的节奏更为悠闲，城内300多座寺庙静静诉说着历史，而遍布街头的咖啡馆和手工艺品店又赋予了它现代文艺的气息。当地人常说，清迈就像一封没拆封的信笺，只有亲自来过，才能读懂它内在的美好。

落地清迈的第一个晚上，我们就很幸运赶上了著名的周日夜市。这确实是我第一次感受泰国的夜市文化。周日夜市正在古城的中心，我们的旅程正好是从著名的塔佩门开始。

塔佩门作为每个清迈旅游攻略的首页，其实是个不大的古城门，因为两边砖红色的城墙保存完好十分好出片所以深受游客喜爱。而我们这次的代步工具却是十分有泰国特色，因为老老小小6个人又有一辆婴儿车，原本两辆出租车都坐不下的情况，却被泰国红色的观光三蹦子给完美解决了。

清迈古城不大，但据说有300多座寺庙，密度之高令人叹为观止。走在古城的任何一条街道上，转角就能遇见一座寺庙，金顶白墙在热带阳光下闪闪发光。这些寺庙大多有数百年历史，有的建于兰纳王朝时期，有的经过历代重修，每一座都有自己的故事。对于行程有限的游客来说，选择参观哪些寺庙反而成了甜蜜的烦恼。

盼道寺虽然很有名，但是因为占地较小，而且坐落于著名的契迪龙寺边上，所以很容易被游客遗漏。这座建于14世纪的寺庙是清迈最古老的寺庙之一，全柚木的建筑在阳光的照射下泛着金黄色的光泽，与常见的石质寺庙截然不同。寺内的 teak wood 建筑工艺精湛，每一处雕刻都诉说着兰纳王朝的辉煌历史。最让人印象深刻的是寺内的僧侣学校，常常能看到年轻僧侣们朗朗读书的场景，为这座古老寺庙增添了浓厚的学术氛围。

寺中央那座巨大的金色佛塔是盼道寺的标志性建筑，这座兰纳风格的钟形佛塔在阳光下熠熠生辉，塔身呈优雅的钟形曲线，与泰国中部常见的方形佛塔形成鲜明对比。佛塔的基座装饰着精美的灰泥浮雕，描绘着佛教故事和兰纳传统图案。塔尖部分层层叠叠，象征着通往佛国净土的阶梯。据说这座佛塔供奉着珍贵的佛骨舍利，是兰纳王朝时期佛教盛行的见证。站在佛塔下仰望，金色的塔身在蓝天白云的映衬下格外庄严神圣，让人不由自主地生出敬畏之心。这种钟形佛塔设计体现了兰纳建筑对美学的独特理解，圆润的线条中蕴含着佛教圆融无碍的智慧。

契迪龙寺（Wat Chedi Luang）是清迈古城内最壮观的寺庙之一，也是这座城市的地标性建筑。”契迪龙”意为”大佛塔”，寺内的巨型兰纳风格佛塔高达80米，始建于1441年，曾是清迈最高的建筑。这座砖红色的佛塔虽然因地震而部分损毁，但依然展现出兰纳王朝鼎盛时期的建筑辉煌。佛塔的基座雕刻着精美的象头雕塑，每一层都有不同风格的佛教装饰，体现了兰纳工匠的精湛技艺。与盼道寺的钟形佛塔不同，契迪龙寺的佛塔呈八角形，层层递减的设计让人联想到通往天国的阶梯。站在佛塔脚下仰望，那种历史的厚重感扑面而来，让人不禁想象当年兰纳王朝的繁荣景象。

历史上，契迪龙寺曾短暂供奉过著名的翡翠佛像，后来辗转迁往老挝，最终安置在曼谷的玉佛寺。1545年的强烈地震使佛塔顶部坍塌，此后经历多次修缮，今天我们看到的保留了“残塔”的独特风貌，更能让人直观感受时间在砖石上的刻痕。寺内还供奉着清迈的城柱（Inthakhin），被视为守护城市的象征，每年都会举行献花祈福的传统节庆，寺庙与城市的命运也因此紧紧相连。

帕辛寺（Wat Phra Singh）是清迈古城最受尊敬的寺庙之一，始建于14世纪兰纳王朝时期。寺内最有名的赖坎殿（Viharn Lai Kham）以精美的柚木结构和金箔装饰闻名，殿中供奉的帕辛佛（Phra Singh）是镇寺之宝，也是宋干节期间会被抬出游行的圣像。寺内壁画细致描绘了兰纳民俗与佛教故事，色彩温润、人物生动。与宏伟的契迪龙寺不同，帕辛寺更显精致典雅，庭院层次分明、雕饰繁复，非常适合慢慢参观，既能近距离感受兰纳工艺，又不至于过度奔波。

在离开帕辛寺的时候，我们很幸运地看到了异国他乡的满月。在杂乱的清迈街头，红色双条车在巷口穿梭，嘟嘟车的引擎声与寺庙的钟声交织。电线像藤蔓一样攀附在杆上，偶尔有纸灯笼在夜风里轻轻摇晃。路边摊冒着热气，烤肉与香料的味道混合着熏香，水果摊上的芒果和榴莲泛着亮光。老屋的柚木屋檐在灯影下显出温润的质感，墙面上的手绘招牌略显褪色却格外有味道。护城河边水光粼粼，满月在波纹里碎成细小的银片；行人步伐不急不缓，流浪猫从巷子口一闪而过。市井与宗教的气息在这里并存，喧闹里有一种慢下来的力量。

素贴山（Doi Suthep）像一道青色的屏风立在清迈城西，山路曲折，红色双条车沿着森林和云雾一路盘旋而上。山巅的双龙寺（Wat Phra That Doi Suthep）因寺前306级台阶两侧的那迦龙身而得名，鳞片在阳光下泛出绿金的光泽。寺内的金色佛塔在高海拔的清风里静静闪耀，钟声与风声交织，远处城景像棋盘一样铺在脚下。根据传说，一头白象驮着佛舍利来到此处，绕行三圈后长鸣而逝，于是便在此建寺，始建年代可追溯到14世纪。天气晴朗时，寺后观景平台能看到古城的方形轮廓与机场跑道的灯带，日落时分更是层层云霞把城与山染成柔和的金色。素贴山是清迈人的精神所在，双龙寺则是这座城的天际信标。虽然从下车点到双龙寺要爬好高的台阶，但确实物有所值。我很喜欢清迈山间朦胧幽静的热带丛林。山林里常有薄雾贴着树冠，潮湿的风带着泥土与叶子的清气。榕树根须缠绕石阶，青苔沿着台阶蔓延，湿润的石面在光影里时暗时亮。藤蔓从树间垂落，偶尔有兰花和野花点缀其间，蝉声与鸟鸣在林间回荡，仿佛织成密密的背景。雨季的痕迹留在树皮与落叶上，抬头能看到层层叠叠的绿意把天光过滤成柔软的颜色。

素贴山脚下就是清迈大学。清迈大学（Chiang Mai University，CMU）建立于1964年，是泰国第一所设立在曼谷之外的公立大学，背靠素贴山森林，校园与城市的边界在山风与湖光之间自然延展。校园学科齐全，涵盖医学、工程、农业、艺术与人文等领域，其中农业与医学长期位居泰北前列，也以兰纳文化研究见长。环校公路串联起教学区与生活区，校车穿梭其间，节奏不急不缓。

自然风景方面，校园内的安卡欧湖（Ang Kaew）是最具代表性的景致，傍晚时分水面如镜，远处是素贴山的树海与落日的金边。榕树与雨林植被在校内随处可见，季风里云影低垂，山色常常把校园染成柔和的绿色。这里既是高等学府，也是清迈最美的公共空间之一，学习与散步可以在同一条林荫道上完成。清迈大学一直是泰国影视剧喜欢的取景地，我很喜欢坐在安卡欧湖边的草坪上，层层叠叠的小坡一直延伸到天边，好像真的就是泰剧里青春的样子。

罗摩利寺（Wat Lok Moli）位于清迈古城北侧的长帕门附近，是一座典型的兰纳风格寺庙。寺内木构大殿沉稳古雅，后方的砖砌佛塔层次分明、线条古拙，保留了清迈早期寺塔的美学。院内常见白线祈福与围塔礼佛，氛围清静而不张扬。

这里有一个颇具特色的祈福习俗：信众会把清水或献供放入系在滑轮与缆绳上的桶中，由僧人或寺工作人员缓缓将其送至塔身高处，再在颂念中洒净。把水“送上去”，象征将心愿与功德递送至更高之处，也寓意净化与祝福随风扩散。我女儿在边上傻乎乎看着我玩这个缆绳游戏，然后一本正经地在边上许愿的样子着实让我想笑。

一场和大象的约会

在清迈，人与大象的关系有着悠久的历史。兰纳时期，大象既是战象也是山林伐木的伙伴；1989年泰国禁伐后，许多大象从伐木营转入救助与保护，清迈周边由此出现了以救援与散养为宗旨的象园与保护区，更多倡导“不骑、不表演”的伦理体验。白象在泰文化中被视为祥瑞，城内寺庙的壁画与雕饰里也常见其身影。如今在清迈与大象相遇，更多意味着学习与守护——在河谷喂食、泥塘洗浴、远距离观察，听护象人讲述它们的名字、来历与性情。

关于清迈北部的象营，坊间流传不少“八卦”：许多营地源自早期同一家族体系，后来因经营理念、路线选择与分工不同而各自独立。有的坚持“不骑不表演”的生态体验，有的主打近距离互动；有的侧重救助与养老，有的强调科普与研究。也常能听到“从老象王家族分出来”的支线与更迭的传闻，但真伪难辨。与其追逐八卦，不如用可观察的指标判断一家象营的素质：饲养密度是否合理、象群是否有自由活动时间与遮阴水源、是否链拴、是否使用铁钩、是否允许游客骑乘、是否公开医疗与救助记录、护象人对个体的熟悉程度与讲述是否真实。把注意力放在福利与伦理上，才是参与者应有的态度。

我一开始是很期待能骑大象的，但因为大象的乘骑表演对其健康伤害很大，因此清迈正规的象营都不提供这个项目了。我这次选择的象营虽然不大，但是对小朋友非常友好，可以投喂，可以给大象做饭团，可以陪大象散步和洗澡。

大象是非常高智商的哺乳动物，但是雄性大象自我意识比较强而且比较敏感，因此象营提供给游客互动的大象只有成年的雌性大象和幼年大象。我很惊讶于大象灵活的鼻子所能做出的精细动作，称之为”大象的手“都不为过。而大象之间也明显能从生活习惯上看出代沟。老一辈的大象从水管里喝水是先吸入鼻子，再通过鼻子注入口腔，而新一辈的幼年大象则是会直接拿水管往自己口腔里注水，十分有趣。而当它们吃饱喝足了也像人类小朋友一样淘气，把水喷的满天都是。

在象营的半天里，时间仿佛变慢了。护象人用极轻的语调和它们交流，手势与口令更像是一种长期的默契。我们保持适度距离，在河谷递上甘蔗和香蕉，看它们从容地取食、下水、扬起泥浆。近距离看见那双深色的眼睛与满身的皱褶，会意识到它们承载的不只是体重，还有山林与岁月。

清迈让我体会到另一种“旅游”的意义：不骑、不表演、不强求互动，把尊重摆在第一位。人与自然的和谐大概就是这样的节制与克制——降低音量，放慢脚步，留出空间。等风起时，象群、山风与河谷像在用同一口呼吸。那一刻，保护不再只是标语，而是每一次轻声、每一次不去打扰的选择。

普吉的风帆和浪花

紧接着我们来到旅行的第二站也是重头戏普吉岛。普吉岛位于安达曼海，是泰国最大的海岛，也是度假与海上运动的集散地。沙滩各有性格：芭东热闹、卡塔适合冲浪、卡伦开阔、卡马拉宁静；外海跳岛则常见皮皮岛、皇帝岛、珊瑚岛等。东北季风（11–4月）海况更稳、能见度好，西南季风（5–10月）浪涌更强、需看旗况与天气。玩法从浮潜、深潜到帆船与日落巡航，海面上的光影切换像翻动的银色书页。岛上老镇保存葡式—中式折衷的街屋与彩色立面，夜市的烟火味与安达曼的海风在同一条街上相逢。普吉并不只是一片海，更是一种节奏：松弛、明亮，但也需要对自然保持敬畏——红旗禁泳、珊瑚与海龟的栖息地需要被好好保护。

因为普吉岛本身交通不便，可以旅行观光的岛屿之间又相距甚远，所以在当地参与散客 tour 是一个很棒的选择。我在国外旅游最大的一个乐趣就在这种散客 tour 里，因为总能遇到来自世界各地的不同人，他们的语言、文化、生活方式都不同，这使得我在旅行里有了更多的机会去了解来自世界各地的人文故事。

但普吉的旅行确实没什么可以介绍的，我们住在芭东海滩，逛了逛周边的夜市和沙滩。我很喜欢珊瑚岛和皇帝岛的跳岛旅行。

皇帝岛附近的浮潜体验非常不错，海水如果冻般透明，能看到6、7米深处的各种大小鱼类，甚至我看到了之前在百科故事中才见过的砗磲（Tridacna）。这种栖息在浅海珊瑚礁的巨型贝类与虫黄藻共生，外套膜色彩绚烂，安达曼海也常见其身影。甚至还有小海蛇在上面徘徊。

游艇回程的路上很幸运看到了雨后如刀剑般笔直的彩虹。

还有金色的晚霞。

如果说珊瑚岛和皇帝岛符合我对普吉的正常期待，那攀牙湾的旅程则完全是惊喜。

攀牙湾是海上的石林。墨绿的海面上，石灰岩柱从水里直刺天空，岛与岛之间像被刀锋切开，形成幽深的海洞与“hong”（泻湖）。小艇钻入海蚀洞的瞬间，光线忽暗忽明，水壁近得能看见贝壳的纹理与钟乳的滴痕；再穿过狭窄的“门”，豁然出现一处被山石环抱的内海，静得只剩桨声。

远处的攀牙湾国家公园与“詹姆斯邦德岛”（Khao Phing Kan / Ko Ta Pu）是名片，但真正打动我的是潮汐与光影。当潮水涨落，红树林露出根系，水面浮起漂来的叶片；午后云层拆散阳光，岩壁的纹理像被重新描过一遍。这里的美是慢的、立体的，需要耐心地看海的呼吸。

其中尤其惊喜的就是坐着独木舟在充满卡斯特溶洞的峡湾里探险。

詹姆斯邦德岛（Ko Ta Pu / Khao Phing Kan）因《007之金枪人》出镜而一举成名。针状的石灰岩柱从海面直立，像一枚钉子扎在安达曼的水面上；对岸的巨岩因地壳运动与海蚀裂成两片，纵横的纹理在潮起潮落间愈发清晰。登岸的游人不少，但只要换几个角度、把脚步放慢，岩柱、云影与水色会拼成不同的画面。这里适合短暂停留与远观：狭窄岸线不宜久留，留意潮汐与风浪，也别踩踏红树林与脆弱的滩涂。

甚至在攀牙湾还能找到几乎没被开发过的完美曳尾白沙滩，一点不逊色于马尔代夫。

曼谷的雨

最后一站曼谷说实话没啥好说的，因为之前来过。但曼谷的堵车居然在近些年也变成了一个景点。

真正让我记住的是雨。曼谷的雨来得很快，云墙在天边合拢，城市立刻像被调低了色温。高架轻轨下滴水成线，嘟嘟车裹着透明雨披穿行，黄绿相间的的士在积水里拉出长长的红色尾灯。堵车在雨里忽然变得有点好看，像一场慢速的灯光秀。

这次只在卧佛寺和大皇宫门口徘徊了一会儿，没有进去参观。

沿着湄南河走，风里带着水汽与河味。郑王庙的塔尖在雨幕里若隐若现，卧佛寺屋檐的滴水像敲着极轻的拍子。雨停得也快，云层抽丝，街道升腾起淡淡的蒸气，城市在水光里显得柔软。曼谷不是只有快与热，它也有一场雨带来的慢与静。郑王庙也如我们7年前参观的一样，留下了许多我和我太太刚刚结婚时候的美好回忆。

Read More

随着大语言模型 (LLM) 在各个领域的广泛应用，如何高效地部署和推理这些模型成为了一个关键挑战。传统的模型推理服务往往面临着内存利用率低、吞吐量受限、延迟不可控等问题，这些瓶颈严重制约了 LLM 在生产环境中的规模化应用。vLLM 作为一个专为 LLM 优化的高性能推理服务框架，通过一系列创新的技术方案，有效解决了这些痛点问题。

本文将深入剖析 vLLM 的核心架构和关键技术实现，从底层的内存管理机制到上层的服务调度策略，全面解析其如何实现高效的 LLM 推理服务。我们将重点探讨以下几个核心技术模块：

PagedAttention 机制：借鉴操作系统中虚拟内存管理的思想，vLLM 提出了 PagedAttention 技术，将 KV Cache 按页进行管理，实现了内存的按需分配和高效利用。这种设计不仅显著降低了内存碎片化问题，还支持了动态序列长度处理，使得内存利用率相比传统方案提升了数倍。

Continuous Batching (连续批处理)：传统的静态批处理方式存在严重的计算资源浪费问题，特别是当批内序列长度差异较大时。vLLM 的连续批处理技术支持序列的动态加入和完成，实现了真正的流水线式处理，大幅提升了系统吞吐量和资源利用效率。

Prefix Caching (前缀缓存)：在实际应用中，很多请求往往共享相同的前缀内容（如系统提示词、模板等）。vLLM 通过智能的前缀缓存机制，能够复用已计算的 KV Cache，避免重复计算，显著降低了推理延迟和计算开销。

Speculative Decoding (推测解码)：为了进一步提升生成速度，vLLM 集成了推测解码技术，通过使用较小的 draft 模型预先生成候选 token，然后由主模型进行验证，实现了在保证输出质量的前提下大幅加速文本生成过程。

分布式架构与多 GPU 协同：面对大模型参数量不断增长的趋势，vLLM 提供了完善的分布式解决方案，支持张量并行、流水线并行等多种并行策略，能够在多 GPU、多节点环境下实现高效的模型推理，满足大规模生产环境的性能需求。

动态扩缩容与服务化：作为一个面向生产的推理框架，vLLM 不仅关注性能优化，还提供了完整的服务化能力，包括请求路由、负载均衡、自动扩缩容等功能，使得用户能够轻松构建高可用、高性能的 LLM 服务集群。

通过对这些关键技术的深入分析，我们将展现 vLLM 如何通过系统性的优化设计，在保证推理质量的前提下，实现了相比传统方案数倍甚至数十倍的性能提升。这些技术创新不仅推动了 LLM 推理服务的发展，也为整个 AI 基础设施领域提供了宝贵的设计思路和实践经验。一共分为五个部分：

• LLM engine 以及 engine core：包含 vLLM 的基础架构（调度、PagedAttention、continuous batching）
• Advanced Features（高级特性）：chunked prefill（分块预填充）、prefix caching（前缀缓存）、guided & speculative decoding（引导解码与推测解码）、disaggregated P/D（Prefill-Decoding 分离）
• Scaling Up：单进程到多进程、多 GPU
• Server Layer：分布式集群服务化部署
• Benchmarks 与 Auto-tuning：平衡延迟与吞吐

LLM Engine & Engine Core

在 vLLM 中，LLM Engine 是最基础的 block，在离线场景中，它本身就支持高吞吐量推理。以下是一个简单的离线推理例子：

from vllm import LLM, SamplingParams

prompts = [
    "Hello, my name is",
    "The president of the United States is",
]

sampling_params = SamplingParams(temperature=0.8, top_p=0.95)

def main():
    llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0")

    outputs = llm.generate(prompts, sampling_params)

if __name__ == "__main__":
    main()

## Environment vars:
##   VLLM_USE_V1="1" ## we're using engine V1
##   VLLM_ENABLE_V1_MULTIPROCESSING="0" ## we're running in a single process

我们调用模型执行器的 execute_model，它会委派给 Worker，而 Worker 又会继续委派给 model runner。

主要步骤如下：

• 更新状态 —— 从 input_batch 中裁剪已完成的请求；更新与前向传播相关的其他 metadata（例如每个请求的 KV cache 块数，用于在分页的 KV cache 内存中建立索引）。
• 准备输入 —— 将缓冲区从 CPU→GPU 复制；计算位置；构建 slot_mapping（示例中会详细说明）；构造注意力 metadata。
• 前向传播 —— 使用自定义的 PagedAttention 内核运行模型。所有序列会被展平并连接为一个长的“超级序列”。位置索引与注意力掩码确保每个序列只关注自己的 token，从而在不使用右侧填充的情况下实现持续批处理。
• 收集最后一个 token 的状态 —— 为每个序列的最终位置提取隐藏状态并计算 logits。
• 采样 —— 按照采样配置（贪心、温度、top-p、top-k 等）从计算出的 logits 中采样 token。

前向步骤本身有两种执行模式：

• Eager 模式（Eager Mode） —— 在启用 eager 执行时运行标准的 PyTorch 前向传播。
• 捕获模式（Capture Mode） —— 在未强制启用 eager 的情况下，执行或回放预先捕获的 CUDA graph（还记得在引擎构建的初始化 KV cache 过程中我们已经捕获了它们）。

这些配置有：

• 离线模式（无 Web 服务或分布式系统架构）；
• 同步执行（所有执行都在单个阻塞进程中进行）；
• 单 GPU（无数据/模型/流水线/专家并行；DP/TP/PP/EP = 1）；
• 使用标准 Transformer 结构（支持像 Jamba 这样的混合模型需要更复杂的混合 KV 缓存内存分配器）。

在这个例子中，我们做了两件事： 1. 实例化了一个 engine； 2. 通过给定的 prompt 来调用 generate 方法进行采样。

LLM Engine constructor

对于engine而言，核心的组成部分有：

• vLLM config：包含模型配置的全部信息、cache、并行策略等；
• processor：通过 validation、tokenization 和 processing 将 raw input -> EngineCoreRequests;
• engine core client：在我们的例子中使用了 InprocClient ，基本上等于 EngineCore ，会逐步搭建成 DPLBAsyncMPClient ，允许大规模提供服务；
• output processor：将 raw EngineCoreOutputs -> RequestOutputs 转换给用户看。

至于 EngineCore 本身由以下组件组成：

• 模型执行器 (Model Executor): 驱动模型的前向传播。我们目前接触的是在单个GPU上使用单个Worker进程的 UniProcExecutor，后续会逐步扩展到支持多GPU的 MultiProcExecutor。
• 结构化输出管理器 (Structured Output Manager): 用于引导式解码（稍后会详细介绍）。
• 调度器 (Scheduler): 决定哪些请求进入下一个引擎步骤，它进一步包含：
  • 策略设置 (policy setting): 可以是FCFS（先到先得）或优先级（高优先级请求优先处理）。
  • 等待和运行队列 (waiting and running queues)。
  • KV缓存管理器 (KV cache manager): PagedAttention机制的核心。

KV Cache Manager 维护了 free_block_queue，也就是可用的 KV Cache blocks组成的资源池；规模往往能到几十万，取决于显存与块大小。当 PagedAttention 执行时，这些块承担索引作用，将各个 token 对应到它们的 KV Cache block。

其中，对于一个标准 Transformer 层（非 MLA）的 block size 可通过以下方式计算：

2 (key/value) * block_size (default=16) * num_kv_heads * head_size * dtype_num_bytes (e.g. 2 for bf16)

当 model executor 构建时，会创建一个 Worker 对象，并执行三个主要步骤（在使用 MultiProcExecutor 时，这些步骤会在不同 GPU 上的每个 worker 进程中独立运行）：

• 初始化设备:
  • 为该 worker 分配 CUDA 设备（e.g. cuda:0 ），并检查模型的 dtype 是否受支持（e.g. bf16 ）
  • 根据设定的 gpu_memory_utilization （e.g. 0.8 → 80% of total VRAM）验证显存是否充足
  • 配置分布式设置（ DP / TP / PP / EP, etc.）
  • 实例化 model_runner （包含采样器、KV cache，以及forward pass的buffers如 input_ids 、 positions, etc.）
  • 实例化 InputBatch 对象（包含 CPU-side forward pass buffering、KV cache indexing、sampling metadata等）
• 加载模型:
  • 实例化模型架构
  • 加载模型权重
  • 调用 model.eval() （PyTorch 的推理模式）
  • 可选：对模型调用 torch.compile()
• 初始化 KV Cache:
  • 获取按层的 KV cache spec。通常为 FullAttentionSpec （同质 Transformer），但在引入混合模型（滑动窗口、Transformer/SSM，如 Jamba）后变得更复杂
  • 运行一次dummy/profiling forward pass，并记录 GPU 内存快照，用于计算在可用显存中能容纳多少 KV cache blocks
  • 为注意力层分配、reshape并绑定 KV cache tensors
  • 准备 attention metadata（如将后端设置为 FlashAttention ），供后续前向过程中的内核使用
  • 若未提供 --enforce-eager，则针对若干预热批大小进行空跑并捕获 CUDA graph。CUDA graph 会把整段 GPU 工作记录为一个 DAG；之后在前向过程中，我们会启动/回放这些预先捕获（预烘焙）的 CUDA graph，削减 kernel 启动开销，因而时延更低。

我们在这里抽象掉了许多底层细节，但以上是后文将反复引用的核心组件与流程。引擎初始化完成后，继续进入 generate 函数。

Generate function

第一步是对请求进行校验并送入 engine 。对于每个 prompt，我们会：

1. 创建一个唯一的请求 ID，并记录其到达时间。
2. 调用输入预处理器对 prompt 进行标记化（tokenize），返回一个字典 dictionary，包含 prompt、prompt_token_ids，以及一个 type（如 text、tokens、embeds, etc.）。
3. 将这些信息打包成一个 EngineCoreRequest，并添加优先级、采样参数及其他 metadata。
4. 将请求传入 engine core，core 会将其包装为一个 Request 对象并将状态设为 WAITING；随后把该请求加入调度器的等待队列（若为先来先服务 FCFS 则使用 append；若为优先级调度则使用 heap-push）。

至此，引擎已经“进料”，执行即可开始。在同步引擎示例中，只会处理这些初始 prompt——运行过程中无法插入新请求。相反，异步引擎支持在运行中注入请求（即“持续批处理” continuous batching）：在每一步之后，同时考虑新请求与已有请求。

前向传播将 batch 扁平化为单序列，配合高效的定制 kernel 处理路径，使得即使在同步引擎中也天然具备 continuous batching 能力。

接下来，只要仍有请求待处理，引擎就会反复调用 step() 函数。每一步包含三个阶段：

• 调度（Schedule）：选择本步要运行的请求（decode，and/or (chunked) prefill）。
• 前向传播（Forward pass）：运行模型并进行 token 采样。
• 后处理（Postprocess）：将采样得到的 token ID 追加到各个 Request，执行反标记化（detokenize），并检查停止条件。若某个请求已完成，则进行清理（例如把它的 KV Cache block 归还到 free_block_queue），并提前返回该请求的输出。

📝 停止条件包括：

• 请求超过长度上限（max_model_length 或其自身的 max_tokens）。
• 采样到 EOS ID（除非启用了 ignore_eos → 在 benchmarking 中可用于强制生成固定数量的输出 token）。
• 采样到的 token 匹配到采样参数中指定的任意 stop_token_ids。
• 输出中出现停止字符串（stop strings）——我们会将输出截断到首次出现停止字符串的位置，并在引擎中终止该请求（注意：stop_token_ids 会保留在输出中，而停止字符串不会保留）。

在流式模式中，我们会在生成过程中实时发送中间 token，但这里暂不展开。接下来，我们将更详细地讨论调度。

Scheduler

推理引擎处理两种主要类型的工作负载：

• Prefill 请求 ： 对所有 prompt token 进行前向传播。这些通常是计算密集型的（阈值取决于硬件和prompt长度）。最后，我们从最终 token 位置的概率分布中采样一个 token。
• Decode 请求 ： 仅对最新的 token 进行前向传播。所有较早的 KV 向量已经被缓存。这些是 memory-bandwidth-bound 的，因为我们仍然需要加载所有 LLM 权重（和 KV cache）来计算一个 token。

V1 scheduler 可以在同一步骤中混合处理两种类型的请求，这得益于更智能的设计选择。相比之下，V0 engine 一次只能处理 prefill 或 decode 中的一种 workload。

Scheduler 优先处理 decode 请求——即那些已经在运行队列中的请求。对于每个这样的请求，它会：

1. 计算要生成的新 token 数量（由于推测解码和异步调度，不总是会在第一步做这些事情，——稍后会详细介绍）。
2. 调用 KV cache manager 的 allocate_slots 函数（详细信息见下文）。
3. 更新 token budget：不断减少第 1 步计算得到的 token 数量。

之后，它处理等待队列中的 prefill 请求：

1. 检索已计算块的数量（如果禁用前缀缓存则返回 0——稍后会介绍）。
2. 调用 KV cache manager 的 allocate_slots 函数。
3. 将请求从等待队列中弹出并移动到运行队列，将其状态设置为 RUNNING。
4. 更新 token budget。

现在让我们看看 allocate_slots 的作用：

1. 计算块数量 — 确定必须分配多少个新的 KV cache 块（n）。每个块默认存储 16 个 token。例如，如果一个 prefill 请求有 17 个新 token，我们需要 ceil(17/16) = 2 个块。
2. 检查可用性 — 如果管理器池中没有足够的块，则提前退出。根据是 decode 还是 prefill 请求，引擎可能会尝试重计算抢占（V0 中支持交换抢占），通过驱逐低优先级请求（调用 kv_cache_manager.free 将 KV 块返回到块池），或者可能跳过调度并继续执行。
3. 分配块 — 通过 KV cache manager 的协调器，从块池（前面提到的 free_block_queue 双向链表）中获取前 n 个块。存储到 req_to_blocks，这是将每个 request_id 映射到其 KV cache block list的字典。

最终，我们准备好做一次前向传递了。

Run Forward pass

我们调用模型执行器的 execute_model，它会委派给 Worker，而 Worker 又进一步委派给 model runner。

主要步骤如下：

• 更新状态 —— 从 input_batch 中裁剪已完成的请求；更新与前向传播相关的其他 metadata（例如每个请求的 KV cache 块数，用于在分页的 KV cache 内存中建立索引）。
• 准备输入 —— 将缓冲区从 CPU→GPU 复制；计算位置；构建 slot_mapping（示例中会详细说明）；构造注意力 metadata。
• 前向传播 —— 使用自定义的 PagedAttention 内核运行模型。所有序列会被展平并拼接为一个长的“超级序列”。位置索引与注意力掩码确保每个序列只关注自身的 token，从而在不进行右侧填充的情况下实现 continuous batching。
• 收集最后一个 token 的状态 —— 为每个序列的最终位置提取隐藏状态并计算 logits。
• 采样 —— 按照采样配置（greedy、temperature、top-p、top-k 等）从计算得到的 logits 中采样 token。

前向步骤本身有两种执行模式：

• Eager 模式（Eager Mode） —— 启用 eager 执行时运行标准的 PyTorch 前向传播。
• 捕获模式（Capture Mode） —— 在未强制启用 eager 的情况下，执行/回放预先捕获的 CUDA graph（还记得我们在引擎构建的初始化 KV cache 过程中已捕获这些 graph）。

下面是一个具体示例，可帮助你更清晰地理解 continuous batching 和 PagedAttention：

Advanced Features — extending the core engine logic

在掌握基本的引擎流程后，我们可以继续了解一些高级特性。

我们已经讨论了抢占（preemption）、PagedAttention 和 continuous batching。

接下来，我们将深入讲解：

• Chunked prefill
• Prefix caching
• Guided decoding
• Speculative decoding
• Disaggregated P/D

Chunked prefill

Chunked prefill（分块式 prefill）是一种通过将长 prompt 的 prefill 步骤拆分为更小的 chunk 来处理长 prompt 的技术。若不使用该方法，一个非常长的请求可能会在某次 engine step 中长时间独占执行，阻止其他 prefill 请求运行，从而推迟所有其他请求并显著提高它们的延迟。

例如，令每个 chunk 包含 n (=8) 个 token，并用小写字母以 “-” 分隔来标记。一个长提示 P 可以表示为 x-y-z，其中 z 是未完成的 chunk（例如仅包含 2 个 tokens）。执行 P 的完整 prefill 至少需要 ≥ 3 个 engine step（如果某一步未被调度执行，还可能需要更多），并且只有在最后一个分块 prefill 步骤中我们才会采样一个新 token。

以下是同一示例的可视化说明：

实现很直接：为每个 engine step 设定“新增 token 数量”的上限。当请求的数量超过 long_prefill_token_threshold 时，将其重置为该阈值。其余流程由底层的索引逻辑（前文已述）自动处理。

在 vLLM V1 中，通过将 long_prefill_token_threshold 设置为正整数即可启用 chunked prefill。（从技术上讲，即使未显式设置也可能发生：若 prompt 长度超过 token 预算，我们会先截断它，并以分块 prefill 的方式运行。）

Prefix Caching

为了解释 prefix caching 的工作原理，可以参考以下代码：

from vllm import LLM, SamplingParams

long_prefix = "<a piece of text that is encoded into more than block_size tokens>"

prompts = [
    "Hello, my name is",
    "The president of the United States is",
]

sampling_params = SamplingParams(temperature=0.8, top_p=0.95)

def main():
    llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0")

    outputs = llm.generate(long_prefix + prompts[0], sampling_params)
    outputs = llm.generate(long_prefix + prompts[1], sampling_params)

if __name__ == "__main__":
    main()

Prefix caching 用于避免对多个 prompt 共享的开头部分重复计算（因此称为 “前缀 Prefix”）。

关键在于 long_prefix：它被定义为长度超过一个 KV cache block 的前缀（默认每块 16 tokens）。为简化示例，假设 long_prefix 的长度恰好为 n × block_size（其中 n ≥ 1）。

也就是说，它必须与块边界完全对齐——否则我们必须重新计算 long_prefix_len % block_size 个 tokens，因为不完整的块无法被缓存。若不使用 prefix caching，每次处理一个具有相同 long_prefix 的新请求时，都要重新计算这 n × block_size 个 tokens。

而使用 prefix caching 时，这些 tokens 只需计算一次（其 KV 存入分页的 KV cache 内存）并被复用，因此仅需处理新的 prompt tokens。这会显著加速 prefill 请求（但对 decode 无帮助）。

那么在 vLLM 中如何工作？

在首次 generate 调用的调度阶段，kv_cache_manager.get_computed_blocks 内，engine 会调用 hash_request_tokens：

• 将 long_prefix + prompts[0] 按 16-token 切分为 chunks。
• 对每个完整 chunk 计算一个 hash（使用内建 hash 或 SHA-256，后者更慢但 hash 冲突更少）。该 hash 组合了上一块的 hash、当前 tokens 以及可选 metadata。可选 metadata 包括：MM hash、LoRA ID、cache salt（注入首块的 hash，保证只有携带该 cache salt 的请求能复用这些块）。
• 每个结果以 BlockHash 对象存储，包含其 hash 与 token IDs；函数返回一个 block hashes 列表。

该列表写入 self.req_to_block_hashes[request_id]。

随后，engine 调用 find_longest_cache_hit，检查这些 hash 是否已存在于 cached_block_hash_to_block 中。对于首个请求，通常不会有命中。

然后我们调用 allocate_slots，它会进一步调用 coordinator.cache_blocks，将新的 BlockHash 条目与已分配的 KV cache blocks 关联，并把映射记录到 cached_block_hash_to_block。

随后，前向传播会在分页的 KV cache 内存中填充对应的 KV，覆盖我们上面分配的这些 KV cache blocks。

在经历多个 engine step 后，系统会继续分配更多 KV cache blocks。但在本示例中这并不重要，因为前缀在 long_prefix 之后就立即发生了差异。

第二次以相同前缀调用 generate 时，前述步骤 1–3 会再次执行，但此时 find_longest_cache_hit 会（通过线性搜索）为全部 n 个块找到命中，engine 可直接复用这些 KV cache blocks。

如果最初的请求仍在运行，这些块的引用计数（reference count）会增加（例如变为 2）。在本例中，第一个请求已经完成，因此这些块被释放回到池中，其引用计数重置为 0。由于我们能够从 cached_block_hash_to_block 取回它们，表明这些块仍然有效（KV cache manager 的逻辑就是这样设计的），所以我们只需再次将它们从 free_block_queue 中移除即可复用。

KV cache blocks 只有在即将从 free_block_queue（从左端弹出）重新分配时才会被判定为“无效”。此时如果发现该块仍有关联的 hash 并存在于 cached_block_hash_to_block 中，我们会清除该块的 hash，并将其从 cached_block_hash_to_block 中移除，以确保它不能再通过 prefix caching 被复用（至少不能用于旧的前缀）。

这就是 prefix caching 的核心：不要重复计算已经见过的前缀——直接复用它们的 KV cache！

如果理解了这个示例，也就理解了 PagedAttention 的工作方式。

Prefix caching 默认启用。若要关闭：将 enable_prefix_caching = False。

Guided Decoding (FSM)

Guided decoding（引导式解码）是一种在每个解码步对 logits 施加约束的技术，约束由基于语法的有限状态机（FSM）定义。这确保了只会采样语法允许的 token。

这是一个非常强大的设定：你可以强制执行从正则语法（Chomsky type-3，例如任意正则表达式模式）到上下文无关语法（type-2，覆盖大多数编程语言）的一切约束。

为使其更具体，我们先从最简单的示例入手，并在先前的代码基础上继续构建：

from vllm import LLM, SamplingParams
from vllm.sampling_params import GuidedDecodingParams

prompts = [
    "This sucks",
    "The weather is beautiful",
]

guided_decoding_params = GuidedDecodingParams(choice=["Positive", "Negative"])
sampling_params = SamplingParams(guided_decoding=guided_decoding_params)

def main():
    llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0")

    outputs = llm.generate(prompts, sampling_params)

if __name__ == "__main__":
    main()

在给出的 toy example 中（假设字符级 tokenization）：在 prefill 阶段，FSM 会对 logits 进行掩蔽，使得只有 “P” 或 “N” 可以被采样。若采样到 “P”，FSM 将转入 “Positive” 分支；下一步仅允许 “o”，依此类推。

vLLM 中的工作方式：

1. 在构造 LLM 引擎时，会创建 StructuredOutputManager；它可访问 tokenizer，并维护一个 _grammar_bitmask 张量；
2. 当添加请求时，状态被设置为 WAITING_FOR_FSM，grammar_init 会选择后端编译器（例如 xgrammar ，这里的大部分复杂度都隐藏在诸如 xgrammar 等第三方库中）；
3. 针对该请求的语法会异步编译；
4. 在调度过程中，如果异步编译已完成，状态切换为 WAITING，并将 request_id 加入 structured_output_request_ids；否则将其放入 skipped_waiting_requests，在下一次引擎步（engine step）重试。
5. 在调度循环结束后（仍处于调度阶段），如果存在 FSM 请求，StructuredOutputManager 会让后端准备/更新 _grammar_bitmask。
6. 当前向传播产生 logits 后，xgr_torch_compile 的函数会将位掩码展开到词表大小（由于使用 32 位整数，展开比例为 32×），并将不允许的 logits 置为 -∞。
7. 在采样下一个 token 之后，通过 accept_tokens 推进该请求的 FSM。直观上我们在 FSM 图上移动到下一个状态。

其中第 6 步值得进一步澄清：

若 vocab_size = 32，_grammar_bitmask 是一个整数；其二进制表示编码了哪些 token 被允许（“1”）与不允许（“0”）。例如，“101…001”会展开为长度为 32 的数组 [1, 0, 1, …, 0, 0, 1]；值为 0 的位置对应的 logits 被置为 -∞。

对更大的词表，会使用多个 32 位字，并按需展开/拼接。后端（例如 xgrammar）负责依据当前 FSM 状态生成这些位模式。

以下还有一个更简单的示例：vocab_size = 8 且使用 8 位整数（适合结合可视化来理解）。

在 vLLM 中，你可以通过传入所需的 guided_decoding 配置来启用该功能。

Speculative Decoding

在自回归生成（autoregressive generation）中，每产生一个新 token 都需要对 LLM 做一次前向传播（forward pass）。这个操作的计算开销非常大，因为每一步都要重新加载并应用全部模型权重，只为计算一个 token！（假设 batch size == 1，更一般的情况是 B）

Speculative decoding 通过引入一个更小的 “Draft LM” 来加速。Draft LM 以更低成本提出 k 个 token 的候选。但我们并不希望最终从这个小模型的结果中直接采样，因为它只是用来猜测可能的续写。我们最终的采样结果仍由大型模型来决定哪些候选是有效的。

具体步骤如下：

1. Draft ：使用小模型在当前上下文上运行，提出 k 个 token；
2. Verify ：使用大模型在“上下文 + k 个草稿 token”上运行一次。这会为这 k 个位置外加一个额外位置产生概率（因此得到 k+1 个候选）；
3. Accept/Reject ：从左到右遍历这 k 个草稿 token：
   • 若大模型对该草稿 token 的概率 ≥ 草稿模型的概率，则接受它；
   • 否则，以 p_large(token) / p_draft(token) 的概率接受它；
   • 在第一次拒绝处停止，或者接受所有 k 个草稿 token；
   • 若所有 k 个草稿 token 都被接受，还可以从大模型“免费”采样额外的第 k+1 个 token（因为我们已经计算了该分布）。
   • 若发生了拒绝，则在该位置构造一个重新平衡的分布（p_large - p_draft，最小值钳制为 0，并归一化为 1），并从中采样最后一个 token。

Why this works：尽管我们使用小模型提出候选，但 accept/reject 规则保证了在期望意义上，序列的分布与逐 token 从大型模型采样的结果完全一致。这意味着 speculative decoding 在统计上等价于标准的自回归解码，但潜在获得更快的 decoding 速度，因为一次大型模型的前向传播即可产出至多 k+1 个 token。

vLLM V1 不支持“LLM draft model”的方法，而是实现了更快但精度较低的提议方案：n-gram、EAGLE 和 Medusa。

三者的一句话概述：

• n-gram：取最后 prompt_lookup_max 个 token；在序列中寻找此前的匹配；若找到，则提出紧随该匹配之后的 k 个 token；否则减小窗口并重试，直到 prompt_lookup_min。
• EAGLE：对 LLM 做一次 “model surgery” ，保留 embeddings 与 LM head，用轻量级 MLP 替换 transformer stack；将其微调为一个廉价draft。
• Medusa：在大型模型的顶端（embeddings before LM head）训练 auxiliary linear head，用于并行预测接下来的 k 个 token；这些 head 能比单独运行一个小 LM 更高效地提出 token 候选。

在 vLLM 中使用 ngram 作为草稿方法来调用 speculative decoding 的方式如下：

from vllm import LLM, SamplingParams

prompts = [
    "Hello, my name is",
    "The president of the United States is",
]

sampling_params = SamplingParams(temperature=0.8, top_p=0.95)

speculative_config={
    "method": "ngram",
    "prompt_lookup_max": 5,
    "prompt_lookup_min": 3,
    "num_speculative_tokens": 3,
}

def main():
    llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0", speculative_config=speculative_config)

    outputs = llm.generate(prompts, sampling_params)

if __name__ == "__main__":
    main()

vLLM 中的工作方式：

设置阶段（引擎构造期间）：

1. 初始化设备：创建一个 drafter（草稿模型，例如 NgramProposer）和一个 rejection_sampler（其部分代码用 Triton 编写）。
2. 加载模型：加载草稿模型权重（对 n-gram 而言为空操作）。

在 generate 函数中的后续步骤（假设收到全新请求）：

1. 用大模型运行常规的 prefill 步骤。
2. 前向传播和标准采样后，调用 propose_draft_token_ids(k) 从 draft model 采样 k 个草稿 token。
3. 将这些存储在 request.spec_token_ids 中（更新请求metadata）。
4. 在下一个 engine step 中，当请求处于运行队列时，将 len(request.spec_token_ids) 加到”新 token”计数中，以便 allocate_slots 为前向传播预留足够的 KV 块。
5. 将 spec_token_ids 复制到 input_batch.token_ids_cpu 中，形成（上下文 + 草稿）token。
6. 通过 _calc_spec_decode_metadata 计算metadata（从 input_batch.token_ids_cpu 复制 token，准备 logits 等），然后在草稿 token 上运行大模型前向传播。
7. 不使用常规的 logits 采样，而是用 rejection_sampler 从左到右接受/拒绝并产生 output_token_ids。
8. 重复步骤 2-7，直到满足停止条件。

理解这一过程的最佳方式是启动调试器并逐步执行代码库，但本节希望能让你对此有所了解。

Disaggregated P/D

我此前已经提到过进行 P/D（prefill/decode）解耦的动机。

在实际推理生产过程中，prefill 与 decode 具有截然不同的性能画像（前者更偏计算受限 compute-bound，后者更偏内存带宽受限 memory-bandwidth-bound），因此将它们的执行拆分是一个合理的设计。这能更紧致地控制延迟——包括 TFTT（time-to-first-token）与 ITL（inter-token latency），其细节在基准测试章节会进一步展开。

在实践中，我们会运行 N 个 vLLM prefill 实例与 M 个 vLLM decode 实例，并根据实时请求的混合情况进行自动扩缩。Prefill worker 会将 KV 写入一个专用的 KV-cache 服务；decode worker 则从中读取。这样可以将长且突发的 prefill 与稳定、对延迟敏感的 decode 有效隔离。

那么在 vLLM 中如何实现？为便于说明，下面的示例使用 SharedStorageConnector：这是一个用于展示机制细节的调试型 connector 实现。其中Connector 是 vLLM 用于在实例之间交换 KV 的抽象。Connector 接口目前尚不稳定，短期内计划进行一些改进，这些改动可能包含不兼容的变更。

我们会启动两个 vLLM 实例（GPU 0 用于 prefill，GPU 1 用于 decode），然后在它们之间传输 KV cache：

import os
import time
from multiprocessing import Event, Process
import multiprocessing as mp

from vllm import LLM, SamplingParams
from vllm.config import KVTransferConfig

prompts = [
    "Hello, my name is",
    "The president of the United States is",
]

def run_prefill(prefill_done):
  os.environ["CUDA_VISIBLE_DEVICES"] = "0"

  sampling_params = SamplingParams(temperature=0, top_p=0.95, max_tokens=1)

  ktc=KVTransferConfig(
      kv_connector="SharedStorageConnector",
      kv_role="kv_both",
      kv_connector_extra_config={"shared_storage_path": "local_storage"},
  )

  llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0", kv_transfer_config=ktc)
  llm.generate(prompts, sampling_params)

  prefill_done.set()  ## notify decode instance that KV cache is ready

  ## To keep the prefill node running in case the decode node is not done;
  ## otherwise, the script might exit prematurely, causing incomplete decoding.
  try:
      while True:
          time.sleep(1)
  except KeyboardInterrupt:
      print("Script stopped by user.")

def run_decode(prefill_done):
  os.environ["CUDA_VISIBLE_DEVICES"] = "1"

  sampling_params = SamplingParams(temperature=0, top_p=0.95)

  ktc=KVTransferConfig(
      kv_connector="SharedStorageConnector",
      kv_role="kv_both",
      kv_connector_extra_config={"shared_storage_path": "local_storage"},
  )

  llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0", kv_transfer_config=ktc)

  prefill_done.wait()  ## block waiting for KV cache from prefill instance

  ## Internally it'll first fetch KV cache before starting the decoding loop
  outputs = llm.generate(prompts, sampling_params)

if __name__ == "__main__":
  prefill_done = Event()
  prefill_process = Process(target=run_prefill, args=(prefill_done,))
  decode_process = Process(target=run_decode, args=(prefill_done,))

  prefill_process.start()
  decode_process.start()

  decode_process.join()
  prefill_process.terminate()

vLLM 中的步骤如下：

1. 实例化（引擎构造期间）, connector 会在两个地方被创建：
   • 在 worker 的设备初始化流程中（init_worker_distributed_environment 下），以角色 “worker” 创建 connector；
   • 在 scheduler 的构造函数中，以角色 “scheduler” 创建 connector；
2. Cache 查找：当 scheduler 从等待队列处理 prefill 请求（在本地 prefix-cache 检查之后），会调用 connector 的 get_num_new_matched_tokens，以检测 KV-cache 服务中是否存在外部缓存的 token。prefill 场景下该值始终为 0；decode 场景下可能命中。结果会在调用 allocate_slots 之前加到本地计数中；
3. 状态更新：scheduler 随后调用 connector.update_state_after_alloc，记录命中 cache 的请求（对 prefill 而言是 no-op）；
4. metadata构建：在调度末尾，scheduler 调用 meta = connector.build_connector_meta：
   • prefill 将所有 is_store=True 的请求加入（用于上传 KV）；
   • decode 将所有 is_store=False 的请求加入（用于获取 KV）；
5. Context Manager：在前向传播之前，engine 进入一个 KV-connector 的 context manager：
   • 进入时：调用 kv_connector.start_load_kv，对 decode 而言，它会从外部服务器加载 KV 并注入到分页内存；对 prefill 而言是 no-op；
   • 退出时：调用 kv_connector.wait_for_save，对 prefill 而言，它会阻塞直至 KV 上传到外部服务器；对 decode 而言是 no-op。

下图给出一个可视化示例：

• 对于 SharedStorageConnector，所谓的“external server”其实只是本地文件系统。
• 依据配置，KV 也可以按层进行传输（在每个注意力层前/后进行加载/保存）。
• decode 仅在其请求的第一步加载一次外部 KV；之后便在本地进行计算与存储。

From UniprocExecutor to MultiProcExecutor

在核心技术已经就位后，我们可以开始讨论扩容（scaling up）。假设你的模型权重已经无法放入单张 GPU 的显存。首选方案是在同一节点上使用张量并行（tensor parallelism，TP），将模型在多块 GPU 之间分片（例如 TP=8）。如果模型仍然无法容纳，下一步是跨节点的流水线并行（pipeline parallelism，PP）。但是在实际操作中，我们注意到几个点：

• 同一节点内的带宽（intranode）显著高于跨节点（internode），这也是为什么在实践中通常优先选择张量并行（TP）而非流水线并行（PP）。同时也成立的是，PP 的通信量通常少于 TP。
• 本文不讨论专家并行（expert parallelism，EP），因为我们聚焦的是标准 Transformer 而非 MoE；也不覆盖序列并行（sequence parallelism），原因是 TP 与 PP 在实践中最为常用。

到了这个阶段，我们需要多个 GPU 进程（workers）以及一个编排层来协调它们。这正是 MultiProcExecutor 所提供的能力。

vLLM 的 MultiProcExecutor 运行机制如下：

1. 初始化阶段：MultiProcExecutor 创建 rpc_broadcast_mq 消息队列（底层以共享内存实现）。
2. 进程派生：构造函数按 world_size 循环（例如 TP=8 ⇒ world_size=8），通过 WorkerProc.make_worker_process 为每个 rank 派生守护进程。
3. 管道建立：对每个 worker，父进程先创建一对 pipe（reader 与 writer）。
4. 子进程入口：新进程运行 WorkerProc.worker_main，在其中实例化 worker，并按 UniProcExecutor 的同样顺序进行 init device、load model 等初始化。
5. 角色判定与队列设置：每个 worker 判断自己是否为驱动（TP 组的 rank 0）或普通 worker；并各自设置两条队列：
   • rpc_broadcast_mq（与父进程共享），用于接收工作；
   • worker_response_mq，用于向父进程发送执行结果。
6. 进程间协调完成：初始化期间，每个子进程通过 pipe 将其 worker_response_mq 句柄发回父进程；父进程在收齐所有句柄后解除阻塞，这一步标志着协调完成。
7. 工作循环：workers 进入忙循环，阻塞在 rpc_broadcast_mq.dequeue；工作项到达后，执行该项（路径与 UniProcExecutor 相同，但内容为 TP/PP 特定的分片任务），并通过 worker_response_mq.enqueue 发送结果。
8. 运行时调度：当请求抵达引擎，MultiProcExecutor 会以非阻塞方式将其广播入队到所有子 worker 的 rpc_broadcast_mq；随后在指定输出 rank 的 worker_response_mq.dequeue 上等待，以收集最终结果。

从引擎视角，所有的接口保持不变。多进程的复杂度被 model executor.execute_model 通过调用 model executor 的 execute_model 函数所抽象：

• UniProcExecutor：execute_model 直接触达单个 worker 的 execute_model；
• MultiProcExecutor：execute_model 通过 rpc_broadcast_mq 间接触达每个 worker 的 execute_model。

借助上述机制，我们可以在不改变引擎接口的前提下，按资源上限运行更大的模型。下一步是继续横向扩展：启用数据并行（DP > 1）在多节点复制模型，加入轻量级 DP 协调层，做跨副本的负载均衡，并在前面部署一个或多个 API 服务器处理进入流量。

Distributed system serving vLLM

在生产环境中，搭建推理服务基础设施的方式有很多。为保持具体，这里举一个例子：假设我们有两台 H100 节点，并希望在它们上运行四个 vLLM 引擎。如果模型需要 TP=4，可以将节点按如下方式进行配置。

在第一台节点上，以 headless mode （不启用 API 服务器）运行引擎，使用如下参数：

vllm serve <model-name>
  --tensor-parallel-size 4
  --data-parallel-size 4
  --data-parallel-size-local 2
  --data-parallel-start-rank 0
  --data-parallel-address <master-ip>
  --data-parallel-rpc-port 13345
  --headless

以同样的命令在另一个节点上运行，但进行两处调整：不使用 --headless，并修改 --data-parallel-start-rank 为 2。

vllm serve <model-name>
  --tensor-parallel-size 4
  --data-parallel-size 4
  --data-parallel-size-local 2
  --data-parallel-start-rank 2
  --data-parallel-address <master-ip>
  --data-parallel-rpc-port 13345

On the headless server node

在 headless 节点上，CoreEngineProcManager 会启动 2 个进程（由 --data-parallel-size-local 指定），每个进程运行 EngineCoreProc.run_engine_core。这些函数各自创建一个 DPEngineCoreProc（引擎核心），随后进入其忙循环。

DPEngineCoreProc 会初始化其父 EngineCoreProc（EngineCore 的子组件），其步骤包括：

1. 创建 input_queue 与 output_queue（queue.Queue）；
2. 使用 DEALER ZMQ 套接字（异步消息库）与另一节点的前端进行初始握手，并接收协调地址信息；
3. 初始化数据并行（DP）通信组（例如使用 NCCL 后端）；
4. 使用 MultiProcExecutor 初始化 EngineCore（如前述，在 4 张 GPU 上配置 TP=4）；
5. 创建 ready_event（threading.Event）；
6. 启动输入守护线程（threading.Thread）运行 process_input_sockets(..., ready_event)，并以类似方式启动输出线程；
7. 在主线程中等待 ready_event，直到跨 2 节点的全部 4 个进程的输入线程完成协调握手，最后执行 ready_event.set()；
8. 一旦解除阻塞，向前端发送携带 metadata 的“ready”消息（例如分页 KV 缓存内存中可用的 num_gpu_blocks）；
9. 随后主线程、输入线程和输出线程分别进入各自的忙循环。

最终会有 4 个子进程（每个对应一个 DP 副本），每个进程都运行主线程、输入线程和输出线程。它们与 DP 协调器和前端完成协调握手后，三个线程将进入稳态的忙循环。

当前稳态运行：

• 输入线程 — 阻塞在 input socket 上，直到从 API 服务器路由来请求；收到请求后，解码载荷，通过 input_queue.put_nowait(...) 将工作项入队，然后返回继续阻塞在套接字上；
• 主线程 — 在 input_queue.get(...) 上被唤醒，将请求传递给引擎；MultiProcExecutor 运行前向传播并将结果入队到 output_queue；
• 输出线程 — 在 output_queue.get(...) 上被唤醒，将结果发送回 API 服务器，然后恢复阻塞状态。

附加机制：

• DP 波次计数器 — 系统跟踪”波次”；当所有引擎变为空闲时它们会静默，当新工作到达时计数器递增（用于协调/指标）。
• 控制消息 — API 服务器可以发送的不仅仅是推理请求（例如中止和实用/控制 RPC）。
• 锁步的虚拟步骤 — 如果任何 DP 副本有工作，所有副本都执行前向步骤；没有请求的副本执行虚拟步骤以参与必需的同步点（避免阻塞活跃副本）。
• 锁步澄清：这实际上只对 MoE 模型是必需的，其中专家层形成 EP 或 TP 组，而注意力层仍然是 DP。目前总是与 DP 一起执行 - 这只是因为”内置”非 MoE DP 的用途有限，因为你可以只运行多个独立的 vLLM 并以正常方式在它们之间进行负载均衡。

现在来看第二部分，API 服务器节点上发生了什么？

On the API server node

在前端（API 服务器）节点，我们实例化一个 AsyncLLM 对象（对 LLM 引擎的 asyncio 封装）。其内部会创建 DPLBAsyncMPClient（数据并行、负载均衡、异步、多进程客户端）。

在 MPClient 的父类中，launch_core_engines 函数会执行，并：

1. 创建用于启动握手的 ZMQ 地址（与 headless 节点上的做法一致）；
2. 启动一个 DPCoordinator 进程；
3. 创建一个 CoreEngineProcManager（与 headless 节点相同）。

在 AsyncMPClient（MPClient 的子类）中，我们：

1. 创建 outputs_queue（asyncio.Queue）；
2. 创建 asyncio 任务 process_outputs_socket，通过输出套接字与所有 4 个 DPEngineCoreProc 的输出线程通信，并写入 outputs_queue；
3. 随后由 AsyncLLM 启动另一个 asyncio 任务 output_handler，从该队列读取，最终将信息发送到 create_completion 函数。

在 DPAsyncMPClient 中，我们创建 asyncio 任务 run_engine_stats_update_task，与 DP 协调器通信。

DP 协调器在前端（API 服务器）与后端（引擎核心）之间进行调解，它：

• 周期性地向前端的 run_engine_stats_update_task 发送负载均衡信息（队列大小、等待/运行中的请求）；
• 处理来自前端的 SCALE_ELASTIC_EP 命令，通过动态改变引擎数量（仅在 Ray 后端可用）；
• 向后端发送 START_DP_WAVE 事件（由前端触发时），并回报波次状态更新。

总结一下，前端（AsyncLLM）运行若干 asyncio 任务（注意：并发而非并行）：

• 一类任务通过 generate 路径处理输入请求（每个新的客户端请求都会生成一个新的 asyncio 任务）；
• 两个任务（process_outputs_socket、output_handler）处理来自底层引擎的输出消息；
• 一个任务（run_engine_stats_update_task）维持与 DP 协调器的通信：发送波次触发、轮询负载均衡状态，以及处理动态扩缩请求。

最后，主服务器进程创建一个 FastAPI 应用并挂载诸如 OpenAIServingCompletion 和 OpenAIServingChat 的端点，提供 /completion、/chat/completion 等接口；整个栈通过 Uvicorn 对外服务。

把这些拼在一起，就是完整的请求生命周期！你会在终端中发送：

curl -X POST http://localhost:8000/v1/completions -H "Content-Type: application/json" -d '{
  "model": "TinyLlama/TinyLlama-1.1B-Chat-v1.0",
  "prompt": "The capital of France is",
  "max_tokens": 50,
  "temperature": 0.7
}'

接下来会发生什么：

1. 请求命中 API 服务器上的 OpenAIServingCompletion.create_completion 路由；
2. 该函数以异步方式对提示进行分词，并准备 metadata（请求 ID、采样参数、时间戳等）；
3. 随后调用 AsyncLLM.generate，其流程与同步引擎一致，最终会触发 DPAsyncMPClient.add_request_async；
4. 该调用进一步执行 get_core_engine_for_request，根据 DP 协调器的状态在各个引擎间做负载均衡（选择得分最低/负载最低的引擎：score = len(waiting) * 4 + len(running)）；
5. 将 ADD 请求发送到选定引擎的 input_socket；
6. 在该引擎上：
   • 输入线程 — 解除阻塞，从输入套接字解码数据，并将工作项放入主线程的 input_queue；
   • 主线程 — 在 input_queue 上解除阻塞，将请求加入引擎，并反复调用 engine_core.step()，在满足停止条件前不断将中间结果入队到 output_queue；
   • 输出线程 — 在 output_queue 上解除阻塞，并通过输出套接字将结果发送回去；
7. 这些结果会触发 AsyncLLM 的输出类 asyncio 任务（process_outputs_socket 与 output_handler），它们把 token 逐步传递回 FastAPI 的 create_completion 路由；
8. FastAPI 将附加 metadata（结束原因、logprobs、使用信息等），并通过 Uvicorn 返回 JSONResponse 到你的终端！

就这样，你的补全结果返回了——整个分布式系统都隐藏在一个简单的 curl 命令背后！当增加更多 API 服务器时，负载均衡主要由操作系统/套接字层处理。从应用视角看，几乎无需改动——复杂性被抽象掉了。而当使用 Ray 作为 DP 后端时，可以暴露一个 URL 端点（/scale_elastic_ep），以实现对引擎副本数量的自动扩缩。

Benchmarks and auto-tuning - latency vs throughput

到目前为止，我们一直在分析 “gas particles” ——请求如何在 engine /系统内部流动的细节。现在是时候拉远视角，整体审视系统，并提出一个问题：如何度量一个推理系统的性能？

在最高层面，有两项彼此 “竞争” 或者说 “冲突” 的指标：

• 延迟（Latency）——从请求提交到返回 tokens 的时间
• 吞吐量（Throughput）——系统每秒能够生成/处理的 tokens/请求数量

延迟在交互式应用中最为重要，因为用户在等待响应。吞吐量在离线工作负载中更重要，例如用于预训练/后训练运行的合成数据生成、数据清洗/处理，以及一般的离线批量推理作业。在解释为何延迟与吞吐量会互相“竞争”之前，先定义一些常见的推理指标：

下面给出一个简化模型来解释这两项指标为何相互“竞争”。假设：权重 I/O 而非 KV 缓存 I/O 占主导，即我们处理的是短序列。

当观察批大小 B 对单次解码步的影响时，这种权衡会变得清晰：当 B ↓ → 1 时，ITL（Inter-Token Latency，token 间延迟）下降——每步的工作更少，且该 token 不再与其他 token “竞争”；当 B ↑ → ∞ 时，ITL 上升，因为每步需要执行更多 FLOPs；但吞吐量会提高（直到达到峰值性能），因为权重 I/O 被更多 token 分摊。

屋顶线（Roofline）模型有助于理解：当批量低于饱和批 B_sat 时，步时由 HBM 带宽主导（逐层将权重流入片上内存），因此步延迟近乎平坦——计算 1 个与 10 个 token 所需时间相近。超过 B_sat 后，内核会转为计算受限，步时大致随 B 增长；每增加一个 token 都会增加 ITL。

更严谨的分析还需要考虑内核自动调优（kernel auto-tuning）：随着批量 B 增大，运行时可能会针对该形状切换到更高效的内核，从而改变实际达到的性能 P_kernel。步延迟可表示为 t = FLOPs_step / P_kernel，其中 FLOPs_step 是该步的计算工作量。可以看到，当 P_kernel 接近峰值性能 P_peak 时，每步的计算量增加会直接导致延迟上升。

How to benchmark in vLLM

vLLM 提供 vllm bench {serve,latency,throughput} 命令行工具（CLI），它封装了 vllm/benchmarks/{server,latency,throughput}.py 三个脚本，便于统一运行与统计。

脚本功能如下：

• latency：使用较短的输入（默认 32 tokens），以较小的批（默认 8）采样 128 个输出 token。脚本会运行多次迭代，并报告该批的端到端（e2e）延迟。
• throughput：一次性提交一组固定的提示（默认：1000 条 ShareGPT 样本），即 QPS=Inf 模式；并在整次运行中统计并报告输入/输出/总 token 以及每秒请求数（RPS）。
• serve：启动一个 vLLM 服务器，并通过从泊松分布（或更一般地，Gamma 分布）抽取请求到达间隔来模拟真实世界负载。在给定时间窗内发送请求，测量前述各项指标；同时可以选择在服务端启用最大并发限制（通过信号量实现，例如限制为 64 个并发请求）。

示例：运行 latency 脚本（其中Benchmark configs used in CI live under .buildkite/nightly-benchmarks/tests.）

vllm bench latency
  --model <model-name>
  --input-tokens 32
  --output-tokens 128
  --batch-size 8

此外，还提供一个自动调优脚本：它通过驱动 serve 基准测试来搜索满足目标 SLO 的参数设置（例如：“在保持 p99 端到端延迟 < 500 ms 的同时最大化吞吐量”），并返回一个建议配置。

Epilogue

我们从基本的引擎核心（UniprocExecutor）开始，添加了投机解码和前缀缓存等高级功能，扩展到 MultiProcExecutor（TP/PP > 1），最后进行横向扩展，将所有内容包装在异步引擎和分布式服务堆栈中——最后介绍了如何测量系统性能。

vLLM 还包含一些被我们略过的专门处理。例如：

• Diverse hardware backends：TPUs、AWS Neuron（Trainium/Inferentia）等；
• Architectures/techniques：MLA、MoE、编码器-解码器（例如 Whisper）、池化/嵌入模型、EPLB、m-RoPE、LoRA、ALiBi、无注意力变体、滑动窗口注意力、多模态 LM 和状态空间模型（例如 Mamba/Mamba-2、Jamba）；
• TP/PP/SP；
• Hybrid KV-cache logic (Jenga)，更复杂的采样方法如束搜索采样等；
• 实验性功能：异步调度。

好的一点是，这些大部分都与上述描述的主要流程正交——你几乎可以将它们视为”插件”（当然，在实践中存在一些耦合）。

References

1. vLLM
2. “Attention Is All You Need”
3. “Efficient Memory Management for Large Language Model Serving with PagedAttention”
4. “DeepSeek-V2: A Strong, Economical, and Efficient Mixture-of-Experts Language Model”
5. “Jenga: Effective Memory Management for Serving LLM with Heterogeneity”
6. “Orca: A Distributed Serving System for Transformer-Based Generative Models”
7. “XGrammar: Flexible and Efficient Structured Generation Engine for Large Language Models”
8. “Accelerating Large Language Model Decoding with Speculative Sampling”
9. “EAGLE: Speculative Sampling Requires Rethinking Feature Uncertainty”
10. “Medusa: Simple LLM Inference Acceleration Framework with Multiple Decoding Heads”
11. LMCache

Read More

在探讨GPU如何工作之前，我们首先要回答一个更根本的问题：为什么AI的发展离不开GPU？

答案的核心在于AI，尤其是深度学习，其本质是大规模的矩阵和向量运算。无论是训练一个复杂的神经网络，还是运行一个大语言模型（LLM）进行推理，其底层都涉及海量的、可以被分解为独立部分的数学计算。例如，在神经网络的前向传播中，每一层的输出都是由前一层的输入与权重矩阵相乘，再加上偏置向量得到的。这个过程需要对成千上万个数字进行重复的、模式相同的乘法和加法运算。

这种计算任务的特点是：计算量巨大，但逻辑简单且高度并行。这正是GPU大显身手的舞台。GPU的设计初衷是为了处理图形渲染，而图形渲染本身也是一种高度并行的任务——屏幕上的每个像素点都可以被独立计算。这种为并行而生的架构，恰好与AI的算力需求不谋而合。因此，当我们谈论AI Infra时，GPU便成为了绕不开的基石。

一、 CPU vs. GPU：不同的设计哲学

要理解GPU，最好的方式就是将它与我们更熟悉的CPU（中央处理器）进行对比。它们虽然都是处理器，但设计理念却截然不同，这决定了它们各自擅长的领域。

CPU：精于串行的“瑞士军刀”

CPU被设计成一个通用的、能够处理各种复杂任务的“多面手”。它的核心特点是：

• 少数强大的核心：一个典型的CPU通常只有几个到几十个核心。
• 复杂的控制逻辑：每个核心都配备了复杂的控制单元、分支预测器和大量的缓存，使其能够快速处理各种复杂的指令和逻辑判断。
• 高时钟频率：CPU追求单个任务的极致执行速度。

你可以把一个CPU核心想象成一位经验丰富、能力全面的老教授。他能处理各种疑难杂症，无论是复杂的逻辑推理还是精密的计算，都能应对自如。但他的精力有限，无法同时处理成千上万个简单问题。因此，CPU非常适合处理操作系统、应用程序等需要复杂逻辑判断和串行执行的任务。

GPU：擅长并行的“人海战术”

与CPU不同，GPU的设计目标非常专一：大规模并行计算。它的核心特点是：

• 成千上万的简单核心：一块GPU芯片上集成了数千个甚至更多的计算核心（如NVIDIA的CUDA核心）。
• 简化的控制逻辑：每个核心的控制逻辑和缓存都非常简单，它们被设计用来高效地执行同一条指令。
• 高内存带宽：为了同时喂饱成千上万个核心，GPU配备了高带宽内存（HBM），确保数据能够被快速地传输。

你可以把GPU想象成一个由成千上万名小学生组成的计算方阵。虽然每个小学生的计算能力有限，也无法处理复杂的逻辑问题，但当你给他们下达一个简单的、统一的计算任务时（比如“所有人都计算1+1”），他们能以惊人的速度同时完成成千上万次计算。这正是深度学习所需要的。

(图注：CPU架构示意图，其特点是拥有强大的ALU（算术逻辑单元）和大量的缓存，但计算核心数量较少;GPU架构示意图，其特点是拥有海量的ALU，控制单元和缓存相对简单)

总而言之，CPU和GPU的设计差异决定了它们在AI计算中的分工：CPU负责整体的逻辑控制和任务调度，而GPU则专注于执行那些计算密集型、高度并行的核心任务。 理解了这种根本性的差异，我们才能更好地深入GPU的内部，探究其并行计算的奥秘。

二、 深入GPU架构：并行计算的奥秘

理解了GPU“人海战术”的设计哲学后，我们来进一步拆解其内部结构，看看这支庞大的“计算军团”是如何被组织和调度的。

SIMD/SIMT：并行计算的灵魂

GPU之所以能实现大规模并行，核心在于其计算模型。你可能听说过两个术语：SIMD（单指令，多数据）和SIMT（单指令，多线程）。

• SIMD（Single Instruction, Multiple Data）：这是并行计算的一种经典模型。它意味着用一条指令同时对多个数据执行相同的操作。想象一下，老师对一个班的学生说：“请大家把手里的数字都加上5”。在这里，“加5”就是单条指令，而每个学生手里的不同数字就是多份数据。

• SIMT（Single Instruction, Multiple Threads）：这是NVIDIA在CUDA架构中提出的模型，可以看作是SIMD在GPU上的升级版和更灵活的实现。它将成千上万的计算任务包装成线程（Thread），然后将32个线程组成一个线程束（Warp）。同一个Warp中的所有线程在同一个时钟周期内执行相同的指令，但每个线程可以处理不同的数据。SIMT模型的美妙之处在于它对开发者更友好，屏蔽了底层硬件的复杂性。你只需要编写单个线程要执行的程序，GPU的硬件调度器会自动将它映射到成千上万个核心上去并行执行。

在传统的标量计算模型中，CPU的一条指令一次只能操作单个数据。例如，一次浮点加法就是 double + double。当处理如图形、音频或科学计算中常见的大规模数据集时，这种“一次一个”的模式效率极低，因为我们需要对海量数据重复执行完全相同的操作，这暴露了标量处理的瓶颈。为了打破这个瓶颈，现代CPU集成了SIMD（单指令，多数据）架构。CPU增加了能容纳多个数据元素的宽向量寄存器（如256位的YMM寄存器），以及能够并行处理这些数据的执行单元。

无论是SIMD还是SIMT，其本质都是用一条指令驱动海量的计算单元，这是GPU实现超高计算吞吐量的根本。

• Warp/Wavefront 尺寸：NVIDIA 典型 warp=32，AMD wavefront=64。算法与 block 配置需对齐 warp 大小以避免“尾线程”浪费。
• 分支发散（Warp Divergence）：同一 warp 内线程走不同分支，会串行化执行各分支路径，直至在重收敛点合流；可通过数据重排、位掩码/谓词化减少发散。
• 内存合并访问（Coalescing）：warp 内连续地址访问可合并为更少事务，显著提升带宽利用；未对齐或跨行访问会降低有效带宽。
• 占用度（Occupancy）与隐藏延迟：足够多的活跃 warp 能隐藏访存与流水线延迟；占用度受限于寄存器/共享内存/线程数配额。
• 参考：
  • CUDA C++ Programming Guide（SIMT/Thread Hierarchy/Memory Coalescing）: https://docs.nvidia.com/cuda/cuda-c-programming-guide/
  • CUDA Best Practices Guide: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
  • Nsight Compute User Guide: https://docs.nvidia.com/nsight-compute/

CUDA Core、Tensor Core与TPU：专业分工的计算单元

GPU的“计算军团”并非由单一兵种构成，而是由不同类型的专业计算单元组成，以应对不同的任务需求。

2007年，NVIDIA正式推出了CUDA平台。CUDA的革命性在于，它提供了一套简单的编程模型，让开发者能用近似C语言的方式，轻松地驾驭GPU内部成百上千个并行核心。 开发者无需再关心复杂的图形接口，可以直接编写在数千个线程上并发执行的程序。至此终结了GPGPU编程的蛮荒时代，让GPU计算真正走下神坛，成为开发者触手可及的强大工具。

• CUDA Core：这是GPU最基本的计算单元，主要负责执行单精度浮点（FP32）和整数运算。你可以把它看作是GPU里的“普通士兵”，负责执行大部分通用的并行计算任务。

• Tensor Core：这是NVIDIA从Volta架构开始引入的、专为深度学习打造的“特种部队”。Tensor Core专门用于执行大规模的矩阵乘加运算（Matrix Multiply-Accumulate, MMA），并且在硬件层面直接支持混合精度（FP16/FP32）和低精度（INT8/INT4）计算。在一次操作中，一个Tensor Core可以完成一个4x4的矩阵乘法，其效率远超CUDA核心。对于大模型训练和推理中无处不在的矩阵运算，Tensor Core能够带来数倍的性能提升。

• 与TPU的对比：Google的TPU（Tensor Processing Unit）是另一个为AI而生的专用处理器。如果说Tensor Core是GPU里的“特种部队”，那TPU就是一支纯粹的“矩阵运算专业军团”，它将整个芯片的设计都聚焦于此，因此在特定任务上能效比极高。而GPU则更像是一个通用平台，既有CUDA核心处理通用并行任务，又有Tensor Core加速AI任务。

进阶要点与参考（数值格式与算子栈）：

• TF32/BF16：Ampere 引入 TF32（10-bit mantissa）在不改动 FP32 模型的前提下大幅提速；Hopper 进一步强化 BF16/FP8 的吞吐。
• 低比特推理：INT8/INT4 常配合量化感知训练（QAT）/后量化（PTQ）与校准；2:4 结构化稀疏在 Ampere 支持硬件加速。
• 编程接口：CUDA 的 wmma/mma.sync/cuBLASLt 提供 tile 级 MMA；TensorRT 在图级融合/格式选择上进一步榨干 Tensor Core 性能。
• 参考：
  • NVIDIA Ampere Architecture WP: https://resources.nvidia.com/en-us-architecture/ampere-architecture-whitepaper
  • NVIDIA Hopper Architecture WP（FP8）: https://resources.nvidia.com/en-us-architecture/hopper-architecture-whitepaper
  • TensorFloat-32 介绍: https://developer.nvidia.com/blog/tensorfloat-32-precision-format/
  • cuBLAS/cuBLASLt 文档: https://docs.nvidia.com/cublas/
  • TensorRT 文档: https://docs.nvidia.com/deeplearning/tensorrt/

计算单元

GPU中的计算单元是GPC（Graphics Processing Cluster），而一个GPC包含多个TPC（Texture Processing Cluster），而一个TPC中则包含多个SM（Streaming Multiprocessor），SM是GPU执行计算任务的核心单元，每个SM都是一个高度独立的计算单元。

如果说CUDA核心是士兵，那么SM（Streaming Multiprocessor，流式多处理器）就是军营里的“指挥官”。它包含了：

• 一定数量的CUDA核心和Tensor Core。
• 自己的指令缓存和调度器。
• 一小块高速的共享内存（Shared Memory），寄存器（Register File）和L1缓存（L1 Data Cache / Instruction Cache）。
• Warp调度器（Warp Scheduler）等关键组件。

SM是GPU执行任务的核心单位。当一个计算任务（在CUDA中称为Kernel）被启动时，它会被分解成一个个线程块（Thread Block），然后这些线程块被分配到不同的SM上去执行。SM内部的调度器再将线程块分解成线程束（Warp），并安排它们在CUDA核心或Tensor Core上执行。

SM就像一个自给自足的计算工厂，它接收任务，管理资源，调度执行，并最终产出结果。正是由成百上千个这样的“工厂”协同工作，才构成了GPU强大的并行处理能力。

(图注：单个SM的架构图)

• 寄存器/共享内存配额：每个 SM 的寄存器数与 Shared Memory 容量是硬上限；单 kernel 的寄存器/SMEM 使用决定了每个 SM 可驻留的 Block/warp 数量。
• 活跃 warp 与吞吐：更高活跃 warp 有助于隐藏访存与流水线延迟，但并非越高越好；请结合 Nsight Compute 的 achieved_occupancy 与 sm_efficiency 评估瓶颈。
• 调度：多 Warp 调度器/发射端口可以每拍发射多个指令；指令依赖、结构冲突与内存等待共同决定实际 IPC。
• 参考：
  • Nsight Compute 指标说明: https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html
  • CUDA Occupancy Calculator: https://developer.nvidia.com/cuda-occupancy-calculator
  • CUDA C++ Programming Guide（Occupancy/Execution Resources）: https://docs.nvidia.com/cuda/cuda-c-programming-guide/

内存层次

想象一下，你拥有一个成千上万人的计算军团，但如果后勤补给（数据供应）跟不上，这支军团的战斗力将大打折扣。这就是所谓的“内存墙”问题——计算单元的速度远超内存访问速度，导致计算单元不得不花费大量时间等待数据。

为了解决这个问题，现代高端GPU普遍采用HBM（High Bandwidth Memory，高带宽内存）。与传统的DDR内存不同，HBM通过以下方式实现了超高的带宽：

• 3D堆叠：HBM将多个DRAM芯片垂直堆叠起来，并通过硅通孔（TSV）技术进行连接，极大地增加了数据传输的并行度。
• 宽位宽接口：HBM拥有极宽的内存接口（如1024-bit或更高），远超DDR内存的64-bit。

你可以把DDR想象成一条普通的双车道公路，而HBM则是一条拥有32条车道的超级高速公路。通过HBM，GPU能够以极高的速度从显存中读取和写入数据，从而“喂饱”其内部成千上万个嗷嗷待哺的计算核心。

(图注：HBM通过3D堆叠和宽位宽接口技术实现超高内存带宽)

在一块GPU中，HBM和L2 Cache是整个GPU共享的，而L1 Cache/Shared Memory则是SM维度独享的。Shared Memory是每个SM内部的一块高速、可编程的片上缓存。同一线程块（Block）内的所有线程都可以访问它，速度远快于访问全局显存（HBM）。它是实现Block内线程高效协作和数据交换的核心，对于矩阵乘法等需要数据复用的算法至关重要。

• 合并访问与对齐：按 32/64/128B 对齐的顺序访问最友好；跨 stride 或散乱访问会放大事务数。
• Shared Memory 冲突：注意 bank 冲突；通过交错/转置/向量化降低冲突，或使用 cp.async（支持的架构）做块状搬运。
• L2/L1 策略：合适的 cache hint（如 ld.global.ca/cg）与预取可改善带宽利用；Host↔Device 侧使用 pinned memory/cudaMallocHost 提升 H2D/D2H 吞吐。
• 参考：
  • CUDA C++ Programming Guide（Memory Hierarchy/Coalescing/Shared Memory）: https://docs.nvidia.com/cuda/cuda-c-programming-guide/
  • Ampere 异步拷贝与 cp.async: https://developer.nvidia.com/blog/boosting-application-performance-with-the-new-nvidia-ampere-architecture/
  • Pinned Memory（Page-Locked）: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#pinned-memory

异构计算

(图注：CPU/GPU异构计算架构)

CPU是整个系统的核心，是总指挥，GPU的任务指令是由CPU分配的。CPU通过PCIe总线给GPU发送指令和数据交互。

• 互联：PCIe 为通用路径；NVLink/NVSwitch 提升 GPU↔GPU/CPU 带宽；封装内 NVLink‑C2C/UMA 进一步降低延迟。而PCIe支持DMA和MMIO两种通讯模式:
  • MMIO（内存映射I/O，Memory Mapping I/O）由CPU直接控制数据读写，操作系统会把设备地址映射到CPU的虚拟空间中，适合小数据量的指令交互
  • DMA（直接内存访问，Direct Memory Access）则允许设备绕过CPU直接访问系统内存，专为大数据块的高效传输设计。 CPU通过IMC和Memory Channel访问内存，为了提升数据传输带宽，高端CPU通常会支持多内存通道，即多IMC和Memory Channel的组合，以满足日益增长的数据处理需求。
• 内存：UVM 提供单一地址空间（配合预取/访问提示）；Pinned/Zero‑copy 降低拷贝开销；GPUDirect（P2P/RDMA/GDS）减少绕行。
• 编程：多 stream 与 events 重叠拷贝与计算；CUDA Graphs 降低小内核启动开销；NCCL 负责多 GPU 集体通信。
• 并行：DP/TP/PP/EP 可组合扩展超大模型。
• 部署：MIG 分片、MPS 并发，容器与 K8s 做拓扑感知调度。

进阶要点与参考（互联/内存/编程）：

• 互联拓扑：多 GPU 拓扑（双路 NVLink 对接、NVSwitch 全互连）影响带宽对称性与 AllReduce 性能；NCCL 会自适应拓扑构建环/树。
• UVM 迁移：按页迁移（page migration）可能带来缺页中断与 TLB 抖动；使用 cudaMemAdvise/cudaMemPrefetchAsync 预取可降低抖动。
• Pinned 与可分页内存：可分页内存通常经隐式 staging；显式 cudaMallocHost 可减少一次拷贝并提升吞吐；多 copy 引擎支持并发 H2D/D2H。
• 默认流语义：建议开启 PTDS（Per-Thread Default Stream） 避免旧式全局同步；利用 流优先级 与 事件 管理核间依赖。
• MPS/MIG：MPS 通过命令队列复用提高小核吞吐；MIG 将 SM/HBM/L2 切片隔离以实现多租户 QoS。
• 参考：
  • NVLink/NVSwitch: https://developer.nvidia.com/nvlink
  • NCCL 文档: https://developer.nvidia.com/nccl
  • Unified Memory/UVM: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-overview
  • Default Stream/PTDS: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#stream-ordered-memory-operations
  • CUDA MPS: https://docs.nvidia.com/deploy/mps/index.html
  • NVIDIA MIG User Guide: https://docs.nvidia.com/datacenter/tesla/mig-user-guide/

三、 简易的一个例子

以下这个demo是实现两个长度为 $2^{30}$ (约10亿) 的浮点数数组的相加。其中，一个数组 $(x)$ 的所有元素初始化为 $1.0$，另一个数组 $(y)$ 的所有元素初始化为 $2.0$，我们计算 $y[i] = x[i] + y[i]$。

CPU实现

#include <iostream>
#include <math.h>
#include <chrono>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<30;

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  auto start = std::chrono::high_resolution_clock::now();

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  auto stop = std::chrono::high_resolution_clock::now();
  auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(stop - start);
  std::cout << "CPU 'add' function execution time: " << duration.count() << " ms" << std::endl;

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));

  std::cout << "Max error: " << maxError << std::endl;

  delete [] x;
  delete [] y;

  return 0;
}

性能表现：

g++ add.cpp -o add
time ./add

CPU 'add' function execution time: 3740 ms
Max error: 0

real 0m21.418s
user 0m15.798s
sys 0m4.400s

• 计算耗时: 核心的add函数耗时 3740毫秒。
• 总耗时: 整个程序从启动到结束（real time）耗时 21.4秒。这额外的时间主要消耗在分配8GB内存（new float[N]）以及初始化数组上。

GPU实现

包含步骤：

1. 分配内存: 分别在CPU（Host）和GPU（Device, cudaMalloc）上分配内存。
2. 数据传输 (H2D): 将CPU上的输入数据 (h_x, h_y) 拷贝到GPU显存 (d_x, d_y)。
3. 执行Kernel函数: 在GPU上启动addKernel函数，利用其大规模并行能力进行计算。
4. 数据传回 (D2H): 将GPU计算完成的结果 (d_y) 拷贝回CPU内存 (h_y) 以便后续使用或验证。

#include <iostream>
#include <math.h>

#define CUDA_CHECK(call) \
do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error in %s at line %d: %s
          

            Read More