一个 32B 的模型在 8 张 H100 上做共置的 GRPO。训练时每张卡上要有 66 GB 的训练状态（16 字节 / 参数均分到 8 卡），还要跑参考模型的前向；生成时每张卡上要有 8 GB 的推理权重副本（TP = 8）加尽量大的 KV 池。两样东西加起来 150 GB，卡只有 80 GB——所以每一步里显存要换两次主人：生成前训练状态搬到 CPU、推理引擎的权重与 KV 池挂回 GPU；训练前反过来。按 PCIe 的带宽算，一步两次切换搬 146 GB 左右、约 7 秒；步时间 5 分钟时是 2%。

2% 不值得写一篇。值得写的是这 7 秒之外的东西：推理引擎的 KV 池释放再建回来，里面的 prefix cache 全部作废；CUDA graph 捕获的是显存地址，地址变了图就废了——vLLM 的 sleep mode 用 CUDA 虚拟内存 API 把物理页摘下来而保留虚拟地址，正是为了不重捕获；训练器的 PyTorch caching allocator 与推理引擎的 CuMemAllocator 在同一块 HBM 上交替，一方保留着没释放的段另一方就挂不回来，所以切换前后要 empty_cache、要开关 expandable_segments；vLLM 启动时按 gpu_memory_utilization 定下的 KV 池大小是一次性的，共置下这个数字通常只有 0.5，推理引擎在共置形态下只拿到半张卡的 KV 池，decode 吞吐比独占一张卡低 10–30%——这才是共置真正的税，而它不出现在”切换耗时”这个指标里。

这一篇把一次切换拆成每个动作：谁释放什么、走哪条链路、要几秒、留下什么隐性代价；然后看 vLLM 的 sleep / wake_up 与 CuMemAllocator、SGLang 的 torch_memory_saver、FSDP 与 Megatron 的 offload 路径各是怎样实现”让渡”的；再看 verl 的 hybrid worker——ActorRolloutRefWorker，一个进程里同时持有训练器、推理引擎与参考模型——怎样把这些动作排成序；最后回答共置的适用边界：模型多大、回答多长时该换分离。

本篇的核心问题：

一个 32B 模型在 8 卡共置。每步开始生成前要把 16 字节/参数的训练状态搬走、把 KV 池建起来，训练前再反过来。每个动作搬多少字节、走哪条链路、要几秒？¹ 步时间 5 分钟时这是 2% 还是 20%？²

版本：vLLM v0.27.1（vllm/device_allocator/cumem.py、vllm/v1/worker/gpu_worker.py）、verl v0.9.0（verl/workers/engine_workers.py、verl/workers/rollout/vllm_rollout/）、PyTorch 2.13 的 FSDP2。硬件按 H100 SXM：80 GB HBM3、NVLink 节点内 450 GB/s 单向、PCIe Gen5 x16 到 CPU 的实测 pinned 拷贝按 25 GB/s 算（标称 64 GB/s）。

一、总览

1. 8 卡共置：每步搬多少字节、走哪条链路

32B（Qwen3-32B 规格，32.8B 参数）、8 卡共置、FSDP 训练、vLLM TP = 8 推理，一步里显存的两次换手：

步时间 5 分钟时约 2%；8B 模型 64 卡每卡 4.5 GB、0.2 秒，可忽略；DeepSeek-V3 规格 256 卡每卡 94 GB、约 4 秒，对几十分钟的步也可忽略。纯搬运时间几乎从来不是放弃共置的理由。

真正的账在别处：

• KV 池只有半张卡：vLLM 按 gpu_memory_utilization（verl 共置默认 0.5）在启动时一次性划定 KV 块数，之后 sleep / wake 只是同一批块的摘下挂回。8B 场景 KV 池从 56 GB 缩到 20 GB，并发从 92 降到 33，每卡 decode 吞吐从 2600 降到 1900 token/s——生成时间多四分之一，远超切换的秒数。开了训练侧的 param / optimizer offload 后可以把它提到 0.7–0.85，但要给训练侧的常驻部分（CUDA context、NCCL 缓冲、CUDA graph 池、碎片）留出余量。
• 常驻的东西：CUDA graph 的显存池不能 offload（verl 配置里明说了这一点，给了 cudagraph_capture_sizes 让你缩小它），两套 NCCL 通信器（训练器的、推理引擎的）各有缓冲，两个 allocator 各自保留的段——加起来几 GB 到十几 GB，两种形态都要让出。
• 主机内存：FSDP 的 offload 与 vLLM level 1 的 sleep 都要 pinned 主机内存——8 卡 32B 每卡 65 GB 训练状态就是 520 GB pinned，一台机器 1–2 TB 的内存被吃掉一半；level 2 的 sleep 不留 CPU 副本正是为省这一块。

2. 让渡的三种做法

在同一块显存上让两方轮流，只有三种做法，三个框架层各用一种：

vLLM 的 sleep mode 做的关键一步，是让”丢弃再重建”不改变虚拟地址：cuMemUnmap 把物理页从虚拟地址上摘下、cuMemRelease 归还物理内存，虚拟地址空间仍然保留；wake_up 时 cuMemCreate 新的物理页再 cuMemMap 到同一个地址。于是所有持有这些地址的东西——CUDA graph、KV 块表、模型的 parameter 对象——都不用重建。第三章讲它。

3. 本文的章节安排

二、显存归属的账

1. 一张卡上的两套东西

共置形态下一张卡在一步里先后要放下（32B、8 卡、FSDP、vLLM TP = 8）：

训练侧的 ①–④ 全放卡上就已经把 80 GB 用光——32B 在 8 卡上即便不共置，FSDP 也要 offload 优化器状态（或者用 Megatron 的分布式优化器加 TP，每卡的状态字节相同但激活更省）。加上推理侧，两侧合计 120 GB 以上，每一步必须换手两次。

8B 的账要宽松得多：训练侧每卡 16 GB + 2 GB，推理侧 16 GB 权重 + KV 池，两侧加起来放得下——但 KV 池想要尽量大，训练的激活也要空间，实际仍然换手，只是搬的少（64 卡时每卡每步 4.5 GB）。

2. 常驻的部分

有几样东西谁也让不出去，两种形态下都占着：

• 两个 CUDA context 与两套 NCCL 通信器：训练器的进程组（FSDP 的 all-gather / reduce-scatter）与推理引擎的进程组（TP 的 all-reduce）是独立的，各有自己的缓冲区与 channel；同一进程里两套通信器还可能互相干扰（第八篇的 hang 一类）。
• CUDA graph 池：vLLM 默认对一组 batch size 捕获 decode 的 CUDA graph，图里引用的中间缓冲在一个专用显存池里，这个池不经过 CuMemAllocator、不能 sleep。verl 的 rollout.cudagraph_capture_sizes 让你只捕获小 batch（如 [1, 2, 4, 8, 16, 32]），或者 enforce_eager: true 干脆不用 CUDA graph——代价是 decode 每步的 kernel 启动开销回来（8B 上每步几毫秒，与 36 毫秒的带宽时间相比是 10% 量级）。
• allocator 的保留段：PyTorch caching allocator 释放张量时并不把显存还给驱动，只是标记为可复用；这些”reserved but not allocated”的段对另一个 allocator是黑洞。所以切换前要 torch.cuda.empty_cache()，verl 包了一层 aggressive_empty_cache(force_sync=True)（gc + synchronize + empty_cache）。

3. gpu_memory_utilization 的含义

vLLM 启动时做一次 profiling：以 gpu_memory_utilization × 总显存 为预算，减去权重、减去一次最大 batch 前向的激活峰值，剩下的全部切成 KV 块。块数从此固定——sleep 是把这些块的物理页摘掉、wake 是挂回同样多的块，不会因为训练器让出了显存而变多。所以共置下 KV 池的大小由这个比例一次定死：0.5 意味着 KV 池最多 40 GB 减权重；训练侧 offload 得越干净，这个比例就能开得越高。

用上一篇的并发公式算它对生成的影响（8B、TP = 1、\(P + \bar L / 2 = 4.6\text{K}\)、\(k_{kv} = 128\) KiB）：

从 0.5 到 0.85，生成时间差 22%——比切换的秒数大一个量级。这是共置形态最该调的一个参数，调的前提是训练侧确实 offload 干净、并且给常驻部分留够余量（一般 8–10 GB）。

三、推理引擎的让渡：vLLM 的 sleep mode

1. 三个 level

vLLM 的 LLM.sleep(level) / wake_up(tags)（v1 引擎里落在 EngineCore.sleep → Worker.sleep）：

• level 0：只暂停调度器（请求照收不处理），显存不动；—— 换权重前”停一下”用
• level 1：权重拷到 CPU（pinned），KV 池丢弃；—— 权重不变、只让 KV 的场景（LoRA、MTP 草稿模型）
• level 2：权重与 KV 都丢弃，只保留 model.named_buffers() 的 CPU 副本；—— 权重马上要被新的覆盖：RL 共置的默认

verl 的 hybrid 模式默认 level 2（_sleep_hybrid()）：新权重反正要从训练器来，level 1 的 CPU 副本是白搬 8 GB、白占 8 GB pinned 内存。只有三种情况退回 level 1——LoRA 只更新 adapter、base 权重不动；MTP 的草稿模型由 vLLM 自己初始化、权重同步覆盖不到它，level 2 丢了就没了；以及 NPU 上不支持 level 2。

2. CuMemAllocator：物理页摘下、虚拟地址不动

普通的 cudaMalloc 把虚拟地址与物理页一次分好，cudaFree 一起收回。vLLM 的 sleep mode 换成 CUDA 的虚拟内存管理 API（cuMemAddressReserve / cuMemCreate / cuMemMap / cuMemUnmap / cuMemRelease），把两步拆开。vllm/device_allocator/cumem.py 用 PyTorch 的 CUDAPluggableAllocator 接进 torch.cuda.MemPool，在 with allocator.use_memory_pool(tag="weights") / tag="kv_cache" 的上下文里分配的所有张量都经它之手，并带一个标签：

# cumem.py（节选）—— sleep：按标签决定备份还是丢弃；两者都 unmap
def sleep(self, offload_tags):
    for ptr, data in self.pointer_to_data.items():
        if data.tag in offload_tags:                     # level 1 的 "weights"
            cpu_backup = torch.empty(size, dtype=uint8, device="cpu", pin_memory=True)
            libcudart.cudaMemcpy(cpu_backup.data_ptr(), ptr, size)
            data.cpu_backup_tensor = cpu_backup
        unmap_and_release(data.handle)                   # 物理页摘下、归还；虚拟地址 ptr 保留
    gc.collect(); torch.cuda.empty_cache()

def wake_up(self, tags):
    for ptr, data in self.pointer_to_data.items():
        if tags is None or data.tag in tags:
            create_and_map(data.handle)                  # 新物理页，映射回同一个 ptr
            if data.cpu_backup_tensor is not None:
                libcudart.cudaMemcpy(ptr, cpu_backup.data_ptr(), size)   # level 1 才有

Worker.sleep(level) 把 level 翻译成标签：level 1 备份 weights、丢 kv_cache；level 2 全丢；wake_up(tags=["weights"]) 与 wake_up(tags=["kv_cache"]) 可以分开唤醒——verl 的权重同步就利用了这一点：先只挂回权重区域、把新权重写进去、再挂回 KV 池，中间这段时间 KV 池的物理内存空着，留给训练器 all-gather 权重用。

%%{init: {"flowchart": {"wrappingWidth": 380}}}%%
%% 图：weights 与 kv_cache 两个标签区域在一步里的映射状态（hybrid 模式，sleep level 2）
flowchart TB
    G["<b>生成中</b><br/>weights：虚拟地址 + 物理页（8 GB）<br/>kv_cache：虚拟地址 + 物理页（~30 GB）"]
    S["<b>sleep(level=2)</b><br/>两个区域都 cuMemUnmap + cuMemRelease<br/>物理页还给驱动，虚拟地址仍预留<br/>训练器拿到这 38 GB 物理显存"]
    W["<b>wake_up(tags=[weights])</b><br/>weights：cuMemCreate + cuMemMap 回原地址，内容未定<br/>kv_cache：仍只有虚拟地址<br/>训练器 all-gather 参数，经 IPC 写进 weights"]
    K["<b>wake_up(tags=[kv_cache])</b><br/>kv_cache：物理页映射回原地址<br/>fp8 KV 清零并重置 scale；reset_prefix_cache"]
    P["持有这些地址的对象不用重建：<br/>CUDA graph · model.parameters() · KV 块表"]
    G --> S --> W --> K -- "下一步生成" --> G
    P -.- W
    P -.- K

因为 ptr 不变，三样东西得以幸存：

• CUDA graph：捕获时记下的是权重与中间缓冲的地址，地址不变图仍有效——这是 vLLM 官方文档里 sleep mode 的第一个设计目标，也是 RL 共置每步不用重新 warm up 几十秒的原因。
• 模型对象：model.parameters() 的每个张量的 data_ptr() 不变，load_weights 直接往里写。
• KV 块表：块的地址不变，调度器不用重建元数据；但块里的内容已经不是原来的——所以要 reset_prefix_cache()。

3. 唤醒后要补的两件事

Worker.wake_up 在 allocator.wake_up 之后做两件事：把 level 2 前保存的 named_buffers()（RoPE 的 cos/sin 表、归一化层的 running 统计一类不属于 state_dict 权重同步的缓冲）拷回去；若 KV 用 fp8，post_kv_cache_wake_up() 把 KV 块清零并把 k_scale / v_scale 重置为 1——新映射的物理页内容是随机的，fp8 的 scale 若留在 0 会让整个 KV 变成零。这两处是”丢弃再重建”的代价：任何不经权重同步的状态都要显式恢复，漏一个就是难查的静默错误（第八篇会回到”权重之外的权重”）。

4. SGLang：torch_memory_saver

SGLang 走同一条路，实现方式不同：torch_memory_saver 是一个 LD_PRELOAD 库，在 with memory_saver.region(tag=...) 上下文内拦截 cudaMalloc，底下同样换成 cuMemCreate + cuMemMap；pause(tag) / resume(tag) 对应 vLLM 的 sleep / wake_up。verl 的 SGLang 路径调用 release_memory_occupation() / resume_memory_occupation(tags=["weights"] | ["kv_cache"])，slime 也是这两个接口（slime/ray/rollout.py 的 offload() / onload(tags)）。机制层面两者一致：按标签分区、物理页摘挂、虚拟地址不动，这是两个推理引擎为 RL 共置提供的核心接口。

四、训练器的让渡：FSDP 与 Megatron 的 offload

1. FSDP 一侧

训练器让出的是 ①–④：参数分片、梯度、优化器状态、参考模型。verl 的 FSDP 引擎（verl/workers/engine/fsdp/）提供三对函数：

• offload_fsdp_model_to_cpu / load_fsdp_model_to_gpu：bf16 参数分片（FSDP1 flat_param；FSDP2 逐 DTensor）
• offload_fsdp_optimizer / load_fsdp_optimizer：fp32 主参数 + Adam 状态
• offload_fsdp_grad / load_fsdp_grad：梯度（一般不搬：optimizer step 后直接释放）

对应配置 actor.fsdp_config.param_offload / optimizer_offload（默认 false；共置的大模型配方几乎都开）。TrainingWorker.to(device, model, optimizer, grad) 是统一入口。v0.9.0 里不是控制器调它，是 train_mode() / eval_mode() 两个上下文进出时自动调（BaseEngineCtx._context_switch）；调用方不传 disable_auto_offload，每次进出都搬：

• old_log_prob：actor 进 eval_mode 时参数 H2D 8 GB，出来时 D2H 8 GB
• ref 前向：ref 是 forward_only，固定走 FSDP2 CPUOffloadPolicy，前向时逐层 H2D，共 8 GB，用完就丢；它的 to() 直接返回，不回搬
• update_actor：进 train_mode 时参数 + 优化器 H2D 57 GB；出来时先 zero_grad 释放梯度，再 D2H 57 GB
• update_weights：参数已经在 CPU。FSDP2 的 get_per_tensor_param() 逐参数 .to(device).full_tensor()，8 GB 分片逐个 H2D；后面那句 to("cpu") 什么也不搬

优化器状态只在 update_actor 期间在卡上——这是把 32B 塞进 8 卡的关键：前向阶段（old_log_prob、ref）只需要 bf16 参数与激活，49 GB 的优化器状态可以留在 CPU。代价是每步多一对 49 GB 的往返（4 秒），以及 pinned 内存。

FSDP2 的 CPUOffloadPolicy 是另一条路：由 FSDP 自己在每次 all-gather 前从 CPU 取、用完就丢，参数常驻 CPU。它把”搬”摊进了训练的每一层，每层的 all-gather 前多一次 H2D；对共置的好处是训练器在 GPU 上的常驻几乎为零，代价是训练 MFU 掉几个点。verl 对它做了兼容（_uses_fsdp2_cpu_offload_policy）：actor 默认走显式 offload，ref 因为只前向，固定用它（CPUOffloadPolicy(pin_memory=True)）。

2. pinned 主机内存

D2H / H2D 要达到 25 GB/s 必须走 pinned 内存，pageable 内存只有几 GB/s。8 卡 32B 共置的一台机器上，pinned 的账：

一台 8 卡 H100 机器通常 1–2 TB 内存，520 GB 的 pinned 分配在 Linux 上需要 ulimit -l 放开、可能触发 NUMA 不均衡（8 张卡挂在两个 socket 上，pinned 缓冲落在哪个 NUMA 节点决定 PCIe 拷贝走不走 UPI）。70B 在 8 卡上 16N = 1.1 TB 的训练状态，pinned 内存本身就放不下——这已经是共置的硬边界（第八章）。

3. Megatron 一侧

Megatron 的训练状态按 TP / PP / DP 切分，分布式优化器（use_distributed_optimizer）把 fp32 主参数与 Adam 状态按 DP 均分。verl 的 Megatron 引擎有对应的 offload_megatron_model_to_cpu / offload_megatron_optimizer / offload_megatron_copy_params，多了一样要搬的东西：分布式优化器持有的参数的 fp32 副本（copy_params）。字节数与 FSDP 相同（都是 16N 均分），搬运方式同为 PCIe。差别在并行布局：Megatron 一侧是 TP × PP 的分片，推理引擎一侧是 TP（可能不同度）× DP 的分片，权重同步的布局映射复杂得多——第四篇。

4. 谁先谁后

一张 80 GB 的卡上，训练 → 生成的顺序不能错：

1. 训练器：optimizer step 完成，退出 train_mode——zero_grad 释放梯度
2. 训练器：同一个退出里 bf16 参数 + 优化器状态 D2H（57 GB）——之后卡上只剩训练器的常驻部分和推理引擎的常驻部分
3. 训练器：aggressive_empty_cache——把 caching allocator 保留的段真正还给驱动
4. 推理引擎：wake_up(tags=[“weights”])——8 GB 的权重区域挂回（空的）
5. 训练器：参数分片逐个 H2D → all-gather → IPC → 推理引擎写入——卡上同时有推理权重 8 GB + 一个参数的完整张量 + 一个 bucket
6. 训练器：to("cpu")——参数早在 CPU 上，什么也不搬；再 empty_cache 一次
7. 推理引擎：wake_up(tags=[“kv_cache”])——最后挂 KV 池，因为它要”尽量大”
8. 推理引擎：reset_prefix_cache

峰值由第 2、4、7 步的先后决定：训练器先把 57 GB 全部让出，再挂权重区域同步，最后才挂 KV 池——KV 池是最大的一块，必须等其他都让出来。这是 v0.9.0 默认配置下 update_actor 加 ActorRolloutRefWorker.update_weights() 的顺序（第五章）。

对照：调用方给 train_mini_batch 传 disable_auto_offload=True 时，train_mode 退出不搬，bf16 参数留在卡上。这时是另一个顺序：优化器状态先手动 D2H（49 GB）→ 挂权重区域 → 直接拿卡上的参数 all-gather 同步（不用 H2D）→ update_weights() 第 3 步的 to("cpu") 才真把 8 GB bf16 参数 D2H → 挂 KV 池。update_weights() 里那句 offload 是为这种用法留的。

五、一次切换的时间线：verl 的 hybrid worker

1. ActorRolloutRefWorker

共置形态在 verl 里的具体形状是一个 worker 类同时持有三样东西——训练器、参考模型、推理引擎，外加一个管权重同步的 checkpoint engine：

sync 模式下（trainer_sync.py）控制器在两个钩子调它：样本取够后 sleep_replicas()，训练完 update_weights()。后者展开就是第四章第 4 节的序列（engine_workers.py，节选）：

async def update_weights(self, global_steps=None, mode="auto"):
    ...
    set_expandable_segments(False)                     # 推理引擎的 cuMem 池与 expandable segments 不共存
    aggressive_empty_cache(force_sync=True)
    # 1. 只挂回权重区域（KV 池仍睡着）
    if self.config.rollout.free_cache_engine:
        await self.rollout.resume(tags=["weights"])
    # 2. 训练器逐参数 all-gather，生成器按需产出完整张量
    per_tensor_param, peft_config = self.actor.engine.get_per_tensor_param(layered_summon=...)
    # 3. 分桶经 CUDA IPC 送进推理引擎（第七章）
    await self.rollout.update_weights(per_tensor_param, peft_config=peft_config, global_steps=global_steps)
    # 4. 训练器的 bf16 参数让出去
    if self.actor.engine.is_param_offload_enabled:
        self.actor.engine.to("cpu", model=True, optimizer=False, grad=False)
    aggressive_empty_cache(force_sync=True)
    # 5. 最后挂回 KV 池
    if self.config.rollout.free_cache_engine:
        await self.rollout.resume(tags=["kv_cache"])
    set_expandable_segments(True)

rollout.resume(tags) 落到 vLLM server 的 wake_up(tags)，后者调 engine.wake_up(tags) 再 reset_prefix_cache(reset_connector=True)；rollout.release() 落到 engine.sleep(level=2)（hybrid）。注意 verl 用的是 engine.wake_up() 而不是 collective_rpc("wake_up")：前者经 DP 协调器广播到全部 EngineCore 进程，后者只到一个 DP 分片内的 TP worker——DP > 1 时用错了会有一半实例的权重没释放、训练反向时 OOM（代码注释里记着这个坑）。

2. 一步的显存时间线

把第二章的字节与本章的顺序合起来，32B、8 卡、一步（sync 模式）：

把这张表按阶段画成一张卡上的显存占用（常驻与激活的高度是示意量级）：

图里有两处要看。一是 update_actor 那一列贴着 80 GB：fp32 主参数加 Adam 状态一项就是 49 GB，所以 32B 在 8 卡上必须开 optimizer_offload，让它只在这一段上卡。二是生成那一列上面空着三十多 GB：训练器已经全部让出，但 KV 池的块数在启动时按 gpu_memory_utilization = 0.5 定死，空出来的显存推理引擎用不上。这一块空白就是第六章第 1 节说的共置真实代价。

搬运合计约 7 秒，其中 4 秒是优化器状态的一对往返。三个可以省的地方：

• 优化器状态不 offload：如果卡放得下（8B、70B 在 64 卡以上），关掉 optimizer_offload，省 4 秒——但 32B 在 8 卡上放不下，没有选择。
• ref 与 actor 共用 base：LoRA 训练时 ref 就是 base 权重，不用单独一份；全参训练下 ref 每步 8 GB 的逐层 H2D 可以改成常驻（若放得下）或与 old_log_prob 前向合并调度。
• 重叠：D2H 可以与前一段计算重叠（用独立 stream、非阻塞拷贝），vLLM 的 wake_up(["weights"]) 可以与优化器状态的 D2H 重叠——verl 目前是顺序执行的，代码里有 TODO。

3. 与步时间的比例

没有一行超过 2%。切换的搬运时间在任何合理配置下都不是瓶颈；把它压到 0 也换不回多少墙钟。共置的代价在下一章。

六、隐性代价

1. KV 池的上限

第二章第 3 节算过：gpu_memory_utilization 从 0.5 到 0.85，8B 场景的生成时间从 745 秒降到 610 秒。这个参数在共置下能开多大，取决于生成阶段卡上还剩什么：

• 常驻（训练器侧）：CUDA context ~0.5 GB · NCCL 缓冲 ~1 GB · caching allocator 残余 ~1–2 GB
• 常驻（推理侧）：CUDA graph 池 1–3 GB · NCCL 缓冲 ~1 GB
• 若 param_offload 关：bf16 参数分片 8 GB（32B/8 卡）或 16 GB（8B/TP=1……即整份）
• 若 optimizer_offload 关：+49 GB —— 此时 KV 池只能开到 0.2–0.3

经验值：训练侧两个 offload 都开、CUDA graph 只捕获小 batch，可以开到 0.8–0.85；只开 param offload 开到 0.6–0.7；都不开就是默认的 0.5 甚至更低。这是共置形态里对墙钟影响最大的一个旋钮，值得在 8 卡上逐档试、看 timing_s/gen 与 OOM 的边界。

2. CUDA graph

vLLM 在启动时按 cudagraph_capture_sizes（默认从 1 到 max_num_seqs 的一组 batch size）为 decode 捕获 CUDA graph，捕获本身几十秒到一两分钟（第一次生成前的 warm up）。有了 CuMemAllocator 的地址保持，sleep / wake 不需要重捕获——这是共置每步只花几秒而不是几分钟的前提。代价是图的中间缓冲池常驻显存（不能 sleep），大 batch 的图占得多；cudagraph_capture_sizes: [1, 2, 4, 8, 16, 32] 一类的配置把池压到 1 GB 以下，代价是并发超过 32 的 decode 步走 eager 模式——而 RL rollout 的并发几乎总是超过 32。所以实际的取舍是：要么给 CUDA graph 池留 2–3 GB、要么接受高并发下 decode 变慢 10% 左右。enforce_eager: true 是后者的极端。

3. prefix cache

权重换了，KV 里的内容全部作废，reset_prefix_cache() 是必须的。它的代价在两处：

• 同一 prompt 的 \(G\) 条回答共享 prompt 的 KV——这是 GRPO 里 prefix cache 的主要用处（上一篇算过它省显存不省 FLOP）。一步之内不受影响，因为 reset 只在步边界。
• 部分 rollout（colocate_async 模式）：被中断的序列在新权重下续接时，已生成的前半段必须重新 prefill——它的 KV 是旧权重算的，不能用。一条 8K 长的序列被中断，续接时先付 8K token 的 prefill（\(2N \times 8\text{K}\)）；一步里有几百条被中断就是几个 PFLOP（peta floating point operations，10^15 次浮点运算），不大，但它是部分 rollout 的固有开销，第五篇再算。

4. 两个 allocator、一块 HBM

训练器用 PyTorch 的 caching allocator，推理引擎的权重与 KV 用 CuMemAllocator（vLLM 自己的激活仍走 caching allocator）。两者在同一块物理显存上交替，问题有三个：

• 保留段：caching allocator 释放的段不还给驱动，wake_up 时 cuMemCreate 拿不到物理页就 OOM——所以 aggressive_empty_cache。
• expandable_segments：PyTorch 的这个选项让 caching allocator 用 cuMemMap 把段按需扩展，对变长序列的训练能显著减少碎片；但它与 CuMemAllocator 在同一进程里各自操作虚拟地址空间，verl 的做法是唤醒推理引擎前关掉、挂回 KV 池后再开（set_expandable_segments(False/True)），让两套 cuMem 操作不在同一时段进行。
• 碎片：训练一步的变长 micro-batch 让 caching allocator 留下大小不一的空洞；一步之后 empty_cache 归零，所以碎片不跨步累积，但一步之内训练的峰值要按”有碎片”估——这是 32B/8 卡这种贴着边跑的配置常见的 OOM 来源：第 1 步能过、第 37 步因为回答变长了不能过。

5. 隐性代价的合计

排在最前面的是 KV 池上限，它不是切换的代价，是”两方常驻的东西加起来”的代价。

七、共置下的权重同步

1. 进程内的路径

共置下训练器与推理引擎在同一台机器、（server 模式下）不同进程但同一张卡。新权重不走网络，走两步：

• 训练器侧 all-gather：FSDP 的参数是分片的 DTensor，get_per_tensor_param() 返回一个生成器，逐参数 .full_tensor()——每个参数在 8 卡间 all-gather 成完整张量。全模型 66 GB 逐个流过每张卡（每卡接收其中 7/8），NVLink 450 GB/s 下约 0.15 秒的通信；生成器是懒的，同一时刻卡上只有一个参数的完整张量。
• CUDA IPC 分桶：完整张量拷进一个固定大小的 bucket（update_weights_bucket_megabytes，默认 512 MB），bucket 的 CUDA IPC handle 经 ZMQ 发给同一张卡上的 vLLM worker 进程；worker 用 rebuild_ipc 打开 handle、按 bucket_meta 里的 offset / shape / dtype 切出每个张量、调 model.load_weights() 写入自己的 TP 分片；回 ACK，训练器填下一个 bucket。

1. 训练器在卡上分配一块 bucket，把它的 CUDA IPC handle 经 ZMQ 发给同卡的 vLLM worker——只在开始时发一次。
2. worker rebuild_ipc(handle) 把这块显存映射进自己的进程，回 ACK。
3. 训练器 for name, param in get_per_tensor_param()：full = param.full_tensor()（NVLink all-gather），bucket[offset:] = full 装桶；下一个参数装不下了，就 synchronize()，经 ZMQ 发 bucket_meta（每个张量的 name / shape / dtype / offset）。
4. worker 按 meta 从映射的 bucket 里切出张量，逐张量 load_weights（取自己的 TP 分片），回 ACK。
5. 训练器收到 ACK，从 offset 0 装下一个 bucket；最后一桶带 is_last=True。

66 GB 经 512 MB 的桶是 130 轮往返，每轮的 IPC 打开与拷贝在同卡内几乎是显存带宽速度，整体约 1 秒。峰值显存是两个 bucket（发送侧一个、接收侧映射一个）加一个完整参数——与模型大小无关，这是分桶的意义。

2. 为什么不直接共享

训练器的分片与推理引擎的分片布局不同（FSDP 按参数的第 0 维切成 8 份，vLLM TP = 8 按注意力头 / 列切），dtype 可能不同（训练 bf16 权重、推理可能 FP8），参数名不同（HF 命名 vs 融合的 QKV）。所以中间必须有一次”完整张量”的形态——all-gather 到完整、再由接收方切出自己的分片。第四篇把这层布局映射与跨机的传输方式展开。

3. 共置下同步的代价

8B 场景全模型 16 GB，all-gather + IPC 约 0.3 秒；32B 约 1 秒；DeepSeek-V3 规格 1.34 TB bf16 要几十秒——这时 all-gather 本身成了瓶颈（每卡接收 1.3 TB 的 31/32），MoE 的专家权重要先按 EP 布局 gather 再重切，第四篇专门讨论。共置的”同步在卡内”这个优点，在几百 B 的 MoE 上打了折扣。

八、适用边界

1. 三条硬边界

• 放得下：训练侧的峰值（bf16 参数 + 优化器状态 + 激活）与推理侧的常驻部分同时在一张卡上时不能超过 80 GB，训练状态可以 offload 但 pinned 内存要放得下。经验公式：\(16N / n_{train} + \text{激活} \le 70\) GB，且 \(16N \le\) 主机内存的一半。70B 要 ≥ 16 卡（每卡 70 GB）实际 ≥ 32 卡；DeepSeek-V3 规格 671B 要 ≥ 256 卡且训练侧 FP8。
• 两侧的并行配置能共存：推理引擎的 TP 度与训练器的 TP × PP 在同一组卡上各建进程组，卡数必须相同；推理侧想用 EP × DP 而训练侧是 TP × PP × EP 时，EP 度不同要在权重同步里重切——可以做，但每一步都做。
• 长尾占比：共置的长尾浪费是全部卡 × 长尾时间（上一篇），\(f > 0.4\) 时应至少换 colocate_async。

2. 一张决策表

3. 何时换分离

共置该换分离的信号，按出现顺序：

1. gpu_memory_utilization 开不到 0.7 且 timing_s/gen 占步时间七成以上——KV 池上限在拖生成；
2. 步与步之间 OOM 时有时无——两个 allocator 的碎片贴着边；
3. 长尾（最长回答 vs 平均回答）占生成时间四成以上——先试 colocate_async，不够再分离；
4. 训练侧想换 Megatron 的 TP × PP × EP、推理侧想用不同的 EP 度——布局每步重切的代价开始接近同步本身；
5. 卡数到 128 以上——共置的长尾浪费按卡数放大，分离的配比与跨机同步开销不随卡数变。

分离不是免费的：它要跨机同步权重（第四篇）、要定配比（上一篇）、要接受至少一步的 off-policy（第五篇）。但它把本篇的全部问题——显存归属、KV 池上限、两个 allocator、pinned 内存——一次性消掉：推理引擎独占自己的卡，KV 池开到 0.9，CUDA graph 常驻无所谓，训练器的 offload 只为放下自己。

九、本文小结

1. 要点回顾

• 共置形态下一张卡在一步里换两次主人：训练时 16N 均分的训练状态 + 参考模型，生成时推理权重副本 + KV 池；32B/8 卡两侧合计 120 GB 以上，必须换手。
• 让渡有三种做法：搬到 CPU（FSDP / Megatron 的 offload、vLLM sleep level 1 的权重）、丢弃再重建（vLLM level 2 的权重、KV 池——新权重反正要来）、不动（CUDA graph 池、NCCL 缓冲、CUDA context）。
• vLLM 的 CuMemAllocator 用 CUDA 虚拟内存 API 把物理页摘下而保留虚拟地址，所以 CUDA graph、模型对象、KV 块表都不用重建；按标签（weights / kv_cache）分区让权重与 KV 池能分开唤醒——verl 先挂权重、同步、offload 训练器参数、最后挂 KV 池。SGLang 的 torch_memory_saver 同一机制。
• 切换的搬运时间在任何合理配置下 < 2%（32B/8 卡约 7 秒，其中 4 秒是优化器状态往返），不是共置的真实代价。
• 真实代价是常驻部分挤掉的 KV 池：gpu_memory_utilization 0.5 让 8B 的生成时间多四分之一；训练侧 offload 干净、CUDA graph 只捕获小 batch，可以开到 0.8–0.85——这是共置最值得调的旋钮。
• 其余隐性代价：prefix cache 每步重置（部分 rollout 要重 prefill 已生成部分）、两个 allocator 的保留段与 expandable_segments 冲突（aggressive_empty_cache、开关切换）、pinned 主机内存（32B/8 卡 520 GB）。
• 共置下权重同步走进程内：逐参数 all-gather → 512 MB bucket → CUDA IPC → 推理 worker 取 TP 分片；峰值两个 bucket，8B 0.3 秒、32B 1 秒、几百 B 的 MoE 几十秒。
• 边界：放得下（\(16N / n \le 70\) GB 且 pinned 内存够）、两侧并行配置能共存、长尾占比 < 0.4；卡数上百后共置的长尾浪费按卡数放大，该分离。

2. 速查表

3. 下一篇

本篇的权重同步是最简单的情形：同一张卡、同一种 dtype、FSDP 到 TP 只要”all-gather 成完整再切”。分离形态下训练器与推理引擎在不同机器，训练侧可能是 Megatron 的 TP × PP × EP、推理侧是 vLLM 的 TP × EP（度不同）、还要在线量化成 FP8；每步 16 GB 到 1.3 TB 的字节要走 NCCL 广播、RDMA 点对点或增量传输。下一篇把”从训练分片到推理分片”这一步的布局映射与传输方式拆开算：

训练器是 Megatron TP=4、PP=2、EP=8 的 671B MoE，推理引擎是 vLLM TP=8、EP=4 的 FP8 副本，两边在不同机器。一次同步要做哪几步映射、传多少字节、走哪条链路、至少几秒？增量同步能省多少？

下一篇：权重同步——从训练分片到推理分片。

实践建议：在 8 卡上用 verl 的 sync 模式跑一个 7B–32B 的 GRPO，打开 VERL_LOGGING_LEVEL=DEBUG 看 log_gpu_memory_usage 在 “Before resume weights / After update_weights / After resume kv_cache” 三个点的输出，与第二章的账对一遍；然后把 gpu_memory_utilization 从 0.5 逐档提到 OOM 的边界，记录每档的 timing_s/gen——这条曲线就是你这个配置下共置的真实代价。

十、自测

1. 32B、8 卡共置，训练 → 生成的切换搬多少字节 / 卡、走什么链路、多久？5 分钟一步时占比多少？

   答案

   优化器状态也 offload 时，update_actor 一结束参数 + 优化器状态 \(14N/8 \approx 57\) GB D2H（梯度直接释放），同步时参数再逐片 H2D 8 GB，走 PCIe pinned 约 25 GB/s，约 2.6 秒，加 NVLink all-gather + IPC 约 1 秒。生成 → 训练还有 old_log_prob 前后各 8 GB、ref 逐层 8 GB、update_actor 前 57 GB。整步约 146 GB、7 秒（4 秒是优化器状态），占 300 秒的 2%。

2. vLLM sleep level 0 / 1 / 2 各做什么？RL 默认用哪个、为什么？

   答案

   0：只停调度、显存不动；1：权重备份到 CPU、KV 丢弃（LoRA、MTP 头这类同步不覆盖的部分需要）；2：权重与 KV 全丢弃（新权重反正每步要同步进来，不必备份）——RL 默认 level 2。

3. CuMemAllocator 为什么要“保留虚拟地址、只摘物理页”？如果直接 del 模型重建会怎样？

   答案

   CUDA graph 里烧的是虚拟地址、KV 块表存的是地址、模型对象持有 tensor——地址不变就都不用重建；重建要重新 profile、重捕获几十张 CUDA graph、重建块表，几十秒到分钟级，且碎片风险。SGLang 的 torch_memory_saver 同理。

4. gpu_memory_utilization 从 0.5 调到 0.85，8B 的生成时间怎么变？什么条件下能开这么高？

   答案

   KV 池扩大近两倍、并发 \(c\) 翻倍、波数减半，生成时间 745 → 610 秒；条件：训练侧状态 offload 干净（生成期间几乎不占显存）、CUDA graph 只捕获小 batch（graph 池小）、NCCL 缓冲与激活峰值留够。

5. 共置形态的边界条件是什么？超出后该怎么办？

   答案

   训练状态放得下：\(16N/n \le 70\) GB 且 pinned 主机内存够（32B / 8 卡要 520 GB）；两侧并行配置能在同一组卡上共存；长尾占比 < 0.4（否则训练器等长尾的浪费太大）；卡数上百后长尾浪费按卡数放大——该转分离 / 异步。

• 上一篇
  RL 后训练基础设施（02）：系统形态——共置、分离与异步
• 下一篇
  RL 后训练基础设施（04）：权重同步——从训练分片到推理分片