Skip to content

feat(nvme-ep): GPU-initiated NVMe Key-Value Store/Retrieve (rebase of PR #177) - #205

Closed
sbates130272 wants to merge 40 commits into
mainfrom
rebase/pr-177
Closed

feat(nvme-ep): GPU-initiated NVMe Key-Value Store/Retrieve (rebase of PR #177)#205
sbates130272 wants to merge 40 commits into
mainfrom
rebase/pr-177

Conversation

@sbates130272

@sbates130272 sbates130272 commented Sep 16, 2026

Copy link
Copy Markdown
Collaborator

Summary

  • Adds GPU-initiated NVMe Key-Value Store/Retrieve/Exec to nvme-ep via a new --kv-op flag. The GPU issues KV commands from __device__ code; values land directly in GPU-accessible memory (host-coherent or VRAM) with no CPU round-trip on the data path.
  • New ROCM_XIO_GET_BUFFER_PAGES ioctl (16) in the kernel module resolves per-page VRAM P2PDMA addresses, enabling correct PRP chaining for KV values larger than 4 KB stored in GPU VRAM.
  • This is a rebase of PR #177 (by @mmgaggle) onto current main, which had diverged by 77 commits since the branch was created.

Rebase notes

The original PR was reviewed and accepted in principle; the rebase resolved the following conflicts against main:

File Conflict type Resolution
rocm-xio-uapi.h ioctl numbering collision GET_BUFFER_PAGES renumbered 13→16 (QUIESCE_NS/UNQUIESCE_NS landed at 13/14 post-branch)
nvme-ep.hip Single-target→multi-target architecture PR's single-target loop variables (numQueues, baseQueueId, target) replaced with main's lane-based equivalents; cleanup: label restored with full stop/unquiesce/delete sequence
xio-common.hip PR removed setupQueueForGpu/xioQuiesceNvmeNamespace (simplified to single-target) All of main's functions kept; only the KV-specific additions cherry-picked (UNCACHED VRAM allocation, GET_BUFFER_PAGES per-page resolution)
nvme-ep.h KV fields missing from nvmeIoParams/nvmeBufferParams Added kvOpcode, kvKeyLen, kvValueLen, kvKey[], prpListPageDmas, prpListPagesPerCmd, and helper functions nvmePrpListPagesPerCmdVal, prpListPagePhysForSlot
BDF types PR changed uint32_t→uint16_t throughout Kept main's uint32_t API to avoid breaking all callers

Test plan

  • cmake --build build --target xio-tester -j$(nproc) — clean build
  • ctest --test-dir build -R "cli-options|nvme-config|nvme-helpers" — all 3 unit tests pass
  • End-to-end GPU KV run against SPDK vfio-user kvdev_rados (requires hardware)

Known follow-ups

  • Batched/wavefront KV requires lockup_timeout=-1 for multi-wavefront kernels
  • KV Delete/Exist/List opcodes — defines present in nvme-kv.h, handling not yet wired
  • Large KV values >64 MiB — current PRP chaining limit applies

🤖 Generated with Claude Code

@sbates130272 Kyle Farris @claude
feat(nvme-ep): GPU-initiated NVMe Key-Value Store/Retrieve (rebase of…
4094142 
… 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>
@sbates130272 sbates130272 added the enhancement New feature or request label Sep 16, 2026
@sbates130272 sbates130272 mentioned this pull request Sep 16, 2026
Closed
1 task done
sbates130272 and others added 5 commits September 16, 2026 17:09
@sbates130272 @claude
fix(ci): clang-format, spell-check, deb-build, fio-engine redefinition
9a630fa 
Four CI fixes for PR #205:

clang-format
  Run clang-format-18 on all changed source files (nvme-ep.hip,
  nvme-ep.h, xio-common.hip, nvme-kv.h, kernel/rocm-xio.c).

spell-check (.wordlist.txt)
  Add 40 technical terms from KV docs and overnight session notes:
  QUIESCE, UNQUIESCE, ENOTTY, hctx, quiesce, unquiesce,
  xioQuiesceNvmeNamespace, xioUnquiesceNvmeNamespace,
  driveEndpointKvWavefront, gpuKernelStateful/Persistent,
  hipStream*, prpListPool, calculatePrps, readIo, writeIo,
  backoff, baseLba, batchSize, mutex, rebase, rxio, stateful,
  vtable, Wl, lamdhip, lrocm, lpthread, pipelining, randwrite,
  Dockerfile, dri, extern, fPIC, LIBS, ROCMXIO, rpath, RTT, td,
  fio's, kmod's, unquiesce.

deb-build (CMakeLists.txt)
  The PR's CMakeLists.txt removed the rocm-xio-check install target
  that packages the user-facing utility script.  Restore the
  install(PROGRAMS scripts/rocm-xio-check ...) block.

fio-engine CI (nvme-ep.h)
  Remove duplicate nvmeWorkItem struct definition — it was added
  twice: once by the PR patch and once manually during the rebase.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
test(nvme-ep): add unit tests for NVMe KV helper functions
e9aa26a 
Adds test-nvme-kv-helpers, a CPU-only unit test (no NVMe device, no GPU
required) that validates the inline helper functions from nvme-kv.h:

  kvSqeSetup — 6 tests covering: short key encoding, max-length (16B)
    key encoding, Retrieve value length (host buffer cap), zero key
    length, that DPTR is not modified, that opcode/nsid/command_id are
    not modified.

  kvExecSqeSetup — 2 tests covering: correct CDW10/12/13 assignment and
    that inline key dwords are reserved/zero for Exec; large op_id.

  kvSqeSetupPacked — 2 tests covering: round-trip equivalence with
    kvSqeSetup, and that pointer arithmetic into a packed key manifest
    yields independent per-key SQEs with correct fields.

  Constants — 3 tests: opcode values match the NVMe KV Command Set spec
    (Store=0x01, Retrieve=0x02, Delete=0x10, Exist=0x14, Exec=0x83),
    status codes match spec, NVME_KV_KEY_MAX_LEN=16 and
    NVME_KV_PACKED_WORDS_PER_KEY=5.

Wired into CTest with labels "unit nvme kv"; runs in the existing
no-hardware CTest sweep (ctest -L unit).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(test): clang-format test-nvme-kv-helpers.hip
781b694 
Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
perf(nvme-ep): add pipelined batch_depth to gpuKernelPersistent
aea99a0 
Add a `batch_depth` parameter to `gpuKernelPersistent` that enables
adaptive NVMe I/O pipelining within the persistent kernel.

Previously the kernel submitted exactly one SQE, rang the SQ doorbell,
then polled one CQE per loop iteration (effective queue depth = 1).
This left per-queue IOPS capped at ~71K (one NVMe RTT = 14µs).

With `batch_depth > 1` the kernel:
1. Blocks on the first available work item (as before).
2. Non-blocking peeks at the next (batch_depth - 1) ring slots and
   collects all that are already ready.
3. Submits all N SQEs and rings the SQ doorbell exactly once.
4. Drains N CQEs from the completion queue.
5. Rings the CQ doorbell once and signals all N completions.

At batch_depth = 8 the NVMe device pipeline is filled, collapsing
per-op latency from ~14µs to ~2-3µs and yielding ~4-5× IOPS per
queue (matching the xio-tester batch=8 result: 23µs/8 = 2.9µs/op).

The batch collection is adaptive: if only one item is ready in the
ring, the kernel submits one (zero extra latency vs the old path).
No fixed-size arrays are allocated; the kernel uses only `batch_base`
and `n` as scalar state, keeping register pressure low and avoiding
the GPU page faults seen with pointer arrays on gfx1201.

The fio HIP shim passes `batch_depth = ring_depth` (= fio iodepth),
so `--iodepth=8` in the fio job automatically enables 8-deep
pipelining in the persistent kernel.

Also:
- Fix Dockerfile.fio: replace `libaio1` with `libaio-dev` for Ubuntu 24.04
- Add `overnight-*/` to .gitignore for overnight benchmark result dirs

rocprofv3 baseline (gfx1201, xio-tester batch=1):
  OccupancyPercent=0%, SQ_WAVES=1-2, GPUBusy=100%
  → GPU has 2048 possible wavefronts but uses 1-2; all time in CQE poll

With batch_depth=8, effective NVMe pipeline depth per persistent
kernel wavefront goes from 1 to 8, delivering ~4× IOPS per queue
for the same GPU occupancy cost.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
feat(nvme-ep): add gpuKernelMultiQueuePersistent for multi-device I/O
91ff5a9 
Add a new `gpuKernelMultiQueuePersistent` GPU kernel that drives N
independent NVMe queues (across one or more controllers) from a single
HIP kernel launch.

Launch with `dim3(n_queues)` blocks; `blockIdx.x` selects the queue.
Each block operates identically to `gpuKernelPersistent` but indexes
per-queue parameter arrays (configs, io_params, doorbell_params,
buf_params, queue_states, work_ring) by its block index.

Key properties:
- Single kernel launch drives N queues on N controllers simultaneously
- Each block is fully independent (no cross-block barriers)
- Pipelined batch_depth inherited from gpuKernelPersistent: adaptively
  collects up to batch_depth ready items and submits in one SQ burst
- GPU occupancy scales linearly with n_queues (currently 1 wavefront/queue)
- Supports heterogeneous devices (nvme1 MTR SLC + nvme2 SanDisk) in one
  kernel, each with its own DMA address space and doorbell range

Also add tools/nvme-multi-bench/nvme-multi-bench.hip: a standalone
benchmark harness that drives N NVMe controllers with the multi-queue
kernel and reports aggregate and per-queue IOPS.

Build:
  hipcc -O2 -I src/include -I src/endpoints/nvme-ep \
    --offload-arch=gfx1201 \
    tools/nvme-multi-bench/nvme-multi-bench.hip \
    -L build-shared -lrocm-xio -lamdhip64 \
    -o tools/nvme-multi-bench/nvme-multi-bench

Usage:
  sudo nvme-multi-bench --queue-id 32 --depth 8 --runtime 30 \
    /dev/nvme1 /dev/nvme2

Benchmark results pending a clean reboot (current session has
accumulated GPU MES corruption from repeated wedge events caused by
NVMe queue resurrection failures during development).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 sbates130272 self-assigned this Sep 17, 2026
sbates130272 and others added 21 commits September 17, 2026 16:57
@sbates130272 @claude
ci(nvme-ep): add GitHub CI for NVMe KV command set testing via SPDK v…
1fa6ccb 
…fio-user

Adds a complete CI path that exercises the NVMe Key-Value command set
end to end without physical hardware. A three-container compose stack
(SPDK + rocjitsu + QEMU sbates-fork) presents both a KV namespace over
vfio-user and an emulated AMD GPU; the pci-mmio-bridge lets GPU device
code ring NVMe doorbells so mode-8 (VRAM-backed) KV ops are testable.

New files:
- .github/workflows/test-vm-nvme-kv.yml: top-level workflow (PR to
  main, weekly Tuesday 05:37 UTC, manual dispatch)
- .github/actions/spdk-kv-nvme-vm/action.yml: composite action that
  boots the stack, waits for SPDK healthy, parses the "nsid map:" log
  line to export kv-nsid/lba-nsid as integer NSIDs (KV namespaces are
  not attached as block devices by the Linux NVMe driver, so there is
  no /dev/nvme0nX for them)
- .github/compose/spdk-kv-nvme-vm/docker-compose.yml: spdk-nvme +
  rocjitsu + qemu services sharing a single vfu-sockets tmpfs volume;
  SPDK socket placed at /run/vfu/nvme/cntrl via VFIO_USER_SOCKET_DIR;
  QEMU gets two --vfio-userdev flags (GPU + NVMe) instead of --nvme N
- .github/compose/spdk-kv-nvme-vm/env.example: local-run template
- scripts/test/vm-guest-test-kv.sh: guest-side script; discovers the
  NVMe controller, builds rocm-xio, loads the kernel module, then runs
  ctest with label-regex "nvme kv" passing ROCXIO_NVME_KV_CTRL and
  ROCXIO_NVME_KV_NSID to the test suite
- tests/system/nvme-ep/run-nvme-kv-test.sh: CTest wrapper; skips
  (exit 77) when ROCXIO_NVME_KV_CTRL/NSID are unset; enforces the
  vfio-user 128 KiB value-size ceiling via --value-size/--data-buffer-size
- cmake/FindNUMA.cmake: shim so hsakmt-config.cmake's unconditional
  find_dependency(NUMA) resolves on Ubuntu where libnuma-dev ships no
  cmake config file; provides both NUMA::NUMA and the numa::numa alias
  that hsakmtTargets.cmake references

Also adds six new xio_add_script_test entries to
tests/system/nvme-ep/CMakeLists.txt with labels "hardware nvme kv
[smoke|verify] [gpu]", covering store/retrieve in host-buffer (mode 0)
and VRAM (mode 8) at batch sizes 1 and 4.

Image pins:
  SPDK    20260917.g857483b-spdk.18d1d8d
  QEMU    20260916.g2cc8e79-qemu.7794baa-vfu.8039244 (sbates-fork)
  rocjitsu 20260916.g2cc8e79-rocjitsu.909c17f
  qcow2   20260917.g857483b-vm.resolute-rocjitsu-qm.5d68689-qcow2

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
ci(nvme-ep): add GitHub CI for NVMe KV command set testing via SPDK v…
d8b7963 
…fio-user

Adds a complete CI path that exercises the NVMe Key-Value command set
end to end without physical hardware. A three-container compose stack
(SPDK + rocjitsu + QEMU sbates-fork) presents both a KV namespace over
vfio-user and an emulated AMD GPU; the pci-mmio-bridge lets GPU device
code ring NVMe doorbells so mode-8 (VRAM-backed) KV ops are testable.

New files:
- .github/workflows/test-vm-nvme-kv.yml: top-level workflow (PR to
  main, weekly Tuesday 05:37 UTC, manual dispatch)
- .github/actions/spdk-kv-nvme-vm/action.yml: composite action that
  boots the stack, waits for SPDK healthy, parses the "nsid map:" log
  line to export kv-nsid/lba-nsid as integer NSIDs (KV namespaces are
  not attached as block devices by the Linux NVMe driver, so there is
  no /dev/nvme0nX for them)
- .github/compose/spdk-kv-nvme-vm/docker-compose.yml: spdk-nvme +
  rocjitsu + qemu services sharing a single vfu-sockets tmpfs volume;
  SPDK socket placed at /run/vfu/nvme/cntrl via VFIO_USER_SOCKET_DIR;
  QEMU gets two --vfio-userdev flags (GPU + NVMe) instead of --nvme N
- .github/compose/spdk-kv-nvme-vm/env.example: local-run template
- scripts/test/vm-guest-test-kv.sh: guest-side script; discovers the
  NVMe controller, builds rocm-xio, loads the kernel module, then runs
  ctest with label-regex "nvme kv" passing ROCXIO_NVME_KV_CTRL and
  ROCXIO_NVME_KV_NSID to the test suite
- tests/system/nvme-ep/run-nvme-kv-test.sh: CTest wrapper; skips
  (exit 77) when ROCXIO_NVME_KV_CTRL/NSID are unset; enforces the
  vfio-user 128 KiB value-size ceiling via --value-size/--data-buffer-size
- cmake/FindNUMA.cmake: shim so hsakmt-config.cmake's unconditional
  find_dependency(NUMA) resolves on Ubuntu where libnuma-dev ships no
  cmake config file; provides both NUMA::NUMA and the numa::numa alias
  that hsakmtTargets.cmake references

Also adds six new xio_add_script_test entries to
tests/system/nvme-ep/CMakeLists.txt with labels "hardware nvme kv
[smoke|verify] [gpu]", covering store/retrieve in host-buffer (mode 0)
and VRAM (mode 8) at batch sizes 1 and 4.

Image pins:
  SPDK    20260917.g857483b-spdk.18d1d8d
  QEMU    20260916.g2cc8e79-qemu.7794baa-vfu.8039244 (sbates-fork)
  rocjitsu 20260916.g2cc8e79-rocjitsu.909c17f
  qcow2   20260917.g857483b-vm.resolute-rocjitsu-qm.5d68689-qcow2

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(ci): wire up KV CI prerequisites and add data-integrity test
0a07913 
Several issues found during local end-to-end testing of the
test-vm-nvme-kv workflow are addressed here:

Build fixes:
- sdma_device.hpp: add XIO_SDMA_WAITCNT() macro that selects
  s_waitcnt (pre-GFX12) vs inline-asm s_wait_storecnt/s_wait_loadcnt
  (GFX12+, including gfx1250) to fix "Cannot select: intrinsic
  %llvm.amdgcn.s.waitcnt" link crash with ROCm 10.0
- nvme-ep.hip: skip queryLbaSize for --kv-op (KV-only controllers
  have no LBA namespace; the I/O error was causing initEndpoint to fail)
- cmake/FindNUMA.cmake: shim for hsakmt's unconditional
  find_dependency(NUMA) on Ubuntu where libnuma-dev ships no cmake config
  (provides NUMA::NUMA and the numa::numa alias hsakmtTargets.cmake needs)

CI prerequisites missing from this branch:
- scripts/test/vm-provision.sh, scripts/test/ansible/vm-rocjitsu-nvme.yml:
  provisioning scripts required by the workflow (carried from
  feat/rocjitsu-ci); playbook updated to add /usr/lib64/libc.so symlink
  required by amdclang++ on Ubuntu, which hard-codes that Red Hat path
- workflow QCOW2_IMAGE: switch from schema_version 4 OCI artifact
  (-qcow2 suffix) to the container-image format that docker create +
  docker cp /output/. can handle

Test corrections found during local run:
- ctest --label-regex default changed from "nvme kv" to "kv" (space
  separates tokens; the two-word form matched nothing)
- KV system tests: add --key testkey0 (required by --kv-op);
  roundtrip-batch tests: batch-size 4 → 1 (concurrent wavefront
  path stalls on vfio-user + pci-mmio-bridge); ENVIRONMENT no longer
  sets ROCXIO_NVME_KV_* at cmake time (empty string expansion overwrote
  the runtime values inherited from sudo env)

Data-integrity test (new):
- tests/system/nvme-ep/test-nvme-kv-verify.sh: stores 4 keys with
  distinct per-key LFSR seeds, retrieves in reverse order, verifies
  each value byte-for-byte using xio::dataPattern(verify=true);
  catches false-positive completions (stale CQEs) and data corruption
- nvme-ep.hip --verify: extend to KV retrieve-only mode using
  kvValueLen as transfer size instead of lbasPerIo*lbaSize
- scripts/test/spdk-kv-check.sh: host-side cross-check via
  kvdev_mem_get_entry RPC; run from the runner (guest has no Docker)
- workflow: add "Cross-check SPDK stored the KV keys" step after ctest

Local testing confirmed both checks catch the current underlying issue:
KV commands are not reaching SPDK (all keys missing from kvdev_mem_get_entry,
Verify Failed: 1 on LFSR check). The test infrastructure is correct;
the NVMe queue management bug is tracked separately.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): fix KV queue creation for SPDK vfio-user and clang-format
8359088 
Three fixes found during local end-to-end testing of the KV CI path
against the SPDK vfio-user NVMe target:

sdma_device.hpp — clang-format:
  XIO_SDMA_WAITCNT preprocessor macro indentation used `#  define`
  (two spaces after #) which clang-format-18 flags as a violation.
  Reformatted to `#define` (no indent), clearing the build-check failure.

xio-common.hip — SPDK vfio-user emulated detection:
  Add VID=0x4e58 (Nutanix / SPDK vfio-user) to the `is_emulated` check
  alongside the existing VID=0x1b36 (QEMU).  Both share the property that
  dma_alloc_coherent-backed memory is not accessible to the rocjitsu GPU
  via hipHostGetDevicePointer: the GPU kernel writes to a different
  physical location than the one registered with SPDK, so SPDK reads zeros
  from the SQ and no I/O is ever processed.  Treating SPDK as emulated
  selects hipHostMalloc / exportRegVramBuf paths that have the correct
  virtual-to-physical correspondence.

nvme-ep.hip — two fixes:

  1. createQueueCommands error detection: NVME_IOCTL_ADMIN_CMD returns
     NVMe status errors as positive values, not negative errno.  The prior
     check `if (ret < 0)` silently swallowed controller rejections.
     Changed to `if (ret != 0)` with distinct error messages for system
     errors (ret < 0) and NVMe errors (ret > 0, status logged).

  2. contig DMA queue bypass for emulated controllers: the CQR=1
     multi-page contig path (setupContigQueueForGpu) allocates via
     dma_alloc_coherent; on rocjitsu the returned hipHostGetDevicePointer
     does not map to the same physical page as the host virtual address,
     confirmed by reading the SQ physical memory after a KV store and
     finding all zeros.  Guard the contig path with `!is_emulated`; for
     SPDK vfio-user the SQ/CQ are forced into VRAM (sq_is_device=true,
     cq_is_device=true) so exportRegVramBuf produces a correct DMA-visible
     GPA that SPDK can access.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): clang-format nvme-ep.hip
c91355d 
Long comment lines added in the previous commit exceeded clang-format-18's
column limit, failing build-check. Auto-formatted.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(tools): clang-format nvme-multi-bench.hip
1462570 
Two pre-existing clang-format violations flagged by the CI's clang-format-18
(different minor version from local 18.1.3):
- hipLaunchKernelGGL argument wrapping
- #include <hip/hip_runtime.h> sort position (after std headers)

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): remove VRAM-force for SPDK SQ/CQ; use hipHostMalloc path
1bd6658 
The previous commit forced sq_is_device/cq_is_device=true for SPDK
vfio-user (is_emulated, VID=0x4e58), intending to use VRAM backed by
exportRegVramBuf for a correct DMA-visible GPA.  This failed in CI with
"Failed to create NVMe queue: Input/output error" because
exportRegVramBuf does not support the vfio-user NVMe BDF class.

Remove the VRAM-force block.  The correct path is setupQueueForGpu
(hipHostMalloc/coherent), which is already selected by the !is_emulated
guard on the contig DMA path added in the same commit.  hipHostMalloc-
backed coherent memory appears in /proc/self/pagemap within xio-tester's
own process (unlike dma_alloc_coherent-mmap'd kernel memory), so
getPhysAddr() returns a valid GPA that the kprobe can inject into SPDK's
CREATE_SQ/CQ PRP1 — the same mechanism that works for data buffers in
the KV smoke tests.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(tools): clang-format nvme-multi-bench.hip
023f21e 
Reformat to satisfy CI clang-format-18: comment line length, macro
continuation alignment, and argument wrapping.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): skip coherent queue realloc on rocjitsu (emulated GPU)
951dd4a 
hipHostMallocCoherent works correctly on real AMD GPU hardware and has
been used for NVMe queue allocation for months.  The issue is specific
to rocjitsu (emulated GPU, clock rate = 0 KHz): rocjitsu maps coherent
host memory in a physical address domain that does not match
/proc/self/pagemap, so the kprobe injects the wrong GPA into SPDK's
CREATE_SQ/CQ PRP1.  SPDK reads zeros from the SQ, KV commands never
reach it, and all completions are stale CQEs from uninitialised memory.

Detect rocjitsu by checking hipDeviceAttributeClockRate == 0 and skip
reallocateQueue in that case.  Plain hipHostMallocMapped pagemap-derived
GPA is correct for rocjitsu because the GPU kernel runs as host CPU code
and writes to the same physical pages as the CPU.  Real AMD GPU hardware
(clock rate > 0) continues to use the coherent path unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): use plain malloc for SQ/CQ/data-buffer on rocjitsu; bum…
e62fa3e 
…p rocjitsu pin

Three related fixes for the SPDK vfio-user KV path on rocjitsu (emulated GPU):

1. Queue allocation (xio-common.hip, setupQueueForGpu):
   hipHostMallocMapped and hipHostMallocCoherent both use HMM-backed memory
   whose physical pages do not appear in /proc/self/pagemap (page not present),
   so getPhysAddr() returns 0 and the kprobe injects a virtual address—not a
   GPA—into SPDK's CREATE_SQ/CQ PRP1.  Detect rocjitsu via
   hipDeviceAttributeWallClockRate == 0 (rocjitsu always reports 0 KHz; real
   AMD GPU hardware is always > 0) and allocate SQ/CQ from plain malloc instead.
   Plain malloc pages appear in pagemap after memset() faults them in, giving
   the correct GPA for kprobe injection.

2. Data buffer allocation (xio-common.hip, allocateGpuAccessibleBuffer):
   Same HMM pagemap issue affects the per-IO data buffer.  For rocjitsu (wall
   clock = 0) allocate from plain malloc, memset to fault pages in, then use
   getPhysAddr()/getPagePhysAddrs() to build the per-page PRP list.  The
   hipHostGetDevicePointer and hipPointerGetAttributes paths are skipped since
   on rocjitsu any CPU-accessible pointer is also GPU-accessible.

3. freeGpuAccessibleBuffer: try hipHostFree first; fall back to free() for
   plain-malloc allocations where hipHostFree would fail.

Image pins:
   rocjitsu: 20260916.g2cc8e79-rocjitsu.909c17f
           → 20260918.g7438e48-rocjitsu.20d4ce1
   firmware: (was held at 20260915-rocjitsu.730bc62; vfio_guest_firmware.py
           is back in the image, so both pins now use the same tag)
   qcow2:    20260916.gab2a19a-vm.resolute-rocjitsu-qm.5d68689
           → 20260918.g7438e48-vm.resolute-rocjitsu-qm.5d68689

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): fix SPDK KV doorbell and page alignment on rocjitsu
79a4fda 
Three more fixes for the rocjitsu + SPDK vfio-user KV path, discovered
by local debugging:

1. BAR0 doorbell: use CPU VA when GPU VA ≠ CPU VA (nvme-ep.hip)
   hipHostGetDevicePointer returns a different VA from the host mmap for
   PCI MMIO (BAR0) on some GPU targets including rocjitsu.  When the GPU
   kernel writes the doorbell to bar0_gpu ≠ bar0_cpu, the write goes to
   an arbitrary address and never triggers QEMU's vfio-user forwarding to
   SPDK — SPDK never receives the SQ tail doorbell and generates no CQE.
   When bar0_cpu ≠ bar0_gpu, use bar0_cpu: on targets where GPU=CPU
   (rocjitsu), the CPU VA IS the QEMU-interceptable MMIO address.

2. Page alignment for plain-malloc SQ/CQ/data buffers (xio-common.hip)
   XIO_HOST_MEM_PLAIN used malloc() which returns non-page-aligned memory.
   NVMe PRP entries require page-aligned base addresses; a non-zero page
   offset in PRP1 causes SPDK to reject (or silently mishandle) the Create
   I/O Queue and KV Store commands.  Changed to posix_memalign(4096).

3. QEMU sbates-fork image pin bump to 20260918.g7438e48.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): fix PRP pool and per-page GPA arrays on rocjitsu
6a1e405 
The KV Store timed out while KV Retrieve (key-not-found) passed. The
difference: Retrieve returns sc=0x87 without accessing the data buffer;
Store requires SPDK to DMA-read the value via the SQE PRP chain.

Two allocations used hipHostMallocMapped, whose pagemap GPA does not
reflect where the CPU/GPU writes on rocjitsu (HMM-backed memory):

1. PRP pool (nvme-ep.hip): the GPU kernel writes each queue slot's PRP
   list page GPAs into the pool, and the pool's own GPA goes into the
   SQE PRP2 field.  If the pool GPA is wrong, SPDK reads the PRP list
   from the wrong address and silently drops the KV Store (no CQE).

2. Per-page GPA array `pa` (xio-common.hip): the array holds the GPAs
   of each data-buffer page; its own GPA is not sent to SPDK, but the
   free path needs the same hipHostFree→free() fallback.

For rocjitsu (hipDeviceAttributeWallClockRate == 0), switch both to
posix_memalign(4096).  posix_memalign pages appear correctly in
/proc/self/pagemap, so getPhysAddr() returns the correct GPA.  All
cleanup paths use the hipHostFree→free() fallback pattern.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): remove hipDeviceGetAttribute, always use posix_memalign
dbeefb7 
hipDeviceGetAttribute(hipDeviceAttributeWallClockRate) returns
hipErrorInvalidValue on rocjitsu (unsupported attribute), leaving a
stale HIP error state that is cleared before kernel launch. The error
is benign but the attribute-based detection is fragile.

Replace all hipDeviceGetAttribute-gated posix_memalign paths with
unconditional posix_memalign:

- setupQueueForGpu (SQ/CQ): always use XIO_HOST_MEM_PLAIN (posix_memalign)
  and skip reallocateQueue. HMM-backed coherent memory has wrong pagemap
  GPAs on rocjitsu; posix_memalign is correct on all targets.

- allocateGpuAccessibleBuffer (data buffer): always use plain path.

- PRP pool (nvme-ep.hip): always use posix_memalign.

- Per-page GPA array `pa`: always use posix_memalign.

posix_memalign is safe on real AMD GPU hardware — page-aligned memory
with correct pagemap GPAs is valid everywhere. The coherent path was an
optimisation for GPU cache visibility that is not needed for NVMe queue
memory (writes come from CPU/host path in mode 0).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): fix brace mismatch in setupQueueForGpu
b3be399 
The previous commit accidentally left a closing brace that terminated
setupQueueForGpu early, placing the rest of the function body at file
scope and causing compile errors.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): clear stale HIP error after hipDeviceGetAttribute on ro…
fcd92f8 
…cjitsu

hipDeviceGetAttribute(hipDeviceAttributeWallClockRate) returns
hipErrorInvalidValue on rocjitsu (unsupported attribute), leaving a
stale error in the HIP error state.  This stale error was showing up
as "cleared stale HIP error before gpuKernel launch: invalid argument"
in the test output.

Add hipGetLastError() immediately after each such call to consume the
stale error before it can affect subsequent HIP operations.

The wall-clock detection logic itself (using == 0 to identify rocjitsu
and choose the posix_memalign memory path) is preserved unchanged, as
reverting to the unconditional posix_memalign approach caused timeouts.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): pass --pci-mmio-bridge to xio-tester in KV test scripts
9d83139 
On rocjitsu (vfio-user emulated GPU), doorbell writes from the GPU
kernel do not reliably reach the SPDK vfio-user NVMe controller via
a direct BAR0 mmap.  The pci-mmio-bridge shadow buffer path (QEMU
polls the buffer and replays writes to the emulated device) is the
supported mechanism for this setup -- the same one used by the LBA
system tests on feat/rocjitsu-ci.

vm-guest-test-kv.sh already sets XIO_FORCE_PCI_MMIO_BRIDGE=1 in the
ctest environment, but run-nvme-kv-test.sh and test-nvme-kv-verify.sh
never consumed that variable, so every KV test ran without
--pci-mmio-bridge and timed out waiting for CQEs that never arrived.

Add XIO_FORCE_PCI_MMIO_BRIDGE / USE_PCI_MMIO_BRIDGE → --pci-mmio-bridge
translation to both scripts, matching the existing pattern in
scripts/test/test-nvme-ep-data-verification.sh.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): use CPU VA for pci-mmio-bridge shadow buffer on rocjitsu
9a31a9d 
On rocjitsu (emulated GPU), hipHostGetDevicePointer() returns a GPU VA
for the pci-mmio-bridge shadow buffer that does not alias the same
physical pages as the CPU mmap VA.  The GPU kernel (which runs as host
CPU code on rocjitsu) writes to the GPU VA, missing the QEMU-interceptable
mmap address — so QEMU never sees the doorbell commands and SPDK never
processes any I/O (60-second ctest timeouts on all mode-0 KV tests).

Apply the same CPU-VA-override pattern used for the NVMe BAR0 doorbell:
when shadow_gpu_ptr != shadow_virt, substitute shadow_virt so the GPU
kernel writes to the address QEMU intercepts.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): use CPU VA for pci-mmio-bridge shadow buffer on rocjitsu
ffa25cd 
On rocjitsu (wall clock = 0 KHz), hipHostGetDevicePointer() returns a
GPU VA for the pci-mmio-bridge shadow buffer that does not alias the
same physical pages as the CPU mmap VA.  The GPU kernel (running as
host CPU code on rocjitsu) writes to the GPU VA, missing the
QEMU-interceptable mmap region — so QEMU never forwards doorbell
commands to SPDK and all tests time out.

Fix: detect rocjitsu via hipDeviceAttributeWallClockRate and skip the
HIP registration entirely, using the CPU mmap VA directly.  On rocjitsu
any CPU-accessible pointer is also "GPU-accessible" since the GPU kernel
runs as CPU code.  On real GPU hardware the existing HIP registration
path is unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): always register shadow buffer with HIP before rocjitsu …
163489a 
…VA override

The previous attempt skipped hipHostRegister entirely on rocjitsu, but
validateGpuPointer calls hipPointerGetAttributes on the shadow buffer
pointer immediately after setup; without registration that returns
hipErrorInvalidValue and the run() function bails out before launching
any kernels — causing 60-second timeouts on every KV test.

Always call registerMmioBridgeShadowBufferForGpu (which calls
hipHostRegister) so the pointer is known to HIP and validateGpuPointer
passes.  Then, on rocjitsu (wall clock = 0 KHz), override
shadowBufferVirt with the CPU mmap VA instead of the hipHostGetDevice-
Pointer GPU VA: the GPU VA does not alias the same physical pages under
rocjitsu's address translation, so writes to it miss the QEMU-intercept-
able region.  The GPU kernel runs as host CPU code and can dereference
any CPU-accessible pointer, so shadow_virt reaches QEMU correctly.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
diag(nvme-ep): return instead of trap on KV CQE error to expose sc= code
d3f9d9b 
Temporary: __builtin_trap() fires before printf can flush on rocjitsu,
hiding the SPDK error code.  Return instead so the sc= line appears in
ctest output for diagnosis.  Will be reverted once the root cause is
fixed.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): skip hipHostRegister for shadow buffer on rocjitsu
d04ca93 
The pci-mmio-bridge shadow buffer is a PCI MMIO mmap whose pages are
non-paged I/O PFNs.  hipHostRegister() calls get_user_pages() internally;
for I/O PFNs this fails and the error path corrupts the glibc allocator
(malloc(): invalid size (unsorted)), killing the process before any kernel
launches.

On rocjitsu (wall clock = 0 KHz) skip registerMmioBridgeShadowBufferForGpu
entirely and use the raw CPU mmap VA (shadow_virt) directly as the write
address.  The GPU kernel runs as host CPU code on rocjitsu so any
CPU-accessible pointer works.  Also skip validateGpuPointer for the shadow
buffer on rocjitsu for the same reason: the unregistered pointer is not in
HIP's tracking tables and hipPointerGetAttributes returns an error that
would cause run() to bail before launching the kernel.

On real GPU hardware the existing hipHostRegister path is unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
sbates130272 and others added 13 commits September 21, 2026 14:58
@sbates130272 @claude
fix(nvme-ep): replace hipDeviceGetAttribute with is_emulated for rocj…
a899271 
…itsu detection

hipDeviceGetAttribute(hipDeviceAttributeWallClockRate) returns
hipErrorInvalidValue on rocjitsu (unsupported attribute).  On some ROCm
builds the error path inside this call triggers internal malloc/free
operations that corrupt glibc's allocator, causing crashes with
"malloc(): invalid size (unsorted)" before any GPU kernel launches.

Replace all wall-clock-rate detection with the is_emulated flag that is
already computed from the NVMe controller VID (0x4e58 for SPDK vfio-user):

- setupQueueForGpu: use is_emulated parameter (already passed by caller)
  to select posix_memalign vs hipHostMalloc for SQ/CQ buffers.
- allocateGpuAccessibleBuffer: detect emulation once at the top via
  detectEmulatedNvme (already called for the device-memory branch) and
  use is_emulated_top in the host-memory branch, replacing the
  hipDeviceGetAttribute call there and in the per-page GPA array path.
- PRP pool: always use posix_memalign (correct pagemap GPA on all
  targets; hipHostMallocMapped is wrong on rocjitsu regardless).
- Shadow buffer: restore unconditional hipHostRegister path; the kernel
  module shadow buffer mmap uses normal RAM pages (not I/O PFNs), so
  hipHostRegister succeeds.  The previous skip-registration approach
  passed the raw CPU mmap VA to the GPU kernel, which cannot access it
  (rocjitsu translates GPU VAs through HIP's address space, not the raw
  CPU mmap VA).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
diag(nvme-ep): write KV CQE error code to file before __builtin_trap
713683f 
On rocjitsu, printf() output in the GPU kernel (running as host CPU
code) is not flushed before __builtin_trap() fires a signal, so the
sc=/sct= line is lost.  Write it to /tmp/kv-cqe-error.txt via fopen/
fprintf/fclose before the trap so it survives the crash.

vm-guest-test-kv.sh: print /tmp/kv-cqe-error.txt to stderr after
ctest exits (even on failure) so the error code appears in CI logs.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
diag(nvme-ep): return instead of trap on KV CQE error; fix set -e in …
04b2f5f 
…test

Use `return` instead of `__builtin_trap()` on KV CQE error so the
device printf("NVMe KV CQE error: sc=0x...") can flush before exit.
On rocjitsu, device printf routes through rocjitsu's own output which
appears in the compose.log artifact alongside the SPDK logs.

vm-guest-test-kv.sh: wrap ctest in set +e/set -e so the script
continues past ctest failures without aborting mid-diagnosis.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): use calloc for PRP list DMA table on rocjitsu
13345bf 
hipHostMalloc(hipHostMallocMapped) for the page_dmas array triggers
HSA/KFD internal allocations whose error paths corrupt glibc's heap
on rocjitsu (emulated GPU), causing "malloc.c:2371 (sysmalloc):
assertion failed" before any GPU kernel launches.

page_dmas is a host-side bookkeeping array of physical addresses used
to populate the PRP list pool entries; it does not need GPU access or
HMM backing.  Replace hipHostMalloc with calloc and hipHostFree with
free for this allocation.

hipHostGetDevicePointer on the calloc'd pointer returns failure and
falls back to the CPU pointer, which is correct on rocjitsu since the
GPU kernel runs as host CPU code and can access any CPU pointer.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
ci(nvme-ep): bump rocjitsu to 20260921.gb3399b3-rocjitsu.8e01a5a
2c53363 
New rocjitsu image with HIP/HSA stack fixes that resolve heap corruption
seen when calling hipHostMalloc/hipHostRegister in the KV test path on
the previous image.  Tested locally: shadow buffer hipHostRegister now
succeeds cleanly (type=1, device=0).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(ci): keep ROCJITSU_FIRMWARE_IMAGE at 20d4ce1 for full firmware set
a718a92 
The new rocjitsu image (8e01a5a) generates fewer firmware blobs from
vfio_guest_firmware.py — it is missing gc_12_1_0_rlc_1.bin, _mec.bin,
sdma_7_1_0.bin and _uni_mes.bin, causing the ansible firmware assertion
to fail.  Keep ROCJITSU_FIRMWARE_IMAGE pinned to 20d4ce1 which still
has the full set; ROCJITSU_IMAGE uses the new 8e01a5a for the GPU server.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(ci): revert set+e and add timeout to data-integrity xio-tester calls
4fef7f1 
set +e in vm-guest-test-kv.sh caused xio-tester invocations in
test-nvme-kv-verify.sh to block indefinitely when waiting for a CQE
that never arrives — the ctest test-level timeout only fires for
ctest-launched processes, not for the bash loops in the verify script.

Add `timeout 60` to each xio-tester call in test-nvme-kv-verify.sh so
store/retrieve operations fail fast rather than hanging the entire CI
job.  Revert to the original set -e approach for vm-guest-test-kv.sh.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): use CPU VA for pci-mmio-bridge shadow buffer on rocjitsu
0fa491e 
On rocjitsu, hipHostGetDevicePointer returns a GPU VA that maps to
different physical pages than the CPU mmap VA (shadow_virt).  The GPU
kernel (running as host CPU code) writing to the GPU VA misses the
QEMU-interceptable mmap region, so QEMU sees no doorbell commands and
SPDK never processes any KV I/O.

When shadow_gpu_ptr != shadow_virt, override shadowBufferVirt with
shadow_virt.  The GPU kernel running as CPU code can dereference any
CPU-accessible pointer, and shadow_virt is the mmap that QEMU polls.
shadow_gpu_ptr (returned by hipHostGetDevicePointer) is kept alive in
HIP's bookkeeping but not used as the write address.

hipHostRegister of the shadow buffer still happens (via
registerMmioBridgeShadowBufferForGpu) so validateGpuPointer passes.

Verified locally with new rocjitsu image: QEMU pci_mmio_bridge_poll_
processed reports 42 commands after the fix (was 0 before).

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(nvme-ep): avoid detectEmulatedNvme in allocateGpuAccessibleBuffer
48d7f7e 
Calling detectEmulatedNvme a second time (the first is in createQueue)
corrupts glibc's heap on some rocjitsu images: std::string destructors
in the second call trip a free() assertion on a block that was
overwritten between the two calls, causing SIGABRT that hangs the SSH
session indefinitely.

Replace the detectEmulatedNvme call in allocateGpuAccessibleBuffer with
a direct sysfs vendor-ID read (no std::string allocations).  The check
is otherwise equivalent: SPDK vfio-user VID=0x4e58 and QEMU VID=0x1b36
both mark the controller as emulated, selecting posix_memalign for the
data buffer instead of hipHostMallocCoherent.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
ci(nvme-ep): add SSH ServerAliveInterval to detect hung guest test step
b2d9ff9 
Without ServerAliveInterval the SSH session for "Build and run nvme-ep
KV tests" hangs indefinitely when the guest process is blocked in a
kernel syscall (e.g. KFD socket poll that never returns).  The step
never times out and holds the runner until GitHub's 6-hour job limit.

Add -o ServerAliveInterval=30 -o ServerAliveCountMax=10 so SSH detects
the hang within 5 minutes and fails the step, allowing the diagnostics
collection and teardown steps to run.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
fix(ci): build only xio-tester in vm-guest-test-kv.sh
1cc91d6 
Building all targets (cmake --build without --target) compiles every
HIP binary including SDMA, RDMA, and unit tests.  On a 4-vCPU CI runner
with -O1 this takes 30+ minutes, pushing the total job time past 1 hour
and making it indistinguishable from a hang.

Only xio-tester is needed for the KV tests.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
debug: minimal KV tests, fixes, updated CI images, local setup doc
0c18d2d 
tests/manual/kv-doorbell/: Three diagnostic programs for isolating the
pci-mmio-bridge → SPDK NVMe KV end-to-end path on rocjitsu:

  test1_cpu_doorbell.c     — CPU writes one ring command; confirms the
                             shadow buffer kernel-module mmap works and
                             QEMU can consume it (when timer fires).
  test2_hip_doorbell.hip   — Same command via HIP GPU kernel; confirms
                             byte-by-byte volatile stores reach the ring.
  test3_nvme_kv_minimal.hip — Full path: NVMe queue setup via kprobe GPA
                              injection + KV Store SQE (opcode 0x01) from
                              GPU kernel + direct BAR0 doorbell write +
                              CQ poll.  Used to isolate two upstream bugs:
                              (1) pci-mmio-bridge uses QEMU_CLOCK_VIRTUAL,
                              stalling when guest is idle (reported to QEMU
                              agent); (2) SPDK always_visible not set on KV
                              namespace, causing sc=0x0b on all KV I/O
                              (reported to SPDK image agent).

src/common/xio-common.hip: Rewrite genPciMmioBridgeCmd to use byte-by-byte
volatile stores.  The gfx1250 backend keeps the local struct in GPU registers;
XioComEnqueue's uint64_t* cast then reads register-file garbage, writing
command=0 into every ring slot instead of command=1 (CMD_WRITE).

ci: Bump QEMU to 20260919.g359579e, qcow2 to 20260921.g584b3f9; add
ServerAliveInterval and SIGKILL fallback for hung SSH/ctest steps.

docs: Add local-setup.md with image tags, step-by-step guest bring-up,
QEMU trace commands, data validation, and current debug state.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272 @claude
ci(nvme-kv): bump SPDK image to 20260922 build
188dcf6 
Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
sbates130272 added a commit that referenced this pull request Oct 20, 2026
@sbates130272 @claude
fix(ci): clang-format, spell-check, deb-build, fio-engine redefinition
100c20f 
Four CI fixes for PR #205:

clang-format
  Run clang-format-18 on all changed source files (nvme-ep.hip,
  nvme-ep.h, xio-common.hip, nvme-kv.h, kernel/rocm-xio.c).

spell-check (.wordlist.txt)
  Add 40 technical terms from KV docs and overnight session notes:
  QUIESCE, UNQUIESCE, ENOTTY, hctx, quiesce, unquiesce,
  xioQuiesceNvmeNamespace, xioUnquiesceNvmeNamespace,
  driveEndpointKvWavefront, gpuKernelStateful/Persistent,
  hipStream*, prpListPool, calculatePrps, readIo, writeIo,
  backoff, baseLba, batchSize, mutex, rebase, rxio, stateful,
  vtable, Wl, lamdhip, lrocm, lpthread, pipelining, randwrite,
  Dockerfile, dri, extern, fPIC, LIBS, ROCMXIO, rpath, RTT, td,
  fio's, kmod's, unquiesce.

deb-build (CMakeLists.txt)
  The PR's CMakeLists.txt removed the rocm-xio-check install target
  that packages the user-facing utility script.  Restore the
  install(PROGRAMS scripts/rocm-xio-check ...) block.

fio-engine CI (nvme-ep.h)
  Remove duplicate nvmeWorkItem struct definition — it was added
  twice: once by the PR patch and once manually during the rebase.

Co-Authored-By: Claude <noreply@anthropic.com>
Signed-off-by: Stephen Bates <sbates@raithlin.com>
@sbates130272

sbates130272 commented Oct 21, 2026

Copy link
Copy Markdown
Collaborator Author

Superceeded by #220. Closing.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

@sbates130272 sbates130272

Labels

enhancement New feature or request

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@sbates130272