这句话描述的是现代 LLM 推理引擎如 SGLang、vLLM、TensorRT-LLM 等为了消除 CPU-GPU 间的同步等待卡顿Pipeline Bubbles在异步调度Async Scheduling和内存管理上采取的精细化设计。核心矛盾在于GPU 还在跑当前 stepCPU 就必须开始排发下一个 step 的任务。此时 CPU 并不确定 GPU 的真实执行结果如生成的具体 Token 或投机解码拒绝了几个 Token因此需要一套“先猜测记账、后异步确认”的机制。1. 前一轮 Token 尚未回到 CPU背景在高性能推理引擎中CPU 主线程通过 CUDA Stream 异步向 GPU 发射 Kernel不会在每个 step 结束时调用cudaStreamSynchronize()强行等待 GPU 返回 Sampling 结果否则 PCIe 传输和同步开销会极大地拉低吞吐。当 CPU 正在调度第N1N1N1步时第NNN步生成的 Token 还在 GPU 上或者正在通过 D2HDevice to Host拷贝回 CPU 的途中。2. Placeholder 与 In-flight 状态占位与流水线标记In-flight飞行中/进行中CPU 维护的状态机中将第NNN步占用的资源如正在计算的 KV block、投机生成的候选 Token标记为In-flight。这意味着“这部分资源已投递给 GPU 跑着呢CPU 暂时不能改动它但也还没拿到最终确认结果”。Placeholder占位符在准备第N1N1N1步的计算输入时由于真实的 Token ID 还没传回 CPUCPU 会用一个 Dummy/Placeholder Token 先把 Tensor 形状和逻辑位置“占住”提前构建好下一轮的 Launch Parameters 并打包送入 Stream。3. 乐观记账Optimistic AccountingCPU 在不知道第NNN步真实输出如投机解码接受了 3 个还是 0 个 Token的情况下先按“最好/预期的结果”提前更新显存分配表假设全中乐观地认为投机 Token 全被接受直接在 CPU 的 Page Table 里给它们分配好相应的 KV Cache Block。事后修正Rollback当第NNN步的真实结果终于异步回调回 CPU 后如果发现只接受了 1 个 TokenCPU 再把多分配的逻辑 Block 额度调回Cancel/Rollback。这种“先记账、错再改”的做法换取了 CPU 流水线的无缝重叠Overlap。4. 通过 Fence 延迟释放安全回收安全隐患当投机解码拒绝了某些候选 Token或者某个 Request 结束时CPU 的逻辑记账会把对应的 KV Block 标为“空闲”。但此时 GPU 可能还在 Stream 里读取或写入这块 KV 内存如果 CPU 立即将这块 Block 分配给另一个并发请求会导致严重的数据踩踏Use-After-Free / Data Corruption。FenceCUDA Event机制当 CPU 决定释放某个 KV Block 时不立即将其推入Free Pool而是插入一个Fence如cudaEventRecord到当前的 CUDA Stream 中。这个 Block 被放入一个Pending Free队列并绑定该 Fence。CPU 后续在分配新内存时轮询检查该 Fence 是否已完成cudaEventQuery。唯有 GPU 确认执行完 Fence 之前的全部 Kernels 后CPU 才真正把该 KV Block 安全地回收给新请求使用。