如图 6 所示，KVDrive 旨在支持高吞吐量的长上下文 LLM 推理，即使在 GPU 内存紧张的情况下也是如此。当 KV 缓存超过 GPU 容量时，它会被卸载到主机 DRAM 或 SSD，并且系统在预填充阶段在 GPU 内存中构建一个索引。在解码期间，每个新 token 都遵循一个三阶段的工作流程：通过索引识别关键 KV 条目（❶）；将选定的条目从 DRAM 或 SSD 获取到 GPU HBM

如图 6 所示，KVDrive 旨在支持高吞吐量的长上下文 LLM 推理，即使在 GPU 内存紧张的情况下也是如此。当 KV 缓存超过 GPU 容量时，它会被卸载到主机 DRAM 或 SSD，并且系统在预填充阶段在 GPU 内存中构建一个索引。在解码期间，每个新 token 都遵循一个三阶段的工作流程：通过索引识别关键 KV 条目（❶）；将选定的条目从 DRAM 或 SSD 获取到 GPU HBM

难点在于划分warp。涉及三次划分warp。全局内存到共享内存如何划分warp?利用TensorCor 即wmma 接口实现16*16 矩阵相乘时，如何划分warp？将计算结果，写回全局内存，怎样划分warp 实现并行？这三次划分warp 可以独立的划分。原因在于，三次操作间已经__syncthreads()同步过了。代码中， 后两次划分warp方式可以保持一致。笔者学习过程中，对数据搬用划分可以

特性传统单步预测MTP 多步预测监督信号稀疏 (Sparse)密集 (Dense)每个Token的任务预测1步未来预测k+1步未来接收的梯度来自1个目标来自k+1个目标数据利用率基础提升k倍模型学习到的能力主要关注局部、短程依赖强制学习短、中、长程依赖一个生动的类比：传统训练（稀疏信号）: 像一个只盯着脚下走路的人。他每走一步，只看下一步要落在哪里。他能走得很稳，但可能对远方的路线规划能力较弱。M

统一计算，然后拆分: 先用self.mtp一口气计算出所有层级的隐状态[H₀...Hₖ]，然后用chunk拆分开。迭代滚动，匹配目标: 在循环中，通过不断地将labels向前滚动，巧妙地为每个MTP模块（H₁H₂, …）匹配上正确的未来目标（next_2next_3, …）。计算损失，但不立即反传: 为每个MTP模块计算出mtp_loss。梯度“嫁接”: 使用这个“自动求导函数钩子”，将mtp_l

完整构建可能需要几分钟到十几分钟，取决于机器性能。

好的，这张图是一张非常经典且信息量巨大的GPU内存层次结构金字塔。它详细展示了NVIDIA H100 SXM5 GPU内部不同层级内存的特性，解释了为什么理解这个层次结构对于编写高性能代码至关重要。让我们从上到下，逐层详细解读这张图。这个金字塔形象地表达了计算机体系结构中的一个基本权衡：高性能编程的目标就是尽可能将需要频繁访问的数据保留在金字塔的顶层，最大限度地减少对底层慢速内存的访问。解读：这是

统一计算，然后拆分: 先用self.mtp一口气计算出所有层级的隐状态[H₀...Hₖ]，然后用chunk拆分开。迭代滚动，匹配目标: 在循环中，通过不断地将labels向前滚动，巧妙地为每个MTP模块（H₁H₂, …）匹配上正确的未来目标（next_2next_3, …）。计算损失，但不立即反传: 为每个MTP模块计算出mtp_loss。梯度“嫁接”: 使用这个“自动求导函数钩子”，将mtp_l

特性传统单步预测MTP 多步预测监督信号稀疏 (Sparse)密集 (Dense)每个Token的任务预测1步未来预测k+1步未来接收的梯度来自1个目标来自k+1个目标数据利用率基础提升k倍模型学习到的能力主要关注局部、短程依赖强制学习短、中、长程依赖一个生动的类比：传统训练（稀疏信号）: 像一个只盯着脚下走路的人。他每走一步，只看下一步要落在哪里。他能走得很稳，但可能对远方的路线规划能力较弱。M

插件化异步生产者-消费者带背压流式统计持久化的 benchmark 执行框架。yield：把一个值交给调用者，并暂停自己await：把控制权交给事件循环，等异步事件完成后再恢复async for：调用方不断恢复这个异步生成器，直到它结束内部 generator不是独立并发协程任务，它只是内部驱动的一个异步生成器。这段代码本质上是：用构造一个带连接池、超时、trace 的异步 HTTP 客户端外壳，