AI 推理时代的 Ceph：让 GPU 直读 Ceph/RADOS KV|ceph|gpu|key|上下文|人工智能模型|字节|寻址

新钛云服已累计为您分享905篇技术干货

AI 推理为什么开始关心存储路径

大模型推理里，大家通常最关心 GPU 算力、显存容量、batch size、prefill、decode、TTFT（time to first token，首 token 延迟）。但当模型上下文越来越长，尤其是 agent 场景里动不动把系统提示词、工具定义、历史上下文、代码仓库片段塞进 prompt，另一个问题会变得越来越突出：KV cache 到底放在哪里，怎么取回来，代价有多大？

KV cache 的价值很直接：prefill 阶段已经算过的 key/value state，如果下次还能复用，就可以少算一大段上下文，降低 TTFT，也降低 GPU 上重复计算的成本。此前 Ceph 社区已经讨论过vLLM KV caching(https://ceph.io/en/news/blog/2025/vllm-kv-caching/) 相关工作，也在 Cephalocon 上分享过这个方向。

真正有意思的问题是：如果 KV cache 不只放在本机内存或本地 SSD，而是放在一个分布式存储系统里，GPU 能不能直接把它读回来？进一步说，能不能绕开 guest CPU 的数据路径，让 GPU 自己发起 I/O，把数据直接放进 GPU memory？

这就是本文的主线：

一块 AMD GPU 从 __device__ code 发起 NVMe Key-Value Retrieve，请求落到 SPDK/vfio-user 暴露出来的 NVMe-KV controller，后端由 Ceph/RADOS object 承载，返回的 value 通过 P2P-DMA 直接进入 GPU memory。

换成更短的话：GPU 不再等 CPU 搬数据，而是自己发起 I/O，从 Ceph/RADOS KV 里把 AI 推理需要的数据取回来。

为什么不能直接用块设备

去年有一篇论文很值得关注：

GPU-Initiated On-Demand High-Throughput Storage Access in the BaM System Architecture

（https://arxiv.org/pdf/2203.04910）。它展示了一种思路：CPU 负责把 kernel 加载到 GPU，随后 GPU 可以像 NVMe initiator 一样自己驱动存储 I/O。

这个方向很吸引人，因为它把 CPU 从 accelerator-to-storage 的数据路径里拿掉了。但如果把它直接套到 KV cache 场景，就会遇到接口不匹配的问题。

block 设备天然关心的是 (device, offset, length)。而 KV cache 更像内容寻址：一个 prompt prefix 或 token block 的 hash 就是 key，对应的 value 是已经计算好的 KV state。用 block 来做这件事，意味着还要维护一张中心化映射表，把序列 hash 映射到具体 offset。这样不仅多一层间接寻址，还会引入一致性和协调成本。

对 KV cache 来说，更自然的抽象是：

value = 对应的 KV state

也就是说，它本来就应该是 Key-Value，而不是 block。

把 RADOS 暴露成 NVMe Key-Value

2025 年，NVMe Key Value Command Set Specification有了第一个正式批准版本。与此同时，Ceph 已经有基于 SPDK 的 NVMe/TCP 实现。于是一个很自然的想法出现了：能不能把 RADOS object 通过 NVMe-KV 的形式暴露出去？

对 Ceph 不熟的读者，可以简单理解为：RADOS 是 Ceph 的底层对象存储接口，能力比传统对象存储更丰富，支持 object 级别的读写、删除、OMAP、watch/notify，以及在 object 上执行代码的 class 机制（`cls`）。

这次实验选择的映射关系非常直接：

NVMe-oF 概念 RADOS 概念 Subsystem Pool Namespace RADOS namespace Key（1-16 字节） Object name（oid） Value Object data

这样一来，一个 KV pair 就是一个 RADOS object。对 AI 推理来说，prefix hash 可以直接变成 key；不同 GPU、不同节点只要算出同一个 key，就会访问同一个 object，不需要中心化 coordinator，也不需要维护 block offset 表。

还有一个值得展开的方向：RADOS 支持在 object 上执行 `cls`。如果以后通过厂商自定义 NVMe command 或控制面 gRPC 服务，把经过 allowlist 的 class 映射成特定 opcode，就可以让存储不仅负责“放字节”，还可以在数据旁边做计算，比如生成缩略图、抽取 embedding，甚至做更贴近 AI 数据管线的预处理。这就是 computational storage 的思路。

rados-nkv：把 Ceph 接到 SPDK NVMe-KV 后面

实验从一个原型开始：rados-nkv。

6 月 4 日下午，有了 NVMe-KV 规范和初步方案后，作者把工作拆成多个任务，让 agent 辅助编码与 review。这个过程本身也很有意思：用 AI 工具去构建服务 AI workload 的底层存储能力，有点“递归造基础设施”的味道。

过程中还有一个意外收获：NVIDIA 最近已经向 SPDK 添加了一个NVMe KV client(https://github.com/spdk/spdk/commit/93cf1ebc57cea99caba6c68f30354f93dce90ffd)，这给原型开发提供了很好的参考。

最终形成的rados-nkv(https://review.spdk.io/c/spdk/spdk/+/28851) 是一个 SPDK controller：它通过vfio-user(https://github.com/nutanix/libvfio-user) 向外暴露 NVMe Key-Value namespace，后端则把 key/value 映射到 RADOS object。端到端测试已经能跑通。

上图展示的是通过 vfio-user loopback 验证的 KV 控制路径：host app 发出 Store/Retrieve/Delete/Exist，SPDK 的 NVMe-KV target 解码命令，rados-nkv kvdev 通过 async librados 访问 RADOS。

为了让它服务 AI 推理，还需要接入上层生态。

例如NIXL plugin(https://github.com/ai-dynamo/nixl/pull/1717)，

因为llm-d-kv-cache(https://github.com/llm-d/llm-d-kv-cache) 和LMCache(https://github.com/lmcache/lmcache) 都会通过 NIXL / kv-connector 从 vLLM offload KV-cache block。另一个补充方向是rados-nkv-weights

(https://github.com/mmgaggle/rados-nkv-weights)：让 vLLM 通过同一个 namespace 流式加载模型权重。

到这里，已经有了一个可用的 Ceph-backed NVMe-KV namespace。但这还只是 CPU 或普通 host 程序访问。真正的目标是：GPU-initiated I/O。

先用块设备路径做验证

要让 GPU 自己访问这个 namespace，实验需要真实硬件。我们准备了一台 Strix Halo Mini PC，里面有 AMD Ryzen AI / Radeon 8060S 这类集成 GPU。系统选择 Fedora 44，主要是为了使用足够新的 kernel，包含 dma-buf 相关改进。

实验环境大致是这样：

host 上运行 SPDK nvmf_tgt，通过 vfio-user 暴露一个模拟 NVMe controller；
QEMU guest 通过 vfio-user-pci 看到一个普通 /dev/nvme0；
iGPU 通过 passthrough 进入 guest；
guest 内运行 ROCm XIO，让 GPU kernel 自己驱动 NVMe queue pair。

不过，第一步并没有直接冲 KV。更稳妥的做法是先跑 block namespace：把一个由 RBD 后端支撑的传统 NVMe block namespace 传给 guest。

原因很简单：GPU-initiated block read 是已经被证明过的方向。如果 block 都跑不通，问题大概率在 passthrough、vfio-user、QEMU bridge 或 ROCm XIO plumbing 上，而不是新写的 KV 逻辑。

这是一条很重要的工程经验：做新东西时，先用已知可行的旧接口验证底层路径。

ROCm XIO 在路径里负责什么

ROCm XIO(https://github.com/ROCm/rocm-xio) 是这次实验里让 GPU 代码驱动 NVMe queue pair 的关键组件。它的 nvme-ep endpoint 会运行一个 __device__ kernel，大致做四件事：

构造 NVMe Submission Queue Entry（SQE，64 字节命令），写入 Submission Queue；
写 SQ-tail doorbell，告诉 controller 有新命令；
在 Completion Queue 上自旋，等待 phase bit 翻转；
写 CQ-head doorbell，表示 completion 已处理。

配套的 kernel module /dev/rocm-xio 负责特权态部分，包括 pin queue memory、处理 admin command、导出 GPU VRAM buffer 供 P2P-DMA 使用等。

ROCm XIO 允许选择 queue、doorbell 和 data buffer 放在 host RAM 还是 GPU VRAM 中。通过 --memory-mode 这个 bitmask 控制：

bit0x10x20x40x8对象 SQ CQ doorbell data buffer

0x0 表示都在 host RAM，是最容易跑起来的 baseline。0x8 表示 data buffer 放到 VRAM，这才是本文真正关心的路径：payload 通过 P2P-DMA 直接进入 GPU memory。

vfio-user 这里也有一个关键细节：SPDK 默认把 NVMe doorbell register（BAR0 offset 0x1000 之后）以 sparse-mmap 共享页的方式暴露出来。client 写 doorbell 时，不必每次都走 socket 往返；SPDK poll loop 会读取共享页里的 tail/head 值。

理论上，这条路径很顺：GPU 写 doorbell，QEMU/vfio-user/SPDK 看见，执行 I/O，再把 completion 写回去。

现实里，第一次运行失败了。

第一次失败：GPU 发出了请求，却等不到响应

第一次真正让 GPU 发起 I/O 时，程序卡在 completion polling 上：

gpuKernel sync failed for queue 8 (719)

现象是：GPU 构造了命令，也写了 doorbell，但 Completion Queue 里迟迟没有结果。`0x1016` 是 GPU kernel 放弃等待后触发的异常，不是根因。

麻烦在于，这台 APU 的 iGPU passthrough 调试成本很高。guest 里的 GPU fault 可能把 SMU（System Management Unit）带到一个无法重新初始化的状态：

amdgpu: hw_init of IP block failed -22

很多时候只有重启 host 才能恢复。换句话说，实验预算很残酷：每次 host reboot 之后，可能只有一次真正的 GPU shot。

在这种环境里，靠猜和反复试错不可持续。调试方法必须换成：不要“试试看”，而要先“证明”。

Bug ：不是 GPU 问题，是地址撞上了 VGA framebuffer

第一步是读源码，把 QEMU bridge、SPDK vfio-user 和 ROCm XIO 的 completion path 逐层跟下来。纸面上看，doorbell offset、phase bit、completion posting 都是对的。

第二步是把 GPU 暂时拿掉，写一个很小的 CPU probe 程序，从 CPU 侧驱动同一个 shadow ring。这个测试不需要 GPU、不需要 passthrough，也不需要重启。它能读回 NVMe controller register：

VERDICT: bridge<->SPDK round-trip WORKS from CPU

这说明 bridge、vfio-user-pci、SPDK 这几层至少在 CPU 路径下是通的。问题范围被砍掉了一半。

接着打开 QEMU bridge trace，结果非常关键：

  0 × pci_mmio_bridge_write

bridge 在 poll 到一堆全零 command。更诡异的是，这种现象甚至发生在 GPU kernel 没有成功运行的 boot 上。也就是说，这不是 GPU 写出来的垃圾。

真正的原因很朴素：shadow-gpa = 0x80000000 和 QEMU standard VGA framebuffer 的 BAR 地址重叠了。

pci 0000:00:01.0: [1234:1111] class 0x030000 BAR 0 [mem 0x80000000-0x80ffffff pref]

bridge 以为自己在 polling shadow ring，实际上读到的是 VGA framebuffer。GPU 写 doorbell 也写进了 framebuffer。

修复方式是把 shadow ring 挪到 4GB 以上、避开 32-bit PCI MMIO window 的空洞中：

+ shadow-gpa=0x100000000

CPU probe 先验证 ring 干净、读写能 round-trip，然后才再消耗一次 GPU boot。这次没有继续把问题归咎于 GPU coherence、MMIO 注册、SMU 这些复杂因素。根因就是一个地址空间重叠。

Bug ：GPU 要 1024-entry queue，controller 只报 256

地址修好后，下一次 GPU shot 终于出现了干净的 doorbell：

pci_mmio_bridge_write BDF=0x0028 BAR=0 offset=0x1040 value=0x1 size=4

0x1040 正是 queue 8 的 SQ-tail doorbell：0x1000 + 2 * 8 * 4。这说明 GPU → shadow ring → bridge → SPDK 的 doorbell 路径已经通了。

但仍然没有 completion。SPDK log 给出了第二个原因：

handle_del_io_q:    *ERROR*: I/O sqid:8 does not exist

前面 CPU probe 读到的 CAP 是：

CAP = 0x80000201e0100ff

低 16 bit 是 MQES（Maximum Queue Entries Supported），0x00ff = 255，表示最大 queue size 是 256。而 ROCm XIO 硬编码使用 1024-entry queue。于是 SPDK 拒绝创建 queue，GPU 后续写到的其实是一个不存在 queue 的 doorbell register。

修复方式是在 SPDK transport 上把 queue depth 提高：

nvmf_create_transport -t VFIOUSER -q 1024 -m 16

同样，先用 CPU probe 验证 CAP 从 ...0100ff 变成 ...0103ff，也就是 MQES 从 255 变成 1023。确认 controller 会接受 1024-entry queue 后，再进入 GPU 测试。

这两个 bug 的共同点很值得记住：都不是那些听起来很高级的问题。一个是 framebuffer overlap，一个是 queue size ceiling。

块设备路径跑通：GPU 从 RADOS 读数据

带着两个修复进入下一轮 boot，第一次 GPU-initiated NVMe read 成功：

Test completed successfully!

然后把后端换成指向 rbd/gpuimg 的 bdev_rbd，Ceph/RADOS 的计数器开始增长：

RD_OPS 65 -> 249     RD 49 KiB -> 2.1 MiB     (during the GPU run)

这说明数据确实来自 Ceph：gfx1151 iGPU 从 __device__ code 发起 NVMe block read，经 QEMU bridge、SPDK、`librbd`，最终访问 RADOS OSD。guest CPU 不在数据路径中。

最初延迟大约 987 us/IO，几乎贴着 bridge 默认 1 ms poll interval。把 poll-interval-ns 从 1000000 降到 10000（10 us）后，结果变成：

mode 0:  Min 105 us · Avg 210 us · Max 367 us

这意味着 bridge 不再是主要瓶颈，真正的存储路径开始显现。

接着启用 --memory-mode 8，让 data buffer 位于 GPU memory：

Iterations: 50   Average: 136.24 us   Test completed successfully!

此时 read buffer 在 GPU device memory 中，ROCm XIO 把它导出为 SPDK 可访问的 guest-physical IOVA，并放进 NVMe read command 的 PRP。SPDK 从 RADOS 取回 payload 后，通过 P2P-DMA 直接写入 GPU memory，不需要 host bounce。

这是 block-first proof：后端是 RBD image，也就是 bdev_rbd`/`librbd。它只是彩排，不是最终目标。

真正的目标：让 GPU 直接访问 KV namespace

block 跑通之后，才进入正题：让同样的 GPU-initiated、CPU-free、payload-into-VRAM 路径，终止在 RADOS object 支撑的 Key-Value namespace 上。

这需要给 ROCm XIO 增加一条 NVMe-KV command path（相关 PR(https://github.com/ROCm/rocm-xio/pull/177)）。原来的 nvme-ep kernel 只会构造 block read：opcode 0x02，带 LBA 和 length。KV Retrieve 也是 64 字节 SQE，但字段语义不同：

key 最多 16 字节，直接嵌在 command 里；
value buffer 仍然通过 PRP/SGL 描述；
Store opcode 是 0x01，Retrieve opcode 是 0x02；
controller 如何解释 opcode，取决于 namespace 的 Command Set Identifier，这里是 Key Value（0x1）。

具体字段上，key 的前 8 字节放在 CDW2/CDW3，后 8 字节放在 CDW14/CDW15；key length 放在 CDW11 [7:0]；请求的 value size 放在 CDW10；value buffer 继续使用 DPTR 里的 PRP1/PRP2。

这也是复用的好处：GPU 侧原本已经知道如何把 block read 的目标 buffer 指向 VRAM dma-buf。KV Retrieve 的 value 落点使用同一套 PRP 机制，所以 `--memory-mode 8` 下的 P2PDMA IOVA 逻辑可以直接复用。新增代码主要是 KV SQE 字段打包。

target 侧变化更小。doorbell → shadow ring → pci-mmio-bridge → vfio-user → SPDK 这条路径不关心 command 是 block read 还是 KV Retrieve；只要后面挂的不是 bdev_rbd，而是 rados-nkv kvdev，Retrieve 就会变成一次 kvdev_rados lookup，再通过 librados 访问一个以 key 命名的 RADOS object。

第一次 GPU 直接发起 KV Retrieve

最终端到端路径跑通：

       rados stat              -> kvpool/6770756b65793031  size 4096

数据链路是：

  -> Ceph/RADOS object

Store 时，GPU 用一个已知 LFSR pattern 填充 value buffer。随后在 Ceph 侧把 object dump 出来，可以看到真实数据，而不是全零占位：

# sha256 = ea1b6431...  (not the all-zeros hash)

真正关键的是 Retrieve。使用 --memory-mode 8 时，GPU 按 key 请求 value，SPDK 从 RADOS object 取回数据，并通过 P2P-DMA 直接写入 GPU memory：

        rocm-xio: Registered buffer ... (P2PDMA attachment kept alive)

这个 140 us/IO 和前面 block 路径的 136 us/IO 基本处在同一量级。用 bridge 默认 1 ms poll 时，KV Retrieve 也会接近 1 ms；把 poll interval 降到 10 us 后，它贴近真实 RADOS read latency。

这就是这次实验最核心的结论：

GPU 从 device code 发起 NVMe-KV Retrieve，value 来自 Ceph/RADOS object，并直接进入 GPU memory。guest CPU 不在数据路径中，也不需要 block-to-object 查找表。

同样的 GPU-initiated path，但后端从 bdev_rbd 换成了 rados-nkv：key 定址，value 来自 RADOS object。

这个组合新在哪里

GPU-initiated storage 本身不是全新概念；NVMe-KV 也不是全新概念。CPU 侧的xNVMe(https://arxiv.org/html/2411.06980v1) 已经在统一不同 NVMe command set，包括 Key-Value。基于 vfio-user 的 KV SSD 模拟也有先例，例如 Simon Lund 的vfio-user-kvssd(https://github.com/safl/vfio-user-kvssd)。

在 AI KV cache 领域，也有一些专有方案已经在做 GPU-initiated key-value request，比如 Pliops XDP LightningAI；NVIDIA 的CMX / DOCA Memos(https://www.nvidia.com/en-us/data-center/ai-storage/cmx/) 则在 BlueField 上用 key-value API 承接 KV-cache context tier。

这次实验的特别之处在于组合方式：

发起方是 AMD GPU；
command 是 GPU-initiated NVMe Key-Value；
target 不是本地 KV SSD，也不是专有 appliance；
后端是分布式 Ceph/RADOS object；
整条路径基于开源软件定义存储，没有定制 DPU 或 ASIC。

这对 AI 推理很重要，因为 KV cache 的天然接口就是 Key-Value。token 序列的 hash 可以直接作为 key。两个 GPU 只要算出相同 prefix，就能访问同一个 RADOS object，不需要先问中心化目录“这个 block 在哪个 offset”。

Wavefront 多 key：一次提交一批 key

单次提交只取一个 key，GPU 利用率并不理想。ROCm XIO 还有一条 wavefront KV path：一个 workgroup 里的线程 1..B 分别打包各自 key 的 SQE，线程 0 统一写一次 SQ-tail doorbell，然后收割 B 个 completion。

例如一次提交四个 key：

    --batch-size 4 --value-size 4096 --data-buffer-size 16384 --pci-mmio-bridge

随后 Ceph 里出现四个 RADOS object，并且 mtime 相同，说明它们来自同一次 batched submission：

shard3 -> kvpool/736861726433  size 4096

这个过程又暴露了第三个 bug。

Bug ：Completion 读取被撕裂

batched Retrieve 时，第一个 key 的 value length 偶尔返回 0，而后面的 key 正常：

KV retrieve OK: key_idx=3 returned value length = 4096 bytes

看 completion path 后，原因很清楚：cqeRead 复制 16 字节 CQE 时，先读低 dword，而 phase bit 在最后一个 dword。poll loop 在单次 cqeRead 上等待 phase 翻转。于是可能读到旧的 cdw0（value length 还是 0），同时又读到新的 phase bit。这就是一次 torn read。

为什么主要影响 key_idx=0？因为第一个 completion 的 polling 和 controller 写 completion 的竞争最激烈；等循环处理到后面的 slot 时，写入已经稳定了。

修复也很直接：观察到 phase flip 后，先做 system fence，再重新读取 CQE，之后才信任 cdw0。

这次修复可以在 guest 中在线编辑、重建、重新运行，不需要重启 host。修复后，single-key 和 batch path 都返回正确的 4096 bytes。

调试经验：复杂系统里，先缩小问题边界

这次实验横跨 GPU、QEMU、vfio-user、SPDK、ROCm、RADOS，看起来很容易掉进“玄学调试”。真正有效的做法反而很朴素：

跨硬件边界做 bisect。CPU-driven shadow-ring probe 是关键，它把 GPU 从问题空间里暂时拿掉，直接验证 bridge 到 SPDK 的路径。
先用便宜测试验证修复，再动昂贵资源。GPA 迁移和 MQES 提升都先用 CPU probe 验证，再消耗 GPU boot。
读源码，不要只猜现象。QEMU、SPDK、ROCm XIO 每一层都要确认真实行为，而不是凭接口文档想象。
打开 instrumentation。一个 QEMU trace 就把“GPU 0x1016 异常”具体化成“338 条全零 command”，并直接排除了 GPU 写坏 ring 的假设。
优先怀疑简单错误。最后的三个 bug 分别是地址重叠、queue depth 不匹配、CQE 读取顺序问题。它们都比“GPU coherence 玄学”普通得多。

对 AI 推理的意义：一次提交覆盖长上下文

ROCm XIO 使用 1024-entry submission queue。wavefront path 每次批量提交最多可以写入 queue_length - 1 个 SQE。结合 KV cache 里常见的 block 聚合粒度，可以粗略算一笔账：

1024 SQEs * 256 tokens per block = 262,144 tokens

也就是说，一次批量提交理论上可以覆盖 26 万 token 级别的 prefix。对 agentic coding、长上下文问答、RAG 后处理、多轮工具调用这类场景，这个数量级已经很有现实意义。

256-token block 不是随手选的。LMCache 和 llm-d-kv-cache 已经会在通过 kv-connector 交给外部系统前，把 vLLM 的 16-token paged block 聚合到这个粒度。换句话说，内容寻址的单位与现有生态正在使用的单位是一致的。

如果以 IBM Granite 4.0 30B 这类模型估算，一个 256-token block 大约是 32 MB。一次批量提交一组 key 后，SPDK 可以把这些请求分发到 async reactor，由 `librados` 异步访问 RADOS，再通过 P2P-DMA 把每个 value 放到 VRAM-backed dma-buf 中。

这不是在做“更快的对象存储 API”那么简单，而是在给 AI 推理准备一个更贴近 workload 语义的存储接口：按内容寻址、按 batch 提交、直接进入 GPU memory。

下一步：从功能原型走向可复现与可评估

现在跑通的是功能性原型，不是正式 benchmark。文中的延迟数据来自单机、warm-cache 场景，主要用来说明 bridge 已经不是瓶颈；真实吞吐、TTFT 收益、冷缓存表现、集群规模影响、hugepage-backed 配置下的结果，还需要后续系统性测试。

后面有几个方向值得继续推进：

接入真实推理框架。通过 NIXL / kv-connector 让 vLLM、LMCache、llm-d-kv-cache 真正使用 rados-nkv。
扩展到模型权重流式加载。rados-nkv-weights 可以让模型权重也走同一套 namespace。
引入受控的 computational storage。通过 gRPC 控制面把 allowlist 里的 RADOS `cls` 映射到 opcode，在数据旁边执行安全、可控的计算。
探索不同部署形态。可以与 hypervisor 共驻，也可以放到 BlueField-4 这类 DPU 上，或者在 storage host 侧结合 RDMA fabric 与 GDS-class NIC。
整理可复现 recipe。把带 rados-nkv 的 SPDK、ROCm XIO、patched QEMU、guest plumbing 串成一份可操作文档，降低复现实验门槛。

结语

定期体检就像给身体定期做的 “安全大检查”。它能帮助我们早期发现癌症的蛛丝马迹。尤其是吸烟人群和有肺癌家族史的人，千万别忽视定期体检的重要性。

这篇文章真正想说明的，不只是“某个 demo 跑通了”。更重要的是，AI 推理的存储路径正在从传统 block/file/object API 走向更贴近模型运行时语义的形态。

KV cache 的 key 是内容 hash，value 是可复用的中间状态；GPU 是真正消费这些数据的地方；Ceph/RADOS 提供分布式、可扩展、可复制、可纠删的对象底座。把这三者用 NVMe-KV 和 GPU-initiated I/O 串起来，就有机会把长上下文推理里的重复计算和数据搬运成本降下来。

所以，NVMe doorbell 在这里不只是一个寄存器细节，它指向的是 AI 存储路径的一种变化：