#45595 改了什么:

之前(旧机制):当所有请求都结束时,build_connector_meta() 会同步阻塞地把所有 in-flight 的 store job 全部 flush 掉(调用 worker.wait())。

现在(新机制):不再阻塞 flush,而是通过 has_pending_push_work() 让引擎持续 stepping,直到 job 自然完成。

相当于之前多了一个在请求结束的时候同步的逻辑,这里去掉可以增加并发度。

为什么旧机制是多余的

旧机制的触发条件是 all(rs.req.is_finished() for rs in self._req_status.values()),只有所有请求都结束了才触发。 生产环境中通常有连续请求,这个条件很少为 True。真正保护数据的一直都是 fence 冲突检测(block 被复用时强制 flush)。

旧机制只兜底了”引擎即将空闲但还有 job 没完成”这个边界情况,新机制用 has_pending_push_work() 同样兜住了,只是不阻塞。

异步传输原理

GPU→CPU 拷贝是异步的,不需要多线程:

  1. CPU 线程调用 cudaMemcpyAsync() 把拷贝任务提交给 GPU 的 DMA 引擎(硬件级别的搬运工)
  2. 拷贝跑在独立的 CUDA stream 上,不阻塞计算 stream(模型推理)
  3. CPU 线程提交完就返回,DMA 引擎在 GPU 硬件层面自己搬数据
  4. 每个 step 的 update_from_output() 调用 end_event.query() 非阻塞轮询完成状态

end_event.query() 原理:GPU 执行到 event 位置时硬件设置一个标志位,CPU 读这个标志位。 query() 读一次就返回(非阻塞),synchronize() 等到 True 才返回(阻塞)。

整个设计是单 CPU 线程 + GPU 硬件异步,三个独立执行单元交错运行:

对测试的影响

旧机制下测试依赖”请求结束 → 自动 flush”,新机制下去掉了这个保证。 测试需要改为 block 复用模式:Run 1 创建 job 留在 in-flight,Run 2 新请求复用同样的 GPU blocks → fence 检测到冲突 → 强制 flush。