diff --git a/docs/kv_cache_page_size.md b/docs/kv_cache_page_size.md
index 76e99a88d7..6508363edb 100644
--- a/docs/kv_cache_page_size.md
+++ b/docs/kv_cache_page_size.md
@@ -20,6 +20,163 @@ KV 存储和请求表，但资源的申请、缓存和释放以完整物理页
 真实推理索引由模型执行层从请求表选取。
 `ReqManager` 只保留请求 ID、请求表及其原有生命周期职责，不提供 page allocator 接口。
 
+## KV 页分配
+
+`KvCacheAllocator(size, page_size)` 的 `size`、申请数量、释放索引和共享剩余容量均使用 token 单位。
+内部 pinned `mem_state` 每页仅保存一个物理 token 页首索引，形如 `[0, page_size, 2 * page_size, ...]`，
+游标按页计数。申请和释放必须覆盖完整页；释放时按 `page_size` 取页首，重新申请时展开回 token 索引。
+分配结果直接写入 token 级 pinned 返回 tensor：`page_size=1` 使用 `copy_`，
+其他页大小使用广播 `torch.add(..., out=...)`，不再分配中间展开 tensor。
+返回缓冲区按需创建，单次申请超过 `1024 * 1024` 个 token 时独立分配；其他非空申请从
+`1024 * 1024` 个 `int32` 元素的 pinned 缓冲区切片，仅保留未使用的尾部，不回绕覆盖旧结果。
+尾部不足时新建缓冲区，旧切片自行持有原存储；空申请不创建大缓冲区。
+`mem_state` 的容量随 `page_size` 缩小，分配器额外持有的返回缓冲区底层存储最多为 4 MiB，
+不再按 KV 总容量预分配三倍 token 表。在途结果及 PyTorch pinned allocator 缓存的存储另计。
+
+## 请求 Token 索引
+
+`ReqManager` 只持有 GPU `req_to_token_indexs`，不再分配全局 CPU 页表。
+每个 `InferReq` 通过成员 `token_index` 持有一个 `ReqTokenIndex`，实现位于 `server/router/model_infer/infer_req_ext1.py`。
+索引操作直接使用 `req.token_index.append_token_indexes(...)` 等接口，`InferReq` 不再提供同名转发方法。
+基类初始化公共对象和状态，并统一维护上传提交协议及索引流等待；
+各子类实现或复用存储、读写和重置逻辑。
+扩展构造函数接收 `infer_req` 和必填的 `index_stream`，仅提取并保存固定的 `req_idx`、GPU 请求表行视图及索引流，
+避免反向持有请求对象形成循环引用；索引流由外部选择，Helper 内部不查询默认流或捕获构造时的当前流。
+CPU 缓存实现同时记录输入长度 `input_len`，helper 在共享请求初始化后、前缀匹配前构造。
+工厂 `create_req_token_index(infer_req, index_stream)` 从请求参数中选择实现，流参数直接传给子类，
+不再单独传 GPU 行和模式。在线服务和静态 benchmark 在统一创建入口传入 `torch.cuda.current_stream()`，
+Helper 保存该流，后续不随调用点的当前流切换：
+
+| `--req_token_index_mode` | 存储与优点 | 代价 |
+| --- | --- | --- |
+| `pin_mem_cache`（默认） | `PinMemCacheReqTokenIndex` 按请求申请普通 CPU token 缓存，上传时仅为待提交区间创建 pinned 快照；CPU 读取无需 D2H，GPU 更新可批量提交 | 每次上传增加一次 CPU 拷贝和临时 pinned 区间申请，需要维护上传进度 |
+| `pin_mem_page_cache` | `PinMemPageCacheReqTokenIndex` 每个 KV 页仅保存一个 pinned CPU 首 token 索引，常驻索引项数约缩小为原来的 `1 / page_size` | CPU 读取和 GPU 上传时需要展开 token 索引，适用于 `page_size > 1` |
+| `pure_gpu` | `PureGpuReqTokenIndex` 直接维护 GPU 请求表，无常驻 CPU 索引缓存 | CPU 读取需要 D2H 并等待结果，频繁的 CPU 索引操作可能更慢 |
+
+三种模式均保留相同的 GPU 请求表，显存占用基本相同。追加后的所需长度为 `end = token_index_len + mem_indexes.numel()`。
+完整 token 模式的 CPU 缓存首次按 `max(input_len + 2K, end)`
+向上对齐到 1K token 申请，为输出预留空间；后续容量不足时，在本次所需长度 `end` 上多预留 2K token，再向上对齐到 1K，
+容量上限为 GPU 请求表行长度。完整模式使用普通 CPU Tensor 保存每个 token，首次分配及扩容均不申请 pinned 内存。
+页模式的常驻存储仍使用 `torch.empty(..., pin_memory=True, requires_grad=False)`，在首次追加时按 GPU 请求表的最大长度
+一次性申请 `ceil(max_sequence_length / page_size)` 个页首索引，后续追加无需扩容或复制旧缓存。
+页模式继承完整模式的 GPU 上传、释放快照和 reset 流程，覆盖 CPU 索引的存储和展开。
+压缩依赖已有的页内连续索引约束；追加完整页，输入从下标 0 按 page_size 取页首。
+页模式读取支持页内任意 token 区间：内部展开覆盖区间的完整页，再按页内偏移截取精确范围。
+CPU 快照独立于压缩缓存，PD 尾段传输和最终 metadata 无需自行处理页对齐。
+缓冲区在首次非空写入时按需申请，重置时保留输入长度，
+丢弃已申请的常驻缓冲区引用。上传使用独立暂存区，其引用和异步传输依赖由上传路径管理。
+完整模式的上传快照和页模式的展开暂存区统一使用基类 `ReqTokenIndex._alloc_return_pinned_tensor()`，
+按 1M token（4 MiB）申请 int32 slab。函数内部通过函数属性 `return_buffer_state` 惰性保存完整 slab 和下一个空闲 offset，
+两种模式共享同一分配进度，每次只创建返回值这一个切片。返回区间互不重叠，但可持有同一块底层 storage。
+offset 只前进，不回绕；空间不足时换用新 slab 并从零开始，请求 reset 不重置共享 offset。
+超出 slab 容量的申请独立分配，空申请不改变 slab 或 offset。旧 slab 的 storage 在所有切片引用和异步传输依赖结束后
+由 PyTorch 回收。
+纯 GPU 模式对普通 CPU 输入使用 `mem_indexes.pin_memory()` 转为独立 pinned 存储，已有 pinned 输入继续复用。
+三种模式均不借用 KV allocator 的返回缓冲区作为请求常驻缓存。
+
+统一接口为 `append_token_indexes(mem_indexes)`、`get_cpu_token_indexes(start, end)`、
+`get_free_token_indexes(start, end)` 和 `get_gpu_token_indexes(start, end)`，均使用 token 单位，
+追加接口将完整输入接到已记录索引的尾部，非空输入在三种模式下均断言为 CPU Tensor；空输入直接返回。
+CPU 缓存模式直接复制到自身的常驻缓冲区，无需先 pin 输入；完整模式为普通 CPU 内存，页模式为 pinned 内存。
+GPU 模式仅在输入未 pinned 时内部转换，已有 pinned 输入直接复用。
+前缀匹配和索引重建路径直接传入 CPU 索引，不再自行 pin。读取接口均支持页内任意区间；提交给 allocator 的释放区间仍须覆盖完整页。
+CPU getter 返回独立快照，缓存、PD 元数据和异步任务可跨请求释放持有；GPU getter 返回当前请求的
+GPU 表视图，不应跨请求槽位复用持有。Helper 保存外部传入的索引流，GPU 索引读写固定在该流，
+调用方保证 token 区间有效；GPU getter 返回前让调用流等待索引流，没有新增上传时也需承接已有写入的依赖。
+调用流就是索引流时不额外等待；这里只排入 GPU 依赖，不阻塞 CPU，也不改变调用方的当前流。
+模型和 KV 传输流在各自入口统一等待默认流的流程继续保留，Helper 不维护逐请求 CUDA 写事件。
+
+释放专用接口 `get_free_token_indexes` 在 `pin_mem_cache` 模式返回 CPU 快照，在 `pure_gpu`
+模式复用 GPU getter 的范围校验和流等待，再返回 GPU 请求表切片，省掉完整 token 索引的 D2H 快照。
+流等待只保证已有索引写入就绪；重写请求槽位前，仍须在同一流按序消费旧视图，或让索引流等待消费流。
+批量释放先合并视图，再归还请求 ID，暂停恢复也遵循此顺序。
+allocator 仍维护 CPU 空闲页表，因此最终按 `page_size` 取页首并阻塞 D2H，返回时才能重新分配；
+此接口不取消该同步，也不能代替缓存和元数据需要的独立 CPU 快照。
+
+### 上传与源内存生命周期
+
+`token_index_len` 是当前连续记录的 token 索引长度，包含页内预留槽位；每次追加从该位置开始，
+并按输入元素个数增长。CPU 缓存扩容时保留
+`[0, token_index_len)` 的已有映射。它与 CPU 缓存容量、已算好 KV 的长度 `cur_kv_len` 不同；
+完整模式的 CPU token 容量为 `_cpu_buffer.numel()`，页模式为 `_cpu_buffer.numel() * page_size`。
+`gpu_index_submitted_len` 是最新 GPU 索引写入已提交的前缀右端点，后续索引流操作可按序使用
+`[0, gpu_index_submitted_len)`。它表示提交进度，不表示 GPU 传输已完成，也不是累计上传量；
+CPU 缓存模式的追加接口只更新 CPU 索引和 `token_index_len`，不修改 GPU 提交进度。
+需要替换前缀时，先保存仍需保留的索引（包括预留尾部），再调用 `reset()` 并按顺序追加，重建完整连续映射。
+首次前缀匹配时索引为空；暂停恢复前已由统一释放流程重置，两处匹配入口直接追加命中前缀，不重复 `reset()`。
+释放或暂停请求时也由索引扩展的 `reset()` 统一清零两个长度。
+CPU 模式重置时将 `_cpu_buffer` 设为 `None`，下一次追加重新申请，随后从零开始提交 GPU 更新。
+单独的 GPU getter 通过 `get_upload_indexes()` 准备独立 pinned 暂存区：完整模式复制待提交区间，页模式展开压缩索引。
+上传源不借用常驻缓存，常驻缓存的扩容和重置不影响在途 H2D，也不需要在 reset 中同步 CUDA。
+PyTorch `copy_(non_blocking=True)` 会登记 pinned 源的流使用，暂存区在 H2D 完成前不会被回收复用。
+
+`get_upload_indexes()` 不接受区间参数，CPU 模式返回内部计算的
+`[gpu_index_submitted_len, token_index_len)` 的 pinned token 索引。完整模式每次通过基类接口申请恰好覆盖该区间的切片，
+同步复制普通 CPU 缓存中的索引后返回；连续调用返回互不重叠的快照，不覆写上次结果，不包含容量预留区间。
+切片的逻辑长度只覆盖待上传索引，其底层 storage 通常为共享的 4 MiB slab。
+页模式按需展开。两种模式均不改变提交进度；无待上传内容时返回空 tensor，
+纯 GPU 模式返回 `None`。CPU 缓存模式的 GPU getter 校验 `end <= token_index_len` 后按需提交完整待上传区间，
+纯 GPU 模式校验 `end <= gpu_index_submitted_len`；读取均不允许推进超出已有索引的提交进度。
+批量入口不会替 GPU 模式补写提交进度；若缺少待上传数据，由后续拷贝操作报错，提交进度保持不变。
+
+例如 CPU 缓存模式追加 64 个索引后，两个字段分别为 `64/0`；提交到 GPU 后为 `64/64`；
+再追加 32 个索引后为 `96/64`，下次仅需提交 `[64,96)`。需要替换映射时，`reset()` 将两者清为 `0/0`，
+重新追加 96 个完整索引后为 `96/0`，再从头提交 GPU 更新。纯 GPU 模式在追加接口中
+直接提交写表，所以追加后两者相同。
+`_get_classed_reqs` 返回前，backend 将本批索引对象交给 `ReqTokenIndex.submit_pending_uploads`，
+并传入批量拷贝函数；backend 不再读写索引流、GPU 行和提交进度。
+统一入口检查所有请求共用一个索引流，准备源、目标视图列表，在索引流提交拷贝，
+仅在整批提交成功后更新提交进度。单请求 GPU getter 复用该协议，默认通过 PyTorch `copy_` 提交；
+backend 路径使用 `cudaMemcpyBatchAsync`，直接提交索引对象准备的 pinned 快照，不再额外聚合或复制。
+返回前仅在调用流不同于索引流时排入等待，没有新上传时也承接已有写入的依赖。
+backend 不再持有上传流，也不逐条交接多条索引流。
+原生 API 不接受 legacy 默认流，适配由 `PinnedH2DBatchCopier` 内部处理：按设备延迟创建并复用非默认上传流，
+先让上传流等待默认流，提交后再让默认流等待上传流；非默认索引流直接提交，PyTorch 回退直接使用当前流。
+这些操作只排入 GPU 依赖，不阻塞 CPU。
+模型入口原有的推理流等待默认流继续保留；普通增量追加不等待上一批模型执行结束。
+
+原生拷贝使用 `cudaMemcpySrcAccessOrderStream`，需要源内存在实际传输结束前有效且不变。
+仅调用 `index_stream.wait_stream(upload_stream)` 不能保护 CPU 源的生命周期，因为 CPU 会立即返回。
+`cuda_wrapper.py` 将原生接口开关、内部上传流、后台线程池和提交/保活方法统一封装在 `PinnedH2DBatchCopier` 中，
+原有 `cuda_memcpy_pinned_h2d_batch_async` 函数保留为委托入口。原生提交成功后，在同一上传流上记录一个
+`torch.cuda.Event(blocking=True)`，交给进程内复用的单线程 `ThreadPoolExecutor` 等待。
+线程池在首次需要保活时加锁创建，工作线程在首次 submit 时启动；模块导入、空批次和 PyTorch 回退均不创建线程池，
+也不因线程池初始化而初始化 CUDA。
+后台任务持有 `list(src_pinned_tensors)`，独立于调用方可变的列表；源切片会保留完整 pinned storage，
+即使调用方立即丢弃源 Tensor 或清空自己的列表，底层内存仍不会被回收复用。
+后台 `event.synchronize()` 休眠等待并释放 GIL，完成后清空源列表，不向 Future 返回 Tensor；
+主推理线程不等待后台任务，默认索引流由封装承接内部上传流的依赖，调用流按需等待索引流。
+索引扩展不维护逐请求事件，backend 不维护保活列表，也不需要额外的小片段 H2D 或 C++ 扩展。
+该方式只保护引用生命周期，不阻止 CPU 主动改写上传源。完整模式的 pinned 快照和页模式的展开暂存区
+均独立于常驻缓存，reset 和扩容不会改写它们；调用方在传输完成前不能主动修改返回的 pinned Tensor。
+单线程按提交顺序处理；多上传流共用线程池时可能延后回收，不改变 GPU 工作顺序。
+正常进程退出时由线程池等待未完成任务，不取消保活任务。
+
+事件创建、记录、后台任务提交及完成等待属于信任域；异常直接传播，不增加同步收尾、源引用隔离或恢复逻辑。
+原生 batch API 不可用时回退到普通 PyTorch `copy_`，由其内部登记 host allocator，无需创建后台保活任务。
+
+GPU 模式在追加接口中直接向构造时指定的索引流提交 GPU 更新，槽位复用后的新旧写入遵循同一流顺序；
+追加后由调用流等待索引流，D2H 读取也固定在索引流并阻塞到 CPU 快照就绪。
+已 pinned 的输入直接用于 H2D，省去中间快照的申请和 CPU 拷贝；普通 CPU 输入在追加接口内部 pin 后上传。
+KV allocator 返回的 pinned 切片不会被后续申请覆盖；前缀匹配和重建路径直接传入 CPU 索引，由 Helper 按需处理。
+PyTorch `copy_(non_blocking=True)` 内部向 host allocator 登记源 storage 的流使用，
+因此调用方可立即释放输入引用，无需手工记录事件；但在 H2D 完成前不能主动改写源内容。
+
+### 请求结束与跨流使用
+
+暂停或释放时重置请求 Helper；前缀重映射通过重置后重新追加，重新建立 CPU 模式的上传进度。
+PD Decode 传输失败时，等待全部任务返回后将 `cur_kv_len` 回退到原共享前缀长度（无共享前缀则为 0），
+保留 `hold_kv_len` 和索引状态，由后续统一释放回收尾部页面，避免失败 KV 进入缓存。
+CPU KV Cache 在卸载流上下文中调用 GPU getter；Helper 在请求索引流上准备数据，再让当前卸载流等待索引流，调用方无需额外同步。
+纯 GPU 模式直接使用请求表视图，不再执行无谓的 D2H/H2D。
+DP 发布前复用统一上传入口，所有 TP rank 通过默认流上的 NCCL barrier 确认索引可见后再发布。
+diverse 模式会重写旧前缀，在重写入口统一让默认流等待推理流，防止覆盖上一批模型仍在读取的映射。
+KV 卸载完成事件和请求延迟释放机制仍保留，它们负责实际 KV 任务和槽位生命周期。
+模型预热不创建请求对象，直接将 allocator 返回的 token 索引拷入 GPU 请求表。
+静态 benchmark 使用包含 `req_idx`、`req_manager`、`args` 和 `shm_req.input_len` 的轻量请求对象创建索引 Helper，
+不进入真实 `InferReq` 的共享内存生命周期。
+HOLD 请求直接在 GPU 表初始化为预留物理页的循环索引。
+
 ## 生命周期
 
 - Prefill：请求通过调度后，按目标 KV 长度将容量向上对齐，新页一次申请并完整写入请求表。模型执行时再选择
@@ -27,6 +184,8 @@ KV 存储和请求表，但资源的申请、缓存和释放以完整物理页
 - Decode：请求通过调度后，若尾页仍有预留槽位则不再访问 allocator；容量不足时先补齐完整页再执行。
 - Prefix Cache 命中：命中长度天然是整页，`cur_kv_len` 与 `hold_kv_len` 同时初始化为命中长度。
 - 完成/暂停：完整逻辑页可进入 Radix Cache；重复页和未完成尾页按物理页展开后释放。
+  混合注意力在 Prefill 尚未完成时，也可插入已保存大页 checkpoint 且 KV 已计算完成的最长前缀；
+  缓存长度不超过目标缓存边界，其余 KV 统一释放；大页 checkpoint 随前缀移交缓存，重复 checkpoint 在插入时回收。
 - Attention：FA3/FlashInfer 页表每项由物理 token 页首索引除以 `page_size` 得到，KV buffer 视为
   `[page_count, page_size, ...]`。
 
@@ -43,3 +202,142 @@ diverse mode 仍只支持 `page_size=1`，启动阶段会对其他取值明确
   `[HOLD, HOLD+1, ..., HOLD+page_size-1]` 循环填充，而不是重复同一个 token 索引。
 - Radix 子节点用“首个完整 token 页”作为键，避免不同序列仅首 token 相同造成页级分支冲突。
 - 非法的 `page_size < 1` 以及尚未支持的功能组合在模型加载前失败。
+
+## 附录：选择批量异步 H2D 的原因
+
+当前实现采用“CPU 写入 pinned 索引 → CUDA 异步 H2D → GPU kernel 读取 device 请求表”的流程。
+选择显式拷贝的一项工程依据是：根据以往的使用经验，CPU 写入 pinned 内存后，直接让 GPU kernel
+读取这块 host 内存的方案，曾在极端时序下观察到读取值异常。因此，本实现通过 CUDA 内存传输接口
+建立明确的数据交接过程，模型 kernel 统一读取 GPU 上的 `req_to_token_indexs`。
+
+排查上述现象时，需要考虑 CPU 存储缓冲、写合并、cache 一致性以及 CPU/GPU 访问顺序对数据可见性的影响。
+“CPU 写入仍滞留于写缓冲，尚未对设备可见”是历史排查中考虑的可能原因，本文没有将其确认为所有平台上的根因。
+具体的 cache 维护与可见性实现取决于硬件和驱动，正确性应依赖 CUDA API 的访问顺序与完成语义，
+不假定每次异步拷贝都会执行固定的 CPU cache flush，或必须先把数据写回 DRAM 才能被设备正确读取。
+CUDA 也支持正确同步的 mapped pinned memory；其文档明确要求应用协调 host/device 访问，避免读写竞争。
+参见 [CUDA 13.0：Mapped Memory](https://docs.nvidia.com/cuda/archive/13.0.0/cuda-c-programming-guide/index.html#mapped-memory)。
+
+本实现的正确性由以下约束共同保证：
+
+1. CPU 完成源索引写入后才提交 H2D；上传完成前不再改写已提交的源区间。
+2. `cudaMemcpyBatchAsync` 使用 `cudaMemcpySrcAccessOrderStream`，源读取遵循上传流顺序；
+   API 返回只表示已经提交，源内存须保持有效到实际传输结束。
+3. GPU 消费流通过已有的流等待承接上传依赖，随后读取 device 请求表。
+4. 原生 batch 路径由完成事件和后台任务持有 pinned 源引用；PyTorch `copy_` 回退路径
+   由其 host allocator 跟踪异步使用，避免传输完成前回收复用源内存。
+
+源访问顺序的定义参见 [CUDA 13.0 Runtime API：Memory Management](https://docs.nvidia.com/cuda/archive/13.0.0/cuda-runtime-api/group__CUDART__MEMORY.html)。
+批量接口用于将多个请求的索引上传合并为一次 API 提交，并直接使用各请求的 pinned 视图，避免额外的 CPU 聚合拷贝；
+这些正确性约束同样适用于普通的异步 H2D，batch 接口本身不能替代源保活、访问排序和禁止在途改写。
+
+PyTorch 文档介绍了使用 pinned CPU Tensor 配合 `to(..., non_blocking=True)` / `cuda(non_blocking=True)` 进行异步传输的方式，
+并说明 pinned 源在传输完成前被 CPU 改写可能导致 GPU 数据错误。本实现遵循这一显式传输范式，
+`cudaMemcpyBatchAsync` 则是多个索引区间的批量提交选择。
+参见 [PyTorch 2.11：Use pinned memory buffers](https://docs.pytorch.org/docs/2.11/notes/cuda.html#use-pinned-memory-buffers)
+和 [PyTorch：non_blocking 与 pin_memory 使用指南](https://docs.pytorch.org/tutorials/intermediate/pinmem_nonblock.html#other-copy-directions-gpu-cpu-cpu-mps)。
+
+## 附录：分离常驻索引与 pinned 上传缓冲区
+
+### 设计动机与容量估算
+
+请求索引有两种不同的生命周期：完整映射需要保存到请求结束或重建，用于前缀匹配、缓存管理和资源释放；
+H2D 源只需保存到本次上传完成。将完整映射长期保存在 pinned 内存中，会让短期传输需要的锁页资源
+随全部存活请求的索引容量增长。在当前离散 GPU 部署中，大量长期锁页会增加主机内存压力；
+NVIDIA 也明确指出，过量分配 page-locked memory 会影响系统整体性能。
+参见 [NVIDIA CUDA FAQ：Host/Device transfer](https://developer.nvidia.com/cuda/faq)。
+
+面向 V4.1 这类模型的后续部署，需要提前考虑上下文长度、并发请求数和推理进程数同时增长的情况。
+这里按部署容量分析，不假定 V4.1 的具体上下文上限、并发量或实测内存占用。
+对每个推理进程，设已分配缓存的请求数为 `R`，请求 `i` 的完整 token 缓存容量为 `C_i`，
+每个索引为 `int32`，则完整索引的逻辑字节数为 `4 × Σ C_i`。
+如果每个请求各自申请 pinned 缓存，还需要计算底层分配块的容量取整。
+本次核对的 PyTorch 2.11 host caching allocator 将申请字节数向上取整到 2 的幂，
+因此这部分常驻 pinned 块约占 `Σ round_pow2(4 × C_i)`，扩容期间的新旧块和空闲缓存另计。
+这是该版本 pinned allocator 的行为，不是所有 CPU Tensor 的通用分配规则。
+参见 [PyTorch 2.11：host_memory_stats](https://docs.pytorch.org/docs/2.11/generated/torch.cuda.memory.host_memory_stats.html)。
+
+例如，假设输入长 `128 × 1024` token，GPU 请求表行长度允许再预留 2K token，
+当前完整缓存策略会申请 `130 × 1024` 个索引，逻辑大小为 520 KiB；若直接申请 pinned tensor，
+对应块会取整为 1 MiB。每进程 128 个这样的存活请求、同机 8 个进程分别保存这些映射时，
+仅常驻索引就会占用约 1 GiB pinned 块，尚未计入上传、旧块和其他模块。
+多进程占用应按各进程的实际请求集合求和；TP 进程可能保存相同映射，DP 进程则应按各自请求数计算。
+
+默认 `pin_mem_cache` 现将这部分长期存储改为普通 CPU Tensor，上述算例的常驻缓存逻辑大小合计为 520 MiB，
+另加临时 pinned 上传区间。Linux 默认 CPU allocator 按请求大小进行对齐分配，
+没有上述 host caching allocator 的统一 2 的幂取整步骤；系统 malloc 仍可能存在对齐、缓存和碎片开销。
+当前代码自己的 1K 对齐、2K 预留以及扩容复制也仍然存在，普通 CPU 缓存不能当作零成本。
+参见 [PyTorch 2.11：alloc_cpu 实现](https://github.com/pytorch/pytorch/blob/v2.11.0/c10/core/impl/alloc_cpu.cpp)。
+
+### 为什么采用普通 CPU 常驻缓存和共享 pinned slab
+
+完整模式的 `get_upload_indexes()` 只复制 `[gpu_index_submitted_len, token_index_len)`，
+已经提交的历史索引和 CPU 缓存预留空间都不会重复进入上传快照。
+这样保留了 CPU 读取无需 D2H 的能力，同时将 pinned 源的生命周期缩短到本次传输。
+代价是每次上传增加一次同步 CPU 拷贝；前缀重建后需要重新上传完整映射，不能总按少量 decode 增量估算。
+
+如果每次快照都单独调用 `torch.empty(..., pin_memory=True)`，小额上传仍会频繁进入 host allocator，
+且不同大小的快照分别受容量取整影响。`_alloc_return_pinned_tensor()` 为常规申请维护一个
+`1024 × 1024` 个 int32 元素的 slab，让多个请求取得互不重叠的区间，摊薄底层分配和取整开销。
+4 MiB 是当前的小额分配与保留粒度取舍，并非所有模型、并发量和上传大小下的最优值。
+空申请不创建 slab，单次超过 slab 容量时独立申请，避免把长期保留的工作块扩成一个大块。
+
+完整模式将索引复制进这个区间，页模式将压缩页索引直接展开进这个区间；
+两者共享相同的分配与生命周期约定。返回值是一次上传使用的快照，调用方不应把它保存为长期索引缓存。
+多个快照可以共享 storage，但不得共享可写区间。offset 只前进，空间不足时换新 slab，
+请求 reset 不回绕 offset，避免后续 CPU 写入覆盖尚未完成的 H2D 源。
+PyTorch `copy_` 的 allocator 跟踪和原生 batch 路径的事件保活，继续负责传输结束前不回收源存储。
+
+### 职责划分与状态归属
+
+| 组件 | 负责的内容 |
+| --- | --- |
+| `PinMemCacheReqTokenIndex` | 普通 CPU 完整映射的追加、容量管理、快照生成和重置 |
+| `PinMemPageCacheReqTokenIndex` | 页首映射的存储及 token 索引展开，复用公共上传协议 |
+| `ReqTokenIndex` | 上传区间与提交进度、索引流依赖、临时 pinned 返回区间分配 |
+| backend | 选择本批请求并调用统一提交入口 |
+| `PinnedH2DBatchCopier` | CUDA 拷贝提交、原生接口适配及异步源保活 |
+
+这种划分让常驻存储策略可以变化，而 backend 无需了解 CPU 缓存布局或维护上传进度。
+`get_upload_indexes()` 只准备数据，`submit_pending_uploads()` 在整批提交成功后推进进度，
+流完成与引用保活由传输层处理；缓存扩容和 reset 无需等待上传完成。
+页模式继承完整模式时覆盖所有依赖压缩布局的追加、读取和分配操作，复用 GPU getter、释放入口及 reset。
+保留 `PinMemCacheReqTokenIndex` 和 `pin_mem_cache` 名称以兼容现有调用和配置；名称中的 PinMem
+现在对应上传缓冲区，常驻完整缓存的实际类型以实现、类说明和 CLI 帮助为准。
+
+`return_buffer_state = (buffer, offset)` 放在分配函数对象上，由函数内部惰性建立和更新，
+使两个字段的使用范围集中在唯一分配入口。它仍是进程内共享的持久状态，函数返回后不会消失，
+移动字段本身不减少内存；主要收益来自常驻缓存改用普通 CPU 内存及上传区间共享 slab。
+该入口依赖同一推理线程串行申请，未提供多线程分配保护；未来若增加并行申请线程，需重新设计同步或状态归属。
+
+### 内存边界与后续评估
+
+4 MiB 仅表示此分配函数当前持有的一块 slab，不能解释为整个进程的 pinned 内存上限。
+上传缓冲区的实际占用还包括旧 slab 的存活切片、在途传输、超过 slab 的独立申请，以及 PyTorch 缓存的空闲块。
+一个小切片也会保留整块 slab；GPU 上传积压、后台回收延迟或调用方长期持有返回值，都会延长整块存储的存活时间。
+当前没有对同时存活的 slab 数量设置硬上限，也没有上传积压的内存背压。
+KV allocator 的返回缓冲区和其他模块的 pinned 申请应独立计入整机预算，不能只统计这个函数。
+
+`pin_mem_page_cache` 的常驻缓存目前仍使用 pinned 内存，并在首次追加时按 GPU 行长度 `S`
+申请 `ceil(S / page_size)` 个页首索引，而不是按本次实际长度增长。
+其常驻 pinned 块约为 `R × round_pow2(4 × ceil(S / page_size))`，上传暂存区另计。
+`page_size=1` 时没有压缩收益；即使页较大，配置很大的最大序列长度且并发很多时，也仍需评估这部分预分配。
+默认完整模式的改进不能代表页模式已经消除了长期锁页。若该项成为瓶颈，后续可评估普通 CPU 页缓存或按需增长，
+并测量对页索引展开耗时的影响。
+
+本次设计主要降低完整索引的长期 pinned 占用。首次长 prefill 或大规模重建时，普通 CPU 缓存与 pinned 快照会同时存在，
+因此不保证主机总内存峰值下降，也不保证推理吞吐提升。
+GPU `req_to_token_indexs` 仍为 token 级表，约占 `4 × (max_request_num + 1) × S` 字节；
+模型权重、KV Cache 和该 GPU 表的显存都不因这项改动缩小。
+
+评估 V4.1 等模型的实际部署时，应同时观察长 prefill、稳态 decode、前缀重建和上传积压场景，
+按主机汇总各推理进程的 CPU RSS、pinned 块占用及峰值，并测量索引准备、上传和端到端延迟。
+PyTorch 2.11 的 `host_memory_stats()` 中，`allocated_bytes.current` 包含活跃块和缓存块，
+`active_bytes.current` 表示尚未变成可复用空闲块的容量，`num_host_alloc` 用于观察底层分配次数；
+这些是进程级 host allocator 指标，不能直接归因到单个索引类，也不能用 GPU 的 `memory_allocated()` 替代。
+各 bucket 的峰值是汇总近似值，取整与指标定义参见上述 PyTorch 文档。
+
+仓库中的 [CPU/pinned 保活分配基准](../test/benchmark/benchmark_cpu_pinned_alloc_keepalive.py)
+保留所有申请结果，可观察不同大小的直接分配耗时和 pinned 块取整；
+[索引分配基准](../test/benchmark/benchmark_req_token_index_alloc.py) 比较直接申请与 slab 切片路径。
+两者都不测量完整上传流程，前者也不写入普通 CPU Tensor 的全部页面，不能据此直接判断真实工作负载的 RSS 或推理性能。
+现有 [索引测试](../unit_tests/common/test_req_token_index.py) 覆盖快照独立性、跨模式共享、换块及在途 H2D 源的生命周期。
diff --git a/lightllm/common/basemodel/basemodel.py b/lightllm/common/basemodel/basemodel.py
index 5a2d5b77d8..2d48237b2b 100755
--- a/lightllm/common/basemodel/basemodel.py
+++ b/lightllm/common/basemodel/basemodel.py
@@ -6,7 +6,6 @@
 import math
 import json
 import torch
-import torch.nn.functional as F
 import triton
 from typing import final, List
 from tqdm import tqdm
@@ -26,6 +25,8 @@
 from lightllm.common.basemodel.prefill_cuda_graph import PrefillCudaGraph
 from lightllm.common.quantization import Quantcfg
 from lightllm.common.basemodel.triton_kernel.gather_token_id import gather_token, gather_token_prefill_decode_mixed
+from lightllm.common.basemodel.triton_kernel.pad_decode_inputs import pad_decode_inputs
+from lightllm.common.basemodel.triton_kernel.pad_prefill_inputs import pad_prefill_inputs
 from lightllm.utils.log_utils import init_logger
 from lightllm.utils.dist_utils import get_dp_world_size
 from lightllm.utils.profile_max_tokens import profile_mtp_weight_memory
@@ -399,6 +400,57 @@ def _create_inferstate(self, model_input: ModelInput, microbatch_index: int = 0)
 
         return infer_state
 
+    def _prepare_input_ids(self, model_input: ModelInput):
+        """普通与 overlap 入口统一准备真实 token，随后再由状态准备流程 padding。"""
+        if model_input.is_prefill:
+            if self.args.enable_prefill_decode_mixed and model_input.input_ids.shape[0] > 0:
+                gather_token_prefill_decode_mixed(
+                    input_ids=model_input.input_ids,
+                    req_to_next_token_ids=self.req_manager.req_sampling_params_manager.req_to_next_token_ids,
+                    b_req_idx=model_input.b_req_idx,
+                    b_mtp_index=model_input.b_mtp_index,
+                    b_is_decode_req=model_input.b_is_decode_req,
+                    b_prefill_start_loc=model_input.b_prefill_start_loc,
+                )
+        elif model_input.input_ids is None:
+            if model_input.batch_size > 0:
+                model_input.input_ids = gather_token(
+                    req_to_next_token_ids=self.req_manager.req_sampling_params_manager.req_to_next_token_ids,
+                    b_req_idx=model_input.b_req_idx,
+                    b_mtp_index=model_input.b_mtp_index,
+                )
+            else:
+                # 空 DP rank 不启动 gather kernel；padding 随后填入 dummy token。
+                model_input.input_ids = torch.empty((0,), dtype=torch.int64, device=model_input.b_req_idx.device)
+
+    def _create_prefill_inferstate(
+        self,
+        model_input: ModelInput,
+        handle_token_num: int,
+        microbatch_index: int = 0,
+        mem_indexes_ready_event=None,
+    ):
+        """统一 prefill padding 和状态初始化，允许在 KV 索引就绪处记录事件。"""
+        model_input = self._create_padded_prefill_model_input(model_input, handle_token_num)
+        infer_state = self._create_inferstate(model_input, microbatch_index)
+        if mem_indexes_ready_event is not None:
+            mem_indexes_ready_event.record()
+        infer_state.init_some_extra_state(self)
+        infer_state.init_att_state()
+        return infer_state
+
+    def _create_decode_inferstate(
+        self, model_input: ModelInput, batch_size: int, need_capture: bool, microbatch_index: int = 0
+    ):
+        """统一 decode padding、graph 标记和初始化顺序。"""
+        model_input = self._create_padded_decode_model_input(model_input, batch_size)
+        infer_state = self._create_inferstate(model_input, microbatch_index)
+        # Attention backend 在初始化时需要知道是否正在 capture。
+        infer_state.is_cuda_graph = need_capture
+        infer_state.init_some_extra_state(self)
+        infer_state.init_att_state()
+        return infer_state
+
     def _create_padded_decode_model_input(self, model_input: ModelInput, new_batch_size: int):
         if model_input.batch_size == new_batch_size:
             return model_input
@@ -410,27 +462,24 @@ def _create_padded_decode_model_input(self, model_input: ModelInput, new_batch_s
         new_model_input.batch_size = new_batch_size
         new_model_input.total_token_num += padded_batch_size * 2
         new_model_input.max_kv_seq_len = max(2, model_input.max_kv_seq_len)
-        new_model_input.input_ids = F.pad(new_model_input.input_ids, (0, padded_batch_size), mode="constant", value=1)
-        new_model_input.b_req_idx = F.pad(
-            new_model_input.b_req_idx, (0, padded_batch_size), mode="constant", value=self.req_manager.HOLD_REQUEST_ID
-        )
-        new_model_input.b_mtp_index = F.pad(
-            new_model_input.b_mtp_index, (0, padded_batch_size), mode="constant", value=0
-        )
-        new_model_input.b_seq_len = F.pad(new_model_input.b_seq_len, (0, padded_batch_size), mode="constant", value=2)
-        if new_model_input.b_position_delta is not None:
-            new_model_input.b_position_delta = F.pad(
-                new_model_input.b_position_delta, (0, padded_batch_size), mode="constant", value=0
-            )
-        new_model_input.multimodal_params = new_model_input.multimodal_params + [
-            {"images": [], "audios": []} for _ in range(padded_batch_size)
-        ]
-
-        new_model_input.b_shared_seq_len = F.pad(
-            new_model_input.b_shared_seq_len, (0, padded_batch_size), mode="constant", value=0
-        )
-        new_model_input.b_shared_radix_node_id = F.pad(
-            new_model_input.b_shared_radix_node_id, (0, padded_batch_size), mode="constant", value=-1
+        (
+            new_model_input.input_ids,
+            new_model_input.b_req_idx,
+            new_model_input.b_mtp_index,
+            new_model_input.b_seq_len,
+            new_model_input.b_position_delta,
+            new_model_input.b_shared_seq_len,
+            new_model_input.b_shared_radix_node_id,
+        ) = pad_decode_inputs(
+            input_ids=model_input.input_ids,
+            b_req_idx=model_input.b_req_idx,
+            b_mtp_index=model_input.b_mtp_index,
+            b_seq_len=model_input.b_seq_len,
+            b_position_delta=model_input.b_position_delta,
+            b_shared_seq_len=model_input.b_shared_seq_len,
+            b_shared_radix_node_id=model_input.b_shared_radix_node_id,
+            new_batch_size=new_batch_size,
+            hold_request_id=self.req_manager.HOLD_REQUEST_ID,
         )
 
         # 特殊模型，特殊模式的特殊变量的特殊 padding
@@ -440,6 +489,10 @@ def _create_padded_decode_model_input(self, model_input: ModelInput, new_batch_s
                 new_batch_size=new_batch_size,
             )
 
+        new_model_input.multimodal_params = new_model_input.multimodal_params + [
+            {"images": [], "audios": []} for _ in range(padded_batch_size)
+        ]
+
         new_model_input.check_input()
         return new_model_input
 
@@ -458,14 +511,23 @@ def _create_padded_prefill_model_input(self, model_input: ModelInput, new_handle
         new_model_input.max_q_seq_len = max(padded_token_num, model_input.max_q_seq_len)
         new_model_input.max_kv_seq_len = max(padded_token_num, model_input.max_kv_seq_len)
         new_model_input.max_cache_len = max(0, model_input.max_cache_len)
-        new_model_input.input_ids = F.pad(new_model_input.input_ids, (0, padded_token_num), mode="constant", value=1)
-        new_model_input.b_req_idx = F.pad(
-            new_model_input.b_req_idx, (0, 1), mode="constant", value=self.req_manager.HOLD_REQUEST_ID
+        (
+            new_model_input.input_ids,
+            new_model_input.b_req_idx,
+            new_model_input.b_mtp_index,
+            new_model_input.b_seq_len,
+            new_model_input.b_ready_cache_len,
+            new_model_input.b_is_decode_req,
+        ) = pad_prefill_inputs(
+            input_ids=model_input.input_ids,
+            b_req_idx=model_input.b_req_idx,
+            b_mtp_index=model_input.b_mtp_index,
+            b_seq_len=model_input.b_seq_len,
+            b_ready_cache_len=model_input.b_ready_cache_len,
+            b_is_decode_req=model_input.b_is_decode_req,
+            new_handle_token_num=new_handle_token_num,
+            hold_request_id=self.req_manager.HOLD_REQUEST_ID,
         )
-        new_model_input.b_mtp_index = F.pad(new_model_input.b_mtp_index, (0, 1), mode="constant", value=0)
-        new_model_input.b_seq_len = F.pad(new_model_input.b_seq_len, (0, 1), mode="constant", value=padded_token_num)
-        new_model_input.b_ready_cache_len = F.pad(new_model_input.b_ready_cache_len, (0, 1), mode="constant", value=0)
-        new_model_input.b_is_decode_req = F.pad(new_model_input.b_is_decode_req, (0, 1), mode="constant", value=False)
         b_q_seq_len = new_model_input.b_seq_len - new_model_input.b_ready_cache_len
         new_model_input.b_prefill_start_loc = b_q_seq_len.cumsum(dim=0, dtype=torch.int32) - b_q_seq_len
         # 构建新的list, 使用 append 可能会让外面使用的数组引用发生变化，导致错误。
@@ -530,15 +592,7 @@ def _prefill(
         self,
         model_input: ModelInput,
     ):
-        if self.args.enable_prefill_decode_mixed and model_input.input_ids.shape[0] > 0:
-            gather_token_prefill_decode_mixed(
-                input_ids=model_input.input_ids,
-                req_to_next_token_ids=self.req_manager.req_sampling_params_manager.req_to_next_token_ids,
-                b_req_idx=model_input.b_req_idx,
-                b_mtp_index=model_input.b_mtp_index,
-                b_is_decode_req=model_input.b_is_decode_req,
-                b_prefill_start_loc=model_input.b_prefill_start_loc,
-            )
+        self._prepare_input_ids(model_input)
 
         origin_handle_token_num = model_input.input_ids.shape[0]
         origin_batch_size = model_input.batch_size
@@ -554,16 +608,10 @@ def _prefill(
                 handle_token_num=infer_handle_token_num
             )
 
-        model_input = self._create_padded_prefill_model_input(
-            model_input=model_input, new_handle_token_num=infer_handle_token_num
-        )
-
-        infer_state = self._create_inferstate(model_input)
         prefill_mem_indexes_ready_event = torch.cuda.Event()
-        prefill_mem_indexes_ready_event.record()
-
-        infer_state.init_some_extra_state(self)
-        infer_state.init_att_state()
+        infer_state = self._create_prefill_inferstate(
+            model_input, infer_handle_token_num, mem_indexes_ready_event=prefill_mem_indexes_ready_event
+        )
         model_output = self._context_forward(infer_state=infer_state)
 
         model_output = self._create_unpad_prefill_model_output(
@@ -578,21 +626,7 @@ def _decode(
         self,
         model_input: ModelInput,
     ) -> ModelOutput:
-        if model_input.input_ids is None:
-            if model_input.batch_size > 0:
-                model_input.input_ids = gather_token(
-                    req_to_next_token_ids=(self.req_manager.req_sampling_params_manager.req_to_next_token_ids),
-                    b_req_idx=model_input.b_req_idx,
-                    b_mtp_index=model_input.b_mtp_index,
-                )
-            else:
-                # 空 DP rank 不启动 gather kernel，但仍将 input_ids 规范化为
-                # CUDA 空 tensor；后续内部 padding 会为 dummy request 填入 token id 1。
-                model_input.input_ids = torch.empty(
-                    (0,),
-                    dtype=torch.int64,
-                    device=model_input.b_req_idx.device,
-                )
+        self._prepare_input_ids(model_input)
 
         origin_batch_size = model_input.batch_size
         # 空 DP rank 也需要完整的 dummy verify group；TP/SP 切分不能破坏 MTP 分组。
@@ -610,13 +644,7 @@ def _decode(
             infer_batch_size = self.graph.find_closest_graph_batch_size(batch_size=infer_batch_size)
             need_capture = self.graph.need_capture(infer_batch_size)
 
-        model_input = self._create_padded_decode_model_input(model_input=model_input, new_batch_size=infer_batch_size)
-        infer_state = self._create_inferstate(model_input)
-        # attention backend 会根据该标记准备 CUDA Graph capture 专用状态，
-        # 因此必须在 init_att_state 之前完成赋值。
-        infer_state.is_cuda_graph = need_capture
-        infer_state.init_some_extra_state(self)
-        infer_state.init_att_state()
+        infer_state = self._create_decode_inferstate(model_input, infer_batch_size, need_capture)
 
         if use_cuda_graph:
             if need_capture:
@@ -728,15 +756,7 @@ def microbatch_overlap_prefill(self, model_input0: ModelInput, model_input1: Mod
 
         for model_input in (model_input0, model_input1):
             model_input.to_cuda()
-            if self.args.enable_prefill_decode_mixed and model_input.input_ids.shape[0] > 0:
-                gather_token_prefill_decode_mixed(
-                    input_ids=model_input.input_ids,
-                    req_to_next_token_ids=(self.req_manager.req_sampling_params_manager.req_to_next_token_ids),
-                    b_req_idx=model_input.b_req_idx,
-                    b_mtp_index=model_input.b_mtp_index,
-                    b_is_decode_req=model_input.b_is_decode_req,
-                    b_prefill_start_loc=model_input.b_prefill_start_loc,
-                )
+            self._prepare_input_ids(model_input)
         return self._microbatch_overlap_prefill_cuda(model_input0, model_input1)
 
     def _microbatch_overlap_prefill_cuda(self, model_input0: ModelInput, model_input1: ModelInput):
@@ -754,20 +774,8 @@ def _microbatch_overlap_prefill_cuda(self, model_input0: ModelInput, model_input
         origin_batch_size0 = model_input0.batch_size
         origin_batch_size1 = model_input1.batch_size
 
-        model_input0 = self._create_padded_prefill_model_input(
-            model_input=model_input0, new_handle_token_num=infer_handle_token_num0
-        )
-        model_input1 = self._create_padded_prefill_model_input(
-            model_input=model_input1, new_handle_token_num=infer_handle_token_num1
-        )
-
-        infer_state0 = self._create_inferstate(model_input0, 0)
-        infer_state0.init_some_extra_state(self)
-        infer_state0.init_att_state()
-
-        infer_state1 = self._create_inferstate(model_input1, 1)
-        infer_state1.init_some_extra_state(self)
-        infer_state1.init_att_state()
+        infer_state0 = self._create_prefill_inferstate(model_input0, infer_handle_token_num0, 0)
+        infer_state1 = self._create_prefill_inferstate(model_input1, infer_handle_token_num1, 1)
 
         prefill_mem_indexes_ready_event = torch.cuda.Event()
         prefill_mem_indexes_ready_event.record()
@@ -797,19 +805,7 @@ def microbatch_overlap_decode(self, model_input0: ModelInput, model_input1: Mode
 
         for model_input in (model_input0, model_input1):
             model_input.to_cuda()
-            if model_input.input_ids is None:
-                if model_input.batch_size > 0:
-                    model_input.input_ids = gather_token(
-                        req_to_next_token_ids=(self.req_manager.req_sampling_params_manager.req_to_next_token_ids),
-                        b_req_idx=model_input.b_req_idx,
-                        b_mtp_index=model_input.b_mtp_index,
-                    )
-                else:
-                    model_input.input_ids = torch.empty(
-                        (0,),
-                        dtype=torch.int64,
-                        device=model_input.b_req_idx.device,
-                    )
+            self._prepare_input_ids(model_input)
         return self._microbatch_overlap_decode_cuda(model_input0, model_input1)
 
     def _microbatch_overlap_decode_cuda(self, model_input0: ModelInput, model_input1: ModelInput):
@@ -820,21 +816,16 @@ def _microbatch_overlap_decode_cuda(self, model_input0: ModelInput, model_input1
         max_len_in_batch = max(2, model_input0.max_kv_seq_len, model_input1.max_kv_seq_len)
         infer_batch_size = self._align_decode_batch_size(max(1, origin_batch_size0, origin_batch_size1))
 
-        if self.graph is not None and self.graph.can_run(infer_batch_size, max_len_in_batch):
+        use_cuda_graph = self.graph is not None and self.graph.can_run(infer_batch_size, max_len_in_batch)
+        need_capture = False
+        if use_cuda_graph:
             infer_batch_size = self.graph.find_closest_graph_batch_size(infer_batch_size)
             need_capture = self.graph.need_capture(infer_batch_size)
-            padded_model_input0 = self._create_padded_decode_model_input(model_input0, infer_batch_size)
-            padded_model_input1 = self._create_padded_decode_model_input(model_input1, infer_batch_size)
-            infer_state0 = self._create_inferstate(padded_model_input0, 0)
-            infer_state0.is_cuda_graph = need_capture
-            infer_state0.init_some_extra_state(self)
-            infer_state0.init_att_state()
-
-            infer_state1 = self._create_inferstate(padded_model_input1, 1)
-            infer_state1.is_cuda_graph = need_capture
-            infer_state1.init_some_extra_state(self)
-            infer_state1.init_att_state()
 
+        infer_state0 = self._create_decode_inferstate(model_input0, infer_batch_size, need_capture, 0)
+        infer_state1 = self._create_decode_inferstate(model_input1, infer_batch_size, need_capture, 1)
+
+        if use_cuda_graph:
             if need_capture:
                 model_output0, model_output1 = self.graph.capture_decode(
                     self._overlap_tpsp_token_forward,
@@ -847,23 +838,11 @@ def _microbatch_overlap_decode_cuda(self, model_input0: ModelInput, model_input1
                     infer_state1=infer_state1,
                 )
 
-            model_output0 = self._create_unpad_decode_model_output(model_output0, origin_batch_size=origin_batch_size0)
-            model_output1 = self._create_unpad_decode_model_output(model_output1, origin_batch_size=origin_batch_size1)
         else:
-            model_input0 = self._create_padded_decode_model_input(model_input0, infer_batch_size)
-            model_input1 = self._create_padded_decode_model_input(model_input1, infer_batch_size)
-            infer_state0 = self._create_inferstate(model_input0, 0)
-            infer_state0.init_some_extra_state(self)
-            infer_state0.init_att_state()
-
-            infer_state1 = self._create_inferstate(model_input1, 1)
-            infer_state1.init_some_extra_state(self)
-            infer_state1.init_att_state()
-
             model_output0, model_output1 = self._overlap_tpsp_token_forward(infer_state0, infer_state1=infer_state1)
-            model_output0 = self._create_unpad_decode_model_output(model_output0, origin_batch_size=origin_batch_size0)
-            model_output1 = self._create_unpad_decode_model_output(model_output1, origin_batch_size=origin_batch_size1)
 
+        model_output0 = self._create_unpad_decode_model_output(model_output0, origin_batch_size=origin_batch_size0)
+        model_output1 = self._create_unpad_decode_model_output(model_output1, origin_batch_size=origin_batch_size1)
         return model_output0, model_output1
 
     @final
@@ -987,8 +966,9 @@ def _check_max_len_infer(self):
             b_req_idx = torch.tensor([self.req_manager.alloc()], dtype=torch.int32, device="cuda")
             page_size = self.mem_manager.page_size
             alloc_token_num = triton.cdiv(len(dummy_input_ids), page_size) * page_size
-            mem_indexes = self.mem_manager.alloc(alloc_token_num).cuda()
-            self.req_manager.req_to_token_indexs[b_req_idx[0], : len(mem_indexes)] = mem_indexes
+            mem_indexes = self.mem_manager.alloc(alloc_token_num)
+            req_idx = int(b_req_idx[0])
+            self.req_manager.req_to_token_indexs[req_idx, :alloc_token_num].copy_(mem_indexes, non_blocking=True)
             b_seq_len = torch.ones(1, dtype=torch.int32, device="cuda")
             b_seq_len[:] = self.batch_max_tokens
             b_ready_cache_len = torch.zeros(1, dtype=torch.int32, device="cuda")
@@ -1068,8 +1048,9 @@ def _autotune_warmup(self):
                 b_req_idx = torch.tensor([self.req_manager.alloc()], dtype=torch.int32, device="cuda")
                 page_size = self.mem_manager.page_size
                 alloc_token_num = triton.cdiv(len(dummy_input_ids), page_size) * page_size
-                mem_indexes = self.mem_manager.alloc(alloc_token_num).cuda()
-                self.req_manager.req_to_token_indexs[b_req_idx[0], : len(mem_indexes)] = mem_indexes
+                mem_indexes = self.mem_manager.alloc(alloc_token_num)
+                req_idx = int(b_req_idx[0])
+                self.req_manager.req_to_token_indexs[req_idx, :alloc_token_num].copy_(mem_indexes, non_blocking=True)
                 b_seq_len = torch.ones(1, dtype=torch.int32, device="cuda")
                 b_seq_len[:] = input_len
                 b_ready_cache_len = torch.zeros(1, dtype=torch.int32, device="cuda")
diff --git a/lightllm/common/basemodel/triton_kernel/pad_decode_inputs.py b/lightllm/common/basemodel/triton_kernel/pad_decode_inputs.py
new file mode 100644
index 0000000000..69ddef61f8
--- /dev/null
+++ b/lightllm/common/basemodel/triton_kernel/pad_decode_inputs.py
@@ -0,0 +1,127 @@
+from typing import Optional
+
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _pad_decode_inputs_kernel(
+    input_ids,
+    b_req_idx,
+    b_mtp_index,
+    b_seq_len,
+    b_position_delta,
+    b_shared_seq_len,
+    b_shared_radix_node_id,
+    padded_input_ids,
+    padded_b_req_idx,
+    padded_b_mtp_index,
+    padded_b_seq_len,
+    padded_b_position_delta,
+    padded_b_shared_seq_len,
+    padded_b_shared_radix_node_id,
+    batch_size,
+    new_batch_size,
+    hold_request_id,
+    HAS_B_POSITION_DELTA: tl.constexpr,
+    BLOCK_SIZE: tl.constexpr,
+):
+    offsets = tl.program_id(0) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
+    source_mask = offsets < batch_size
+    output_mask = offsets < new_batch_size
+
+    # 保留已有请求，新增 HOLD 请求使用 token id 1、seq_len 2。
+    token_ids = tl.load(input_ids + offsets, mask=source_mask, other=1)
+    tl.store(padded_input_ids + offsets, token_ids, mask=output_mask)
+
+    req_idx = tl.load(b_req_idx + offsets, mask=source_mask, other=hold_request_id)
+    tl.store(padded_b_req_idx + offsets, req_idx, mask=output_mask)
+
+    mtp_index = tl.load(b_mtp_index + offsets, mask=source_mask, other=0)
+    tl.store(padded_b_mtp_index + offsets, mtp_index, mask=output_mask)
+
+    seq_len = tl.load(b_seq_len + offsets, mask=source_mask, other=2)
+    tl.store(padded_b_seq_len + offsets, seq_len, mask=output_mask)
+
+    # 位置偏移按需补齐；HOLD 请求不携带位置偏移和共享缓存。
+    if HAS_B_POSITION_DELTA:
+        position_delta = tl.load(b_position_delta + offsets, mask=source_mask, other=0)
+        tl.store(padded_b_position_delta + offsets, position_delta, mask=output_mask)
+
+    shared_seq_len = tl.load(b_shared_seq_len + offsets, mask=source_mask, other=0)
+    tl.store(padded_b_shared_seq_len + offsets, shared_seq_len, mask=output_mask)
+
+    shared_radix_node_id = tl.load(b_shared_radix_node_id + offsets, mask=source_mask, other=-1)
+    tl.store(padded_b_shared_radix_node_id + offsets, shared_radix_node_id, mask=output_mask)
+
+
+@torch.no_grad()
+def pad_decode_inputs(
+    input_ids: torch.Tensor,
+    b_req_idx: torch.Tensor,
+    b_mtp_index: torch.Tensor,
+    b_seq_len: torch.Tensor,
+    b_position_delta: Optional[torch.Tensor],
+    b_shared_seq_len: torch.Tensor,
+    b_shared_radix_node_id: torch.Tensor,
+    new_batch_size: int,
+    hold_request_id: int,
+):
+    """一次 kernel 调用将 decode Token 和请求元数据补齐到 new_batch_size。
+
+    调用方保证输入为同一 CUDA 设备上的连续一维 Tensor，各字段按请求对齐，
+    且 new_batch_size 大于现有请求数。b_position_delta 为 None 时保持 None。
+    返回顺序与七个 Tensor 输入一致。
+    """
+    batch_size = b_req_idx.shape[0]
+
+    padded_input_ids = torch.empty((new_batch_size,), dtype=input_ids.dtype, device=input_ids.device)
+    padded_b_req_idx = torch.empty((new_batch_size,), dtype=b_req_idx.dtype, device=b_req_idx.device)
+    padded_b_mtp_index = torch.empty((new_batch_size,), dtype=b_mtp_index.dtype, device=b_mtp_index.device)
+    padded_b_seq_len = torch.empty((new_batch_size,), dtype=b_seq_len.dtype, device=b_seq_len.device)
+    padded_b_position_delta = None
+    if b_position_delta is not None:
+        padded_b_position_delta = torch.empty(
+            (new_batch_size,), dtype=b_position_delta.dtype, device=b_position_delta.device
+        )
+    padded_b_shared_seq_len = torch.empty(
+        (new_batch_size,), dtype=b_shared_seq_len.dtype, device=b_shared_seq_len.device
+    )
+    padded_b_shared_radix_node_id = torch.empty(
+        (new_batch_size,), dtype=b_shared_radix_node_id.dtype, device=b_shared_radix_node_id.device
+    )
+
+    block_size = 256
+    grid = (triton.cdiv(new_batch_size, block_size),)
+    _pad_decode_inputs_kernel[grid](
+        input_ids=input_ids,
+        b_req_idx=b_req_idx,
+        b_mtp_index=b_mtp_index,
+        b_seq_len=b_seq_len,
+        b_position_delta=b_position_delta,
+        b_shared_seq_len=b_shared_seq_len,
+        b_shared_radix_node_id=b_shared_radix_node_id,
+        padded_input_ids=padded_input_ids,
+        padded_b_req_idx=padded_b_req_idx,
+        padded_b_mtp_index=padded_b_mtp_index,
+        padded_b_seq_len=padded_b_seq_len,
+        padded_b_position_delta=padded_b_position_delta,
+        padded_b_shared_seq_len=padded_b_shared_seq_len,
+        padded_b_shared_radix_node_id=padded_b_shared_radix_node_id,
+        batch_size=batch_size,
+        new_batch_size=new_batch_size,
+        hold_request_id=hold_request_id,
+        HAS_B_POSITION_DELTA=b_position_delta is not None,
+        BLOCK_SIZE=block_size,
+        num_warps=4,
+    )
+    return (
+        padded_input_ids,
+        padded_b_req_idx,
+        padded_b_mtp_index,
+        padded_b_seq_len,
+        padded_b_position_delta,
+        padded_b_shared_seq_len,
+        padded_b_shared_radix_node_id,
+    )
diff --git a/lightllm/common/basemodel/triton_kernel/pad_prefill_inputs.py b/lightllm/common/basemodel/triton_kernel/pad_prefill_inputs.py
new file mode 100644
index 0000000000..c2828e251d
--- /dev/null
+++ b/lightllm/common/basemodel/triton_kernel/pad_prefill_inputs.py
@@ -0,0 +1,113 @@
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _pad_prefill_inputs_kernel(
+    input_ids,
+    b_req_idx,
+    b_mtp_index,
+    b_seq_len,
+    b_ready_cache_len,
+    b_is_decode_req,
+    padded_input_ids,
+    padded_b_req_idx,
+    padded_b_mtp_index,
+    padded_b_seq_len,
+    padded_b_ready_cache_len,
+    padded_b_is_decode_req,
+    handle_token_num,
+    padded_token_num,
+    batch_size,
+    hold_request_id,
+    BLOCK_SIZE: tl.constexpr,
+):
+    block_start = tl.program_id(0) * BLOCK_SIZE
+    offsets = block_start + tl.arange(0, BLOCK_SIZE)
+
+    # Token 按实际补齐数量扩展，新增位置填入 token id 1。
+    token_ids = tl.load(input_ids + offsets, mask=offsets < handle_token_num, other=1)
+    tl.store(padded_input_ids + offsets, token_ids, mask=offsets < handle_token_num + padded_token_num)
+
+    # 请求元数据只追加一个 HOLD 请求；后续 Token block 不再处理元数据。
+    if block_start < batch_size + 1:
+        source_mask = offsets < batch_size
+        output_mask = offsets < batch_size + 1
+
+        req_idx = tl.load(b_req_idx + offsets, mask=source_mask, other=hold_request_id)
+        tl.store(padded_b_req_idx + offsets, req_idx, mask=output_mask)
+
+        mtp_index = tl.load(b_mtp_index + offsets, mask=source_mask, other=0)
+        tl.store(padded_b_mtp_index + offsets, mtp_index, mask=output_mask)
+
+        seq_len = tl.load(b_seq_len + offsets, mask=source_mask, other=padded_token_num)
+        tl.store(padded_b_seq_len + offsets, seq_len, mask=output_mask)
+
+        ready_cache_len = tl.load(b_ready_cache_len + offsets, mask=source_mask, other=0)
+        tl.store(padded_b_ready_cache_len + offsets, ready_cache_len, mask=output_mask)
+
+        is_decode_req = tl.load(b_is_decode_req + offsets, mask=source_mask, other=False)
+        tl.store(padded_b_is_decode_req + offsets, is_decode_req, mask=output_mask)
+
+
+@torch.no_grad()
+def pad_prefill_inputs(
+    input_ids: torch.Tensor,
+    b_req_idx: torch.Tensor,
+    b_mtp_index: torch.Tensor,
+    b_seq_len: torch.Tensor,
+    b_ready_cache_len: torch.Tensor,
+    b_is_decode_req: torch.Tensor,
+    new_handle_token_num: int,
+    hold_request_id: int,
+):
+    """一次 kernel 调用补齐 prefill Token，并追加一个 HOLD 请求。
+
+    调用方保证输入为同一 CUDA 设备上的连续一维 Tensor，元数据按请求对齐，
+    且 new_handle_token_num 大于现有 Token 数。返回顺序与六个 Tensor 输入一致。
+    """
+    handle_token_num = input_ids.shape[0]
+    padded_token_num = new_handle_token_num - handle_token_num
+    batch_size = b_req_idx.shape[0]
+    new_batch_size = batch_size + 1
+
+    padded_input_ids = torch.empty((new_handle_token_num,), dtype=input_ids.dtype, device=input_ids.device)
+    padded_b_req_idx = torch.empty((new_batch_size,), dtype=b_req_idx.dtype, device=b_req_idx.device)
+    padded_b_mtp_index = torch.empty((new_batch_size,), dtype=b_mtp_index.dtype, device=b_mtp_index.device)
+    padded_b_seq_len = torch.empty((new_batch_size,), dtype=b_seq_len.dtype, device=b_seq_len.device)
+    padded_b_ready_cache_len = torch.empty(
+        (new_batch_size,), dtype=b_ready_cache_len.dtype, device=b_ready_cache_len.device
+    )
+    padded_b_is_decode_req = torch.empty((new_batch_size,), dtype=b_is_decode_req.dtype, device=b_is_decode_req.device)
+
+    block_size = 256
+    grid = (triton.cdiv(max(new_handle_token_num, new_batch_size), block_size),)
+    _pad_prefill_inputs_kernel[grid](
+        input_ids=input_ids,
+        b_req_idx=b_req_idx,
+        b_mtp_index=b_mtp_index,
+        b_seq_len=b_seq_len,
+        b_ready_cache_len=b_ready_cache_len,
+        b_is_decode_req=b_is_decode_req,
+        padded_input_ids=padded_input_ids,
+        padded_b_req_idx=padded_b_req_idx,
+        padded_b_mtp_index=padded_b_mtp_index,
+        padded_b_seq_len=padded_b_seq_len,
+        padded_b_ready_cache_len=padded_b_ready_cache_len,
+        padded_b_is_decode_req=padded_b_is_decode_req,
+        handle_token_num=handle_token_num,
+        padded_token_num=padded_token_num,
+        batch_size=batch_size,
+        hold_request_id=hold_request_id,
+        BLOCK_SIZE=block_size,
+        num_warps=4,
+    )
+    return (
+        padded_input_ids,
+        padded_b_req_idx,
+        padded_b_mtp_index,
+        padded_b_seq_len,
+        padded_b_ready_cache_len,
+        padded_b_is_decode_req,
+    )
diff --git a/lightllm/common/cuda_wrapper.py b/lightllm/common/cuda_wrapper.py
index 32823f1766..5d37e8327f 100644
--- a/lightllm/common/cuda_wrapper.py
+++ b/lightllm/common/cuda_wrapper.py
@@ -16,13 +16,12 @@
 # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 # See the License for the specific language governing permissions and
 # limitations under the License.
-"""This file is a pure Python wrapper for the cudart library.
-It avoids the need to compile a separate shared library, and is
-convenient for use when we just need to call a few functions.
-"""
+"""Python cudart wrapper with asynchronous pinned-source lifetime management."""
 
 import ctypes
+from concurrent.futures import ThreadPoolExecutor
 from dataclasses import dataclass
+from threading import Lock
 from typing import Any, Dict, List, Optional
 
 # this line makes it possible to directly load `libcudart.so` using `ctypes`
@@ -40,6 +39,125 @@
 cudaMemcpyKind = ctypes.c_int
 
 
+class PinnedH2DBatchCopier:
+    """每个推理进程共享批量 H2D 状态，统一管理异步源保活与异常回退。"""
+
+    _native_batch_available = True
+    # 原生 batch 禁止 legacy 默认流，仅在默认流调用时按设备延迟创建上传流。
+    _upload_streams: Dict[torch.device, torch.cuda.Stream] = {}
+    # 只在原生 H2D 成功提交、需要后台保活时创建线程池，导入和 copy_ 回退均不创建。
+    _release_pool: Optional[ThreadPoolExecutor] = None
+    _release_pool_lock = Lock()
+
+    @classmethod
+    def copy_async(cls, dst_gpu_tensors: List[torch.Tensor], src_pinned_tensors: List[torch.Tensor]):
+        """按当前流顺序批量异步 H2D，并保护 pinned 源直到传输结束。
+
+        源由 PyTorch pinned allocator 申请，各对 Tensor 连续且形状和 dtype 一致。
+        默认流的原生 batch 适配封装在内部；调用方只遵循当前流顺序，不阻塞 CPU。
+        上传完成前不能主动改写源，目标须保持有效到传输结束。
+        """
+        if not dst_gpu_tensors:
+            return
+        stream = torch.cuda.current_stream()
+        if cls._native_batch_available and stream.cuda_stream == 0:
+            # 原生 API 不接受 legacy 默认流；上传流先等待当前流，提交后再交回依赖。
+            upload_stream = cls._get_upload_stream(stream.device)
+            upload_stream.wait_stream(stream)
+            with torch.cuda.stream(upload_stream):
+                cls._copy_on_stream(dst_gpu_tensors, src_pinned_tensors, upload_stream)
+            stream.wait_stream(upload_stream)
+        else:
+            cls._copy_on_stream(dst_gpu_tensors, src_pinned_tensors, stream)
+
+    @classmethod
+    def _get_upload_stream(cls, device: torch.device) -> torch.cuda.Stream:
+        if device not in cls._upload_streams:
+            cls._upload_streams[device] = torch.cuda.Stream(device=device)
+        return cls._upload_streams[device]
+
+    @classmethod
+    def _copy_on_stream(
+        cls, dst_gpu_tensors: List[torch.Tensor], src_pinned_tensors: List[torch.Tensor], stream: torch.cuda.Stream
+    ):
+        """在已选定的当前流提交，沿用原生保活和 PyTorch 回退。"""
+        if cls._native_batch_available and cls._try_native_batch_h2d(dst_gpu_tensors, src_pinned_tensors, stream):
+            cls._protect_pinned_sources(src_pinned_tensors, stream)
+            return
+
+        # PyTorch copy_ 已经登记 host allocator，回退路径无需额外保活或登记。
+        for dst, src in zip(dst_gpu_tensors, src_pinned_tensors):
+            dst.copy_(src, non_blocking=True)
+
+    @classmethod
+    def _try_native_batch_h2d(
+        cls, dst_gpu_tensors: List[torch.Tensor], src_pinned_tensors: List[torch.Tensor], stream: torch.cuda.Stream
+    ) -> bool:
+        """尝试原生批量提交；接口准备或提交失败时禁用该路径，让调用方回退到 copy_。"""
+        try:
+            from cuda.bindings import runtime as cuda_runtime
+
+            attrs = cuda_runtime.cudaMemcpyAttributes()
+            # Stream 模式让源读取遵循上传流顺序；提交成功后由完成事件和后台任务保活。
+            attrs.srcAccessOrder = cuda_runtime.cudaMemcpySrcAccessOrder.cudaMemcpySrcAccessOrderStream
+            dst_ptrs = [tensor.data_ptr() for tensor in dst_gpu_tensors]
+            src_ptrs = [tensor.data_ptr() for tensor in src_pinned_tensors]
+            byte_sizes = [tensor.numel() * tensor.element_size() for tensor in src_pinned_tensors]
+            batch_copy = cuda_runtime.cudaMemcpyBatchAsync
+        except Exception as error:
+            cls._native_batch_available = False
+            logger.warning(
+                "Native batch H2D API unavailable; "
+                f"disabling native batch copy and falling back to PyTorch copy_: {error}"
+            )
+            return False
+
+        try:
+            (error,) = batch_copy(dst_ptrs, src_ptrs, byte_sizes, len(byte_sizes), [attrs], [0], 1, stream.cuda_stream)
+            if error != cuda_runtime.cudaError_t.cudaSuccess:
+                raise RuntimeError(f"cudaMemcpyBatchAsync failed: {error.name}")
+        except Exception as error:
+            cls._native_batch_available = False
+            logger.warning(
+                "Native batch H2D submission failed; "
+                f"disabling native batch copy and falling back to PyTorch copy_: {error}"
+            )
+            return False
+        return True
+
+    @classmethod
+    def _protect_pinned_sources(cls, src_pinned_tensors: List[torch.Tensor], stream: torch.cuda.Stream):
+        """每批 raw H2D 在上传流记录一个完成事件，由复用线程池接管源 Tensor 引用。"""
+        # blocking=True 让后台线程休眠等待；Event.synchronize 会释放 GIL，不阻塞主推理线程。
+        # 事件排在 raw batch 之后，覆盖此前所有源读取；调用方的 wait_stream 负责 GPU 数据依赖。
+        event = torch.cuda.Event(blocking=True)
+        event.record(stream)
+        # 必须复制容器：调用方可能立即清空自己的列表，切片本身仍持有整块 pinned storage。
+        cls._get_release_pool().submit(cls._wait_and_release_pinned_sources, event, list(src_pinned_tensors))
+
+    @classmethod
+    def _get_release_pool(cls) -> ThreadPoolExecutor:
+        if cls._release_pool is None:
+            # 首次创建加锁，避免多个提交线程同时建立等待线程池；后续调用直接复用。
+            with cls._release_pool_lock:
+                if cls._release_pool is None:
+                    cls._release_pool = ThreadPoolExecutor(max_workers=1, thread_name_prefix="lightllm-pinned-release")
+        return cls._release_pool
+
+    @classmethod
+    def _wait_and_release_pinned_sources(cls, event: torch.cuda.Event, sources: List[torch.Tensor]):
+        """只管理 CPU 源的生命周期，不修改请求状态，也不向流提交新的 GPU 工作。"""
+        event.synchronize()
+        # 普通工作线程可安全释放 Tensor；不受 CUDA host callback 禁止调用 CUDA API 的限制。
+        # 不把 sources 返回给 Future，避免任务结束后仍被结果对象持有。
+        sources.clear()
+
+
+def cuda_memcpy_pinned_h2d_batch_async(dst_gpu_tensors: List[torch.Tensor], src_pinned_tensors: List[torch.Tensor]):
+    """按当前流顺序批量异步 H2D，默认流适配、源内存保活和回退由封装统一管理。"""
+    PinnedH2DBatchCopier.copy_async(dst_gpu_tensors, src_pinned_tensors)
+
+
 class cudaIpcMemHandle_t(ctypes.Structure):
     _fields_ = [("internal", ctypes.c_byte * 128)]
 
diff --git a/lightllm/common/kv_cache_mem_manager/allocator.py b/lightllm/common/kv_cache_mem_manager/allocator.py
index 1331943b1a..839aa49a33 100644
--- a/lightllm/common/kv_cache_mem_manager/allocator.py
+++ b/lightllm/common/kv_cache_mem_manager/allocator.py
@@ -8,20 +8,13 @@
 
 
 class KvCacheAllocator:
-    def __init__(self, size: int) -> None:
-        self.size = size
-        self.mem_state = torch.arange(
-            0, self.size, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
-        )
-        self.mark_start = 0
-        self.mark_end = self.size
+    """内部按页维护空闲索引，对外的容量和申请、释放参数均以 token 为单位。"""
 
-        self._mem_state_return = torch.arange(
-            0, self.size * 3, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
-        )
-        self._return_start = 0
-
-        self.can_use_mem_size = self.size
+    def __init__(self, size: int, page_size: int) -> None:
+        assert page_size >= 1, "page_size must be positive"
+        self.page_size = page_size
+        self.page_offsets = torch.arange(page_size, dtype=torch.int32, device="cpu", requires_grad=False)
+        self._init_buffers(size)
 
         rank_in_node = get_current_rank_in_node()
         # 用共享内存进行共享，router 模块读取进行精确的调度估计；基础层会统一添加服务前缀以防止实例冲突。
@@ -29,75 +22,87 @@ def __init__(self, size: int) -> None:
         self.shared_can_use_token_num.set_value(self.can_use_mem_size)
         return
 
-    def alloc(self, need_size) -> torch.Tensor:
-        if need_size > self.mark_end - self.mark_start:
+    def alloc(self, need_size: int) -> torch.Tensor:
+        assert need_size >= 0 and need_size % self.page_size == 0, "KV cache allocations must cover complete pages"
+        need_page_num = need_size // self.page_size
+        if need_page_num > self.mark_end - self.mark_start:
             logger.error(f"warn no enough cache need_size {need_size} left_size {self.can_use_mem_size}")
             assert False, "error alloc state"
 
-        start = self.mark_start
-        end = self.mark_start + need_size
-        self.mark_start += need_size
+        page_start = self.mark_start
+        page_end = page_start + need_page_num
 
+        dst = self._alloc_return_tensor(need_size)
+        page_starts = self.mem_state[page_start:page_end]
+        if self.page_size == 1:
+            dst.copy_(page_starts)
+        else:
+            # 将每个 page 的页首索引展开为页内连续的 token 级别索引，直接写入 dst。
+            torch.add(page_starts[:, None], self.page_offsets, out=dst.view(-1, self.page_size))
+        self.mark_start = page_end
         self.can_use_mem_size -= need_size
         self.shared_can_use_token_num.set_value(self.can_use_mem_size)
-
-        # 利用缓冲区返回，避免异步情况下的内存竞争
-        if self._return_start + need_size > self._mem_state_return.shape[0]:
-            self._return_start = 0
-        ans = self._mem_state_return[self._return_start : self._return_start + need_size]
-        ans.copy_(self.mem_state[start:end])
-        self._return_start += need_size
-        return ans
+        return dst
 
     def free(self, free_index: Union[torch.Tensor, List[int]]):
-        """_summary_
-
-        Args:
-            free_index (torch.Tensor): _description_
-        """
-
-        end = self.mark_start
-        start = self.mark_start - len(free_index)
-        assert start >= 0, f"error free state start: {self.mark_start} free len {len(free_index)}"
-
+        """接收按完整页展开的 token 索引，每页仅回收一个页首索引。"""
+        free_token_num = len(free_index)
+        assert free_token_num % self.page_size == 0, "KV cache releases must cover complete pages"
+        free_page_num = free_token_num // self.page_size
+        page_end = self.mark_start
+        page_start = page_end - free_page_num
+        assert page_start >= 0, f"error free state start: {self.mark_start} free len {len(free_index)}"
+
+        page_starts = free_index[:: self.page_size]
         if isinstance(free_index, list):
-            self.mem_state.numpy()[start:end] = free_index
+            self.mem_state.numpy()[page_start:page_end] = page_starts
         else:
-            # 从 gpu 到 cpu 的拷贝操作是流内阻塞操作
-            self.mem_state[start:end] = free_index
-
-        self.mark_start -= len(free_index)
+            # GPU 输入在默认流按页取首索引后阻塞 D2H；返回时 CPU 空闲页表已可重新分配。
+            self.mem_state[page_start:page_end].copy_(page_starts)
 
-        self.can_use_mem_size += len(free_index)
+        self.mark_start = page_start
+        self.can_use_mem_size += free_token_num
         self.shared_can_use_token_num.set_value(self.can_use_mem_size)
 
-        if self.can_use_mem_size == len(self.mem_state):
+        if self.can_use_mem_size == self.size:
             logger.debug(f"freed all gpu mem size {self.can_use_mem_size}")
         return
 
     def free_all(self):
-        self.mem_state.numpy()[:] = list(range(0, len(self.mem_state)))
+        self.mem_state = torch.arange(0, self.size, self.page_size, out=self.mem_state, requires_grad=False)
         self.mark_start = 0
-        self.mark_end = len(self.mem_state)
-        self.can_use_mem_size = len(self.mem_state)
+        self.mark_end = self.mem_state.numel()
+        self.can_use_mem_size = self.size
         self.shared_can_use_token_num.set_value(self.can_use_mem_size)
         return
 
     def resize(self, new_size: int) -> None:
-        """
-        just for test code
-        """
-        self.size = new_size
+        """仅供测试：按新的 token 容量重新初始化空闲页表。"""
+        self._init_buffers(new_size)
+        self.shared_can_use_token_num.set_value(self.can_use_mem_size)
+
+    def _init_buffers(self, size: int) -> None:
+        assert size >= 0 and size % self.page_size == 0, "KV cache size must cover complete pages"
+        self.size = size
+        # 每项保存一个物理页的首 token 索引；[mark_start, mark_end) 是空闲页区间。
         self.mem_state = torch.arange(
-            0, self.size, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
+            0, size, self.page_size, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
         )
         self.mark_start = 0
-        self.mark_end = self.size
-
-        self._mem_state_return = torch.arange(
-            0, self.size * 3, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
-        )
-        self._return_start = 0
-
-        self.can_use_mem_size = self.size
-        self.shared_can_use_token_num.set_value(self.can_use_mem_size)
+        self.mark_end = self.mem_state.numel()
+        self._return_buffer = None
+        self.can_use_mem_size = size
+
+    def _alloc_return_tensor(self, need_size: int) -> torch.Tensor:
+        """返回后续申请不会覆写的 pinned 区间，仅保留 slab 尚未使用的尾部。"""
+        buffer_token_num = 1024 * 1024
+        if need_size == 0 or need_size > buffer_token_num:
+            return torch.empty(need_size, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True)
+        if self._return_buffer is None or self._return_buffer.numel() < need_size:
+            self._return_buffer = torch.empty(
+                buffer_token_num, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
+            )
+        # 只保留未使用的尾部，已返回的切片不会被后续分配覆盖。
+        dst = self._return_buffer[:need_size]
+        self._return_buffer = self._return_buffer[need_size:]
+        return dst
diff --git a/lightllm/common/kv_cache_mem_manager/mem_manager.py b/lightllm/common/kv_cache_mem_manager/mem_manager.py
index 38c870ce4e..1f7dd600e9 100755
--- a/lightllm/common/kv_cache_mem_manager/mem_manager.py
+++ b/lightllm/common/kv_cache_mem_manager/mem_manager.py
@@ -47,7 +47,7 @@ def __init__(self, size, dtype, head_num, head_dim, layer_num, always_copy=False
         self.page_size = get_env_start_args().page_size
         self.size = self.size // self.page_size * self.page_size
 
-        self.allocator = KvCacheAllocator(self.size)
+        self.allocator = KvCacheAllocator(self.size, page_size=self.page_size)
 
         self._init_buffers(
             self.size,
diff --git a/lightllm/common/req_manager/base.py b/lightllm/common/req_manager/base.py
index cfe5d716ee..2b13c70107 100644
--- a/lightllm/common/req_manager/base.py
+++ b/lightllm/common/req_manager/base.py
@@ -1,3 +1,4 @@
+from collections import deque
 from typing import List
 
 import torch
@@ -10,39 +11,27 @@
 logger = init_logger("lightllm.common.req_manager")
 
 
-class _ReqNode:
-    def __init__(self, index):
-        self.index = index
-        self.next: "_ReqNode" = None
-
-
-class _ReqLinkedList:
+class _ReqIndexPool:
     def __init__(self, max_request_num):
-        self.nodes = [_ReqNode(i) for i in range(max_request_num)]
+        self._free_indexes = deque(range(max_request_num))
         self.marks = [0 for _ in range(max_request_num)]
-        self.root_node = _ReqNode(-1)
-        for i in range(0, max_request_num - 1):
-            self.nodes[i].next = self.nodes[i + 1]
-        self.root_node.next = self.nodes[0]
         self.can_alloc_size = max_request_num
         return
 
     def alloc(self):
-        if self.root_node.next is None:
+        if not self._free_indexes:
             logger.warning("alloc req index fail")
             return None
-        get_node = self.root_node.next
-        self.root_node.next = self.root_node.next.next
-        assert self.marks[get_node.index] == 0
-        self.marks[get_node.index] = 1
+        index = self._free_indexes.popleft()
+        assert self.marks[index] == 0
+        self.marks[index] = 1
         self.can_alloc_size -= 1
-        return get_node.index
+        return index
 
     def free(self, index):
         assert self.marks[index] == 1
-        node = self.nodes[index]
-        node.next = self.root_node.next
-        self.root_node.next = node
+        # 保持原有头插行为，优先复用最近释放的请求索引。
+        self._free_indexes.appendleft(index)
         self.marks[index] = 0
         self.can_alloc_size += 1
         return
@@ -59,11 +48,12 @@ def __init__(self, max_request_num, max_sequence_length, mem_manager: MemoryMana
         # 的这个请求管理 id， 主要是为了兼容 DP 运行模式下，让各个 DP 能 padding 到 DP 中最大
         # 的那个batch size 进行运行，所有 padding 的请求都会使用预留的这个请求管理 id 进行处理
         # 这样让 DP 的实现更为简化一些。
-        self.req_list = _ReqLinkedList(max_request_num)
+        self.req_list = _ReqIndexPool(max_request_num)
         page_size = get_env_start_args().page_size
-        max_sequence_length = (max_sequence_length + page_size - 1) // page_size * page_size
+        self.page_size = page_size
+        aligned_sequence_length = (max_sequence_length + page_size - 1) // page_size * page_size
         self.req_to_token_indexs = torch.zeros(
-            (max_request_num + 1, max_sequence_length), dtype=torch.int32, device="cuda"
+            (max_request_num + 1, aligned_sequence_length), dtype=torch.int32, device="cuda"
         )
         self.req_sampling_params_manager = ReqSamplingParamsManager(max_request_num)
         self.max_request_num = max_request_num
@@ -86,13 +76,13 @@ def init_hold_request_indexs(self):
         # HOLD_REQUEST_ID 对应的请求行供 DP padding、overlap microbatch 等占位请求使用。将该行
         # 按 page_size 划分后，每一页都映射到 mem_manager 额外保留的同一个物理页；这样占位请求
         # 无论访问哪一个逻辑位置，都会落到合法且不会参与正常分配的 KV cache 地址上。
+        assert self.mem_manager.page_size == self.page_size
         hold_row = self.req_to_token_indexs[self.HOLD_REQUEST_ID]
-        hold_page = torch.tensor(
-            self.mem_manager.HOLD_TOKEN_MEMINDEXES,
-            dtype=hold_row.dtype,
-            device=hold_row.device,
+        token_offsets = torch.arange(
+            hold_row.numel(), dtype=hold_row.dtype, device=hold_row.device, requires_grad=False
         )
-        hold_row.view(-1, self.mem_manager.page_size).copy_(hold_page)
+        hold_page_start = self.mem_manager.HOLD_TOKEN_MEMINDEXES[0]
+        hold_row.copy_(token_offsets.remainder(self.page_size).add_(hold_page_start))
         return
 
     def alloc(self):
@@ -117,5 +107,5 @@ def free_token(self, free_token_index):
         return
 
     def free_all(self):
-        self.req_list = _ReqLinkedList(self.max_request_num)
+        self.req_list = _ReqIndexPool(self.max_request_num)
         return
diff --git a/lightllm/common/req_manager/req_sampling_params.py b/lightllm/common/req_manager/req_sampling_params.py
index 87271b6145..476569f36d 100644
--- a/lightllm/common/req_manager/req_sampling_params.py
+++ b/lightllm/common/req_manager/req_sampling_params.py
@@ -81,15 +81,23 @@ def init_req_sampling_params(self, req: "InferReq"):
         else:
             self.req_to_out_token_id_counter[req.req_idx].fill_(0)
             if req.sampling_param.shm_param.input_penalty and req.need_out_token_id_statistics:
-                prompt_ids = g_pin_mem_manager.gen_from_list(
-                    key="prompt_ids_for_penalty",
-                    data=req.shm_req.get_prompt_ids_numpy(),
+                # 每次独立申请 pinned 内存，避免缓冲池循环复用时覆盖尚未完成的异步上传源。
+                # 上传源释放后由 PyTorch pinned allocator 跟踪拷贝完成再复用；上传和计数
+                # 在同一流上按序执行，后续采样由推理入口的流依赖保证，无需在此阻塞 CPU。
+                prompt_ids_numpy = req.shm_req.get_prompt_ids_numpy()
+                prompt_ids_cpu = torch.empty(
+                    prompt_ids_numpy.shape,
                     dtype=torch.int32,
-                ).cuda(non_blocking=True)
+                    device="cpu",
+                    pin_memory=True,
+                )
+                prompt_ids_cpu.numpy()[:] = prompt_ids_numpy
+                # 当前流的 GPU 暂存区覆盖最大请求长度，仅上传实际 prompt 部分。
+                prompt_ids = self._get_prompt_ids_gpu_buffer()[: prompt_ids_cpu.numel()]
+                prompt_ids.copy_(prompt_ids_cpu, non_blocking=True)
                 token_id_counter(
                     prompt_ids=prompt_ids, out_token_id_counter=self.req_to_out_token_id_counter[req.req_idx]
                 )
-                torch.cuda.current_stream().synchronize()
 
         return
 
@@ -149,3 +157,18 @@ def gen_cpu_out_token_counter_sampling_params(self, req_objs: List["InferReq"]):
             p_token_counts_tensor.cuda(non_blocking=True),
             p_cumsum_seq_len_tensor.cuda(non_blocking=True),
         )
+
+    def _get_prompt_ids_gpu_buffer(self) -> torch.Tensor:
+        """获取当前 CUDA 流的 int32 暂存区，容量为最大请求长度加 64。"""
+        buffers = getattr(self, "_prompt_ids_gpu_buffers", None)
+        if buffers is None:
+            buffers = {}
+            self._prompt_ids_gpu_buffers = buffers
+        stream = torch.cuda.current_stream()
+        buffer = buffers.get(stream)
+        if buffer is None:
+            buffer = torch.empty(get_env_start_args().max_req_total_len + 64, dtype=torch.int32, device=stream.device)
+            buffers[stream] = buffer
+        # 每个流最多缓存一个 GPU buffer；使用者须在当前流上提交拷贝和读取，
+        # 下一次覆盖会排在此前读取之后，不同流使用独立 buffer，避免互相覆盖。
+        return buffer
diff --git a/lightllm/server/api_cli.py b/lightllm/server/api_cli.py
index d4c6504cce..b4984088f3 100644
--- a/lightllm/server/api_cli.py
+++ b/lightllm/server/api_cli.py
@@ -497,6 +497,19 @@ def add_cli_args(parser: argparse.ArgumentParser) -> argparse.ArgumentParser:
         help="""KV cache page size in tokens. Values greater than 1 make each request reserve
         page-aligned contiguous KV slots and make paged attention/cache reuse operate on full pages.""",
     )
+    parser.add_argument(
+        "--req_token_index_mode",
+        type=str,
+        choices=["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"],
+        default="pin_mem_cache",
+        help="""Storage mode for request token indexes. pin_mem_cache (default): allocates a pageable CPU cache
+        per request and copies pending indexes to a fresh pinned slab slice for each upload, allowing CPU reads without
+        D2H transfers and batched GPU uploads. pin_mem_page_cache: stores one pinned CPU index per KV page,
+        reducing persistent index memory but expanding indexes for token reads and GPU uploads; page_size=1
+        provides no compression. pure_gpu: maintains the GPU request table directly without
+        a persistent CPU index cache; CPU reads require D2H transfers and may incur synchronization waits.
+        All modes retain the same GPU request table and have approximately the same GPU memory usage.""",
+    )
     parser.add_argument(
         "--vit_att_backend",
         type=str,
diff --git a/lightllm/server/core/objs/start_args_type.py b/lightllm/server/core/objs/start_args_type.py
index ae84204fec..5af19daa02 100644
--- a/lightllm/server/core/objs/start_args_type.py
+++ b/lightllm/server/core/objs/start_args_type.py
@@ -176,6 +176,9 @@ class StartArgs:
         default_factory=lambda: ["auto"], metadata={"choices": ["auto", "triton", "fa3", "flashinfer"]}
     )
     page_size: int = field(default=1)
+    req_token_index_mode: str = field(
+        default="pin_mem_cache", metadata={"choices": ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"]}
+    )
     vit_att_backend: List[str] = field(
         default_factory=lambda: ["auto"], metadata={"choices": ["auto", "triton", "fa3", "sdpa", "xformers"]}
     )
diff --git a/lightllm/server/router/model_infer/infer_batch.py b/lightllm/server/router/model_infer/infer_batch.py
index ebc2035877..c447f79f67 100644
--- a/lightllm/server/router/model_infer/infer_batch.py
+++ b/lightllm/server/router/model_infer/infer_batch.py
@@ -19,12 +19,12 @@
 from lightllm.utils.log_utils import init_logger
 from lightllm.server.req_id_generator import convert_sub_id_to_group_id
 from lightllm.server.multimodal_params import MultimodalParams
-from lightllm.utils.custom_kernel_utis import custom_cat
 from lightllm.utils.envs_utils import get_env_start_args
 from lightllm.server.pd_io_struct import PDDecodeNodeInfo
 from lightllm.server.embed_cache.embed_cache_client import CpuEmbedCacheClient
 from lightllm.server.multi_level_kv_cache import CachePlacementController, CacheTier
 from lightllm.server.router.model_infer.infer_req_ext import FinalTokenMetadataExt, PromptSelectedLogprobsExt
+from lightllm.server.router.model_infer.infer_req_ext1 import ReqTokenIndex, create_req_token_index
 
 if TYPE_CHECKING:
     from lightllm.server.router.model_infer.mode_backend.base_backend import ModeBackend
@@ -129,7 +129,7 @@ def add_reqs(self, requests: List[Tuple[int, int, Any, int]], init_prefix_cache:
 
     def free_a_req_mem(self, free_token_index: List, req: "InferReq"):
         if self.radix_cache is None:
-            free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][0 : req.hold_kv_len])
+            free_token_index.append(req.token_index.get_free_token_indexes(0, req.hold_kv_len))
         elif CacheTier.GPU not in req.cache_tiers:
             self._free_req_mem_without_radix_insert(free_token_index=free_token_index, req=req)
         else:
@@ -140,12 +140,14 @@ def free_a_req_mem(self, free_token_index: List, req: "InferReq"):
                 assert len(req.hybrid_len_to_big_page_id) == 0
         req.cur_kv_len = 0
         req.hold_kv_len = 0
+        # 待释放列表已持有 CPU 快照或 GPU 行视图；reset 清理扩展状态，不改写 GPU 表。
+        req.token_index.reset()
         req.shm_req.shm_cur_kv_len = req.cur_kv_len
         return
 
     def _free_req_mem_without_radix_insert(self, free_token_index: List, req: "InferReq"):
         shared_kv_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
-        free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][shared_kv_len : req.hold_kv_len])
+        free_token_index.append(req.token_index.get_free_token_indexes(shared_kv_len, req.hold_kv_len))
 
         if self.is_hybrid_att_model:
             # 释放请求尾部尚未移交给 radix cache 的 hybrid attention 小页状态。
@@ -166,32 +168,23 @@ def _free_req_mem_without_radix_insert(self, free_token_index: List, req: "Infer
     def _full_att_free_req(self, free_token_index: List, req: "InferReq"):
         page_size = self.args.page_size
         cache_kv_len = req.cur_kv_len // page_size * page_size
-        # radix cache 只保存完整页，因此不足一页的有效 KV 和页内预留空间都不会参与插入。
-        # 先释放截断后的 [cache_kv_len, hold_kv_len) 区间，再处理可缓存前缀的插入和去重。
-        free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][cache_kv_len : req.hold_kv_len])
-        input_token_ids = req.get_input_token_ids()
-        key = torch.tensor(input_token_ids[0:cache_kv_len], dtype=torch.int64, device="cpu")
-        # .cpu() 是 流内阻塞操作
-        value = self.req_manager.req_to_token_indexs[req.req_idx][:cache_kv_len].detach().cpu()
-
+        # 只缓存完整页，不足一页的有效 KV 和预留尾部均归还 allocator。
+        # 待释放索引可以借用 GPU 行；缓存值须独立于请求槽位及其后续复用。
+        free_token_index.append(req.token_index.get_free_token_indexes(cache_kv_len, req.hold_kv_len))
+        key = torch.tensor(req.get_input_token_ids()[:cache_kv_len], dtype=torch.int64, device="cpu")
+        value = req.token_index.get_cpu_token_indexes(0, cache_kv_len)
         prefix_len, _ = self.radix_cache.insert(key, value)
-        old_prefix_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
-        free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][old_prefix_len:prefix_len])
-        if req.shared_kv_node is not None:
-            assert req.shared_kv_node.node_prefix_total_len <= prefix_len
-            self.radix_cache.dec_node_ref_counter(req.shared_kv_node)
-            req.shared_kv_node = None
-        return
+        self._release_duplicate_prefix(free_token_index, req, prefix_len)
 
     def _hybrid_att_free_req(self, free_token_index: List, req: "InferReq"):
         assert g_infer_context.is_hybrid_att_model is True
         args = get_env_start_args()
         hash_page_size = args.linear_att_hash_page_size
         big_page_num = args.linear_att_page_block_num
+        big_page_token_num = hash_page_size * big_page_num
         shared_kv_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
-        tail_big_page_token_num = (
-            req.hybrid_cache_len // (hash_page_size * big_page_num) * (hash_page_size * big_page_num)
-        )
+        completed_kv_len = min(req.cur_kv_len, req.hybrid_cache_len)
+        completed_big_page_token_num = completed_kv_len // big_page_token_num * big_page_token_num
         page_num = req.hybrid_cache_len // hash_page_size
         assert req.hybrid_cache_len >= shared_kv_len
         if req.tail_small_page_buffer_id is not None:
@@ -203,83 +196,58 @@ def _hybrid_att_free_req(self, free_token_index: List, req: "InferReq"):
         if req.hybrid_cache_len <= req.cur_kv_len and req.tail_small_page_buffer_id is not None:
             # 只有小页可以有 tail_small_page_buffer_id，然后进行小页插入。
             assert page_num % big_page_num != 0
-            free_token_index.append(
-                self.req_manager.req_to_token_indexs[req.req_idx][req.hybrid_cache_len : req.hold_kv_len]
-            )
-            req.cur_kv_len = req.hybrid_cache_len
-            input_token_ids = req.get_input_token_ids()
-            key = torch.tensor(input_token_ids[0 : req.cur_kv_len], dtype=torch.int64, device="cpu")
-            value = self.req_manager.req_to_token_indexs[req.req_idx][: req.cur_kv_len].detach().cpu()
-            block_hashs = req.shm_req.hybrid_token_hash_list.get_all()[:page_num]
-            state_idxs = [None for _ in range(page_num)]
-            state_idxs[-1] = req.tail_small_page_buffer_id
-            req.tail_small_page_buffer_id = None
-            prefix_len, _ = self.radix_cache.insert(
-                key,
-                value,
-                block_hashs=block_hashs,
-                block_state_idxs=state_idxs,
-                len_to_big_page_id=req.hybrid_len_to_big_page_id,
-            )
-            old_prefix_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
-            free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][old_prefix_len:prefix_len])
-            if req.shared_kv_node is not None:
-                assert req.shared_kv_node.node_prefix_total_len <= prefix_len
-                self.radix_cache.dec_node_ref_counter(req.shared_kv_node)
-                req.shared_kv_node = None
-            return
-
-        if shared_kv_len < tail_big_page_token_num <= req.cur_kv_len:
-            free_token_index.append(
-                self.req_manager.req_to_token_indexs[req.req_idx][tail_big_page_token_num : req.hold_kv_len]
-            )
-            req.cur_kv_len = tail_big_page_token_num
-
+            cache_kv_len = req.hybrid_cache_len
+        elif completed_big_page_token_num > shared_kv_len:
+            # Prefill 在大页边界保存 checkpoint，因此按已完成长度对齐即可。
             assert req.tail_small_page_buffer_id is None
-            input_token_ids = req.get_input_token_ids()
-            key = torch.tensor(input_token_ids[0 : req.cur_kv_len], dtype=torch.int64, device="cpu")
-            value = self.req_manager.req_to_token_indexs[req.req_idx][: req.cur_kv_len].detach().cpu()
-            cur_page_num = tail_big_page_token_num // hash_page_size
-            assert tail_big_page_token_num % hash_page_size == 0
-            block_hashs = req.shm_req.hybrid_token_hash_list.get_all()[:cur_page_num]
-            state_idxs = [None for _ in range(cur_page_num)]
-            prefix_len, _ = self.radix_cache.insert(
-                key,
-                value,
-                block_hashs=block_hashs,
-                block_state_idxs=state_idxs,
-                len_to_big_page_id=req.hybrid_len_to_big_page_id,
-            )
-            old_prefix_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
-            free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][old_prefix_len:prefix_len])
-            if req.shared_kv_node is not None:
-                assert req.shared_kv_node.node_prefix_total_len <= prefix_len
-                self.radix_cache.dec_node_ref_counter(req.shared_kv_node)
-                req.shared_kv_node = None
-            return
-
-        if shared_kv_len <= req.cur_kv_len:
-            free_token_index.append(self.req_manager.req_to_token_indexs[req.req_idx][shared_kv_len : req.hold_kv_len])
-            # 该分支不会把 prefill 阶段累积的 big page id 插入 radix cache（典型为 pause/abort
-            # 在 prefill 跨过 big page 边界后、到达末尾前触发），需在此显式释放，避免泄漏。
-
-            # 释放本请求 prefill 阶段在 big page 边界上申请、但尚未插入 radix cache 的 big page
-            # state buffer。仅当请求未走 insert 分支(小页/大页插入)就被释放时才会有残留，典型场景：
-            # big page 模式下请求在 prefill 跨过 big page 边界后、到达末尾前被 pause / abort。
-            # 若不释放，会泄漏 big page state slot，并触发 free_a_req_mem 中 dict 为空的断言。
-            if req.hybrid_len_to_big_page_id:
-                self.radix_cache.big_page_buffers.free_state_cache(list(req.hybrid_len_to_big_page_id.values()))
-                req.hybrid_len_to_big_page_id.clear()
-
-            req.cur_kv_len = shared_kv_len
+            assert len(req.hybrid_len_to_big_page_id) > 0
+            cache_kv_len = completed_big_page_token_num
+        elif shared_kv_len <= req.cur_kv_len:
+            # 没有新 checkpoint，区分共享小页、共享大页和空前缀。
             assert req.tail_small_page_buffer_id is None
-            if req.shared_kv_node is not None:
-                assert req.shared_kv_node.node_prefix_total_len == req.cur_kv_len
-                self.radix_cache.dec_node_ref_counter(req.shared_kv_node)
-                req.shared_kv_node = None
-            return
+            assert len(req.hybrid_len_to_big_page_id) == 0
+            if req.shared_kv_node is not None and not req.shared_kv_node.is_big_page_node():
+                # 大页关闭时，前缀匹配会直接保留共享小页。请求只持有节点引用，
+                # 没有可移交给 insert 的新小页 checkpoint；节点引用只保护 KV，
+                # 其 CPU page buffer 仍可能被独立淘汰，因此不能依赖它仍然存在。
+                # 即使 buffer 仍在，也不能借用其 id 重新插入，否则重复插入会将
+                # 缓存自身持有的槽位当作重复 checkpoint 释放。
+                # 这里只回收请求私有的尾部 KV 并解除引用，共享 KV 和 checkpoint
+                # 继续由 radix cache 管理，避免改变原有的小页插入协议。
+                self._free_req_mem_without_radix_insert(free_token_index, req)
+                return
+            # 共享大页可在没有新 checkpoint 时重复插入；空前缀也可直接插入。
+            cache_kv_len = shared_kv_len
+        else:
+            assert False, f"error state: cur_kv_len: {req.cur_kv_len}"
+
+        free_token_index.append(req.token_index.get_free_token_indexes(cache_kv_len, req.hold_kv_len))
+        req.cur_kv_len = cache_kv_len
+        key = torch.tensor(req.get_input_token_ids()[:cache_kv_len], dtype=torch.int64, device="cpu")
+        value = req.token_index.get_cpu_token_indexes(0, cache_kv_len)
+        page_num = cache_kv_len // hash_page_size
+        state_idxs = [None] * page_num
+        if req.tail_small_page_buffer_id is not None:
+            state_idxs[-1] = req.tail_small_page_buffer_id
+            req.tail_small_page_buffer_id = None
+        prefix_len, _ = self.radix_cache.insert(
+            key,
+            value,
+            block_hashs=req.shm_req.hybrid_token_hash_list.get_all()[:page_num],
+            block_state_idxs=state_idxs,
+            len_to_big_page_id=req.hybrid_len_to_big_page_id,
+        )
+        assert len(req.hybrid_len_to_big_page_id) == 0
+        self._release_duplicate_prefix(free_token_index, req, prefix_len)
 
-        assert False, f"error state: cur_kv_len: {req.cur_kv_len}"
+    def _release_duplicate_prefix(self, free_token_index: List, req: "InferReq", prefix_len: int):
+        """回收插入时发现的重复 KV 页，并解除请求对共享节点的引用。"""
+        old_prefix_len = 0 if req.shared_kv_node is None else req.shared_kv_node.node_prefix_total_len
+        free_token_index.append(req.token_index.get_free_token_indexes(old_prefix_len, prefix_len))
+        if req.shared_kv_node is not None:
+            assert req.shared_kv_node.node_prefix_total_len <= prefix_len
+            self.radix_cache.dec_node_ref_counter(req.shared_kv_node)
+            req.shared_kv_node = None
 
     def _save_promptcache_kvbuffer(self):
         """
@@ -323,7 +291,8 @@ def _filter(self, finished_request_ids: List[int], modify_shm_finish_state: bool
             self.shm_req_manager.put_back_req_obj(req.shm_req)
 
         if free_token_index:
-            free_token_index = custom_cat(free_token_index)
+            # 同一服务的索引模式一致；pure_gpu 在默认流合并 GPU 视图，随后才归还请求槽位。
+            free_token_index = torch.cat(free_token_index)
         self.req_manager.free(free_req_index, free_token_index)
 
         finished_req_ids_set = set(finished_request_ids)
@@ -363,7 +332,8 @@ def pause_reqs(self, pause_reqs: List["InferReq"], is_master_in_dp: bool):
                     logger.debug(f"infer paused req id {req.req_id}")
 
             if len(free_token_index) != 0:
-                free_token_index = custom_cat(free_token_index)
+                # 恢复请求会重写请求表，必须先在默认流消费暂停时借用的 GPU 视图。
+                free_token_index = torch.cat(free_token_index)
                 self.req_manager.free_token(free_token_index)
         return self
 
@@ -457,7 +427,8 @@ def __init__(
         vocab_size: int,
     ) -> None:
         self.shm_param = shm_req.sample_params
-        self.disable_prompt_cache = self.shm_param.disable_prompt_cache
+        # Cached KV skips prompt logits, so logprobs requests must recompute the full prompt.
+        self.disable_prompt_cache = self.shm_param.disable_prompt_cache or self.shm_param.prompt_logprobs >= 0
         if self.shm_param.top_k == -1:
             self.shm_param.top_k = vocab_size
 
@@ -537,6 +508,7 @@ def __init__(
         self.args = get_env_start_args()
         self.req_id = req_id
         self.req_idx = req_idx
+        self.req_manager = g_infer_context.req_manager
         self.shm_index = shm_index
         self.multimodal_params = multimodal_params
         self.vocab_size = vocab_size
@@ -594,6 +566,7 @@ def __init__(
             self.get_chuncked_input_token_ids = self.get_chuncked_input_token_ids_for_hybrid_att
 
         self._init_all_state()
+        self.token_index: ReqTokenIndex = create_req_token_index(self, index_stream=torch.cuda.current_stream())
 
         self.generator = None
         if self.sampling_param.shm_param.seed != -1:
@@ -650,19 +623,13 @@ def _match_radix_cache(self):
             key = key[0 : len(key) - 1]  # 最后一个不需要，因为需要一个额外的token，让其在prefill的时候输出下一个token的值
             share_node, kv_len, value_tensor = g_infer_context.radix_cache.match_prefix(key, update_refs=True)
             if share_node is not None:
-                self.shared_kv_node = share_node
-                ready_cache_len = share_node.node_prefix_total_len
-                # 从 cpu 到 gpu 是流内阻塞操作
-                g_infer_context.req_manager.req_to_token_indexs[self.req_idx, 0:ready_cache_len] = value_tensor
-                self.cur_kv_len = int(ready_cache_len)  # 序列化问题, 该对象可能为numpy.int64，用 int(*)转换
-                self.hold_kv_len = self.cur_kv_len
-                assert self.hold_kv_len % self.args.page_size == 0
-                self.shm_req.prompt_cache_len = self.cur_kv_len  # 记录 prompt cache 的命中长度
+                self._set_cache_prefix(share_node, value_tensor)
 
         self.shm_req.shm_cur_kv_len = self.cur_kv_len
         return
 
     def _hybrid_match_radix_cache(self):
+        """匹配完整 hash 页，恢复命中的 KV 和 checkpoint，再发布请求长度。"""
         assert (
             g_infer_context.is_hybrid_att_model is True
         ), "current _hybrid_match_radix_cache only supports hybrid attention models, to do..."
@@ -675,103 +642,15 @@ def _hybrid_match_radix_cache(self):
         match_block_num = match_tokens // hash_page_size
         block_hashs = block_hashs[:match_block_num]
         assert len(block_hashs) == self.shm_req.hybrid_token_hash_list.size
-        big_page_token_num = hash_page_size * self.args.linear_att_page_block_num
-        big_page_is_disable = big_page_token_num > self.args.max_req_total_len
         if enable_prompt_cache and match_tokens > 1 and len(block_hashs) > 0 and self.cur_kv_len == 0:
             input_token_ids = self.shm_req.shm_prompt_ids.arr[0 : self.get_cur_total_len()]
             key = torch.tensor(input_token_ids[0:match_tokens], dtype=torch.int64, device="cpu")
             assert len(key) == len(block_hashs) * hash_page_size
-            share_node, kv_len, value_tensor = g_infer_context.radix_cache.match_prefix(
+            shared_node, _, token_indexes = g_infer_context.radix_cache.match_prefix(
                 key, block_hashs=block_hashs, update_refs=True
             )
-            if share_node is not None:
-                assert self.tail_small_page_buffer_id is None
-                if share_node.is_big_page_node():
-                    # 大页匹配
-                    self.shared_kv_node = share_node
-                    ready_cache_len = share_node.node_prefix_total_len
-                    # 从 cpu 到 gpu 是流内阻塞操作
-                    g_infer_context.req_manager.req_to_token_indexs[self.req_idx, 0:ready_cache_len] = value_tensor
-                    self.cur_kv_len = int(ready_cache_len)  # 序列化问题, 该对象可能为numpy.int64，用 int(*)转换
-                    self.shm_req.prompt_cache_len = self.cur_kv_len  # 记录 prompt cache 的命中长度
-                    assert self.tail_small_page_buffer_id is None
-                    # 恢复 hybrid checkpoint
-                    g_infer_context.req_manager.restore_big_page_state(
-                        big_page_buffer_idx=share_node.big_page_buffer_idx, req=self
-                    )
-                else:
-                    # 小页匹配
-                    if big_page_is_disable:
-                        # 如果 大页本质是被禁用的，可以直接使用小页的匹配结果
-                        self.shared_kv_node = share_node
-                        ready_cache_len = share_node.node_prefix_total_len
-                        # 从 cpu 到 gpu 是流内阻塞操作
-                        g_infer_context.req_manager.req_to_token_indexs[self.req_idx, 0:ready_cache_len] = value_tensor
-                        self.cur_kv_len = int(ready_cache_len)  # 序列化问题, 该对象可能为numpy.int64，用 int(*)转换
-                        self.shm_req.prompt_cache_len = self.cur_kv_len  # 记录 prompt cache 的命中长度
-                        assert self.tail_small_page_buffer_id is None
-                        # 恢复 hybrid checkpoint
-                        g_infer_context.req_manager.restore_small_page_state(
-                            req=self,
-                        )
-                    else:
-                        # 如果 大页本质是被启用的，则需要使用小页的匹配结果, 将小页的kv 复制到的新申请的kv位置，同时释放
-                        # 对应的小页对应的节点，递归找到对应最近的大叶节点进行返回,然后赋值到req.shared_node 对象上
-                        shared_kv_len = share_node.node_prefix_total_len
-                        cur_big_page_tokens = (shared_kv_len // big_page_token_num) * big_page_token_num
-                        need_tokens = shared_kv_len - cur_big_page_tokens
-                        radix_cache = g_infer_context.radix_cache
-                        if g_infer_context.get_can_alloc_token_num() > need_tokens:
-                            # 有充足的token 容量时
-                            assert need_tokens % self.args.page_size == 0
-                            radix_cache.free_radix_cache_to_get_enough_token(need_token_num=need_tokens)
-                            tail_mems = radix_cache.mem_manager.alloc(need_size=need_tokens)
-                            g_infer_context.req_manager.req_to_token_indexs[
-                                self.req_idx, 0:cur_big_page_tokens
-                            ] = value_tensor[0:cur_big_page_tokens]
-                            g_infer_context.req_manager.req_to_token_indexs[
-                                self.req_idx, cur_big_page_tokens:shared_kv_len
-                            ] = tail_mems
-
-                            # 将 对应的 value_tensors 中的 kv 数据 拷贝到 tail_mems 中对应的数据去
-                            radix_cache.mem_manager.operator.copy_mem_to_mem(
-                                value_tensor[cur_big_page_tokens:shared_kv_len], tail_mems
-                            )
-                            # 尾部 KV 换到新 mem 后，同步拷贝已捕获的 top-k prompt logprobs。
-                            self.prompt_selected_logprobs.copy_capture_slots_if_needed(
-                                source_indexes=value_tensor[cur_big_page_tokens:shared_kv_len],
-                                destination_indexes=tail_mems,
-                            )
-
-                            self.shared_kv_node = share_node  # 只是为了保证 restore_small_page_state 正确调用
-                            g_infer_context.req_manager.restore_small_page_state(
-                                req=self,
-                            )
-                            self.shared_kv_node = None
-
-                            big_page_shared_node = radix_cache.deref_to_first_big_page_node(node=share_node)
-                            self.shared_kv_node = big_page_shared_node
-                            self.cur_kv_len = int(shared_kv_len)  # 序列化问题, 该对象可能为numpy.int64，用 int(*)转换
-                            self.shm_req.prompt_cache_len = self.cur_kv_len  # 记录 prompt cache 的命中长度
-                        else:
-                            # 没有充足的token 容量时， 直接找到最接近的大页，进行大页恢复
-                            share_node = radix_cache.deref_to_first_big_page_node(node=share_node)
-                            if share_node is not None:
-                                assert share_node.is_big_page_node()
-                                # 大页匹配
-                                self.shared_kv_node = share_node
-                                ready_cache_len = share_node.node_prefix_total_len
-                                # 从 cpu 到 gpu 是流内阻塞操作
-                                g_infer_context.req_manager.req_to_token_indexs[
-                                    self.req_idx, 0:ready_cache_len
-                                ] = value_tensor[0:ready_cache_len]
-                                self.cur_kv_len = int(ready_cache_len)  # 序列化问题, 该对象可能为numpy.int64，用 int(*)转换
-                                self.shm_req.prompt_cache_len = self.cur_kv_len  # 记录 prompt cache 的命中长度
-                                assert self.tail_small_page_buffer_id is None
-                                # 恢复 hybrid checkpoint
-                                g_infer_context.req_manager.restore_big_page_state(
-                                    big_page_buffer_idx=share_node.big_page_buffer_idx, req=self
-                                )
+            if shared_node is not None:
+                self._restore_hybrid_cache_match(shared_node, token_indexes)
 
         self.shm_req.shm_cur_kv_len = self.cur_kv_len
 
@@ -782,10 +661,105 @@ def _hybrid_match_radix_cache(self):
         assert self.hold_kv_len % self.args.page_size == 0
 
         if self.cur_kv_len == 0:
-            # 说明没有任何命中
             g_infer_context.req_manager.init_hybrid_attention_state(req=self)
         return
 
+    def _restore_hybrid_cache_match(self, shared_node: HybridAttPagedTreeNode, token_indexes: torch.Tensor):
+        """直接复用共享前缀，或将小页尾部转为私有 KV；容量不足时回退到大页。"""
+        # 新请求或已完成暂停释放的请求，其索引和私有 checkpoint 状态为空。
+        assert self.tail_small_page_buffer_id is None
+        big_page_token_num = self.args.linear_att_hash_page_size * self.args.linear_att_page_block_num
+        big_pages_disabled = big_page_token_num > self.args.max_req_total_len
+        if shared_node.is_big_page_node() or big_pages_disabled:
+            self._restore_hybrid_shared_prefix(shared_node, token_indexes)
+            return
+
+        radix_cache = g_infer_context.radix_cache
+        shared_kv_len = shared_node.node_prefix_total_len
+        big_page_prefix_len = shared_kv_len // big_page_token_num * big_page_token_num
+        private_tail_len = shared_kv_len - big_page_prefix_len
+
+        # 保持原有容量条件：只有可用量严格大于尾部需求时才重建小页尾部。
+        if g_infer_context.get_can_alloc_token_num() <= private_tail_len:
+            big_page_node = radix_cache.deref_to_first_big_page_node(node=shared_node)
+            if big_page_node is not None:
+                assert big_page_node.is_big_page_node()
+                self._restore_hybrid_shared_prefix(big_page_node, token_indexes[: big_page_node.node_prefix_total_len])
+            return
+
+        assert private_tail_len % self.args.page_size == 0
+        radix_cache.free_radix_cache_to_get_enough_token(need_token_num=private_tail_len)
+        private_tail_indexes = radix_cache.mem_manager.alloc(need_size=private_tail_len)
+        cached_tail_indexes = token_indexes[big_page_prefix_len:shared_kv_len]
+        radix_cache.mem_manager.operator.copy_mem_to_mem(cached_tail_indexes, private_tail_indexes)
+        self.prompt_selected_logprobs.copy_capture_slots_if_needed(
+            source_indexes=cached_tail_indexes, destination_indexes=private_tail_indexes
+        )
+
+        # 小页恢复接口从请求读取节点；必须先恢复 checkpoint，再解除小页引用。
+        self.shared_kv_node = shared_node
+        g_infer_context.req_manager.restore_small_page_state(req=self)
+        self.shared_kv_node = None
+        big_page_node = radix_cache.deref_to_first_big_page_node(node=shared_node)
+        self._set_cache_prefix(
+            shared_node=big_page_node,
+            shared_prefix_indexes=token_indexes[:big_page_prefix_len],
+            private_tail_indexes=private_tail_indexes,
+        )
+
+    def _restore_hybrid_shared_prefix(self, shared_node: HybridAttPagedTreeNode, token_indexes: torch.Tensor):
+        """登记直接共享的 KV 前缀，并恢复同一节点的 checkpoint。"""
+        self._set_cache_prefix(shared_node, token_indexes)
+        assert self.tail_small_page_buffer_id is None
+        if shared_node.is_big_page_node():
+            g_infer_context.req_manager.restore_big_page_state(
+                big_page_buffer_idx=shared_node.big_page_buffer_idx, req=self
+            )
+        else:
+            g_infer_context.req_manager.restore_small_page_state(req=self)
+
+    def _set_cache_prefix(
+        self,
+        shared_node: Optional[Union[TreeNode, HybridAttPagedTreeNode]],
+        shared_prefix_indexes: torch.Tensor,
+        private_tail_indexes: Optional[torch.Tensor] = None,
+    ):
+        """登记已命中的 KV 前缀，统一更新请求索引、持有长度和缓存命中信息。
+
+        Args:
+            shared_node: 请求最终保留引用的 radix cache 节点。普通匹配时为命中节点；
+                hybrid 小页尾部复制到私有 KV 后，为回退到的最近大页节点。
+                若没有可保留的共享大页，则为 None。节点引用计数由调用方的匹配、回退流程维护。
+            shared_prefix_indexes: 共享前缀对应的物理 KV 槽位索引，为按逻辑 token 顺序排列的
+                一维 CPU Tensor，每个元素代表一个 token 槽位，覆盖请求从位置 0 开始的前缀。
+                长度对应 shared_node 的完整前缀长度；没有共享节点时传空 Tensor。
+            private_tail_indexes: 紧接共享前缀的请求私有 KV 槽位索引，同样是一维 CPU Tensor。
+                用于 hybrid 小页命中后，将尾部 KV 复制到新分配槽位的场景；这里传入复制目标
+                的索引。其 KV 由当前请求独占，不计入 shared_node 的共享长度。无需私有尾部时
+                使用 None。两段索引均使用 token 单位，多 token 页模式下须覆盖完整物理页。
+
+        使用场景：
+            普通前缀命中、hybrid 大页命中，以及大页禁用时直接复用小页，都只传共享前缀。
+            hybrid 小页命中且允许重建尾部时，调用方先提交尾部 KV 拷贝、恢复小页 checkpoint，
+            再回退共享节点，最后分别传入保留的共享前缀和私有尾部索引，无需先拼接 Tensor。
+            容量不足、只能回退到大页时，只传该大页覆盖的共享前缀。
+
+        调用时请求索引必须为空，适用于新请求或已完成暂停释放的请求。
+        cur_kv_len、hold_kv_len 和 prompt_cache_len 均设为两段索引的总长度，因为私有尾部
+        也来自已计算的缓存 KV。例如共享前缀为 8、私有尾部为 4 时，命中长度为 12，
+        共享节点仍只保护前 8 个 token。shm_cur_kv_len 由外层匹配方法统一发布。
+        """
+        # 新请求或已完成暂停释放的请求，其索引状态为空，无需再次 reset。
+        self.shared_kv_node = shared_node
+        self.token_index.append_token_indexes(shared_prefix_indexes)
+        cache_len = shared_prefix_indexes.numel()
+        if private_tail_indexes is not None:
+            self.token_index.append_token_indexes(private_tail_indexes)
+            cache_len += private_tail_indexes.numel()
+        self.cur_kv_len = self.hold_kv_len = int(cache_len)
+        assert self.hold_kv_len % self.args.page_size == 0
+        self.shm_req.prompt_cache_len = self.cur_kv_len
+
     def is_master_req(self):
         """
         diverse 模式下，判断当前请求是否为独立主请求，其进行prefill后，将
@@ -971,6 +945,20 @@ def _kv_cache_alloc_need(self, target_kv_len: int) -> int:
         assert alloc_token_num % page_size == 0
         return alloc_token_num
 
+    def alloc_kv_mem(self, alloc_token_num: int) -> Optional[torch.Tensor]:
+        """按完整页追加 KV 容量，更新请求索引和持有长度；零申请返回 None。"""
+        if alloc_token_num == 0:
+            return None
+
+        assert alloc_token_num > 0 and alloc_token_num % self.args.page_size == 0
+        if g_infer_context.radix_cache is not None:
+            g_infer_context.radix_cache.free_radix_cache_to_get_enough_token(alloc_token_num)
+
+        mem_indexes = self.req_manager.mem_manager.alloc(alloc_token_num)
+        self.token_index.append_token_indexes(mem_indexes)
+        self.hold_kv_len += alloc_token_num
+        return mem_indexes
+
 
 class InferReqUpdatePack:
     """
diff --git a/lightllm/server/router/model_infer/infer_req_ext.py b/lightllm/server/router/model_infer/infer_req_ext.py
index 14a13ff66a..cf96d57c51 100644
--- a/lightllm/server/router/model_infer/infer_req_ext.py
+++ b/lightllm/server/router/model_infer/infer_req_ext.py
@@ -148,12 +148,6 @@ class FinalTokenMetadataExt:
     def __init__(self, req: "InferReq") -> None:
         self._req = req
 
-    def _mem_indexes(self) -> torch.Tensor:
-        # 延迟导入，避免与 infer_batch 形成模块级循环依赖。
-        from lightllm.server.router.model_infer.infer_batch import g_infer_context
-
-        return g_infer_context.req_manager.req_to_token_indexs[self._req.req_idx]
-
     def collect_prompt_logprobs(self) -> Optional[Tuple[np.ndarray, np.ndarray]]:
         """从 PromptLogprobsCaptureManager 提取 ``prompt_logprobs>0`` 的 top-k。
 
@@ -171,7 +165,7 @@ def collect_prompt_logprobs(self) -> Optional[Tuple[np.ndarray, np.ndarray]]:
             return None
 
         mgr = PromptLogprobsCaptureManager.get_instance()
-        mem_indexes = self._mem_indexes()[: req.shm_req.input_len - 1]
+        mem_indexes = req.token_index.get_cpu_token_indexes(0, req.shm_req.input_len - 1)
         return mgr.extract(mem_indexes, topk)
 
     def collect_routed_experts(self) -> Optional[np.ndarray]:
@@ -183,7 +177,7 @@ def collect_routed_experts(self) -> Optional[np.ndarray]:
         if capture_len <= 0:
             return None
 
-        mem_indexes = self._mem_indexes()[0:capture_len]
+        mem_indexes = req.token_index.get_cpu_token_indexes(0, capture_len)
         return MoeRouteInfoManager.get_instance().extract(mem_indexes)
 
     def dump(self) -> None:
diff --git a/lightllm/server/router/model_infer/infer_req_ext1.py b/lightllm/server/router/model_infer/infer_req_ext1.py
new file mode 100644
index 0000000000..7235bab546
--- /dev/null
+++ b/lightllm/server/router/model_infer/infer_req_ext1.py
@@ -0,0 +1,344 @@
+"""InferReq 的 token 索引扩展，仅保存请求槽位和索引状态。"""
+
+from abc import ABC, abstractmethod
+from typing import TYPE_CHECKING, Callable, List, Optional, Sequence
+
+import torch
+
+if TYPE_CHECKING:
+    from lightllm.server.router.model_infer.infer_batch import InferReq
+
+
+class ReqTokenIndex(ABC):
+    """单个请求的 token 索引，使用调用方指定的流维护；区间右端点不包含在内。
+
+    调用方保证 token 区间有效，并在构造时传入统一的索引流；在线服务在请求构造入口传入当前流。
+    GPU 索引读写固定在索引流，返回 GPU 视图前让调用流等待，Helper 不维护逐请求 CUDA 事件。
+    """
+
+    def __init__(self, infer_req: "InferReq", index_stream: torch.cuda.Stream):
+        # 仅保存固定槽位编号和 GPU 行视图，避免反向持有 InferReq 形成循环引用。
+        self.req_idx = infer_req.req_idx
+        self.gpu_indexes = infer_req.req_manager.req_to_token_indexs[self.req_idx]
+        # 索引流由外部选择，后续 GPU 索引操作不随调用流切换。
+        self._index_stream = index_stream
+        # 当前连续记录的 token 索引长度，包含页内预留槽位；
+        # CPU 扩容时复制 [0, token_index_len)，它不是缓存容量，也不是已算好 KV 的长度。
+        self.token_index_len = 0
+        # [0, gpu_index_submitted_len) 的最新索引写入已提交，后续索引流操作可按序使用。
+        # 前缀重建前由 reset 清零；它不是累计上传量，也不表示 GPU 传输已经完成。
+        self.gpu_index_submitted_len = 0
+
+    @abstractmethod
+    def append_token_indexes(self, mem_indexes: torch.Tensor):
+        """将 token 索引追加到当前记录的末尾，追加长度由 mem_indexes.numel() 决定。
+
+        非空输入必须是 CPU Tensor；CPU 缓存模式直接复制到自己的常驻缓存。
+        pure_gpu 按需 pin 后异步上传；已 pinned 的输入在 H2D 完成前不得改写，但可释放引用。
+        pin_mem_page_cache 仅接受从页边界追加完整页，调用方保证每页内物理 token 索引连续。
+        前缀替换先 reset，再按顺序追加重建。
+        """
+        pass
+
+    @abstractmethod
+    def get_cpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """读取有效长度内任意 token 区间，返回可跨请求生命周期持有的独立 CPU 快照。"""
+        pass
+
+    @abstractmethod
+    def get_gpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """在索引流准备 GPU 表视图；跨请求槽位复用持有时需另取快照。"""
+        pass
+
+    @abstractmethod
+    def get_upload_indexes(self) -> Optional[torch.Tensor]:
+        """返回 [gpu_index_submitted_len, token_index_len) 的 pinned CPU token 索引，不更新提交进度。
+
+        完整 CPU 缓存每次复制到新申请的 pinned 区间，页缓存按需展开到独立暂存区。
+        返回值不借用常驻缓存，调用方可直接批量上传；GPU 模式返回 None。
+        """
+        pass
+
+    @abstractmethod
+    def get_free_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """返回仅用于释放的 token 索引，可能是 CPU 或 GPU Tensor。
+
+        两种 CPU 缓存模式返回 CPU 快照；pure_gpu 返回借用的 GPU 请求表视图。
+
+        GPU 视图借用请求表，返回前让调用流等待索引流。重写请求槽位前，须在同一流
+        按序消费，或让索引流等待消费流。此接口不提供跨槽位复用的快照语义，
+        缓存和元数据仍使用 CPU getter。
+        """
+        pass
+
+    @abstractmethod
+    def reset(self):
+        """重置有效索引长度和 GPU 提交进度，释放 CPU 缓存引用。"""
+        pass
+
+    @staticmethod
+    def submit_pending_uploads(
+        token_indexes: Sequence["ReqTokenIndex"],
+        copy_batch: Optional[Callable[[List[torch.Tensor], List[torch.Tensor]], None]] = None,
+    ) -> None:
+        """统一准备待上传区间、提交拷贝并确认进度，随后让调用流承接依赖。
+
+        copy_batch 接收 GPU 目标和 pinned CPU 源列表，必须在当前索引流提交拷贝，
+        并保护异步源的生命周期；未提供时使用 PyTorch copy_。
+        只有整批提交成功才推进进度；提交不表示传输完成，无新上传也需等待已有写入。
+        """
+        if not token_indexes:
+            return
+        index_stream = token_indexes[0]._index_stream
+        assert all(
+            token_index._index_stream == index_stream for token_index in token_indexes
+        ), "All requests in a batch must use the same index stream"
+        cpu_slices, gpu_slices, pending_uploads = [], [], []
+        for token_index in token_indexes:
+            start, end = token_index.gpu_index_submitted_len, token_index.token_index_len
+            if start < end:
+                source = token_index.get_upload_indexes()
+                cpu_slices.append(source)
+                gpu_slices.append(token_index.gpu_indexes[start:end])
+                pending_uploads.append((token_index, end))
+        if pending_uploads:
+            with torch.cuda.stream(index_stream):
+                if copy_batch is None:
+                    for destination, source in zip(gpu_slices, cpu_slices):
+                        destination.copy_(source, non_blocking=True)
+                else:
+                    copy_batch(gpu_slices, cpu_slices)
+            for token_index, end in pending_uploads:
+                token_index._mark_gpu_indexes_submitted(end)
+        token_indexes[0]._wait_for_index_stream()
+
+    def _mark_gpu_indexes_submitted(self, end: int) -> None:
+        """只确认已成功提交的连续前缀，不把 GPU 完成状态混入长度计数。"""
+        assert self.gpu_index_submitted_len <= end <= self.token_index_len
+        self.gpu_index_submitted_len = end
+
+    def _wait_for_index_stream(self) -> None:
+        current_stream = torch.cuda.current_stream(self._index_stream.device)
+        if current_stream != self._index_stream:
+            current_stream.wait_stream(self._index_stream)
+
+    @staticmethod
+    def _alloc_return_pinned_tensor(need_size: int) -> torch.Tensor:
+        """申请 int32 token 索引的 pinned 区间，slab 容量为 1M token（4 MiB）。"""
+        buffer_token_num = 1024 * 1024
+        if need_size == 0 or need_size > buffer_token_num:
+            return torch.empty(need_size, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True)
+        # 在函数对象上惰性保存 slab 和 offset，所有索引实例共享，由推理线程串行申请。
+        allocator = ReqTokenIndex._alloc_return_pinned_tensor
+        buffer, start = getattr(allocator, "return_buffer_state", (None, 0))
+        end = start + need_size
+        if buffer is None or end > buffer_token_num:
+            buffer = torch.empty(
+                buffer_token_num, dtype=torch.int32, device="cpu", requires_grad=False, pin_memory=True
+            )
+            start = 0
+            end = need_size
+        # 每次只创建返回值这一个视图；offset 只前进，换新 slab 时才归零，不覆写在途 H2D 的源。
+        allocator.return_buffer_state = (buffer, end)
+        return buffer[start:end]
+
+
+class PinMemCacheReqTokenIndex(ReqTokenIndex):
+    """按请求保存普通 CPU token 缓存，上传时复制待提交区间到独立 pinned 区间。"""
+
+    def __init__(self, infer_req: "InferReq", index_stream: torch.cuda.Stream):
+        super().__init__(infer_req, index_stream)
+        self.input_len = infer_req.shm_req.input_len
+        self._cpu_buffer: Optional[torch.Tensor] = None
+
+    def append_token_indexes(self, mem_indexes: torch.Tensor):
+        """在 CPU 索引末尾追加，GPU 提交进度由上传流程维护。"""
+        append_len = mem_indexes.numel()
+        if append_len == 0:
+            return
+        assert mem_indexes.device.type == "cpu", "mem_indexes must be a CPU Tensor"
+        start = self.token_index_len
+        end = start + append_len
+        self._reserve(end)
+        self._cpu_buffer[start:end].copy_(mem_indexes)
+        self.token_index_len = end
+
+    def get_cpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        assert end <= self.token_index_len, "Read range exceeds recorded token indexes"
+        if start == end:
+            return torch.empty(0, dtype=self.gpu_indexes.dtype, device="cpu", requires_grad=False)
+        return self._cpu_buffer[start:end].clone()
+
+    def get_gpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """在索引流提交全部待上传索引，让调用流承接依赖后返回 GPU 视图。"""
+        assert end <= self.token_index_len, "Read range exceeds recorded token indexes"
+        if self.gpu_index_submitted_len < end:
+            self.submit_pending_uploads([self])
+        else:
+            self._wait_for_index_stream()
+        return self.gpu_indexes[start:end]
+
+    def get_upload_indexes(self) -> torch.Tensor:
+        """每次为待上传区间创建 pinned 快照，不保留常驻缓存或更新提交进度。"""
+        start, end = self.gpu_index_submitted_len, self.token_index_len
+        upload_indexes = self._alloc_return_pinned_tensor(end - start)
+        if start < end:
+            upload_indexes.copy_(self._cpu_buffer[start:end])
+        return upload_indexes
+
+    def get_free_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """返回用于释放的 CPU token 索引快照。"""
+        return self.get_cpu_token_indexes(start, end)
+
+    def reset(self):
+        # 上传使用独立 pinned 暂存区，常驻缓存可直接释放，无需等待在途 H2D。
+        self.token_index_len = self.gpu_index_submitted_len = 0
+        self._cpu_buffer = None
+
+    def _reserve(self, end: int):
+        """将 CPU 缓存扩至至少 end 个 token，只复制已有的有效索引。"""
+        old_capacity = 0 if self._cpu_buffer is None else self._cpu_buffer.numel()
+        if old_capacity >= end:
+            return
+        if old_capacity == 0:
+            # 首次按完整输入长度额外预留 2K 输出 token，同时覆盖本次写入区间。
+            reserve_len = max(self.input_len + 2 * 1024, end)
+        else:
+            # 容量不足时在所需长度上额外预留 2K token，减少后续 decode 扩容。
+            reserve_len = end + 2 * 1024
+        # 按 1K token 向上对齐，最多申请到 GPU 请求表允许的长度。
+        new_capacity = min(self.gpu_indexes.numel(), (reserve_len + 1023) // 1024 * 1024)
+        new_buffer = torch.empty(new_capacity, dtype=self.gpu_indexes.dtype, device="cpu", requires_grad=False)
+        if self.token_index_len:
+            new_buffer[: self.token_index_len].copy_(self._cpu_buffer[: self.token_index_len])
+        self._cpu_buffer = new_buffer
+
+
+class PinMemPageCacheReqTokenIndex(PinMemCacheReqTokenIndex):
+    """每个逻辑页仅保存首 token 的 pinned CPU 索引，对外仍使用 token 单位。"""
+
+    # 同一进程使用统一的 page_size，所有请求共享首次初始化的页内偏移。
+    _page_offsets: Optional[torch.Tensor] = None
+
+    def __init__(self, infer_req: "InferReq", index_stream: torch.cuda.Stream):
+        super().__init__(infer_req, index_stream)
+        self.page_size = infer_req.req_manager.page_size
+        if PinMemPageCacheReqTokenIndex._page_offsets is None:
+            PinMemPageCacheReqTokenIndex._page_offsets = torch.arange(
+                self.page_size, dtype=self.gpu_indexes.dtype, device="cpu", requires_grad=False
+            )
+
+    def append_token_indexes(self, mem_indexes: torch.Tensor):
+        """从页边界追加完整页；调用方保证同一逻辑页内物理 token 索引连续。"""
+        append_len = mem_indexes.numel()
+        if append_len == 0:
+            return
+        assert mem_indexes.device.type == "cpu", "mem_indexes must be a CPU Tensor"
+        start = self.token_index_len
+        end = start + append_len
+        assert start % self.page_size == 0, "Append start must be page-aligned"
+        assert end % self.page_size == 0, "Append end must be page-aligned"
+        self._reserve(end)
+        # 输入从页首开始，每隔 page_size 个 token 保存一个页首索引。
+        first_page = start // self.page_size
+        last_page = (end + self.page_size - 1) // self.page_size
+        self._cpu_buffer[first_page:last_page].copy_(mem_indexes[:: self.page_size])
+        self.token_index_len = end
+
+    def get_cpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        assert end <= self.token_index_len, "Read range exceeds recorded token indexes"
+        return self._expand_token_indexes(start, end, pin_memory=False)
+
+    def get_upload_indexes(self) -> torch.Tensor:
+        return self._expand_token_indexes(self.gpu_index_submitted_len, self.token_index_len, pin_memory=True)
+
+    def _reserve(self, end: int):
+        """首次按 GPU 请求表的最大长度申请全部页索引，后续追加直接复用。"""
+        if self._cpu_buffer is not None:
+            return
+        page_capacity = (self.gpu_indexes.numel() + self.page_size - 1) // self.page_size
+        self._cpu_buffer = torch.empty(
+            page_capacity, dtype=self.gpu_indexes.dtype, device="cpu", pin_memory=True, requires_grad=False
+        )
+
+    def _expand_token_indexes(self, start: int, end: int, pin_memory: bool) -> torch.Tensor:
+        """展开覆盖读取区间的完整页，再截取精确 token 范围，返回值不借用压缩缓存。"""
+        if start == end:
+            return torch.empty(
+                0, dtype=self.gpu_indexes.dtype, device="cpu", pin_memory=pin_memory, requires_grad=False
+            )
+        first_page = start // self.page_size
+        last_page = (end + self.page_size - 1) // self.page_size
+        token_num = (last_page - first_page) * self.page_size
+        if pin_memory:
+            # 上传使用一次性的 slab 区间，后续追加、扩容或 reset 均不会覆写在途 H2D 的源。
+            token_indexes = self._alloc_return_pinned_tensor(token_num)
+        else:
+            token_indexes = torch.empty(token_num, dtype=self.gpu_indexes.dtype, device="cpu", requires_grad=False)
+        page_starts = self._cpu_buffer[first_page:last_page]
+        # 将每个 page 的首索引加上页内偏移，直接展开到 token 级别的目标缓冲区。
+        torch.add(page_starts[:, None], self._page_offsets, out=token_indexes.view(-1, self.page_size))
+        # 读取可从页内任意位置开始或结束，页内偏移只影响返回范围，不改变压缩页缓存。
+        offset = start % self.page_size
+        return token_indexes[offset : offset + end - start]
+
+
+class PureGpuReqTokenIndex(ReqTokenIndex):
+    """直接维护 GPU 请求表，不持有 CPU 索引缓存；CPU 读取会等待 D2H 完成。"""
+
+    def append_token_indexes(self, mem_indexes: torch.Tensor):
+        append_len = mem_indexes.numel()
+        if append_len == 0:
+            return
+        assert mem_indexes.device.type == "cpu", "mem_indexes must be a CPU Tensor"
+        # 普通 CPU 输入在这里转为独立 pinned storage；已有 pinned 输入直接复用。
+        if not mem_indexes.is_pinned():
+            mem_indexes = mem_indexes.pin_memory()
+        start = self.token_index_len
+        end = start + append_len
+        # copy_ 会向 PyTorch pinned allocator 登记流使用，H2D 完成前不会回收源内存。
+        # 普通 CPU 输入已形成独立 pinned 副本；直接复用的 pinned 输入在传输完成前不能改写。
+        with torch.cuda.stream(self._index_stream):
+            self.gpu_indexes[start:end].copy_(mem_indexes, non_blocking=True)
+        self.token_index_len = end
+        self._mark_gpu_indexes_submitted(end)
+        self._wait_for_index_stream()
+
+    def get_cpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        assert end <= self.token_index_len, "Read range exceeds recorded token indexes"
+        # 阻塞到 D2H 完成，GPU 索引拷回后得到独立 CPU 快照。
+        with torch.cuda.stream(self._index_stream):
+            return self.gpu_indexes[start:end].detach().cpu()
+
+    def get_gpu_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        # append 已在索引流提交写入，读取只校验范围，不推进提交进度。
+        assert end <= self.gpu_index_submitted_len, "Read range exceeds submitted GPU indexes"
+        self._wait_for_index_stream()
+        return self.gpu_indexes[start:end]
+
+    def get_upload_indexes(self) -> None:
+        return None
+
+    def get_free_token_indexes(self, start: int, end: int) -> torch.Tensor:
+        """返回用于释放的 GPU token 索引视图，让调用流先等待索引写入。"""
+        # 复用 GPU getter 的范围校验和流等待，避免跨流释放读取尚未就绪的索引。
+        # allocator.free 会按页取首索引并阻塞 D2H；借用视图仍须在槽位重写前消费。
+        return self.get_gpu_token_indexes(start, end)
+
+    def reset(self):
+        # 保留 GPU 行内容，统一释放流程还需在槽位重写前消费借用的索引视图。
+        self.token_index_len = self.gpu_index_submitted_len = 0
+
+
+def create_req_token_index(infer_req: "InferReq", index_stream: torch.cuda.Stream) -> ReqTokenIndex:
+    """从请求参数选择实现；索引流由调用方显式指定，内部不选择默认流或当前流。"""
+    mode = infer_req.args.req_token_index_mode
+    if mode == "pin_mem_cache":
+        return PinMemCacheReqTokenIndex(infer_req, index_stream)
+    if mode == "pin_mem_page_cache":
+        return PinMemPageCacheReqTokenIndex(infer_req, index_stream)
+    if mode == "pure_gpu":
+        return PureGpuReqTokenIndex(infer_req, index_stream)
+    raise ValueError(
+        f"Unknown request token index mode: {mode}; expected pin_mem_cache, pin_mem_page_cache or pure_gpu"
+    )
diff --git a/lightllm/server/router/model_infer/mode_backend/base_backend.py b/lightllm/server/router/model_infer/mode_backend/base_backend.py
index 22f3cbe5dc..243427f1f5 100644
--- a/lightllm/server/router/model_infer/mode_backend/base_backend.py
+++ b/lightllm/server/router/model_infer/mode_backend/base_backend.py
@@ -11,8 +11,10 @@
 from lightllm.utils.log_utils import init_logger
 from lightllm.models import get_draft_model_class, get_model
 from lightllm.server.router.model_infer.infer_batch import InferReq, InferReqUpdatePack
+from lightllm.server.router.model_infer.infer_req_ext1 import ReqTokenIndex
 from lightllm.server.router.token_load import TokenLoad
 from lightllm.common.basemodel.basemodel import TpPartBaseModel
+from lightllm.common.cuda_wrapper import cuda_memcpy_pinned_h2d_batch_async
 from lightllm.common.basemodel.logprobs_manager import PromptLogprobsCaptureManager
 from lightllm.common.basemodel.moe_route_info_manager import MoeRouteInfoManager
 from lightllm.common.req_manager import HybridAttentionReqManager
@@ -555,9 +557,9 @@ def _read_pd_trans_io_buffer_and_update_req_status(self):
                                 req_to_next_token_ids = (
                                     self.model.req_manager.req_sampling_params_manager.req_to_next_token_ids
                                 )
-                                # to do 这个地方是否需要加流同步
                                 req_to_next_token_ids[req.req_idx, 0:1].fill_(obj.first_gen_token_id)
-                                torch.cuda.current_stream().synchronize()
+                                # handle 只更新 CPU 状态；推理入口通过 overlap_stream.wait_stream(current_stream)
+                                # 保证读取首 token 前写入已完成，无需在此阻塞 CPU。修改写入或推理流时须保留此依赖。
                                 InferReqUpdatePack(req_obj=req, output_len=req.cur_output_len).handle(
                                     next_token_id=obj.first_gen_token_id,
                                     next_token_logprob=obj.first_gen_token_logprob,
@@ -644,28 +646,11 @@ def remaining_prefill_tokens(req: InferReq) -> int:
         return ready_reqs
 
     # 一些可以复用的通用功能函数
-    def _alloc_req_kv_mem(
-        self,
-        req_obj: InferReq,
-        alloc_token_num: int,
-        no_blcoking_copy: bool = False,
-    ) -> Optional[torch.Tensor]:
-        if alloc_token_num == 0:
-            return None
-
-        assert alloc_token_num > 0 and alloc_token_num % self.args.page_size == 0
-        if g_infer_context.radix_cache is not None:
-            g_infer_context.radix_cache.free_radix_cache_to_get_enough_token(alloc_token_num)
-
-        old_hold_kv_len = req_obj.hold_kv_len
-        new_hold_kv_len = old_hold_kv_len + alloc_token_num
-        mem_indexes = g_infer_context.req_manager.mem_manager.alloc(alloc_token_num)
-        # 高频调度路径允许异步写入请求索引表，其他调用方默认保持原有的同步拷贝语义。
-        g_infer_context.req_manager.req_to_token_indexs[req_obj.req_idx, old_hold_kv_len:new_hold_kv_len].copy_(
-            mem_indexes, non_blocking=no_blcoking_copy
+    def _copy_req_token_indexes_to_gpu(self, run_reqs: List[InferReq]):
+        """选择本批请求；索引对象统一管理上传区间、提交进度和流依赖。"""
+        ReqTokenIndex.submit_pending_uploads(
+            [req.token_index for req in run_reqs], copy_batch=cuda_memcpy_pinned_h2d_batch_async
         )
-        req_obj.hold_kv_len = new_hold_kv_len
-        return mem_indexes
 
     def _get_classed_reqs(
         self,
@@ -691,6 +676,8 @@ def _get_classed_reqs(
         3. finished_reqs 需要释放的请求, 包含正常结束和aborted退出的请求。
         4. prefill_reqs 需要进行prefill操作的请求
         5. decode_reqs 需要进行decode操作的请求
+
+        返回前按所选模式准备两组请求的 GPU token 索引，并更新 token_index.gpu_index_submitted_len。
         """
         # 定期对 radix cache 进行 merge，防止查询插入的操作效率下降
         self._timer_merge_radix_tree()
@@ -714,6 +701,7 @@ def _get_classed_reqs(
         finished_reqs = []
         prefill_reqs = []
         decode_reqs = []
+        decode_allocations = []
 
         # 单轮最多处理少量因 token 容量不足而无法继续的请求，避免一次性影响大量请求。
         # 普通 Decode 请求进入暂停队列等待恢复；PD Decode 请求则强制提前结束并进入清理流程。
@@ -756,12 +744,10 @@ def _get_classed_reqs(
             if is_decode:
                 # KV 容量检查使用额外分配量，已有页的剩余容量可以覆盖部分或全部 decode 需求。
                 _, alloc_token_num = req_obj.decode_need_token_num()
-                # page_size 较小时，decode 会频繁触发 KV 内存分配。此处额外预申请不超过 8 个 token，
-                # 并将数量向下对齐到 page_size 的整数倍，以减少 alloc 调用次数并保持分页分配约束。
-                if alloc_token_num > 0 and self.args.page_size < 8:
-                    alloc_token_num += 8 // self.args.page_size * self.args.page_size
                 if alloc_token_num <= can_alloc_token_num:
-                    self._alloc_req_kv_mem(req_obj, alloc_token_num, no_blcoking_copy=True)
+                    # 只扣除必需容量；等本轮分类结束后，再从剩余预算中分配可选预留。
+                    if alloc_token_num > 0:
+                        decode_allocations.append((req_obj, alloc_token_num))
                     decode_reqs.append(req_obj)
                     can_alloc_token_num -= alloc_token_num
                 else:
@@ -798,7 +784,7 @@ def _get_classed_reqs(
                 if prefill_tokens + token_num > self.batch_max_tokens:
                     continue
                 if alloc_token_num <= can_alloc_token_num:
-                    self._alloc_req_kv_mem(req_obj, alloc_token_num, no_blcoking_copy=True)
+                    req_obj.alloc_kv_mem(alloc_token_num)
                     prefill_tokens += token_num
                     prefill_reqs.append(req_obj)
                     can_alloc_token_num -= alloc_token_num
@@ -807,6 +793,17 @@ def _get_classed_reqs(
                         req_obj.wait_pause = True
                         wait_pause_count += 1
 
+        # 所有接纳请求的必需容量均已计入预算，额外预留不能挤占后续 decode 或 prefill。
+        # page_size 较小时，decode 频繁分配 mem index 会增加 CPU overhead；因此在剩余预算允许时，
+        # 为 decode 请求额外预留一些 token，供后续步骤复用，减少分配次数并降低 CPU overhead。
+        for req_obj, alloc_token_num in decode_allocations:
+            if self.args.page_size < 8:
+                extra_token_num = min(8, can_alloc_token_num) // self.args.page_size * self.args.page_size
+                alloc_token_num += extra_token_num
+                can_alloc_token_num -= extra_token_num
+            # 将必需容量和可选预留合并申请，保留每个请求一次 alloc 的行为。
+            req_obj.alloc_kv_mem(alloc_token_num)
+
         # 先由控制器确定请求需要写入的缓存层级，再按是否包含 CPU cache 决定是否发起 offload。
         cache_controller = g_infer_context.cache_placement_controller
         new_finished_reqs = [req for req in finished_reqs if req.cpu_cache_task_status.is_not_started()]
@@ -850,6 +847,8 @@ def _get_classed_reqs(
                     prefill_reqs.append(decode_req)
                 decode_reqs = []
 
+        if prefill_reqs or decode_reqs:
+            self._copy_req_token_indexes_to_gpu(prefill_reqs + decode_reqs)
         return prefill_reqs, decode_reqs
 
     # 一些可以复用的通用功能函数
diff --git a/lightllm/server/router/model_infer/mode_backend/diverse_backend/impl.py b/lightllm/server/router/model_infer/mode_backend/diverse_backend/impl.py
index 46586959f1..1d8bf4b51e 100644
--- a/lightllm/server/router/model_infer/mode_backend/diverse_backend/impl.py
+++ b/lightllm/server/router/model_infer/mode_backend/diverse_backend/impl.py
@@ -89,6 +89,7 @@ def beam_prefill(self, event_pack: OverlapEventPack, prefill_reqs: List[InferReq
 
         # 第二阶段
         event_pack.notify_post_handle_and_wait_pre_post_handle()
+        # TODO: 分析 prefill 元数据事件同步的必要性，以及对 CPU/GPU overlap 的影响。
         model_output.prefill_mem_indexes_ready_event.synchronize()
         update_packs = self._diverse_pre_post_handle(run_reqs, is_chuncked_mode=not self.disable_chunked_prefill)
         # 第三阶段
@@ -128,6 +129,10 @@ def _diverse_copy(
         # 一些可以复用的通用功能函数
 
     def _diverse_pre_post_handle(self, run_reqs: List[InferReq], is_chuncked_mode: bool) -> List[InferReqUpdatePack]:
+        # diverse 会把 master/slave 已有前缀重新指向共享 KV，属于旧索引重写。
+        # 先在默认流等待上一轮模型读取结束，随后 CPU 缓存模式的批量上传、纯 GPU
+        # 模式的直接写入都沿默认流排序；普通请求仅追加新索引，不需要这一等待。
+        torch.cuda.default_stream().wait_stream(g_infer_context.get_overlap_stream())
         update_func_objs: List[InferReqUpdatePack] = []
         # 通用状态预先填充
         is_master_in_dp = self.is_master_in_dp
@@ -172,19 +177,21 @@ def _diverse_pre_post_handle(self, run_reqs: List[InferReq], is_chuncked_mode: b
                 pack = InferReqUpdatePack(req_obj=req_obj, output_len=pre_master_req_pack.output_len)
                 update_func_objs.append(pack)
 
+        # TODO: 分析 master/slave 状态更新后同步当前流的必要性，以及对 overlap 并行的影响。
         torch.cuda.current_stream().synchronize()
         return update_func_objs
 
     def _master_req_to_radix_cache(self, master_req: InferReq):
         key = master_req.get_input_token_ids()[0 : master_req.cur_kv_len]
         key = torch.tensor(key, dtype=torch.int64, device="cpu")
-        value = self.model.req_manager.req_to_token_indexs[master_req.req_idx][: master_req.cur_kv_len].detach().cpu()
-        prefix_len, new_shared_kv_node = self.radix_cache.insert(key, value)
+        # 一次性保存全部已持有索引。重建共享前缀后，页内预留的尾部仍由 master 持有，
+        # 不能随 reset 丢掉；独立 CPU 快照也避免缓存复用时覆盖尚未追加的尾部。
+        token_indexes = master_req.token_index.get_cpu_token_indexes(0, master_req.hold_kv_len)
+        prefix_len, new_shared_kv_node = self.radix_cache.insert(key, token_indexes[: master_req.cur_kv_len])
         old_prefix_len = 0 if master_req.shared_kv_node is None else master_req.shared_kv_node.node_prefix_total_len
         assert old_prefix_len <= master_req.cur_kv_len
-        self.model.mem_manager.free(
-            self.model.req_manager.req_to_token_indexs[master_req.req_idx][old_prefix_len:prefix_len]
-        )
+        # 释放在默认流同步消费；纯 GPU 模式可借用索引表视图，避免整段 token 索引 D2H。
+        self.model.mem_manager.free(master_req.token_index.get_free_token_indexes(old_prefix_len, prefix_len))
 
         # 将原有共享节点替换为新共享节点，新共享节点对应的长度为当前的cur_kv_len
         self.radix_cache.dec_node_ref_counter(master_req.shared_kv_node)
@@ -196,7 +203,9 @@ def _master_req_to_radix_cache(self, master_req: InferReq):
 
         share_node, kv_len, value = self.radix_cache.match_prefix(key, update_refs=False)
         assert share_node == new_shared_kv_node and kv_len == master_req.cur_kv_len
-        self.model.req_manager.req_to_token_indexs[master_req.req_idx][0 : master_req.cur_kv_len] = value
+        master_req.token_index.reset()
+        master_req.token_index.append_token_indexes(value)
+        master_req.token_index.append_token_indexes(token_indexes[master_req.cur_kv_len :])
         return
 
     def _copy_master_req_to_slave_req(self, slave_req: InferReq):
@@ -208,10 +217,9 @@ def _copy_master_req_to_slave_req(self, slave_req: InferReq):
 
         kv_len = master_req.cur_kv_len
 
-        self.model.req_manager.req_to_token_indexs[slave_req.req_idx][
-            0:kv_len
-        ] = self.model.req_manager.req_to_token_indexs[master_req.req_idx][0:kv_len]
-        # torch.cuda.current_stream().synchronize()
+        token_indexes = master_req.token_index.get_cpu_token_indexes(0, kv_len)
+        slave_req.token_index.reset()
+        slave_req.token_index.append_token_indexes(token_indexes)
         slave_req.shared_kv_node = master_req.shared_kv_node
         slave_req.cur_kv_len = kv_len
         # slave 在 prefill 阶段没有独立分配 KV，共享完成后需要同步其已持有的 KV 长度。
diff --git a/lightllm/server/router/model_infer/mode_backend/dp_backend/dp_shared_kv_trans.py b/lightllm/server/router/model_infer/mode_backend/dp_backend/dp_shared_kv_trans.py
index dedd3bb4a6..ba569eb352 100644
--- a/lightllm/server/router/model_infer/mode_backend/dp_backend/dp_shared_kv_trans.py
+++ b/lightllm/server/router/model_infer/mode_backend/dp_backend/dp_shared_kv_trans.py
@@ -24,7 +24,7 @@ def __init__(self, max_req_num: int, dp_size_in_node: int, backend):
         self.backend: DPChunkedPrefillBackend = backend
         self.max_req_num = max_req_num
 
-        # 0 代表 kv_len, 1 代表 radix_cache_len
+        # 每个请求、DP rank 对应两列：kv_len 和 req_idx。
         self.shared_req_infos = ShmArray(
             name="dp_shared_req_infos",
             shape=(self.max_req_num, dp_size_in_node, 2),
@@ -38,6 +38,11 @@ def fill_reqs_info(self, reqs: List[InferReq]):
         """
         填充请求的 kv 信息到共享内存中
         """
+        # 请求索引统一在默认流维护，原生 batch 的默认流适配由拷贝封装处理。
+        # DP 发布复用统一入口，不另建索引流或逐请求事件。
+        self.backend._copy_req_token_indexes_to_gpu(reqs)
+        # 同步 NCCL barrier 会等待默认流上排入的上传依赖完成；所有 TP rank
+        # 到达此屏障后，其他 DP rank 才能读取已发布的 GPU 索引和共享请求元数据。
         dist.barrier(group=self.backend.node_nccl_group)
         if self.backend.is_master_in_dp:
             self.shared_req_infos.arr[0 : len(reqs), self.dp_rank_in_node, self._KV_LEN_INDEX] = [
@@ -75,10 +80,10 @@ def build_shared_kv_trans_tasks(
 
             if is_current_dp_handle and trans_size > 0 and alloc_token_num <= g_infer_context.get_can_alloc_token_num():
                 assert req.hold_kv_len == req.cur_kv_len
-                mem_indexes = self.backend._alloc_req_kv_mem(req, alloc_token_num)
+                mem_indexes = req.alloc_kv_mem(alloc_token_num)
                 assert mem_indexes is not None
                 # mem_indexes 只描述需要复制的逻辑 KV；页尾预留槽位已经由
-                # _alloc_req_kv_mem 写入请求表，但不参与本次跨 DP 传输。
+                # alloc_kv_mem 记录到请求索引，但不参与本次跨 DP 传输。
                 mem_indexes = mem_indexes[:trans_size]
                 max_kv_len_dp_rank = self.shared_req_infos.arr[req_index, :, self._KV_LEN_INDEX].argmax()
                 max_kv_len_req_idx = int(self.shared_req_infos.arr[req_index, max_kv_len_dp_rank, self._REQ_IDX_INDEX])
@@ -122,7 +127,6 @@ def kv_trans(self, trans_tasks: List["TransTask"]):
                 dp_size_in_node=self.backend.dp_size_in_node,
                 rank_in_dp=self.backend.rank_in_dp,
             )
-            self.backend.logger.info(f"dp_i {self.dp_rank_in_node} transfer kv tokens num: {len(mem_indexes_tensor)}")
 
         for trans_task in trans_tasks:
             trans_task.req.cur_kv_len += len(trans_task.mem_indexes)
diff --git a/lightllm/server/router/model_infer/mode_backend/generic_pre_process.py b/lightllm/server/router/model_infer/mode_backend/generic_pre_process.py
index 7d3f34c4a7..b7f5632923 100644
--- a/lightllm/server/router/model_infer/mode_backend/generic_pre_process.py
+++ b/lightllm/server/router/model_infer/mode_backend/generic_pre_process.py
@@ -10,6 +10,7 @@
 def prepare_prefill_inputs(req_objs: List[InferReq], is_chuncked_mode: bool) -> Tuple[ModelInput, List[InferReq]]:
     run_reqs = []
     total_token_num = 0
+    total_input_token_num = 0
     input_ids = []
     b_req_idx = []
     b_seq_len = []
@@ -40,6 +41,7 @@ def prepare_prefill_inputs(req_objs: List[InferReq], is_chuncked_mode: bool) ->
         b_seq_len.append(seq_len)
         b_q_seq_len.append(input_token_len)
         input_ids.append(input_id)
+        total_input_token_num += input_token_len
         total_token_num += seq_len
         b_ready_cache_len.append(req.cur_kv_len)
         b_mtp_index.append(0)
@@ -55,14 +57,15 @@ def prepare_prefill_inputs(req_objs: List[InferReq], is_chuncked_mode: bool) ->
     max_cache_len = max(b_ready_cache_len, default=0)
     max_q_seq_len = max(b_q_seq_len, default=0)
 
-    input_ids = np.concatenate(input_ids, dtype=np.int64) if input_ids else np.empty((0,), dtype=np.int64)
-    input_ids = torch.tensor(input_ids, dtype=torch.int64, device="cpu")
-    b_req_idx = torch.tensor(b_req_idx, dtype=torch.int32, device="cpu")
-    b_seq_len = torch.tensor(b_seq_len, dtype=torch.int32, device="cpu")
+    token_ids = torch.empty(total_input_token_num, dtype=torch.int64, device="cpu", pin_memory=True)
+    if input_ids:
+        np.concatenate(input_ids, out=token_ids.numpy())
+    input_ids = token_ids
+
+    b_req_idx, b_seq_len, b_mtp_index, b_ready_cache_len, b_q_seq_len = torch.from_numpy(
+        np.asarray((b_req_idx, b_seq_len, b_mtp_index, b_ready_cache_len, b_q_seq_len), dtype=np.int32)
+    ).unbind()
     b_is_decode_req = torch.tensor(b_is_decode_req, dtype=torch.bool, device="cpu")
-    b_mtp_index = torch.tensor(b_mtp_index, dtype=torch.int32, device="cpu")
-    b_ready_cache_len = torch.tensor(b_ready_cache_len, dtype=torch.int32, device="cpu")
-    b_q_seq_len = torch.tensor(b_q_seq_len, dtype=torch.int32, device="cpu")
     b_prefill_start_loc = b_q_seq_len.cumsum(dim=0, dtype=torch.int32) - b_q_seq_len
 
     model_input = ModelInput(
@@ -120,14 +123,12 @@ def prepare_decode_inputs(req_objs: List[InferReq]) -> Tuple[ModelInput, List[In
     max_kv_seq_len = max(b_seq_len, default=0)
     max_q_seq_len = max(b_q_seq_len, default=1)
 
-    b_req_idx = torch.tensor(b_req_idx, dtype=torch.int32, device="cpu")
-    b_seq_len = torch.tensor(b_seq_len, dtype=torch.int32, device="cpu")
-    b_mtp_index = torch.tensor(b_mtp_index, dtype=torch.int32, device="cpu")
+    b_shared_seq_len = [req.get_radix_cache_shared_len() for req in run_reqs]
+    b_req_idx, b_seq_len, b_mtp_index, b_shared_seq_len = torch.from_numpy(
+        np.asarray((b_req_idx, b_seq_len, b_mtp_index, b_shared_seq_len), dtype=np.int32)
+    ).unbind()
     b_position_delta = build_b_position_delta(multimodal_params)
 
-    b_shared_seq_len = torch.tensor(
-        [req.get_radix_cache_shared_len() for req in run_reqs], dtype=torch.int32, device="cpu"
-    )
     b_shared_radix_node_id = torch.tensor(
         [-1 if req.shared_kv_node is None else req.shared_kv_node.time_id % INT64_MAX for req in run_reqs],
         dtype=torch.int64,
diff --git a/lightllm/server/router/model_infer/mode_backend/multi_level_kv_cache.py b/lightllm/server/router/model_infer/mode_backend/multi_level_kv_cache.py
index 12a2e9a048..a1db9c740b 100644
--- a/lightllm/server/router/model_infer/mode_backend/multi_level_kv_cache.py
+++ b/lightllm/server/router/model_infer/mode_backend/multi_level_kv_cache.py
@@ -103,8 +103,8 @@ def load_cpu_cache_to_reqs(self, reqs: List[InferReq]):
                     assert ready_page_num <= len(page_list)
                     need_pages = page_list[ready_page_num:]  # 只取需要的页面
 
-                    mem_indexes = self.backend._alloc_req_kv_mem(req, need_token_num)
-                    assert mem_indexes is not None
+                    allocated_mem_indexes = req.alloc_kv_mem(need_token_num)
+                    assert allocated_mem_indexes is not None
 
                     if self.need_sync_compute_stream():
                         # TODO fa3 现在必须使用同步模式, 未来需要移除
@@ -112,21 +112,17 @@ def load_cpu_cache_to_reqs(self, reqs: List[InferReq]):
                         # g_infer_context.get_overlap_stream().synchronize()
 
                     mem_manager = self.backend.model.mem_manager
-                    req_manager = self.backend.model.req_manager
-
-                    mem_indexes_cuda = mem_indexes.cuda(non_blocking=True)
-                    page_indexes_cuda = torch.tensor(need_pages, dtype=torch.int32, device="cpu").cuda(
+                    page_indexes_cuda = torch.tensor(need_pages, dtype=torch.int32, device="cpu", pin_memory=True).cuda(
                         non_blocking=True
                     )
-                    # hybrid 页面加载必须按完整 page 处理，否则可能缺失恢复运行态所需的 checkpoint，所以
-                    # 这里需要进行pad操作，使操作的页面是完整的。
+                    # hybrid 加载区间必须覆盖完整缓存页，包含恢复运行态所需的 checkpoint。
                     _start = page_len_start_list[ready_page_num]
 
-                    _end = req.cur_kv_len
-                    assert 0 <= _start <= _end, f"invalid pad range [{_start}, {_end}]"
-                    mem_indexes_cuda = torch.cat(
-                        [req_manager.req_to_token_indexs[req.req_idx, _start:_end], mem_indexes_cuda]
-                    )
+                    _end = req.hold_kv_len
+                    assert 0 <= _start <= req.cur_kv_len <= _end, f"invalid pad range [{_start}, {_end}]"
+                    # 索引准备与 KV 加载都在默认流，按提交顺序即可保证索引就绪。
+                    # 纯 GPU 模式直接借用请求表；请求槽位在 KV 加载完成前保持有效。
+                    mem_indexes_cuda = req.token_index.get_gpu_token_indexes(_start, _end)
 
                     assert (
                         len(mem_indexes_cuda) == page_len_list[len(page_list) - 1] - page_len_start_list[ready_page_num]
@@ -278,7 +274,9 @@ def _start_kv_cache_offload_task(
 
             move_token_num = page_len_list[item_size - 1]
             assert req.cur_kv_len >= move_token_num
-            token_indexes = self.backend.model.req_manager.req_to_token_indexs[req.req_idx, 0:move_token_num]
+            # getter 在请求索引流准备数据，并让当前卸载流等待，无需调用方额外同步。
+            token_indexes = req.token_index.get_gpu_token_indexes(0, move_token_num)
+            # 请求的 KV 和 req_idx 仍由卸载任务持有，直到下面的完成事件确认后才释放。
 
             mem_manager = self.backend.model.mem_manager
 
@@ -290,7 +288,7 @@ def _start_kv_cache_offload_task(
                 req=req,
             )
 
-            # 这个操作只是为了在offload 对应的cuda stream中，同步标记下对应的kv cache offload 操作已经完成，
+            # 多 DP rank 先在卸载流会合，再记录任务完成事件，供后续确认请求可释放。
             if self.backend.dp_world_size > 1:
                 dist.all_reduce(self.offload_sync_tensor, op=dist.ReduceOp.MAX, group=self.offload_sync_group)
 
diff --git a/lightllm/server/router/model_infer/mode_backend/pd/decode_node_impl/decode_impl.py b/lightllm/server/router/model_infer/mode_backend/pd/decode_node_impl/decode_impl.py
index 786c60b30a..3306f6ba9d 100644
--- a/lightllm/server/router/model_infer/mode_backend/pd/decode_node_impl/decode_impl.py
+++ b/lightllm/server/router/model_infer/mode_backend/pd/decode_node_impl/decode_impl.py
@@ -99,23 +99,12 @@ def _filter_not_ready_reqs(self, req_ids: List[int]) -> List[InferReq]:
                             f"(reason={req_obj.finish_status.get_finish_reason()}), kv transfer error"
                         )
 
-                # 提前释放有问题的 mem_index。cur_kv_len 只表示已经传输完成的
-                # 逻辑长度，hold_kv_len 还包含最后一个页面中尚未使用的预留槽位。
+                # 只有原共享前缀仍是可信 KV，失败尾部不能参与缓存插入。
+                # 保留 hold_kv_len 和索引状态，后续统一释放流程会回收全部持有的尾页。
                 old_prefix_len = 0 if req_obj.shared_kv_node is None else req_obj.shared_kv_node.node_prefix_total_len
-                error_mem_len = req_obj.hold_kv_len - old_prefix_len
-                if error_mem_len > 0:
-                    mem_indexes = (
-                        self.model.req_manager.req_to_token_indexs[
-                            req_obj.req_idx, old_prefix_len : req_obj.hold_kv_len
-                        ]
-                        .detach()
-                        .cpu()
-                    )
-                    self.model.mem_manager.free(mem_indexes)
-                    req_obj.cur_kv_len = old_prefix_len
-                    req_obj.hold_kv_len = old_prefix_len
-                    if self.is_master_in_dp:
-                        req_obj.shm_req.shm_cur_kv_len = req_obj.cur_kv_len
+                req_obj.cur_kv_len = old_prefix_len
+                if self.is_master_in_dp:
+                    req_obj.shm_req.shm_cur_kv_len = req_obj.cur_kv_len
 
             ans_list.append(req_obj)
         return ans_list
@@ -134,7 +123,7 @@ def _decode_node_gen_trans_tasks(self, req_obj: InferReq):
             assert req_obj.hold_kv_len == req_obj.cur_kv_len
             need_mem_size = req_obj._kv_cache_alloc_need(input_len)
 
-            mem_indexes = self._alloc_req_kv_mem(req_obj, need_mem_size)
+            mem_indexes = req_obj.alloc_kv_mem(need_mem_size)
             assert mem_indexes is not None
             # 传输只覆盖真实 KV；最后一个模型页面中尚未使用的部分继续留在
             # req_to_token_indexs 中，供后续 decode 直接切片使用。
diff --git a/lightllm/server/router/model_infer/mode_backend/pd/prefill_node_impl/prefill_impl.py b/lightllm/server/router/model_infer/mode_backend/pd/prefill_node_impl/prefill_impl.py
index 8641e3b69c..451afe4fa9 100644
--- a/lightllm/server/router/model_infer/mode_backend/pd/prefill_node_impl/prefill_impl.py
+++ b/lightllm/server/router/model_infer/mode_backend/pd/prefill_node_impl/prefill_impl.py
@@ -122,12 +122,7 @@ def _create_pd_trans_task(
 
         pd_decode_node_info = req_obj.sampling_param.pd_decode_node
         if page_kind == "kv":
-            mem_indexes = (
-                self.model.req_manager.req_to_token_indexs[req_obj.req_idx, kv_start_index:kv_end_index]
-                .detach()
-                .cpu()
-                .tolist()
-            )
+            mem_indexes = req_obj.token_index.get_cpu_token_indexes(kv_start_index, kv_end_index).tolist()
             req_idx = None
         elif page_kind == "att_state":
             mem_indexes = []
diff --git a/lightllm/utils/custom_kernel_utis.py b/lightllm/utils/custom_kernel_utis.py
index 1a0b46ca0d..6ecfffbea6 100644
--- a/lightllm/utils/custom_kernel_utis.py
+++ b/lightllm/utils/custom_kernel_utis.py
@@ -21,6 +21,7 @@ def custom_cat(tensors):
     for t, size in zip(tensors, sizes):
         out_tensor[start_loc : (start_loc + size)].copy_(t, non_blocking=True)
         start_loc += size
+    # TODO: 分析 D2H 完成等待的阻塞开销，以及是否能缩小当前流的同步范围。
     torch.cuda.current_stream().synchronize()
 
     return out_tensor
diff --git a/test/benchmark/benchmark_cpu_pinned_alloc_keepalive.py b/test/benchmark/benchmark_cpu_pinned_alloc_keepalive.py
new file mode 100644
index 0000000000..02e7bad583
--- /dev/null
+++ b/test/benchmark/benchmark_cpu_pinned_alloc_keepalive.py
@@ -0,0 +1,289 @@
+"""Compare torch.empty allocation latency while retaining every tensor.
+
+Example (select an idle GPU and a CPU on the same NUMA node):
+    CUDA_VISIBLE_DEVICES=4 python test/benchmark/benchmark_cpu_pinned_alloc_keepalive.py \
+        --cpu 48 --malloc-usable-size --output-dir /tmp/cpu_pinned_alloc
+
+Each mode/repeat runs in a fresh process. Warmup and measured tensors, across
+all sizes, remain alive until process exit. No fill, H2D, or release is timed.
+Ordinary CPU empty allocations need not fault in all physical pages. This is
+an allocation API benchmark, not a comparison of ready-to-use memory bandwidth.
+"""
+
+import argparse
+import csv
+import ctypes
+import gc
+import json
+import os
+from pathlib import Path
+import platform
+import random
+import statistics
+import subprocess
+import sys
+import time
+
+
+MIB = 1024 * 1024
+DEFAULT_SIZES = [
+    64,
+    256,
+    1024,
+    4096,
+    4097,
+    16384,
+    65536,
+    262144,
+    MIB,
+    MIB + 1,
+    4 * MIB,
+    16 * MIB,
+    64 * MIB,
+    64 * MIB + 1,
+]
+
+
+def percentile(values, fraction):
+    ordered = sorted(values)
+    position = (len(ordered) - 1) * fraction
+    lower = int(position)
+    upper = min(lower + 1, len(ordered) - 1)
+    return ordered[lower] + (ordered[upper] - ordered[lower]) * (position - lower)
+
+
+def rss_bytes():
+    return int(Path("/proc/self/statm").read_text().split()[1]) * os.sysconf("SC_PAGE_SIZE")
+
+
+def sample_count(size, args):
+    # Both modes use the same count. Budget using pinned allocator rounding.
+    rounded = 1 << (size - 1).bit_length()
+    return max(args.min_count, min(args.max_count, args.case_mib * MIB // rounded))
+
+
+def run_worker(args):
+    if args.cpu is not None:
+        os.sched_setaffinity(0, {args.cpu})
+    import torch
+
+    torch.set_num_threads(1)
+    torch.cuda.init()
+    pinned = args.mode == "pinned"
+    # Retain warmup too: initialization is excluded without returning any blocks.
+    warmup = [torch.empty(4096, dtype=torch.uint8, device="cpu", pin_memory=pinned) for _ in range(32)]
+    usable_size = None
+    if args.malloc_usable_size and not pinned:
+        if platform.libc_ver()[0] != "glibc" or os.environ.get("LD_PRELOAD"):
+            raise RuntimeError("--malloc-usable-size requires the default glibc CPU allocator without LD_PRELOAD")
+        # Diagnostic only; never call this on CUDA host allocations. The caller
+        # must use a PyTorch build with the default system CPU allocator.
+        usable_size = ctypes.CDLL(None).malloc_usable_size
+        usable_size.argtypes = [ctypes.c_void_p]
+        usable_size.restype = ctypes.c_size_t
+
+    sizes = list(args.sizes)
+    random.Random(args.repeat).shuffle(sizes)
+    kept = []
+    rows = []
+    gc.disable()
+    for size in sizes:
+        count = sample_count(size, args)
+        tensors = [None] * count
+        durations_ns = [0] * count
+        before_rss = rss_bytes()
+        before_stats = torch.cuda.memory.host_memory_stats() if pinned else None
+        for index in range(count):
+            start = time.perf_counter_ns()
+            tensor = torch.empty(size, dtype=torch.uint8, device="cpu", pin_memory=pinned, requires_grad=False)
+            end = time.perf_counter_ns()
+            tensors[index] = tensor
+            durations_ns[index] = end - start
+        kept.append(tensors)
+        after_stats = torch.cuda.memory.host_memory_stats() if pinned else None
+        after_rss = rss_bytes()
+        # Validation and diagnostics are outside the timed region.
+        assert all(t.nbytes == size and t.untyped_storage().nbytes() == size for t in tensors)
+        assert len({t.data_ptr() for t in tensors}) == count
+        if pinned:
+            actual_allocs = after_stats["num_host_alloc"] - before_stats["num_host_alloc"]
+            assert actual_allocs == count, (size, count, actual_allocs)
+            block_bytes = (after_stats["allocated_bytes.current"] - before_stats["allocated_bytes.current"]) // count
+            assert block_bytes == 1 << (size - 1).bit_length(), (size, block_bytes)
+        else:
+            actual_allocs = block_bytes = None
+        usable = [usable_size(t.data_ptr()) for t in tensors] if usable_size else []
+        assert not usable or min(usable) >= size
+        rows.append(
+            {
+                "mode": args.mode,
+                "repeat": args.repeat,
+                "requested_bytes": size,
+                "count": count,
+                "latency_ns": durations_ns,
+                "median_us": statistics.median(durations_ns) / 1000,
+                "mean_us": statistics.mean(durations_ns) / 1000,
+                "p95_us": percentile(durations_ns, 0.95) / 1000,
+                "rss_delta_bytes": after_rss - before_rss,
+                "pinned_block_bytes": block_bytes,
+                "new_host_allocations": actual_allocs,
+                "cpu_malloc_usable_min": min(usable) if usable else None,
+                "cpu_malloc_usable_median": statistics.median(usable) if usable else None,
+                "cpu_malloc_usable_max": max(usable) if usable else None,
+            }
+        )
+
+    metadata = {
+        "python": platform.python_version(),
+        "torch": torch.__version__,
+        "torch_commit": torch.version.git_version,
+        "cuda": torch.version.cuda,
+        "gpu": torch.cuda.get_device_name(0),
+        "cuda_visible_devices": os.environ.get("CUDA_VISIBLE_DEVICES"),
+        "cpu_affinity": sorted(os.sched_getaffinity(0)),
+        "libc": platform.libc_ver(),
+        "cpu_model": next(
+            line.split(":", 1)[1].strip()
+            for line in Path("/proc/cpuinfo").read_text().splitlines()
+            if line.startswith("model name")
+        ),
+        "final_rss_bytes": rss_bytes(),
+        "warmup_count": len(warmup),
+        "retained_measured_tensors": sum(len(group) for group in kept),
+        "size_order": sizes,
+    }
+    Path(args.worker_output).write_text(json.dumps({"metadata": metadata, "rows": rows}, indent=2) + "\n")
+    # Exit while warmup/kept still own their storages; never release measured
+    # tensors inside this process or let local-variable teardown recycle them.
+    os._exit(0)
+
+
+def summarize(trials, sizes):
+    summary = []
+    for size in sorted(sizes):
+        row = {"requested_bytes": size}
+        for mode in ("cpu", "pinned"):
+            selected = [
+                r for trial in trials for r in trial["rows"] if r["mode"] == mode and r["requested_bytes"] == size
+            ]
+            times = [value for item in selected for value in item["latency_ns"]]
+            row[f"{mode}_count_per_repeat"] = selected[0]["count"]
+            row[f"{mode}_median_us"] = statistics.median(times) / 1000
+            row[f"{mode}_mean_us"] = statistics.mean(times) / 1000
+            row[f"{mode}_p95_us"] = percentile(times, 0.95) / 1000
+            row[f"{mode}_repeat_median_min_us"] = min(r["median_us"] for r in selected)
+            row[f"{mode}_repeat_median_max_us"] = max(r["median_us"] for r in selected)
+            if mode == "cpu":
+                usable = [r for r in selected if r["cpu_malloc_usable_min"] is not None]
+                row["cpu_malloc_usable_min"] = min(r["cpu_malloc_usable_min"] for r in usable) if usable else None
+                row["cpu_malloc_usable_max"] = max(r["cpu_malloc_usable_max"] for r in usable) if usable else None
+            else:
+                row["pinned_block_bytes"] = selected[0]["pinned_block_bytes"]
+        row["pinned_over_cpu_median"] = row["pinned_median_us"] / row["cpu_median_us"]
+        summary.append(row)
+    return summary
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--output-dir", type=Path, default=Path("/tmp/cpu_pinned_alloc_keepalive"))
+    parser.add_argument("--sizes", nargs="+", type=int, default=DEFAULT_SIZES, help="Allocation sizes in bytes")
+    parser.add_argument("--repeats", type=int, default=3)
+    parser.add_argument(
+        "--case-mib", type=int, default=128, help="Target retained pinned MiB per size; min-count may exceed it"
+    )
+    parser.add_argument("--min-count", type=int, default=8)
+    parser.add_argument("--max-count", type=int, default=2048)
+    parser.add_argument("--cpu", type=int, help="Bind workers to this logical CPU")
+    parser.add_argument(
+        "--malloc-usable-size", action="store_true", help="Probe CPU malloc sizes; only with default glibc allocator"
+    )
+    parser.add_argument("--worker", action="store_true", help=argparse.SUPPRESS)
+    parser.add_argument("--mode", choices=["cpu", "pinned"], help=argparse.SUPPRESS)
+    parser.add_argument("--repeat", type=int, default=0, help=argparse.SUPPRESS)
+    parser.add_argument("--worker-output", help=argparse.SUPPRESS)
+    args = parser.parse_args()
+    if min(args.sizes + [args.repeats, args.case_mib, args.min_count]) < 1 or args.max_count < args.min_count:
+        parser.error("Sizes/counts/repeats must be positive and max-count >= min-count")
+    if len(set(args.sizes)) != len(args.sizes):
+        parser.error("Sizes must be unique")
+    for name in ("PYTORCH_ALLOC_CONF", "PYTORCH_CUDA_ALLOC_CONF"):
+        if os.environ.get(name):
+            parser.error(f"Unset {name} to measure default allocator behavior")
+    if args.worker:
+        run_worker(args)
+
+    args.output_dir.mkdir(parents=True, exist_ok=False)
+    trials = []
+    estimated_pinned = sum(sample_count(size, args) * (1 << (size - 1).bit_length()) for size in args.sizes)
+    print(f"Estimated retained pinned payload per worker: {estimated_pinned / MIB:.1f} MiB", flush=True)
+    for repeat in range(args.repeats):
+        # Alternate mode order; both modes use the same shuffled size order.
+        for mode in ("cpu", "pinned") if repeat % 2 == 0 else ("pinned", "cpu"):
+            path = args.output_dir / f"{mode}_{repeat}.json"
+            command = [
+                sys.executable,
+                str(Path(__file__).resolve()),
+                "--worker",
+                "--mode",
+                mode,
+                "--repeat",
+                str(repeat),
+                "--worker-output",
+                str(path),
+                "--case-mib",
+                str(args.case_mib),
+                "--min-count",
+                str(args.min_count),
+                "--max-count",
+                str(args.max_count),
+                "--sizes",
+                *map(str, args.sizes),
+            ]
+            if args.cpu is not None:
+                command.extend(["--cpu", str(args.cpu)])
+            if args.malloc_usable_size:
+                command.append("--malloc-usable-size")
+            subprocess.run(command, check=True)
+            trial = json.loads(path.read_text())
+            trials.append(trial)
+            print(
+                f'{mode} repeat {repeat + 1}: retained {trial["metadata"]["retained_measured_tensors"]} tensors; '
+                f'RSS {trial["metadata"]["final_rss_bytes"] / MIB:.1f} MiB',
+                flush=True,
+            )
+
+    summary = summarize(trials, args.sizes)
+    with (args.output_dir / "summary.csv").open("w", newline="") as output:
+        writer = csv.DictWriter(output, fieldnames=list(summary[0]))
+        writer.writeheader()
+        writer.writerows(summary)
+    (args.output_dir / "results.json").write_text(
+        json.dumps(
+            {
+                "parameters": {
+                    key: str(value) if isinstance(value, Path) else value for key, value in vars(args).items()
+                },
+                "method": "torch.empty(uint8); no fill/copy; "
+                "all warmup and measured tensors retained until process exit; "
+                "CUDA initialization excluded; GC disabled; per-allocation timings; fresh process per mode/repeat",
+                "estimated_retained_pinned_bytes": estimated_pinned,
+                "summary": summary,
+                "trials": trials,
+            },
+            indent=2,
+        )
+        + "\n"
+    )
+    print("bytes,count/repeat,cpu median us,pinned median us,ratio,pinned block bytes,cpu malloc min/max")
+    for row in summary:
+        print(
+            f'{row["requested_bytes"]},{row["cpu_count_per_repeat"]},{row["cpu_median_us"]:.3f},'
+            f'{row["pinned_median_us"]:.3f},{row["pinned_over_cpu_median"]:.2f},'
+            f'{row["pinned_block_bytes"]},{row["cpu_malloc_usable_min"]}/{row["cpu_malloc_usable_max"]}'
+        )
+    print(f"Results: {args.output_dir}")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/benchmark/benchmark_event_synchronize_gil.py b/test/benchmark/benchmark_event_synchronize_gil.py
new file mode 100644
index 0000000000..669bc7eb3a
--- /dev/null
+++ b/test/benchmark/benchmark_event_synchronize_gil.py
@@ -0,0 +1,237 @@
+"""实测 CUDA Event 等待是否释放 GIL，以及服务保活线程等待时主线程能否执行 Python。
+
+示例：CUDA_VISIBLE_DEVICES=6 PYTHONPATH=. python test/benchmark/benchmark_event_synchronize_gil.py
+所有等待都在 cuda_wrapper 的同一个后台线程池执行，主线程只做 Python 计数和时间采样。
+PyDLL.usleep 是持有 GIL 的负对照，CDLL.usleep 是释放 GIL 的正对照。
+"""
+
+import argparse
+import ctypes
+from datetime import datetime, timezone
+import json
+import os
+from pathlib import Path
+import statistics
+import sys
+import threading
+import time
+
+
+METHODS = (
+    "pydll_holds_gil",
+    "cdll_releases_gil",
+    "event_blocking_true",
+    "event_blocking_false",
+    "service_release_worker",
+)
+COUNTER_CHUNK = 4096
+
+
+def calibrate_delay(torch, stream, wait_ms):
+    """先预热，再按实际 GPU 时间估算约 wait_ms 的延时，避免依赖空闲时的 SM 频率。"""
+    reference_cycles = 50_000_000
+    start = torch.cuda.Event(enable_timing=True)
+    end = torch.cuda.Event(enable_timing=True)
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(reference_cycles)
+    stream.synchronize()
+    durations = []
+    for _ in range(3):
+        with torch.cuda.stream(stream):
+            start.record(stream)
+            torch.cuda._sleep(reference_cycles)
+            end.record(stream)
+        end.synchronize()
+        durations.append(start.elapsed_time(end))
+    return int(reference_cycles * wait_ms / statistics.median(durations)), durations
+
+
+def timed_wait(wait, state):
+    state["thread"] = threading.current_thread().name
+    state["start_ns"] = time.perf_counter_ns()
+    try:
+        wait()
+    finally:
+        state["end_ns"] = time.perf_counter_ns()
+        state["finished"] = True
+
+
+def count_python_work(state):
+    # 此循环不调用 torch、CUDA、sleep、锁或 Future；只有 Python 整数加法和时间采样。
+    # 每 4096 次采样一次，既能验证持续推进，也避免保存过多时间戳。
+    samples = []
+    counter = 0
+    while not state["finished"]:
+        for _ in range(COUNTER_CHUNK):
+            counter += 1
+        samples.append((time.perf_counter_ns(), counter))
+    return samples
+
+
+def run_trial(torch, wrapper, stream, host_waits, cycles, method, repeat, wait_ms, margin_ms):
+    event = None
+    timing_start, timing_end = None, None
+    sources = None
+    pending_before_wait = None
+    if method in host_waits:
+        wait = lambda: host_waits[method](int(wait_ms * 1000))
+    else:
+        # 被测 Event 与服务保持相同的 enable_timing=False 配置；计时用独立事件。
+        event = torch.cuda.Event(blocking=method != "event_blocking_false")
+        timing_start = torch.cuda.Event(enable_timing=True)
+        timing_end = torch.cuda.Event(enable_timing=True)
+        if method == "service_release_worker":
+            sources = [torch.empty(16, dtype=torch.int32, device="cpu", pin_memory=True, requires_grad=False)]
+        with torch.cuda.stream(stream):
+            timing_start.record(stream)
+            torch.cuda._sleep(cycles)
+            event.record(stream)
+            timing_end.record(stream)
+        pending_before_wait = not event.query()
+        if method == "service_release_worker":
+            wait = lambda: wrapper.PinnedH2DBatchCopier._wait_and_release_pinned_sources(event, sources)
+        else:
+            wait = event.synchronize
+
+    state = {"finished": False}
+    future = wrapper.PinnedH2DBatchCopier._get_release_pool().submit(timed_wait, wait, state)
+    samples = count_python_work(state)
+    future.result(timeout=10)
+    gpu_delay_ms = None
+    if event is not None:
+        stream.synchronize()
+        gpu_delay_ms = timing_start.elapsed_time(timing_end)
+
+    # 严格截取等待内部区间，排除进入原生调用之前和返回 Python 之后的调度窗口。
+    margin_ns = int(margin_ms * 1_000_000)
+    window_start = state["start_ns"] + margin_ns
+    window_end = state["end_ns"] - margin_ns
+    interior_samples = [(stamp, count) for stamp, count in samples if window_start < stamp < window_end]
+    increments = interior_samples[-1][1] - interior_samples[0][1] if len(interior_samples) > 1 else 0
+    progress_span_ms = (
+        (interior_samples[-1][0] - interior_samples[0][0]) / 1_000_000 if len(interior_samples) > 1 else 0.0
+    )
+    duration_ms = (state["end_ns"] - state["start_ns"]) / 1_000_000
+    # 等待不足目标的 60% 或内部窗口不足 150ms 时不接受结果，避免过短等待误判 GIL 行为。
+    valid_wait = duration_ms >= wait_ms * 0.6 and window_end - window_start >= 150_000_000
+    expects_gil_release = method != "pydll_holds_gil"
+    if expects_gil_release:
+        # 正对照要求内部窗口持续推进至少 100ms；负对照要求内部完全没有 Python 采样。
+        progress_matches = increments > 0 and progress_span_ms >= 100
+    else:
+        progress_matches = len(interior_samples) == 0
+    passed = valid_wait and progress_matches and pending_before_wait is not False
+    if sources is not None:
+        passed = passed and len(sources) == 0
+
+    return {
+        "method": method,
+        "repeat": repeat,
+        "wait_ms": duration_ms,
+        "gpu_delay_ms": gpu_delay_ms,
+        "event_pending_before_wait": pending_before_wait,
+        "interior_window_ms": max(0.0, (window_end - window_start) / 1_000_000),
+        "interior_samples": len(interior_samples),
+        "python_increments": increments,
+        "progress_span_ms": progress_span_ms,
+        "worker_thread": state["thread"],
+        "sources_cleared": len(sources) == 0 if sources is not None else None,
+        "valid_wait": valid_wait,
+        "passed": passed,
+    }
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--output-dir", type=Path)
+    parser.add_argument("--repeats", type=int, default=5)
+    parser.add_argument("--wait-ms", type=float, default=500)
+    parser.add_argument("--margin-ms", type=float, default=50)
+    args = parser.parse_args()
+    visible = os.environ.get("CUDA_VISIBLE_DEVICES", "")
+    if not visible or "," in visible or visible == "3":
+        parser.error("Set CUDA_VISIBLE_DEVICES to one idle GPU; GPU 3 must not be used")
+    if args.repeats < 1 or args.wait_ms < 400 or not 0 < args.margin_ms < args.wait_ms / 4:
+        parser.error(
+            "Repeats must be positive, wait must be at least 400 ms, and margin must be between 0 and wait / 4"
+        )
+
+    import torch
+    from lightllm.common import cuda_wrapper
+
+    torch.set_num_threads(1)
+    torch.cuda.set_device(0)
+    stream = torch.cuda.Stream()
+    # ctypes 的两种加载方式调用同一个 libc 函数，唯一关注的区别是是否在原生等待中释放 GIL。
+    libraries = {"pydll_holds_gil": ctypes.PyDLL(None), "cdll_releases_gil": ctypes.CDLL(None)}
+    host_waits = {}
+    for method, library in libraries.items():
+        function = library.usleep
+        function.argtypes = [ctypes.c_uint]
+        function.restype = ctypes.c_int
+        host_waits[method] = function
+    cuda_wrapper.PinnedH2DBatchCopier._get_release_pool().submit(lambda: None).result(timeout=10)
+    cycles, calibration_ms = calibrate_delay(torch, stream, args.wait_ms)
+    root = args.output_dir or Path("/tmp") / (
+        "event_synchronize_gil_" + datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ")
+    )
+    root.mkdir(parents=True, exist_ok=False)
+    print(f"Results directory: {root}", flush=True)
+    rows = []
+    try:
+        for repeat in range(args.repeats):
+            for method in METHODS:
+                rows.append(
+                    run_trial(
+                        torch, cuda_wrapper, stream, host_waits, cycles, method, repeat, args.wait_ms, args.margin_ms
+                    )
+                )
+            print(f"Completed round {repeat + 1}/{args.repeats}", flush=True)
+    finally:
+        stream.synchronize()
+        cuda_wrapper.PinnedH2DBatchCopier._get_release_pool().submit(lambda: None).result(timeout=10)
+    summaries = []
+    for method in METHODS:
+        cases = [row for row in rows if row["method"] == method]
+        summary = {
+            "method": method,
+            "passed": sum(row["passed"] for row in cases),
+            "trials": len(cases),
+            "wait_ms_median": statistics.median(row["wait_ms"] for row in cases),
+            "python_increments_min": min(row["python_increments"] for row in cases),
+            "python_increments_median": statistics.median(row["python_increments"] for row in cases),
+            "python_increments_max": max(row["python_increments"] for row in cases),
+            "progress_span_ms_min": min(row["progress_span_ms"] for row in cases),
+        }
+        summaries.append(summary)
+        print(json.dumps(summary, ensure_ascii=False), flush=True)
+    evidence = {
+        "environment": {
+            "python": sys.version.split()[0],
+            "torch": torch.__version__,
+            "cuda": torch.version.cuda,
+            "cuda_visible_devices": visible,
+            "gpu": torch.cuda.get_device_name(0),
+            "python_switch_interval_s": sys.getswitchinterval(),
+        },
+        "config": {
+            "repeats": args.repeats,
+            "target_wait_ms": args.wait_ms,
+            "margin_ms": args.margin_ms,
+            "counter_chunk": COUNTER_CHUNK,
+            "gpu_delay_cycles": cycles,
+            "calibration_ms": calibration_ms,
+        },
+        "summary": summaries,
+        "trials": rows,
+        "passed": all(row["passed"] for row in rows),
+    }
+    (root / "result.json").write_text(json.dumps(evidence, ensure_ascii=False, indent=2) + "\n")
+    if not evidence["passed"]:
+        raise SystemExit(
+            "Experiment criteria were not fully met; check wait durations and window progress in result.json"
+        )
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/benchmark/benchmark_pinned_allocation_rounding.py b/test/benchmark/benchmark_pinned_allocation_rounding.py
new file mode 100644
index 0000000000..d77bb032e6
--- /dev/null
+++ b/test/benchmark/benchmark_pinned_allocation_rounding.py
@@ -0,0 +1,127 @@
+"""Verify pinned allocation rounding and cache reuse in a standalone CUDA process.
+
+Example:
+    CUDA_VISIBLE_DEVICES=4 python test/benchmark/benchmark_pinned_allocation_rounding.py \
+        --output /tmp/pinned_allocation_rounding.json
+"""
+
+import argparse
+import json
+import os
+from pathlib import Path
+import weakref
+
+import torch
+
+
+MIB = 1024 * 1024
+STAT_KEYS = ("allocated_bytes.current", "active_bytes.current", "num_host_alloc", "num_host_free")
+
+
+def snapshot():
+    stats = torch.cuda.memory.host_memory_stats()
+    return {key: stats[key] for key in STAT_KEYS}
+
+
+def measure(label, numel, dtype):
+    # 仅清理本独立进程的 host 缓存，确保首次申请无法命中已有尺寸桶。
+    torch._C._host_emptyCache()
+    before = snapshot()
+    tensor = torch.empty(numel, dtype=dtype, device="cpu", pin_memory=True, requires_grad=False)
+    allocated = snapshot()
+    requested_bytes = tensor.nbytes
+    expected_block_bytes = 1 << (requested_bytes - 1).bit_length()
+    block_bytes = allocated["allocated_bytes.current"] - before["allocated_bytes.current"]
+    pointer = tensor.data_ptr()
+    storage_bytes = tensor.untyped_storage().nbytes()
+    assert tensor.is_pinned()
+    assert block_bytes == expected_block_bytes, (label, block_bytes, expected_block_bytes)
+    assert allocated["active_bytes.current"] - before["active_bytes.current"] == block_bytes
+    assert allocated["num_host_alloc"] - before["num_host_alloc"] == 1
+    assert storage_bytes == requested_bytes
+
+    # 没有 H2D 或共享视图；删除最后一个引用后，整块内存应进入空闲缓存。
+    tensor_ref = weakref.ref(tensor)
+    del tensor
+    assert tensor_ref() is None
+    released = snapshot()
+    # 此安装版本的无事件 free 路径未扣减 active 计数；用地址复用和实际分配次数验证回收。
+    assert released["allocated_bytes.current"] == allocated["allocated_bytes.current"]
+
+    # 按桶的完整大小再申请，验证与原申请复用同一块内存，不再调用底层申请。
+    reused = torch.empty(block_bytes, dtype=torch.uint8, device="cpu", pin_memory=True, requires_grad=False)
+    reuse_stats = snapshot()
+    same_pointer = reused.data_ptr() == pointer
+    assert same_pointer
+    assert reuse_stats["num_host_alloc"] == allocated["num_host_alloc"]
+    del reused
+    torch._C._host_emptyCache()
+    cleared = snapshot()
+    assert cleared["allocated_bytes.current"] == before["allocated_bytes.current"]
+    assert cleared["num_host_free"] - before["num_host_free"] == 1
+    return {
+        "case": label,
+        "dtype": str(dtype),
+        "numel": numel,
+        "tensor_bytes": requested_bytes,
+        "storage_bytes": storage_bytes,
+        "allocator_block_bytes": block_bytes,
+        "tensor_mib": requested_bytes / MIB,
+        "allocator_block_mib": block_bytes / MIB,
+        "reuse_request_bytes": block_bytes,
+        "reuse_same_pointer": same_pointer,
+        "reuse_new_host_allocations": reuse_stats["num_host_alloc"] - allocated["num_host_alloc"],
+        "stats": {
+            "before": before,
+            "allocated": allocated,
+            "released": released,
+            "reused": reuse_stats,
+            "cache_cleared": cleared,
+        },
+    }
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--output", type=Path, required=True)
+    args = parser.parse_args()
+    # reserve segment 会改变缓存清理行为，本原型验证默认 host allocator 的独立分配路径。
+    for name in ("PYTORCH_ALLOC_CONF", "PYTORCH_CUDA_ALLOC_CONF"):
+        assert not os.environ.get(name), f"Run this standalone prototype with {name} unset"
+    torch.cuda.init()
+    warmup = torch.empty(1, dtype=torch.uint8, pin_memory=True)
+    del warmup
+    cases = [
+        ("64 MiB uint8", 64 * MIB, torch.uint8),
+        ("64 MiB + 1 byte uint8", 64 * MIB + 1, torch.uint8),
+        ("65 MiB uint8", 65 * MIB, torch.uint8),
+        ("128 MiB uint8", 128 * MIB, torch.uint8),
+        ("65 million bytes (decimal)", 65_000_000, torch.uint8),
+        ("65M int32 elements (binary M)", 65 * MIB, torch.int32),
+    ]
+    results = [measure(*case) for case in cases]
+    output = {
+        "torch": torch.__version__,
+        "torch_commit": torch.version.git_version,
+        "cuda_visible_devices": os.environ.get("CUDA_VISIBLE_DEVICES"),
+        "allocator_environment": {
+            name: os.environ.get(name) for name in ("PYTORCH_ALLOC_CONF", "PYTORCH_CUDA_ALLOC_CONF")
+        },
+        "active_counter_note": "This build does not decrement active counters on event-free release; "
+        "reuse is verified by pointer equality and unchanged num_host_alloc.",
+        "results": results,
+    }
+    args.output.parent.mkdir(parents=True, exist_ok=True)
+    args.output.write_text(json.dumps(output, indent=2) + "\n")
+    for result in results:
+        print(
+            f"{result['case']}: tensor={result['tensor_mib']:.6f} MiB, "
+            f"allocator={result['allocator_block_mib']:.0f} MiB, "
+            f"reuse_same_pointer={result['reuse_same_pointer']}, "
+            f"reuse_new_host_allocations={result['reuse_new_host_allocations']}"
+        )
+    print(f"All assertions passed. Results: {args.output}")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/benchmark/benchmark_pinned_empty_h2d.py b/test/benchmark/benchmark_pinned_empty_h2d.py
new file mode 100644
index 0000000000..f0cfc9f892
--- /dev/null
+++ b/test/benchmark/benchmark_pinned_empty_h2d.py
@@ -0,0 +1,248 @@
+"""实测空 H2D 能否像单元素 H2D 一样保护 raw batch 使用的 pinned storage。
+
+示例：CUDA_VISIBLE_DEVICES=6 python test/benchmark/benchmark_pinned_empty_h2d.py
+每种方法和规模在独立子进程运行，避免其他案例留下的 allocator 状态影响观察。
+只操作实验自身申请的内存，不启动模型、不修改服务的生命周期保护实现。
+"""
+
+import argparse
+import csv
+import json
+import os
+from pathlib import Path
+import subprocess
+import sys
+from datetime import datetime, timezone
+
+
+METHODS = ("none", "empty_copy", "empty_to", "one_element_copy", "empty_copy_keep_view")
+
+
+def protect_source(target, source, method):
+    if method == "empty_copy":
+        target[:0].copy_(source[:0], non_blocking=True)
+    elif method == "empty_to":
+        # 与 empty_copy 一样丢弃临时视图和结果；只有 keep_view 对照延长源引用。
+        source[:0].to(target.device, non_blocking=True)
+    elif method == "one_element_copy":
+        target[:1].copy_(source[:1], non_blocking=True)
+    elif method == "empty_copy_keep_view":
+        # 此对照有意保留 CPU 空视图：即使 numel 为 0，它仍持有整个原 storage。
+        # 如果这组成功，不能据此认为空 H2D 向 allocator 登记了流。
+        view = source[:0]
+        target[:0].copy_(view, non_blocking=True)
+        return view
+    return None
+
+
+def submit_raw_batch(runtime, target, source, stream):
+    attrs = runtime.cudaMemcpyAttributes()
+    # 与服务相同的 Stream 模式：真正读取源数据的时间遵循上传流顺序，
+    # raw batch 不会替调用方向 PyTorch host allocator 登记源内存。
+    attrs.srcAccessOrder = runtime.cudaMemcpySrcAccessOrder.cudaMemcpySrcAccessOrderStream
+    (error,) = runtime.cudaMemcpyBatchAsync(
+        [target.data_ptr()],
+        [source.data_ptr()],
+        [source.numel() * source.element_size()],
+        1,
+        [attrs],
+        [0],
+        1,
+        stream.cuda_stream,
+    )
+    if error != runtime.cudaError_t.cudaSuccess:
+        raise RuntimeError(f"cudaMemcpyBatchAsync failed: {error}")
+
+
+def run_trial(torch, runtime, stream, method, token_count, repeat, delay_cycles, attempts):
+    offset = 64
+    allocation_size = token_count + offset * 2
+    expected_value, overwritten_value = repeat + 100, -repeat - 1000
+    owner = torch.full((allocation_size,), expected_value, dtype=torch.int32, pin_memory=True, requires_grad=False)
+    source = owner[offset : offset + token_count]
+    source_storage_ptr = owner.untyped_storage().data_ptr()
+    empty_view = source[:0]
+    empty_metadata = {
+        "numel": empty_view.numel(),
+        "data_ptr": empty_view.data_ptr(),
+        "shares_source_storage": empty_view.untyped_storage().data_ptr() == source_storage_ptr,
+    }
+    del empty_view
+    target = torch.empty(token_count, dtype=torch.int32, device="cuda", requires_grad=False)
+    stream.wait_stream(torch.cuda.default_stream())
+    # 同一流先延时，再提交真实 raw H2D 和待验证的保护操作。
+    # 这样 CPU 能在 GPU 尚未读取源数据时释放原 Tensor、触发 allocator 复用。
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(delay_cycles)
+        submit_raw_batch(runtime, target, source, stream)
+        kept_view = protect_source(target, source, method)
+    del owner, source
+
+    held_allocations = []
+    reused_while_pending = False
+    reuse_seen = False
+    overwrite_finished_while_pending = False
+    pending_checks = 0
+    try:
+        for _ in range(attempts):
+            if stream.query():
+                break
+            replacement = torch.empty(allocation_size, dtype=torch.int32, pin_memory=True, requires_grad=False)
+            # 保留所有新分配块。若旧块被复用并改写，也让它活到上传流结束，
+            # 将风险局限为本实验的数据被覆盖，不让 DMA 访问已经归还驱动的内存。
+            held_allocations.append(replacement)
+            still_pending = not stream.query()
+            pending_checks += int(still_pending)
+            if replacement.untyped_storage().data_ptr() == source_storage_ptr:
+                reuse_seen = True
+                reused_while_pending = still_pending
+                if still_pending:
+                    replacement.fill_(overwritten_value)
+                    overwrite_finished_while_pending = not stream.query()
+                break
+    finally:
+        # 包括异常路径也先等待本实验的上传流，再释放 replacement 和保留的空视图。
+        stream.synchronize()
+    del kept_view
+    result = target.cpu()
+    matching_tokens = int(result.eq(expected_value).sum().item())
+    overwritten_tokens = int(result.eq(overwritten_value).sum().item())
+
+    reused_after_completion = None
+    if not reuse_seen:
+        # 正向对照还应证明上传完成后原分配块可以再次使用，而非被永久保留。
+        reused_after_completion = False
+        for _ in range(attempts):
+            replacement = torch.empty(allocation_size, dtype=torch.int32, pin_memory=True, requires_grad=False)
+            held_allocations.append(replacement)
+            if replacement.untyped_storage().data_ptr() == source_storage_ptr:
+                reused_after_completion = True
+                break
+    return {
+        "repeat": repeat,
+        "method": method,
+        "token_count": token_count,
+        "payload_bytes": token_count * 4,
+        "source_storage_ptr": source_storage_ptr,
+        "empty_view": empty_metadata,
+        "pending_allocation_checks": pending_checks,
+        "source_reused_while_pending": reused_while_pending,
+        "overwrite_finished_while_pending": overwrite_finished_while_pending,
+        "matching_tokens": matching_tokens,
+        "overwritten_tokens": overwritten_tokens,
+        "output_correct": matching_tokens == token_count,
+        "source_reused_after_completion": reused_after_completion,
+    }
+
+
+def run_worker(args):
+    import torch
+    from cuda.bindings import runtime
+
+    torch.set_num_threads(1)
+    torch.cuda.set_device(0)
+    stream = torch.cuda.Stream()
+    # 首次初始化必须在延时试验前完成，不能让导入/驱动初始化耗时掩盖复用窗口。
+    source = torch.ones(256, dtype=torch.int32, pin_memory=True, requires_grad=False)
+    target = torch.empty(256, dtype=torch.int32, device="cuda", requires_grad=False)
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(1)
+        submit_raw_batch(runtime, target, source, stream)
+        protect_source(target, source, args.method)
+    stream.synchronize()
+    del source, target
+    rows = [
+        run_trial(torch, runtime, stream, args.method, args.token_count, repeat, args.delay_cycles, args.attempts)
+        for repeat in range(args.repeats)
+    ]
+    print(json.dumps({"torch_version": torch.__version__, "cuda_version": torch.version.cuda, "trials": rows}))
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--output-dir", type=Path)
+    parser.add_argument("--repeats", type=int, default=10)
+    parser.add_argument("--sizes", type=int, nargs="+", default=[8192, 262144], help="展开后的 int32 token 数")
+    parser.add_argument("--delay-cycles", type=int, default=200_000_000)
+    parser.add_argument("--attempts", type=int, default=64)
+    parser.add_argument("--worker", action="store_true", help=argparse.SUPPRESS)
+    parser.add_argument("--method", choices=METHODS, help=argparse.SUPPRESS)
+    parser.add_argument("--token-count", type=int, help=argparse.SUPPRESS)
+    args = parser.parse_args()
+    # 显式限制到一张卡，防止实验误用用户禁止的 GPU 3 或其他运行中的服务。
+    visible = os.environ.get("CUDA_VISIBLE_DEVICES", "")
+    if not visible or "," in visible or visible == "3":
+        parser.error("Set CUDA_VISIBLE_DEVICES to one idle GPU; GPU 3 must not be used")
+    if args.repeats < 1 or min(args.sizes) < 1 or args.delay_cycles < 1 or args.attempts < 1:
+        parser.error("Repeats, sizes, delay cycles and allocation attempts must be positive")
+    if args.worker:
+        run_worker(args)
+        return
+
+    root = args.output_dir or Path("/tmp") / (
+        "pinned_empty_h2d_" + datetime.now(timezone.utc).strftime("%Y%m%dT%H%M%SZ")
+    )
+    root.mkdir(parents=True, exist_ok=False)
+    all_trials, summaries = [], []
+    version = None
+    for token_count in args.sizes:
+        for method in METHODS:
+            command = [
+                sys.executable,
+                str(Path(__file__).resolve()),
+                "--worker",
+                "--method",
+                method,
+                "--token-count",
+                str(token_count),
+                "--repeats",
+                str(args.repeats),
+                "--delay-cycles",
+                str(args.delay_cycles),
+                "--attempts",
+                str(args.attempts),
+            ]
+            child = subprocess.run(command, capture_output=True, text=True, check=True, timeout=180)
+            (root / f"{token_count}_{method}.json").write_text(child.stdout)
+            if child.stderr:
+                (root / f"{token_count}_{method}.stderr").write_text(child.stderr)
+            data = json.loads(child.stdout)
+            version = {key: data[key] for key in ("torch_version", "cuda_version")}
+            rows = data["trials"]
+            all_trials.extend(rows)
+            summary = {
+                "token_count": token_count,
+                "method": method,
+                "repeats": len(rows),
+                "pending_window_trials": sum(row["pending_allocation_checks"] > 0 for row in rows),
+                "reused_while_pending": sum(row["source_reused_while_pending"] for row in rows),
+                "overwrite_finished_while_pending": sum(row["overwrite_finished_while_pending"] for row in rows),
+                "correct_outputs": sum(row["output_correct"] for row in rows),
+                "fully_overwritten_outputs": sum(row["overwritten_tokens"] == token_count for row in rows),
+                "reusable_after_completion": sum(row["source_reused_after_completion"] is True for row in rows),
+            }
+            summaries.append(summary)
+            print(json.dumps(summary, ensure_ascii=False), flush=True)
+
+    evidence = {
+        **version,
+        "cuda_visible_devices": visible,
+        "delay_cycles": args.delay_cycles,
+        "allocation_attempts": args.attempts,
+        "source_slice_offset_tokens": 64,
+        "raw_copy": "cudaMemcpyBatchAsync with cudaMemcpySrcAccessOrderStream",
+        "isolated_subprocess_per_method_and_size": True,
+        "replacement_allocations_retained_until_stream_completion": True,
+        "summary": summaries,
+        "trials": all_trials,
+    }
+    (root / "result.json").write_text(json.dumps(evidence, ensure_ascii=False, indent=2))
+    with (root / "summary.csv").open("w", newline="") as output:
+        writer = csv.DictWriter(output, fieldnames=list(summaries[0]))
+        writer.writeheader()
+        writer.writerows(summaries)
+    print(f"Experiment results: {root}", flush=True)
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/benchmark/benchmark_req_token_index_alloc.py b/test/benchmark/benchmark_req_token_index_alloc.py
new file mode 100644
index 0000000000..f7ccc0cb0b
--- /dev/null
+++ b/test/benchmark/benchmark_req_token_index_alloc.py
@@ -0,0 +1,199 @@
+"""Measure shared token index pinned allocation, excluding data writes and H2D.
+
+Example:
+    CUDA_VISIBLE_DEVICES=4 PYTHONPATH=. python test/benchmark/benchmark_req_token_index_alloc.py \
+        --cpu 72 --output-dir /tmp/req_token_index_alloc
+"""
+
+import argparse
+import csv
+import gc
+import hashlib
+import json
+import math
+import os
+from pathlib import Path
+import platform
+import statistics
+import time
+
+import torch
+
+from lightllm.server.router.model_infer.infer_req_ext1 import ReqTokenIndex
+
+
+SLAB_TOKENS = 1024 * 1024
+CLOCK = time.perf_counter_ns
+
+
+def allocate_direct(token_num):
+    return torch.empty(token_num, dtype=torch.int32, device="cpu", pin_memory=True, requires_grad=False)
+
+
+def summarize(samples):
+    ordered = sorted(samples)
+    return {
+        "count": len(samples),
+        "min_us": min(samples),
+        "max_us": max(samples),
+        "mean_us": statistics.mean(samples),
+        "median_us": statistics.median(samples),
+        "p95_us": ordered[int(0.95 * (len(ordered) - 1))],
+    }
+
+
+def measure_cached(mode, token_num, repeats, backing):
+    allocate = allocate_direct if mode == "direct_cached" else ReqTokenIndex._alloc_return_pinned_tensor
+    samples = []
+    for _ in range(repeats):
+        if mode == "slab_slice" and 0 < token_num <= SLAB_TOKENS:
+            # 仅测切片路径：当前没有存活的返回值或异步传输，允许在计时外重置测试 backing。
+            # 生产接口始终向后切分；slab_cached 模式保留其真实换块行为。
+            buffer, offset = getattr(allocate, "return_buffer_state", (None, 0))
+            if buffer is None or offset + token_num > SLAB_TOKENS:
+                allocate.return_buffer_state = (backing, 0)
+        start = CLOCK()
+        tensor = allocate(token_num)
+        samples.append((CLOCK() - start) / 1000)
+        del tensor
+    return samples
+
+
+def measure_growth(mode, token_num, live_count):
+    # 仅清理 HOST 空闲缓存；用于构造缓存未命中，不用于生产服务。
+    ReqTokenIndex._alloc_return_pinned_tensor.return_buffer_state = (None, 0)
+    gc.collect()
+    torch.cuda.synchronize()
+    torch._C._host_emptyCache()
+    before = torch.cuda.memory.host_memory_stats()
+    allocate = allocate_direct if mode == "direct" else ReqTokenIndex._alloc_return_pinned_tensor
+    samples, held = [], []
+    for _ in range(live_count):
+        start = CLOCK()
+        tensor = allocate(token_num)
+        samples.append((CLOCK() - start) / 1000)
+        held.append(tensor)
+    after = torch.cuda.memory.host_memory_stats()
+    host_calls = after["num_host_alloc"] - before["num_host_alloc"]
+    expected_calls = live_count
+    if mode == "slab" and token_num <= SLAB_TOKENS:
+        expected_calls = math.ceil(live_count / (SLAB_TOKENS // token_num))
+    assert host_calls == expected_calls, (mode, token_num, host_calls, expected_calls)
+    assert all(tensor.is_pinned() and tensor.numel() == token_num for tensor in held)
+    return {
+        "samples_us": samples,
+        "summary": summarize(samples),
+        "host_allocations": host_calls,
+        "allocated_bytes": after["allocated_bytes.current"] - before["allocated_bytes.current"],
+        "host_alloc_time_us": after["host_alloc_time.total"] - before["host_alloc_time.total"],
+    }
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--output-dir", type=Path, required=True)
+    parser.add_argument("--cpu", type=int)
+    parser.add_argument("--rounds", type=int, default=10)
+    parser.add_argument("--repeats", type=int, default=2000)
+    args = parser.parse_args()
+    if args.cpu is not None:
+        os.sched_setaffinity(0, {args.cpu})
+    args.output_dir.mkdir(parents=True, exist_ok=True)
+    torch.set_num_threads(1)
+    torch.cuda.init()
+    seed = allocate_direct(1)
+    del seed
+    torch._C._host_emptyCache()
+    start = CLOCK()
+    seed = ReqTokenIndex._alloc_return_pinned_tensor(1)
+    initial_slab_us = (CLOCK() - start) / 1000
+    backing = seed._base
+    assert backing.numel() == SLAB_TOKENS and backing.is_pinned()
+    del seed
+    gc.collect()
+    gc.disable()
+
+    sizes = (0, 128, 1024, 2048, 8192, 32768, 262144, SLAB_TOKENS // 2, SLAB_TOKENS, 2 * SLAB_TOKENS)
+    modes = ("direct_cached", "slab_slice", "slab_cached")
+    result = {
+        "metadata": {
+            "torch": torch.__version__,
+            "torch_git": torch.version.git_version,
+            "python": platform.python_version(),
+            "gpu": torch.cuda.get_device_name(0),
+            "cuda_visible_devices": os.environ.get("CUDA_VISIBLE_DEVICES"),
+            "cpu_affinity": sorted(os.sched_getaffinity(0)),
+            "torch_threads": torch.get_num_threads(),
+            "allocator_env": {name: os.environ.get(name) for name in ("PYTORCH_ALLOC_CONF", "PYTORCH_CUDA_ALLOC_CONF")},
+            "source_sha256": hashlib.sha256(
+                Path(ReqTokenIndex._alloc_return_pinned_tensor.__code__.co_filename).read_bytes()
+            ).hexdigest(),
+            "slab_tokens": SLAB_TOKENS,
+            "slab_bytes": SLAB_TOKENS * 4,
+            "scope": "CPU allocation wall time; excludes release, fill, H2D, imports, CUDA init and GC.",
+            "cached_modes": {
+                "direct_cached": "torch.empty with warm HOST size buckets",
+                "slab_slice": "slab slicing; backing reset outside timer when exhausted, no live or in-flight data",
+                "slab_cached": "actual shared allocator, including warm HOST-cache slab replacements",
+            },
+            "growth": "Retain every returned tensor; clear unused HOST cache before each cohort; include fresh slabs.",
+        },
+        "initial_slab_allocation_us": initial_slab_us,
+        "cached": {str(size): {mode: {"rounds_us": []} for mode in modes} for size in sizes},
+        "growth": {},
+    }
+
+    # 预热所有直接申请桶，再准备两个可交替复用的 slab；验证计时区间没有调用 host 分配 API。
+    for token_num in sizes:
+        seed = allocate_direct(token_num)
+        del seed
+    warm = [allocate_direct(SLAB_TOKENS) for _ in range(2)]
+    del warm
+    before = torch.cuda.memory.host_memory_stats()["num_host_alloc"]
+    for repeat in range(args.rounds):
+        for token_num in sizes if repeat % 2 == 0 else reversed(sizes):
+            for mode in modes if repeat % 2 == 0 else reversed(modes):
+                samples = measure_cached(mode, token_num, args.repeats, backing)
+                result["cached"][str(token_num)][mode]["rounds_us"].append(samples)
+        print(f"Cached round {repeat + 1}/{args.rounds} complete", flush=True)
+    result["cached_host_allocations"] = torch.cuda.memory.host_memory_stats()["num_host_alloc"] - before
+    assert result["cached_host_allocations"] == 0
+    for data in result["cached"].values():
+        for measured in data.values():
+            rounds = measured["rounds_us"]
+            measured["summary"] = summarize([sample for samples in rounds for sample in samples])
+            measured["round_medians_us"] = [statistics.median(samples) for samples in rounds]
+
+    for token_num in (1024, 8192, 32768, 262144, 1048576, 2 * SLAB_TOKENS):
+        live_count = min(256, max(16, 16 * 1024 * 1024 // token_num))
+        data = {mode: {"rounds": []} for mode in ("direct", "slab")}
+        data["live_count"] = live_count
+        for repeat in range(args.rounds):
+            for mode in ("direct", "slab") if repeat % 2 == 0 else ("slab", "direct"):
+                data[mode]["rounds"].append(measure_growth(mode, token_num, live_count))
+        for mode in ("direct", "slab"):
+            rounds = data[mode]["rounds"]
+            data[mode]["summary"] = summarize([sample for batch in rounds for sample in batch["samples_us"]])
+            data[mode]["round_mean_us"] = [batch["summary"]["mean_us"] for batch in rounds]
+        result["growth"][str(token_num)] = data
+        print(f"Uncached growth for {token_num} tokens x {live_count} complete", flush=True)
+
+    ReqTokenIndex._alloc_return_pinned_tensor.return_buffer_state = (None, 0)
+    gc.enable()
+    (args.output_dir / "result.json").write_text(json.dumps(result, indent=2) + "\n")
+    with (args.output_dir / "timings.csv").open("w", newline="") as output:
+        writer = csv.writer(output)
+        fields = ("count", "min_us", "max_us", "mean_us", "median_us", "p95_us")
+        writer.writerow(("scenario", "tokens", "bytes", "mode", *fields))
+        for scenario in ("cached", "growth"):
+            for tokens, data in result[scenario].items():
+                for mode, measured in data.items():
+                    if mode != "live_count":
+                        writer.writerow(
+                            (scenario, tokens, int(tokens) * 4, mode, *[measured["summary"][f] for f in fields])
+                        )
+    print(f"Results: {args.output_dir}", flush=True)
+
+
+if __name__ == "__main__":
+    main()
diff --git a/test/benchmark/static_inference/static_benchmark.py b/test/benchmark/static_inference/static_benchmark.py
index 279840d5af..ce19f3bc61 100644
--- a/test/benchmark/static_inference/static_benchmark.py
+++ b/test/benchmark/static_inference/static_benchmark.py
@@ -28,6 +28,7 @@
     sys.path.append(str(REPO_ROOT))
 
 from lightllm.common.basemodel.batch_objs import ModelInput, ModelMtpOutputCollector, ModelOutput
+from lightllm.server.router.model_infer.infer_req_ext1 import ReqTokenIndex, create_req_token_index
 from lightllm.common.basemodel.triton_kernel.gen_mtp_prefill_params import gen_mtp_new_input_ids
 from lightllm.models import get_draft_model_class, get_model
 from lightllm.server.api_cli import make_argument_parser
@@ -158,6 +159,7 @@ def __init__(
         self.draft_models = draft_models
         self.token_source = token_source
         self._hold_kv_len_by_req: Dict[int, int] = {}
+        self._req_token_indexes: Dict[int, ReqTokenIndex] = {}
 
     def _case_iters(self, warmup: bool) -> int:
         return self.args.warmup_iters if warmup else self.args.bench_iters
@@ -558,28 +560,43 @@ def _make_decode_input(
 
     def _ensure_req_kv_capacity(self, req_idx: torch.Tensor, target_kv_lens: torch.Tensor) -> None:
         """按请求预分配完整 KV 页，并将物理索引写入请求索引表。"""
-        req_ids = req_idx.detach().cpu().tolist()
+        req_indexes = req_idx.detach().cpu().tolist()
         target_lens = target_kv_lens.detach().cpu().tolist()
-        assert len(req_ids) == len(target_lens)
+        assert len(req_indexes) == len(target_lens)
 
-        # MTP decode 会为同一个请求展开多行，只需按其中最大的逻辑长度扩容一次。
-        max_target_len_by_req = {}
-        for req_id, target_len in zip(req_ids, target_lens):
-            req_id = int(req_id)
-            max_target_len_by_req[req_id] = max(max_target_len_by_req.get(req_id, 0), int(target_len))
+        # 按 ReqManager 槽位合并 MTP 展开的多行，只按最大逻辑 token 长度扩容一次。
+        max_target_len_by_req_index = {}
+        for req_index, target_len in zip(req_indexes, target_lens):
+            req_index = int(req_index)
+            max_target_len_by_req_index[req_index] = max(max_target_len_by_req_index.get(req_index, 0), int(target_len))
 
         mem_manager = self.model.req_manager.mem_manager
         page_size = mem_manager.page_size
-        for req_id, target_len in max_target_len_by_req.items():
+        req_manager = self.model.req_manager
+        for req_index, target_len in max_target_len_by_req_index.items():
             new_hold_kv_len = align_up(target_len, page_size)
-            old_hold_kv_len = self._hold_kv_len_by_req.get(req_id, 0)
+            old_hold_kv_len = self._hold_kv_len_by_req.get(req_index, 0)
             if new_hold_kv_len <= old_hold_kv_len:
                 continue
 
             alloc_token_num = new_hold_kv_len - old_hold_kv_len
             mem_indexes = mem_manager.alloc(alloc_token_num)
-            self.model.req_manager.req_to_token_indexs[req_id, old_hold_kv_len:new_hold_kv_len].copy_(mem_indexes)
-            self._hold_kv_len_by_req[req_id] = new_hold_kv_len
+            if req_index not in self._req_token_indexes:
+                # 静态 benchmark 不进入真实 InferReq 的共享内存生命周期，轻量对象
+                # 提供请求槽位、启动参数和输入长度，复用在线请求的索引工厂。
+                benchmark_req = SimpleNamespace(
+                    req_idx=req_index,
+                    req_manager=req_manager,
+                    args=self.args,
+                    shm_req=SimpleNamespace(input_len=target_len),
+                )
+                self._req_token_indexes[req_index] = create_req_token_index(
+                    benchmark_req, index_stream=torch.cuda.current_stream()
+                )
+            token_index = self._req_token_indexes[req_index]
+            token_index.append_token_indexes(mem_indexes)
+            token_index.get_gpu_token_indexes(old_hold_kv_len, new_hold_kv_len)
+            self._hold_kv_len_by_req[req_index] = new_hold_kv_len
 
     def _forward_prefill_input(self, model_input: ModelInput, allow_overlap: bool) -> ModelOutput:
         if allow_overlap and self.args.enable_prefill_microbatch_overlap and model_input.batch_size > 1:
@@ -732,6 +749,7 @@ def _reset_model_cache(self):
         self.model.mem_manager.free_all()
         self.model.req_manager.free_all()
         self._hold_kv_len_by_req.clear()
+        self._req_token_indexes.clear()
         torch.cuda.synchronize()
         torch.cuda.empty_cache()
 
diff --git a/unit_tests/common/basemodel/test_model_input.py b/unit_tests/common/basemodel/test_model_input.py
index 6672cccc31..68e02717f3 100644
--- a/unit_tests/common/basemodel/test_model_input.py
+++ b/unit_tests/common/basemodel/test_model_input.py
@@ -104,6 +104,7 @@ def test_prefill_rejects_position_delta():
         model_input.check_input()
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
 def test_padded_prefill_adds_non_decode_request_marker():
     model_input = ModelInput(
         batch_size=1,
@@ -127,6 +128,7 @@ def test_padded_prefill_adds_non_decode_request_marker():
         req_manager=SimpleNamespace(HOLD_REQUEST_ID=-1),
     )
 
+    model_input.to_cuda()
     padded_input = TpPartBaseModel._create_padded_prefill_model_input(
         model,
         model_input=model_input,
@@ -137,6 +139,7 @@ def test_padded_prefill_adds_non_decode_request_marker():
     assert padded_input.b_is_decode_req.tolist() == [True, False]
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
 def test_padded_prefill_builds_internal_request_for_empty_input():
     model_input = ModelInput(
         batch_size=0,
@@ -161,6 +164,7 @@ def test_padded_prefill_builds_internal_request_for_empty_input():
         req_manager=SimpleNamespace(HOLD_REQUEST_ID=88),
     )
 
+    model_input.to_cuda()
     padded_input = TpPartBaseModel._create_padded_prefill_model_input(
         model,
         model_input=model_input,
@@ -176,6 +180,7 @@ def test_padded_prefill_builds_internal_request_for_empty_input():
     assert torch.equal(padded_input.mtp_draft_input_hiddens, torch.zeros((1, 4)))
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
 def test_padded_decode_builds_internal_request_from_empty_token_tensor():
     model_input = ModelInput(
         batch_size=0,
@@ -197,6 +202,7 @@ def test_padded_decode_builds_internal_request_from_empty_token_tensor():
         req_manager=SimpleNamespace(HOLD_REQUEST_ID=88),
     )
 
+    model_input.to_cuda()
     padded_input = TpPartBaseModel._create_padded_decode_model_input(
         model,
         model_input=model_input,
diff --git a/unit_tests/common/basemodel/test_model_output.py b/unit_tests/common/basemodel/test_model_output.py
index 3f6664d031..388378bb6f 100644
--- a/unit_tests/common/basemodel/test_model_output.py
+++ b/unit_tests/common/basemodel/test_model_output.py
@@ -77,7 +77,7 @@ def test_prefill_unpad_restores_empty_output():
 
 
 def _create_empty_decode_input():
-    return ModelInput(
+    model_input = ModelInput(
         batch_size=0,
         total_token_num=0,
         max_q_seq_len=1,
@@ -92,9 +92,12 @@ def _create_empty_decode_input():
         is_prefill=False,
         multimodal_params=[],
     )
+    model_input.to_cuda()
+    return model_input
 
 
 @torch.no_grad()
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
 def test_decode_pads_only_once_after_selecting_execution_path(monkeypatch):
     execution_configs = (
         # eager 普通模式：空 batch 只补一个 dummy request。
@@ -131,7 +134,7 @@ def pad_once(model_input, new_batch_size):
 
         graph_flags_at_att_init = []
 
-        def create_infer_state(model_input):
+        def create_infer_state(model_input, microbatch_index=0):
             infer_state = SimpleNamespace(
                 b_req_idx=model_input.b_req_idx,
                 b_seq_len=model_input.b_seq_len,
@@ -164,6 +167,7 @@ def create_infer_state(model_input):
         assert output.logits.shape == (0, 4)
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
 @pytest.mark.parametrize("overlap", [False, True])
 @pytest.mark.parametrize("dynamic", [False, True])
 @pytest.mark.parametrize("graph_mode", ["eager", "capture", "replay"])
@@ -218,6 +222,7 @@ def forward(state):
         model_input.b_req_idx = torch.tensor([req for req, _ in request_rows], dtype=torch.int32)
         model_input.b_mtp_index = torch.tensor([step for _, step in request_rows], dtype=torch.int32)
         model_input.b_seq_len = 10 + model_input.b_mtp_index
+    model_input.to_cuda()
     expected = torch.stack((model_input.b_req_idx, model_input.b_seq_len), dim=1)
     if overlap:
         output, empty_output = model._microbatch_overlap_decode_cuda(model_input, _create_empty_decode_input())
@@ -229,3 +234,59 @@ def forward(state):
     assert seen_sizes == [expected_size] * (2 if overlap else 1)
     torch.testing.assert_close(output.logits, expected)
     torch.testing.assert_close(torch.stack((model_input.b_req_idx, model_input.b_seq_len), dim=1), expected)
+
+
+@pytest.mark.parametrize("overlap", [False, True])
+def test_prefill_preparation_preserves_index_event_and_initialization_order(monkeypatch, overlap):
+    model = TpPartBaseModel.__new__(TpPartBaseModel)
+    model.args = SimpleNamespace(enable_prefill_decode_mixed=False, enable_tpsp_mix_mode=True)
+    model.tp_world_size_ = 2
+    model.prefill_graph = None
+    calls = []
+
+    class ReadyEvent:
+        def record(self):
+            calls.append("ready")
+
+    monkeypatch.setattr(basemodel.torch.cuda, "Event", ReadyEvent)
+    monkeypatch.setattr(basemodel.dist_group_manager, "clear_deepep_buffer", lambda: None)
+
+    def pad_input(model_input, handle_token_num):
+        assert handle_token_num == 4
+        return model_input
+
+    def create_state(model_input, microbatch_index=0):
+        calls.append(("select_indexes", microbatch_index))
+        return SimpleNamespace(
+            init_some_extra_state=lambda _: calls.append(("extra", microbatch_index)),
+            init_att_state=lambda: calls.append(("attention", microbatch_index)),
+        )
+
+    def forward(*args, **kwargs):
+        calls.append("forward")
+        return ModelOutput(logits=torch.ones(1, 4))
+
+    model._create_padded_prefill_model_input = pad_input
+    model._create_inferstate = create_state
+    model._context_forward = forward
+    model._overlap_tpsp_context_forward = lambda *args, **kwargs: (forward(), forward())
+    model_input = SimpleNamespace(is_prefill=True, input_ids=torch.arange(3), batch_size=1)
+    if overlap:
+        outputs = model._microbatch_overlap_prefill_cuda(model_input, model_input)
+        expected = [
+            ("select_indexes", 0),
+            ("extra", 0),
+            ("attention", 0),
+            ("select_indexes", 1),
+            ("extra", 1),
+            ("attention", 1),
+            "ready",
+            "forward",
+            "forward",
+        ]
+        assert outputs[0].prefill_mem_indexes_ready_event is outputs[1].prefill_mem_indexes_ready_event
+    else:
+        outputs = (model._prefill(model_input),)
+        expected = [("select_indexes", 0), "ready", ("extra", 0), ("attention", 0), "forward"]
+    assert calls == expected
+    assert all(isinstance(output.prefill_mem_indexes_ready_event, ReadyEvent) for output in outputs)
diff --git a/unit_tests/common/basemodel/triton_kernel/test_pad_decode_inputs.py b/unit_tests/common/basemodel/triton_kernel/test_pad_decode_inputs.py
new file mode 100644
index 0000000000..15302ef9ad
--- /dev/null
+++ b/unit_tests/common/basemodel/triton_kernel/test_pad_decode_inputs.py
@@ -0,0 +1,92 @@
+import pytest
+import torch
+import torch.nn.functional as F
+
+from lightllm.common.basemodel.triton_kernel.pad_decode_inputs import pad_decode_inputs
+
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
+
+
+def _make_inputs(batch_size, has_position_delta):
+    return dict(
+        input_ids=torch.arange(batch_size, dtype=torch.int64, device="cuda") + 2 ** 40,
+        b_req_idx=torch.arange(batch_size, dtype=torch.int32, device="cuda") + 10,
+        b_mtp_index=torch.arange(batch_size, dtype=torch.int32, device="cuda") % 4,
+        b_seq_len=torch.arange(batch_size, dtype=torch.int32, device="cuda") + 5,
+        b_position_delta=(
+            torch.arange(batch_size, dtype=torch.int32, device="cuda") - 9 if has_position_delta else None
+        ),
+        b_shared_seq_len=torch.arange(batch_size, dtype=torch.int32, device="cuda") % 3,
+        b_shared_radix_node_id=torch.arange(batch_size, dtype=torch.int64, device="cuda") + 2 ** 40,
+    )
+
+
+def _reference(inputs, padded_batch_size, hold_request_id):
+    return (
+        F.pad(inputs["input_ids"], (0, padded_batch_size), value=1),
+        F.pad(inputs["b_req_idx"], (0, padded_batch_size), value=hold_request_id),
+        F.pad(inputs["b_mtp_index"], (0, padded_batch_size), value=0),
+        F.pad(inputs["b_seq_len"], (0, padded_batch_size), value=2),
+        (
+            F.pad(inputs["b_position_delta"], (0, padded_batch_size), value=0)
+            if inputs["b_position_delta"] is not None
+            else None
+        ),
+        F.pad(inputs["b_shared_seq_len"], (0, padded_batch_size), value=0),
+        F.pad(inputs["b_shared_radix_node_id"], (0, padded_batch_size), value=-1),
+    )
+
+
+@pytest.mark.parametrize("hold_request_id", [-1, 88])
+@pytest.mark.parametrize("has_position_delta", [False, True])
+@pytest.mark.parametrize(
+    "batch_size,new_batch_size",
+    [(0, 1), (0, 257), (1, 2), (3, 8), (255, 256), (256, 257), (257, 1025), (1025, 4096)],
+)
+def test_pad_decode_inputs_matches_torch(batch_size, new_batch_size, has_position_delta, hold_request_id):
+    inputs = _make_inputs(batch_size, has_position_delta)
+    originals = {name: value.clone() for name, value in inputs.items() if value is not None}
+    expected = _reference(inputs, new_batch_size - batch_size, hold_request_id)
+
+    actual = pad_decode_inputs(**inputs, new_batch_size=new_batch_size, hold_request_id=hold_request_id)
+
+    for output, reference in zip(actual, expected):
+        if reference is None:
+            assert output is None
+        else:
+            assert output.dtype == reference.dtype
+            assert output.device == reference.device
+            assert output.shape == reference.shape
+            assert torch.equal(output, reference)
+    for name, original in originals.items():
+        assert torch.equal(inputs[name], original)
+
+
+@pytest.mark.parametrize("has_position_delta", [False, True])
+def test_pad_decode_inputs_cuda_graph_replays_updated_inputs(has_position_delta):
+    inputs = _make_inputs(batch_size=3, has_position_delta=has_position_delta)
+    kwargs = dict(**inputs, new_batch_size=513, hold_request_id=88)
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        for _ in range(3):
+            pad_decode_inputs(**kwargs)
+    torch.cuda.current_stream().wait_stream(stream)
+
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph, stream=stream):
+        outputs = pad_decode_inputs(**kwargs)
+
+    inputs["input_ids"].add_(7)
+    inputs["b_req_idx"].add_(2)
+    inputs["b_shared_radix_node_id"].add_(11)
+    if has_position_delta:
+        inputs["b_position_delta"].sub_(3)
+    graph.replay()
+
+    for output, reference in zip(outputs, _reference(inputs, padded_batch_size=510, hold_request_id=88)):
+        if reference is None:
+            assert output is None
+        else:
+            assert torch.equal(output, reference)
diff --git a/unit_tests/common/basemodel/triton_kernel/test_pad_prefill_inputs.py b/unit_tests/common/basemodel/triton_kernel/test_pad_prefill_inputs.py
new file mode 100644
index 0000000000..17ab85aee1
--- /dev/null
+++ b/unit_tests/common/basemodel/triton_kernel/test_pad_prefill_inputs.py
@@ -0,0 +1,86 @@
+import pytest
+import torch
+import torch.nn.functional as F
+
+from lightllm.common.basemodel.triton_kernel.pad_prefill_inputs import pad_prefill_inputs
+
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA is required")
+
+
+def _make_inputs(batch_size, handle_token_num):
+    return dict(
+        input_ids=torch.arange(handle_token_num, dtype=torch.int64, device="cuda") + 2 ** 40,
+        b_req_idx=torch.arange(batch_size, dtype=torch.int32, device="cuda") + 10,
+        b_mtp_index=torch.arange(batch_size, dtype=torch.int32, device="cuda") % 4,
+        b_seq_len=torch.arange(batch_size, dtype=torch.int32, device="cuda") + 5,
+        b_ready_cache_len=torch.arange(batch_size, dtype=torch.int32, device="cuda") % 3,
+        b_is_decode_req=torch.arange(batch_size, device="cuda") % 2 == 0,
+    )
+
+
+def _reference(inputs, padded_token_num, hold_request_id):
+    return (
+        F.pad(inputs["input_ids"], (0, padded_token_num), value=1),
+        F.pad(inputs["b_req_idx"], (0, 1), value=hold_request_id),
+        F.pad(inputs["b_mtp_index"], (0, 1), value=0),
+        F.pad(inputs["b_seq_len"], (0, 1), value=padded_token_num),
+        F.pad(inputs["b_ready_cache_len"], (0, 1), value=0),
+        F.pad(inputs["b_is_decode_req"], (0, 1), value=False),
+    )
+
+
+@pytest.mark.parametrize("hold_request_id", [-1, 88])
+@pytest.mark.parametrize(
+    "batch_size,handle_token_num,padded_token_num",
+    [
+        (0, 0, 1),
+        (0, 0, 129),
+        (1, 1, 3),
+        (3, 17, 1),
+        (257, 2051, 11),
+        (1025, 4099, 257),
+        (3, 65537, 7),
+    ],
+)
+def test_pad_prefill_inputs_matches_torch(batch_size, handle_token_num, padded_token_num, hold_request_id):
+    inputs = _make_inputs(batch_size, handle_token_num)
+    originals = {name: value.clone() for name, value in inputs.items()}
+    expected = _reference(inputs, padded_token_num, hold_request_id)
+
+    actual = pad_prefill_inputs(
+        **inputs,
+        new_handle_token_num=handle_token_num + padded_token_num,
+        hold_request_id=hold_request_id,
+    )
+
+    for output, reference in zip(actual, expected):
+        assert output.dtype == reference.dtype
+        assert output.device == reference.device
+        assert output.shape == reference.shape
+        assert torch.equal(output, reference)
+    for name, original in originals.items():
+        assert torch.equal(inputs[name], original)
+
+
+def test_pad_prefill_inputs_cuda_graph_replays_updated_inputs():
+    inputs = _make_inputs(batch_size=3, handle_token_num=17)
+    kwargs = dict(**inputs, new_handle_token_num=24, hold_request_id=88)
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        for _ in range(3):
+            pad_prefill_inputs(**kwargs)
+    torch.cuda.current_stream().wait_stream(stream)
+
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph, stream=stream):
+        outputs = pad_prefill_inputs(**kwargs)
+
+    inputs["input_ids"].add_(7)
+    inputs["b_req_idx"].add_(2)
+    inputs["b_is_decode_req"].logical_not_()
+    graph.replay()
+
+    for output, reference in zip(outputs, _reference(inputs, padded_token_num=7, hold_request_id=88)):
+        assert torch.equal(output, reference)
diff --git a/unit_tests/common/cpu_index_helpers.py b/unit_tests/common/cpu_index_helpers.py
new file mode 100644
index 0000000000..36ac5f0b60
--- /dev/null
+++ b/unit_tests/common/cpu_index_helpers.py
@@ -0,0 +1,69 @@
+"""在轻量生命周期测试中复用生产 ReqTokenIndex；矩阵仅作为测试数据源。"""
+from types import SimpleNamespace
+
+import torch
+
+
+def make_cpu_index_manager(token_table, page_size=1, **attributes):
+    from lightllm.common.req_manager.base import ReqManager
+    from lightllm.server.router.model_infer.infer_req_ext1 import PinMemCacheReqTokenIndex
+
+    class CpuOnlyTokenIndex(PinMemCacheReqTokenIndex):
+        def _reserve(self, end):
+            # CPU 生命周期测试只模拟缓冲区分配，追加、快照和 reset 仍复用生产实现。
+            # 按测试请求表容量申请普通内存，不依赖 CUDA 或 pinned allocator。
+            if self._cpu_buffer is not None and self._cpu_buffer.numel() >= end:
+                return
+            buffer = torch.empty_like(self.gpu_indexes, device="cpu", requires_grad=False)
+            if self.token_index_len:
+                buffer[: self.token_index_len].copy_(self._cpu_buffer[: self.token_index_len])
+            self._cpu_buffer = buffer
+
+    class TestReqManager(ReqManager):
+        def get_req_token_index(self, req_idx):
+            if req_idx not in self._req_token_indexes:
+                req = SimpleNamespace(
+                    req_idx=req_idx,
+                    req_manager=self,
+                    args=SimpleNamespace(req_token_index_mode="pin_mem_cache"),
+                    shm_req=SimpleNamespace(input_len=self.token_indexes_cpu.shape[1]),
+                )
+                if self.req_to_token_indexs.is_cuda:
+                    helper = PinMemCacheReqTokenIndex(
+                        req, index_stream=torch.cuda.default_stream(self.req_to_token_indexs.device)
+                    )
+                else:
+                    # CPU 测试不调用 GPU 接口，无需构造索引流。
+                    helper = CpuOnlyTokenIndex(req, index_stream=None)
+                # 既有底层 CUDA 测试直接布置 source；生产实现按请求独立申请。
+                helper._cpu_buffer = self.token_indexes_cpu[req_idx]
+                self._req_token_indexes[req_idx] = helper
+            return self._req_token_indexes[req_idx]
+
+        def get_cpu_token_indexes(self, req_idx, start, end):
+            return self.get_req_token_index(req_idx).get_cpu_token_indexes(start, end)
+
+        def append_token_indexes(self, req_idx, indexes):
+            self.get_req_token_index(req_idx).append_token_indexes(indexes)
+
+    manager = TestReqManager.__new__(TestReqManager)
+    manager.page_size = page_size
+    manager.token_indexes_cpu = token_table
+    manager.req_to_token_indexs = torch.full_like(token_table, -1)
+    manager._req_token_indexes = {}
+    for name, value in attributes.items():
+        setattr(manager, name, value)
+    return manager
+
+
+def bind_cpu_indexes(req, manager, gpu_index_submitted_len=0, token_index_len=None):
+    req.req_manager = manager
+    if not hasattr(req, "args"):
+        req.args = SimpleNamespace(req_token_index_mode="pin_mem_cache")
+    req.token_index = manager.get_req_token_index(req.req_idx)
+    # 矩阵只提供容量和预置数据；追加位置取请求实际持有的长度，不能使用整行容量。
+    if token_index_len is None:
+        token_index_len = req.hold_kv_len if hasattr(req, "hold_kv_len") else 0
+    req.token_index.token_index_len = token_index_len
+    req.token_index.gpu_index_submitted_len = gpu_index_submitted_len
+    return req
diff --git a/unit_tests/common/test_infer_batch_cache_release.py b/unit_tests/common/test_infer_batch_cache_release.py
new file mode 100644
index 0000000000..7b0f6d97f7
--- /dev/null
+++ b/unit_tests/common/test_infer_batch_cache_release.py
@@ -0,0 +1,591 @@
+"""Exercise request release with real radix trees, KV allocation and checkpoint slot pools."""
+
+from types import SimpleNamespace
+from unittest.mock import Mock
+
+import numpy as np
+import pytest
+import torch
+from sortedcontainers import SortedDict
+
+from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+from lightllm.common.kv_cache_mem_manager import allocator as allocator_impl
+from lightllm.common.state_cache_manager.base import StateCacheManager
+from lightllm.server.multi_level_kv_cache import CacheTier
+from lightllm.server.router.dynamic_prompt import hybrid_att_radix_cache, radix_cache
+from lightllm.server.router.model_infer import infer_batch
+from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
+
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="requires pinned memory and CUDA index tables")
+
+
+class _LocalSharedArray:
+    def __init__(self, name, shape, dtype):
+        self.arr = np.zeros(shape, dtype=dtype)
+
+
+class _StateSlots(StateCacheManager):
+    def __init__(self):
+        super().__init__(size=16)
+        self.states = torch.empty((self.size, 1))
+
+    def get_state_cache(self, buffer_idx):
+        return self.states[buffer_idx]
+
+
+@pytest.fixture
+def make_release_case(monkeypatch):
+    # Only replace interprocess counters; allocation, tree operations and slot accounting are production code.
+    monkeypatch.setattr(radix_cache, "SharedArray", _LocalSharedArray)
+    monkeypatch.setattr(hybrid_att_radix_cache, "SharedArray", _LocalSharedArray)
+    monkeypatch.setattr(allocator_impl, "SharedInt", lambda name: SimpleNamespace(set_value=lambda value: None))
+    monkeypatch.setattr(allocator_impl, "get_current_rank_in_node", lambda: 0)
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+
+    def make_case(attention, mode, page_size, cache_case):
+        hybrid = attention != "full"
+        partial_prefill = attention == "hybrid_partial"
+        allocator = allocator_impl.KvCacheAllocator(64, page_size)
+        small_slots, big_slots = _StateSlots(), _StateSlots()
+        memory = SimpleNamespace(
+            page_size=page_size,
+            allocator=allocator,
+            alloc=allocator.alloc,
+            free=allocator.free,
+            big_page_buffers=big_slots,
+        )
+        if hybrid:
+            cache = hybrid_att_radix_cache.HybridAttPagedRadixCache(
+                64,
+                0,
+                hash_page_size=4,
+                big_page_num=2,
+                kv_cache_mem_manager=memory,
+                small_page_buffers=small_slots,
+            )
+        else:
+            cache = radix_cache.RadixCache(64, 0, mem_manager=memory, page_size=page_size)
+        key = torch.arange(32, dtype=torch.int64)
+        seeded_indexes = None
+        if cache_case != "uncached":
+            seeded_indexes = allocator.alloc(8).clone()
+            if hybrid:
+                cache.insert(
+                    key[:8],
+                    seeded_indexes,
+                    block_hashs=[1, 2],
+                    block_state_idxs=[None, None],
+                    len_to_big_page_id=SortedDict({8: big_slots.alloc_one_state_cache()}),
+                )
+            else:
+                cache.insert(key[:8], seeded_indexes)
+
+        shared_node, shared_indexes = None, None
+        if cache_case == "shared":
+            if hybrid:
+                shared_node, shared_len, shared_indexes = cache.match_prefix(
+                    key[:8], block_hashs=[1, 2], update_refs=True
+                )
+            else:
+                shared_node, shared_len, shared_indexes = cache.match_prefix(key[:8], update_refs=True)
+            assert shared_len == cache.get_refed_tokens_num() == 8
+        shared_len = 0 if shared_node is None else shared_node.node_prefix_total_len
+        context = infer_batch.InferenceContext()
+        context.args = SimpleNamespace(page_size=page_size, linear_att_hash_page_size=4, linear_att_page_block_num=2)
+        context.is_hybrid_att_model = hybrid
+        context.radix_cache = cache
+        context.req_manager = SimpleNamespace(
+            page_size=page_size, req_to_token_indexs=torch.full((1, 32), -1, dtype=torch.int32, device="cuda")
+        )
+        req = infer_batch.InferReq.__new__(infer_batch.InferReq)
+        req.__dict__.update(
+            req_idx=0,
+            req_manager=context.req_manager,
+            args=SimpleNamespace(req_token_index_mode=mode),
+            cur_kv_len=12 if partial_prefill else 13,
+            hold_kv_len=16,
+            cache_tiers=(CacheTier.GPU,),
+            shared_kv_node=shared_node,
+            hybrid_cache_len=24 if partial_prefill else 12,
+            tail_small_page_buffer_id=small_slots.alloc_one_state_cache() if attention == "hybrid_small" else None,
+            hybrid_len_to_big_page_id=SortedDict(),
+            get_input_token_ids=lambda: key.tolist(),
+            shm_req=SimpleNamespace(
+                input_len=25 if partial_prefill else 13,
+                shm_cur_kv_len=12 if partial_prefill else 13,
+                hybrid_token_hash_list=SimpleNamespace(get_all=lambda: [1, 2, 3, 4, 5, 6]),
+            ),
+        )
+        if hybrid and shared_len < 8:
+            req.hybrid_len_to_big_page_id[8] = big_slots.alloc_one_state_cache()
+        req.token_index = create_req_token_index(req, torch.cuda.current_stream())
+        if shared_indexes is not None:
+            req.token_index.append_token_indexes(shared_indexes)
+        req.token_index.append_token_indexes(allocator.alloc(req.hold_kv_len - shared_len))
+        cache_len = (13 // page_size * page_size) if not hybrid else (12 if attention == "hybrid_small" else 8)
+        expected_indexes = req.token_index.get_cpu_token_indexes(0, cache_len)
+        if seeded_indexes is not None:
+            expected_indexes[:8] = seeded_indexes
+        monkeypatch.setattr(infer_batch, "g_infer_context", context)
+        monkeypatch.setattr(infer_batch, "get_env_start_args", lambda: context.args)
+        return context, req, allocator, small_slots, big_slots, expected_indexes
+
+    return make_case
+
+
+@pytest.mark.parametrize("attention", ["full", "hybrid_small", "hybrid_big", "hybrid_partial"])
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("page_size", [1, 4])
+@pytest.mark.parametrize("cache_case", ["uncached", "duplicate", "shared"])
+def test_release_and_eviction_conserve_kv_pages_and_checkpoint_slots(
+    make_release_case, attention, mode, page_size, cache_case
+):
+    context, req, allocator, small_slots, big_slots, expected_indexes = make_release_case(
+        attention, mode, page_size, cache_case
+    )
+    cache = context.radix_cache
+    released = []
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    # Reusing the request row must not change cached values, including pure_gpu release views.
+    context.req_manager.req_to_token_indexs.fill_(-99)
+    assert req.cur_kv_len == req.hold_kv_len == req.shm_req.shm_cur_kv_len == 0
+    assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+    assert req.shared_kv_node is None and req.tail_small_page_buffer_id is None
+    assert not req.hybrid_len_to_big_page_id
+    cache_len = expected_indexes.numel()
+    assert cache.get_refed_tokens_num() == 0
+    assert cache.root_node.ref_counter == 1
+    assert cache.get_tree_total_tokens_num() == cache_len
+    assert allocator.can_use_mem_size == allocator.size - cache_len
+    assert big_slots.get_used_cache_num() == (1 if attention != "full" and cache_len > 0 else 0)
+    assert small_slots.get_used_cache_num() == (1 if attention == "hybrid_small" else 0)
+
+    key = torch.arange(cache_len, dtype=torch.int64)
+    if attention == "full":
+        node, matched, indexes = cache.match_prefix(key, update_refs=True)
+    else:
+        node, matched, indexes = cache.match_prefix(key, block_hashs=[1, 2, 3][: cache_len // 4], update_refs=True)
+    assert matched == cache_len
+    if cache_len:
+        torch.testing.assert_close(indexes, expected_indexes, check_dtype=False)
+        cache.dec_node_ref_counter(node)
+    else:
+        assert node is None and indexes is None
+
+    free_indexes = allocator.alloc(allocator.can_use_mem_size)
+    assert len(set(free_indexes.tolist())) == free_indexes.numel()
+    assert set(free_indexes.tolist()).isdisjoint(expected_indexes.tolist())
+    allocator.free(free_indexes)
+    # Real eviction must return every physical KV page and both kinds of checkpoint slot exactly once.
+    cache.free_radix_cache_to_get_enough_token(allocator.size)
+    assert cache.get_tree_total_tokens_num() == cache.get_refed_tokens_num() == 0
+    assert allocator.can_use_mem_size == allocator.size
+    assert sorted(allocator.alloc(allocator.size).tolist()) == list(range(allocator.size))
+    for slots in (small_slots, big_slots):
+        assert slots.get_used_cache_num() == 0
+        assert sorted(slots.free_list) == list(range(slots.size))
+
+
+@pytest.mark.parametrize(
+    "checkpoint_lens,computed_len,target_len,inserted_len",
+    [
+        ([8, 16], 20, 24, 16),
+        ([8, 16, 24], 20, 24, 16),
+        ([8, 16], 16, 24, 16),
+        ([8], 7, 24, 0),
+        ([], 7, 24, 0),
+        ([8, 16], 20, 12, 8),
+        ([8, 16, 24], 24, 24, 24),
+        ([8], 20, 12, 8),
+    ],
+    ids=[
+        "multiple",
+        "future",
+        "at_boundary",
+        "future_only",
+        "none",
+        "target_limit",
+        "final_big",
+        "decode_after_target",
+    ],
+)
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("cache_case", ["uncached", "duplicate", "shared"])
+def test_partial_prefill_validates_checkpoints_and_caches_completed_prefix(
+    make_release_case, monkeypatch, checkpoint_lens, computed_len, target_len, inserted_len, mode, cache_case
+):
+    context, req, allocator, small_slots, big_slots, first_page_indexes = make_release_case(
+        "hybrid_partial", mode, 4, cache_case
+    )
+    cache = context.radix_cache
+    shared_len = 8 if cache_case == "shared" else 0
+    req.cur_kv_len = max(computed_len, shared_len)
+    req.hybrid_cache_len = target_len
+    req.token_index.append_token_indexes(allocator.alloc(8))
+    req.hold_kv_len = 24
+    for length in list(req.hybrid_len_to_big_page_id):
+        if length not in checkpoint_lens:
+            big_slots.free_state_cache([req.hybrid_len_to_big_page_id.pop(length)])
+    for length in checkpoint_lens:
+        if length > shared_len and length not in req.hybrid_len_to_big_page_id:
+            req.hybrid_len_to_big_page_id[length] = big_slots.alloc_one_state_cache()
+
+    cache_len = max(inserted_len, 8 if cache_case != "uncached" else 0)
+    expected_indexes = req.token_index.get_cpu_token_indexes(0, cache_len)
+    if cache_case != "uncached":
+        expected_indexes[:8] = first_page_indexes
+    checkpoint = req.hybrid_len_to_big_page_id.get(inserted_len)
+    insert = Mock(wraps=cache.insert)
+    monkeypatch.setattr(cache, "insert", insert)
+    released = []
+    no_new_big_page = inserted_len <= shared_len
+    out_of_range_checkpoint = any(length > max(shared_len, inserted_len) for length in req.hybrid_len_to_big_page_id)
+    if req.hybrid_len_to_big_page_id and (no_new_big_page or out_of_range_checkpoint):
+        # 无新大页时持有 checkpoint，或 checkpoint 超出插入前缀，均违反请求状态约束。
+        checkpoints = dict(req.hybrid_len_to_big_page_id)
+        available_tokens = allocator.can_use_mem_size
+        cached_tokens = cache.get_tree_total_tokens_num()
+        with pytest.raises(AssertionError):
+            context.free_a_req_mem(released, req)
+        if no_new_big_page:
+            insert.assert_not_called()
+            assert released == []
+            assert req.cur_kv_len == max(computed_len, shared_len)
+            assert req.hybrid_len_to_big_page_id == checkpoints
+            assert cache.get_tree_total_tokens_num() == cached_tokens
+        else:
+            assert insert.call_count == 1
+            assert len(released) == 1
+            assert req.cur_kv_len == inserted_len
+            # insert 可能已接管有效 checkpoint，未消费的条目仍由请求持有。
+            assert set(req.hybrid_len_to_big_page_id.items()) <= set(checkpoints.items())
+            assert cache.get_tree_total_tokens_num() <= inserted_len
+        assert req.token_index.token_index_len == req.hold_kv_len == 24
+        assert allocator.can_use_mem_size == available_tokens
+        assert cache.get_refed_tokens_num() == shared_len
+        assert big_slots.get_used_cache_num() == cache.get_tree_total_tokens_num() // 8 + len(
+            req.hybrid_len_to_big_page_id
+        )
+        return
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    context.req_manager.req_to_token_indexs.fill_(-99)
+
+    assert insert.call_count == 1
+    assert insert.call_args.args[0].numel() == max(shared_len, inserted_len)
+    assert req.cur_kv_len == req.hold_kv_len == req.token_index.token_index_len == 0
+    assert req.shared_kv_node is None and not req.hybrid_len_to_big_page_id
+    assert cache.get_tree_total_tokens_num() == cache_len
+    assert cache.get_refed_tokens_num() == 0
+    assert allocator.can_use_mem_size == allocator.size - cache_len
+    assert small_slots.get_used_cache_num() == 0
+    assert big_slots.get_used_cache_num() == cache_len // 8
+    node, matched, indexes = cache.match_prefix(
+        torch.arange(24, dtype=torch.int64), block_hashs=[1, 2, 3, 4, 5, 6], update_refs=True
+    )
+    assert matched == cache_len
+    if cache_len:
+        torch.testing.assert_close(indexes, expected_indexes, check_dtype=False)
+        if cache_case == "uncached" or cache_len > 8:
+            assert node.big_page_buffer_idx == checkpoint
+        cache.dec_node_ref_counter(node)
+    else:
+        assert node is None and indexes is None
+
+    cache.free_radix_cache_to_get_enough_token(allocator.size)
+    assert allocator.can_use_mem_size == allocator.size
+    assert sorted(allocator.alloc(allocator.size).tolist()) == list(range(allocator.size))
+    assert big_slots.get_used_cache_num() == 0
+    assert sorted(big_slots.free_list) == list(range(big_slots.size))
+
+
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("reserved_len", [0, 8])
+def test_hybrid_release_before_first_forward_handles_reserved_capacity(
+    make_release_case, monkeypatch, mode, reserved_len
+):
+    context, req, allocator, small_slots, big_slots, _ = make_release_case("hybrid_partial", mode, 4, "uncached")
+    allocator.free(req.token_index.get_free_token_indexes(0, req.hold_kv_len))
+    req.token_index.reset()
+    big_slots.free_state_cache(list(req.hybrid_len_to_big_page_id.values()))
+    req.hybrid_len_to_big_page_id.clear()
+    req.cur_kv_len = 0
+    req.hold_kv_len = reserved_len
+    if reserved_len:
+        req.token_index.append_token_indexes(allocator.alloc(reserved_len))
+    insert = Mock(wraps=context.radix_cache.insert)
+    monkeypatch.setattr(context.radix_cache, "insert", insert)
+
+    released = []
+    context.free_a_req_mem(released, req)
+    if reserved_len:
+        allocator.free(torch.cat(released))
+        assert insert.call_count == 1
+        assert insert.call_args.args[0].numel() == insert.call_args.args[1].numel() == 0
+    else:
+        insert.assert_not_called()
+        assert released == []
+    assert context.radix_cache.get_tree_total_tokens_num() == context.radix_cache.get_refed_tokens_num() == 0
+    assert context.radix_cache.root_node.ref_counter == 1
+    assert req.cur_kv_len == req.hold_kv_len == req.token_index.token_index_len == 0
+    assert allocator.can_use_mem_size == allocator.size
+    assert sorted(allocator.alloc(allocator.size).tolist()) == list(range(allocator.size))
+    assert small_slots.get_used_cache_num() == big_slots.get_used_cache_num() == 0
+
+
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("shared_prefix_len", [0, 8])
+@pytest.mark.parametrize("available_tokens", [3, 4, 5])
+def test_hybrid_small_page_match_preserves_tail_capacity_and_checkpoint_order(
+    make_release_case, monkeypatch, mode, shared_prefix_len, available_tokens
+):
+    context, req, allocator, small_slots, big_slots, _ = make_release_case("hybrid_small", mode, 4, "uncached")
+    matched_len = shared_prefix_len + 4
+    if shared_prefix_len == 0:
+        big_slots.free_state_cache(list(req.hybrid_len_to_big_page_id.values()))
+        req.hybrid_len_to_big_page_id.clear()
+    req.cur_kv_len = req.hybrid_cache_len = matched_len
+    small_checkpoint = req.tail_small_page_buffer_id
+    cached_indexes = req.token_index.get_cpu_token_indexes(0, matched_len)
+    released = []
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    free_tokens_before_match = allocator.can_use_mem_size
+
+    req.args = SimpleNamespace(
+        page_size=4,
+        linear_att_hash_page_size=4,
+        linear_att_page_block_num=2,
+        max_req_total_len=32,
+    )
+    req.sampling_param = SimpleNamespace(disable_prompt_cache=False)
+    req.get_cur_total_len = lambda: matched_len + 1
+    req.shm_req.shm_prompt_ids = SimpleNamespace(arr=np.arange(matched_len + 1))
+    block_hashs = list(range(1, matched_len // 4 + 1))
+    req.shm_req.hybrid_token_hash_list = SimpleNamespace(size=len(block_hashs), get_all=lambda: block_hashs)
+    req.shm_req.prompt_cache_len = 0
+    monkeypatch.setattr(context, "get_can_alloc_token_num", lambda: available_tokens)
+
+    operations = []
+    copy_kv = Mock(side_effect=lambda *args: operations.append("copy_kv"))
+    context.radix_cache.mem_manager.operator = SimpleNamespace(copy_mem_to_mem=copy_kv)
+    copy_logprobs = Mock(side_effect=lambda **kwargs: operations.append("copy_logprobs"))
+    req.prompt_selected_logprobs = SimpleNamespace(copy_capture_slots_if_needed=copy_logprobs)
+
+    def restore_small_page(req):
+        assert req.shared_kv_node.small_page_buffer_idx == small_checkpoint
+        assert req.shared_kv_node.ref_counter > 0
+        operations.append("restore_small")
+
+    context.req_manager.restore_small_page_state = restore_small_page
+    context.req_manager.restore_big_page_state = Mock()
+    context.req_manager.init_hybrid_attention_state = Mock()
+    deref = context.radix_cache.deref_to_first_big_page_node
+
+    def release_small_page_reference(node):
+        operations.append("deref")
+        return deref(node)
+
+    monkeypatch.setattr(context.radix_cache, "deref_to_first_big_page_node", release_small_page_reference)
+    req._hybrid_match_radix_cache()
+
+    rebuild_tail = available_tokens > 4
+    expected_len = matched_len if rebuild_tail else shared_prefix_len
+    expected_indexes = cached_indexes[:shared_prefix_len]
+    if rebuild_tail:
+        source_indexes, private_indexes = copy_kv.call_args.args
+        torch.testing.assert_close(source_indexes, cached_indexes[shared_prefix_len:])
+        assert private_indexes.numel() == 4
+        assert not torch.isin(private_indexes, cached_indexes).any()
+        copy_logprobs.assert_called_once()
+        torch.testing.assert_close(copy_logprobs.call_args.kwargs["source_indexes"], source_indexes)
+        torch.testing.assert_close(copy_logprobs.call_args.kwargs["destination_indexes"], private_indexes)
+        expected_indexes = torch.cat([expected_indexes, private_indexes])
+        assert operations == ["copy_kv", "copy_logprobs", "restore_small", "deref"]
+    else:
+        copy_kv.assert_not_called()
+        copy_logprobs.assert_not_called()
+        assert operations == ["deref"]
+
+    assert req.cur_kv_len == req.hold_kv_len == req.shm_req.shm_cur_kv_len == expected_len
+    assert req.shm_req.prompt_cache_len == req.token_index.token_index_len == expected_len
+    torch.testing.assert_close(req.token_index.get_cpu_token_indexes(0, expected_len), expected_indexes)
+    assert allocator.can_use_mem_size == free_tokens_before_match - (4 if rebuild_tail else 0)
+    assert context.radix_cache.get_refed_tokens_num() == shared_prefix_len
+    if shared_prefix_len:
+        assert req.shared_kv_node.is_big_page_node()
+        assert req.shared_kv_node.node_prefix_total_len == shared_prefix_len
+    else:
+        assert req.shared_kv_node is None
+    if not rebuild_tail and shared_prefix_len:
+        context.req_manager.restore_big_page_state.assert_called_once_with(
+            big_page_buffer_idx=req.shared_kv_node.big_page_buffer_idx, req=req
+        )
+    else:
+        context.req_manager.restore_big_page_state.assert_not_called()
+    if expected_len == 0:
+        context.req_manager.init_hybrid_attention_state.assert_called_once_with(req=req)
+    else:
+        context.req_manager.init_hybrid_attention_state.assert_not_called()
+
+    released = []
+    context.free_a_req_mem(released, req)
+    if released:
+        allocator.free(torch.cat(released))
+    context.radix_cache.clear_tree_nodes()
+    assert allocator.can_use_mem_size == allocator.size
+    assert small_slots.get_used_cache_num() == big_slots.get_used_cache_num() == 0
+
+
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("new_tokens", [0, 1])
+@pytest.mark.parametrize("evict_checkpoint", [False, True])
+def test_hybrid_release_skips_insert_for_shared_small_prefix(
+    make_release_case, monkeypatch, mode, new_tokens, evict_checkpoint
+):
+    context, req, allocator, small_slots, big_slots, expected_indexes = make_release_case(
+        "hybrid_small", mode, 4, "uncached"
+    )
+    # Use the real match path with big pages disabled, which retains a shared small node.
+    big_slots.free_state_cache(list(req.hybrid_len_to_big_page_id.values()))
+    req.hybrid_len_to_big_page_id.clear()
+    context.args.linear_att_page_block_num = 10000000
+    context.radix_cache = hybrid_att_radix_cache.HybridAttPagedRadixCache(
+        64,
+        0,
+        hash_page_size=4,
+        big_page_num=context.args.linear_att_page_block_num,
+        kv_cache_mem_manager=context.radix_cache.mem_manager,
+        small_page_buffers=small_slots,
+    )
+    req.args = SimpleNamespace(
+        req_token_index_mode=mode,
+        page_size=4,
+        linear_att_hash_page_size=4,
+        linear_att_page_block_num=context.args.linear_att_page_block_num,
+        max_req_total_len=32,
+    )
+    req.sampling_param = SimpleNamespace(disable_prompt_cache=False)
+    req.get_cur_total_len = lambda: 13
+    req.shm_req.shm_prompt_ids = SimpleNamespace(arr=np.arange(13))
+    req.shm_req.hybrid_token_hash_list = SimpleNamespace(size=3, get_all=lambda: [1, 2, 3])
+    restored_checkpoints = []
+    context.req_manager.restore_small_page_state = lambda req: restored_checkpoints.append(
+        req.shared_kv_node.small_page_buffer_idx
+    )
+    released = []
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    cache = context.radix_cache
+    req._hybrid_match_radix_cache()
+    node = req.shared_kv_node
+    shared_len = req.cur_kv_len
+    checkpoint = node.small_page_buffer_idx
+    assert shared_len == 12 and checkpoint is not None
+    assert not node.is_big_page_node()
+    assert restored_checkpoints == [checkpoint]
+    reserved_checkpoints = []
+    if evict_checkpoint:
+        # Exhaust the slot pool, then recycle the checkpoint while its node is still referenced.
+        reserved_checkpoints = small_slots.alloc_state_cache(small_slots.get_free_cache_num())
+        cache.free_one_small_page_buffer()
+        assert node.ref_counter > 0 and node.small_page_buffer_idx is None
+        reused_checkpoint = small_slots.alloc_one_state_cache()
+        assert reused_checkpoint == checkpoint
+        reserved_checkpoints.append(reused_checkpoint)
+        small_slots.states[reused_checkpoint].fill_(42)
+    req.cur_kv_len = shared_len + new_tokens
+    req.hold_kv_len = 16
+    req.token_index.append_token_indexes(allocator.alloc(req.hold_kv_len - shared_len))
+    insert = Mock(wraps=cache.insert)
+    monkeypatch.setattr(cache, "insert", insert)
+
+    released = []
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    context.req_manager.req_to_token_indexs.fill_(-99)
+
+    insert.assert_not_called()
+    assert req.shared_kv_node is None
+    assert req.cur_kv_len == req.hold_kv_len == req.token_index.token_index_len == 0
+    assert cache.get_tree_total_tokens_num() == shared_len
+    assert cache.get_refed_tokens_num() == 0
+    assert allocator.can_use_mem_size == allocator.size - shared_len
+    assert big_slots.get_used_cache_num() == 0
+    assert small_slots.get_used_cache_num() == (small_slots.size if evict_checkpoint else 1)
+    assert node.small_page_buffer_idx == (None if evict_checkpoint else checkpoint)
+    if evict_checkpoint:
+        assert small_slots.states[reused_checkpoint].item() == 42
+    node, matched, indexes = cache.match_prefix(
+        torch.arange(12, dtype=torch.int64), block_hashs=[1, 2, 3], update_refs=True
+    )
+    if evict_checkpoint:
+        assert node is None and indexes is None and matched == 0
+        small_slots.free_state_cache(reserved_checkpoints)
+    else:
+        assert matched == shared_len
+        torch.testing.assert_close(indexes, expected_indexes, check_dtype=False)
+        cache.dec_node_ref_counter(node)
+    cache.free_radix_cache_to_get_enough_token(allocator.size)
+    assert allocator.can_use_mem_size == allocator.size
+    assert sorted(allocator.alloc(allocator.size).tolist()) == list(range(allocator.size))
+    for slots in (small_slots, big_slots):
+        assert slots.get_used_cache_num() == 0
+        assert sorted(slots.free_list) == list(range(slots.size))
+
+
+@pytest.mark.parametrize("prefix_len", [4, 12])
+def test_hybrid_new_small_prefix_still_requires_checkpoint(make_release_case, prefix_len):
+    context, req, _, _, big_slots, _ = make_release_case("hybrid_partial", "pin_mem_cache", 4, "uncached")
+    indexes = req.token_index.get_cpu_token_indexes(0, prefix_len)
+    big_checkpoints = req.hybrid_len_to_big_page_id if prefix_len == 12 else SortedDict()
+    original_checkpoints = dict(big_checkpoints)
+    with pytest.raises(AssertionError):
+        context.radix_cache.insert(
+            torch.arange(prefix_len, dtype=torch.int64),
+            indexes,
+            block_hashs=[1, 2, 3][: prefix_len // 4],
+            block_state_idxs=[None] * (prefix_len // 4),
+            len_to_big_page_id=big_checkpoints,
+        )
+    assert context.radix_cache.get_tree_total_tokens_num() == 0
+    assert context.radix_cache.root_node.children == {}
+    assert big_checkpoints == original_checkpoints
+    assert big_slots.get_used_cache_num() == 1
+
+
+@pytest.mark.parametrize("attention", ["hybrid_small", "hybrid_partial"])
+@pytest.mark.parametrize("failed_range", ["tail", "cached_prefix"])
+def test_hybrid_snapshot_failure_keeps_checkpoint_owned_by_request(
+    make_release_case, monkeypatch, attention, failed_range
+):
+    context, req, allocator, small_slots, big_slots, _ = make_release_case(attention, "pin_mem_cache", 4, "uncached")
+    original_kv_len = req.cur_kv_len
+    cache_len = 12 if attention == "hybrid_small" else 8
+    checkpoint = req.tail_small_page_buffer_id
+    big_checkpoints = dict(req.hybrid_len_to_big_page_id)
+    get_indexes = req.token_index.get_cpu_token_indexes
+
+    def failed_snapshot(start, end):
+        if (start == 0) == (failed_range == "cached_prefix"):
+            raise RuntimeError("snapshot failed")
+        return get_indexes(start, end)
+
+    with monkeypatch.context() as failure:
+        failure.setattr(req.token_index, "get_cpu_token_indexes", failed_snapshot)
+        with pytest.raises(RuntimeError, match="snapshot failed"):
+            context.free_a_req_mem([], req)
+    assert req.tail_small_page_buffer_id == checkpoint
+    assert req.hybrid_len_to_big_page_id == big_checkpoints
+    assert req.cur_kv_len == (original_kv_len if failed_range == "tail" else cache_len)
+    assert small_slots.get_used_cache_num() == (1 if attention == "hybrid_small" else 0)
+    assert big_slots.get_used_cache_num() == len(big_checkpoints)
+    assert context.radix_cache.get_tree_total_tokens_num() == 0
+    released = []
+    context.free_a_req_mem(released, req)
+    allocator.free(torch.cat(released))
+    context.radix_cache.free_radix_cache_to_get_enough_token(allocator.size)
+    assert allocator.can_use_mem_size == allocator.size
+    assert small_slots.get_used_cache_num() == big_slots.get_used_cache_num() == 0
diff --git a/unit_tests/common/test_kv_cache_allocator.py b/unit_tests/common/test_kv_cache_allocator.py
new file mode 100644
index 0000000000..f2fcfc4b41
--- /dev/null
+++ b/unit_tests/common/test_kv_cache_allocator.py
@@ -0,0 +1,334 @@
+import random
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from lightllm.common.kv_cache_mem_manager import allocator as impl
+from lightllm.common.kv_cache_mem_manager import mem_manager as mem_impl
+
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="requires pinned CPU allocations")
+
+
+class _SharedInt:
+    def __init__(self, name):
+        self.name = name
+        self.value = None
+
+    def set_value(self, value):
+        self.value = value
+
+    def get_value(self):
+        return self.value
+
+
+@pytest.fixture
+def make_allocator(monkeypatch):
+    monkeypatch.setattr(impl, "SharedInt", _SharedInt)
+    monkeypatch.setattr(impl, "get_current_rank_in_node", lambda: 0)
+    return impl.KvCacheAllocator
+
+
+def assert_capacity(allocator, available_tokens):
+    assert allocator.can_use_mem_size == available_tokens
+    assert allocator.shared_can_use_token_num.get_value() == available_tokens
+    assert (allocator.mark_end - allocator.mark_start) * allocator.page_size == available_tokens
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_compressed_page_state_returns_complete_pinned_token_indexes(make_allocator, page_size):
+    allocator = make_allocator(5 * page_size, page_size)
+    assert allocator.mem_state.tolist() == [i * page_size for i in range(5)]
+    assert allocator.mem_state.dtype == torch.int32 and allocator.mem_state.is_pinned()
+    assert allocator.mem_state.numel() == allocator.size // page_size
+    assert allocator._return_buffer is None
+    assert_capacity(allocator, 5 * page_size)
+
+    first = allocator.alloc(2 * page_size)
+    second = allocator.alloc(page_size)
+    assert first.tolist() == list(range(2 * page_size))
+    assert second.tolist() == list(range(2 * page_size, 3 * page_size))
+    assert first.dtype == torch.int32 and first.is_pinned()
+    assert first.untyped_storage().data_ptr() == allocator._return_buffer.untyped_storage().data_ptr()
+    assert first.untyped_storage().nbytes() == 1024 * 1024 * first.element_size()
+    assert allocator._return_buffer.numel() == 1024 * 1024 - 3 * page_size
+    assert second.data_ptr() == first.data_ptr() + first.numel() * first.element_size()
+    assert first.untyped_storage().data_ptr() != allocator.mem_state.untyped_storage().data_ptr()
+    assert_capacity(allocator, 2 * page_size)
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+@pytest.mark.parametrize("index_kind", ["list", "cpu", "cuda"])
+def test_nonadjacent_pages_are_recycled_from_token_indexes(make_allocator, page_size, index_kind):
+    allocator = make_allocator(5 * page_size, page_size)
+    indexes = allocator.alloc(5 * page_size).clone()
+    # Recycle physical pages 3 and 1, in that order, using full token-level indexes.
+    released = torch.cat([indexes[3 * page_size : 4 * page_size], indexes[page_size : 2 * page_size]])
+    expected = released.clone()
+    if index_kind == "list":
+        released = released.tolist()
+    elif index_kind == "cuda":
+        released = released.cuda()
+    allocator.free(released)
+    assert_capacity(allocator, 2 * page_size)
+    assert allocator.mem_state[allocator.mark_start :].tolist() == [3 * page_size, page_size]
+
+    # free() owns the compressed indexes even if the caller reuses its input buffer.
+    if isinstance(released, torch.Tensor):
+        released.fill_(-1)
+    else:
+        released[:] = [-1] * len(released)
+    recycled = allocator.alloc(2 * page_size)
+    assert torch.equal(recycled, expected)
+    assert_capacity(allocator, 0)
+    with pytest.raises(AssertionError, match="error alloc state"):
+        allocator.alloc(page_size)
+    assert_capacity(allocator, 0)
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 64])
+def test_return_buffer_replacement_and_free_preserve_retained_results(make_allocator, page_size):
+    # A full slab for aligned page sizes; page_size=3 leaves one unusable element.
+    size = 1024 * 1024 // page_size * page_size
+    allocator = make_allocator(size, page_size)
+    pointers, retained = [], []
+    for iteration in range(4):
+        indexes = allocator.alloc(size)
+        pointers.append(indexes.data_ptr())
+        for old, snapshot in retained:
+            assert torch.equal(old, snapshot)
+        snapshot = indexes.clone()
+        retained.append((indexes, snapshot))
+        allocator.free(indexes.view(-1, page_size).flip(0).reshape(-1))
+        assert torch.equal(indexes, snapshot)
+        assert_capacity(allocator, size)
+    assert len(set(pointers)) == 4
+    assert allocator._return_buffer.numel() == 1024 * 1024 - size
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 64])
+def test_large_requests_allocate_separately_and_preserve_small_buffer(make_allocator, page_size):
+    large_size = (1024 * 1024 // page_size + 1) * page_size
+    allocator = make_allocator(large_size + page_size, page_size)
+    large = allocator.alloc(large_size)
+    snapshot = large.clone()
+    assert allocator._return_buffer is None
+    assert large.is_pinned() and large.dtype == torch.int32
+    assert large.untyped_storage().nbytes() == large_size * large.element_size()
+    assert large.tolist() == list(range(large_size))
+
+    small = allocator.alloc(page_size)
+    remaining = allocator._return_buffer
+    assert remaining.numel() == 1024 * 1024 - page_size
+    assert small.tolist() == list(range(large_size, large_size + page_size))
+    allocator.free(large.view(-1, page_size).flip(0).reshape(-1))
+    recycled = allocator.alloc(large_size)
+    assert recycled.data_ptr() != large.data_ptr()
+    assert allocator._return_buffer is remaining
+    assert torch.equal(recycled, snapshot.view(-1, page_size).flip(0).reshape(-1))
+    assert torch.equal(large, snapshot)
+    assert_capacity(allocator, 0)
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_free_all_restores_page_stride_and_token_capacity(make_allocator, page_size):
+    allocator = make_allocator(6 * page_size, page_size)
+    indexes = allocator.alloc(6 * page_size)
+    allocator.free(indexes.view(-1, page_size).flip(0).reshape(-1))
+    allocator.alloc(page_size)
+    remaining = allocator._return_buffer
+    allocator.free_all()
+    assert allocator.mem_state.tolist() == [i * page_size for i in range(6)]
+    assert allocator._return_buffer is remaining
+    assert_capacity(allocator, 6 * page_size)
+    assert allocator.alloc(6 * page_size).tolist() == list(range(6 * page_size))
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_resize_and_empty_operations_preserve_page_size(make_allocator, page_size):
+    allocator = make_allocator(3 * page_size, page_size)
+    previous = allocator.alloc(page_size)
+    snapshot = previous.clone()
+    for page_count in [7, 2, 0, 4]:
+        allocator.resize(page_count * page_size)
+        assert allocator.mem_state.tolist() == [i * page_size for i in range(page_count)]
+        assert allocator._return_buffer is None
+        assert torch.equal(previous, snapshot)
+        assert_capacity(allocator, page_count * page_size)
+        assert allocator.alloc(0).numel() == 0
+        assert allocator._return_buffer is None
+        allocator.free([])
+        allocator.free(torch.empty(0, dtype=torch.int32))
+        assert_capacity(allocator, page_count * page_size)
+        assert allocator.alloc(page_count * page_size).tolist() == list(range(page_count * page_size))
+        assert_capacity(allocator, 0)
+        allocator.free_all()
+        assert_capacity(allocator, page_count * page_size)
+
+
+@pytest.mark.parametrize("size,page_size", [(4, 0), (4, -1), (-4, 4), (5, 4)])
+def test_invalid_constructor_sizes_are_rejected(make_allocator, size, page_size):
+    with pytest.raises(AssertionError):
+        make_allocator(size, page_size)
+
+
+def test_invalid_operations_leave_allocation_state_unchanged(make_allocator):
+    allocator = make_allocator(16, 4)
+    allocator.alloc(4)
+    for operation in [
+        lambda: allocator.alloc(-4),
+        lambda: allocator.alloc(1),
+        lambda: allocator.alloc(16),
+        lambda: allocator.free([0, 1, 2]),
+        lambda: allocator.free(list(range(8))),
+        lambda: allocator.resize(17),
+    ]:
+        before = allocator.mem_state.clone()
+        remaining = allocator._return_buffer
+        with pytest.raises(AssertionError):
+            operation()
+        assert torch.equal(allocator.mem_state, before)
+        assert allocator._return_buffer is remaining
+        assert allocator.size == 16
+        assert_capacity(allocator, 12)
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_randomized_alloc_free_never_overlaps_live_pages(make_allocator, page_size):
+    page_count = 23
+    allocator = make_allocator(page_count * page_size, page_size)
+    rng = random.Random(7341 + page_size)
+    live_pages = set()
+    for _ in range(300):
+        available = page_count - len(live_pages)
+        if available and (not live_pages or rng.random() < 0.6):
+            count = rng.randint(1, min(available, 5))
+            indexes = allocator.alloc(count * page_size).clone()
+            blocks = indexes.view(-1, page_size)
+            bases = blocks[:, 0].tolist()
+            assert len(set(bases)) == count
+            assert all(0 <= base < allocator.size and base % page_size == 0 for base in bases)
+            assert live_pages.isdisjoint(bases)
+            assert torch.equal(blocks, blocks[:, :1] + torch.arange(page_size, dtype=torch.int32))
+            live_pages.update(bases)
+        else:
+            bases = rng.sample(sorted(live_pages), rng.randint(1, min(len(live_pages), 5)))
+            indexes = [base + offset for base in bases for offset in range(page_size)]
+            allocator.free(indexes if rng.random() < 0.5 else torch.tensor(indexes, dtype=torch.int32))
+            live_pages.difference_update(bases)
+        assert_capacity(allocator, (page_count - len(live_pages)) * page_size)
+    allocator.free([base + offset for base in sorted(live_pages) for offset in range(page_size)])
+    assert_capacity(allocator, allocator.size)
+    assert sorted(allocator.alloc(allocator.size).tolist()) == list(range(allocator.size))
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_memory_manager_passes_page_size_and_resizes_token_capacity(make_allocator, monkeypatch, page_size):
+    monkeypatch.setattr(mem_impl, "get_env_start_args", lambda: SimpleNamespace(page_size=page_size))
+    manager = mem_impl.MemoryManager(5 * page_size + page_size - 1, torch.float16, 1, 8, 1)
+    assert manager.size == 5 * page_size
+    assert manager.allocator.page_size == page_size
+    assert manager.allocator.mem_state.tolist() == [i * page_size for i in range(5)]
+    assert manager.HOLD_TOKEN_MEMINDEXES == tuple(range(5 * page_size, 6 * page_size))
+    indexes = manager.alloc(2 * page_size)
+    manager.free(indexes)
+    assert_capacity(manager.allocator, manager.size)
+
+    manager.resize_mem(3 * page_size + page_size - 1)
+    assert manager.size == 3 * page_size
+    assert manager.allocator.mem_state.tolist() == [i * page_size for i in range(3)]
+    assert manager.HOLD_TOKEN_MEMINDEXES == tuple(range(3 * page_size, 4 * page_size))
+    indexes = manager.alloc(manager.size)
+    assert indexes.tolist() == list(range(manager.size))
+    manager.free(indexes.cuda())
+    assert_capacity(manager.allocator, manager.size)
+
+
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("page_size", [1, 64])
+def test_request_free_indexes_recycle_pages_after_reset_batch_and_slot_reuse(
+    make_allocator, monkeypatch, mode, page_size
+):
+    # 延迟导入，沿用在线推理初始化顺序，不增加另一套请求/allocator 夹具。
+    from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+    from lightllm.server.router.model_infer.infer_batch import InferReq, InferenceContext
+    from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
+
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+    allocator = make_allocator(6 * page_size, page_size)
+    allocated = allocator.alloc(allocator.size)
+    first = torch.cat([allocated[4 * page_size : 5 * page_size], allocated[page_size : 2 * page_size]]).pin_memory()
+    second = allocated[3 * page_size : 4 * page_size]
+    expected = torch.cat([first, second])
+    table = torch.full((2, 2 * page_size), -1, dtype=torch.int32, device="cuda")
+    manager = SimpleNamespace(req_to_token_indexs=table, page_size=page_size)
+    context = InferenceContext.__new__(InferenceContext)
+    context.req_manager = manager
+    context.radix_cache = None
+    reqs = []
+
+    def unexpected_snapshot(*args):
+        pytest.fail("The pure_gpu release path must not copy a full token index snapshot to CPU first")
+
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        for row, indexes in enumerate((first, second)):
+            req = InferReq.__new__(InferReq)
+            req.req_idx = row
+            req.req_manager = manager
+            req.args = SimpleNamespace(req_token_index_mode=mode)
+            req.cur_kv_len = req.hold_kv_len = indexes.numel()
+            req.shm_req = SimpleNamespace(input_len=req.cur_kv_len, shm_cur_kv_len=req.cur_kv_len)
+            req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+            req.token_index.append_token_indexes(indexes)
+            if mode == "pure_gpu":
+                monkeypatch.setattr(req.token_index, "get_cpu_token_indexes", unexpected_snapshot)
+            reqs.append(req)
+
+        # 混合空区间和非空区间；GPU 模式的空区间也必须留在 GPU，才能直接 cat。
+        torch.cuda._sleep(10_000_000)
+        released = [reqs[0].token_index.get_free_token_indexes(0, 0)]
+        for req in reqs:
+            context.free_a_req_mem(released, req)
+            assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+        assert released[0].numel() == 0
+        assert all(indexes.device.type == ("cuda" if mode == "pure_gpu" else "cpu") for indexes in released)
+        merged = torch.cat(released)
+
+        # 默认流先合并旧视图，再重写请求槽位；allocator 只能回收已合并的旧索引。
+        # 不在 free 前同步或 D2H，真实 free 必须等待默认流并按 page_size 抽取页首。
+        table.fill_(-99)
+        allocator.free(merged)
+        assert_capacity(allocator, 3 * page_size)
+        assert allocator.mem_state[allocator.mark_start :].tolist() == [4 * page_size, page_size, 3 * page_size]
+        assert allocator.alloc(3 * page_size).tolist() == expected.tolist()
+        assert_capacity(allocator, 0)
+
+
+@pytest.mark.parametrize("page_size", [1, 64])
+def test_pure_gpu_free_waits_for_index_stream_before_allocator_d2h(make_allocator, page_size):
+    from lightllm.server.router.model_infer.infer_req_ext1 import PureGpuReqTokenIndex
+
+    allocator = make_allocator(4 * page_size, page_size)
+    indexes = allocator.alloc(2 * page_size)
+    table = torch.full((1, indexes.numel()), -1, dtype=torch.int32, device="cuda")
+    index_stream = torch.cuda.Stream()
+    index_stream.wait_stream(torch.cuda.current_stream())
+    req = SimpleNamespace(req_idx=0, req_manager=SimpleNamespace(req_to_token_indexs=table))
+    helper = PureGpuReqTokenIndex(req, index_stream=index_stream)
+
+    with torch.cuda.stream(index_stream):
+        # 延迟索引写入；主线程不等待，默认流必须通过释放 getter 承接写入依赖。
+        torch.cuda._sleep(100_000_000)
+        helper.append_token_indexes(indexes)
+
+    free_indexes = helper.get_free_token_indexes(0, indexes.numel())
+    assert free_indexes.data_ptr() == table[0].data_ptr()
+    # 不先同步索引流：真实 allocator.free 在当前流按页取首索引并阻塞 D2H。
+    allocator.free(free_indexes)
+    helper.reset()
+    with torch.cuda.stream(index_stream):
+        table.fill_(-99)
+    assert_capacity(allocator, allocator.size)
+    assert allocator.alloc(indexes.numel()).tolist() == indexes.tolist()
+    index_stream.synchronize()
diff --git a/unit_tests/common/test_pinned_memory_stream.py b/unit_tests/common/test_pinned_memory_stream.py
new file mode 100644
index 0000000000..f87ee97b51
--- /dev/null
+++ b/unit_tests/common/test_pinned_memory_stream.py
@@ -0,0 +1,353 @@
+"""验证 raw H2D 的源引用由完成事件和后台线程保活，不依赖服务启动。"""
+
+import gc
+from concurrent.futures import ThreadPoolExecutor
+import threading
+from types import SimpleNamespace
+import weakref
+
+import pytest
+import torch
+
+from lightllm.common import cuda_wrapper
+
+requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA and pinned CPU memory")
+
+
+@pytest.fixture(autouse=True)
+def release_futures(monkeypatch):
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", True)
+    pool_factory = cuda_wrapper.ThreadPoolExecutor
+    pools = {}
+    futures = []
+
+    def track_pool(pool):
+        original_submit = pool.submit
+        pools[pool] = original_submit
+
+        def record_submit(*args, **kwargs):
+            future = original_submit(*args, **kwargs)
+            futures.append(future)
+            return future
+
+        monkeypatch.setattr(pool, "submit", record_submit)
+
+    def create_pool(*args, **kwargs):
+        pool = pool_factory(*args, **kwargs)
+        track_pool(pool)
+        return pool
+
+    # 观察已有或后续实际创建的线程池，夹具自身不触发延迟初始化。
+    pool = cuda_wrapper.PinnedH2DBatchCopier._release_pool
+    if pool is not None:
+        track_pool(pool)
+    monkeypatch.setattr(cuda_wrapper, "ThreadPoolExecutor", create_pool)
+    yield futures
+    # 先收尾所有真实等待任务，再撤销本例 monkeypatch；生产池可被后续测试继续复用。
+    for future in futures:
+        future.result(timeout=10)
+    for submit in pools.values():
+        submit(lambda: None).result(timeout=10)
+
+
+def test_empty_batch_and_fallback_do_not_create_release_pool(monkeypatch):
+    copier = cuda_wrapper.PinnedH2DBatchCopier
+    monkeypatch.setattr(copier, "_release_pool", None)
+    monkeypatch.setattr(copier, "_native_batch_available", False)
+    monkeypatch.setattr(torch.cuda, "current_stream", lambda: SimpleNamespace(cuda_stream=0))
+
+    def unexpected_pool(*args, **kwargs):
+        pytest.fail("Empty batches and copy_ fallback must not create the source release pool")
+
+    monkeypatch.setattr(cuda_wrapper, "ThreadPoolExecutor", unexpected_pool)
+    cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([], [])
+    source = torch.arange(8, dtype=torch.int32)
+    target = torch.empty_like(source)
+    cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target], [source])
+    assert copier._release_pool is None
+    assert torch.equal(target, source)
+
+
+def test_release_pool_is_initialized_once_for_concurrent_callers(monkeypatch):
+    copier = cuda_wrapper.PinnedH2DBatchCopier
+    monkeypatch.setattr(copier, "_release_pool", None)
+    entered = threading.Event()
+    allow_creation = threading.Event()
+    start = threading.Barrier(8)
+    shared_pool = object()
+    creations = []
+
+    def create_pool(**kwargs):
+        creations.append(kwargs)
+        entered.set()
+        assert allow_creation.wait(timeout=5)
+        return shared_pool
+
+    def get_pool():
+        start.wait(timeout=5)
+        return copier._get_release_pool()
+
+    monkeypatch.setattr(cuda_wrapper, "ThreadPoolExecutor", create_pool)
+    with ThreadPoolExecutor(max_workers=8) as callers:
+        futures = [callers.submit(get_pool) for _ in range(8)]
+        try:
+            assert entered.wait(timeout=5)
+        finally:
+            allow_creation.set()
+        assert all(future.result(timeout=5) is shared_pool for future in futures)
+    assert creations == [{"max_workers": 1, "thread_name_prefix": "lightllm-pinned-release"}]
+
+
+@requires_cuda
+@pytest.mark.parametrize("native_batch", [False, True])
+def test_default_stream_copy_orders_reads_and_reuses_upload_stream(monkeypatch, native_batch):
+    runtime = pytest.importorskip("cuda.bindings.runtime")
+    copier = cuda_wrapper.PinnedH2DBatchCopier
+    monkeypatch.setattr(copier, "_native_batch_available", native_batch)
+    stream = torch.cuda.default_stream()
+    source = torch.arange(8192, dtype=torch.int32, pin_memory=True)
+    target = torch.empty(8192, dtype=torch.int32, device="cuda")
+    native_copy = runtime.cudaMemcpyBatchAsync
+    submitted_streams, observed = [], []
+
+    def record_batch(*args):
+        submitted_streams.append(args[-1])
+        return native_copy(*args)
+
+    monkeypatch.setattr(runtime, "cudaMemcpyBatchAsync", record_batch)
+    with torch.cuda.stream(stream):
+        for _ in range(2):
+            torch.cuda._sleep(20_000_000)
+            target.fill_(-1)
+            cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target], [source])
+            assert torch.cuda.current_stream() == stream
+            observed.append(target.clone())
+    for result in observed:
+        assert torch.equal(result.cpu(), source)
+    if native_batch:
+        assert len(submitted_streams) == 2
+        assert submitted_streams[0] == submitted_streams[1] != 0
+    else:
+        assert submitted_streams == []
+
+
+@requires_cuda
+def test_raw_batch_records_one_completion_event_without_small_copy(monkeypatch, release_futures):
+    runtime = pytest.importorskip("cuda.bindings.runtime")
+    source = torch.arange(2048, dtype=torch.int32, pin_memory=True)
+    another_source = torch.arange(2048, dtype=torch.int32, pin_memory=True).add_(10000)
+    sources = [source[128:384], source[1000:1256], another_source[17:273]]
+    target = torch.full((1024,), -1, dtype=torch.int32, device="cuda")
+    targets = [target[2:258], target[300:556], target[600:856]]
+    expected = torch.full((1024,), -1, dtype=torch.int32)
+    expected[2:258], expected[300:556], expected[600:856] = sources
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+    calls = []
+    native_copy, event_record = runtime.cudaMemcpyBatchAsync, torch.cuda.Event.record
+
+    def record_native(*args):
+        calls.append(("batch", args[3], args[-1]))
+        return native_copy(*args)
+
+    def record_event(event, on_stream=None):
+        calls.append(("event", 1, on_stream.cuda_stream))
+        return event_record(event, on_stream)
+
+    def unexpected_copy(*args, **kwargs):
+        pytest.fail("Raw batch source protection must not submit a small copy_ or enter per-tensor fallback")
+
+    def unexpected_sync(*args, **kwargs):
+        pytest.fail("Normal raw batches must not synchronize a CUDA stream on the submitting thread")
+
+    monkeypatch.setattr(runtime, "cudaMemcpyBatchAsync", record_native)
+    monkeypatch.setattr(torch.cuda.Event, "record", record_event)
+    monkeypatch.setattr(torch.Tensor, "copy_", unexpected_copy)
+    with torch.cuda.stream(stream), monkeypatch.context() as no_sync:
+        no_sync.setattr(torch.cuda.Stream, "synchronize", unexpected_sync)
+        no_sync.setattr(torch.cuda, "synchronize", unexpected_sync)
+        cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async(targets, sources)
+        assert len(release_futures) == 1
+        release_futures[0].result(timeout=10)
+    assert cuda_wrapper.PinnedH2DBatchCopier._native_batch_available
+    assert calls == [("batch", 3, stream.cuda_stream), ("event", 1, stream.cuda_stream)]
+    assert torch.equal(target.cpu(), expected)
+
+
+@requires_cuda
+def test_background_wait_holds_tensor_references_without_blocking_submitter(monkeypatch, release_futures):
+    pytest.importorskip("cuda.bindings.runtime")
+    source = torch.arange(8192, dtype=torch.int32, pin_memory=True)
+    source_views = [source[64:1088], source[1088:4160]]
+    expected = torch.cat(source_views)
+    source_ref = weakref.ref(source)
+    view_refs = [weakref.ref(view) for view in source_views]
+    target = torch.empty(4096, dtype=torch.int32, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+    worker_wait_started = threading.Event()
+    allow_worker_finish = threading.Event()
+    main_thread = threading.get_ident()
+    event_synchronize = torch.cuda.Event.synchronize
+
+    def controlled_wait(event):
+        assert threading.get_ident() != main_thread, "Only the background worker may synchronize completion events"
+        worker_wait_started.set()
+        assert allow_worker_finish.wait(timeout=10), "The test did not allow the background wait to finish"
+        event_synchronize(event)
+
+    monkeypatch.setattr(torch.cuda.Event, "synchronize", controlled_wait)
+    try:
+        with torch.cuda.stream(stream):
+            torch.cuda._sleep(100_000_000)
+            cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target[:1024], target[1024:]], source_views)
+        # gate 仍关闭时提交调用已返回，且后台任务正在等待，而非阻塞主线程。
+        assert worker_wait_started.wait(timeout=5)
+        assert len(release_futures) == 1 and not release_futures[0].done()
+        source_views.clear()
+        del source_views, source
+        gc.collect()
+        assert source_ref() is not None and all(reference() is not None for reference in view_refs)
+    finally:
+        allow_worker_finish.set()
+        for future in release_futures:
+            assert future.result(timeout=10) is None
+        stream.synchronize()
+    gc.collect()
+    assert source_ref() is None and all(reference() is None for reference in view_refs)
+    assert torch.equal(target.cpu(), expected)
+
+
+@requires_cuda
+@pytest.mark.parametrize("native_batch", [True, False])
+@pytest.mark.parametrize("shared_storage", [True, False])
+def test_async_copy_keeps_freed_source_storage_out_of_reuse(monkeypatch, native_batch, shared_storage):
+    if native_batch:
+        pytest.importorskip("cuda.bindings.runtime")
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", native_batch)
+    source = torch.arange(8192, dtype=torch.int32, pin_memory=True)
+    another_source = source if shared_storage else torch.arange(8192, dtype=torch.int32, pin_memory=True).add_(10000)
+    source_ptrs = {source.data_ptr(), another_source.data_ptr()}
+    source_views = [source[64:1088], another_source[1088:4160]]
+    expected = torch.cat(source_views)
+    target = torch.empty(4096, dtype=torch.int32, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(200_000_000)
+        cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target[:1024], target[1024:]], source_views)
+    assert cuda_wrapper.PinnedH2DBatchCopier._native_batch_available is native_batch
+    del source_views, source, another_source
+
+    pending_checks = 0
+    for _ in range(64):
+        still_pending = not stream.query()
+        replacement = torch.empty(8192, dtype=torch.int32, pin_memory=True)
+        if still_pending:
+            assert replacement.data_ptr() not in source_ptrs
+            pending_checks += 1
+        replacement.fill_(-1)
+        del replacement
+    stream.synchronize()
+    assert pending_checks > 0, "Delayed transfers must cover source release and pinned memory reuse checks"
+    assert torch.equal(target.cpu(), expected)
+
+
+@requires_cuda
+@pytest.mark.parametrize("failure", ["record", "pool", "submit"])
+def test_event_setup_failure_propagates_without_sync_or_fallback(monkeypatch, failure):
+    pytest.importorskip("cuda.bindings.runtime")
+    source = torch.arange(1024, dtype=torch.int32, pin_memory=True)
+    target = torch.empty_like(source, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+    synchronize = torch.cuda.Stream.synchronize
+
+    def fail_setup(*args, **kwargs):
+        raise RuntimeError(f"test {failure} failure")
+
+    def unexpected_sync(*args, **kwargs):
+        pytest.fail("Source protection setup errors must propagate without synchronous cleanup")
+
+    def unexpected_copy(*args, **kwargs):
+        pytest.fail("copy_ fallback must not hide source protection failure after raw batch submission")
+
+    if failure == "record":
+        monkeypatch.setattr(torch.cuda.Event, "record", fail_setup)
+    elif failure == "pool":
+        monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_release_pool", None)
+        monkeypatch.setattr(cuda_wrapper, "ThreadPoolExecutor", fail_setup)
+    else:
+        monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier._get_release_pool(), "submit", fail_setup)
+    monkeypatch.setattr(torch.cuda.Stream, "synchronize", unexpected_sync)
+    monkeypatch.setattr(torch.Tensor, "copy_", unexpected_copy)
+    try:
+        with torch.cuda.stream(stream):
+            torch.cuda._sleep(20_000_000)
+            with pytest.raises(RuntimeError, match=f"test {failure} failure"):
+                cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target], [source])
+    finally:
+        # 测试捕获注入异常后仍要继续运行，自行持有源并等待在途传输结束。
+        synchronize(stream)
+    assert cuda_wrapper.PinnedH2DBatchCopier._native_batch_available
+    assert torch.equal(target.cpu(), source)
+
+
+@requires_cuda
+def test_pytorch_fallback_needs_no_completion_event_or_pool_task(monkeypatch, release_futures):
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", False)
+    source = torch.arange(1024, dtype=torch.int32, pin_memory=True)
+    target = torch.empty_like(source, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+    copy_sizes = []
+    copy = torch.Tensor.copy_
+
+    def unexpected_event(*args, **kwargs):
+        pytest.fail("PyTorch copy_ already protects source storage and must not create an extra completion event")
+
+    def record_copy(dst, src, *args, **kwargs):
+        copy_sizes.append(src.numel() * src.element_size())
+        return copy(dst, src, *args, **kwargs)
+
+    monkeypatch.setattr(torch.cuda, "Event", unexpected_event)
+    monkeypatch.setattr(torch.Tensor, "copy_", record_copy)
+    with torch.cuda.stream(stream):
+        cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target], [source])
+    stream.synchronize()
+    assert not cuda_wrapper.PinnedH2DBatchCopier._native_batch_available
+    assert release_futures == []
+    assert copy_sizes == [source.numel() * source.element_size()]
+    assert torch.equal(target.cpu(), source)
+
+
+@requires_cuda
+def test_source_release_pool_is_reused_across_batches(release_futures):
+    pytest.importorskip("cuda.bindings.runtime")
+    pool = cuda_wrapper.PinnedH2DBatchCopier._get_release_pool()
+    source = torch.arange(1024, dtype=torch.int32, pin_memory=True)
+    target = torch.empty((2, 1024), dtype=torch.int32, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.default_stream())
+    with torch.cuda.stream(stream):
+        for row in range(2):
+            cuda_wrapper.cuda_memcpy_pinned_h2d_batch_async([target[row]], [source])
+    assert cuda_wrapper.PinnedH2DBatchCopier._get_release_pool() is pool
+    assert len(release_futures) == 2
+    for future in release_futures:
+        assert future.result(timeout=10) is None
+    assert torch.equal(target.cpu(), source.expand(2, -1))
+
+
+def test_background_wait_failure_propagates():
+    source = torch.arange(8, dtype=torch.int32)
+    sources = [source]
+
+    def failed_wait():
+        raise RuntimeError("test background wait failure")
+
+    with pytest.raises(RuntimeError, match="test background wait failure"):
+        cuda_wrapper.PinnedH2DBatchCopier._wait_and_release_pinned_sources(
+            SimpleNamespace(synchronize=failed_wait), sources
+        )
diff --git a/unit_tests/common/test_req_manager_cpu.py b/unit_tests/common/test_req_manager_cpu.py
new file mode 100644
index 0000000000..30137cb04f
--- /dev/null
+++ b/unit_tests/common/test_req_manager_cpu.py
@@ -0,0 +1,724 @@
+from types import SimpleNamespace
+from contextlib import nullcontext
+import builtins
+
+import pytest
+import torch
+
+# Use the production import order to avoid the existing package initialization cycle.
+from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+from lightllm.common.basemodel.basemodel import TpPartBaseModel
+from lightllm.common.req_manager.base import ReqManager, _ReqIndexPool
+from lightllm.common.cuda_wrapper import cuda_memcpy_pinned_h2d_batch_async
+from lightllm.server.router.model_infer.infer_batch import InferReq, InferenceContext
+from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
+from lightllm.server.router.model_infer import infer_batch
+from lightllm.server.router.model_infer.mode_backend.base_backend import ModeBackend
+from unit_tests.common.cpu_index_helpers import bind_cpu_indexes, make_cpu_index_manager
+
+
+@pytest.fixture(autouse=True)
+def reset_batch_copy_availability(monkeypatch):
+    from lightllm.common import cuda_wrapper
+
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", True)
+    monkeypatch.setattr(infer_batch.g_infer_context, "req_manager", None)
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+
+
+def _make_manager(device="cpu"):
+    manager = make_cpu_index_manager(
+        torch.full((3, 16), -1, dtype=torch.int32, pin_memory=torch.device(device).type == "cuda"),
+        req_to_token_indexs=torch.full((3, 16), -1, dtype=torch.int32, device=device),
+    )
+    infer_batch.g_infer_context.req_manager = manager
+    manager.max_request_num = 2
+    manager.req_list = _ReqIndexPool(2)
+    return manager
+
+
+def test_req_index_allocation_preserves_order_capacity_and_reuse():
+    indexes = _ReqIndexPool(3)
+    assert indexes.is_all_free() and indexes.can_alloc_size == 3
+    assert [indexes.alloc() for _ in range(3)] == [0, 1, 2]
+    assert indexes.can_alloc_size == 0 and not indexes.is_all_free()
+    assert indexes.alloc() is None
+
+    indexes.free(1)
+    indexes.free(0)
+    assert indexes.can_alloc_size == 2
+    assert [indexes.alloc(), indexes.alloc()] == [0, 1]
+    assert indexes.alloc() is None
+    for index in range(3):
+        indexes.free(index)
+    assert indexes.is_all_free() and indexes.can_alloc_size == 3
+
+
+def test_req_index_free_rejects_unallocated_and_duplicate_indexes():
+    indexes = _ReqIndexPool(2)
+    with pytest.raises(AssertionError):
+        indexes.free(0)
+    assert indexes.alloc() == 0
+    indexes.free(0)
+    with pytest.raises(AssertionError):
+        indexes.free(0)
+    assert indexes.can_alloc_size == 2
+    assert [indexes.alloc(), indexes.alloc()] == [0, 1]
+    assert indexes.alloc() is None
+
+
+def test_cpu_updates_own_allocator_values_and_defer_gpu_publication():
+    manager = _make_manager()
+    indexes = torch.tensor([20, 21, 22, 23], dtype=torch.int32)
+    manager.token_indexes_cpu[0, 4:8].copy_(indexes)
+    indexes.fill_(99)
+    manager.token_indexes_cpu[0, :4].copy_(torch.arange(4, dtype=torch.int32))
+
+    assert manager.token_indexes_cpu[0, :8].tolist() == [0, 1, 2, 3, 20, 21, 22, 23]
+    assert manager.req_to_token_indexs.eq(-1).all()
+
+
+def test_releasing_request_memory_resets_uploaded_length():
+    manager = _make_manager()
+    manager.token_indexes_cpu[0, :8].copy_(torch.arange(8, dtype=torch.int32))
+    context = InferenceContext()
+    context.req_manager = manager
+    req = SimpleNamespace(
+        req_idx=0,
+        cur_kv_len=6,
+        hold_kv_len=8,
+        shm_req=SimpleNamespace(shm_cur_kv_len=6),
+    )
+    bind_cpu_indexes(req, manager, gpu_index_submitted_len=8)
+    released = []
+    context.free_a_req_mem(released, req)
+
+    assert released[0].tolist() == list(range(8))
+    assert req.cur_kv_len == req.hold_kv_len == req.token_index.gpu_index_submitted_len == 0
+    assert req.token_index.token_index_len == 0
+    assert req.token_index._cpu_buffer is None
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_shared_export_keeps_cpu_mirror_local_and_pinned(monkeypatch):
+    from lightllm.common.kv_cache_mem_manager import mem_manager as impl
+
+    req_manager = _make_manager("cuda")
+    mem_manager = impl.MemoryManager.__new__(impl.MemoryManager)
+    source = req_manager.token_indexes_cpu
+    source_ptr = source.data_ptr()
+    exports = []
+
+    def serialize(manager):
+        assert manager.req_to_token_indexs is req_manager.req_to_token_indexs
+        assert not hasattr(manager, "token_indexes_cpu")
+        exports.append(manager)
+        return memoryview(b"test")
+
+    def forbid_cpu_export(*args, **kwargs):
+        pytest.fail("CPU mirror export and host registration must not run")
+
+    monkeypatch.setattr(torch.Tensor, "share_memory_", forbid_cpu_export)
+    monkeypatch.setattr(torch.cuda, "cudart", forbid_cpu_export)
+    monkeypatch.setattr(impl, "ForkingPickler", SimpleNamespace(dumps=serialize))
+    monkeypatch.setattr(impl, "kv_trans_use_p2p", lambda: False)
+    monkeypatch.setattr(impl, "FileLock", lambda path: nullcontext())
+    monkeypatch.setattr(impl, "get_node_world_size", lambda: 1)
+    monkeypatch.setattr(impl, "get_current_rank_in_node", lambda: 0)
+    monkeypatch.setattr(impl, "get_service_shm_name", lambda name: name)
+    monkeypatch.setattr(
+        impl,
+        "create_or_link_shm",
+        lambda name, expected_size, **kwargs: SimpleNamespace(
+            name=name, size=expected_size, buf=bytearray(expected_size)
+        ),
+    )
+
+    mem_manager.write_to_shm(req_manager)
+    mem_manager.write_to_shm(req_manager)
+
+    assert len(exports) == 4
+    assert req_manager.token_indexes_cpu is source
+    assert source.data_ptr() == source_ptr
+    assert source.is_pinned() and not source.is_shared()
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("hybrid", [False, True])
+@pytest.mark.parametrize("resuming", [False, True])
+def test_prefix_match_appends_to_initialized_or_released_index_state(monkeypatch, mode, hybrid, resuming):
+    manager = _make_manager("cuda")
+    manager.page_size = 4
+    req = InferReq.__new__(InferReq)
+    req.req_idx = 0
+    req.req_manager = manager
+    req.args = SimpleNamespace(
+        page_size=4,
+        req_token_index_mode=mode,
+        linear_att_hash_page_size=2,
+        linear_att_page_block_num=2,
+        max_req_total_len=16,
+    )
+    req.cur_kv_len = req.hold_kv_len = 0
+    req.shared_kv_node = req.tail_small_page_buffer_id = None
+    req.sampling_param = SimpleNamespace(disable_prompt_cache=False)
+    req.shm_req = SimpleNamespace(
+        input_len=8,
+        shm_prompt_ids=SimpleNamespace(arr=list(range(8))),
+        hybrid_token_hash_list=SimpleNamespace(size=3, get_all=lambda: [1, 2, 3]),
+    )
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.current_stream())
+    req.get_cur_total_len = lambda: 8
+    if resuming:
+        req.token_index.append_token_indexes(torch.arange(8, dtype=torch.int32))
+        req.token_index.get_gpu_token_indexes(0, 8)
+        req.cur_kv_len, req.hold_kv_len = 6, 8
+        release_context = InferenceContext()
+        released = []
+        release_context.free_a_req_mem(released, req)
+        assert released[0].cpu().tolist() == list(range(8))
+        assert req.cur_kv_len == req.hold_kv_len == 0
+        assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+
+    node = SimpleNamespace(node_prefix_total_len=4, big_page_buffer_idx=0, is_big_page_node=lambda: True)
+    indexes = torch.arange(32, 36, dtype=torch.int32)
+    manager.restore_big_page_state = lambda **kwargs: None
+    context = SimpleNamespace(
+        is_hybrid_att_model=hybrid,
+        req_manager=manager,
+        radix_cache=SimpleNamespace(match_prefix=lambda key, **kwargs: (node, 4, indexes)),
+    )
+    monkeypatch.setattr(infer_batch, "g_infer_context", context)
+
+    def unexpected_reset():
+        pytest.fail("Prefix matching starts with empty indexes and must not reset them again")
+
+    monkeypatch.setattr(req.token_index, "reset", unexpected_reset)
+    if hybrid:
+        req._hybrid_match_radix_cache()
+    else:
+        req._match_radix_cache()
+
+    assert req.token_index.token_index_len == 4
+    assert req.token_index.gpu_index_submitted_len == (4 if mode == "pure_gpu" else 0)
+    assert req.cur_kv_len == req.hold_kv_len == 4
+    assert req.shared_kv_node is node
+    assert req.shm_req.prompt_cache_len == req.shm_req.shm_cur_kv_len == 4
+    assert req.token_index.get_cpu_token_indexes(0, 4).tolist() == [32, 33, 34, 35]
+    assert req.token_index.get_gpu_token_indexes(0, 4).cpu().tolist() == [32, 33, 34, 35]
+
+
+def test_radix_insertion_owns_cpu_indexes_before_request_row_reuse():
+    manager = _make_manager()
+    manager.token_indexes_cpu[0, :8].copy_(torch.arange(8, dtype=torch.int32))
+    cached = []
+
+    def insert(key, indexes):
+        cached.append(indexes)
+        return 0, None
+
+    context = InferenceContext()
+    context.args = SimpleNamespace(page_size=4)
+    context.req_manager = manager
+    context.radix_cache = SimpleNamespace(insert=insert)
+    req = SimpleNamespace(
+        req_idx=0,
+        cur_kv_len=6,
+        hold_kv_len=8,
+        shared_kv_node=None,
+        get_input_token_ids=lambda: list(range(6)),
+    )
+    bind_cpu_indexes(req, manager)
+    released = []
+    context._full_att_free_req(released, req)
+    assert released[0].tolist() == [4, 5, 6, 7]
+    manager.token_indexes_cpu[0, :8].fill_(99)
+
+    assert cached[0].tolist() == [0, 1, 2, 3]
+    assert manager.req_to_token_indexs.eq(-1).all()
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_async_batch_reads_source_after_prior_stream_writes_and_preserves_untouched_ranges():
+    manager = _make_manager("cuda")
+    incoming_indexes = torch.tensor([40, 41, 42, 43], dtype=torch.int32, device="cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(2_000_000)
+        manager.token_indexes_cpu[0, 4:8].copy_(incoming_indexes, non_blocking=True)
+        cuda_memcpy_pinned_h2d_batch_async(
+            [manager.req_to_token_indexs[0, 4:8]],
+            [manager.token_indexes_cpu[0, 4:8]],
+        )
+    stream.synchronize()
+
+    assert manager.req_to_token_indexs[0, 4:8].cpu().tolist() == [40, 41, 42, 43]
+    assert manager.req_to_token_indexs[0, :4].eq(-1).all()
+    assert manager.req_to_token_indexs[0, 8:].eq(-1).all()
+    assert manager.req_to_token_indexs[1:].eq(-1).all()
+    manager.token_indexes_cpu[0, 4:8].copy_(torch.tensor([80, 81, 82, 83], dtype=torch.int32))
+    with torch.cuda.stream(stream):
+        cuda_memcpy_pinned_h2d_batch_async(
+            [manager.req_to_token_indexs[0, 4:8]],
+            [manager.token_indexes_cpu[0, 4:8]],
+        )
+    stream.synchronize()
+    assert manager.req_to_token_indexs[0, 4:8].cpu().tolist() == [80, 81, 82, 83]
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("default_stream", [False, True])
+def test_native_batch_uses_stream_order_and_handles_disjoint_ranges(monkeypatch, default_stream):
+    from cuda.bindings import runtime as cuda_runtime
+
+    manager = _make_manager("cuda")
+    manager.token_indexes_cpu.copy_(torch.arange(48, dtype=torch.int32).reshape(3, 16))
+    ranges = [(0, 4, 8), (1, 1, 5), (0, 12, 16)]
+    expected = torch.full_like(manager.token_indexes_cpu, -1)
+    for row, start, end in ranges:
+        expected[row, start:end].copy_(manager.token_indexes_cpu[row, start:end])
+    calls = []
+    native_copy = cuda_runtime.cudaMemcpyBatchAsync
+
+    def record_native(*args):
+        calls.append((args[3], args[2], args[4][0].srcAccessOrder))
+        return native_copy(*args)
+
+    monkeypatch.setattr(cuda_runtime, "cudaMemcpyBatchAsync", record_native)
+    stream = torch.cuda.default_stream() if default_stream else torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(2_000_000)
+        cuda_memcpy_pinned_h2d_batch_async(
+            [manager.req_to_token_indexs[row, start:end] for row, start, end in ranges],
+            [manager.token_indexes_cpu[row, start:end] for row, start, end in ranges],
+        )
+        assert torch.cuda.current_stream() == stream
+        result = manager.req_to_token_indexs.clone()
+    stream.synchronize()
+
+    assert calls == [(3, [16, 16, 16], cuda_runtime.cudaMemcpySrcAccessOrder.cudaMemcpySrcAccessOrderStream)]
+    assert torch.equal(result.cpu(), expected)
+
+
+def test_empty_batch_does_not_call_cuda(monkeypatch):
+    original_import = builtins.__import__
+
+    def check_import(name, *args, **kwargs):
+        if name == "cuda.bindings":
+            pytest.fail("empty batch must not import CUDA bindings")
+        return original_import(name, *args, **kwargs)
+
+    monkeypatch.setattr(builtins, "__import__", check_import)
+    cuda_memcpy_pinned_h2d_batch_async([], [])
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("failure", ["missing_binding", "missing_api", "native_error", "exception"])
+def test_batch_fallback_uses_current_stream_without_synchronizing(monkeypatch, failure):
+    from cuda.bindings import runtime as cuda_runtime
+    from lightllm.common import cuda_wrapper as base
+
+    manager = _make_manager("cuda")
+    manager.token_indexes_cpu.copy_(torch.arange(48, dtype=torch.int32).reshape(3, 16))
+    ranges = [(0, 4, 8), (1, 1, 5), (0, 12, 16)]
+    expected = torch.full_like(manager.token_indexes_cpu, -1)
+    for row, start, end in ranges:
+        expected[row, start:end].copy_(manager.token_indexes_cpu[row, start:end])
+
+    batch_attempts = []
+    if failure == "missing_binding":
+        original_import = builtins.__import__
+
+        def missing_import(name, *args, **kwargs):
+            if name == "cuda.bindings":
+                batch_attempts.append(name)
+                raise ModuleNotFoundError("cuda.bindings is unavailable")
+            return original_import(name, *args, **kwargs)
+
+        monkeypatch.setattr(builtins, "__import__", missing_import)
+    elif failure == "missing_api":
+        monkeypatch.delattr(cuda_runtime, "cudaMemcpyBatchAsync")
+    else:
+
+        def failed_batch(*args):
+            batch_attempts.append(args)
+            if failure == "exception":
+                raise RuntimeError("batch API could not be loaded")
+            return (cuda_runtime.cudaError_t.cudaErrorNotSupported,)
+
+        monkeypatch.setattr(cuda_runtime, "cudaMemcpyBatchAsync", failed_batch)
+
+    loop_calls = []
+    torch_copy = torch.Tensor.copy_
+
+    def record_copy(dst, src, *args, **kwargs):
+        if dst.is_cuda and src.device.type == "cpu":
+            assert kwargs["non_blocking"] is True
+            loop_calls.append((dst.numel() * dst.element_size(), torch.cuda.current_stream(dst.device).cuda_stream))
+        return torch_copy(dst, src, *args, **kwargs)
+
+    monkeypatch.setattr(torch.Tensor, "copy_", record_copy)
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    dsts = [manager.req_to_token_indexs[row, start:end] for row, start, end in ranges]
+    srcs = [manager.token_indexes_cpu[row, start:end] for row, start, end in ranges]
+
+    def unexpected_sync(*args, **kwargs):
+        pytest.fail("copy submission must not synchronize")
+
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(2_000_000)
+        manager.req_to_token_indexs.fill_(-1)
+        with monkeypatch.context() as no_sync:
+            no_sync.setattr(torch.cuda.Stream, "synchronize", unexpected_sync)
+            no_sync.setattr(torch.cuda, "synchronize", unexpected_sync)
+            cuda_memcpy_pinned_h2d_batch_async(dsts, srcs)
+        result = manager.req_to_token_indexs.clone()
+    stream.synchronize()
+
+    assert torch.equal(result.cpu(), expected)
+    assert [size for size, _ in loop_calls] == [16] * 3
+    assert all(copy_stream == stream.cuda_stream for _, copy_stream in loop_calls)
+    assert not base.PinnedH2DBatchCopier._native_batch_available
+    attempts_after_failure = len(batch_attempts)
+    manager.token_indexes_cpu.fill_(99)
+    with torch.cuda.stream(stream):
+        cuda_memcpy_pinned_h2d_batch_async(dsts, srcs)
+    stream.synchronize()
+    assert len(batch_attempts) == attempts_after_failure
+    assert len(loop_calls) == 6
+    assert manager.req_to_token_indexs[0, 4:8].cpu().tolist() == [99] * 4
+
+
+@pytest.mark.parametrize("entry", ["forward", "microbatch_overlap_prefill", "microbatch_overlap_decode"])
+def test_model_entries_do_not_manage_request_upload_progress(entry):
+    model = TpPartBaseModel.__new__(TpPartBaseModel)
+    events = []
+    model.req_manager = SimpleNamespace()
+    model.args = SimpleNamespace(enable_prefill_decode_mixed=False)
+    model._prefill = lambda model_input: events.append("infer")
+    model._decode = lambda model_input: events.append("infer")
+    model._microbatch_overlap_prefill_cuda = lambda *inputs: events.append("infer")
+    model._microbatch_overlap_decode_cuda = lambda *inputs: events.append("infer")
+    model_input = SimpleNamespace(
+        is_prefill=True,
+        input_ids=torch.ones(1, dtype=torch.int64),
+        to_cuda=lambda: events.append("input"),
+    )
+
+    if entry == "forward":
+        model.forward(model_input)
+        assert events == ["input", "infer"]
+    else:
+        getattr(model, entry)(model_input, model_input)
+        assert events == ["input", "input", "infer"]
+
+
+def _make_scheduler(monkeypatch, manager, reqs, budget=16, mixed=False):
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+
+    args = SimpleNamespace(page_size=4, enable_cpu_cache=False, enable_prefill_decode_mixed=mixed, run_mode="normal")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    context = SimpleNamespace(
+        req_manager=manager,
+        radix_cache=None,
+        get_can_alloc_token_num=lambda: 32,
+        get_overlap_stream=lambda: stream,
+        cache_placement_controller=SimpleNamespace(set_req_cache_way=lambda reqs: None),
+        filter_reqs=lambda finished_reqs: None,
+        pause_reqs=lambda reqs, is_master_in_dp: None,
+    )
+    monkeypatch.setattr(base_backend, "g_infer_context", context)
+    manager.mem_manager = SimpleNamespace(
+        alloc=lambda size: torch.arange(100, 100 + size, dtype=torch.int32, pin_memory=True)
+    )
+    backend = ModeBackend.__new__(ModeBackend)
+    backend.args = args
+    backend.support_overlap = False
+    backend.disable_chunked_prefill = False
+    backend.batch_max_tokens = budget
+    backend.is_master_in_dp = True
+    backend._timer_merge_radix_tree = lambda: None
+    backend._filter_not_ready_reqs = lambda req_ids: reqs
+    backend._reorder_pd_high_priority_reqs = lambda reqs: reqs
+    backend._reorder_long_prefill_reqs = lambda reqs: reqs
+    for req in reqs:
+        req.args = args
+    copies = []
+    copy_indexes = backend._copy_req_token_indexes_to_gpu
+
+    def record_copy(run_reqs):
+        ranges = [
+            (req.req_idx, req.token_index.gpu_index_submitted_len, req.hold_kv_len)
+            for req in run_reqs
+            if req.token_index.gpu_index_submitted_len < req.hold_kv_len
+        ]
+        if ranges:
+            copies.append(ranges)
+        copy_indexes(run_reqs)
+
+    backend._copy_req_token_indexes_to_gpu = record_copy
+    return backend, stream, copies
+
+
+def _make_scheduled_req(row, total_len=12, gpu_index_submitted_len=4, hold_kv_len=4):
+    req = InferReq.__new__(InferReq)
+    req.req_idx = req.req_id = row
+    req.cur_kv_len = 3
+    req.hold_kv_len = hold_kv_len
+    bind_cpu_indexes(req, infer_batch.g_infer_context.req_manager, gpu_index_submitted_len=gpu_index_submitted_len)
+    req.mtp_step = 0
+    req.filter_mark = req.wait_pause = req.paused = req.infer_aborted = False
+    req.finish_status = infer_batch.FinishStatus()
+    req.shm_req = SimpleNamespace(input_len=total_len)
+    req.is_slave_req = lambda: False
+    req.get_cur_total_len = lambda: total_len
+    req.get_chuncked_input_token_len = lambda: min(total_len, req.cur_kv_len + 3)
+    return req
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_scheduler_submits_native_batch_on_upload_stream_and_returns_to_default_stream(monkeypatch):
+    from cuda.bindings import runtime as cuda_runtime
+
+    manager = _make_manager("cuda")
+    req = _make_scheduled_req(0, gpu_index_submitted_len=0)
+    backend, stream, _ = _make_scheduler(monkeypatch, manager, [req])
+    submitted_streams = []
+    native_copy = cuda_runtime.cudaMemcpyBatchAsync
+
+    def record_batch(*args):
+        submitted_streams.append(args[-1])
+        return native_copy(*args)
+
+    monkeypatch.setattr(cuda_runtime, "cudaMemcpyBatchAsync", record_batch)
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        backend._get_classed_reqs(req_ids=[0])
+        assert torch.cuda.current_stream().cuda_stream == 0
+    torch.cuda.default_stream().synchronize()
+
+    assert len(submitted_streams) == 1
+    assert submitted_streams[0] not in (0, stream.cuda_stream)
+    assert req.token_index.gpu_index_submitted_len == req.hold_kv_len
+    assert torch.equal(
+        manager.req_to_token_indexs[0, : req.hold_kv_len].cpu(), manager.token_indexes_cpu[0, : req.hold_kv_len]
+    )
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_scheduler_uploads_only_new_tail_and_reuses_uploaded_page(monkeypatch):
+    manager = _make_manager("cuda")
+    manager.token_indexes_cpu[0, :4].copy_(torch.arange(4, dtype=torch.int32))
+    manager.req_to_token_indexs[0, :4].copy_(manager.token_indexes_cpu[0, :4])
+    req = _make_scheduled_req(0)
+    backend, stream, copies = _make_scheduler(monkeypatch, manager, [req])
+
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        prefill, decode = backend._get_classed_reqs(req_ids=[0])
+    torch.cuda.default_stream().synchronize()
+    assert prefill == [req] and decode == []
+    assert copies == [[(0, 4, 8)]]
+    assert req.token_index.gpu_index_submitted_len == req.hold_kv_len == 8
+    assert manager.req_to_token_indexs[0, :8].cpu().tolist() == [0, 1, 2, 3, 100, 101, 102, 103]
+
+    req.cur_kv_len = 5  # The next chunk still fits into the reserved page.
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        backend._get_classed_reqs(req_ids=[0])
+    assert copies == [[(0, 4, 8)]]
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_scheduler_leaves_budget_skipped_request_unpublished(monkeypatch):
+    manager = _make_manager("cuda")
+    manager.token_indexes_cpu[0, :8].copy_(torch.arange(8, dtype=torch.int32))
+    req = _make_scheduled_req(0, gpu_index_submitted_len=0, hold_kv_len=8)
+    backend, stream, copies = _make_scheduler(monkeypatch, manager, [req], budget=2)
+
+    assert backend._get_classed_reqs(req_ids=[0]) == ([], [])
+    assert req.token_index.gpu_index_submitted_len == 0
+    assert copies == []
+    assert manager.req_to_token_indexs.eq(-1).all()
+
+    backend.batch_max_tokens = 3
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        backend._get_classed_reqs(req_ids=[0])
+    torch.cuda.default_stream().synchronize()
+    assert copies == [[(0, 0, 8)]]
+    assert req.token_index.gpu_index_submitted_len == 8
+
+
+@pytest.mark.parametrize("shared_prefix_len", [0, 4])
+@pytest.mark.parametrize("is_master_in_dp", [False, True])
+def test_pd_transfer_failure_defers_owned_tail_cleanup(monkeypatch, shared_prefix_len, is_master_in_dp):
+    from lightllm.server.router.model_infer.mode_backend.pd.decode_node_impl import decode_impl
+
+    manager = _make_manager()
+    manager.token_indexes_cpu[0, :12].copy_(torch.arange(12, dtype=torch.int32))
+    shared_node = SimpleNamespace(node_prefix_total_len=shared_prefix_len) if shared_prefix_len else None
+    simulated_finishes = []
+    req = SimpleNamespace(
+        req_idx=0,
+        req_id=0,
+        infer_aborted=False,
+        pd_task_num=2,
+        pd_task_failed_num=1,
+        pd_task_success_num=0,
+        finish_status=infer_batch.FinishStatus(),
+        shared_kv_node=shared_node,
+        cur_kv_len=10,
+        hold_kv_len=12,
+        cur_output_len=3,
+        cache_tiers=(infer_batch.CacheTier.GPU,),
+        shm_req=SimpleNamespace(
+            shm_cur_kv_len=10,
+            mark_simulated_finished=lambda status, output_len: simulated_finishes.append((status, output_len)),
+        ),
+        get_input_token_ids=lambda: list(range(12)),
+    )
+    bind_cpu_indexes(req, manager, gpu_index_submitted_len=12)
+    original_token_index_len = req.token_index.token_index_len
+    original_buffer = req.token_index._cpu_buffer
+    monkeypatch.setattr(decode_impl, "g_infer_context", SimpleNamespace(requests_mapping={0: req}))
+    freed = []
+    manager.mem_manager = SimpleNamespace(free=lambda indexes: freed.extend(indexes.tolist()))
+    backend = decode_impl.PDDecodeNode.__new__(decode_impl.PDDecodeNode)
+    backend.is_master_in_dp = is_master_in_dp
+    backend.model = SimpleNamespace(req_manager=manager, mem_manager=manager.mem_manager)
+
+    # 任意传输任务仍未结束时，不能结束请求、回退逻辑长度或释放物理容量。
+    assert backend._filter_not_ready_reqs([0]) == []
+    assert not req.finish_status.is_finished()
+    assert req.cur_kv_len == req.shm_req.shm_cur_kv_len == 10
+    assert freed == simulated_finishes == []
+
+    req.pd_task_success_num = 1
+    for _ in range(2):
+        assert backend._filter_not_ready_reqs([0]) == [req]
+        assert req.finish_status.get_status() == infer_batch.FinishStatus.FINISHED_ERROR
+        assert req.cur_kv_len == shared_prefix_len
+        assert req.shm_req.shm_cur_kv_len == (shared_prefix_len if is_master_in_dp else 10)
+        assert req.hold_kv_len == req.token_index.gpu_index_submitted_len == 12
+        assert req.token_index.token_index_len == original_token_index_len
+        assert req.token_index._cpu_buffer is original_buffer
+        assert freed == []
+    assert simulated_finishes == ([(infer_batch.FinishStatus.FINISHED_ERROR, 3)] if is_master_in_dp else [])
+
+    # 走真实统一释放路径：只保留可信共享前缀，失败的完整页和页内预留尾部都应回收。
+    cached, released_refs = [], []
+
+    def insert(key, indexes):
+        cached.append((key, indexes))
+        return len(key), shared_node
+
+    context = InferenceContext()
+    context.args = SimpleNamespace(page_size=4)
+    context.req_manager = manager
+    context.radix_cache = SimpleNamespace(insert=insert, dec_node_ref_counter=released_refs.append)
+    released = []
+    context.free_a_req_mem(released, req)
+    manager.free_token(torch.cat(released))
+
+    assert freed == list(range(shared_prefix_len, 12))
+    assert len(cached) == 1
+    assert cached[0][0].tolist() == cached[0][1].tolist() == list(range(shared_prefix_len))
+    assert released_refs == ([shared_node] if shared_node is not None else [])
+    assert req.shared_kv_node is None
+    assert req.cur_kv_len == req.hold_kv_len == req.shm_req.shm_cur_kv_len == 0
+    assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+    assert req.token_index._cpu_buffer is None
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("mixed", [False, True])
+def test_scheduler_publishes_prefill_and_decode_before_return(monkeypatch, mixed):
+    manager = _make_manager("cuda")
+    for row in (0, 1):
+        manager.token_indexes_cpu[row, :4].copy_(torch.arange(row * 4, row * 4 + 4, dtype=torch.int32))
+    prefill_req = _make_scheduled_req(0, gpu_index_submitted_len=0)
+    decode_req = _make_scheduled_req(1, total_len=4, gpu_index_submitted_len=0)
+    backend, stream, copies = _make_scheduler(monkeypatch, manager, [prefill_req, decode_req], mixed=mixed)
+
+    with torch.cuda.stream(torch.cuda.default_stream()):
+        prefill, decode = backend._get_classed_reqs(req_ids=[0, 1])
+    torch.cuda.default_stream().synchronize()
+
+    assert copies == [[(0, 0, 8), (1, 0, 4)]]
+    assert prefill_req.token_index.gpu_index_submitted_len == 8
+    assert decode_req.token_index.gpu_index_submitted_len == 4
+    assert prefill == ([prefill_req, decode_req] if mixed else [prefill_req])
+    assert decode == ([] if mixed else [decode_req])
+    for req in (prefill_req, decode_req):
+        assert torch.equal(
+            manager.req_to_token_indexs[req.req_idx, : req.hold_kv_len].cpu(),
+            manager.token_indexes_cpu[req.req_idx, : req.hold_kv_len],
+        )
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("loop_fails", [False, True])
+def test_request_upload_lengths_follow_fallback_copy_result(monkeypatch, loop_fails):
+    from cuda.bindings import runtime as cuda_runtime
+
+    manager = _make_manager("cuda")
+    reqs = [
+        _make_scheduled_req(0, gpu_index_submitted_len=0),
+        _make_scheduled_req(1, total_len=4, gpu_index_submitted_len=0),
+    ]
+    backend, stream, _ = _make_scheduler(monkeypatch, manager, reqs)
+    monkeypatch.setattr(
+        cuda_runtime,
+        "cudaMemcpyBatchAsync",
+        lambda *args: (cuda_runtime.cudaError_t.cudaErrorInvalidValue,),
+    )
+
+    if loop_fails:
+        torch_copy = torch.Tensor.copy_
+
+        def failed_copy(dst, src, *args, **kwargs):
+            if dst.is_cuda and src.device.type == "cpu":
+                raise RuntimeError("individual copy failed")
+            return torch_copy(dst, src, *args, **kwargs)
+
+        monkeypatch.setattr(torch.Tensor, "copy_", failed_copy)
+        with torch.cuda.stream(torch.cuda.default_stream()), pytest.raises(
+            RuntimeError, match="individual copy failed"
+        ):
+            backend._get_classed_reqs(req_ids=[0, 1])
+        assert [req.token_index.gpu_index_submitted_len for req in reqs] == [0, 0]
+    else:
+        with torch.cuda.stream(torch.cuda.default_stream()):
+            backend._get_classed_reqs(req_ids=[0, 1])
+        torch.cuda.default_stream().synchronize()
+        for req in reqs:
+            assert req.token_index.gpu_index_submitted_len == req.hold_kv_len
+            assert torch.equal(
+                manager.req_to_token_indexs[req.req_idx, : req.hold_kv_len].cpu(),
+                manager.token_indexes_cpu[req.req_idx, : req.hold_kv_len],
+            )
+
+
+@pytest.mark.skipif(torch.cuda.device_count() < 2, reason="requires two CUDA devices")
+@pytest.mark.parametrize("batch_available", [False, True])
+def test_copy_uses_current_stream_on_current_device(monkeypatch, batch_available):
+    from lightllm.common import cuda_wrapper as base
+
+    monkeypatch.setattr(base.PinnedH2DBatchCopier, "_native_batch_available", batch_available)
+    manager = _make_manager("cuda:1")
+    manager.token_indexes_cpu[0, :4].copy_(torch.arange(4, dtype=torch.int32))
+    stream = torch.cuda.Stream(device=1)
+    stream.wait_stream(torch.cuda.current_stream(1))
+    with torch.cuda.device(0):
+        with torch.cuda.stream(stream):
+            cuda_memcpy_pinned_h2d_batch_async(
+                [manager.req_to_token_indexs[0, :4]],
+                [manager.token_indexes_cpu[0, :4]],
+            )
+            assert torch.cuda.current_device() == 1
+            assert torch.cuda.current_stream().cuda_stream == stream.cuda_stream
+        assert torch.cuda.current_device() == 0
+    stream.synchronize()
+    assert manager.req_to_token_indexs[0, :4].cpu().tolist() == list(range(4))
diff --git a/unit_tests/common/test_req_manager_page.py b/unit_tests/common/test_req_manager_page.py
index f45ac26ceb..3bc4314847 100644
--- a/unit_tests/common/test_req_manager_page.py
+++ b/unit_tests/common/test_req_manager_page.py
@@ -3,11 +3,13 @@
 import pytest
 import torch
 
-from lightllm.common.req_manager import ReqManager
 from lightllm.server.router.model_infer.mode_backend import generic_pre_process
+from lightllm.common.req_manager import ReqManager
 from lightllm.server.router.model_infer import infer_batch
 from lightllm.server.router.model_infer.infer_batch import InferReq, InferenceContext
 from lightllm.server.router.model_infer.mode_backend import base_backend
+from lightllm.server.multi_level_kv_cache import CacheTier
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
 
 
 class _FakeMemManager:
@@ -28,6 +30,7 @@ def test_bind_mem_manager_initializes_hold_request_with_reserved_page():
     req_manager = ReqManager.__new__(ReqManager)
     req_manager.HOLD_REQUEST_ID = 2
     req_manager.req_list = SimpleNamespace(is_all_free=lambda: True)
+    req_manager.page_size = 4
     req_manager.req_to_token_indexs = torch.full((3, 8), -1, dtype=torch.int32)
     mem_manager = SimpleNamespace(page_size=4, HOLD_TOKEN_MEMINDEXES=(100, 101, 102, 103))
 
@@ -48,94 +51,89 @@ def test_hold_request_indexes_require_all_requests_released():
 
 def _make_context(monkeypatch):
     mem_manager = _FakeMemManager()
-    req_manager = SimpleNamespace(
-        mem_manager=mem_manager,
-        req_to_token_indexs=torch.full((2, 16), -1, dtype=torch.int32),
-    )
+    req_manager = make_cpu_index_manager(torch.full((2, 16), -1, dtype=torch.int32), page_size=4)
+    req_manager.mem_manager = mem_manager
+    req_manager.max_request_num = 2
+    req_manager.page_size = 4
+
     context = SimpleNamespace(
         args=SimpleNamespace(page_size=4),
         req_manager=req_manager,
         radix_cache=None,
+        get_overlap_stream=lambda: None,
     )
     monkeypatch.setattr(generic_pre_process, "g_infer_context", context)
     monkeypatch.setattr(base_backend, "g_infer_context", context)
+    monkeypatch.setattr(infer_batch, "g_infer_context", context)
     backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
     backend.args = context.args
+
+    def copy_indexes(run_reqs):
+        for req in run_reqs:
+            start, end = req.token_index.gpu_index_submitted_len, req.hold_kv_len
+            req_manager.req_to_token_indexs[req.req_idx, start:end].copy_(
+                req.token_index.get_cpu_token_indexes(start, end)
+            )
+            req.token_index.gpu_index_submitted_len = end
+
+    backend._copy_req_token_indexes_to_gpu = copy_indexes
     return context, backend
 
 
 def _make_req(req_idx):
-    req = SimpleNamespace(req_idx=req_idx, cur_kv_len=0, hold_kv_len=0)
+    req = InferReq.__new__(InferReq)
+    req.req_idx = req_idx
+    req.args = infer_batch.g_infer_context.args
+    bind_cpu_indexes(req, infer_batch.g_infer_context.req_manager)
+    req.cur_kv_len = req.hold_kv_len = 0
     req._kv_cache_alloc_need = lambda target_len: InferReq._kv_cache_alloc_need(req, target_len)
     return req
 
 
 def test_request_reuses_reserved_page_tail_before_allocating_next_page(monkeypatch):
-    context, backend = _make_context(monkeypatch)
+    context, _ = _make_context(monkeypatch)
     req = _make_req(0)
 
-    mem_indexes = backend._alloc_req_kv_mem(req, alloc_token_num=4)
+    mem_indexes = req.alloc_kv_mem(alloc_token_num=4)
     assert mem_indexes.tolist() == [0, 1, 2, 3]
     assert req.hold_kv_len == 4
     assert context.req_manager.mem_manager.alloc_sizes == [4]
-    assert context.req_manager.req_to_token_indexs[0, :4].tolist() == [0, 1, 2, 3]
+    assert context.req_manager.get_cpu_token_indexes(0, 0, 4).tolist() == [0, 1, 2, 3]
 
     req.cur_kv_len = 3
-    assert backend._alloc_req_kv_mem(req, alloc_token_num=0) is None
+    assert req.alloc_kv_mem(alloc_token_num=0) is None
     assert context.req_manager.mem_manager.alloc_sizes == [4]
 
     req.cur_kv_len = 4
-    backend._alloc_req_kv_mem(req, alloc_token_num=4)
+    req.alloc_kv_mem(alloc_token_num=4)
     assert req.hold_kv_len == 8
     assert context.req_manager.mem_manager.alloc_sizes == [4, 4]
-    assert context.req_manager.req_to_token_indexs[0, :8].tolist() == list(range(8))
+    assert context.req_manager.get_cpu_token_indexes(0, 0, 8).tolist() == list(range(8))
 
 
 def test_reservation_fills_each_request_table_row(monkeypatch):
-    _, backend = _make_context(monkeypatch)
+    _make_context(monkeypatch)
     req0 = _make_req(0)
     req1 = _make_req(1)
 
-    backend._alloc_req_kv_mem(req0, alloc_token_num=4)
-    backend._alloc_req_kv_mem(req1, alloc_token_num=4)
+    req0.alloc_kv_mem(alloc_token_num=4)
+    req1.alloc_kv_mem(alloc_token_num=4)
 
-    assert generic_pre_process.g_infer_context.req_manager.req_to_token_indexs[0, :4].tolist() == [0, 1, 2, 3]
-    assert generic_pre_process.g_infer_context.req_manager.req_to_token_indexs[1, :4].tolist() == [4, 5, 6, 7]
+    assert generic_pre_process.g_infer_context.req_manager.get_cpu_token_indexes(0, 0, 4).tolist() == [0, 1, 2, 3]
+    assert generic_pre_process.g_infer_context.req_manager.get_cpu_token_indexes(1, 0, 4).tolist() == [4, 5, 6, 7]
     assert req0.hold_kv_len == req1.hold_kv_len == 4
 
 
-def test_alloc_req_kv_mem_forwards_non_blocking_copy_option(monkeypatch):
-    class _CopyTarget:
-        def __init__(self):
-            self.non_blocking_values = []
-
-        def copy_(self, source, non_blocking=False):
-            self.non_blocking_values.append(non_blocking)
-
-    class _ReqToTokenIndexes:
-        def __init__(self, target):
-            self.target = target
-
-        def __getitem__(self, key):
-            return self.target
-
-    target = _CopyTarget()
-    context = SimpleNamespace(
-        req_manager=SimpleNamespace(
-            mem_manager=_FakeMemManager(),
-            req_to_token_indexs=_ReqToTokenIndexes(target),
-        ),
-        radix_cache=None,
-    )
-    monkeypatch.setattr(base_backend, "g_infer_context", context)
-    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
-    backend.args = SimpleNamespace(page_size=4)
+def test_alloc_kv_mem_defers_gpu_table_update(monkeypatch):
+    context, _ = _make_context(monkeypatch)
     req = _make_req(0)
 
-    backend._alloc_req_kv_mem(req, alloc_token_num=4)
-    backend._alloc_req_kv_mem(req, alloc_token_num=4, no_blcoking_copy=True)
+    req.alloc_kv_mem(alloc_token_num=4)
+    req.alloc_kv_mem(alloc_token_num=4)
 
-    assert target.non_blocking_values == [False, True]
+    assert context.req_manager.get_cpu_token_indexes(0, 0, 8).tolist() == list(range(8))
+    assert context.req_manager.req_to_token_indexs.eq(-1).all()
+    assert req.token_index.gpu_index_submitted_len == 0
 
 
 def test_need_token_num_distinguishes_compute_and_page_allocation():
@@ -207,24 +205,11 @@ def test_prefill_scheduler_checks_compute_and_kv_budgets_separately(
         req.get_chuncked_input_token_len = lambda: 6
         req.prefill_need_token_num = MethodType(InferReq.prefill_need_token_num, req)
         start_index = req.req_idx * hold_kv_len
-        context.req_manager.req_to_token_indexs[req.req_idx, :hold_kv_len] = torch.arange(
-            start_index, start_index + hold_kv_len, dtype=torch.int32
+        context.req_manager.append_token_indexes(
+            req.req_idx, torch.arange(start_index, start_index + hold_kv_len, dtype=torch.int32)
         )
     context.req_manager.mem_manager.next_index = 2 * hold_kv_len
     backend._filter_not_ready_reqs = lambda req_ids: reqs
-    copy_modes = []
-    original_alloc_req_kv_mem = backend._alloc_req_kv_mem
-
-    def _record_alloc(req_obj, alloc_token_num, no_blcoking_copy=False):
-        copy_modes.append(no_blcoking_copy)
-        return original_alloc_req_kv_mem(
-            req_obj,
-            alloc_token_num,
-            no_blcoking_copy=no_blcoking_copy,
-        )
-
-    backend._alloc_req_kv_mem = _record_alloc
-
     prefill_reqs, decode_reqs = backend._get_classed_reqs(req_ids=[0, 1])
 
     assert prefill_reqs == [reqs[0]]
@@ -234,10 +219,9 @@ def _record_alloc(req_obj, alloc_token_num, no_blcoking_copy=False):
     assert not reqs[0].wait_pause
     assert reqs[1].hold_kv_len == hold_kv_len
     assert reqs[1].wait_pause is second_req_wait_pause
-    assert copy_modes == [True]
 
 
-def test_decode_scheduler_uses_non_blocking_request_table_copy(monkeypatch):
+def test_decode_scheduler_reserves_cpu_request_table(monkeypatch):
     context, backend = _make_context(monkeypatch)
     backend.args.enable_cpu_cache = False
     backend.args.enable_prefill_decode_mixed = False
@@ -265,30 +249,113 @@ def test_decode_scheduler_uses_non_blocking_request_table_copy(monkeypatch):
     req.finish_status = infer_batch.FinishStatus()
     req.get_cur_total_len = lambda: 5
     req.decode_need_token_num = MethodType(InferReq.decode_need_token_num, req)
-    context.req_manager.req_to_token_indexs[0, :4] = torch.arange(4, dtype=torch.int32)
+    context.req_manager.append_token_indexes(0, torch.arange(4, dtype=torch.int32))
     context.req_manager.mem_manager.next_index = 4
     backend._filter_not_ready_reqs = lambda req_ids: [req]
 
-    copy_modes = []
-    original_alloc_req_kv_mem = backend._alloc_req_kv_mem
-
-    def _record_alloc(req_obj, alloc_token_num, no_blcoking_copy=False):
-        copy_modes.append(no_blcoking_copy)
-        return original_alloc_req_kv_mem(
-            req_obj,
-            alloc_token_num,
-            no_blcoking_copy=no_blcoking_copy,
-        )
-
-    backend._alloc_req_kv_mem = _record_alloc
-
     prefill_reqs, decode_reqs = backend._get_classed_reqs(req_ids=[0])
 
     assert prefill_reqs == []
     assert decode_reqs == [req]
     assert req.hold_kv_len == 16
     assert context.req_manager.mem_manager.alloc_sizes == [12]
-    assert copy_modes == [True]
+
+
+def _make_capacity_scheduler(monkeypatch, page_size, available, run_mode, support_overlap=False):
+    context, backend = _make_context(monkeypatch)
+    context.args.page_size = context.req_manager.page_size = page_size
+    context.req_manager.token_indexes_cpu = torch.full((2, 256), -1, dtype=torch.int32)
+    context.req_manager.req_to_token_indexs = torch.full((2, 256), -1, dtype=torch.int32)
+    backend.args.enable_cpu_cache = False
+    backend.args.enable_prefill_decode_mixed = False
+    backend.args.run_mode = run_mode
+    backend.support_overlap = support_overlap
+    backend.disable_chunked_prefill = True
+    backend.batch_max_tokens = 256
+    backend.is_master_in_dp = True
+    backend.logger = SimpleNamespace(info=lambda message: None)
+    backend._timer_merge_radix_tree = lambda: None
+    backend._reorder_pd_high_priority_reqs = lambda reqs: reqs
+    backend._reorder_long_prefill_reqs = lambda reqs: reqs
+    context.get_can_alloc_token_num = lambda: available
+    context.cache_placement_controller = SimpleNamespace(set_req_cache_way=lambda reqs: None)
+    context.finished_reqs = []
+    context.filter_reqs = lambda finished_reqs: context.finished_reqs.extend(finished_reqs)
+    context.pause_reqs = lambda reqs, is_master_in_dp: None
+    return context, backend
+
+
+def _make_capacity_req(req_idx, context, is_decode=True):
+    req = _make_req(req_idx)
+    req.req_id = req_idx
+    req.args = context.args
+    req.cur_kv_len = req.hold_kv_len = context.args.page_size
+    req.mtp_step = 0
+    req.filter_mark = req.wait_pause = req.paused = req.infer_aborted = False
+    req.finish_status = infer_batch.FinishStatus()
+    req.cpu_cache_task_status = SimpleNamespace(is_not_started=lambda: True)
+    req.is_slave_req = lambda: False
+    req.get_cur_total_len = lambda: req.cur_kv_len + (1 if is_decode else 2 * context.args.page_size)
+    req.decode_need_token_num = MethodType(InferReq.decode_need_token_num, req)
+    req.prefill_need_token_num = MethodType(InferReq.prefill_need_token_num, req)
+    context.req_manager.append_token_indexes(req_idx, torch.arange(req.hold_kv_len, dtype=torch.int32))
+    return req
+
+
+@pytest.mark.parametrize("page_size", [1, 3, 4, 7, 8, 64])
+@pytest.mark.parametrize("run_mode", ["normal", "decode"])
+@pytest.mark.parametrize("extra_budget", [0, 1, 8])
+@pytest.mark.parametrize("request_types", [(True,), (True, True), (True, False), (False, True)])
+def test_decode_headroom_only_uses_budget_after_all_required_allocations(
+    monkeypatch, page_size, run_mode, extra_budget, request_types
+):
+    required = sum(page_size if is_decode else 2 * page_size for is_decode in request_types)
+    context, backend = _make_capacity_scheduler(monkeypatch, page_size, required + extra_budget, run_mode)
+    reqs = [_make_capacity_req(index, context, is_decode) for index, is_decode in enumerate(request_types)]
+    context.req_manager.mem_manager.next_index = len(reqs) * page_size
+    backend._filter_not_ready_reqs = lambda req_ids: reqs
+
+    prefill, decode = backend._get_classed_reqs(req_ids=list(range(len(reqs))))
+
+    assert prefill == [req for req, is_decode in zip(reqs, request_types) if not is_decode]
+    assert decode == [req for req, is_decode in zip(reqs, request_types) if is_decode]
+    assert context.finished_reqs == []
+    remaining = extra_budget
+    for req, is_decode in zip(reqs, request_types):
+        extra = min(8, remaining) // page_size * page_size if is_decode and page_size < 8 else 0
+        remaining -= extra
+        assert req.hold_kv_len == page_size + (page_size if is_decode else 2 * page_size) + extra
+        assert req.token_index.token_index_len == req.hold_kv_len
+        assert req.token_index.gpu_index_submitted_len == req.hold_kv_len
+        assert not req.wait_pause and not req.filter_mark
+        assert not hasattr(req, "finished_by_pd_decode_capacity")
+    # 必需量和可选预留合并申请，不额外增加分配次数，且不会超出总预算。
+    allocations = context.req_manager.mem_manager.alloc_sizes
+    assert len(allocations) == len(reqs)
+    assert sum(allocations) == required + extra_budget - remaining
+
+
+@pytest.mark.parametrize("page_size", [1, 4, 64])
+@pytest.mark.parametrize("run_mode", ["normal", "decode"])
+@pytest.mark.parametrize("support_overlap", [False, True])
+def test_decode_still_rejects_requests_when_required_capacity_is_insufficient(
+    monkeypatch, page_size, run_mode, support_overlap
+):
+    context, backend = _make_capacity_scheduler(monkeypatch, page_size, page_size - 1, run_mode, support_overlap)
+    req = _make_capacity_req(0, context)
+    backend._filter_not_ready_reqs = lambda req_ids: [req]
+
+    assert backend._get_classed_reqs(req_ids=[0]) == ([], [])
+
+    assert context.req_manager.mem_manager.alloc_sizes == []
+    assert req.hold_kv_len == page_size
+    if run_mode == "normal":
+        assert req.wait_pause
+        assert context.finished_reqs == []
+    else:
+        assert req.finished_by_pd_decode_capacity
+        assert req.filter_mark is support_overlap
+        assert context.finished_reqs == ([] if support_overlap else [req])
 
 
 @pytest.mark.parametrize(
@@ -300,6 +367,8 @@ def test_decode_small_page_preallocation_reuses_one_allocation_for_multiple_step
 ):
     context, backend = _make_context(monkeypatch)
     context.args.page_size = page_size
+    context.req_manager.page_size = page_size
+    context.req_manager.token_indexes_cpu = torch.full((2, (64 + page_size - 1) // page_size), -1, dtype=torch.int32)
     context.req_manager.req_to_token_indexs = torch.full((2, 64), -1, dtype=torch.int32)
     backend.args.enable_cpu_cache = False
     backend.args.enable_prefill_decode_mixed = False
@@ -327,7 +396,7 @@ def test_decode_small_page_preallocation_reuses_one_allocation_for_multiple_step
     req.finish_status = infer_batch.FinishStatus()
     req.get_cur_total_len = lambda: req.cur_kv_len + 1
     req.decode_need_token_num = MethodType(InferReq.decode_need_token_num, req)
-    context.req_manager.req_to_token_indexs[0, :page_size] = torch.arange(page_size, dtype=torch.int32)
+    context.req_manager.append_token_indexes(0, torch.arange(page_size, dtype=torch.int32))
     context.req_manager.mem_manager.next_index = page_size
     backend._filter_not_ready_reqs = lambda req_ids: [req]
 
@@ -341,7 +410,7 @@ def test_decode_small_page_preallocation_reuses_one_allocation_for_multiple_step
 
 
 def test_decode_reserves_mtp_headroom(monkeypatch):
-    context, backend = _make_context(monkeypatch)
+    context, _ = _make_context(monkeypatch)
     req = _make_req(0)
     req.cur_kv_len = 3
     req.hold_kv_len = 4
@@ -352,12 +421,12 @@ def test_decode_reserves_mtp_headroom(monkeypatch):
     req.get_radix_cache_shared_len = lambda: 0
     req.args = context.args
 
-    context.req_manager.req_to_token_indexs[0, :4] = torch.arange(4, dtype=torch.int32)
+    context.req_manager.append_token_indexes(0, torch.arange(4, dtype=torch.int32))
     context.req_manager.mem_manager.next_index = 4
     token_num, alloc_token_num = InferReq.decode_need_token_num(req)
     assert token_num == 6
     assert alloc_token_num == 8
-    backend._alloc_req_kv_mem(req, alloc_token_num)
+    req.alloc_kv_mem(alloc_token_num)
 
     model_input, run_reqs = generic_pre_process.prepare_decode_inputs([req])
 
@@ -365,12 +434,14 @@ def test_decode_reserves_mtp_headroom(monkeypatch):
     assert model_input.b_seq_len.tolist() == [4, 5, 6]
     assert req.hold_kv_len == 12
     assert context.req_manager.mem_manager.alloc_sizes == [8]
-    assert context.req_manager.req_to_token_indexs[0, :12].tolist() == list(range(12))
+    assert context.req_manager.get_cpu_token_indexes(0, 0, 12).tolist() == list(range(12))
 
 
 def test_page_size_one_uses_the_same_scheduler_preallocation(monkeypatch):
-    context, backend = _make_context(monkeypatch)
+    context, _ = _make_context(monkeypatch)
     context.args.page_size = 1
+    context.req_manager.page_size = 1
+    context.req_manager.token_indexes_cpu = torch.full((2, 16), -1, dtype=torch.int32)
     req = _make_req(0)
     req.cur_kv_len = 3
     req.hold_kv_len = 3
@@ -380,43 +451,149 @@ def test_page_size_one_uses_the_same_scheduler_preallocation(monkeypatch):
     req.get_cur_total_len = lambda: 4
     req.get_radix_cache_shared_len = lambda: 0
     req.args = context.args
-    context.req_manager.req_to_token_indexs[0, :3] = torch.arange(3, dtype=torch.int32)
+    context.req_manager.append_token_indexes(0, torch.arange(3, dtype=torch.int32))
     context.req_manager.mem_manager.next_index = 3
     token_num, alloc_token_num = InferReq.decode_need_token_num(req)
     assert token_num == alloc_token_num == 6
-    backend._alloc_req_kv_mem(req, alloc_token_num)
+    req.alloc_kv_mem(alloc_token_num)
 
     model_input, _ = generic_pre_process.prepare_decode_inputs([req])
 
     assert req.hold_kv_len == 9
     assert context.req_manager.mem_manager.alloc_sizes == [6]
-    assert context.req_manager.req_to_token_indexs[0, :9].tolist() == list(range(9))
+    assert context.req_manager.get_cpu_token_indexes(0, 0, 9).tolist() == list(range(9))
 
 
 def test_page_size_one_frees_all_preallocated_indexes():
     infer_context = InferenceContext.__new__(InferenceContext)
     infer_context.args = SimpleNamespace(page_size=1)
     infer_context.radix_cache = None
-    infer_context.req_manager = SimpleNamespace(req_to_token_indexs=torch.arange(16, dtype=torch.int32)[None, :])
+    infer_context.req_manager = make_cpu_index_manager(
+        token_table=torch.arange(16, dtype=torch.int32)[None, :], page_size=1
+    )
     req = SimpleNamespace(
         req_idx=0,
         cur_kv_len=3,
         hold_kv_len=12,
         shm_req=SimpleNamespace(shm_cur_kv_len=3),
     )
+    bind_cpu_indexes(req, infer_context.req_manager)
     free_token_indexes = []
 
     infer_context.free_a_req_mem(free_token_indexes, req)
 
     assert free_token_indexes[0].tolist() == list(range(12))
     assert req.cur_kv_len == req.hold_kv_len == 0
+    assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+
+
+@pytest.mark.parametrize("attention", ["full", "hybrid_small", "hybrid_big"])
+def test_cache_release_preserves_snapshots_and_reclaims_tail_duplicates_and_reference(monkeypatch, attention):
+    hybrid = attention != "full"
+    context = InferenceContext()
+    context.args = SimpleNamespace(page_size=4, linear_att_hash_page_size=4, linear_att_page_block_num=2)
+    context.is_hybrid_att_model = hybrid
+    context.req_manager = make_cpu_index_manager(torch.arange(16, dtype=torch.int32)[None, :], page_size=4)
+    cached = []
+    released_refs = []
+
+    def insert(key, indexes, *, block_hashs=None, block_state_idxs=None, len_to_big_page_id=None):
+        cached.append((key, indexes, block_hashs, block_state_idxs))
+        # 模拟 radix cache 接管请求已生成的大页 checkpoint。
+        if hybrid:
+            len_to_big_page_id.clear()
+        return 8, None
+
+    context.radix_cache = SimpleNamespace(insert=insert, dec_node_ref_counter=released_refs.append)
+    shared_node = SimpleNamespace(node_prefix_total_len=4)
+    req = SimpleNamespace(
+        req_idx=0,
+        cur_kv_len=13 if hybrid else 11,
+        hold_kv_len=16,
+        cache_tiers=(CacheTier.GPU,),
+        shared_kv_node=shared_node,
+        hybrid_cache_len=12,
+        tail_small_page_buffer_id=7 if attention == "hybrid_small" else None,
+        hybrid_len_to_big_page_id={8: 9} if hybrid else {},
+        get_input_token_ids=lambda: list(range(13)),
+        shm_req=SimpleNamespace(
+            shm_cur_kv_len=13,
+            hybrid_token_hash_list=SimpleNamespace(get_all=lambda: [100, 200, 300]),
+        ),
+    )
+    bind_cpu_indexes(req, context.req_manager)
+    monkeypatch.setattr(infer_batch, "g_infer_context", context)
+    monkeypatch.setattr(infer_batch, "get_env_start_args", lambda: context.args)
+    released = []
+
+    context.free_a_req_mem(released, req)
+
+    cache_len = 12 if attention == "hybrid_small" else 8
+    assert [indexes.tolist() for indexes in released] == [list(range(cache_len, 16)), list(range(4, 8))]
+    assert released_refs == [shared_node]
+    assert req.shared_kv_node is None
+    assert req.cur_kv_len == req.hold_kv_len == req.shm_req.shm_cur_kv_len == 0
+    assert req.token_index.token_index_len == req.token_index.gpu_index_submitted_len == 0
+    assert req.tail_small_page_buffer_id is None and req.hybrid_len_to_big_page_id == {}
+    context.req_manager.token_indexes_cpu.fill_(-99)
+    key, indexes, block_hashs, block_state_idxs = cached[0]
+    assert key.tolist() == indexes.tolist() == list(range(cache_len))
+    if hybrid:
+        assert block_hashs == [100, 200, 300][: cache_len // 4]
+        assert block_state_idxs == ([None, None, 7] if attention == "hybrid_small" else [None, None])
+    else:
+        assert block_hashs is None and block_state_idxs is None
+
+
+def _insert_empty_hybrid_prefix(key, indexes, *, block_hashs, block_state_idxs, len_to_big_page_id):
+    assert key.numel() == indexes.numel() == 0
+    assert block_hashs == block_state_idxs == []
+    assert not len_to_big_page_id
+    return 0, None
+
+
+def test_hybrid_release_without_new_big_page_rejects_pending_checkpoints(monkeypatch):
+    context = InferenceContext()
+    context.args = SimpleNamespace(linear_att_hash_page_size=4, linear_att_page_block_num=2)
+    context.is_hybrid_att_model = True
+    context.req_manager = make_cpu_index_manager(torch.arange(8, dtype=torch.int32)[None, :])
+    freed_states = []
+    context.radix_cache = SimpleNamespace(
+        insert=_insert_empty_hybrid_prefix,
+        big_page_buffers=SimpleNamespace(free_state_cache=freed_states.extend),
+    )
+    req = SimpleNamespace(
+        req_idx=0,
+        cur_kv_len=5,
+        hold_kv_len=8,
+        hybrid_cache_len=0,
+        tail_small_page_buffer_id=None,
+        hybrid_len_to_big_page_id={8: 9},
+        shared_kv_node=None,
+        get_input_token_ids=lambda: list(range(8)),
+        shm_req=SimpleNamespace(hybrid_token_hash_list=SimpleNamespace(get_all=lambda: [])),
+    )
+    bind_cpu_indexes(req, context.req_manager)
+    monkeypatch.setattr(infer_batch, "g_infer_context", context)
+    monkeypatch.setattr(infer_batch, "get_env_start_args", lambda: context.args)
+    released = []
+
+    with pytest.raises(AssertionError):
+        context._hybrid_att_free_req(released, req)
+
+    assert released == []
+    assert freed_states == []
+    assert req.hybrid_len_to_big_page_id == {8: 9}
+    assert req.cur_kv_len == 5 and req.hold_kv_len == 8
 
 
 def test_linear_attention_frees_reserved_page_tail(monkeypatch):
     infer_context = InferenceContext.__new__(InferenceContext)
     infer_context.args = SimpleNamespace(linear_att_hash_page_size=4, linear_att_page_block_num=2)
-    infer_context.radix_cache = SimpleNamespace()
-    infer_context.req_manager = SimpleNamespace(req_to_token_indexs=torch.arange(16, dtype=torch.int32)[None, :])
+    infer_context.radix_cache = SimpleNamespace(insert=_insert_empty_hybrid_prefix)
+    infer_context.req_manager = make_cpu_index_manager(
+        token_table=torch.arange(16, dtype=torch.int32)[None, :], page_size=1
+    )
     req = SimpleNamespace(
         req_idx=0,
         cur_kv_len=5,
@@ -425,7 +602,10 @@ def test_linear_attention_frees_reserved_page_tail(monkeypatch):
         tail_small_page_buffer_id=None,
         hybrid_len_to_big_page_id={},
         shared_kv_node=None,
+        get_input_token_ids=lambda: list(range(8)),
+        shm_req=SimpleNamespace(hybrid_token_hash_list=SimpleNamespace(get_all=lambda: [])),
     )
+    bind_cpu_indexes(req, infer_context.req_manager)
     free_token_indexes = []
     monkeypatch.setattr(infer_batch.g_infer_context, "is_hybrid_att_model", True)
     monkeypatch.setattr(infer_batch, "get_env_start_args", lambda: infer_context.args)
@@ -438,8 +618,10 @@ def test_linear_attention_frees_reserved_page_tail(monkeypatch):
 def test_linear_attention_frees_preallocated_pages_before_first_forward(monkeypatch):
     infer_context = InferenceContext.__new__(InferenceContext)
     infer_context.args = SimpleNamespace(linear_att_hash_page_size=4, linear_att_page_block_num=2)
-    infer_context.radix_cache = SimpleNamespace()
-    infer_context.req_manager = SimpleNamespace(req_to_token_indexs=torch.arange(8, dtype=torch.int32)[None, :])
+    infer_context.radix_cache = SimpleNamespace(insert=_insert_empty_hybrid_prefix)
+    infer_context.req_manager = make_cpu_index_manager(
+        token_table=torch.arange(8, dtype=torch.int32)[None, :], page_size=1
+    )
     req = SimpleNamespace(
         req_idx=0,
         cur_kv_len=0,
@@ -448,7 +630,10 @@ def test_linear_attention_frees_preallocated_pages_before_first_forward(monkeypa
         tail_small_page_buffer_id=None,
         hybrid_len_to_big_page_id={},
         shared_kv_node=None,
+        get_input_token_ids=lambda: list(range(8)),
+        shm_req=SimpleNamespace(hybrid_token_hash_list=SimpleNamespace(get_all=lambda: [])),
     )
+    bind_cpu_indexes(req, infer_context.req_manager)
     free_token_indexes = []
     monkeypatch.setattr(infer_batch.g_infer_context, "is_hybrid_att_model", True)
     monkeypatch.setattr(infer_batch, "get_env_start_args", lambda: infer_context.args)
@@ -493,15 +678,15 @@ def test_recover_paused_reqs_uses_page_allocation_need(monkeypatch):
     infer_context.backend = SimpleNamespace(disable_chunked_prefill=False)
     infer_context.radix_cache = None
     freed_indexes = []
-    infer_context.req_manager = SimpleNamespace(
-        req_to_token_indexs=torch.arange(16, dtype=torch.int32).reshape(2, 8),
+    infer_context.req_manager = make_cpu_index_manager(
+        token_table=torch.arange(16, dtype=torch.int32).reshape(2, 8),
         free_token=lambda indexes: freed_indexes.extend(indexes.tolist()),
+        page_size=1,
     )
     infer_context.get_can_alloc_token_num = lambda: 5
     large_req = _make_paused_req(req_idx=0, target_kv_len=5, page_size=4)
     small_req = _make_paused_req(req_idx=1, target_kv_len=4, page_size=4)
     monkeypatch.setattr(infer_batch.g_infer_context, "is_hybrid_att_model", False)
-    monkeypatch.setattr(infer_batch, "custom_cat", lambda tensors: torch.cat(tensors))
 
     infer_context.recover_paused_reqs([large_req, small_req], is_master_in_dp=True)
 
@@ -518,14 +703,12 @@ def test_recover_paused_reqs_requires_capacity_for_the_full_sequence(monkeypatch
     infer_context = InferenceContext.__new__(InferenceContext)
     infer_context.args = SimpleNamespace(page_size=4)
     infer_context.radix_cache = None
-    infer_context.req_manager = SimpleNamespace(
-        req_to_token_indexs=torch.arange(8, dtype=torch.int32)[None, :],
-        free_token=lambda indexes: None,
+    infer_context.req_manager = make_cpu_index_manager(
+        token_table=torch.arange(8, dtype=torch.int32)[None, :], free_token=lambda indexes: None, page_size=1
     )
     infer_context.get_can_alloc_token_num = lambda: 4
     req = _make_paused_req(req_idx=0, target_kv_len=8, page_size=4, chunk_kv_len=4)
     monkeypatch.setattr(infer_batch.g_infer_context, "is_hybrid_att_model", False)
-    monkeypatch.setattr(infer_batch, "custom_cat", lambda tensors: torch.cat(tensors))
 
     infer_context.recover_paused_reqs([req], is_master_in_dp=True)
 
diff --git a/unit_tests/common/test_req_sampling_params.py b/unit_tests/common/test_req_sampling_params.py
new file mode 100644
index 0000000000..c9b7688ca5
--- /dev/null
+++ b/unit_tests/common/test_req_sampling_params.py
@@ -0,0 +1,109 @@
+from types import SimpleNamespace
+
+import numpy as np
+import pytest
+import torch
+
+# 沿用生产导入顺序，避免 basemodel 与 req_manager 的模块级循环导入。
+from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+from lightllm.common.req_manager import req_sampling_params
+from lightllm.common.req_manager.req_sampling_params import ReqSamplingParamsManager
+from lightllm.common.basemodel.triton_kernel.gen_sampling_params import token_id_counter
+
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+
+
+def test_prompt_ids_gpu_buffer_reuses_per_stream_and_isolates_async_copies(monkeypatch):
+    monkeypatch.setattr(req_sampling_params, "get_env_start_args", lambda: SimpleNamespace(max_req_total_len=32))
+    manager = ReqSamplingParamsManager.__new__(ReqSamplingParamsManager)
+    stream0, stream1 = torch.cuda.Stream(), torch.cuda.Stream()
+    source0 = torch.full((96,), 11, dtype=torch.int32, pin_memory=True)
+    source1 = torch.full((96,), 22, dtype=torch.int32, pin_memory=True)
+    assert not hasattr(manager, "_prompt_ids_gpu_buffers")
+    with torch.cuda.stream(stream0):
+        first = manager._get_prompt_ids_gpu_buffer()
+        assert manager._get_prompt_ids_gpu_buffer() is first
+        assert first.shape == (96,)
+        assert first.dtype == torch.int32 and first.device == stream0.device
+        torch.cuda._sleep(100_000_000)
+        first.copy_(source0, non_blocking=True)
+        completed0 = torch.cuda.Event()
+        completed0.record()
+    with torch.cuda.stream(stream1):
+        other = manager._get_prompt_ids_gpu_buffer()
+        assert manager._get_prompt_ids_gpu_buffer() is other
+        assert other.data_ptr() != first.data_ptr()
+        assert other.shape == (96,)
+        torch.cuda._sleep(100_000_000)
+        other.copy_(source1, non_blocking=True)
+        completed1 = torch.cuda.Event()
+        completed1.record()
+    with torch.cuda.stream(stream0):
+        assert manager._get_prompt_ids_gpu_buffer() is first
+    completed0.synchronize()
+    completed1.synchronize()
+    assert torch.equal(first.cpu(), source0)
+    assert torch.equal(other.cpu(), source1)
+    assert len(manager._prompt_ids_gpu_buffers) == 2
+
+
+@pytest.mark.parametrize("mode", ["gpu_counter", "pin_mem_counter"])
+def test_prompt_counter_initialization_reuses_stream_buffer(monkeypatch, mode):
+    monkeypatch.setattr(req_sampling_params, "get_env_start_args", lambda: SimpleNamespace(max_req_total_len=8192))
+    manager = ReqSamplingParamsManager.__new__(ReqSamplingParamsManager)
+    manager.penalty_counter_mode = mode
+    manager.req_to_next_token_ids = torch.zeros((8, 1), dtype=torch.int64, device="cuda")
+    manager.req_to_next_token_scores = None
+    for name in (
+        "req_to_presence_penalty",
+        "req_to_frequency_penalty",
+        "req_to_repetition_penalty",
+        "req_to_exponential_decay_length_penalty",
+    ):
+        setattr(manager, name, torch.zeros(8, device="cuda"))
+    manager.req_to_out_token_id_counter = torch.zeros(
+        (8, 16),
+        dtype=torch.int32,
+        device="cuda" if mode == "gpu_counter" else "cpu",
+        pin_memory=mode == "pin_mem_counter",
+    )
+    # 提前编译计数 kernel，随后延迟 GPU，检验连续复用暂存区时的异步执行顺序。
+    token_id_counter(torch.zeros(1, dtype=torch.int32, device="cuda"), manager.req_to_out_token_id_counter[0])
+    torch.cuda.synchronize()
+
+    def forbid_synchronize(*args, **kwargs):
+        pytest.fail("request initialization must not synchronize its CUDA stream")
+
+    monkeypatch.setattr(torch.cuda.Stream, "synchronize", forbid_synchronize)
+    stream = torch.cuda.Stream()
+    buffer_ptrs = []
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(100_000_000)
+        for index in range(8):
+            params = SimpleNamespace(
+                presence_penalty=1.0,
+                frequency_penalty=0.0,
+                repetition_penalty=1.0,
+                input_penalty=True,
+                exponential_decay_length_penalty=SimpleNamespace(to_tuple=lambda: (0, 1.0)),
+            )
+            prompt = np.full(4096, index + 1, dtype=np.int64)
+            req = SimpleNamespace(
+                req_idx=index,
+                sampling_param=SimpleNamespace(shm_param=params),
+                shm_req=SimpleNamespace(
+                    alloc_shm_numpy_len=8192,
+                    get_prompt_ids_numpy=lambda prompt=prompt: prompt,
+                ),
+                get_last_gen_token=lambda: 0,
+            )
+            manager.init_req_sampling_params(req)
+            buffer_ptrs.append(manager._get_prompt_ids_gpu_buffer().data_ptr())
+        completed = torch.cuda.Event()
+        completed.record()
+    completed.synchronize()
+    expected = torch.zeros((8, 16), dtype=torch.int32)
+    expected[torch.arange(8), torch.arange(1, 9)] = 4096
+    assert torch.equal(manager.req_to_out_token_id_counter.cpu(), expected)
+    assert len(set(buffer_ptrs)) == len(manager._prompt_ids_gpu_buffers) == 1
diff --git a/unit_tests/common/test_req_token_index.py b/unit_tests/common/test_req_token_index.py
new file mode 100644
index 0000000000..29117bb4cf
--- /dev/null
+++ b/unit_tests/common/test_req_token_index.py
@@ -0,0 +1,823 @@
+from types import SimpleNamespace
+import weakref
+
+import pytest
+import torch
+
+from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+from lightllm.common.req_manager import base
+from lightllm.server.router.model_infer.infer_req_ext1 import (
+    ReqTokenIndex,
+    PinMemCacheReqTokenIndex,
+    PinMemPageCacheReqTokenIndex,
+    PureGpuReqTokenIndex,
+    create_req_token_index,
+)
+from lightllm.server.router.model_infer.infer_batch import InferReq
+from lightllm.server.api_cli import make_argument_parser
+from lightllm.server.core.objs.start_args_type import StartArgs
+
+
+MODES = ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"]
+requires_cuda = pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA and pinned CPU memory")
+
+
+@pytest.fixture(autouse=True)
+def reset_page_offsets(monkeypatch):
+    # 各用例模拟不同启动配置；隔离进程级页偏移，同一用例内的请求仍共享。
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+
+
+def make_req(token_table, mode="pin_mem_cache", req_idx=0, input_len=0, page_size=1):
+    req = InferReq.__new__(InferReq)
+    req.req_idx = req_idx
+    req.req_manager = SimpleNamespace(req_to_token_indexs=token_table, page_size=page_size)
+    req.args = (
+        StartArgs(page_size=page_size) if mode is None else StartArgs(req_token_index_mode=mode, page_size=page_size)
+    )
+    req.shm_req = SimpleNamespace(input_len=input_len)
+    return req
+
+
+@pytest.mark.parametrize(
+    ("mode", "expected_type"),
+    [
+        (None, PinMemCacheReqTokenIndex),
+        ("pin_mem_cache", PinMemCacheReqTokenIndex),
+        ("pin_mem_page_cache", PinMemPageCacheReqTokenIndex),
+        ("pure_gpu", PureGpuReqTokenIndex),
+    ],
+)
+def test_factory_uses_request_mode_and_row_view_without_retaining_request(mode, expected_type):
+    table = torch.arange(24, dtype=torch.int32).view(3, 8)
+    req = make_req(table, mode, req_idx=1)
+    # 构造只保存传入的流，不执行 CUDA 操作；用占位对象验证透传即可。
+    index_stream = object()
+    helper = create_req_token_index(req, index_stream=index_stream)
+
+    assert not hasattr(InferReq, "move_to_gpu_index_len")
+    assert isinstance(helper, expected_type)
+    assert helper._alloc_return_pinned_tensor is ReqTokenIndex._alloc_return_pinned_tensor
+    assert helper._index_stream is index_stream
+    assert helper.req_idx == req.req_idx
+    assert helper.gpu_indexes.data_ptr() == table[req.req_idx].data_ptr()
+    expected = table.clone()
+    expected[req.req_idx, 3] = -99
+    helper.gpu_indexes[3] = -99
+    assert torch.equal(table, expected)
+    # 模拟真实持有关系；保留 Helper 时，请求也应能立即释放，无需循环 GC。
+    req.token_index = helper
+    request_ref = weakref.ref(req)
+    del req
+    assert request_ref() is None
+
+
+@pytest.mark.parametrize("mode", MODES)
+def test_start_args_and_cli_select_mode(mode):
+    parser = make_argument_parser()
+    assert parser.parse_args([]).req_token_index_mode == "pin_mem_cache"
+    assert StartArgs().req_token_index_mode == "pin_mem_cache"
+    assert parser.parse_args(["--req_token_index_mode", mode]).req_token_index_mode == mode
+    assert StartArgs(req_token_index_mode=mode).req_token_index_mode == mode
+    with pytest.raises(SystemExit):
+        parser.parse_args(["--req_token_index_mode", "unknown"])
+    with pytest.raises(ValueError, match="Unknown request token index mode"):
+        create_req_token_index(make_req(torch.empty((1, 8), dtype=torch.int32), "unknown"), index_stream=object())
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+@pytest.mark.parametrize("pinned", [False, True])
+def test_append_accepts_cpu_input_and_uses_selected_storage_mode(monkeypatch, mode, pinned):
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    helper = create_req_token_index(make_req(table, mode), index_stream=torch.cuda.current_stream())
+    indexes = torch.arange(4, dtype=torch.int32, pin_memory=pinned)
+    alloc_pinned = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor
+    pinned_sizes = []
+
+    def record_alloc_pinned(need_size):
+        pinned_sizes.append(need_size)
+        return alloc_pinned(need_size)
+
+    monkeypatch.setattr(ReqTokenIndex, "_alloc_return_pinned_tensor", staticmethod(record_alloc_pinned))
+    helper.append_token_indexes(indexes)
+    # 普通 CPU 输入已经复制到自有 storage，可在 append 返回后立即改写。
+    if not pinned:
+        indexes.fill_(-99)
+
+    assert pinned_sizes == []
+    assert helper.token_index_len == 4
+    assert helper.get_cpu_token_indexes(0, 4).tolist() == [0, 1, 2, 3]
+    assert helper.get_gpu_token_indexes(0, 4).cpu().tolist() == [0, 1, 2, 3]
+    if mode != "pure_gpu":
+        assert helper._cpu_buffer.is_pinned() == (mode == "pin_mem_page_cache")
+        assert helper._cpu_buffer.untyped_storage().nbytes() == helper._cpu_buffer.nbytes
+
+
+@requires_cuda
+def test_upload_snapshots_own_only_pending_indexes_across_cache_changes():
+    table = torch.full((1, 8192), -1, dtype=torch.int32, device="cuda")
+    helper = create_req_token_index(make_req(table), index_stream=torch.cuda.default_stream())
+    helper.append_token_indexes(torch.arange(4, dtype=torch.int32))
+    helper.get_gpu_token_indexes(0, 4)
+    helper.append_token_indexes(torch.arange(4, 8, dtype=torch.int32))
+    buffer_ref = weakref.ref(helper._cpu_buffer)
+
+    first = helper.get_upload_indexes()
+    second = helper.get_upload_indexes()
+    assert helper.gpu_index_submitted_len == 4
+    assert first.is_pinned() and second.is_pinned()
+    assert first.dtype == second.dtype == table.dtype
+    assert first.tolist() == second.tolist() == [4, 5, 6, 7]
+    assert first.nbytes == second.nbytes == 4 * table.element_size()
+    assert first.data_ptr() != second.data_ptr()
+    assert not helper._cpu_buffer.is_pinned()
+
+    # 修改常驻缓存、扩容和 reset 均不能改变已经返回的上传快照。
+    helper._cpu_buffer[4:8].fill_(-99)
+    helper.append_token_indexes(torch.arange(8, 4096, dtype=torch.int32))
+    assert buffer_ref() is None
+    assert not helper._cpu_buffer.is_pinned()
+    helper.reset()
+    helper.append_token_indexes(torch.arange(100, 104, dtype=torch.int32))
+    assert first.tolist() == second.tolist() == [4, 5, 6, 7]
+    assert helper.get_upload_indexes().tolist() == [100, 101, 102, 103]
+
+
+@requires_cuda
+def test_shared_pinned_allocator_preserves_slices_across_slab_replacement(monkeypatch):
+    allocator = ReqTokenIndex._alloc_return_pinned_tensor
+    monkeypatch.setattr(allocator, "return_buffer_state", (None, 0), raising=False)
+    buffer_token_num = 1024 * 1024
+    first = PinMemCacheReqTokenIndex._alloc_return_pinned_tensor(5)
+    second = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(3)
+    first.fill_(11)
+    second.fill_(22)
+    assert second.data_ptr() == first.data_ptr() + first.nbytes
+    assert first.is_pinned() and second.is_pinned()
+    assert first.dtype == second.dtype == torch.int32
+    assert not first.requires_grad and not second.requires_grad
+
+    buffer, offset = allocator.return_buffer_state
+    assert offset == 8
+    empty = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(0)
+    oversized = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(buffer_token_num + 1)
+    assert empty.numel() == 0 and oversized.numel() == buffer_token_num + 1 and oversized.is_pinned()
+    assert allocator.return_buffer_state[0] is buffer
+    assert allocator.return_buffer_state[1] == 8
+    assert oversized.untyped_storage().data_ptr() != first.untyped_storage().data_ptr()
+
+    replacement = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(buffer_token_num - 7)
+    replacement.fill_(33)
+    assert replacement.untyped_storage().data_ptr() != first.untyped_storage().data_ptr()
+    assert first.tolist() == [11] * 5 and second.tolist() == [22] * 3
+    assert allocator.return_buffer_state[0].numel() == buffer_token_num
+    assert allocator.return_buffer_state[1] == buffer_token_num - 7
+
+
+@requires_cuda
+def test_shared_pinned_allocator_rollover_keeps_released_h2d_source_valid(monkeypatch):
+    monkeypatch.setattr(ReqTokenIndex._alloc_return_pinned_tensor, "return_buffer_state", (None, 0), raising=False)
+    buffer_token_num = 1024 * 1024
+    source = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(buffer_token_num)
+    source.fill_(11)
+    source_ref = weakref.ref(source)
+    target = torch.empty(buffer_token_num, dtype=torch.int32, device="cuda")
+    stream = torch.cuda.Stream()
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(200_000_000)
+        target.copy_(source, non_blocking=True)
+    del source
+    assert source_ref() is None and not stream.query()
+
+    # 连续换掉多个 slab，旧 storage 释放后仍须等在途 H2D 结束才能被重新使用。
+    for value in range(22, 26):
+        replacement = PinMemPageCacheReqTokenIndex._alloc_return_pinned_tensor(buffer_token_num)
+        replacement.fill_(value)
+        del replacement
+    stream.synchronize()
+    assert target.eq(11).all().item()
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_append_rejects_gpu_input_without_changing_index_state(mode):
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    helper = create_req_token_index(make_req(table, mode), index_stream=torch.cuda.current_stream())
+
+    with pytest.raises(AssertionError, match="mem_indexes must be a CPU Tensor"):
+        helper.append_token_indexes(torch.arange(4, dtype=torch.int32, device="cuda"))
+
+    assert helper.token_index_len == helper.gpu_index_submitted_len == 0
+    assert table.eq(-1).all()
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+@pytest.mark.parametrize("getter_name", ["get_cpu_token_indexes", "get_gpu_token_indexes", "get_free_token_indexes"])
+@pytest.mark.parametrize("start,end", [(0, 65), (65, 65)])
+def test_read_getters_reject_unrecorded_indexes_without_advancing_progress(mode, getter_name, start, end):
+    table = torch.full((1, 128), -7, dtype=torch.int32, device="cuda")
+    helper = create_req_token_index(make_req(table, mode, page_size=64), index_stream=torch.cuda.current_stream())
+    helper.append_token_indexes(torch.arange(64, dtype=torch.int32, pin_memory=True))
+    submitted_len = helper.gpu_index_submitted_len
+    getter = getattr(helper, getter_name)
+
+    with pytest.raises(AssertionError, match="Read range exceeds"):
+        getter(start, end)
+
+    assert helper.token_index_len == 64
+    assert helper.gpu_index_submitted_len == submitted_len
+    assert getter(1, 3).cpu().tolist() == [1, 2]
+    assert getter(64, 64).numel() == 0
+
+
+@pytest.mark.parametrize("start,append_len,boundary", [(0, 63, "end"), (1, 63, "start")])
+def test_page_cache_rejects_partial_page_appends_before_mutating_state(start, append_len, boundary):
+    table = torch.full((1, 128), -7, dtype=torch.int32)
+    helper = create_req_token_index(make_req(table, "pin_mem_page_cache", page_size=64), index_stream=object())
+    # 模拟不对齐的已有长度，验证从页内追加也会在分配或写入之前失败。
+    helper.token_index_len = start
+
+    with pytest.raises(AssertionError, match=f"Append {boundary} must be page-aligned"):
+        helper.append_token_indexes(torch.arange(append_len, dtype=torch.int32))
+
+    assert helper.token_index_len == start
+    assert helper.gpu_index_submitted_len == 0
+    assert helper._cpu_buffer is None
+    assert table.eq(-7).all()
+
+
+@requires_cuda
+def test_batch_upload_rejects_missing_gpu_submission_without_advancing_progress(monkeypatch):
+    from lightllm.common import cuda_wrapper
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+
+    # 故障注入会禁用原生 batch 路径；测试结束后恢复，避免影响其他上传用例。
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", True)
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    req = make_req(table, "pure_gpu")
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.current_stream())
+    # 模拟不一致状态：记录了索引，但没有提交 GPU 写入；批量入口不能替它补写进度。
+    req.token_index.token_index_len = 4
+    monkeypatch.setattr(base_backend, "g_infer_context", SimpleNamespace(req_manager=req.req_manager))
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+
+    with pytest.raises(TypeError):
+        backend._copy_req_token_indexes_to_gpu([req])
+
+    assert req.token_index.gpu_index_submitted_len == 0
+    assert not hasattr(backend, "_index_upload_stream")
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_backend_rejects_requests_with_different_index_streams(monkeypatch, mode):
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+
+    table = torch.full((2, 8), -1, dtype=torch.int32, device="cuda")
+    reqs = []
+    for row in range(2):
+        index_stream = torch.cuda.Stream()
+        index_stream.wait_stream(torch.cuda.current_stream())
+        req = make_req(table, mode, req_idx=row)
+        req.token_index = create_req_token_index(req, index_stream=index_stream)
+        req.token_index.append_token_indexes(torch.arange(4, dtype=torch.int32, pin_memory=True))
+        reqs.append(req)
+
+    def unexpected_upload(*args):
+        pytest.fail("A batch with mismatched index streams must be rejected before upload")
+
+    monkeypatch.setattr(base_backend, "cuda_memcpy_pinned_h2d_batch_async", unexpected_upload)
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+    with pytest.raises(AssertionError, match="All requests in a batch must use the same index stream"):
+        backend._copy_req_token_indexes_to_gpu(reqs)
+    assert [req.token_index.gpu_index_submitted_len for req in reqs] == [4 if mode == "pure_gpu" else 0] * 2
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache"])
+def test_partial_batch_submission_can_retry_and_share_progress_with_gpu_getter(mode):
+    table = torch.full((2, 16), -1, dtype=torch.int32, device="cuda")
+    helpers = [
+        create_req_token_index(make_req(table, mode, req_idx=row, page_size=4), torch.cuda.current_stream())
+        for row in range(2)
+    ]
+    for row, helper in enumerate(helpers):
+        helper.append_token_indexes(torch.arange(row * 8, row * 8 + 4, dtype=torch.int32))
+
+    def partial_copy(destinations, sources):
+        destinations[0].copy_(sources[0], non_blocking=True)
+        raise RuntimeError("partial submission failed")
+
+    with pytest.raises(RuntimeError, match="partial submission failed"):
+        ReqTokenIndex.submit_pending_uploads(helpers, copy_batch=partial_copy)
+    assert [helper.gpu_index_submitted_len for helper in helpers] == [0, 0]
+
+    # 重试须覆盖完整待提交区间；单请求 getter 随后只提交新追加部分。
+    ReqTokenIndex.submit_pending_uploads(helpers)
+    helpers[0].append_token_indexes(torch.arange(4, 8, dtype=torch.int32))
+    assert helpers[0].get_gpu_token_indexes(0, 8).cpu().tolist() == list(range(8))
+    assert helpers[1].get_gpu_token_indexes(0, 4).cpu().tolist() == list(range(8, 12))
+    assert [helper.gpu_index_submitted_len for helper in helpers] == [8, 4]
+
+    def unexpected_copy(*args):
+        pytest.fail("Getter and batch submission must share the same upload progress")
+
+    ReqTokenIndex.submit_pending_uploads(helpers, copy_batch=unexpected_copy)
+    assert table[0, 8:].eq(-1).all() and table[1, 4:].eq(-1).all()
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_token_units_and_snapshot_ownership(mode, page_size):
+    row = torch.full((4 * page_size,), -1, dtype=torch.int32, device="cuda")
+    req = make_req(row.unsqueeze(0), mode, page_size=page_size)
+    helper = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    req.token_index = helper
+    count = row.numel()
+    positions = torch.arange(count, dtype=torch.int32)
+    expected = positions.div(page_size, rounding_mode="floor") * (3 * page_size) + positions.remainder(page_size)
+    original = expected.pin_memory()
+    split = page_size if mode == "pin_mem_page_cache" else page_size + 1
+    req.token_index.append_token_indexes(original[:split])
+    req.token_index.append_token_indexes(torch.empty(0, dtype=torch.int32))
+    assert helper.token_index_len == split
+    req.token_index.append_token_indexes(original[split:])
+    assert helper.token_index_len == count
+    ranges = [(0, count), (1, count - 1), (page_size, page_size + 2), (count, count)]
+    for start, end in ranges:
+        assert torch.equal(req.token_index.get_cpu_token_indexes(start, end), expected[start:end])
+    snapshot = req.token_index.get_cpu_token_indexes(0, count)
+    # pure_gpu 的阻塞 CPU getter 已等待 H2D 完成，此后才可主动改写上传源。
+    original.fill_(99)
+    # GPU 读取区间只决定返回视图；首次读取同时提交内部记录的全部待上传索引。
+    assert torch.equal(req.token_index.get_gpu_token_indexes(1, 3).cpu(), expected[1:3])
+    assert helper.gpu_index_submitted_len == count
+    gpu = req.token_index.get_gpu_token_indexes(0, count)
+    assert gpu.data_ptr() == row.data_ptr()
+    assert torch.equal(gpu.cpu(), expected)
+    if mode != "pure_gpu":
+        cpu_buffer = helper._cpu_buffer
+    helper.reset()
+    assert helper.token_index_len == helper.gpu_index_submitted_len == 0
+    if mode != "pure_gpu":
+        assert helper._cpu_buffer is None
+    else:
+        assert not hasattr(helper, "_cpu_buffer")
+    replacement = torch.arange(count, dtype=torch.int32, pin_memory=True)
+    helper.append_token_indexes(replacement)
+    if mode != "pure_gpu":
+        assert helper._cpu_buffer is not cpu_buffer
+    assert torch.equal(helper.get_gpu_token_indexes(0, count).cpu(), replacement)
+    assert torch.equal(snapshot, expected)
+
+
+@requires_cuda
+@pytest.mark.parametrize("page_size", [1, 3, 4, 64])
+def test_page_cache_compresses_page_starts_and_reads_partial_token_ranges(page_size):
+    table = torch.full((1, 8 * page_size), -1, dtype=torch.int32, device="cuda")
+    req = make_req(table, "pin_mem_page_cache", input_len=4 * page_size, page_size=page_size)
+    helper = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    page_starts = torch.tensor([9, 2, 15, 4], dtype=torch.int32) * page_size
+    expected = (page_starts[:, None] + torch.arange(page_size, dtype=torch.int32)).reshape(-1)
+    source = expected.clone()
+    start = 0
+    for end in (page_size, 2 * page_size, 3 * page_size, expected.numel()):
+        helper.append_token_indexes(source[start:end])
+        start = end
+    source.fill_(-99)
+
+    assert helper._cpu_buffer.is_pinned() and helper._cpu_buffer.numel() == 8
+    assert helper._cpu_buffer.nbytes * page_size == table.nbytes
+    assert torch.equal(helper._cpu_buffer[:4], page_starts)
+    for start, end in (
+        (0, expected.numel()),
+        (1, expected.numel() - 1),
+        (page_size - 1, 2 * page_size + 1),
+        (1, 2),
+        (1, 1),
+        (page_size, page_size),
+    ):
+        assert torch.equal(helper.get_cpu_token_indexes(start, end), expected[start:end])
+        assert torch.equal(helper.get_free_token_indexes(start, end), expected[start:end])
+
+    snapshot = helper.get_cpu_token_indexes(0, expected.numel())
+    upload = helper.get_upload_indexes()
+    assert upload.is_pinned() and torch.equal(upload, expected)
+    assert helper.gpu_index_submitted_len == 0
+    # 模拟已提交前缀结束于页内，待上传接口仍返回精确的剩余 token 区间。
+    submitted = page_size + 1
+    table[0, :submitted].copy_(upload[:submitted], non_blocking=True)
+    helper.gpu_index_submitted_len = submitted
+    pending = helper.get_upload_indexes()
+    assert pending.is_pinned() and torch.equal(pending, expected[submitted:])
+    assert helper.gpu_index_submitted_len == submitted
+    assert torch.equal(helper.get_gpu_token_indexes(0, expected.numel()).cpu(), expected)
+    assert helper.get_upload_indexes().numel() == 0
+    assert table[0, expected.numel() :].eq(-1).all()
+
+    helper.reset()
+    helper.append_token_indexes(torch.arange(expected.numel(), dtype=torch.int32))
+    assert torch.equal(snapshot, expected) and torch.equal(upload, expected)
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_free_getter_borrows_gpu_indexes_but_metadata_getter_keeps_cpu_snapshot(monkeypatch, mode):
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    req = make_req(table, mode)
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    req.token_index.append_token_indexes(torch.arange(6, dtype=torch.int32, pin_memory=True))
+    snapshot = req.token_index.get_cpu_token_indexes(1, 5)
+
+    def unexpected_snapshot(*args):
+        pytest.fail("The pure_gpu free getter must not call the CPU snapshot getter")
+
+    with monkeypatch.context() as no_snapshot:
+        if mode == "pure_gpu":
+            no_snapshot.setattr(req.token_index, "get_cpu_token_indexes", unexpected_snapshot)
+        free_indexes = req.token_index.get_free_token_indexes(1, 5)
+        empty = req.token_index.get_free_token_indexes(3, 3)
+
+    expected_device = "cuda" if mode == "pure_gpu" else "cpu"
+    assert free_indexes.device.type == empty.device.type == expected_device
+    assert empty.numel() == 0
+    if mode == "pure_gpu":
+        assert free_indexes.data_ptr() == table[0, 1:5].data_ptr()
+    updated = torch.arange(101, 105, dtype=torch.int32)
+    replacement = torch.arange(6, dtype=torch.int32, pin_memory=True)
+    replacement[1:5] = updated
+    req.token_index.reset()
+    req.token_index.append_token_indexes(replacement)
+
+    assert snapshot.device.type == "cpu"
+    assert snapshot.tolist() == [1, 2, 3, 4]
+    assert free_indexes.cpu().tolist() == (updated.tolist() if mode == "pure_gpu" else [1, 2, 3, 4])
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_reset_rebuild_keeps_cpu_appends_and_gpu_submission_separate(mode):
+    helper = create_req_token_index(
+        make_req(torch.zeros((1, 8), dtype=torch.int32, device="cuda"), mode), index_stream=torch.cuda.default_stream()
+    )
+    pending = helper.get_upload_indexes()
+    if mode == "pure_gpu":
+        assert pending is None
+    else:
+        assert pending.numel() == 0
+    helper.append_token_indexes(torch.arange(4, dtype=torch.int32, pin_memory=True))
+    pending = helper.get_upload_indexes()
+    if mode == "pure_gpu":
+        assert pending is None
+    else:
+        assert pending.tolist() == [0, 1, 2, 3]
+    assert helper.gpu_index_submitted_len == (4 if mode == "pure_gpu" else 0)
+    before = helper.get_cpu_token_indexes(0, 4)
+    helper.get_gpu_token_indexes(0, 4)
+    pending = helper.get_upload_indexes()
+    if mode == "pure_gpu":
+        assert pending is None
+    else:
+        assert pending.numel() == 0
+    helper.reset()
+    helper.append_token_indexes(torch.tensor([0, 31], dtype=torch.int32, pin_memory=True))
+    assert helper.token_index_len == 2
+    assert helper.gpu_index_submitted_len == (2 if mode == "pure_gpu" else 0)
+    if mode != "pure_gpu":
+        assert helper.gpu_indexes[:4].cpu().tolist() == [0, 1, 2, 3]
+    helper.append_token_indexes(torch.tensor([32, 3], dtype=torch.int32, pin_memory=True))
+    assert helper.token_index_len == 4
+    assert helper.gpu_index_submitted_len == (4 if mode == "pure_gpu" else 0)
+    assert helper.get_gpu_token_indexes(0, 4).cpu().tolist() == [0, 31, 32, 3]
+    assert helper.gpu_index_submitted_len == 4
+    assert before.tolist() == [0, 1, 2, 3]
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache"])
+@pytest.mark.parametrize("page_size", [1, 64])
+@pytest.mark.parametrize(
+    ("input_len", "max_length", "initial_capacity", "next_end", "grown_capacity"),
+    [
+        (1023, 16384, 3072, 3073, 6144),
+        (1024, 16384, 3072, 3073, 6144),
+        (1025, 16384, 4096, 4097, 7168),
+        (2048, 16384, 4096, 5000, 7168),
+        (1025, 5500, 4096, 5000, 5500),
+    ],
+)
+def test_cpu_buffer_capacity_and_lifecycle(
+    mode, page_size, input_len, max_length, initial_capacity, next_end, grown_capacity
+):
+    req = make_req(
+        torch.empty((1, max_length), dtype=torch.int32, device="cuda"), mode, input_len=input_len, page_size=page_size
+    )
+    helper = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    storage_page_size = page_size if mode == "pin_mem_page_cache" else 1
+    if mode == "pin_mem_page_cache":
+        next_end = (next_end + page_size - 1) // page_size * page_size
+        initial_elements = grown_elements = (max_length + page_size - 1) // page_size
+    else:
+        initial_elements = initial_capacity
+        grown_elements = grown_capacity
+    assert helper.input_len == input_len
+    assert helper._cpu_buffer is None
+    helper.append_token_indexes(torch.empty(0, dtype=torch.int32))
+    assert helper._cpu_buffer is None
+    # 首次追加最小存储单位；完整模式预留 2K 输出，页模式一次申请到最大长度。
+    first_index = torch.arange(storage_page_size, dtype=torch.int32, pin_memory=True)
+    helper.append_token_indexes(first_index)
+    assert helper._cpu_buffer.numel() == initial_elements
+    assert helper._cpu_buffer.is_pinned() == (mode == "pin_mem_page_cache")
+    assert helper._cpu_buffer.data_ptr() != first_index.data_ptr()
+    helper.append_token_indexes(torch.arange(storage_page_size, initial_capacity, dtype=torch.int32, pin_memory=True))
+    old_buffer = helper._cpu_buffer
+    helper.append_token_indexes(torch.arange(initial_capacity, next_end, dtype=torch.int32, pin_memory=True))
+    assert helper._cpu_buffer.numel() == grown_elements
+    assert helper._cpu_buffer.is_pinned() == (mode == "pin_mem_page_cache")
+    if mode == "pin_mem_page_cache":
+        assert helper._cpu_buffer is old_buffer
+    else:
+        assert helper._cpu_buffer.data_ptr() != old_buffer.data_ptr()
+    snapshot = helper.get_cpu_token_indexes(0, next_end)
+    assert snapshot.tolist() == list(range(next_end))
+    grown_buffer = helper._cpu_buffer
+    helper.reset()
+    assert helper.input_len == input_len and helper._cpu_buffer is None
+    assert helper.token_index_len == helper.gpu_index_submitted_len == 0
+    replacement = torch.arange(99, 99 + storage_page_size, dtype=torch.int32, pin_memory=True)
+    helper.append_token_indexes(replacement)
+    assert helper._cpu_buffer is not grown_buffer
+    assert helper._cpu_buffer.numel() == initial_elements
+    assert helper.token_index_len == storage_page_size
+    assert torch.equal(helper.get_cpu_token_indexes(0, helper.token_index_len), replacement)
+    assert snapshot.tolist() == list(range(next_end))
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache"])
+@pytest.mark.parametrize("page_size", [1, 64])
+@pytest.mark.parametrize("use_default_stream", [False, True])
+def test_cpu_buffer_reset_during_upload_keeps_old_data_until_h2d_completes(mode, page_size, use_default_stream):
+    count = 8192
+    row = torch.full((count,), -1, dtype=torch.int32, device="cuda")
+    index_stream = torch.cuda.default_stream() if use_default_stream else torch.cuda.Stream()
+    index_stream.wait_stream(torch.cuda.current_stream())
+    helper = create_req_token_index(
+        make_req(row.unsqueeze(0), mode, input_len=count, page_size=page_size), index_stream=index_stream
+    )
+    original = torch.arange(count, dtype=torch.int32, pin_memory=True)
+    replacement = (torch.arange(count, dtype=torch.int32).remainder(page_size) + count * 2).pin_memory()
+    helper.append_token_indexes(original)
+    source_ref = weakref.ref(helper._cpu_buffer)
+
+    with torch.cuda.stream(index_stream):
+        torch.cuda._sleep(200_000_000)
+    helper.get_gpu_token_indexes(0, count)
+    assert not index_stream.query(), "Delayed H2D must remain pending through reset and append"
+    helper.reset()
+    assert helper._cpu_buffer is None and source_ref() is None
+    helper.append_token_indexes(replacement)
+
+    # reset 已释放旧 Tensor 的 Python 引用；新缓存写入不能污染仍在传输的旧数据。
+    index_stream.synchronize()
+    assert torch.equal(row.cpu(), original)
+    assert torch.equal(helper.get_gpu_token_indexes(0, count).cpu(), replacement)
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+@pytest.mark.parametrize("pre_upload", [False, True])
+@pytest.mark.parametrize("use_default_stream", [False, True])
+def test_gpu_getter_uses_injected_stream_and_orders_calling_stream(monkeypatch, mode, pre_upload, use_default_stream):
+    row = torch.full((16,), -1, dtype=torch.int32, device="cuda")
+    index_stream = torch.cuda.default_stream() if use_default_stream else torch.cuda.Stream()
+    index_stream.wait_stream(torch.cuda.current_stream())
+    reader = torch.cuda.Stream()
+    with torch.cuda.stream(reader):
+        # 构造时的当前流与传入的索引流不同，Helper 必须始终使用外部指定的流。
+        helper = create_req_token_index(make_req(row.unsqueeze(0), mode), index_stream=index_stream)
+        assert helper._index_stream is index_stream
+        assert torch.cuda.current_stream() == reader
+    copy = torch.Tensor.copy_
+    cpu = torch.Tensor.cpu
+    uploads = []
+
+    def record_copy(dst, src, *args, **kwargs):
+        if dst.is_cuda and src.device.type == "cpu":
+            assert torch.cuda.current_stream() == index_stream
+            uploads.append(src.numel())
+        return copy(dst, src, *args, **kwargs)
+
+    def check_cpu_read(src, *args, **kwargs):
+        if src.is_cuda:
+            assert torch.cuda.current_stream() == index_stream
+        return cpu(src, *args, **kwargs)
+
+    monkeypatch.setattr(torch.Tensor, "copy_", record_copy)
+    monkeypatch.setattr(torch.Tensor, "cpu", check_cpu_read)
+    with torch.cuda.stream(index_stream):
+        torch.cuda._sleep(30_000_000)
+    with torch.cuda.stream(reader):
+        helper.append_token_indexes(torch.arange(8, dtype=torch.int32, pin_memory=True))
+        assert torch.cuda.current_stream() == reader
+    if pre_upload:
+        # 索引流已提交上传，但尚未完成；另一个流读取时仍必须承接依赖。
+        helper.get_gpu_token_indexes(0, 8)
+    with torch.cuda.stream(reader):
+        indexes = helper.get_gpu_token_indexes(0, 8)
+        assert torch.cuda.current_stream() == reader
+        actual = indexes.clone()
+        assert helper.get_cpu_token_indexes(0, 8).tolist() == list(range(8))
+        assert torch.cuda.current_stream() == reader
+    reader.synchronize()
+    assert uploads == [8]
+    with torch.cuda.stream(index_stream):
+        assert actual.cpu().tolist() == list(range(8))
+
+
+@requires_cuda
+def test_pure_gpu_slot_reuse_keeps_old_and_new_helper_writes_in_order():
+    row = torch.full((8,), -1, dtype=torch.int32, device="cuda")
+    torch.cuda._sleep(30_000_000)
+    old = create_req_token_index(make_req(row.unsqueeze(0), "pure_gpu"), index_stream=torch.cuda.default_stream())
+    old.append_token_indexes(torch.full((8,), 11, dtype=torch.int32, pin_memory=True))
+    # 源 Tensor 引用已释放；先保留旧写入结果，验证后续 pinned 分配不会覆盖在途 H2D。
+    first_write = row.clone()
+    new = create_req_token_index(make_req(row.unsqueeze(0), "pure_gpu"), index_stream=torch.cuda.default_stream())
+    new.append_token_indexes(torch.full((8,), 22, dtype=torch.int32, pin_memory=True))
+    assert new.get_cpu_token_indexes(0, 8).tolist() == [22] * 8
+    assert first_write.cpu().tolist() == [11] * 8
+
+
+@requires_cuda
+@pytest.mark.parametrize("page_size", [1, 4, 64])
+def test_manager_owns_only_gpu_table_and_initializes_hold_page(monkeypatch, page_size):
+    from lightllm.common.req_manager import req_sampling_params
+
+    monkeypatch.setattr(base, "get_env_start_args", lambda: SimpleNamespace(page_size=page_size))
+    monkeypatch.setattr(req_sampling_params, "ReqSamplingParamsManager", lambda count: None)
+    mem_manager = SimpleNamespace(page_size=page_size, HOLD_TOKEN_MEMINDEXES=tuple(range(page_size * 8, page_size * 9)))
+    manager = base.ReqManager(2, page_size * 2 + 1, mem_manager)
+    assert not hasattr(manager, "req_token_index_mode")
+    assert not hasattr(manager, "req_to_page_indexs_cpu")
+    assert not hasattr(manager, "get_cpu_token_indexes")
+    assert manager.req_to_token_indexs.shape == (3, 3 * page_size)
+    assert (
+        manager.req_to_token_indexs[manager.HOLD_REQUEST_ID].cpu().tolist()
+        == list(mem_manager.HOLD_TOKEN_MEMINDEXES) * 3
+    )
+
+
+@requires_cuda
+@pytest.mark.parametrize("batch_available", [False, True])
+@pytest.mark.parametrize("mode", MODES)
+@pytest.mark.parametrize("page_size", [1, 64])
+@pytest.mark.parametrize("index_stream_kind", ["default", "shared"])
+def test_backend_uploads_pinned_snapshots_without_request_events_or_keepalive_list(
+    monkeypatch, mode, page_size, batch_available, index_stream_kind
+):
+    from lightllm.common import cuda_wrapper
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+
+    count = max(8, page_size)
+    table = torch.full((2, 2 * count), -1, dtype=torch.int32, device="cuda")
+    expected = (torch.arange(2 * count, dtype=torch.int32).view(2, count) + 32).pin_memory()
+    index_streams = [torch.cuda.default_stream()] * 2
+    if index_stream_kind == "shared":
+        index_streams = [torch.cuda.Stream()] * 2
+    for index_stream in set(index_streams):
+        index_stream.wait_stream(torch.cuda.default_stream())
+    reqs = []
+    for row, index_stream in enumerate(index_streams):
+        req = make_req(table, mode, req_idx=row, page_size=page_size)
+        req.token_index = create_req_token_index(req, index_stream=index_stream)
+        with torch.cuda.stream(index_stream):
+            # 批量上传须先等待旧写入；pure_gpu 在本流追加，调用流尚未承接依赖。
+            torch.cuda._sleep(60_000_000)
+            table[row, :count].fill_(-2)
+            req.token_index.append_token_indexes(expected[row])
+        req.hold_kv_len = count
+        reqs.append(req)
+    monkeypatch.setattr(base_backend, "g_infer_context", SimpleNamespace(req_manager=reqs[0].req_manager))
+    monkeypatch.setattr(cuda_wrapper.PinnedH2DBatchCopier, "_native_batch_available", batch_available)
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+    stream = index_streams[0]
+    stream.wait_stream(torch.cuda.default_stream())
+    with torch.cuda.stream(stream):
+        torch.cuda._sleep(30_000_000)
+    events = []
+    batch_calls = []
+    copy_batch = base_backend.cuda_memcpy_pinned_h2d_batch_async
+    event_type = torch.cuda.Event
+
+    def record_batch(dst_tensors, src_tensors):
+        # 完整缓存复制待上传区间，页缓存先展开；两者都应批量提交独立的 token 索引快照。
+        assert all(tensor.is_pinned() for tensor in src_tensors)
+        assert all(torch.equal(tensor, expected[row]) for row, tensor in enumerate(src_tensors))
+        if mode == "pin_mem_cache":
+            for req, source in zip(reqs, src_tensors):
+                assert source.untyped_storage().data_ptr() != req.token_index._cpu_buffer.data_ptr()
+                # 即使上传尚未开始，改写原缓存也不能改变已经准备好的 pinned 源。
+                req.token_index._cpu_buffer[:count].fill_(-99)
+            assert all(torch.equal(tensor, expected[row]) for row, tensor in enumerate(src_tensors))
+        batch_calls.append(len(src_tensors))
+        return copy_batch(dst_tensors, src_tensors)
+
+    def record_event(*args, **kwargs):
+        event = event_type(*args, **kwargs)
+        events.append(event)
+        return event
+
+    with monkeypatch.context() as batch_events:
+        batch_events.setattr(torch.cuda, "Event", record_event)
+        batch_events.setattr(base_backend, "cuda_memcpy_pinned_h2d_batch_async", record_batch)
+        backend._copy_req_token_indexes_to_gpu(reqs)
+        observed = []
+        for req, index_stream in zip(reqs, index_streams):
+            with torch.cuda.stream(index_stream):
+                # 不等待后台线程：已标记提交的索引必须能在请求自己的索引流立即使用。
+                observed.append(req.token_index.get_gpu_token_indexes(0, count).clone())
+        if mode != "pure_gpu":
+            for req in reqs:
+                req.token_index.reset()
+                req.token_index.append_token_indexes(torch.arange(count, dtype=torch.int32, pin_memory=True))
+        # 两个请求只建立一个 batch 保活事件；pure_gpu 与 PyTorch 回退无需该事件。
+        assert len(events) == (1 if mode != "pure_gpu" and batch_available else 0)
+        assert batch_calls == ([] if mode == "pure_gpu" else [2])
+        pool = cuda_wrapper.PinnedH2DBatchCopier._release_pool
+        if pool is not None:
+            pool.submit(lambda: None).result(timeout=10)
+    assert [req.token_index.gpu_index_submitted_len for req in reqs] == [count if mode == "pure_gpu" else 0] * 2
+    assert not hasattr(backend, "_pending_cpu_index_uploads")
+    assert not hasattr(backend, "_index_upload_stream")
+    assert all(not hasattr(req.token_index, "_gpu_write_event") for req in reqs)
+    assert torch.equal(table[:, :count].cpu(), expected)
+    for actual, expected_row, index_stream in zip(observed, expected, index_streams):
+        with torch.cuda.stream(index_stream):
+            assert torch.equal(actual.cpu(), expected_row)
+    backend._copy_req_token_indexes_to_gpu([])
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_backend_waits_for_submitted_indexes_without_pending_cpu_uploads(monkeypatch, mode):
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    index_stream = torch.cuda.default_stream()
+    reader_stream = torch.cuda.Stream()
+    reader_stream.wait_stream(index_stream)
+    req = make_req(table, mode)
+    req.token_index = create_req_token_index(req, index_stream=index_stream)
+    with torch.cuda.stream(index_stream):
+        torch.cuda._sleep(200_000_000)
+        req.token_index.append_token_indexes(torch.arange(8, dtype=torch.int32, pin_memory=True))
+        req.token_index.get_gpu_token_indexes(0, 8)
+    assert req.token_index.gpu_index_submitted_len == req.token_index.token_index_len == 8
+    assert not index_stream.query(), "Submitted indexes must still be pending with no new CPU upload range"
+
+    def unexpected_upload(*args):
+        pytest.fail("No copy should be submitted when there are no pending upload ranges")
+
+    monkeypatch.setattr(base_backend, "cuda_memcpy_pinned_h2d_batch_async", unexpected_upload)
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+    with torch.cuda.stream(reader_stream):
+        backend._copy_req_token_indexes_to_gpu([req])
+        # 直接读取请求表，避免 GPU getter 自带的等待掩盖 batch 入口遗漏的依赖。
+        observed = table.clone()
+    reader_stream.synchronize()
+    assert observed.cpu().tolist() == [list(range(8))]
+
+
+@requires_cuda
+@pytest.mark.parametrize("mode", MODES)
+def test_dp_publication_waits_for_the_selected_index_mode(monkeypatch, mode):
+    from lightllm.server.router.model_infer.mode_backend import base_backend
+    from lightllm.server.router.model_infer.mode_backend.dp_backend import dp_shared_kv_trans
+
+    table = torch.full((1, 8), -1, dtype=torch.int32, device="cuda")
+    torch.cuda._sleep(30_000_000)
+    table.fill_(-2)
+    req = make_req(table, mode)
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    req.token_index.append_token_indexes(torch.arange(8, dtype=torch.int32, pin_memory=True))
+    req.cur_kv_len = req.hold_kv_len = 8
+    context = SimpleNamespace(req_manager=req.req_manager)
+    monkeypatch.setattr(base_backend, "g_infer_context", context)
+    monkeypatch.setattr(dp_shared_kv_trans, "g_infer_context", context)
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+    backend.node_nccl_group = object()
+    backend.is_master_in_dp = False
+
+    def barrier(group):
+        torch.cuda.current_stream().synchronize()
+        assert table.cpu().tolist() == [list(range(8))]
+
+    monkeypatch.setattr(dp_shared_kv_trans.dist, "barrier", barrier)
+    module = dp_shared_kv_trans.DPKVSharedMoudle.__new__(dp_shared_kv_trans.DPKVSharedMoudle)
+    module.backend = backend
+    module.fill_reqs_info([req])
diff --git a/unit_tests/server/router/dynamic_prompt/test_radix_cache.py b/unit_tests/server/router/dynamic_prompt/test_radix_cache.py
index 80265ac40d..c5d9623e62 100644
--- a/unit_tests/server/router/dynamic_prompt/test_radix_cache.py
+++ b/unit_tests/server/router/dynamic_prompt/test_radix_cache.py
@@ -255,6 +255,8 @@ def test_case10():
     )
     assert tree_node is not None
     assert size == 3
+    tree.dec_node_ref_counter(tree_node)
+    tree.mem_manager = SimpleNamespace(allocator=SimpleNamespace(can_use_mem_size=95), free=lambda indexes: None)
     tree.flush_cache()
     tree_node, size, values = tree.match_prefix(
         torch.tensor([1, 2, 3], dtype=torch.int64, device="cpu"), update_refs=True
diff --git a/unit_tests/server/router/model_infer/mode_backend/test_diverse_backend.py b/unit_tests/server/router/model_infer/mode_backend/test_diverse_backend.py
index ab02d5892c..105e59231a 100644
--- a/unit_tests/server/router/model_infer/mode_backend/test_diverse_backend.py
+++ b/unit_tests/server/router/model_infer/mode_backend/test_diverse_backend.py
@@ -1,11 +1,14 @@
 from types import SimpleNamespace
 
 import torch
+import pytest
 
 # 沿用生产代码的模块加载顺序，避免独立运行本测试文件时形成循环导入。
 from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
 from lightllm.server.router.model_infer.infer_batch import InferReq
 from lightllm.server.router.model_infer.mode_backend.diverse_backend.impl import DiversehBackend
+from lightllm.server.router.model_infer.mode_backend.diverse_backend import impl as diverse_impl
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
 
 
 class _FakeRadixCache:
@@ -36,13 +39,14 @@ def test_diverse_copy_skips_slaves_waiting_for_overlap_cleanup():
 
 
 def test_copy_master_kv_updates_slave_hold_kv_len():
-    req_to_token_indexs = torch.full((2, 8), -1, dtype=torch.int32)
-    req_to_token_indexs[0, :4] = torch.tensor([10, 11, 12, 13], dtype=torch.int32)
+    token_indexes_cpu = torch.full((2, 8), -1, dtype=torch.int32)
+    token_indexes_cpu[0, :4] = torch.tensor([10, 11, 12, 13], dtype=torch.int32)
     old_shared_node = object()
     new_shared_node = object()
     master_req = SimpleNamespace(
         req_idx=0,
         cur_kv_len=4,
+        hold_kv_len=4,
         cur_output_len=1,
         shared_kv_node=new_shared_node,
     )
@@ -58,17 +62,99 @@ def test_copy_master_kv_updates_slave_hold_kv_len():
     )
     backend = DiversehBackend.__new__(DiversehBackend)
     backend.model = SimpleNamespace(
-        req_manager=SimpleNamespace(req_to_token_indexs=req_to_token_indexs),
+        req_manager=make_cpu_index_manager(token_table=token_indexes_cpu, page_size=1),
     )
     backend.radix_cache = _FakeRadixCache()
     backend.is_master_in_dp = False
 
+    bind_cpu_indexes(master_req, backend.model.req_manager)
+    bind_cpu_indexes(slave_req, backend.model.req_manager, gpu_index_submitted_len=2)
     backend._copy_master_req_to_slave_req(slave_req)
 
-    assert req_to_token_indexs[1, :4].tolist() == [10, 11, 12, 13]
+    assert slave_req.token_index.get_cpu_token_indexes(0, 4).tolist() == [10, 11, 12, 13]
     assert slave_req.shared_kv_node is new_shared_node
     assert slave_req.cur_kv_len == 4
     assert slave_req.hold_kv_len == 4
+    assert slave_req.token_index.token_index_len == 4
+    assert slave_req.token_index.gpu_index_submitted_len == 0
     assert InferReq._kv_cache_alloc_need(slave_req, target_kv_len=5) == 1
     assert backend.radix_cache.dec_nodes == [old_shared_node]
     assert backend.radix_cache.add_nodes == [new_shared_node]
+
+
+def test_master_shared_prefix_rebuild_preserves_reserved_token_indexes():
+    token_indexes_cpu = torch.tensor([[10, 11, 12, 13, 20, 21, 22, 23]], dtype=torch.int32)
+    shared_indexes = torch.tensor([100, 101, 102, 103], dtype=torch.int32)
+    new_shared_node = SimpleNamespace(node_prefix_total_len=4)
+    master_req = SimpleNamespace(
+        req_idx=0,
+        cur_kv_len=4,
+        hold_kv_len=8,
+        shared_kv_node=None,
+        get_input_token_ids=lambda: [1, 2, 3, 4],
+        args=SimpleNamespace(page_size=4),
+    )
+    freed_indexes = []
+
+    class RadixCache(_FakeRadixCache):
+        def insert(self, key, value):
+            assert key.tolist() == [1, 2, 3, 4]
+            assert value.tolist() == [10, 11, 12, 13]
+            return 4, new_shared_node
+
+        def match_prefix(self, key, update_refs):
+            assert not update_refs
+            return new_shared_node, 4, shared_indexes
+
+    manager = make_cpu_index_manager(token_table=token_indexes_cpu, page_size=4)
+    bind_cpu_indexes(master_req, manager, gpu_index_submitted_len=8)
+    original_buffer = master_req.token_index._cpu_buffer
+    backend = DiversehBackend.__new__(DiversehBackend)
+    backend.model = SimpleNamespace(mem_manager=SimpleNamespace(free=lambda indexes: freed_indexes.append(indexes)))
+    backend.radix_cache = RadixCache()
+
+    backend._master_req_to_radix_cache(master_req)
+
+    assert master_req.token_index.get_cpu_token_indexes(0, 8).tolist() == [100, 101, 102, 103, 20, 21, 22, 23]
+    assert master_req.token_index._cpu_buffer is not original_buffer
+    assert master_req.token_index.token_index_len == master_req.hold_kv_len == 8
+    assert master_req.token_index.gpu_index_submitted_len == 0
+    assert InferReq._kv_cache_alloc_need(master_req, target_kv_len=8) == 0
+    assert InferReq._kv_cache_alloc_need(master_req, target_kv_len=9) == 4
+    assert [indexes.tolist() for indexes in freed_indexes] == [[10, 11, 12, 13]]
+    assert master_req.shared_kv_node is new_shared_node
+
+
+@pytest.mark.parametrize("is_chuncked_mode", [False, True])
+def test_diverse_waits_for_model_before_rewriting_shared_prefix(monkeypatch, is_chuncked_mode):
+    order = []
+    overlap_stream = object()
+
+    class DefaultStream:
+        def wait_stream(self, source):
+            assert source is overlap_stream
+            order.append("wait_model")
+
+        def synchronize(self):
+            order.append("synchronize")
+
+    default_stream = DefaultStream()
+    monkeypatch.setattr(torch.cuda, "default_stream", lambda: default_stream)
+    monkeypatch.setattr(torch.cuda, "current_stream", lambda: default_stream)
+    monkeypatch.setattr(diverse_impl.g_infer_context, "get_overlap_stream", lambda: overlap_stream)
+    master = SimpleNamespace(
+        is_master_req=lambda: True,
+        get_chuncked_input_token_len=lambda: 4,
+        get_cur_total_len=lambda: 4,
+        cur_output_len=0,
+        slave_reqs=[object()],
+    )
+    backend = DiversehBackend.__new__(DiversehBackend)
+    backend.is_master_in_dp = False
+    backend._master_req_to_radix_cache = lambda master_req: order.append("rewrite_indexes")
+
+    packs = backend._diverse_pre_post_handle([master], is_chuncked_mode)
+
+    assert order == ["wait_model", "rewrite_indexes", "synchronize"]
+    assert packs[0].req_obj is master
+    assert packs[0].output_len == 1
diff --git a/unit_tests/server/router/model_infer/mode_backend/test_multi_level_kv_cache.py b/unit_tests/server/router/model_infer/mode_backend/test_multi_level_kv_cache.py
index 260dbe1065..24e8650f3e 100644
--- a/unit_tests/server/router/model_infer/mode_backend/test_multi_level_kv_cache.py
+++ b/unit_tests/server/router/model_infer/mode_backend/test_multi_level_kv_cache.py
@@ -11,7 +11,9 @@
 )
 from lightllm.server.router.model_infer.mode_backend.multi_level_kv_cache import MultiLevelKvCacheModule
 from lightllm.server.router.model_infer.mode_backend import multi_level_kv_cache as multi_level_kv_cache_impl
-from lightllm.server.router.model_infer.infer_batch import InferenceContext
+from lightllm.server.router.model_infer.infer_batch import InferReq, InferenceContext
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
+from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
 
 
 def test_cache_tiers_reassignment_is_rejected():
@@ -42,10 +44,11 @@ def test_non_gpu_cache_tiers_release_owned_tokens_without_radix_insert():
     released_refs = []
     context = InferenceContext()
     context.is_hybrid_att_model = False
-    context.req_manager = SimpleNamespace(req_to_token_indexs=torch.tensor([[10, 11, 12, 13, 14, 15, 16]]))
+    context.req_manager = make_cpu_index_manager(token_table=torch.tensor([[10, 11, 12, 13, 14, 15, 16]]), page_size=1)
     context.radix_cache = SimpleNamespace(dec_node_ref_counter=released_refs.append)
     shared_node = SimpleNamespace(node_prefix_total_len=2)
     req = SimpleNamespace(req_idx=0, cur_kv_len=5, hold_kv_len=7, shared_kv_node=shared_node)
+    bind_cpu_indexes(req, context.req_manager)
     free_token_indexes = []
 
     context._free_req_mem_without_radix_insert(free_token_indexes, req)
@@ -59,14 +62,18 @@ def test_legacy_cache_tiers_still_insert_gpu_radix_cache():
     context = InferenceContext()
     context.radix_cache = object()
     context.is_hybrid_att_model = False
+    context.req_manager = make_cpu_index_manager(token_table=torch.full((1, 4), -1, dtype=torch.int32))
     inserted_reqs = []
     context._full_att_free_req = lambda free_token_index, req: inserted_reqs.append(req)
     req = SimpleNamespace(
+        req_idx=0,
         cache_tiers=(CacheTier.GPU, CacheTier.CPU, CacheTier.DISK),
         cur_kv_len=3,
         shm_req=SimpleNamespace(shm_cur_kv_len=3),
     )
 
+    bind_cpu_indexes(req, context.req_manager)
+
     context.free_a_req_mem([], req)
 
     assert inserted_reqs == [req]
@@ -136,7 +143,7 @@ def test_non_gpu_linear_cache_tiers_release_pending_state_pages():
     freed_big_pages = []
     context = InferenceContext()
     context.is_hybrid_att_model = True
-    context.req_manager = SimpleNamespace(req_to_token_indexs=torch.tensor([[10, 11, 12]]))
+    context.req_manager = make_cpu_index_manager(token_table=torch.tensor([[10, 11, 12]]), page_size=1)
     context.radix_cache = SimpleNamespace(
         small_page_buffers=SimpleNamespace(free_state_cache=freed_small_pages.extend),
         big_page_buffers=SimpleNamespace(free_state_cache=freed_big_pages.extend),
@@ -149,6 +156,7 @@ def test_non_gpu_linear_cache_tiers_release_pending_state_pages():
         tail_small_page_buffer_id=7,
         hybrid_len_to_big_page_id={128: 8, 256: 9},
     )
+    bind_cpu_indexes(req, context.req_manager)
     free_token_indexes = []
 
     context._free_req_mem_without_radix_insert(free_token_indexes, req)
@@ -160,8 +168,12 @@ def test_non_gpu_linear_cache_tiers_release_pending_state_pages():
     assert req.hybrid_len_to_big_page_id == {}
 
 
-def test_cpu_cache_load_uses_exact_aligned_size(monkeypatch):
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires pinned CPU buffers")
+def test_cpu_cache_load_uses_exact_aligned_size(monkeypatch, mode):
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
     table = torch.full((1, 136), -1, dtype=torch.int32)
+    gpu_table = torch.full_like(table, -2)
     table[0, :4] = torch.arange(4, dtype=torch.int32)
     next_index = 4
     alloc_sizes = []
@@ -170,28 +182,30 @@ def test_cpu_cache_load_uses_exact_aligned_size(monkeypatch):
     def alloc(need_size):
         nonlocal next_index
         alloc_sizes.append(need_size)
-        mem_indexes = torch.arange(next_index, next_index + need_size, dtype=torch.int32)
+        mem_indexes = torch.arange(next_index, next_index + need_size, dtype=torch.int32, pin_memory=True)
         next_index += need_size
         return mem_indexes
 
-    def alloc_req_kv_mem(req, alloc_token_num):
+    def alloc_kv_mem(alloc_token_num):
         mem_indexes = alloc(alloc_token_num)
         new_hold_kv_len = req.hold_kv_len + alloc_token_num
+        req.token_index.append_token_indexes(mem_indexes)
         table[req.req_idx, req.hold_kv_len : new_hold_kv_len] = mem_indexes
         req.hold_kv_len = new_hold_kv_len
         return mem_indexes
 
-    operator = SimpleNamespace(
-        load_cpu_cache_to_gpu=lambda mem_indexes, **kwargs: loaded_indexes.extend(mem_indexes.tolist())
-    )
+    def load(mem_indexes, **kwargs):
+        assert mem_indexes.data_ptr() == gpu_table[0, :132].data_ptr()
+        loaded_indexes.extend(mem_indexes.tolist())
+
+    operator = SimpleNamespace(load_cpu_cache_to_gpu=load)
     backend = SimpleNamespace(
         is_master_in_dp=False,
         radix_cache=None,
         model=SimpleNamespace(
             mem_manager=SimpleNamespace(operator=operator),
-            req_manager=SimpleNamespace(req_to_token_indexs=table),
+            req_manager=make_cpu_index_manager(token_table=table, req_to_token_indexs=gpu_table, page_size=4),
         ),
-        _alloc_req_kv_mem=alloc_req_kv_mem,
     )
     module = MultiLevelKvCacheModule.__new__(MultiLevelKvCacheModule)
     module.args = SimpleNamespace(page_size=4, cpu_cache_token_page_size=132)
@@ -200,7 +214,7 @@ def alloc_req_kv_mem(req, alloc_token_num):
     module.need_sync_compute_stream = lambda: False
     module.cpu_cache_client = SimpleNamespace()
     context = SimpleNamespace(
-        req_manager=SimpleNamespace(req_to_token_indexs=table, mem_manager=SimpleNamespace(alloc=alloc)),
+        req_manager=make_cpu_index_manager(token_table=table, mem_manager=SimpleNamespace(alloc=alloc), page_size=1),
         get_can_alloc_token_num=lambda: 200,
     )
     req = SimpleNamespace(
@@ -214,15 +228,116 @@ def alloc_req_kv_mem(req, alloc_token_num):
             token_hash_page_len_list=SimpleNamespace(get_all=lambda: [132]),
         ),
         sampling_param=SimpleNamespace(shm_param=SimpleNamespace(prompt_logprobs=-1)),
+        alloc_kv_mem=alloc_kv_mem,
     )
+    bind_cpu_indexes(req, backend.model.req_manager)
+    req.args.req_token_index_mode = mode
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    req.token_index.append_token_indexes(table[0, :4].pin_memory())
+    backend.model.req_manager._req_token_indexes[0] = req.token_index
+    context.req_manager = backend.model.req_manager
     monkeypatch.setattr(multi_level_kv_cache_impl, "g_infer_context", context)
     monkeypatch.setattr(multi_level_kv_cache_impl.dist, "barrier", lambda group: None)
-    monkeypatch.setattr(torch.Tensor, "cuda", lambda self, **kwargs: self)
+    uploaded_indexes = []
+
+    def upload(tensor, **kwargs):
+        uploaded_indexes.append(tensor)
+        return tensor
+
+    def unexpected_copy(*args, **kwargs):
+        pytest.fail("CPU page-index upload must not concatenate GPU tensors")
+
+    monkeypatch.setattr(torch.Tensor, "cuda", upload)
+    monkeypatch.setattr(torch.Tensor, "pin_memory", lambda tensor: tensor)
+    monkeypatch.setattr(torch, "cat", unexpected_copy)
 
     module.load_cpu_cache_to_reqs([req])
 
+    assert len(uploaded_indexes) == 1  # 仅页面编号使用 .cuda()；索引上传到已有 GPU 表。
     assert alloc_sizes == [128]
     assert req.cur_kv_len == 132
     assert req.hold_kv_len == 132
     assert loaded_indexes == list(range(132))
     assert table[0, :132].tolist() == list(range(132))
+    assert gpu_table[0, :132].tolist() == list(range(132))
+    assert gpu_table[0, 132:].eq(-2).all()
+
+
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("use_default_index_stream", [False, True])
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_cpu_cache_offload_getter_orders_transfer_stream(monkeypatch, mode, use_default_index_stream):
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+    default_stream = torch.cuda.default_stream()
+    index_stream = default_stream if use_default_index_stream else torch.cuda.Stream()
+    index_stream.wait_stream(default_stream)
+    transfer_stream = torch.cuda.Stream()
+    gpu_indexes = torch.full((64,), -1, dtype=torch.int32, device="cuda")
+    req = SimpleNamespace(
+        req_idx=0,
+        req_manager=SimpleNamespace(req_to_token_indexs=gpu_indexes.unsqueeze(0), page_size=64),
+        args=SimpleNamespace(req_token_index_mode=mode),
+        cache_tiers=(CacheTier.CPU,),
+        cur_kv_len=64,
+        shm_req=SimpleNamespace(
+            input_len=64,
+            token_hash_list=SimpleNamespace(get_all=lambda: [123]),
+            token_hash_page_len_list=SimpleNamespace(get_all=lambda: [64]),
+        ),
+    )
+    helper = create_req_token_index(req, index_stream=index_stream)
+    req.token_index = helper
+    expected = torch.arange(64, dtype=torch.int32, pin_memory=True)
+    # 延迟索引流上的写入，验证 getter 在卸载流调用时也会承接依赖。
+    with torch.cuda.stream(index_stream):
+        torch.cuda._sleep(20_000_000)
+    helper.append_token_indexes(expected)
+    order = []
+    observed_indexes = []
+    original_get_gpu_token_indexes = helper.get_gpu_token_indexes
+
+    def prepare_indexes(start, end):
+        assert torch.cuda.current_stream() == transfer_stream
+        order.append("prepare")
+        return original_get_gpu_token_indexes(start, end)
+
+    original_wait_stream = torch.cuda.Stream.wait_stream
+
+    def wait_stream(waiting_stream, source_stream):
+        assert waiting_stream == transfer_stream and source_stream == index_stream
+        order.append("transfer_waits_index")
+        original_wait_stream(waiting_stream, source_stream)
+
+    def offload(mem_indexes, **kwargs):
+        assert torch.cuda.current_stream() == transfer_stream
+        order.append("offload")
+        observed_indexes.append(mem_indexes.clone())
+
+    def broadcast(data, **kwargs):
+        data[0] = {"item_size": 1, "page_list": [3], "ready_list": [False]}
+
+    req.token_index.get_gpu_token_indexes = prepare_indexes
+    module = MultiLevelKvCacheModule.__new__(MultiLevelKvCacheModule)
+    module.backend = SimpleNamespace(
+        is_master_in_dp=False,
+        dp_world_size=1,
+        model=SimpleNamespace(
+            mem_manager=SimpleNamespace(operator=SimpleNamespace(offload_gpu_kv_to_cpu_cache=offload))
+        ),
+    )
+    module.gloo_group = object()
+    module.cpu_cache_client = object()
+    module.page_index_buffer = torch.empty(8, dtype=torch.int32, device="cuda")
+    module.page_ready_buffer = torch.empty(8, dtype=torch.bool, device="cuda")
+    monkeypatch.setattr(torch.cuda.Stream, "wait_stream", wait_stream)
+    monkeypatch.setattr(multi_level_kv_cache_impl.dist, "broadcast_object_list", broadcast)
+
+    task = module._start_kv_cache_offload_task(req, transfer_stream)
+    task.sync_event.synchronize()
+
+    assert order == ["prepare", "transfer_waits_index", "offload"]
+    assert torch.cuda.current_stream() == default_stream
+    assert torch.equal(observed_indexes[0].cpu(), expected)
+    assert task.req_obj is req
+    assert task.move_token_num == 64
+    assert req.cpu_cache_task_status == InferReq._CpuCacheTaskStatus.RUNNING
diff --git a/unit_tests/server/router/model_infer/mode_backend/test_paged_side_allocations.py b/unit_tests/server/router/model_infer/mode_backend/test_paged_side_allocations.py
index 89e15fc93d..d31f2f1e17 100644
--- a/unit_tests/server/router/model_infer/mode_backend/test_paged_side_allocations.py
+++ b/unit_tests/server/router/model_infer/mode_backend/test_paged_side_allocations.py
@@ -8,9 +8,12 @@
 # 在独立运行本测试文件时形成模块级循环导入。
 from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
 from lightllm.server.router.model_infer.infer_batch import InferReq
+from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
+from lightllm.server.router.model_infer.mode_backend import base_backend
 from lightllm.server.router.model_infer.mode_backend.dp_backend import dp_shared_kv_trans
 from lightllm.server.router.model_infer.mode_backend.pd.decode_node_impl import decode_impl
 from lightllm.server.router.model_infer.mode_backend.pd.prefill_node_impl import prefill_impl
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
 
 
 def _bind_alloc_need(req, page_size):
@@ -40,9 +43,8 @@ def test_pd_decode_reserves_model_pages_but_transfers_only_logical_kv(monkeypatc
     tasks = []
     backend = decode_impl.PDDecodeNode.__new__(decode_impl.PDDecodeNode)
     backend.args = SimpleNamespace(pd_kv_page_size=4, page_size=4)
-    backend.model = SimpleNamespace(req_manager=SimpleNamespace(req_to_token_indexs=table))
+    backend.model = SimpleNamespace(req_manager=make_cpu_index_manager(token_table=table, page_size=1))
     backend.is_master_in_dp = False
-    backend._alloc_req_kv_mem = lambda req, size: _reserve_into_table(table, req, alloc_sizes, next_index, size)
 
     def create_task(**kwargs):
         tasks.append(
@@ -68,6 +70,8 @@ def create_task(**kwargs):
         page_size=4,
     )
 
+    req.alloc_kv_mem = lambda size: _reserve_into_table(table, req, alloc_sizes, next_index, size)
+
     backend._decode_node_gen_trans_tasks(req)
 
     assert alloc_sizes == [8]
@@ -83,6 +87,55 @@ def create_task(**kwargs):
     assert table[0, :12].tolist() == list(range(12))
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA and pinned CPU memory")
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("input_len", [128, 130])
+def test_pd_prefill_transfers_exact_logical_tail_with_real_index_helper(monkeypatch, mode, input_len):
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+    monkeypatch.setattr(prefill_impl.g_infer_context, "is_hybrid_att_model", False)
+    page_size = 64
+    token_num = (input_len + page_size - 1) // page_size * page_size
+    req = SimpleNamespace(
+        req_idx=0,
+        req_id=1,
+        req_manager=SimpleNamespace(
+            page_size=page_size, req_to_token_indexs=torch.empty((1, token_num), dtype=torch.int32, device="cuda")
+        ),
+        args=SimpleNamespace(req_token_index_mode=mode),
+        shm_req=SimpleNamespace(input_len=input_len),
+        sampling_param=SimpleNamespace(
+            pd_master_node_id=1,
+            pd_decode_node=SimpleNamespace(agent_name="decode", agent_metadata=None, num_pages=1, page_reg_desc=None),
+        ),
+        cur_kv_len=input_len,
+        hold_kv_len=token_num,
+        pd_trans_kv_start_index=0,
+        pd_trans_device_id=0,
+        pd_task_num=0,
+    )
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    page_starts = torch.tensor([9, 2, 15], dtype=torch.int32)[: token_num // page_size] * page_size
+    expected = (page_starts[:, None] + torch.arange(page_size, dtype=torch.int32)).reshape(-1)
+    req.token_index.append_token_indexes(expected)
+    backend = prefill_impl.PDChunkedPrefillForPrefillNode.__new__(prefill_impl.PDChunkedPrefillForPrefillNode)
+    backend.args = SimpleNamespace(page_size=page_size, pd_kv_page_size=128)
+    backend.is_master_in_dp = True
+    backend.dp_rank_in_node = 0
+    tasks = []
+    backend.info_queue = SimpleNamespace(put=tasks.append)
+
+    backend._prefill_chuncked_handle_func(req, next_token_id=7, next_token_prob=-0.5, output_len=1)
+
+    assert req.pd_task_num == len(tasks)
+    assert req.pd_trans_kv_start_index == input_len
+    assert [(task.start_kv_index, task.end_kv_index) for task in tasks] == (
+        [(0, 128)] if input_len == 128 else [(0, 128), (128, 130)]
+    )
+    assert [index for task in tasks for index in task.mem_indexes] == expected[:input_len].tolist()
+    assert tasks[-1].first_gen_token_id == 7
+    assert tasks[-1].first_gen_token_logprob == -0.5
+
+
 def test_pd_decode_rejects_unaligned_transfer_page_size():
     backend = decode_impl.PDDecodeNode.__new__(decode_impl.PDDecodeNode)
     backend.args = SimpleNamespace(pd_kv_page_size=3, page_size=4)
@@ -101,6 +154,81 @@ def test_pd_prefill_rejects_unaligned_transfer_page_size():
         backend._prefill_chuncked_handle_func(req, next_token_id=1, next_token_prob=1.0, output_len=0)
 
 
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("is_master_in_dp", [False, True])
+def test_dp_cache_fetch_publishes_gpu_indexes_before_barrier_returns(monkeypatch, is_master_in_dp):
+    cpu_table = torch.empty((2, 16), dtype=torch.int32, pin_memory=True)
+    torch.arange(32, out=cpu_table.view(-1))
+    gpu_table = torch.full((2, 16), -1, dtype=torch.int32, device="cuda")
+    gpu_table[1, :2].copy_(cpu_table[1, :2], non_blocking=True)
+    req_manager = make_cpu_index_manager(token_table=cpu_table, req_to_token_indexs=gpu_table, page_size=4)
+    reqs = [
+        SimpleNamespace(req_idx=1, cur_kv_len=4, hold_kv_len=4),
+        SimpleNamespace(req_idx=0, cur_kv_len=4, hold_kv_len=8),
+    ]
+    for req, gpu_index_submitted_len in zip(reqs, (2, 0)):
+        bind_cpu_indexes(req, req_manager, gpu_index_submitted_len=gpu_index_submitted_len)
+    default_stream = torch.cuda.default_stream()
+    # 索引旧写入统一在默认流；上传流必须等待它，随后默认流再等待上传完成。
+    torch.cuda._sleep(20_000_000)
+    gpu_table[0, :8].fill_(-2)
+    context = SimpleNamespace(req_manager=req_manager)
+    monkeypatch.setattr(dp_shared_kv_trans, "g_infer_context", context)
+    monkeypatch.setattr(base_backend, "g_infer_context", context)
+    backend = SimpleNamespace(node_nccl_group=object(), is_master_in_dp=is_master_in_dp)
+    order = []
+
+    def upload(run_reqs):
+        assert torch.cuda.current_stream() == default_stream
+        order.append("upload")
+        base_backend.ModeBackend._copy_req_token_indexes_to_gpu(backend, run_reqs)
+
+    original_synchronize = torch.cuda.Stream.synchronize
+    original_wait_stream = torch.cuda.Stream.wait_stream
+
+    def forbid_synchronize(*args):
+        pytest.fail("DP shared must rely on NCCL barrier rather than explicitly synchronizing a stream")
+
+    def wait_stream(waiting_stream, source_stream):
+        if waiting_stream != default_stream:
+            assert source_stream == default_stream
+            order.append("upload_waits_default")
+        else:
+            assert source_stream != default_stream
+            order.append("default_waits_upload")
+        original_wait_stream(waiting_stream, source_stream)
+
+    def barrier(group):
+        assert group is backend.node_nccl_group
+        assert order == ["upload", "upload_waits_default", "default_waits_upload"]
+        # 模拟同步 NCCL barrier 在 C++ 层等待默认流完成。
+        original_synchronize(default_stream)
+        assert torch.equal(gpu_table[1, :4].cpu(), cpu_table[1, :4])
+        assert torch.equal(gpu_table[0, :8].cpu(), cpu_table[0, :8])
+        assert gpu_table[0, 8:].eq(-1).all()
+        assert gpu_table[1, 4:].eq(-1).all()
+        order.append("barrier")
+
+    backend._copy_req_token_indexes_to_gpu = upload
+    monkeypatch.setattr(torch.cuda.Stream, "synchronize", forbid_synchronize)
+    monkeypatch.setattr(torch.cuda.Stream, "wait_stream", wait_stream)
+    monkeypatch.setattr(dp_shared_kv_trans.dist, "barrier", barrier)
+    module = dp_shared_kv_trans.DPKVSharedMoudle.__new__(dp_shared_kv_trans.DPKVSharedMoudle)
+    module.backend = backend
+    module.dp_rank_in_node = 0
+    module.shared_req_infos = SimpleNamespace(arr=np.full((2, 2, 2), -1, dtype=np.int64))
+
+    module.fill_reqs_info(reqs)
+
+    assert order == ["upload", "upload_waits_default", "default_waits_upload", "barrier"]
+    assert [req.token_index.gpu_index_submitted_len for req in reqs] == [4, 8]
+    if is_master_in_dp:
+        assert module.shared_req_infos.arr[:, 0].tolist() == [[4, 1], [4, 0]]
+    else:
+        assert (module.shared_req_infos.arr == -1).all()
+    assert cpu_table.is_pinned() and not cpu_table.is_shared()
+
+
 def test_dp_cache_fetch_reserves_pages_and_keeps_logical_transfer_size(monkeypatch):
     table = torch.full((1, 16), -1, dtype=torch.int32)
     table[0, :4] = torch.arange(4, dtype=torch.int32)
@@ -123,13 +251,13 @@ def test_dp_cache_fetch_reserves_pages_and_keeps_logical_transfer_size(monkeypat
         node_nccl_group=None,
         mem_managers=[SimpleNamespace(req_to_token_indexs=source_table)] * 2,
     )
-    backend._alloc_req_kv_mem = lambda cur_req, size: _reserve_into_table(table, cur_req, alloc_sizes, next_index, size)
+    req.alloc_kv_mem = lambda size: _reserve_into_table(table, req, alloc_sizes, next_index, size)
     module = dp_shared_kv_trans.DPKVSharedMoudle.__new__(dp_shared_kv_trans.DPKVSharedMoudle)
     module.backend = backend
     module.dp_rank_in_node = 0
     module.shared_req_infos = SimpleNamespace(arr=np.array([[[4, 0], [10, 0]]], dtype=np.int64))
     context = SimpleNamespace(
-        req_manager=SimpleNamespace(req_to_token_indexs=table),
+        req_manager=make_cpu_index_manager(token_table=table, page_size=1),
         get_can_alloc_token_num=lambda: 100,
     )
     monkeypatch.setattr(dp_shared_kv_trans, "g_infer_context", context)
diff --git a/unit_tests/server/router/model_infer/mode_backend/test_prompt_logprobs_cache.py b/unit_tests/server/router/model_infer/mode_backend/test_prompt_logprobs_cache.py
new file mode 100644
index 0000000000..7450e23dd4
--- /dev/null
+++ b/unit_tests/server/router/model_infer/mode_backend/test_prompt_logprobs_cache.py
@@ -0,0 +1,172 @@
+from types import SimpleNamespace
+from unittest.mock import Mock
+
+import pytest
+import torch
+
+# Follow the production import order to avoid the existing package initialization cycle.
+from lightllm.server.router.model_infer.mode_backend import generic_pre_process  # noqa: F401
+from lightllm.server.core.objs.sampling_params import SamplingParams
+from lightllm.server.router.model_infer import infer_batch
+from lightllm.server.router.model_infer.infer_batch import InferReq, InferSamplingParams
+from lightllm.server.router.model_infer.infer_req_ext1 import PinMemPageCacheReqTokenIndex, create_req_token_index
+from lightllm.server.router.model_infer.mode_backend import base_backend
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
+
+
+def _make_req_with_cached_prefix(monkeypatch, prompt_logprobs, disable_prompt_cache=False, hybrid=False):
+    sample_params = SamplingParams()
+    sample_params.top_k = -1
+    sample_params.prompt_logprobs = prompt_logprobs
+    sample_params.disable_prompt_cache = disable_prompt_cache
+    req = InferReq.__new__(InferReq)
+    req.req_idx = 0
+    req.args = SimpleNamespace(
+        page_size=2, linear_att_hash_page_size=2, linear_att_page_block_num=2, max_req_total_len=16
+    )
+    req.cur_kv_len = req.hold_kv_len = 0
+    req.shared_kv_node = req.tail_small_page_buffer_id = None
+    req.shm_req = SimpleNamespace(
+        sample_params=sample_params,
+        input_len=6,
+        prompt_cache_len=0,
+        shm_prompt_ids=SimpleNamespace(arr=torch.arange(6).numpy()),
+        hybrid_token_hash_list=SimpleNamespace(size=2, get_all=lambda: [10, 20]),
+    )
+    req.sampling_param = InferSamplingParams(req.shm_req, vocab_size=8)
+    req.get_cur_total_len = lambda: 6
+    node = SimpleNamespace(node_prefix_total_len=4, big_page_buffer_idx=0, is_big_page_node=lambda: True)
+    match_prefix = Mock(return_value=(node, 4, torch.arange(4, dtype=torch.int32)))
+    manager = make_cpu_index_manager(
+        token_table=torch.zeros((1, 8), dtype=torch.int32),
+        restore_big_page_state=Mock(),
+        init_hybrid_attention_state=Mock(),
+        page_size=2,
+    )
+    bind_cpu_indexes(req, manager)
+    monkeypatch.setattr(
+        infer_batch,
+        "g_infer_context",
+        SimpleNamespace(
+            is_hybrid_att_model=hybrid, req_manager=manager, radix_cache=SimpleNamespace(match_prefix=match_prefix)
+        ),
+    )
+    return req, match_prefix
+
+
+@pytest.mark.parametrize("hybrid", [False, True])
+@pytest.mark.parametrize(
+    "prompt_logprobs, disable_prompt_cache, expected_prefix_len",
+    [(-1, False, 4), (0, False, 0), (3, False, 0), (-1, True, 0)],
+)
+def test_prompt_logprobs_requests_recompute_cached_prefix(
+    monkeypatch, hybrid, prompt_logprobs, disable_prompt_cache, expected_prefix_len
+):
+    req, match_prefix = _make_req_with_cached_prefix(monkeypatch, prompt_logprobs, disable_prompt_cache, hybrid)
+
+    if hybrid:
+        req._hybrid_match_radix_cache()
+    else:
+        req._match_radix_cache()
+
+    assert req.cur_kv_len == req.hold_kv_len == expected_prefix_len
+    assert req.shm_req.prompt_cache_len == expected_prefix_len
+    assert req.shm_req.shm_cur_kv_len == expected_prefix_len
+    assert match_prefix.call_count == (1 if expected_prefix_len else 0)
+
+
+@pytest.mark.parametrize("prompt_logprobs", [0, 3])
+def test_cached_prompt_still_captures_logprobs_for_every_prompt_position(monkeypatch, prompt_logprobs):
+    req, _ = _make_req_with_cached_prefix(monkeypatch, prompt_logprobs)
+    req._match_radix_cache()
+    req.prompt_selected_logprobs = SimpleNamespace(add_chunk=Mock())
+    backend = base_backend.ModeBackend.__new__(base_backend.ModeBackend)
+    backend.disable_chunked_prefill = True
+    backend.is_master_in_dp = True
+    q_len = req.get_cur_total_len() - req.cur_kv_len
+    indexes = torch.arange(q_len, dtype=torch.int32)
+    backend.model = SimpleNamespace(return_all_prompt_logics=True, _select_mem_indexes=lambda _: indexes)
+    capture_manager = SimpleNamespace(is_buffer_initialized=lambda: True, capture=Mock())
+    monkeypatch.setattr(base_backend.PromptLogprobsCaptureManager, "get_instance", lambda: capture_manager)
+    logits = torch.arange(q_len * 8, dtype=torch.float32).reshape(q_len, 8)
+
+    backend._capture_prompt_logprobs_if_needed(model_input=None, run_reqs=[req], prompt_logits=logits)
+
+    if prompt_logprobs == 0:
+        req.prompt_selected_logprobs.add_chunk.assert_called_once()
+        start, end, logprobs, ranks = req.prompt_selected_logprobs.add_chunk.call_args.args
+        assert (start, end) == (1, req.shm_req.input_len)
+        assert logprobs.shape == ranks.shape == (5,)
+        expected = logits[:5].log_softmax(dim=-1)[torch.arange(5), torch.arange(1, 6)]
+        torch.testing.assert_close(logprobs, expected)
+    else:
+        capture_manager.capture.assert_called_once()
+        captured = capture_manager.capture.call_args.kwargs
+        assert captured["mem_indexes"].tolist() == list(range(5))
+        assert captured["top_token_ids"].shape == captured["top_logprobs"].shape == (5, prompt_logprobs)
+
+
+@pytest.mark.parametrize("kind", ["prompt_logprobs", "routed_experts"])
+def test_final_metadata_extracts_owned_token_indexes_through_req_helper(monkeypatch, kind):
+    from lightllm.server.router.model_infer import infer_req_ext
+
+    req, _ = _make_req_with_cached_prefix(monkeypatch, prompt_logprobs=3)
+    req.cur_kv_len = 6
+    req.cur_output_len = 0
+    req.token_index.append_token_indexes(torch.tensor([8, 9, 2, 3, 14, 15], dtype=torch.int32))
+    manager = SimpleNamespace(extract=Mock(return_value="metadata"))
+    metadata = infer_req_ext.FinalTokenMetadataExt(req)
+    if kind == "prompt_logprobs":
+        monkeypatch.setattr(infer_req_ext.PromptLogprobsCaptureManager, "get_instance", lambda: manager)
+        result = metadata.collect_prompt_logprobs()
+    else:
+        monkeypatch.setattr(infer_req_ext.MoeRouteInfoManager, "get_instance", lambda: manager)
+        result = metadata.collect_routed_experts()
+    assert result == "metadata"
+    indexes = manager.extract.call_args.args[0]
+    assert indexes.tolist() == [8, 9, 2, 3, 14]
+    req.token_index.reset()
+    req.token_index.append_token_indexes(torch.arange(6, dtype=torch.int32))
+    assert indexes.tolist() == [8, 9, 2, 3, 14]
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA and pinned CPU memory")
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pin_mem_page_cache", "pure_gpu"])
+@pytest.mark.parametrize("input_len", [128, 130])
+@pytest.mark.parametrize("kind", ["prompt_logprobs", "routed_experts"])
+def test_final_metadata_reads_partial_pages_with_real_index_helper(monkeypatch, mode, input_len, kind):
+    from lightllm.server.router.model_infer import infer_req_ext
+
+    monkeypatch.setattr(PinMemPageCacheReqTokenIndex, "_page_offsets", None)
+    page_size = 64
+    token_num = (input_len + page_size - 1) // page_size * page_size
+    req = SimpleNamespace(
+        req_idx=0,
+        req_manager=SimpleNamespace(
+            page_size=page_size, req_to_token_indexs=torch.empty((1, token_num), dtype=torch.int32, device="cuda")
+        ),
+        args=SimpleNamespace(req_token_index_mode=mode),
+        shm_req=SimpleNamespace(input_len=input_len),
+        sampling_param=SimpleNamespace(shm_param=SimpleNamespace(prompt_logprobs=3)),
+        cur_kv_len=input_len,
+        cur_output_len=0,
+    )
+    req.token_index = create_req_token_index(req, index_stream=torch.cuda.default_stream())
+    page_starts = torch.tensor([9, 2, 15], dtype=torch.int32)[: token_num // page_size] * page_size
+    expected = (page_starts[:, None] + torch.arange(page_size, dtype=torch.int32)).reshape(-1)
+    req.token_index.append_token_indexes(expected)
+    manager = SimpleNamespace(extract=Mock(return_value="metadata"))
+    metadata = infer_req_ext.FinalTokenMetadataExt(req)
+    if kind == "prompt_logprobs":
+        monkeypatch.setattr(infer_req_ext.PromptLogprobsCaptureManager, "get_instance", lambda: manager)
+        result = metadata.collect_prompt_logprobs()
+    else:
+        monkeypatch.setattr(infer_req_ext.MoeRouteInfoManager, "get_instance", lambda: manager)
+        result = metadata.collect_routed_experts()
+
+    assert result == "metadata"
+    indexes = manager.extract.call_args.args[0]
+    assert torch.equal(indexes, expected[: input_len - 1])
+    req.token_index.reset()
+    req.token_index.append_token_indexes(torch.arange(token_num, dtype=torch.int32))
+    assert torch.equal(indexes, expected[: input_len - 1])
diff --git a/unit_tests/test_static_benchmark_paged_kv.py b/unit_tests/test_static_benchmark_paged_kv.py
index 6b57270125..eba292cda4 100644
--- a/unit_tests/test_static_benchmark_paged_kv.py
+++ b/unit_tests/test_static_benchmark_paged_kv.py
@@ -3,6 +3,7 @@
 from types import SimpleNamespace
 
 import numpy as np
+import pytest
 import torch
 
 
@@ -12,12 +13,15 @@
     if str(path) not in sys.path:
         sys.path.append(str(path))
 
+import static_benchmark
+
 from static_benchmark import (
     BenchmarkCase,
     StaticBenchmarkExecutor,
     decode_profile_batch_divisor,
     resolve_batch_max_prefill_cases,
 )
+from unit_tests.common.cpu_index_helpers import make_cpu_index_manager, bind_cpu_indexes
 
 
 class FakeMemManager:
@@ -29,20 +33,23 @@ def __init__(self, page_size: int):
     def alloc(self, size: int) -> torch.Tensor:
         assert size % self.page_size == 0
         self.alloc_sizes.append(size)
-        indexes = torch.arange(self.next_index, self.next_index + size, dtype=torch.int32)
+        indexes = torch.arange(self.next_index, self.next_index + size, dtype=torch.int32, pin_memory=True)
         self.next_index += size
         return indexes
 
 
-def make_executor(page_size: int = 4):
+def make_executor(monkeypatch, page_size: int = 4, mode="pin_mem_cache"):
     mem_manager = FakeMemManager(page_size=page_size)
-    req_manager = SimpleNamespace(
+    req_manager = make_cpu_index_manager(
         mem_manager=mem_manager,
+        token_table=torch.full((4, 32), -1, dtype=torch.int32),
         req_to_token_indexs=torch.full((4, 32), -1, dtype=torch.int32),
+        page_size=page_size,
     )
     model = SimpleNamespace(mem_manager=mem_manager, req_manager=req_manager)
-    args = SimpleNamespace(mtp_mode=None, mtp_step=0)
+    args = SimpleNamespace(mtp_mode=None, mtp_step=0, req_token_index_mode=mode)
     executor = StaticBenchmarkExecutor(args=args, model=model, draft_models=[], token_source=None)
+    req_manager._req_token_indexes = executor._req_token_indexes
     return executor, mem_manager, req_manager
 
 
@@ -52,8 +59,10 @@ def assert_complete_physical_pages(indexes: torch.Tensor, page_size: int):
         assert torch.all(page // page_size == page[0] // page_size)
 
 
-def test_prefill_allocates_and_reuses_complete_pages_per_request():
-    executor, mem_manager, req_manager = make_executor(page_size=4)
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pure_gpu"])
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires pinned CPU buffers")
+def test_prefill_allocates_and_reuses_complete_pages_per_request(monkeypatch, mode):
+    executor, mem_manager, req_manager = make_executor(monkeypatch, page_size=4, mode=mode)
     req_idx = torch.tensor([0, 1], dtype=torch.int32)
 
     model_input = executor._make_prefill_input(
@@ -63,10 +72,16 @@ def test_prefill_allocates_and_reuses_complete_pages_per_request():
     )
 
     assert mem_manager.alloc_sizes == [4, 4]
+    if mode == "pin_mem_cache":
+        assert executor._req_token_indexes[0]._cpu_buffer.is_pinned()
+    else:
+        assert not hasattr(executor._req_token_indexes[0], "_cpu_buffer")
+    for row in (0, 1):
+        assert torch.equal(req_manager.req_to_token_indexs[row, :4], req_manager.get_cpu_token_indexes(row, 0, 4))
     assert model_input.b_is_decode_req.tolist() == [False, False]
     assert executor._hold_kv_len_by_req == {0: 4, 1: 4}
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[0, :4], page_size=4)
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[1, :4], page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(0, 0, 4), page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(1, 0, 4), page_size=4)
 
     # 两个请求都还在已有页内继续 decode，不应产生新的分配。
     executor._make_decode_input(
@@ -90,12 +105,14 @@ def test_prefill_allocates_and_reuses_complete_pages_per_request():
     )
     assert mem_manager.alloc_sizes == [4, 4, 4, 4]
     assert executor._hold_kv_len_by_req == {0: 8, 1: 8}
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[0, :8], page_size=4)
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[1, :8], page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(0, 0, 8), page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(1, 0, 8), page_size=4)
 
 
-def test_mtp_expanded_rows_allocate_once_per_request():
-    executor, mem_manager, req_manager = make_executor(page_size=4)
+@pytest.mark.parametrize("mode", ["pin_mem_cache", "pure_gpu"])
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires pinned CPU buffers")
+def test_mtp_expanded_rows_allocate_once_per_request(monkeypatch, mode):
+    executor, mem_manager, req_manager = make_executor(monkeypatch, page_size=4, mode=mode)
     executor.args.mtp_mode = "eagle_with_att"
     executor.args.mtp_step = 3
 
@@ -112,8 +129,8 @@ def test_mtp_expanded_rows_allocate_once_per_request():
     assert model_input.batch_size == 4
     assert mem_manager.alloc_sizes == [8, 8]
     assert executor._hold_kv_len_by_req == {0: 8, 1: 8}
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[0, :8], page_size=4)
-    assert_complete_physical_pages(req_manager.req_to_token_indexs[1, :8], page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(0, 0, 8), page_size=4)
+    assert_complete_physical_pages(req_manager.get_cpu_token_indexes(1, 0, 8), page_size=4)
 
 
 def test_profile_capacity_uses_page_aligned_per_request_cost():