Repository navigation
Conversation
Add a GPU-initiated NVMe Key-Value path to the nvme-ep endpoint: the GPU
issues KV Store then Retrieve from __device__ code via a new
`nvme-ep --kv-op` mode, with the value landing directly in GPU memory.
- src/include/nvme-kv.h: KV command/opcode definitions
- src/endpoints/nvme-ep/nvme-ep.{h,hip}: KV Store/Retrieve handling
- src/tester/xio-cli-options.cpp: --kv-op CLI option
- docs/nvme-kv.md: design + usage notes
- examples/: KV CPU probe and GPU-over-RADOS demo harness
|
@mmgaggle thanks for this awesome contribution. Give me a few days to review and get you some feedback! |
|
@john00003 can I ask that you take this branch from @mmgaggle for a spin on your VM setup? And provide feedback on any steps that did not work for you or were not clear? |
There was a problem hiding this comment.
Pull request overview
Adds an NVMe Key-Value (KV) command path (Store/Retrieve) to the existing nvme-ep endpoint so GPU device code can issue KV ops and DMA values directly into GPU-accessible memory, plus CLI/docs/examples to exercise the feature.
Changes:
- Add KV SQE encoding helpers/opcodes/status codes (
nvme-kv.h) and wire KV fields intonvme-epconfigs. - Implement a serial single-thread GPU KV driver path (
driveEndpointKv) and gate block-only queries when KV mode is enabled. - Add
--kv-op/--key/--value-sizeCLI options and provide docs + example scripts/programs for CPU/GPU validation.
Reviewed changes
Copilot reviewed 9 out of 9 changed files in this pull request and generated 6 comments.
Show a summary per file
| File | Description |
|---|---|
| src/tester/xio-cli-options.cpp | Adds CLI flags for enabling KV mode and setting key/value size. |
| src/include/nvme-kv.h | New header defining KV opcodes/status and kvSqeSetup() SQE field encoder. |
| src/endpoints/nvme-ep/nvme-ep.hip | Implements GPU KV path, dispatch, and host-side KV config handling. |
| src/endpoints/nvme-ep/nvme-ep.h | Plumbs KV fields into nvmeIoParams and host config struct; includes KV header. |
| examples/stage2_kv_rados_gpu.sh | End-to-end GPU KV Store/Retrieve demo against SPDK kvdev_rados via vfio-user. |
| examples/probe_kv_cpu.sh | GPU-free guest validation harness using CPU passthrough KV probe. |
| examples/kv_cpu_probe.c | CPU passthrough KV Store/Retrieve probe program to validate SQE encoding. |
| examples/create_rbd.c | Helper example to create/write an RBD image (ancillary tooling). |
| docs/nvme-kv.md | Design/usage notes for KV mode and on-the-wire field mapping. |
Comments suppressed due to low confidence (1)
src/endpoints/nvme-ep/nvme-ep.hip:1573
- In KV mode, the batch buffer sizing check still uses
eff_batch * (lbasPerIo*lbaSize). The KV device path is serial and always uses batch index 0, so this check can both (1) under-estimate the needed bytes if lbasPerIo/lbaSize don't reflect kvValueLen, and (2) over-reject configs by multiplying by eff_batch even though KV doesn't consume per-batch slices. It should size/validate against kvValueLen with an effective batch of 1 in KV mode.
uint32_t eff_batch = ioParams.batchSize;
if (eff_batch == 0) {
eff_batch = ioParams.infiniteMode
? (uint32_t)(queue_size - 1)
: (uint32_t)(ioParams.readIo + ioParams.writeIo);
}
if (eff_batch > (uint32_t)(queue_size - 1)) {
eff_batch = queue_size - 1;
}
uint64_t xfer = (uint64_t)ioParams.lbasPerIo * ioParams.lbaSize;
uint64_t needed = (uint64_t)eff_batch * xfer;
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
|
@mmgaggle can you review the copilot and John's comments and push an update please? |
Wavefront/batched Key-Value: - driveEndpointKvWavefront(): cooperative multi-key Store/Retrieve sibling of the block driveEndpointWavefront(). Threads 1..B each encode one key's SQE into its own value-buffer slot; thread 0 rings the doorbell and polls the B completions. Dispatched when kvOpcode != 0 && batch_size > 1; the serial driveEndpointKv() still drives batch_size 1. - Keys travel in a device array packed NVME_KV_PACKED_WORDS_PER_KEY uint32 per key (16-byte LE image + length); new kvSqeSetupPacked() helper. - --keys CLI manifest, host kvKeys plumbing, KV-aware effBatch/kernelThreads and buffer sizing, validation, device upload + cleanup. - examples/stage3_kv_wavefront_gpu.sh: end-to-end batched Store + Retrieve over SPDK kvdev_rados with a per-key rados stat verify loop. Fixes from adversarial review and Copilot/reviewer comments on PR ROCm#177: - Size the PRP-list pool from kvValueLen in KV mode (was lbasPerIo*lbaSize), so a multi-page value no longer hits a null PRP list on the device. [Copilot] - Validate the KV op direction against the io counts (store requires --write-io, retrieve requires --read-io) so the needed value buffer is always allocated. [Copilot] - Pass the clamped effBatch to the device in KV mode so a single --key with --batch-size > 1 collapses to the serial path instead of ringing the doorbell over a zeroed SQE; also fixes a spurious buffer-sizing rejection. - Gate per-completion KV status prints to the first pass over the manifest to avoid a printf flood in infinite mode. - Install src/include/nvme-kv.h so example-nvme-ep-info builds. [john00003] - Add CLI unit tests for --kv-op/--key/--keys/--value-size. [Copilot] - examples: reject KV keys >16 bytes and null-check aligned_alloc in kv_cpu_probe.c; add an argc/usage guard to create_rbd.c. [Copilot] Build clean; 112/112 runnable ctest pass (hardware tests skipped). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
Pushed Reviewer fixes
Also in this update: wavefront/batched KV (the follow-up noted in
Build is clean; 112/112 runnable ctests pass (hardware-gated tests skipped). End-to-end on real NVMe-KV hardware via |
The KV completion poll loops (driveEndpointKv serial and driveEndpointKvWavefront) spin on a single cqeRead() until the phase bit flips, then read the returned value length from cqe.cdw0. cqeRead() copies the 16-byte CQE low dword first and the phase bit is the last dword, so a single read can capture a stale cdw0 (still 0) together with the freshly-flipped phase bit -- a torn read. It bit the first / most-tightly-raced completion: KV Retrieve reported a value length of 0 for key_idx 0 (and for the single-key serial path), while later wavefront slots returned the correct length. Once the phase flip is observed, __threadfence_system() then re-read the CQE so cdw0 is coherent before it is used. Verified on a gfx1151 guest over the pci-mmio-bridge against an SPDK kvdev_rados target: all four wavefront keys and the serial single-key path now report the true 4096-byte value length. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
- calculatePrps (nvme-ep.h): build a chained PRP list (511 data entries + 1 chain pointer per non-final list page) so the GPU can describe multi-page values beyond a single 512-entry list page. The host PRP pool reserves listPagesPerCmd whole list pages per command and records each page's phys (new prpListPagesPerCmd); per-command cap raised to 128 list pages. - xio-common.hip: heap-allocate the per-page phys scratch table (was a stack[1024] = 4 MiB cap) so mode-0 host buffers resolve up to 64 MiB. - Per-page VRAM P2PDMA (candidate fix for an amdgpu KFD interrupt crash on >16 MiB GPU-initiated transfers): new ioctl ROCM_XIO_GET_BUFFER_PAGES walks the kept-alive dma-buf sg_table and returns per-4KiB-page DMA addresses; the mode-8 VRAM path consumes them and emits a chained PRP describing the true (possibly non-contiguous) layout instead of assuming one contiguous run from a single base address. WIP: chaining + resolver validated e2e to 128 KiB (GPU) and 64 MiB (host-side); GPU-initiated retrieve clean to 16 MiB; per-page-VRAM fix compiles (kernel module syntax-checked vs 6.19) but is untested on the GPU. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
sbates130272
left a comment
There was a problem hiding this comment.
@mmgaggle looks good. Some comments. Also are your examples wired into CTests? The main thing would be if we can auto-pull nvme-kv.h from somewhere rather than maintain a copy. That would be lovely.
Also did you not have a fix for a __threadfence_system issue? Did you want to do that as a simpler and seperate PR?
| @@ -0,0 +1,87 @@ | |||
| #!/usr/bin/env bash | |||
There was a problem hiding this comment.
This might be better off in scripts.
There was a problem hiding this comment.
Done in 59ea56f — moved to scripts/test/probe_kv_cpu.sh (and fixed the kv_cpu_probe.c path it copies, now examples/nvme-kv/).
| * NVMe Definitions for rocm-xio nvme-ep Endpoint | ||
| */ | ||
| #include "nvme-ep-generated.h" | ||
| #include "nvme-kv.h" |
There was a problem hiding this comment.
Can we pull this into nvme-ep-generated? Is there a public location we can curl this in from like we do for the other nvme header files?
There was a problem hiding this comment.
Good idea, and the standard defs are upstreamable. The KV opcodes/structs are the public NVMe Key-Value Command Set spec and are already in mainline spdk/spdk (include/spdk/nvme_spec.h — enum spdk_nvme_kv_opcode, spdk_nvme_kv_ns_data, plus a lib/nvme/nvme_kv.c client) from the NVIDIA KV-client contribution. So two options:
- Keep
nvme-kv.hvendored (current). It's small and spec-stable, and it carries logic a fetch can't give us: thekvSqeSetup()encoder (key packed into CDW2/3/14/15) and our vendor KV Exec opcode (0x83, ADR-0005/0014), which is not upstream in SPDK or the kernel. - Generate the standard defs like our other NVMe headers. A
fetch-spdk-kv-headers.sh+ extract step (à lafetch-nvme-headers.sh/extract-nvme-defines.sh) pullingspdk_nvme_kv_opcode/structs intonvme-ep-generated.h, leaving onlykvSqeSetup()+ the vendor exec bits hand-written. Note the kernel header we already fetch (linux/nvme.h) has no KV uAPI, so this would be SPDK-sourced.
Either is fine by me — which do you prefer? Happy to add the fetch script in this PR if you'd rather generate them.
There was a problem hiding this comment.
@mmgaggle can you implement option 2 and fetch the headers. This avoids any mismatching in the future.
| uint32_t batchSize; /**< SQEs per doorbell; 1=serial, 0=all. */ | ||
| /* KV mode (empty kvOp => normal block mode). Trailing defaulted members so | ||
| * the brace-init in the constructor stays valid. */ | ||
| std::string kvOp = ""; /**< "store", "retrieve", or "" (block). */ |
There was a problem hiding this comment.
Why are these strings?
There was a problem hiding this comment.
These are the host-side CLI config fields: kvOp holds the raw --kv-op value ("store"/"retrieve"/"exec") and kvKey the --key string, mirroring the existing std::string accessPattern ("sequential"/"random") convention. They're validated on the host (IsMember for the op) and then lowered to the numeric device-side fields the GPU path actually uses — kvOpcode (uint8_t), kvKey[4] / kvKeyLen. Happy to switch kvOp to an enum if you'd prefer stronger host-side typing.
| @@ -0,0 +1,118 @@ | |||
| /* Copyright (c) 2026 IBM Corporation | |||
There was a problem hiding this comment.
Is there a way we can pull this from the SPDK or Ceph code base?
There was a problem hiding this comment.
Yes for SPDK — see my detailed note on the nvme-ep.h #include thread. Short version: the standard KV opcodes/structs are in mainline spdk/spdk (enum spdk_nvme_kv_opcode in nvme_spec.h) and could be generated like our other NVMe headers. What stays hand-written either way is the kvSqeSetup() encoder and the vendor KV Exec opcode (0x83), which isn't upstream. Ceph isn't a source for these (it doesn't define the NVMe KV command opcodes). I can keep this vendored or add a fetch-spdk-kv-headers.sh — your call.
There was a problem hiding this comment.
@mmgaggle as noted above let's implement a fetch please.
|
@mmgaggle note that @john00003 is working on a new |
The GPU-initiated KV Store/Exec data path delivered only a partial prefix of bulk payloads (correct up to a size-dependent cliff, then zeros). Root cause: the GPU-side data buffer was allocated with coherence too weak for a reader that observes the doorbell while the producing kernel is still running. - VRAM data buffer (-m DATA_DEVICE / P2PDMA): was plain hipMalloc = coarse-grained device memory, whose stores are only guaranteed system-visible at KERNEL COMPLETION. The GPU-initiated path rings the NVMe doorbell from within the still-running kernel, so the controller's P2PDMA read saw a partial prefix. Switch to UNCACHED device memory (hipExtMallocWithFlags / hipDeviceMallocUncached): writes bypass the GPU cache and are immediately P2P-visible, and unlike a fine-grained HSA region it stays dma-buf/P2P-exportable (hsa export rejects fine-grained with status 4100). Verified: fixes bulk Store to >=4 MiB. A separate ~4 MiB P2P-mapping cap remains for >4 MiB (distinct issue). - Host data buffer (-m 0): COHERENT (fine-grained) host memory; producer uses non-temporal stores + __threadfence_system before the doorbell. - KV Exec SQE path, NT-staged input, CLI options/validation. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
ROCM_XIO_GET_BUFFER_PAGES counted pages by iterating each sg segment of the (rounded-up) VRAM allocation in page-size chunks, so the total could exceed the caller's max_pages = ceil(transfer/4096) and return -E2BIG. Userspace then fell back to contiguous-from-base addressing, correct only for the first physical segment, so bulk VRAM KV Stores zeroed past ~4 MiB (e.g. 5 MiB failed at the 4 MiB cliff). Walk the sg list as one concatenated byte stream, emit only full page_size pages, and stop once count == max_pages, returning exactly the transfer's ceil(size/4096) pages (data lives at the allocation base). Add a straddle guard (-EINVAL) for the case where a logical page would span two physically-discontiguous segments. Userspace unchanged. Validated on the live GPU-passthrough VM: GPU-initiated VRAM KV Store now bit-exact at 3/4/5/7/12/16/20/64 MiB. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Apply reviewer (@sbates130272) feedback on file layout and docs: - Group the KV C examples under examples/nvme-kv/ (create_rbd.c, kv_cpu_probe.c) instead of the flat examples/ root, matching the per-example subdirectory convention. - Move the end-to-end shell harnesses to scripts/test/ (probe_kv_cpu.sh, stage2_kv_rados_gpu.sh, stage3_kv_wavefront_gpu.sh); fix the kv_cpu_probe.c path probe_kv_cpu.sh copies. - Convert docs/nvme-kv.md to reStructuredText (docs/nvme-kv.rst) for consistency with the rest of docs/, and wire it into the Sphinx TOC under Conceptual; update script paths to scripts/test/. - create_rbd.c: document Ceph library/package and version requirements, drop the unused <string.h>, and run clang-format. - kv_cpu_probe.c: drop "in guest" wording — the probe runs anywhere /dev/nvme0 is a KV controller (VM guest today, bare metal/DPU later). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Address PR ROCm#177 review (@sbates130272): make the KV examples part of the build validation. - Add examples/nvme-kv/CMakeLists.txt: a standalone C project that always builds kv_cpu_probe and builds create_rbd only when librados/librbd are found, so it configures cleanly on hosts without Ceph. - Register it compile-only in tests/integration/CMakeLists.txt via xio_add_install_test (running the probe needs a KV NVMe controller and a Ceph cluster, so CI builds but does not run it). - Document the new example in docs/reference/examples.rst. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Address PR ROCm#177 review (@sbates130272): explain why the probe redefines the KV opcodes instead of sharing the header. kv_cpu_probe is a dependency-free standalone gcc build and cannot include src/include/nvme-kv.h (HIP + nvme-ep-generated.h); the two literals are the public NVMe KV spec values and match nvme_kv_cmd_store/retrieve (and SPDK's spdk_nvme_kv_opcode). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Address the six failing PR ROCm#177 checks without touching functional code: - clang-format (v18): reformat the KV-touched sources (xio-cli-options.cpp, nvme-ep.h/.hip, nvme-kv.h, kv_cpu_probe.c, xio-common.hip). - doxygen/sphinx: document calculatePrps() prpListPagePhys and prpListPageCount parameters (fail-on-warning gate). - doc8 (--max-line-length 80): wrap six over-length lines in docs/nvme-kv.rst. - codespell: ignore "hda" (GPU HD-audio PCI function .6), not a typo. - spell-check: add KV/Ceph/SPDK/NVMe-oF terms to .wordlist.txt. - shellcheck (SC2034): rename unused loop counter i to _ in the KV test scripts. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
@mmgaggle @sbates130272 This is great stuff. I would love to explore the potential linking of the unique hash/namespace of a KV cache donor identified by SemBlend (https://github.com/WorldFlowAI/semblend) with a 64 byte NVMe command. Would love to work on a proof of concept here that uses some of the work we are doing with SGLang.. |
|
@mmgaggle i suggest we get this part merged. Just be sure to have some reserved fields in the command set for semantic fields we might want to add. |
|
@zbennett10 awesome. Let's get this PR merged. Perhaps you can open an issue so we can track this. I'd love your thoughts on how to merge SemBlend semantics into a NVMe KV command set! |
|
One approach would to be have a vector database that stores Merkle hashes of a cache block's token prefix alongside an embedding from the prefix + cache block sequence. Then you could do a top-k search with an embedding from a new prompt to find one that is closest match. Then instead of initiating a receive for the real Markel hash, you'd use the close-to hash. If there is a way to extract embeddings from the KV data itself then that would be cool too. You potentially extract the embeddings and build a vector index by chaining store/execute/receive calls together, but it could also just be an adjacent vectordb of some sort. |
|
@sbates130272 opened a tracking issue here: #182 My short take, just thinking through this... I would keep the NVMe-KV layer storage-oriented and put the SemBlend semantics in the metadata/manifest above it, at least for the first pass. |
@mmgaggle I think this is worth exploring from an answer quality perspective - that's where we have to vet it. We have proven that you can do this (different approach, think cuVS CAGRA for vector search in the GPU, or HNSW on CPU process, local to the inference engine) and find good quality matches. Other things to think about:
SemBlend does work to help hand inference engines a plan that helps handle a lot of these considerations |
|
@zbennett10 @mmgaggle can ee move this discussion to the issue for SemBlend so we don't lose it or conflate it with this PR. |
|
@copilot can you fix the CI check failures on this branch? |
|
@mmgaggle I think the final thing is the fetch from SPDK for the nvme-kv header files. Can we also rebase to clean up the total number of commits to just 3-4? I can then give it one last review before merging and then we can add in the docker PR from @john00003! |
Add a GPU-initiated NVMe Key-Value path to the nvme-ep endpoint: the GPU
issues KV Store/Retrieve (serial and wavefront/batched) from __device__
code via a new `nvme-ep --kv-op` mode, with the value landing directly in
GPU (VRAM) or host memory. Supports values up to 64 MiB via chained PRP
lists and per-page VRAM P2PDMA.
- src/include/nvme-kv.h: KV command/opcode definitions
- src/endpoints/nvme-ep/nvme-ep.{h,hip}: serial + wavefront KV Store/
Retrieve, PRP chaining (calculatePrps), buffer coherence handling
- src/common/xio-common.hip: heap-allocated per-page phys scratch table
- src/tester/xio-cli-options.cpp: --kv-op/--key/--keys/--value-size CLI
- kernel/rocm-xio: ROCM_XIO_GET_BUFFER_PAGES ioctl + multi-segment VRAM
page-count fix
- docs/nvme-kv.rst, docs/reference/examples.rst: design + usage notes
- examples/nvme-kv/ (create_rbd.c, kv_cpu_probe.c) + CMake wiring into
install-integration CTests
- scripts/test/: probe_kv_cpu.sh, stage2_kv_rados_gpu.sh,
stage3_kv_wavefront_gpu.sh end-to-end harnesses
Squashed from the NVMe-KV development series; incorporates PR ROCm#177 and
Copilot review feedback (KV opcode literals, example/doc reorg, CTest
wiring, torn-CQE fix, bulk data coherence, PRP-pool sizing, KV op
direction validation, CLI unit tests, example input guards).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
@mmgaggle some more issues my AI agent noticed on this branch: What's Missing / Needs Attention Before Merge
|
|
@mmgaggle can you sign your commits with your pgp key? |
… PR #177) Adds a GPU-initiated NVMe Key-Value Command Set path to nvme-ep. The GPU issues KV Store and Retrieve from __device__ code via a new --kv-op mode; the value transfers directly to/from GPU-accessible memory with no CPU round-trip on the data path. New GPU device functions: driveEndpointKv() — single-thread serial KV Store/Retrieve/Exec driveEndpointKvWavefront() — cooperative batched KV (--batch-size >1) New host additions: src/include/nvme-kv.h — KV opcode/status defines and SQE helpers kernel/rocm-xio/rocm-xio.c — ROCM_XIO_GET_BUFFER_PAGES ioctl (16) for resolving per-page VRAM P2PDMA addresses; needed for large KV values (>4 KB) stored in GPU VRAM src/common/xio-common.hip — UNCACHED VRAM allocation for KV data buffers (fixes GPU→NVMe P2P data visibility); per-page VRAM address resolution via GET_BUFFER_PAGES Rebase notes: - Rebased 10 commits from mmgaggle/nvme-kv (branched 77 commits before current main) onto main. - ROCM_XIO_GET_BUFFER_PAGES renumbered ioctl 13→16 to avoid collision with QUIESCE_NS (13) and UNQUIESCE_NS (14) added in main. - BDF type kept as uint32_t throughout (PR used uint16_t; keeping main's API to avoid breaking callers). - All of main's multi-target nvme-ep architecture preserved (NvmeTarget struct, targets vector, lanes-based dispatch). - KV mode skips namespace capacity query and LBA detection (KV namespaces have no block format). - xioQuiesceNvmeNamespace/xioUnquiesceNvmeNamespace retained from main (PR pre-dated their addition). Unit tests: test-nvme-config, test-nvme-helpers, test-xio-cli-options all pass. Known follow-ups (from PR description): - Batched/wavefront KV (current path is serial single-thread for iodepth=1; cooperative path exists but needs lockup_timeout=-1) - KV Delete/Exist/List opcodes (defines present; handling not wired) - Large KV (>64 MiB) — current PRP chaining limits apply Co-Authored-By: Kyle Farris <kfarris@amd.com> Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: Stephen Bates <sbates@raithlin.com>
Add a GPU-initiated NVMe Key-Value path to the nvme-ep endpoint: the GPU issues KV Store then Retrieve from device code via a new nvme-ep --kv-op mode, with the value landing directly in GPU memory.
Motivation
nvme-ep today drives block I/O from the GPU. This adds the NVMe Key-Value Command Set (Store/Retrieve) so device code can put/get values keyed by name, with the value transferred straight into GPU-accessible memory (host or VRAM/P2PDMA) — no CPU round-trip on the data path. The motivating workload is GPU-initiated fetch of key-addressed blobs (e.g. weight shards) directly into VRAM. The feature is additive: it's a flag on the existing endpoint, and block behavior is unchanged when --kv-op is not set.
Technical Details
Test Plan
Test Result
Submission Checklist