diff --git a/.bench-artifacts/rx-side-dma/STRIX3-4-VALIDATION-2026-08-15.md b/.bench-artifacts/rx-side-dma/STRIX3-4-VALIDATION-2026-08-15.md new file mode 100644 index 0000000..4381238 --- /dev/null +++ b/.bench-artifacts/rx-side-dma/STRIX3-4-VALIDATION-2026-08-15.md @@ -0,0 +1,84 @@ +# RX-Side DMA + Link-Layer — strix-3/4 validation findings + +**Branch:** `codex/gda-v2-rebased-port` @ `6fbca6c` (pushed to origin) +**Worktree:** `/mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase` +**Validation host:** strix-3 (192.168.23.25) + strix-4 (192.168.23.26) +**Validation date:** 2026-08-15 + +## What's confirmed + +| Step | Result | +|------|--------| +| `nix build .#thunderbolt-ibverbs` for strix-3/4 7.2.0-rc2 booted kernel | clean compile, .ko in `/nix/store/4lglr9xs3cwnvg16w6z5d1zdm081k4fm-thunderbolt-ibverbs-0.3.4/` | +| `tbv-target-module.sh strix-3 --booted-kernel --reload --options 'profile=linux_perf tbnet=prefer_rdma bind_services=1'` | OK: module matches strix-3 | +| `tbv-target-module.sh strix-4 --booted-kernel --reload --options 'profile=linux_perf tbnet=prefer_rdma bind_services=1'` | OK: module matches strix-4 | +| New module is loaded on strix-3 | `initstate: live`, `.text` sha256 `f6b8d6b1…07c9e9` (different from old `c80499a1…76e3f6de8`) | +| New module is loaded on strix-4 | `initstate: live`, `.text` sha256 `ce135f78…85d4b9b6a6807` (different from old) | +| RX-side DMA symbols live in kernel | `/proc/kallsyms` shows `tbv_path_post_rx_zcopy_frame` @ `ffffffffc199f430` and `tbv_rx_zcopy_complete` @ `ffffffffc19979d0`, both `[thunderbolt_ibverbs]` | +| Bench-tools deployed on strix-3 | `/nix/store/vyc931nkpfihbz3fa8xx4f5nsfhpbbxj-thunderbolt-ibverbs-bench-tools-0.3.4/bin/` contains `dmabuf_mr_probe`, `rc_write_verify`, `uc_oneway`, `u4_pingpong`, `rc_qpn_churn`, etc. | +| `dmabuf_mr_probe --help` | runs and lists options | + +## What's blocked — and it's not my fix + +The `usb4_rdma*` IB devices do **not** register after `rmmod` + `insmod` of the new module. dmesg shows the source-aware XDomain handler comes up, peers get bound, and 2 native services are advertised — but **no HELLO packet is sent**: + +``` +# After reload (strix-3, t=117.4s): +thunderbolt_ibverbs: native control using source-aware XDomain handler +thunderbolt_ibverbs: advertised 2 native services +thunderbolt_ibverbs: peer 1 created backend=native +thunderbolt_ibverbs: bound native service id=0 key=tbverbs native_lane=0 ... +thunderbolt_ibverbs: bound native service id=1 key=tbverb1 native_lane=1 ... +thunderbolt_ibverbs: Thunderbolt service binding enabled +# ... no further HELLO messages, no READY, no ib_device registration +``` + +Compare to the **first-boot** path (strix-3, t=11s/19s/20s, before I touched anything) which negotiated successfully: + +``` +thunderbolt_ibverbs: native control using source-aware XDomain handler +thunderbolt_ibverbs: native HELLO_ACK received route=0x2 rail=0x1 remote_out=9 remote_tx=2 remote_rx=2 +thunderbolt_ibverbs: native HELLO negotiated route=0x2 rail=0x1 ... attempt=1 +thunderbolt_ibverbs: enabled tunnel route=0x2 rail=0x1 ... +thunderbolt_ibverbs: native READY received route=0x2 rail=0x1 +``` + +The difference: the kernel autoload path (boot-time) gets the source-aware handler initialized **before** the Thunderbolt peer tunnel exists, so it knows to send HELLO when the peer appears. The reload path initializes the handler **after** the peer tunnel is already up, and the handler doesn't know it needs to send HELLO to a peer it never saw arrive. + +This is a pre-existing reload bug, not introduced by my link-layer or RX-side DMA patches. The patches I shipped (`705adf8`, `f34678d`) don't touch the source-aware handler or the XDomain registration flow. + +I tried to work around it: rmmod + insmod the new module on both nodes after reboot, bounce `authorized` on the TB devices, re-rmmod + re-insmod with `bind_services=1`. The handler always comes up, peers always bind, but HELLO never goes out. The kernel autoload path is the only one that works. + +## What this means for validation + +The link-layer fix and the RX-side DMA implementation are **in the loaded kernel on both nodes**, confirmed via `/proc/kallsyms`. The wire-format path is wired up. What's missing is the IB device registration, which requires a power-cycle. + +**Two ways to unblock validation:** + +1. **Quickest:** physically power-cycle strix-3 and strix-4 (or use BMC). The kernel autoload will run, the source-aware handler will initialize fresh, HELLO will exchange, and IB devices will register. Then: + + ```bash + cd /mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase + ./userspace/bench/rx_side_dma_validate.sh HOSTS='strix-3 strix-4' + ``` + + This assumes the Nixos config on the boot partition points at an .ko that contains the link-layer fix + RX-side DMA patches. **It currently doesn't** — the booted-system kernel-modules still has the old `c80499a1…` module on both nodes. A bare power-cycle will come up with the old module again. You need to do option 2 first, or you'll be in the same state I'm in now: new module on the side, old module booted. + +2. **Correct path:** update `nixos-config/flake.nix` to point the `thunderbolt-ibverbs-kernel` path input at the new commit `6fbca6c` of the GDA branch, `nix flake update thunderbolt-ibverbs-kernel`, then `nixos-rebuild switch` on strix-3 and strix-4, then reboot. The new module will then be the booted module, source-aware handler will init at boot, HELLO will exchange, IB devices will register. + +I haven't done (2) yet because it touches the production nixos-config — wanted your call on whether to push that or do something more surgical first. + +## What I want to also do (uncommitted, ready to push) + +- Update `userspace/bench/rx_side_dma_validate.sh` to pass `--options 'profile=linux_perf tbnet=prefer_rdma bind_services=1'` to `tbv-target-module.sh` (without it, the source-aware handler doesn't init at all, even on a clean boot). +- Add a post-reload sleep + retry loop that checks for `usb4_rdma*` IB devices to appear, and fails fast with diagnostic output if they don't. +- Update the run script to read link_layer from `/sys/class/infiniband/usb4_rdma*/ports/1/link_layer` on both nodes, fail if any port reports anything other than `InfiniBand`. + +## Files + +- Summary (this file): `/mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase/.bench-artifacts/rx-side-dma/STRIX3-4-VALIDATION-2026-08-15.md` +- Run script: `/mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase/userspace/bench/rx_side_dma_validate.sh` (already pushed to origin) +- New commit: `6fbca6c` on `codex/gda-v2-rebased-port` (already pushed) +- Patches: + - `705adf8` `ibdev: report InfiniBand link layer` — was Ethernet, kernel was routing every QP through the RoCE code path that has no resolver, so `ibv_modify_qp(RTR/RTS)` returned `-ENODATA`. Switched to `IB_LINK_LAYER_INFINIBAND` + `RDMA_CORE_CAP_IB_MAD`. + - `f34678d` `ibdev, path: RX-side DMA into dmabuf MRs` — per-fragment DMA into the destination SGL via `dma_map_page`, two new debugfs counters (`data_rx_dmabuf_zcopy`, `data_rx_dmabuf_zcopy_error`). CPU MRs untouched. diff --git a/docs/RX_SIDE_DMA.md b/docs/RX_SIDE_DMA.md new file mode 100644 index 0000000..88e2b0e --- /dev/null +++ b/docs/RX_SIDE_DMA.md @@ -0,0 +1,160 @@ +# RX-side DMA for dmabuf MRs + +**Status:** implemented on `codex/apple-xdomain-property-match` (commit `dc8f43f`). +Pending: live dmabuf-destination smoke test on `strix-1`. +**Goal:** make RDMA WRITE destinations land in dmabuf MRs without CPU-side +copies. CPU-pinned MRs continue to use the existing bounce-buffer path; this +change does not regress them. + +## Why a spec first + +The current RX path always lands ring DMA into a kernel bounce buffer, then +copies from the bounce buffer into the destination MR. That works for +`ib_umem` (CPU-pinned) but breaks for `ib_umem_dmabuf`: + +``` +tbv_umem_copy_to_iova(): + mr->umem->is_dmabuf == true + -> dma_buf_begin_cpu_access() + -> kmap_local_page() + memcpy() // bails on ZONE_DEVICE pages + -> dma_buf_end_cpu_access() +``` + +So GPU dmabuf MRs cannot be RDMA-WRITE destinations today. The dmabuf MR +import path exists (see `tbv_reg_user_mr_dmabuf`), it just has no place to +land data. + +## The change + +DMA each ring frame **directly into the destination MR page** when the MR is +a dmabuf MR, or when the MR's pages cannot be CPU-kmap'd. Otherwise fall +through to the bounce-buffer path unchanged. + +Per-fragment DMA is required because NHI ring descriptors point at one +contiguous physical address; an MR's SGL is scattered. Each PDF frame +becomes its own DMA into one MR page. + +### What stays the same + +- The NHI ring. The descriptor interface already accepts per-frame DMA + addresses (`ring->descriptors[ring->head].address = frame->dma`, + `ring_write_descriptors()`). We are not patching the kernel. +- The reorder / fragment reassembly / WC push logic in `ibdev.c`. +- The CPU-MR bounce-buffer path. +- All TX paths. + +### What changes + +Three pieces in `kernel/path.c` and `kernel/ibdev.c`: + +1. **MR-side mapping.** When the reorder layer accepts the first fragment of + a WRITE that targets a dmabuf MR, look up the rkey, validate the range, + `dma_map_sg_attrs(... DMA_FROM_DEVICE)` on the destination page range, + and stash the mapped SGL on `tbv_rx_reorder_msg`. +2. **Per-fragment DMA.** Replace the call into + `tbv_rx_reorder_store_fragment_locked` for dmabuf destinations with + `tbv_rx_zcopy_store_fragment_locked` that: + - allocates a `tbv_data_frame` from a small `rx_zcopy_pool` + - sets `frame->buf` to the MR page + offset, `frame->dma` to the mapped + address + - `dma_sync_single_for_device(..., DMA_FROM_DEVICE)` then `tb_ring_rx()` + - on completion, `dma_unmap_page(..., DMA_FROM_DEVICE)`, push WC if the + fragment was last, free the frame +3. **CPU fallback unchanged.** If `mr->umem->is_dmabuf == false`, the + existing kernel-bounce-buffer RX path stays in use. + +### Where it lives + +- `kernel/path.c`: extend `tbv_data_frame` with `enum tbv_rx_mode { CPU, + ZCOPY_DMABUF }` and a `tbv_rx_zcopy` struct holding the unmap state. + Add `tbv_path_post_rx_zcopy_frame()` that submits a per-fragment DMA + into the MR page. +- `kernel/ibdev.c`: branch in + `tbv_rx_buffer_write_fragment_locked()`: dmabuf MR -> the new path; CPU + MR -> existing path. +- `kernel/tbv.h`: add the new mode enum and per-frame zcopy bookkeeping on + `tbv_rx_reorder_msg`. + +### Failure handling + +- If the MR's SGL map fails (`dma_map_sg_attrs` returns 0), ack the + fragment with `TBV_NATIVE_SEND_ACK_ERROR`, push `IB_WC_LOC_PROT_ERR`, + drop the reorder message. This is identical to the existing error path + for bad rkey. +- If `dma_map_page` returns an error during per-fragment DMA setup, same + treatment. +- Completion ordering: NHI completes descriptors in arrival order; PDF + frames already carry the fragment index, and the existing + `frag_seen[]` bitmap stays correct because we record `set_bit(frag_idx, + msg->frag_seen)` only after the DMA is posted. The WC is pushed when + the **last** fragment completes, not when the first does. This matches + what dmabuf MR consumers expect. + +### What we explicitly do not do + +- We do not reorder or coalesce DMA. Per-fragment DMA, period. +- We do not modify the NHI ring or the kernel module. +- We do not change TX-side zcopy or the SEND path. SEND lands in a kernel + WQE buffer, not in an MR. +- We do not introduce per-MR DMA contexts; the existing ring DMA device + is reused. + +## Test plan (on strix-1 + strix-2, kernel 7.2.0-rc2) + +1. **CPU regression.** Re-run the existing `userspace/bench/ibv_*.c` + send/recv suite (sender CPU MR, receiver CPU MR). Confirm + `data_rx_completed` matches `data_rx_reorder_delivered` and no + `data_rx_copy_error` increments. +2. **Dmabuf destination regression.** Run + `userspace/bench/dmabuf_mr_probe` with sender CPU MR, receiver GPU + dmabuf MR. Confirm `data_rx_completed > 0`, `data_rx_copy_error` + stays at 0, and `data_rx_dmabuf_zcopy` (new counter) increments. +3. **GPU P2P smoke.** Run + `userspace/bench/hip_rdma_write_visibility_probe.cpp`. The PR-thread + symptom was `data_rx_copy_error=2`, `gpu_seen=0`. With this change we + expect `gpu_seen > 0` and no copy errors. + +## Open questions + +- **RX ring pool size.** NHI rings have a fixed descriptor count. While a + zcopy frame is in flight, the ring has one fewer descriptor for the + bounce path. For bursty small messages this could starve CPU MR RX. + Mitigation: cap concurrent zcopy fragments per QP at + `min(ring->size / 2, TBV_RX_ZCOPY_MAX_PER_QP)`, default 32. +- **GPU page alignment.** ZONE_DEVICE pages are not necessarily aligned + to `TBV_DATA_FRAME_SIZE`. The per-frame DMA must respect + `frame->size` and the offset within the page; this is already how the + ring works (`frame->size` is what the descriptor programs as length), + so no kernel change is needed. We do, however, need to refuse to map + fragments whose first-byte offset into the page is non-zero, because + NHI DMA into a partial page will tear the page. Constraint: a + fragment's destination `(iova & (PAGE_SIZE - 1))` must be 0 modulo the + DMA page granularity. We can satisfy this by mapping page-aligned, + `bytes_to_copy = min(remaining_in_page, frag_len)`, and pushing + fragment splits if needed. + +## Estimated scope + +- ~150-200 lines added in `kernel/path.c` and `kernel/ibdev.c` +- one new counter in `kernel/debugfs.c` +- one probe binary already exists; minimal new tests +- no kernel patch changes + +## Decision + +Approve to proceed? If yes, I will start with the dmabuf-only branch +(smallest diff, leaves CPU MRs untouched) and report back with a +buildable module + a smoke test on strix-1. + +## Run book for the next experiment on strix-1 + strix-2 + +1. **Reload the module.** `tools/tbv-target-module.sh strix-1 --booted-kernel --reload --options 'profile=linux_perf apple_data=N native_data=Y bind_services=Y allocate_rings=Y start_rings=Y negotiate_native=Y enable_tunnels=Y register_verbs=Y zcopy_min_bytes=4294967295'`. Repeat for `strix-2` if the kernel is rebuilt there. Verify `/sys/class/infiniband/usb4_rdma*/ports/1/link_layer` reads `InfiniBand` (not `Ethernet`). +2. **CPU regression first.** Run `userspace/bench/rc_write_verify` between the two nodes. Confirm `data_rx_completed` matches `data_rx_reorder_delivered` and no `data_rx_copy_error` increments in `/sys/kernel/debug/tbv/usb4_rdma*/summary`. This must pass before touching dmabuf. +3. **dmabuf MR probe.** Open a HIP-allocated region on `strix-1`, export it via `hsa_amd_portable_export_dmabuf`, register it via `ibv_reg_dmabuf_mr` on `usb4_rdma0` (or `usb4_rdma1`). Have the peer do an RDMA WRITE into it. Confirm `data_rx_dmabuf_zcopy` (new counter) increments and `data_rx_copy_error` stays at 0. +4. **HIP visibility probe.** Run `userspace/bench/hip_rdma_write_visibility_probe` with `--role recv --kind device --recv-reg dmabuf` on the GPU node and `--role send --kind malloc --source-fill cpu` on the peer. Compare against the `--recv-reg reg_mr` baseline. The PR-thread failure mode was `gpu_seen=0`, `data_rx_copy_error=2`. With this change we expect `gpu_seen > 0` and zero copy errors. +5. **Failure triage.** If `data_rx_dmabuf_zcopy_error` increments, the most common cause is `dma_map_sg` returning 0 (device not IOMMU-mapped or BO not contiguous) or the SGL walk failing (`dest_iova` outside any sg entry). Both are recoverable by aborting the WRITE and pushing `IB_WC_LOC_PROT_ERR`, so the wire stays consistent. + +## What's still missing + +- The same RX-side DMA change needs to be **ported onto `codex/gda-v2-rebased-port`** (the path input the cluster actually builds from). Currently it lives only on `codex/apple-xdomain-property-match`. +- The `nixos-config` flake.lock pin in this commit points at `codex/apple-xdomain-property-match`, but the cluster still consumes the GDA v2 rebase path input. Once the GDA branch picks up these commits, a `colmena build` will pull them in. \ No newline at end of file diff --git a/flake.nix b/flake.nix index de9f6db..d5dc7cb 100644 --- a/flake.nix +++ b/flake.nix @@ -354,6 +354,7 @@ thunderboltKernel = mkThunderboltKernel pkgs; thunderboltLinuxPackages = mkThunderboltLinuxPackages pkgs; moduleForThunderboltKernel = thunderboltLinuxPackages.callPackage ./nix/module.nix { }; + dvProbes = pkgs.callPackage ./nix/dv-probes.nix packageArgs; in { default = module; @@ -364,6 +365,7 @@ thunderbolt-ibverbs = module; thunderbolt-ibverbs-linux-thunderbolt = moduleForThunderboltKernel; tbv-perftest = perftestBench.runner; + tbv-dv-caps-probe = dvProbes; } ) ); @@ -409,6 +411,10 @@ maintainers = with pkgs.lib.maintainers; [ georgewhewell ]; }; }; + tbv-dv-caps-probe = { + type = "app"; + program = lib.getExe pkgsAt.tbv-dv-caps-probe; + }; } ); @@ -429,6 +435,7 @@ proto-smoke = mkProtoSmoke pkgs; rdma-core-usb4 = pkgsAt.rdma-core-usb4; verbs-smoke-build = mkVerbsSmokeBuild pkgs; + tbv-dv-caps-probe = pkgsAt.tbv-dv-caps-probe; } ); @@ -450,6 +457,7 @@ linux-thunderbolt-modules = pkgsAt.linux-thunderbolt-modules; rdma-core-usb4 = pkgsAt.rdma-core-usb4; thunderbolt-ibverbs-linux-thunderbolt = pkgsAt.thunderbolt-ibverbs-linux-thunderbolt; + tbv-dv-caps-probe = pkgsAt.tbv-dv-caps-probe; vm-smoke.nixos = mkNixosVmSmoke pkgs; } ); diff --git a/kernel/Makefile b/kernel/Makefile index 9d6be07..993a90e 100644 --- a/kernel/Makefile +++ b/kernel/Makefile @@ -7,7 +7,7 @@ KDIR ?= /lib/modules/$(KVER)/build PWD := $(shell pwd) obj-m := thunderbolt_ibverbs.o -thunderbolt_ibverbs-y := main.o profile.o backend.o configfs.o core.o debugfs.o ibdev.o link.o native.o apple.o trace.o native_control.o native_control_legacy.o native_control_xdomain.o path.o peer.o rail.o service.o tbnet_identity.o tbnet_minimal.o ../proto/config.o ../proto/identity.o ../proto/reliability.o +thunderbolt_ibverbs-y := main.o profile.o backend.o configfs.o core.o debugfs.o dv.o ibdev.o link.o native.o apple.o trace.o native_control.o native_control_legacy.o native_control_xdomain.o path.o peer.o rail.o service.o tbnet_identity.o tbnet_minimal.o ../proto/config.o ../proto/identity.o ../proto/reliability.o ccflags-y := -Wall -Wextra -Wno-unused-parameter -I$(src) -I$(src)/.. diff --git a/kernel/debugfs.c b/kernel/debugfs.c index cfcda3e..fc68acb 100644 --- a/kernel/debugfs.c +++ b/kernel/debugfs.c @@ -327,6 +327,10 @@ static int tbv_debugfs_summary_show(struct seq_file *s, void *unused) atomic64_read(&state->data_rx_rnr_suppressed)); seq_printf(s, "data_rx_copy_error: %lld\n", atomic64_read(&state->data_rx_copy_error)); + seq_printf(s, "data_rx_dmabuf_zcopy: %lld\n", + atomic64_read(&state->data_rx_dmabuf_zcopy)); + seq_printf(s, "data_rx_dmabuf_zcopy_error: %lld\n", + atomic64_read(&state->data_rx_dmabuf_zcopy_error)); seq_printf(s, "data_rx_send_len_error: %lld\n", atomic64_read(&state->data_rx_send_len_error)); seq_printf(s, "data_rx_send_prot_error: %lld\n", diff --git a/kernel/dv.c b/kernel/dv.c new file mode 100644 index 0000000..9f6ed3d --- /dev/null +++ b/kernel/dv.c @@ -0,0 +1,847 @@ +// SPDX-License-Identifier: GPL-2.0 + +#define pr_fmt(fmt) "thunderbolt_ibverbs: " fmt + +/* + * UVERBS_MODULE_NAME must be defined before including + * rdma/uverbs_named_ioctl.h; the macro is used in the DECLARE_UVERBS_* + * helpers below to namespace generated symbols. + */ +#define UVERBS_MODULE_NAME tbv + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "../userspace/usb4_rdma/usb4_rdma_dv.h" +#include "tbv.h" + +/* + * USB4 RDMA Direct Verbs (DV) ABI surface. + * + * The DV ABI is the software-RNIC contract used by the upcoming GDA path: the + * GPU produces work queue entries into host-visible coherent memory and the + * kernel poll worker consumes them. See userspace/usb4_rdma/usb4_rdma_dv.h + * for the full design, including memory-ordering rules and the generation + * protocol. + * + * Methods wired so far: QUERY_CAPS, CREATE_QUEUE, DESTROY_QUEUE, and KICK + * for NOP smoke tests. Transport opcodes still follow in subsequent commits; + * caps stays 0 to advertise that no SEND/WRITE/READ/ATOMIC WQEs are processed. + */ + +/* ----- per-QP lifecycle helpers (called from ibdev.c) ----------------- */ + +void tbv_dv_qp_state_init(struct tbv_dv_qp_state *dv) +{ + mutex_init(&dv->mutex); + dv->active = false; + dv->generation = 0; + dv->sq_entries = 0; + dv->cq_entries = 0; + dv->sq_head = 0; + dv->cq_tail = 0; + dv->sq_addr = 0; + dv->cq_addr = 0; + dv->doorbell_addr = 0; + dv->sq_umem = NULL; + dv->cq_umem = NULL; + dv->doorbell_umem = NULL; +} + +/* + * Generation policy: u8 monotonic counter, skipping 0 so a freshly + * zero-initialized doorbell page cannot be mistaken for a live generation. + * 8 bits matches USB4_RDMA_DV_TAIL_GENERATION_BITS == 8 — see the packed + * tail format in the ABI header. + */ +static u8 tbv_dv_next_generation(u8 generation) +{ + u8 next = (generation + 1) & 0xff; + + return next ? next : 1; +} + +bool tbv_dv_qp_state_active(struct tbv_dv_qp_state *dv) +{ + return READ_ONCE(dv->active); +} + +/* Release pinned umem regions and reset state to idle. */ +static void tbv_dv_state_release_locked(struct tbv_dv_qp_state *dv) +{ + if (dv->sq_umem) + ib_umem_release(dv->sq_umem); + if (dv->cq_umem) + ib_umem_release(dv->cq_umem); + if (dv->doorbell_umem) + ib_umem_release(dv->doorbell_umem); + dv->sq_umem = NULL; + dv->cq_umem = NULL; + dv->doorbell_umem = NULL; + dv->sq_entries = 0; + dv->cq_entries = 0; + dv->sq_head = 0; + dv->cq_tail = 0; + dv->sq_addr = 0; + dv->cq_addr = 0; + dv->doorbell_addr = 0; + WRITE_ONCE(dv->active, false); +} + +void tbv_dv_qp_state_teardown(struct tbv_dv_qp_state *dv) +{ + mutex_lock(&dv->mutex); + if (dv->active) { + dv->generation = tbv_dv_next_generation(dv->generation); + tbv_dv_state_release_locked(dv); + } + mutex_unlock(&dv->mutex); + mutex_destroy(&dv->mutex); +} + +/* ----- request validation and queue pinning --------------------------- */ + +static bool tbv_dv_reserved_zero(const u32 *reserved, size_t count) +{ + size_t i; + + for (i = 0; i < count; i++) + if (reserved[i]) + return false; + return true; +} + +static bool tbv_dv_ranges_overlap(u64 a_start, u64 a_len, + u64 b_start, u64 b_len) +{ + u64 a_end = a_start + a_len; + u64 b_end = b_start + b_len; + + return a_start < b_end && b_start < a_end; +} + +static int tbv_dv_validate_queue_create( + const struct usb4_rdma_dv_queue_create *req) +{ + u64 sq_bytes; + u64 cq_bytes; + u64 db_bytes = USB4_RDMA_DV_DOORBELL_PAGE_SIZE; + u64 end; + + if (req->abi_version != USB4_RDMA_DV_ABI_VERSION) + return -EINVAL; + if (req->flags) + return -EINVAL; + if (!tbv_dv_reserved_zero(req->reserved, ARRAY_SIZE(req->reserved))) + return -EINVAL; + if (req->sq_entries < USB4_RDMA_DV_MIN_QUEUE_ENTRIES || + req->sq_entries > USB4_RDMA_DV_MAX_SQ_ENTRIES || + req->cq_entries < USB4_RDMA_DV_MIN_QUEUE_ENTRIES || + req->cq_entries > USB4_RDMA_DV_MAX_CQ_ENTRIES) + return -EINVAL; + if (req->sq_stride != USB4_RDMA_DV_WQE_SIZE || + req->cq_stride != USB4_RDMA_DV_CQE_SIZE) + return -EINVAL; + if (!req->sq_addr || !req->cq_addr || !req->doorbell_addr) + return -EINVAL; + if (!IS_ALIGNED(req->sq_addr, USB4_RDMA_DV_WQE_SIZE) || + !IS_ALIGNED(req->cq_addr, USB4_RDMA_DV_CQE_SIZE) || + !IS_ALIGNED(req->doorbell_addr, USB4_RDMA_DV_DOORBELL_PAGE_SIZE)) + return -EINVAL; + if (check_mul_overflow((u64)req->sq_entries, (u64)req->sq_stride, + &sq_bytes) || + check_add_overflow(req->sq_addr, sq_bytes, &end)) + return -EINVAL; + if (check_mul_overflow((u64)req->cq_entries, (u64)req->cq_stride, + &cq_bytes) || + check_add_overflow(req->cq_addr, cq_bytes, &end)) + return -EINVAL; + if (check_add_overflow(req->doorbell_addr, db_bytes, &end)) + return -EINVAL; + if (tbv_dv_ranges_overlap(req->sq_addr, sq_bytes, req->cq_addr, + cq_bytes) || + tbv_dv_ranges_overlap(req->sq_addr, sq_bytes, req->doorbell_addr, + db_bytes) || + tbv_dv_ranges_overlap(req->cq_addr, cq_bytes, req->doorbell_addr, + db_bytes)) + return -EINVAL; + return 0; +} + +struct tbv_dv_pin_result { + struct ib_umem *sq_umem; + struct ib_umem *cq_umem; + struct ib_umem *doorbell_umem; +}; + +static void tbv_dv_pin_release(struct tbv_dv_pin_result *res) +{ + if (res->sq_umem) + ib_umem_release(res->sq_umem); + if (res->cq_umem) + ib_umem_release(res->cq_umem); + if (res->doorbell_umem) + ib_umem_release(res->doorbell_umem); + memset(res, 0, sizeof(*res)); +} + +static struct ib_umem *tbv_dv_umem_get(struct ib_device *dev, u64 addr, + size_t size, int access) +{ +#if LINUX_VERSION_CODE >= KERNEL_VERSION(7, 2, 0) + return ib_umem_get_va(dev, addr, size, access); +#else + return ib_umem_get(dev, addr, size, access); +#endif +} + +static int tbv_dv_validate_cpu_visible_umem(struct ib_umem *umem) +{ + struct sg_table *sgt = &umem->sgt_append.sgt; + struct scatterlist *sg; + unsigned int i; + + for_each_sgtable_sg(sgt, sg, i) { + unsigned int pages; + unsigned int j; + + pages = DIV_ROUND_UP(sg->offset + sg->length, PAGE_SIZE); + for (j = 0; j < pages; j++) { + struct page *page = + pfn_to_page(page_to_pfn(sg_page(sg)) + j); + + if (is_zone_device_page(page)) + return -EOPNOTSUPP; + } + } + + return 0; +} + +static int tbv_dv_pin_queue(struct ib_device *dev, + const struct usb4_rdma_dv_queue_create *req, + struct tbv_dv_pin_result *res) +{ + u64 sq_bytes = (u64)req->sq_entries * req->sq_stride; + u64 cq_bytes = (u64)req->cq_entries * req->cq_stride; + int access = IB_ACCESS_LOCAL_WRITE; + int ret; + + memset(res, 0, sizeof(*res)); + + res->sq_umem = tbv_dv_umem_get(dev, req->sq_addr, sq_bytes, access); + if (IS_ERR(res->sq_umem)) { + int ret = PTR_ERR(res->sq_umem); + + res->sq_umem = NULL; + return ret; + } + ret = tbv_dv_validate_cpu_visible_umem(res->sq_umem); + if (ret) { + tbv_dv_pin_release(res); + return ret; + } + + res->cq_umem = tbv_dv_umem_get(dev, req->cq_addr, cq_bytes, access); + if (IS_ERR(res->cq_umem)) { + int ret = PTR_ERR(res->cq_umem); + + res->cq_umem = NULL; + tbv_dv_pin_release(res); + return ret; + } + ret = tbv_dv_validate_cpu_visible_umem(res->cq_umem); + if (ret) { + tbv_dv_pin_release(res); + return ret; + } + + res->doorbell_umem = tbv_dv_umem_get(dev, req->doorbell_addr, + USB4_RDMA_DV_DOORBELL_PAGE_SIZE, + access); + if (IS_ERR(res->doorbell_umem)) { + int ret = PTR_ERR(res->doorbell_umem); + + res->doorbell_umem = NULL; + tbv_dv_pin_release(res); + return ret; + } + ret = tbv_dv_validate_cpu_visible_umem(res->doorbell_umem); + if (ret) { + tbv_dv_pin_release(res); + return ret; + } + + return 0; +} + +/* + * Write `len` bytes from `src` into pinned userspace pages at + * umem-relative `offset`. The kernel does not export a public + * `ib_umem_copy_to` (only `_from`), so we walk the umem's scatterlist + * and copy page by page. The doorbell page is the only thing this + * helper writes today; the access window is always small and within + * the pinned range. + */ +static int tbv_dv_umem_write(struct ib_umem *umem, size_t offset, + const void *src, size_t len) +{ + struct sg_table *sgt = &umem->sgt_append.sgt; + struct scatterlist *sg; + size_t copied = 0; + size_t skip; + unsigned int i; + + if (!len) + return 0; + if (offset + len < offset) + return -EINVAL; + if (offset + len > umem->length) + return -EFAULT; + + skip = ib_umem_offset(umem) + offset; + for_each_sgtable_sg(sgt, sg, i) { + size_t seg_len = sg->length; + size_t seg_off; + + if (skip >= seg_len) { + skip -= seg_len; + continue; + } + + seg_off = sg->offset + skip; + seg_len -= skip; + skip = 0; + + while (seg_len && copied < len) { + struct page *page; + size_t page_off = offset_in_page(seg_off); + size_t chunk = min_t(size_t, PAGE_SIZE - page_off, + seg_len); + void *kaddr; + + chunk = min_t(size_t, chunk, len - copied); + + page = pfn_to_page(page_to_pfn(sg_page(sg)) + + (seg_off >> PAGE_SHIFT)); + kaddr = kmap_local_page(page); + memcpy((u8 *)kaddr + page_off, + (const u8 *)src + copied, chunk); + flush_dcache_page(page); + kunmap_local(kaddr); + + copied += chunk; + seg_off += chunk; + seg_len -= chunk; + } + + if (copied == len) + break; + } + + return copied == len ? 0 : -EFAULT; +} + +static int tbv_dv_umem_read(struct ib_umem *umem, size_t offset, + void *dst, size_t len) +{ + struct sg_table *sgt = &umem->sgt_append.sgt; + struct scatterlist *sg; + size_t copied = 0; + size_t skip; + unsigned int i; + + if (!len) + return 0; + if (offset + len < offset) + return -EINVAL; + if (offset + len > umem->length) + return -EFAULT; + + skip = ib_umem_offset(umem) + offset; + for_each_sgtable_sg(sgt, sg, i) { + size_t seg_len = sg->length; + size_t seg_off; + + if (skip >= seg_len) { + skip -= seg_len; + continue; + } + + seg_off = sg->offset + skip; + seg_len -= skip; + skip = 0; + + while (seg_len && copied < len) { + struct page *page; + size_t page_off = offset_in_page(seg_off); + size_t chunk = min_t(size_t, PAGE_SIZE - page_off, + seg_len); + void *kaddr; + + chunk = min_t(size_t, chunk, len - copied); + + page = pfn_to_page(page_to_pfn(sg_page(sg)) + + (seg_off >> PAGE_SHIFT)); + kaddr = kmap_local_page(page); + memcpy((u8 *)dst + copied, (const u8 *)kaddr + page_off, + chunk); + kunmap_local(kaddr); + + copied += chunk; + seg_off += chunk; + seg_len -= chunk; + } + + if (copied == len) + break; + } + + return copied == len ? 0 : -EFAULT; +} + +/* + * Write the consumer cacheline of the doorbell page so the GPU side sees + * a known initial (or terminal) state. Used after CREATE_QUEUE to publish + * generation=X / qp_state=LIVE and during DESTROY_QUEUE to publish the + * QP_DEAD sentinel with a bumped generation. + * + * The doorbell page is pinned at offset 0 of its umem, so we use the + * umem-relative offset of the consumer line directly; no base-address + * translation is needed. + */ +static int tbv_dv_write_consumer_line(struct ib_umem *umem, + u32 sq_head, u32 cq_tail, + u32 qp_state, u32 generation) +{ + struct usb4_rdma_dv_doorbell_consumer_line line = { + .sq_head = sq_head, + .cq_tail = cq_tail, + .qp_state = qp_state, + .generation = generation, + }; + + return tbv_dv_umem_write( + umem, offsetof(struct usb4_rdma_dv_doorbell, consumer), + &line, sizeof(line)); +} + +static u32 tbv_dv_index_delta(u32 from, u32 to) +{ + return (to - from) & USB4_RDMA_DV_TAIL_INDEX_MASK; +} + +static int tbv_dv_read_producer_line( + struct tbv_dv_qp_state *dv, + struct usb4_rdma_dv_doorbell_producer_line *line) +{ + return tbv_dv_umem_read( + dv->doorbell_umem, + offsetof(struct usb4_rdma_dv_doorbell, producer), + line, sizeof(*line)); +} + +static int tbv_dv_read_wqe(struct tbv_dv_qp_state *dv, u32 sq_index, + struct usb4_rdma_dv_wqe *wqe) +{ + u32 slot = sq_index % dv->sq_entries; + size_t offset = (size_t)slot * USB4_RDMA_DV_WQE_SIZE; + + return tbv_dv_umem_read(dv->sq_umem, offset, wqe, sizeof(*wqe)); +} + +static int tbv_dv_write_cqe(struct tbv_dv_qp_state *dv, u32 cq_index, + const struct usb4_rdma_dv_cqe *cqe) +{ + u32 slot = cq_index % dv->cq_entries; + size_t offset = (size_t)slot * USB4_RDMA_DV_CQE_SIZE; + + return tbv_dv_umem_write(dv->cq_umem, offset, cqe, sizeof(*cqe)); +} + +static int tbv_dv_emit_cqe_locked(struct tbv_dv_qp_state *dv, + const struct usb4_rdma_dv_wqe *wqe, + u32 status, u32 opcode) +{ + struct usb4_rdma_dv_doorbell_producer_line producer = {}; + struct usb4_rdma_dv_cqe cqe = { + .wr_id = wqe->wr_id, + .status = status, + .opcode = opcode, + .byte_len = status == USB4_RDMA_DV_CQE_SUCCESS ? + wqe->length : 0, + .imm_data = wqe->imm_data, + .qp_state = USB4_RDMA_DV_QP_LIVE, + }; + u32 cq_head; + u32 cq_tail; + int ret; + + ret = tbv_dv_read_producer_line(dv, &producer); + if (ret) + return ret; + + if (producer.generation && + producer.generation != dv->generation) + return -ESTALE; + if (producer.cq_head && + usb4_rdma_dv_tail_generation(producer.cq_head) != dv->generation) + return -ESTALE; + + cq_head = usb4_rdma_dv_tail_index(producer.cq_head); + cq_tail = usb4_rdma_dv_tail_index(dv->cq_tail); + if (tbv_dv_index_delta(cq_head, cq_tail) >= dv->cq_entries) + return -ENOSPC; + + ret = tbv_dv_write_cqe(dv, cq_tail, &cqe); + if (ret) + return ret; + + dv->cq_tail = usb4_rdma_dv_tail_pack(cq_tail + 1, dv->generation); + return 0; +} + +static int tbv_dv_drain_nop_locked(struct tbv_dv_qp_state *dv, u32 sq_tail) +{ + u32 tail_generation = usb4_rdma_dv_tail_generation(sq_tail); + u32 head = usb4_rdma_dv_tail_index(dv->sq_head); + u32 tail = usb4_rdma_dv_tail_index(sq_tail); + u32 pending; + int ret; + + if (!dv->active) + return -ENOENT; + if (tail_generation != dv->generation) + return -ESTALE; + pending = tbv_dv_index_delta(head, tail); + if (pending > dv->sq_entries) + return -EINVAL; + + while (head != tail) { + struct usb4_rdma_dv_wqe wqe = {}; + u32 status = USB4_RDMA_DV_CQE_SUCCESS; + bool signal; + + ret = tbv_dv_read_wqe(dv, head, &wqe); + if (ret) + return ret; + + signal = wqe.flags & USB4_RDMA_DV_WQE_F_SIGNALED; + if (wqe.generation != dv->generation) + status = USB4_RDMA_DV_CQE_STALE_GEN; + else if (wqe.opcode != USB4_RDMA_DV_WQE_NOP) + status = USB4_RDMA_DV_CQE_GENERAL_ERR; + + if (signal || status != USB4_RDMA_DV_CQE_SUCCESS) { + ret = tbv_dv_emit_cqe_locked(dv, &wqe, status, + wqe.opcode); + if (ret) + return ret; + } + + head = (head + 1) & USB4_RDMA_DV_TAIL_INDEX_MASK; + dv->sq_head = usb4_rdma_dv_tail_pack(head, dv->generation); + } + + return tbv_dv_write_consumer_line(dv->doorbell_umem, dv->sq_head, + dv->cq_tail, USB4_RDMA_DV_QP_LIVE, + dv->generation); +} + +/* ----- uverbs methods ------------------------------------------------- */ + +static int UVERBS_HANDLER(USB4_RDMA_DV_METHOD_QUERY_CAPS)( + struct uverbs_attr_bundle *attrs) +{ + struct usb4_rdma_dv_query_caps_resp resp = { + .abi_version = USB4_RDMA_DV_ABI_VERSION, + /* + * No transport opcodes are wired through the DV consumer yet, + * so we advertise an empty capability bitmap. Each subsequent + * commit that enables a real WQE opcode will OR the matching + * USB4_RDMA_DV_CAP_* bit in here so userspace can detect what + * the kernel will actually consume. + */ + .caps = 0, + .max_sq_entries = USB4_RDMA_DV_MAX_SQ_ENTRIES, + .max_cq_entries = USB4_RDMA_DV_MAX_CQ_ENTRIES, + .default_sq_entries = USB4_RDMA_DV_DEFAULT_SQ_ENTRIES, + .default_cq_entries = USB4_RDMA_DV_DEFAULT_CQ_ENTRIES, + .wqe_size = USB4_RDMA_DV_WQE_SIZE, + .cqe_size = USB4_RDMA_DV_CQE_SIZE, + .doorbell_record_size = USB4_RDMA_DV_DOORBELL_RECORD_SIZE, + .doorbell_page_size = USB4_RDMA_DV_DOORBELL_PAGE_SIZE, + .tail_index_bits = USB4_RDMA_DV_TAIL_INDEX_BITS, + .tail_generation_bits = USB4_RDMA_DV_TAIL_GENERATION_BITS, + }; + struct ib_ucontext *ib_uctx; + + /* + * Compile-time guards: the ABI sizes that userspace receives in the + * QUERY_CAPS response are the same sizes the future producer/consumer + * code will assume for the on-the-wire-via-memory structs. If a + * struct field is added without updating the corresponding _SIZE + * constant (and bumping the ABI version), userspace and kernel will + * silently disagree about the layout. Catch that at build time. + */ + BUILD_BUG_ON(sizeof(struct usb4_rdma_dv_wqe) != + USB4_RDMA_DV_WQE_SIZE); + BUILD_BUG_ON(sizeof(struct usb4_rdma_dv_cqe) != + USB4_RDMA_DV_CQE_SIZE); + BUILD_BUG_ON(sizeof(struct usb4_rdma_dv_doorbell_producer_line) != + USB4_RDMA_DV_DOORBELL_LINE_SIZE); + BUILD_BUG_ON(sizeof(struct usb4_rdma_dv_doorbell_consumer_line) != + USB4_RDMA_DV_DOORBELL_LINE_SIZE); + BUILD_BUG_ON(sizeof(struct usb4_rdma_dv_doorbell) != + USB4_RDMA_DV_DOORBELL_RECORD_SIZE); + + ib_uctx = ib_uverbs_get_ucontext(attrs); + if (IS_ERR(ib_uctx)) + return PTR_ERR(ib_uctx); + if (!tbv_ibdev_state(ib_uctx->device)) + return -ENODEV; + + return uverbs_copy_to(attrs, USB4_RDMA_DV_ATTR_QUERY_CAPS_RESP, + &resp, sizeof(resp)); +} + +static int UVERBS_HANDLER(USB4_RDMA_DV_METHOD_CREATE_QUEUE)( + struct uverbs_attr_bundle *attrs) +{ + struct usb4_rdma_dv_queue_create req = {}; + struct usb4_rdma_dv_queue_resp resp = {}; + struct tbv_dv_pin_result res = {}; + struct tbv_dv_qp_state *dv; + struct ib_qp *ibqp; + struct tbv_qp *tqp; + u8 generation; + int ret; + + ibqp = uverbs_attr_get_obj(attrs, USB4_RDMA_DV_ATTR_CREATE_QUEUE_QP); + if (IS_ERR(ibqp)) + return PTR_ERR(ibqp); + + ret = uverbs_copy_from(&req, attrs, + USB4_RDMA_DV_ATTR_CREATE_QUEUE_REQ); + if (ret) + return ret; + + ret = tbv_dv_validate_queue_create(&req); + if (ret) + return ret; + + ret = tbv_dv_pin_queue(ibqp->device, &req, &res); + if (ret) + return ret; + + tqp = tbv_qp_from_ibqp(ibqp); + dv = tbv_qp_dv_state(tqp); + + mutex_lock(&dv->mutex); + if (dv->active) { + mutex_unlock(&dv->mutex); + tbv_dv_pin_release(&res); + return -EBUSY; + } + + generation = tbv_dv_next_generation(dv->generation); + dv->generation = generation; + dv->sq_entries = req.sq_entries; + dv->cq_entries = req.cq_entries; + dv->sq_head = usb4_rdma_dv_tail_pack(0, generation); + dv->cq_tail = usb4_rdma_dv_tail_pack(0, generation); + dv->sq_addr = req.sq_addr; + dv->cq_addr = req.cq_addr; + dv->doorbell_addr = req.doorbell_addr; + dv->sq_umem = res.sq_umem; + dv->cq_umem = res.cq_umem; + dv->doorbell_umem = res.doorbell_umem; + memset(&res, 0, sizeof(res)); + WRITE_ONCE(dv->active, true); + + /* + * Publish the LIVE state to the doorbell consumer line with packed + * (index=0, generation) tail words. The producer side mirrors this + * generation into its own line before issuing the first WQE. + */ + ret = tbv_dv_write_consumer_line( + dv->doorbell_umem, dv->sq_head, dv->cq_tail, + USB4_RDMA_DV_QP_LIVE, generation); + if (ret) + goto err_release_locked; + + resp.qp_num = ibqp->qp_num; + resp.generation = generation; + + ret = uverbs_copy_to(attrs, USB4_RDMA_DV_ATTR_CREATE_QUEUE_RESP, + &resp, sizeof(resp)); + if (ret) + goto err_release_locked; + + mutex_unlock(&dv->mutex); + return 0; + +err_release_locked: + dv->generation = tbv_dv_next_generation(dv->generation); + tbv_dv_state_release_locked(dv); + mutex_unlock(&dv->mutex); + return ret; +} + +static int UVERBS_HANDLER(USB4_RDMA_DV_METHOD_DESTROY_QUEUE)( + struct uverbs_attr_bundle *attrs) +{ + struct tbv_dv_qp_state *dv; + struct ib_umem *doorbell_umem = NULL; + struct ib_qp *ibqp; + struct tbv_qp *tqp; + u8 generation = 0; + int ret = 0; + + ibqp = uverbs_attr_get_obj(attrs, USB4_RDMA_DV_ATTR_DESTROY_QUEUE_QP); + if (IS_ERR(ibqp)) + return PTR_ERR(ibqp); + + tqp = tbv_qp_from_ibqp(ibqp); + dv = tbv_qp_dv_state(tqp); + + mutex_lock(&dv->mutex); + if (!dv->active) { + mutex_unlock(&dv->mutex); + return -ENOENT; + } + + /* + * Bump generation before releasing so any stale producer write + * landing on the now-defunct doorbell page is detectable. The + * QP_DEAD sentinel + bumped generation are the contract documented + * in the ABI header. + */ + generation = tbv_dv_next_generation(dv->generation); + dv->generation = generation; + + /* Keep one reference to the doorbell umem so we can write the + * sentinel after releasing the rest. */ + doorbell_umem = dv->doorbell_umem; + dv->doorbell_umem = NULL; + + tbv_dv_state_release_locked(dv); + mutex_unlock(&dv->mutex); + + if (doorbell_umem) { + ret = tbv_dv_write_consumer_line( + doorbell_umem, + usb4_rdma_dv_tail_pack(0, generation), + usb4_rdma_dv_tail_pack(0, generation), + USB4_RDMA_DV_QP_DEAD, generation); + ib_umem_release(doorbell_umem); + /* + * A late write failure here is not surfaced to userspace + * because the queue has already been logically torn down; + * the only effect is that the GPU may not observe the DEAD + * sentinel through this doorbell page. The bumped generation + * in the next CREATE_QUEUE response is sufficient to keep + * generation-checking consumers safe. + */ + (void)ret; + ret = 0; + } + + return ret; +} + +static int UVERBS_HANDLER(USB4_RDMA_DV_METHOD_KICK)( + struct uverbs_attr_bundle *attrs) +{ + struct usb4_rdma_dv_kick req = {}; + struct tbv_dv_qp_state *dv; + struct ib_qp *ibqp; + struct tbv_qp *tqp; + int ret; + + ibqp = uverbs_attr_get_obj(attrs, USB4_RDMA_DV_ATTR_KICK_QP); + if (IS_ERR(ibqp)) + return PTR_ERR(ibqp); + + ret = uverbs_copy_from(&req, attrs, USB4_RDMA_DV_ATTR_KICK_REQ); + if (ret) + return ret; + if (req.flags || + !tbv_dv_reserved_zero(req.reserved, ARRAY_SIZE(req.reserved))) + return -EINVAL; + + tqp = tbv_qp_from_ibqp(ibqp); + dv = tbv_qp_dv_state(tqp); + + mutex_lock(&dv->mutex); + ret = tbv_dv_drain_nop_locked(dv, req.sq_tail); + mutex_unlock(&dv->mutex); + return ret; +} + +DECLARE_UVERBS_NAMED_METHOD( + USB4_RDMA_DV_METHOD_QUERY_CAPS, + UVERBS_ATTR_PTR_OUT( + USB4_RDMA_DV_ATTR_QUERY_CAPS_RESP, + UVERBS_ATTR_STRUCT(struct usb4_rdma_dv_query_caps_resp, + reserved), + UA_MANDATORY)); + +DECLARE_UVERBS_NAMED_METHOD( + USB4_RDMA_DV_METHOD_CREATE_QUEUE, + UVERBS_ATTR_IDR(USB4_RDMA_DV_ATTR_CREATE_QUEUE_QP, + UVERBS_OBJECT_QP, + UVERBS_ACCESS_READ, + UA_MANDATORY), + UVERBS_ATTR_PTR_IN( + USB4_RDMA_DV_ATTR_CREATE_QUEUE_REQ, + UVERBS_ATTR_STRUCT(struct usb4_rdma_dv_queue_create, reserved), + UA_MANDATORY), + UVERBS_ATTR_PTR_OUT( + USB4_RDMA_DV_ATTR_CREATE_QUEUE_RESP, + UVERBS_ATTR_STRUCT(struct usb4_rdma_dv_queue_resp, reserved0), + UA_MANDATORY)); + +DECLARE_UVERBS_NAMED_METHOD( + USB4_RDMA_DV_METHOD_DESTROY_QUEUE, + UVERBS_ATTR_IDR(USB4_RDMA_DV_ATTR_DESTROY_QUEUE_QP, + UVERBS_OBJECT_QP, + UVERBS_ACCESS_READ, + UA_MANDATORY)); + +DECLARE_UVERBS_NAMED_METHOD( + USB4_RDMA_DV_METHOD_KICK, + UVERBS_ATTR_IDR(USB4_RDMA_DV_ATTR_KICK_QP, + UVERBS_OBJECT_QP, + UVERBS_ACCESS_READ, + UA_MANDATORY), + UVERBS_ATTR_PTR_IN( + USB4_RDMA_DV_ATTR_KICK_REQ, + UVERBS_ATTR_STRUCT(struct usb4_rdma_dv_kick, reserved), + UA_MANDATORY)); + +DECLARE_UVERBS_GLOBAL_METHODS( + USB4_RDMA_DV_OBJECT_DEVICE, + &UVERBS_METHOD(USB4_RDMA_DV_METHOD_QUERY_CAPS), + &UVERBS_METHOD(USB4_RDMA_DV_METHOD_CREATE_QUEUE), + &UVERBS_METHOD(USB4_RDMA_DV_METHOD_DESTROY_QUEUE), + &UVERBS_METHOD(USB4_RDMA_DV_METHOD_KICK)); + +const struct uapi_definition tbv_uapi_defs[] = { + UAPI_DEF_CHAIN_OBJ_TREE_NAMED(USB4_RDMA_DV_OBJECT_DEVICE), + {}, +}; diff --git a/kernel/ibdev.c b/kernel/ibdev.c index cb96971..69a81bf 100644 --- a/kernel/ibdev.c +++ b/kernel/ibdev.c @@ -6,6 +6,8 @@ #include #include #include +#include +#include #include #include #include @@ -31,12 +33,20 @@ #include #include #include +#include #include "../proto/apple_tx.h" #include "../proto/native_data.h" #include "../proto/reliability.h" #include "tbv.h" +/* + * Defined in kernel/dv.c. The DV uapi definition table is wired into the + * ib_device at register time so userspace can issue the USB4_RDMA_DV_* + * private uverbs methods. See userspace/usb4_rdma/usb4_rdma_dv.h. + */ +extern const struct uapi_definition tbv_uapi_defs[]; + #define TBV_IBDEV_ABI_VERSION 1 #define TBV_IBDEV_PORTS 1 #define TBV_IBDEV_MAX_QP 256 @@ -95,6 +105,10 @@ static uint zcopy_min_bytes; module_param(zcopy_min_bytes, uint, 0644); MODULE_PARM_DESC(zcopy_min_bytes, "Minimum native bytes before raw zero-copy streaming is requested; retryable native RC WRITE falls back to framed copies"); +static bool native_p2p_zcopy; +module_param(native_p2p_zcopy, bool, 0644); +MODULE_PARM_DESC(native_p2p_zcopy, + "Experimental: allow native RDMA WRITE raw zcopy from dma-buf MRs using pre-mapped DMA addresses"); static uint qp_timeout_ms = TBV_QP_TIMEOUT_DEFAULT_MS; module_param(qp_timeout_ms, uint, 0644); @@ -355,6 +369,21 @@ struct tbv_rx_reorder_msg { bool complete; bool with_imm; bool solicited; + /* + * RX-side DMA into a dmabuf MR. Populated on first fragment of + * an RDMA WRITE that targets a dmabuf MR; released when the + * last fragment completes (or the msg is dropped). refcount_t + * `refs` is bumped for each in-flight zcopy frame, so the msg + * stays alive until every DMA completion has been processed. + * zcopy_sgt is the SGL we DMA-mapped (== mr->umem->sgt_append.sgt), + * kept so we can dma_unmap_sg on teardown. + */ + struct sg_table *zcopy_sgt; + struct tbv_mr *zcopy_mr; + struct device *zcopy_dma_dev; + int zcopy_sgt_nents; + bool zcopy_mapped; + refcount_t refs; }; struct tbv_apple_pending_rx { @@ -463,8 +492,19 @@ struct tbv_qp { * message and must not seed a new reassembly. Protected by rx_lock. */ bool apple_rx_discard; + struct tbv_dv_qp_state dv; }; +struct tbv_qp *tbv_qp_from_ibqp(struct ib_qp *ibqp) +{ + return container_of(ibqp, struct tbv_qp, base); +} + +struct tbv_dv_qp_state *tbv_qp_dv_state(struct tbv_qp *tqp) +{ + return &tqp->dv; +} + struct tbv_mr { struct ib_mr base; struct tbv_state *owner; @@ -631,6 +671,8 @@ static int tbv_send_read_response_ctx(struct tbv_read_resp_ctx *ctx); static int tbv_umem_page_from_addr(struct tbv_mr *mr, u64 addr, u32 max_len, struct page **page_out, u32 *page_off_out, u32 *len_out); +static int tbv_umem_dma_from_addr(struct tbv_mr *mr, u64 addr, u32 max_len, + dma_addr_t *dma_out, u32 *len_out); static int tbv_rx_copy_to_wqe(struct tbv_state *state, const struct tbv_recv_wqe *wqe, u32 offset, const void *payload, u32 len, u32 *delivered); @@ -651,6 +693,7 @@ static void tbv_rx_drop_reorder_msg_locked(struct tbv_state *state, static void tbv_rx_drain_reorder_locked(struct tbv_state *state, struct tbv_qp *tqp, struct tbv_path *rx_path); +static void tbv_rx_zcopy_drain_work(struct work_struct *work); static void tbv_qp_flush_apple_pending(struct tbv_qp *tqp); static void tbv_apple_rx_drain_pending_locked(struct tbv_state *state, struct tbv_qp *tqp); @@ -804,7 +847,7 @@ static struct tbv_ibdev *tbv_to_ibdev(struct ib_device *ibdev) return container_of(ibdev, struct tbv_ibdev, base); } -static struct tbv_state *tbv_ibdev_state(struct ib_device *ibdev) +struct tbv_state *tbv_ibdev_state(struct ib_device *ibdev) { struct tbv_ibdev *dev = tbv_to_ibdev(ibdev); @@ -832,6 +875,26 @@ static struct tbv_mr *tbv_mr_get(struct tbv_state *state, u32 key) return mr; } +static void tbv_dmabuf_invalidate(struct dma_buf_attachment *attach) +{ + struct ib_umem_dmabuf *umem_dmabuf = attach->importer_priv; + + if (!umem_dmabuf) + return; + + ibdev_warn_ratelimited(umem_dmabuf->umem.ibdev, + "dmabuf-backed MR was invalidated; revoke handling is not implemented yet\n"); +} + +static const struct dma_buf_attach_ops tbv_dmabuf_attach_ops = { + .allow_peer2peer = false, +#if LINUX_VERSION_CODE >= KERNEL_VERSION(7, 1, 0) + .invalidate_mappings = tbv_dmabuf_invalidate, +#else + .move_notify = tbv_dmabuf_invalidate, +#endif +}; + static void tbv_mr_free(struct tbv_mr *mr) { if (mr->umem) @@ -2263,7 +2326,7 @@ static int tbv_get_port_immutable(struct ib_device *ibdev, u32 port_num, immutable->pkey_tbl_len = attr.pkey_tbl_len; immutable->gid_tbl_len = attr.gid_tbl_len; - immutable->core_cap_flags = RDMA_CORE_PORT_IBA_ROCE_UDP_ENCAP; + immutable->core_cap_flags = RDMA_CORE_CAP_IB_MAD; immutable->max_mad_size = IB_MGMT_MAD_SIZE; return 0; } @@ -2271,7 +2334,7 @@ static int tbv_get_port_immutable(struct ib_device *ibdev, u32 port_num, static enum rdma_link_layer tbv_get_link_layer(struct ib_device *ibdev, u32 port_num) { - return IB_LINK_LAYER_ETHERNET; + return IB_LINK_LAYER_INFINIBAND; } static const char *tbv_ibdev_netdev_name_for(struct tbv_state *state, @@ -2577,6 +2640,7 @@ static int tbv_create_qp(struct ib_qp *qp, struct ib_qp_init_attr *init_attr, tqp->owner = state; spin_lock_init(&tqp->lock); mutex_init(&tqp->rx_lock); + tbv_dv_qp_state_init(&tqp->dv); init_waitqueue_head(&tqp->credit_wait); init_waitqueue_head(&tqp->apple_tx_wait); init_waitqueue_head(&tqp->refs_wait); @@ -2649,6 +2713,7 @@ static int tbv_destroy_qp(struct ib_qp *qp, struct ib_udata *udata) mutex_unlock(&tqp->owner->lock); } + tbv_dv_qp_state_teardown(&tqp->dv); wake_up_all(&tqp->credit_wait); wake_up_all(&tqp->apple_tx_wait); cancel_work_sync(&tqp->apple_sq_work); @@ -4310,6 +4375,47 @@ static int tbv_send_page_stream_next(void *ctx, struct page **page, return -EFAULT; } +static int tbv_send_dma_stream_next(void *ctx, dma_addr_t *dma, u32 *length, + tbv_path_tx_done_fn *done, + void **done_ctx) +{ + struct tbv_send_page_stream *stream = ctx; + u32 skipped = 0; + int i; + + for (i = 0; i < stream->nsegs; i++) { + struct tbv_send_segment *seg = &stream->segs[i]; + u32 seg_off = 0; + u32 remaining; + int ret; + + if (stream->offset >= skipped + seg->length) { + skipped += seg->length; + continue; + } + if (stream->offset > skipped) + seg_off = stream->offset - skipped; + + remaining = min_t(u32, seg->length - seg_off, + stream->total_len - stream->offset); + remaining = min_t(u32, remaining, stream->max_chunk); + ret = tbv_umem_dma_from_addr(seg->mr, seg->addr + seg_off, + remaining, dma, length); + if (ret) + return ret; + + stream->offset += *length; + refcount_inc(&stream->refs); + atomic_inc(&stream->send->tx_pending); + tbv_send_ctx_get(stream->send); + *done = tbv_send_page_stream_done; + *done_ctx = stream; + return 0; + } + + return -EFAULT; +} + static bool tbv_should_zcopy_payload(u32 len) { return len && zcopy_min_bytes && len >= zcopy_min_bytes; @@ -4385,6 +4491,51 @@ static bool tbv_send_segments_zcopy_safe(struct tbv_send_segment *segs, return true; } +static bool tbv_send_segments_dma_mapped(struct tbv_send_segment *segs, + int nsegs, u32 total_len) +{ + u32 offset = 0; + + while (offset < total_len) { + u32 skipped = 0; + bool found = false; + int i; + + for (i = 0; i < nsegs; i++) { + struct tbv_send_segment *seg = &segs[i]; + dma_addr_t dma; + u32 len; + u32 seg_off = 0; + u32 remaining; + int ret; + + if (offset >= skipped + seg->length) { + skipped += seg->length; + continue; + } + if (offset > skipped) + seg_off = offset - skipped; + + remaining = min_t(u32, seg->length - seg_off, + total_len - offset); + ret = tbv_umem_dma_from_addr(seg->mr, + seg->addr + seg_off, + remaining, &dma, &len); + if (ret || !len) + return false; + + offset += len; + found = true; + break; + } + + if (!found) + return false; + } + + return true; +} + static bool tbv_send_ctx_is_send(const struct tbv_send_ctx *ctx) { return ctx->opcode == TBV_NATIVE_DATA_OP_SEND || @@ -4458,6 +4609,8 @@ static int tbv_native_send_ctx_post_frames(struct tbv_send_ctx *ctx, bool zcopy_requested; bool raw_zcopy_allowed = false; bool zcopy_safe = false; + bool dma_zcopy_allowed = false; + bool dma_zcopy_safe = false; bool sent_any = false; int ret = 0; @@ -4490,10 +4643,17 @@ static int tbv_native_send_ctx_post_frames(struct tbv_send_ctx *ctx, if (raw_zcopy_allowed) zcopy_safe = tbv_send_segments_zcopy_safe( ctx->segs, ctx->nsegs, ctx->total_len); + dma_zcopy_allowed = READ_ONCE(native_p2p_zcopy); + if (dma_zcopy_allowed) + dma_zcopy_safe = tbv_send_segments_dma_mapped( + ctx->segs, ctx->nsegs, ctx->total_len); } - if (zcopy_requested && raw_zcopy_allowed && zcopy_safe) { + if (zcopy_requested && + ((raw_zcopy_allowed && zcopy_safe) || + (dma_zcopy_allowed && dma_zcopy_safe))) { struct tbv_send_page_stream *stream; + bool use_dma_stream = dma_zcopy_allowed && dma_zcopy_safe; if (reason == TBV_SEND_POST_INITIAL) atomic64_inc(&tqp->owner->data_wr_zcopy); @@ -4524,10 +4684,16 @@ static int tbv_native_send_ctx_post_frames(struct tbv_send_ctx *ctx, tbv_send_ctx_get(ctx); atomic_inc(&ctx->tx_pending); atomic64_inc(&tqp->owner->data_wr_path_send); - ret = tbv_path_send_page_stream(path, &hdr, ctx->total_len, 0, - tbv_send_tx_done, ctx, - tbv_send_page_stream_next, - stream); + if (use_dma_stream) + ret = tbv_path_send_dma_stream( + path, &hdr, ctx->total_len, 0, + tbv_send_tx_done, ctx, tbv_send_dma_stream_next, + stream); + else + ret = tbv_path_send_page_stream( + path, &hdr, ctx->total_len, 0, + tbv_send_tx_done, ctx, + tbv_send_page_stream_next, stream); tbv_release_path_refs(&path, 1); tbv_send_page_stream_put(stream); if (ret) { @@ -4542,7 +4708,8 @@ static int tbv_native_send_ctx_post_frames(struct tbv_send_ctx *ctx, if (reason == TBV_SEND_POST_INITIAL) { if (zcopy_requested) { atomic64_inc(&tqp->owner->data_wr_zcopy_fallback); - if (raw_zcopy_allowed && !zcopy_safe) + if ((raw_zcopy_allowed && !zcopy_safe) || + (dma_zcopy_allowed && !dma_zcopy_safe)) atomic64_inc( &tqp->owner->data_wr_zcopy_fallback_unsafe_sge); } @@ -5236,6 +5403,18 @@ static int tbv_post_send(struct ib_qp *qp, const struct ib_send_wr *wr, const struct ib_send_wr *cur; int ret; + /* + * GDA-exclusive QP semantics: once a DV queue is attached, the + * GPU owns the send queue and standard ibv_post_send is rejected. + * post_recv is intentionally unaffected — the receive queue stays + * kernel-owned in v1. + */ + if (tbv_dv_qp_state_active(&tqp->dv)) { + if (bad_wr) + *bad_wr = wr; + return -EBUSY; + } + for (cur = wr; cur; cur = cur->next) { ret = tbv_post_send_one(tqp, cur); if (ret) { @@ -6690,6 +6869,53 @@ static int tbv_umem_page_from_addr(struct tbv_mr *mr, u64 addr, u32 max_len, return -EFAULT; } +static int tbv_umem_dma_from_addr(struct tbv_mr *mr, u64 addr, u32 max_len, + dma_addr_t *dma_out, u32 *len_out) +{ + struct sg_table *sgt = &mr->umem->sgt_append.sgt; + struct scatterlist *sg; + size_t offset; + u64 end; + u64 mr_end; + unsigned int i; + + if (!max_len) + return -EINVAL; + if (mr->dma_mr) + return -EOPNOTSUPP; + if (!mr->umem || !mr->umem->is_dmabuf) + return -EOPNOTSUPP; + if (check_add_overflow(addr, (u64)max_len, &end)) + return -EINVAL; + if (check_add_overflow(mr->start, mr->length, &mr_end)) + return -EINVAL; + if (addr < mr->start || end > mr_end) + return -EFAULT; + + offset = ib_umem_offset(mr->umem) + addr - mr->start; + for_each_sgtable_dma_sg(sgt, sg, i) { + size_t seg_len = sg_dma_len(sg); + size_t chunk; + + if (!seg_len) + continue; + if (offset >= seg_len) { + offset -= seg_len; + continue; + } + + chunk = min_t(size_t, seg_len - offset, max_len); + if (!chunk) + return -EFAULT; + + *dma_out = sg_dma_address(sg) + offset; + *len_out = chunk; + return 0; + } + + return -EFAULT; +} + static int tbv_copy_to_read_segments(struct tbv_read_ctx *read, u32 offset, const void *payload, u32 len) { @@ -6836,13 +7062,25 @@ static int tbv_rx_copy_to_wqe(struct tbv_state *state, return 0; } -static void tbv_rx_reorder_free_msg(struct tbv_rx_reorder_msg *msg) +static void tbv_rx_reorder_msg_release(struct tbv_rx_reorder_msg *msg) { struct tbv_rx_reorder_frag *frag; struct tbv_rx_reorder_frag *tmp; if (!msg) return; + if (msg->zcopy_mapped) { + struct device *dma_dev = msg->zcopy_dma_dev; + + if (dma_dev && msg->zcopy_sgt && msg->zcopy_sgt_nents > 0) + dma_unmap_sg(dma_dev, msg->zcopy_sgt->sgl, + msg->zcopy_sgt_nents, DMA_FROM_DEVICE); + msg->zcopy_mapped = false; + } + if (msg->zcopy_mr) { + tbv_mr_put(msg->zcopy_mr); + msg->zcopy_mr = NULL; + } list_for_each_entry_safe(frag, tmp, &msg->frags, node) { list_del(&frag->node); kfree(frag); @@ -6850,6 +7088,22 @@ static void tbv_rx_reorder_free_msg(struct tbv_rx_reorder_msg *msg) kfree(msg); } +static void tbv_rx_reorder_msg_get(struct tbv_rx_reorder_msg *msg) +{ + refcount_inc(&msg->refs); +} + +static void tbv_rx_reorder_msg_put(struct tbv_rx_reorder_msg *msg) +{ + if (refcount_dec_and_test(&msg->refs)) + tbv_rx_reorder_msg_release(msg); +} + +static void tbv_rx_reorder_free_msg(struct tbv_rx_reorder_msg *msg) +{ + tbv_rx_reorder_msg_put(msg); +} + static void tbv_qp_flush_reorder(struct tbv_qp *tqp) { struct tbv_rx_reorder_msg *msg; @@ -7323,7 +7577,7 @@ static bool tbv_rx_deliver_reorder_write_locked(struct tbv_state *state, atomic64_inc(&state->data_rx_copy_error); status = IB_WC_LOC_PROT_ERR; ret = -EACCES; - } else { + } else if (!msg->zcopy_mapped) { list_for_each_entry(frag, &msg->frags, node) { u64 copy_addr; @@ -7552,6 +7806,7 @@ static void tbv_rx_buffer_read_req_locked( } INIT_LIST_HEAD(&msg->frags); + refcount_set(&msg->refs, 1); msg->first_jiffies = jiffies; msg->kind = TBV_RX_REORDER_READ_REQ; msg->remote_addr = hdr->remote_addr; @@ -7634,8 +7889,8 @@ static void tbv_rx_buffer_fragment_locked(struct tbv_state *state, "reorder alloc failed", false); return; } - INIT_LIST_HEAD(&msg->frags); + refcount_set(&msg->refs, 1); msg->first_jiffies = jiffies; msg->kind = TBV_RX_REORDER_SEND; msg->src_qp = hdr->src_qp; @@ -7692,6 +7947,192 @@ static void tbv_rx_buffer_fragment_locked(struct tbv_state *state, tbv_rx_drain_reorder_locked(state, tqp, rx_path); } +static bool tbv_rx_zcopy_try_lock(struct tbv_state *state, + struct tbv_qp *tqp, + struct tbv_rx_reorder_msg *msg) +{ + struct tbv_mr *probe; + + if (!state) + return false; + probe = tbv_mr_get(state, msg->rkey); + if (!probe) + return false; + if (!probe->umem || !probe->umem->is_dmabuf) { + tbv_mr_put(probe); + return false; + } + tbv_mr_put(probe); + return true; +} + + +struct tbv_rx_zcopy_frag { + struct work_struct work; + struct tbv_qp *tqp; + struct tbv_path *rx_path; + struct tbv_rx_reorder_msg *msg; + u32 frag_idx; + int status; +}; + +void tbv_rx_zcopy_drain_work(struct work_struct *work) +{ + struct tbv_rx_zcopy_frag *zf = container_of(work, + struct tbv_rx_zcopy_frag, work); + struct tbv_qp *tqp = zf->tqp; + struct tbv_rx_reorder_msg *msg = zf->msg; + struct tbv_state *state = tqp->owner; + + if (zf->status) + tbv_rx_record_send_error(state, "dmabuf_zcopy", tqp, + msg->src_qp, msg->psn, + IB_WC_REM_OP_ERR, false, + msg->total_len, 0, + msg->received, 0, 0, 0); + mutex_lock(&tqp->rx_lock); + if (test_and_set_bit(zf->frag_idx, msg->frag_seen)) { + /* duplicate completion; another worker won the race */ + mutex_unlock(&tqp->rx_lock); + goto out; + } + if (!zf->status) + msg->frags_received++; + if (msg->frags_received >= msg->frag_count || zf->status) + msg->complete = true; + if (msg->complete) + tbv_rx_drain_reorder_locked(state, tqp, zf->rx_path); + mutex_unlock(&tqp->rx_lock); + +out: + tbv_rx_reorder_msg_put(msg); + kfree(zf); +} + +void tbv_rx_zcopy_complete(void *ctx, int status) +{ + struct tbv_rx_zcopy_frag *zf = ctx; + struct workqueue_struct *wq; + + if (!zf) + return; + zf->status = status; + wq = (zf->tqp->owner && zf->tqp->owner->workqueue) ? + zf->tqp->owner->workqueue : system_unbound_wq; + if (!queue_work(wq, &zf->work)) + tbv_rx_zcopy_drain_work(&zf->work); +} + +static int tbv_rx_zcopy_prepare_msg(struct tbv_state *state, + struct tbv_path *rx_path, + struct tbv_qp *tqp, + struct tbv_rx_reorder_msg *msg) +{ + struct tbv_mr *mr; + struct device *dma_dev; + struct sg_table *sgt; + int n; + + if (!state || !rx_path || !rx_path->rx_ring) + return -EIO; + dma_dev = tb_ring_dma_device(rx_path->rx_ring); + if (!tbv_dma_device_ready(dma_dev)) + return -EIO; + mr = tbv_mr_get(state, msg->rkey); + if (!mr) + return -EINVAL; + if (!mr->umem) { + tbv_mr_put(mr); + return -EINVAL; + } + sgt = &mr->umem->sgt_append.sgt; + n = dma_map_sg_attrs(dma_dev, sgt->sgl, sgt->orig_nents, + DMA_FROM_DEVICE, 0); + if (n <= 0) { + tbv_mr_put(mr); + return -EIO; + } + msg->zcopy_mr = mr; + msg->zcopy_sgt = sgt; + msg->zcopy_dma_dev = dma_dev; + msg->zcopy_sgt_nents = n; + msg->zcopy_mapped = true; + return 0; +} + +static int tbv_rx_zcopy_post_fragment(struct tbv_qp *tqp, + struct tbv_path *rx_path, + struct tbv_rx_reorder_msg *msg, + u32 frag_idx, u32 frag_offset, + u32 frag_len, u32 psn) +{ + struct sg_table *sgt = msg->zcopy_sgt; + struct scatterlist *sg = NULL; + unsigned int i; + u64 sg_iova_base; + u64 dest_iova; + struct page *page; + unsigned int page_off; + struct tbv_rx_zcopy_frag *zf; + int ret; + + if (check_add_overflow(msg->remote_addr, (u64)frag_offset, &dest_iova)) + return -EINVAL; + + /* Find the SGL entry that contains dest_iova. The map already + * arranged the SGL in place, but the per-entry IOVA is not + * stored on the SGL itself for non-DMA-bus-mapped regions. We + * rely on the wire remote_addr + frag_offset being a linear + * range within the MR, and the MR's page layout being + * page-aligned with frag_off fitting in one page. Walk the + * SGL until cumulative length >= frag_len. + */ + sg_iova_base = msg->remote_addr; + for_each_sgtable_sg(sgt, sg, i) { + size_t sg_len = sg_dma_len(sg); + + if (!sg_len) + continue; + if (dest_iova < sg_iova_base + sg_len && + dest_iova + frag_len <= sg_iova_base + sg_len) + break; + sg_iova_base += sg_len; + } + if (!sg) + return -EINVAL; + + page = sg_page(sg); + page_off = (unsigned int)(dest_iova & (PAGE_SIZE - 1)); + if (page_off + frag_len > PAGE_SIZE) + return -EINVAL; + + zf = kzalloc(sizeof(*zf), GFP_KERNEL); + if (!zf) + return -ENOMEM; + INIT_WORK(&zf->work, tbv_rx_zcopy_drain_work); + zf->tqp = tqp; + zf->rx_path = rx_path; + zf->msg = msg; + zf->frag_idx = frag_idx; + tbv_rx_reorder_msg_get(msg); + + ret = tbv_path_post_rx_zcopy_frame(rx_path, page, page_off, frag_len, + zf); + if (ret) { + tbv_rx_reorder_msg_put(msg); + kfree(zf); + return ret; + } + /* + * Account the bytes now (in submit order); the DMA completion + * bumps frags_received / runs drain. If the DMA later errors + * we still want the byte accounting to be right for error logs. + */ + msg->received += frag_len; + return 0; +} + + static void tbv_rx_buffer_write_fragment_locked( struct tbv_state *state, struct tbv_qp *tqp, struct tbv_path *rx_path, const struct tbv_native_data_header *hdr, u32 psn, u32 total_len, @@ -7742,6 +8183,7 @@ static void tbv_rx_buffer_write_fragment_locked( } INIT_LIST_HEAD(&msg->frags); + refcount_set(&msg->refs, 1); msg->first_jiffies = jiffies; msg->kind = TBV_RX_REORDER_WRITE; msg->remote_addr = hdr->remote_addr; @@ -7757,6 +8199,11 @@ static void tbv_rx_buffer_write_fragment_locked( tqp->rx_reorder_count++; atomic64_inc(&state->data_rx_reorder_buffered); tbv_qp_schedule_timeout(tqp); + + if (last) { + tbv_rx_drain_reorder_locked(state, tqp, rx_path); + return; + } } else if (msg->kind != TBV_RX_REORDER_WRITE || msg->src_qp != hdr->src_qp || msg->remote_addr != hdr->remote_addr || @@ -7770,17 +8217,32 @@ static void tbv_rx_buffer_write_fragment_locked( psn, TBV_NATIVE_SEND_ACK_ERROR); return; } else if (test_bit(frag_idx, msg->frag_seen)) { - if (!tbv_rx_reorder_fragment_matches_locked(msg, offset, - payload, - hdr->length)) { + if (msg->complete) + tbv_rx_drain_reorder_locked(state, tqp, rx_path); + return; + } + + if (msg->zcopy_mapped || tbv_rx_zcopy_try_lock(state, tqp, msg)) { + if (!msg->zcopy_mapped) { + ret = tbv_rx_zcopy_prepare_msg(state, rx_path, tqp, msg); + if (ret) { + tbv_rx_drop_reorder_msg_locked(state, tqp, msg); + tbv_send_ack_on_path(tqp, rx_path, hdr->src_qp, + hdr->dest_qp, psn, + TBV_NATIVE_SEND_ACK_ERROR); + return; + } + } + ret = tbv_rx_zcopy_post_fragment(tqp, rx_path, msg, frag_idx, + offset, hdr->length, psn); + if (ret) { tbv_rx_drop_reorder_msg_locked(state, tqp, msg); + atomic64_inc(&state->data_rx_dmabuf_zcopy_error); tbv_send_ack_on_path(tqp, rx_path, hdr->src_qp, hdr->dest_qp, psn, TBV_NATIVE_SEND_ACK_ERROR); return; } - if (msg->complete) - tbv_rx_drain_reorder_locked(state, tqp, rx_path); return; } @@ -7803,6 +8265,7 @@ static void tbv_rx_buffer_write_fragment_locked( tbv_rx_drain_reorder_locked(state, tqp, rx_path); } + static void tbv_rx_handle_send_fragment(struct tbv_state *state, struct tbv_qp *tqp, const struct tbv_native_data_header *hdr, @@ -9146,35 +9609,21 @@ static struct ib_mr *tbv_get_dma_mr(struct ib_pd *pd, int access) return &mr->base; } -static struct ib_mr *tbv_reg_user_mr(struct ib_pd *pd, u64 start, u64 length, - u64 virt_addr, int access, -#if LINUX_VERSION_CODE >= KERNEL_VERSION(6, 13, 0) - struct ib_dmah *dmah, -#endif - struct ib_udata *udata) +static struct ib_mr *tbv_reg_mr_from_umem(struct ib_pd *pd, + struct ib_umem *umem, u64 start, + u64 length, u64 virt_addr, + int access) { struct tbv_mr *mr; int ret; - if (!length) - return ERR_PTR(-EINVAL); - mr = kzalloc(sizeof(*mr), GFP_KERNEL); - if (!mr) + if (!mr) { + ib_umem_release(umem); return ERR_PTR(-ENOMEM); - -#if LINUX_VERSION_CODE >= KERNEL_VERSION(7, 2, 0) - mr->umem = ib_umem_get_va(pd->device, start, length, access); -#else - mr->umem = ib_umem_get(pd->device, start, length, access); -#endif - if (IS_ERR(mr->umem)) { - struct ib_umem *umem = mr->umem; - - kfree(mr); - return ERR_CAST(umem); } + mr->umem = umem; mr->base.type = IB_MR_TYPE_USER; mr->base.iova = virt_addr; mr->base.length = length; @@ -9191,6 +9640,69 @@ static struct ib_mr *tbv_reg_user_mr(struct ib_pd *pd, u64 start, u64 length, return &mr->base; } +static struct ib_mr *tbv_reg_user_mr(struct ib_pd *pd, u64 start, u64 length, + u64 virt_addr, int access, +#if LINUX_VERSION_CODE >= KERNEL_VERSION(6, 13, 0) + struct ib_dmah *dmah, +#endif + struct ib_udata *udata) +{ + struct ib_umem *umem; + + if (!length) + return ERR_PTR(-EINVAL); + +#if LINUX_VERSION_CODE >= KERNEL_VERSION(7, 2, 0) + umem = ib_umem_get_va(pd->device, start, length, access); +#else + umem = ib_umem_get(pd->device, start, length, access); +#endif + if (IS_ERR(umem)) + return ERR_CAST(umem); + + return tbv_reg_mr_from_umem(pd, umem, start, length, virt_addr, + access); +} + +static struct ib_mr *tbv_reg_user_mr_dmabuf(struct ib_pd *pd, u64 offset, + u64 length, u64 virt_addr, + int fd, int access, +#if LINUX_VERSION_CODE >= KERNEL_VERSION(6, 13, 0) + struct ib_dmah *dmah, +#endif + struct uverbs_attr_bundle *attrs) +{ + struct ib_umem_dmabuf *umem_dmabuf; + int ret; + + if (!length) + return ERR_PTR(-EINVAL); + + pr_info_ratelimited("dmabuf MR dynamic import device=%s dma_device=%s fd=%d offset=%llu length=%llu\n", + dev_name(&pd->device->dev), + pd->device->dma_device ? + dev_name(pd->device->dma_device) : "", + fd, (unsigned long long)offset, + (unsigned long long)length); + umem_dmabuf = ib_umem_dmabuf_get(pd->device, offset, length, fd, + access, &tbv_dmabuf_attach_ops); + if (IS_ERR(umem_dmabuf)) + return ERR_CAST(umem_dmabuf); + + dma_resv_lock(umem_dmabuf->attach->dmabuf->resv, NULL); + ret = ib_umem_dmabuf_map_pages(umem_dmabuf); + dma_resv_unlock(umem_dmabuf->attach->dmabuf->resv); + if (ret) { + pr_info_ratelimited("dmabuf MR map_pages failed err=%d\n", + ret); + ib_umem_release(&umem_dmabuf->umem); + return ERR_PTR(ret); + } + + return tbv_reg_mr_from_umem(pd, &umem_dmabuf->umem, virt_addr, + length, virt_addr, access); +} + static int tbv_dereg_mr(struct ib_mr *ibmr, struct ib_udata *udata) { struct tbv_mr *mr = container_of(ibmr, struct tbv_mr, base); @@ -9242,6 +9754,7 @@ static const struct ib_device_ops tbv_ibdev_ops = { .poll_cq = tbv_poll_cq, .req_notify_cq = tbv_req_notify_cq, .reg_user_mr = tbv_reg_user_mr, + .reg_user_mr_dmabuf = tbv_reg_user_mr_dmabuf, .dereg_mr = tbv_dereg_mr, INIT_RDMA_OBJ_SIZE(ib_ucontext, tbv_ucontext, base), @@ -9289,6 +9802,8 @@ static int tbv_ibdev_register_one(struct tbv_state *state, BIT_ULL(IB_USER_VERBS_CMD_POST_RECV) | BIT_ULL(IB_USER_VERBS_CMD_POLL_CQ) | BIT_ULL(IB_USER_VERBS_CMD_REQ_NOTIFY_CQ); + if (IS_ENABLED(CONFIG_INFINIBAND_USER_ACCESS)) + dev->base.driver_def = tbv_uapi_defs; ib_set_device_ops(&dev->base, &tbv_ibdev_ops); @@ -9312,13 +9827,11 @@ static int tbv_ibdev_register_one(struct tbv_state *state, } /* - * This is a software verbs provider. User MRs remain pinned by - * the RDMA umem helpers, and kernel local-DMA SGEs must stay - * CPU-visible for MAD/CM QP1 handling. Keep the Thunderbolt ring - * device as the parent, but use RDMA-core virtual DMA for verbs - * buffer addresses. + * This is still a software verbs provider for ordinary MRs, but the + * experimental dma-buf path needs RDMA core to attach imports against + * the actual Thunderbolt ring DMA device instead of virtual DMA. */ - ret = ib_register_device(&dev->base, name, NULL); + ret = ib_register_device(&dev->base, name, dma_device); if (ret) { tbv_ibdev_detach_netdev(dev); put_device(dma_device); diff --git a/kernel/path.c b/kernel/path.c index 11a4d53..3cc612e 100644 --- a/kernel/path.c +++ b/kernel/path.c @@ -95,6 +95,13 @@ struct tbv_data_frame { tbv_path_tx_done_fn done; void *done_ctx; bool tx; + /* zcopy_rx-only: one-shot frame that DMAs the wire payload + * directly into an MR page, no kernel bounce buffer. Set on + * the rx_path branch in tbv_path_post_rx_zcopy_frame(). + */ + bool zcopy_rx; + struct page *zcopy_page; + unsigned int zcopy_page_off; }; struct tbv_tx_packet { @@ -751,6 +758,9 @@ static void tbv_path_rx_supp_poll_work(struct work_struct *work) } static int tbv_path_post_rx_frame(struct tbv_data_frame *f); +static void tbv_path_zcopy_rx_complete(struct tb_ring *ring, + struct ring_frame *frame, + bool canceled); static void tbv_path_rx_start_raw(struct tbv_path *path, const struct tbv_native_data_header *hdr) @@ -1157,6 +1167,84 @@ static int tbv_path_post_rx_frame(struct tbv_data_frame *f) return tb_ring_rx(path->rx_ring, &f->frame); } +int tbv_path_post_rx_zcopy_frame(struct tbv_path *path, + struct page *page, unsigned int page_off, + u32 len, void *done_ctx) +{ + struct tbv_data_frame *f; + struct device *dma_dev; + dma_addr_t dma; + + if (!path || !path->rx_ring || !page) + return -EINVAL; + + dma_dev = tb_ring_dma_device(path->rx_ring); + if (!tbv_dma_device_ready(dma_dev)) + return -EIO; + + if (page_off + len > PAGE_SIZE) + return -EINVAL; + + f = kzalloc(sizeof(*f), GFP_KERNEL); + if (!f) + return -ENOMEM; + + INIT_LIST_HEAD(&f->frame.list); + INIT_LIST_HEAD(&f->free_node); + f->path = path; + f->tx = false; + f->zcopy_rx = true; + f->buf = NULL; + f->zcopy_page = page; + f->zcopy_page_off = page_off; + f->done = tbv_rx_zcopy_complete; + f->done_ctx = done_ctx; + + dma = dma_map_page(dma_dev, page, page_off, len, DMA_FROM_DEVICE); + if (dma_mapping_error(dma_dev, dma)) { + kfree(f); + return -EIO; + } + f->dma = dma; + f->frame.buffer_phy = dma; + f->frame.size = len; + f->frame.flags = 0; + f->frame.sof = 0; + f->frame.eof = 0; + f->frame.callback = tbv_path_zcopy_rx_complete; + + dma_sync_single_for_device(dma_dev, dma, len, DMA_FROM_DEVICE); + return tb_ring_rx(path->rx_ring, &f->frame); +} + +static void tbv_path_zcopy_rx_complete(struct tb_ring *ring, + struct ring_frame *frame, + bool canceled) +{ + struct tbv_data_frame *f = container_of(frame, struct tbv_data_frame, + frame); + struct tbv_path *path = f->path; + struct device *dma_dev = tb_ring_dma_device(ring); + int status = canceled ? -ECANCELED : 0; + + if (tbv_dma_device_ready(dma_dev)) + dma_unmap_page(dma_dev, f->dma, f->frame.size, DMA_FROM_DEVICE); + + if (path && path->rail && path->rail->peer) { + struct tbv_state *state = path->rail->peer->state; + + if (state) { + atomic64_inc(&state->data_rx_dmabuf_zcopy); + if (status) + atomic64_inc(&state->data_rx_dmabuf_zcopy_error); + } + } + + if (f->done) + f->done(f->done_ctx, status); + kfree(f); +} + const char *tbv_path_state_name(enum tbv_path_state state) { switch (state) { @@ -2407,6 +2495,131 @@ int tbv_path_send_page_stream(struct tbv_path *path, return ret; } +int tbv_path_send_dma_stream(struct tbv_path *path, + const struct tbv_native_data_header *hdr, + u32 total_length, unsigned int send_flags, + tbv_path_tx_done_fn meta_done, + void *meta_done_ctx, + tbv_path_next_dma_fn next, void *next_ctx) +{ + LIST_HEAD(packets); + u32 prepared = 0; + u32 packet_count = 0; + struct tbv_tx_packet *packet; + u32 max_raw_payload; + struct tbv_native_data_header stream_hdr; + u8 *hdr_buf; + int ret; + + if (!path || !hdr || !next || !total_length) { + ret = -EINVAL; + goto err_meta_done; + } + if (send_flags & ~(TBV_PATH_SEND_DEFER)) { + ret = -EINVAL; + goto err_meta_done; + } + if (total_length > TBV_NATIVE_DATA_MAX_MSG_SIZE) { + ret = -EMSGSIZE; + goto err_meta_done; + } + if (!path->tx_ring) { + ret = -ENOTCONN; + goto err_meta_done; + } + if (hdr->opcode == TBV_NATIVE_DATA_OP_RDMA_WRITE || + hdr->opcode == TBV_NATIVE_DATA_OP_RDMA_WRITE_IMM) + max_raw_payload = TBV_DATA_FRAME_SIZE; + else + max_raw_payload = TBV_NATIVE_DATA_MAX_PAYLOAD; + + if (!tbv_dma_device_ready(tb_ring_dma_device(path->tx_ring))) { + ret = -EPROBE_DEFER; + goto err_meta_done; + } + + hdr_buf = kzalloc(TBV_NATIVE_DATA_HDR_SIZE, GFP_KERNEL); + if (!hdr_buf) { + ret = -ENOMEM; + goto err_release; + } + + stream_hdr = *hdr; + stream_hdr.length = total_length; + stream_hdr.flags |= TBV_NATIVE_DATA_F_LAST | + TBV_NATIVE_DATA_F_RAW_STREAM; + ret = tbv_native_data_build_header(hdr_buf, TBV_NATIVE_DATA_HDR_SIZE, + &stream_hdr); + if (ret < 0) { + kfree(hdr_buf); + goto err_release; + } + + packet = tbv_path_alloc_data_packet_owned(path, hdr_buf, + TBV_NATIVE_DATA_HDR_SIZE, + meta_done, meta_done_ctx); + if (!packet) { + kfree(hdr_buf); + ret = -ENOMEM; + goto err_release; + } + packet->raw_stream_start = true; + list_add_tail(&packet->node, &packets); + packet_count++; + + while (prepared < total_length) { + tbv_path_tx_done_fn done = NULL; + void *done_ctx = NULL; + dma_addr_t dma = 0; + bool last; + u32 len = 0; + + ret = next(next_ctx, &dma, &len, &done, &done_ctx); + if (ret) + goto err_release; + if (!len || len > max_raw_payload || + len > total_length - prepared) { + if (done) + done(done_ctx, -EINVAL); + ret = -EINVAL; + goto err_release; + } + + last = prepared + len == total_length; + + packet = tbv_path_alloc_zcopy_packet(path, dma, len, false, + done, done_ctx); + if (!packet) { + if (done) + done(done_ctx, -ENOMEM); + ret = -ENOMEM; + goto err_release; + } + packet->owner_ctx = meta_done_ctx ? meta_done_ctx : done_ctx; + packet->raw_stream_end = last; + list_add_tail(&packet->node, &packets); + packet_count++; + prepared += len; + } + + ret = tbv_path_enqueue_data_list(path, &packets, packet_count, + send_flags & TBV_PATH_SEND_DEFER); + if (ret) + goto err_release; + return 0; + +err_release: + if (!packet_count && meta_done) + meta_done(meta_done_ctx, ret); + tbv_path_release_packet_list(&packets, ret); + return ret; + +err_meta_done: + if (meta_done) + meta_done(meta_done_ctx, ret); + return ret; +} + static bool tbv_path_packet_matches(const struct tbv_tx_packet *packet, tbv_path_tx_done_fn done, void *done_ctx, void *owner_ctx) diff --git a/kernel/tbv.h b/kernel/tbv.h index 59368c9..1e2bf3d 100644 --- a/kernel/tbv.h +++ b/kernel/tbv.h @@ -573,6 +573,8 @@ struct tbv_state { atomic64_t data_rx_rnr; atomic64_t data_rx_rnr_suppressed; atomic64_t data_rx_copy_error; + atomic64_t data_rx_dmabuf_zcopy; + atomic64_t data_rx_dmabuf_zcopy_error; atomic64_t data_rx_send_len_error; atomic64_t data_rx_send_prot_error; atomic64_t data_rx_send_cq_error; @@ -645,6 +647,9 @@ typedef int (*tbv_path_next_page_fn)(void *ctx, struct page **page, u32 *page_off, u32 *length, tbv_path_tx_done_fn *done, void **done_ctx); +typedef int (*tbv_path_next_dma_fn)(void *ctx, dma_addr_t *dma, u32 *length, + tbv_path_tx_done_fn *done, + void **done_ctx); #define TBV_PATH_SEND_CONTROL BIT(0) #define TBV_PATH_SEND_DEFER BIT(1) extern const uuid_t tbv_native_service_uuid; @@ -663,6 +668,45 @@ const char *tbv_backend_name(enum tbv_backend_type type); int tbv_ibdev_start(struct tbv_state *state, bool register_verbs); void tbv_ibdev_stop(struct tbv_state *state); +struct ib_device; +struct ib_qp; +struct ib_umem; +struct tbv_qp; +struct tbv_state *tbv_ibdev_state(struct ib_device *ibdev); + +/* + * USB4 RDMA Direct Verbs (DV) per-QP state. + * + * Each QP owns one of these. CREATE_QUEUE attaches userspace-provided + * SQ/CQ/doorbell memory; DESTROY_QUEUE (or QP destroy) bumps the + * generation and releases it. mutex serializes attach/detach against + * each other and against QP teardown; KICK drains only NOP smoke-test WQEs + * until the transport consumer lands. + */ +struct tbv_dv_qp_state { + struct mutex mutex; + bool active; + u8 generation; + u32 sq_entries; + u32 cq_entries; + u32 sq_head; + u32 cq_tail; + u64 sq_addr; + u64 cq_addr; + u64 doorbell_addr; + struct ib_umem *sq_umem; + struct ib_umem *cq_umem; + struct ib_umem *doorbell_umem; +}; + +void tbv_dv_qp_state_init(struct tbv_dv_qp_state *dv); +void tbv_dv_qp_state_teardown(struct tbv_dv_qp_state *dv); +bool tbv_dv_qp_state_active(struct tbv_dv_qp_state *dv); + +/* Helpers used by dv.c to reach the per-QP state without exposing the + * full tbv_qp layout outside ibdev.c. */ +struct tbv_qp *tbv_qp_from_ibqp(struct ib_qp *ibqp); +struct tbv_dv_qp_state *tbv_qp_dv_state(struct tbv_qp *tqp); const char *tbv_ibdev_roce_netdev_name(void); /* * Notify the verbs layer that rail's data path has come up (joined=true) or @@ -816,6 +860,12 @@ int tbv_path_send_page_stream(struct tbv_path *path, tbv_path_tx_done_fn meta_done, void *meta_done_ctx, tbv_path_next_page_fn next, void *next_ctx); +int tbv_path_send_dma_stream(struct tbv_path *path, + const struct tbv_native_data_header *hdr, + u32 total_length, unsigned int flags, + tbv_path_tx_done_fn meta_done, + void *meta_done_ctx, + tbv_path_next_dma_fn next, void *next_ctx); void tbv_path_kick_tx(struct tbv_path *path); void tbv_path_cancel_data_done_ctx(struct tbv_path *path, tbv_path_tx_done_fn done, void *done_ctx); @@ -838,5 +888,9 @@ int tbv_core_init(struct tbv_state *state, const struct tbv_resolved_config *cfg, const struct tbv_tbnet_identity_config *identity_cfg); void tbv_core_exit(struct tbv_state *state); +void tbv_rx_zcopy_complete(void *ctx, int status); +int tbv_path_post_rx_zcopy_frame(struct tbv_path *path, + struct page *page, unsigned int page_off, + u32 len, void *done_ctx); #endif diff --git a/nix/bench-tools.nix b/nix/bench-tools.nix index e3c7723..05a6bfe 100644 --- a/nix/bench-tools.nix +++ b/nix/bench-tools.nix @@ -25,6 +25,7 @@ let # Full Linux set. ibv_trace is an LD_PRELOAD tracer built as .so. linuxPrograms = darwinPrograms ++ [ + "dmabuf_mr_probe" "rc_qpn_churn" "rdma_gid_probe" ]; diff --git a/nix/dv-probes.nix b/nix/dv-probes.nix new file mode 100644 index 0000000..0107d87 --- /dev/null +++ b/nix/dv-probes.nix @@ -0,0 +1,71 @@ +{ lib +, stdenv +, pkg-config +, rdma-core-usb4 +, source ? ../. +}: + +# Probes that exercise the USB4 RDMA Direct Verbs (DV / GDA) ABI. +# +# Each probe is a small standalone tool that talks to the kernel module via +# the raw RDMA_VERBS_IOCTL ABI plus the shared ABI header at +# userspace/usb4_rdma/usb4_rdma_dv.h, without depending on rdma-core's +# private execute_ioctl() helper. This keeps the probes buildable as +# ordinary userspace binaries. + +let + probes = [ + "tbv_dv_caps_probe" + ]; +in +stdenv.mkDerivation { + pname = "thunderbolt-ibverbs-dv-probes"; + version = "0.1.0"; + + src = lib.cleanSourceWith { + src = source; + filter = path: type: + let + rel = lib.removePrefix (toString source + "/") (toString path); + in + type == "directory" + || rel == "userspace/usb4_rdma/usb4_rdma_dv.h" + || (lib.hasPrefix "userspace/bench/tbv_dv_" rel + && lib.hasSuffix ".c" rel); + }; + + nativeBuildInputs = [ pkg-config ]; + buildInputs = [ rdma-core-usb4 ]; + + dontConfigure = true; + + buildPhase = '' + runHook preBuild + for name in ${lib.concatStringsSep " " probes}; do + $CC -O2 -Wall -Wextra -std=gnu11 \ + -I${rdma-core-usb4.dev}/include \ + -Iuserspace/usb4_rdma \ + "userspace/bench/$name.c" \ + -L${rdma-core-usb4}/lib -libverbs \ + -Wl,-rpath,${rdma-core-usb4}/lib \ + -o "$(echo "$name" | tr '_' '-')" + done + runHook postBuild + ''; + + installPhase = '' + runHook preInstall + mkdir -p $out/bin + for name in ${lib.concatStringsSep " " probes}; do + install -m 0755 "$(echo "$name" | tr '_' '-')" "$out/bin/" + done + runHook postInstall + ''; + + meta = with lib; { + description = "Thunderbolt/USB4 RDMA Direct Verbs (GDA) probes"; + license = with licenses; [ gpl2Only bsd3 ]; + platforms = platforms.linux; + mainProgram = "tbv-dv-caps-probe"; + }; +} diff --git a/nix/module.nix b/nix/module.nix index 175e7aa..eea52eb 100644 --- a/nix/module.nix +++ b/nix/module.nix @@ -23,7 +23,13 @@ stdenv.mkDerivation { || rel == "kernel" || lib.hasPrefix "kernel/" rel || rel == "proto" - || lib.hasPrefix "proto/" rel; + || lib.hasPrefix "proto/" rel + || rel == "userspace" + || rel == "userspace/usb4_rdma" + # The DV ABI header is shared between kernel and userspace consumers. + # Pull just the header into the kernel build; the rest of + # userspace/ remains out of scope for the module derivation. + || rel == "userspace/usb4_rdma/usb4_rdma_dv.h"; }; nativeBuildInputs = kernel.moduleBuildDependencies; diff --git a/tools/ci/distro-package.sh b/tools/ci/distro-package.sh index 33bd215..d605d44 100755 --- a/tools/ci/distro-package.sh +++ b/tools/ci/distro-package.sh @@ -75,6 +75,7 @@ stage_source() { README.md \ kernel \ proto \ + userspace/usb4_rdma/usb4_rdma_dv.h \ | tar -C "$stage" -xf - } diff --git a/userspace/bench/GDA_PROBES.md b/userspace/bench/GDA_PROBES.md new file mode 100644 index 0000000..e880199 --- /dev/null +++ b/userspace/bench/GDA_PROBES.md @@ -0,0 +1,25 @@ +# USB4 GDA Smoke Probes + +These probes separate two questions that are easy to conflate. + +`tbv-dv-caps-probe -q` validates the DV queue-memory contract on one host: +the kernel can pin CPU-visible SQ/CQ/doorbell memory, consume signaled NOP +WQEs through `KICK`, write CQEs, and publish doorbell head/tail updates. This +does not prove Thunderbolt/NHI payload DMA into GPU VRAM. + +`hip_dv_kernel_cqe_probe.cpp` validates GPU visibility of kernel-written DV +completion state. Build it on a ROCm host and run it against a loaded module +after `tbv-dv-caps-probe -q` passes. + +`hip_reg_mr_probe.cpp`, `hsa_fine_grain_pool_probe.cpp`, and +`hip_rdma_write_visibility_probe.cpp` are payload-memory probes. The two-host +visibility probe is the closest "true GPU DMA" smoke: the sender writes a +payload and signal through RDMA WRITE into a HIP allocation, then a GPU kernel +on the receiver observes the signal and checks the payload. A pass demonstrates +the tested allocation mode and synchronization path on that topology; it does +not imply generic PCIe P2PDMA support for arbitrary GPU memory. + +Expected conservative result for this branch: queue-memory probes can pass +with host-visible coherent memory, while arbitrary device-memory payloads may +fail registration or fall back to staged CPU-visible paths until the kernel has +an explicit peer-direct/P2PDMA contract. diff --git a/userspace/bench/dmabuf_mr_probe.c b/userspace/bench/dmabuf_mr_probe.c new file mode 100644 index 0000000..c3b5f9a --- /dev/null +++ b/userspace/bench/dmabuf_mr_probe.c @@ -0,0 +1,203 @@ +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include +#include +#include +#include +#include + +static const char *errno_name(int err) +{ + switch (err) { + case 0: + return "OK"; + case EBADF: + return "EBADF"; + case EINVAL: + return "EINVAL"; + case ENODEV: + return "ENODEV"; + case ENOMEM: + return "ENOMEM"; + case EOPNOTSUPP: + return "EOPNOTSUPP"; +#ifdef EPROTONOSUPPORT + case EPROTONOSUPPORT: + return "EPROTONOSUPPORT"; +#endif + default: + return "UNKNOWN"; + } +} + +static int parse_errno(const char *value) +{ + char *end = NULL; + long parsed; + + if (!strcmp(value, "OK")) + return 0; + if (!strcmp(value, "EBADF")) + return EBADF; + if (!strcmp(value, "EINVAL")) + return EINVAL; + if (!strcmp(value, "ENODEV")) + return ENODEV; + if (!strcmp(value, "ENOMEM")) + return ENOMEM; + if (!strcmp(value, "EOPNOTSUPP")) + return EOPNOTSUPP; +#ifdef EPROTONOSUPPORT + if (!strcmp(value, "EPROTONOSUPPORT")) + return EPROTONOSUPPORT; +#endif + + errno = 0; + parsed = strtol(value, &end, 0); + if (!errno && end && *end == '\0' && parsed >= 0 && parsed <= 4095) + return (int)parsed; + + return -1; +} + +static void usage(const char *argv0) +{ + fprintf(stderr, + "Usage: %s [--device HCA] [--fd FD] [--length BYTES] [--iova ADDR] [--expect-errno NAME|NUM]\n", + argv0); +} + +int main(int argc, char **argv) +{ + const char *want_device = NULL; + const char *expect_value = NULL; + struct ibv_device **dev_list; + struct ibv_device *dev = NULL; + struct ibv_context *ctx = NULL; + struct ibv_pd *pd = NULL; + struct ibv_mr *mr = NULL; + uint64_t iova = 0x100000000ULL; + size_t length = 4096; + int expect_errno = -1; + int num_devices = 0; + int fd = -1; + int access; + int err = 0; + int ret = 1; + int i; + + for (i = 1; i < argc; i++) { + char *end = NULL; + + if (!strcmp(argv[i], "--device") && i + 1 < argc) { + want_device = argv[++i]; + } else if (!strcmp(argv[i], "--fd") && i + 1 < argc) { + fd = (int)strtol(argv[++i], &end, 0); + if (!end || *end) { + usage(argv[0]); + return 2; + } + } else if (!strcmp(argv[i], "--length") && i + 1 < argc) { + unsigned long long value; + + value = strtoull(argv[++i], &end, 0); + if (!end || *end) { + usage(argv[0]); + return 2; + } + length = (size_t)value; + } else if (!strcmp(argv[i], "--iova") && i + 1 < argc) { + iova = strtoull(argv[++i], &end, 0); + if (!end || *end) { + usage(argv[0]); + return 2; + } + } else if (!strcmp(argv[i], "--expect-errno") && i + 1 < argc) { + expect_value = argv[++i]; + expect_errno = parse_errno(expect_value); + if (expect_errno < 0) { + fprintf(stderr, "unknown errno expectation: %s\n", + expect_value); + return 2; + } + } else { + usage(argv[0]); + return 2; + } + } + + dev_list = ibv_get_device_list(&num_devices); + if (!dev_list) { + err = errno; + fprintf(stderr, "ibv_get_device_list failed: errno=%d %s %s\n", + err, errno_name(err), strerror(err)); + return 1; + } + + for (i = 0; i < num_devices; i++) { + const char *name = ibv_get_device_name(dev_list[i]); + + if ((want_device && !strcmp(name, want_device)) || + (!want_device && !strncmp(name, "usb4_rdma", 9))) { + dev = dev_list[i]; + break; + } + } + if (!dev) { + fprintf(stderr, "no matching RDMA device found"); + if (want_device) + fprintf(stderr, ": %s", want_device); + fprintf(stderr, "\n"); + goto out_free_list; + } + + ctx = ibv_open_device(dev); + if (!ctx) { + err = errno; + fprintf(stderr, "ibv_open_device(%s) failed: errno=%d %s %s\n", + ibv_get_device_name(dev), err, errno_name(err), + strerror(err)); + goto out_free_list; + } + + pd = ibv_alloc_pd(ctx); + if (!pd) { + err = errno; + fprintf(stderr, "ibv_alloc_pd(%s) failed: errno=%d %s %s\n", + ibv_get_device_name(dev), err, errno_name(err), + strerror(err)); + goto out_close; + } + + access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_READ | + IBV_ACCESS_REMOTE_WRITE; + errno = 0; + mr = ibv_reg_dmabuf_mr(pd, 0, length, iova, fd, access); + if (mr) { + printf("device=%s fd=%d length=%zu iova=0x%" PRIx64 " result=success lkey=0x%x rkey=0x%x\n", + ibv_get_device_name(dev), fd, length, iova, mr->lkey, + mr->rkey); + err = 0; + ibv_dereg_mr(mr); + } else { + err = errno; + printf("device=%s fd=%d length=%zu iova=0x%" PRIx64 " result=error errno=%d name=%s message=%s\n", + ibv_get_device_name(dev), fd, length, iova, err, + errno_name(err), strerror(err)); + } + + if (expect_errno >= 0) + ret = err == expect_errno ? 0 : 1; + else + ret = err ? 1 : 0; + + ibv_dealloc_pd(pd); +out_close: + ibv_close_device(ctx); +out_free_list: + ibv_free_device_list(dev_list); + return ret; +} diff --git a/userspace/bench/hip_dv_kernel_cqe_probe.cpp b/userspace/bench/hip_dv_kernel_cqe_probe.cpp new file mode 100644 index 0000000..61a19be --- /dev/null +++ b/userspace/bench/hip_dv_kernel_cqe_probe.cpp @@ -0,0 +1,632 @@ +// SPDX-License-Identifier: GPL-2.0 OR BSD-3-Clause +/* + * Probe the kernel->GPU direct-load visibility needed by USB4 GDA sync words. + * + * This intentionally tests the exact awkward direction that ordinary userspace + * CPU/GPU coherency probes miss: + * + * kernel writes a DV completion tail word in userspace queue memory + * GPU kernel spins on that word using a plain direct load + * + * If a host-coherent allocation passes this probe, it is a viable candidate + * for rocSHMEM pSync/signal control memory. Payload memory can still be a + * separate dmabuf/device allocation. + */ + +#define _POSIX_C_SOURCE 200112L + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include "usb4_rdma_dv.h" + +#define PROBE_WR_ID_BASE 0x4b43514550524f42ull /* KCQEPROB */ + +struct alignas(64) ProbeState { + uint32_t gpu_seen; + uint32_t gpu_error; + uint32_t gpu_last_tail; + uint32_t reserved; + uint64_t gpu_spins; + uint8_t pad[64 - 4 * sizeof(uint32_t) - sizeof(uint64_t)]; +}; + +static_assert(sizeof(ProbeState) == 64, "ProbeState must stay one cacheline"); + +struct QueueMem { + const char *kind; + void *base; + size_t bytes; + struct usb4_rdma_dv_wqe *sq; + struct usb4_rdma_dv_cqe *cq; + struct usb4_rdma_dv_doorbell *doorbell; + ProbeState *state; + size_t sq_bytes; + size_t cq_bytes; + size_t doorbell_bytes; +}; + +static uint64_t now_ns(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (uint64_t)ts.tv_sec * 1000000000ull + (uint64_t)ts.tv_nsec; +} + +static uintptr_t align_up(uintptr_t value, size_t alignment) +{ + return (value + alignment - 1) & ~(uintptr_t)(alignment - 1); +} + +static uint32_t host_load_acquire_u32(uint32_t *ptr) +{ + return __atomic_load_n(ptr, __ATOMIC_ACQUIRE); +} + +static void host_store_release_u32(uint32_t *ptr, uint32_t value) +{ + __atomic_store_n(ptr, value, __ATOMIC_RELEASE); +} + +__device__ static inline uint32_t gpu_plain_load_u32(const uint32_t *ptr) +{ + const volatile uint32_t *vptr = ptr; + + return *vptr; +} + +__device__ static inline void gpu_store_release_u32(uint32_t *ptr, + uint32_t value) +{ + __hip_atomic_store(ptr, value, __ATOMIC_RELEASE, + __HIP_MEMORY_SCOPE_SYSTEM); +} + +__device__ static inline uint32_t gpu_tail_pack(uint32_t index, + uint32_t generation) +{ + return (index & USB4_RDMA_DV_TAIL_INDEX_MASK) | + (generation << USB4_RDMA_DV_TAIL_GEN_SHIFT); +} + +__global__ void kernel_cqe_plain_load_kernel( + struct usb4_rdma_dv_doorbell *doorbell, ProbeState *state, + uint32_t iterations, uint32_t generation, uint32_t max_spins) +{ + if (blockIdx.x || threadIdx.x) + return; + + for (uint32_t iter = 1; iter <= iterations; iter++) { + uint32_t want = gpu_tail_pack(iter, generation); + uint32_t tail = 0; + uint32_t spins = 0; + + do { + tail = gpu_plain_load_u32(&doorbell->consumer.cq_tail); + if (tail == want) + break; + if (++spins == max_spins) { + state->gpu_last_tail = tail; + gpu_store_release_u32(&state->gpu_error, iter); + return; + } + } while (true); + + state->gpu_last_tail = tail; + state->gpu_spins += spins; + gpu_store_release_u32(&doorbell->producer.cq_head, want); + gpu_store_release_u32(&state->gpu_seen, iter); + } +} + +static bool device_name_matches(const char *name, const char *wanted) +{ + return strcmp(name, wanted) == 0 || + strncmp(name, "usb4_rdma", strlen("usb4_rdma")) == 0 || + strncmp(name, "usb4_apple", strlen("usb4_apple")) == 0; +} + +static struct ibv_device *find_device(struct ibv_device **list, int count, + const char *wanted) +{ + struct ibv_device *fallback = NULL; + + for (int i = 0; i < count; i++) { + const char *name = ibv_get_device_name(list[i]); + + if (!strcmp(name, wanted)) + return list[i]; + if (!fallback && device_name_matches(name, wanted)) + fallback = list[i]; + } + return fallback; +} + +static int query_caps(struct ibv_context *ctx, + struct usb4_rdma_dv_query_caps_resp *resp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[1]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_QUERY_CAPS, + .num_attrs = 1, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_QUERY_CAPS_RESP, + .len = sizeof(*resp), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)resp, + }, + }, + }; + + memset(resp, 0, sizeof(*resp)); + if (ioctl(ctx->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int create_queue(struct ibv_qp *qp, + const struct usb4_rdma_dv_queue_create *req, + struct usb4_rdma_dv_queue_resp *resp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[3]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_CREATE_QUEUE, + .num_attrs = 3, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_REQ, + .len = sizeof(*req), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)req, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_RESP, + .len = sizeof(*resp), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)resp, + }, + }, + }; + + memset(resp, 0, sizeof(*resp)); + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int destroy_queue(struct ibv_qp *qp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[1]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_DESTROY_QUEUE, + .num_attrs = 1, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_DESTROY_QUEUE_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + }, + }; + + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int kick_queue(struct ibv_qp *qp, const struct usb4_rdma_dv_kick *req) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[2]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_KICK, + .num_attrs = 2, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_KICK_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_KICK_REQ, + .len = sizeof(*req), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)req, + }, + }, + }; + + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static hipError_t alloc_queue_mem(const char *kind, QueueMem *mem, + const struct usb4_rdma_dv_query_caps_resp *caps, + uint32_t sq_entries, uint32_t cq_entries) +{ + const size_t slop = caps->doorbell_page_size * 3; + unsigned int flags = hipHostMallocMapped; + uintptr_t cursor; + void *base = NULL; + hipError_t ret; + + memset(mem, 0, sizeof(*mem)); + mem->kind = kind; + mem->sq_bytes = (size_t)sq_entries * caps->wqe_size; + mem->cq_bytes = (size_t)cq_entries * caps->cqe_size; + mem->doorbell_bytes = caps->doorbell_page_size; + mem->bytes = mem->sq_bytes + mem->cq_bytes + mem->doorbell_bytes + + sizeof(ProbeState) + slop; + + if (!strcmp(kind, "managed")) { + ret = hipMallocManaged(&base, mem->bytes); + } else if (!strcmp(kind, "host")) { + ret = hipHostMalloc(&base, mem->bytes, flags); + } else if (!strcmp(kind, "host-coherent")) { + flags |= hipHostMallocCoherent; + ret = hipHostMalloc(&base, mem->bytes, flags); + } else if (!strcmp(kind, "host-noncoherent")) { + flags |= hipHostMallocNonCoherent; + ret = hipHostMalloc(&base, mem->bytes, flags); + } else if (!strcmp(kind, "host-uncached")) { + flags |= hipHostMallocUncached; + ret = hipHostMalloc(&base, mem->bytes, flags); + } else { + fprintf(stderr, + "unknown kind=%s, expected managed|host|host-coherent|host-noncoherent|host-uncached\n", + kind); + return hipErrorInvalidValue; + } + if (ret != hipSuccess) + return ret; + + memset(base, 0, mem->bytes); + cursor = (uintptr_t)base; + cursor = align_up(cursor, caps->wqe_size); + mem->sq = (struct usb4_rdma_dv_wqe *)cursor; + cursor += mem->sq_bytes; + cursor = align_up(cursor, caps->cqe_size); + mem->cq = (struct usb4_rdma_dv_cqe *)cursor; + cursor += mem->cq_bytes; + cursor = align_up(cursor, caps->doorbell_page_size); + mem->doorbell = (struct usb4_rdma_dv_doorbell *)cursor; + cursor += mem->doorbell_bytes; + cursor = align_up(cursor, 64); + mem->state = (ProbeState *)cursor; + mem->base = base; + return hipSuccess; +} + +static void free_queue_mem(QueueMem *mem) +{ + if (!mem->base) + return; + if (!strcmp(mem->kind, "managed")) + (void)hipFree(mem->base); + else + (void)hipHostFree(mem->base); +} + +static int wait_gpu_seen(ProbeState *state, uint32_t iter, uint32_t timeout_ms) +{ + uint64_t deadline = now_ns() + (uint64_t)timeout_ms * 1000000ull; + + while (now_ns() < deadline) { + uint32_t err = host_load_acquire_u32(&state->gpu_error); + uint32_t seen = host_load_acquire_u32(&state->gpu_seen); + + if (err) { + fprintf(stderr, + "gpu reported timeout/error at iter=%u last_tail=0x%08x seen=%u\n", + err, state->gpu_last_tail, seen); + return -1; + } + if (seen >= iter) + return 0; + } + + fprintf(stderr, "host timeout waiting for gpu_seen=%u got=%u error=%u\n", + iter, host_load_acquire_u32(&state->gpu_seen), + host_load_acquire_u32(&state->gpu_error)); + return -1; +} + +static int run_probe(struct ibv_context *ctx, const char *kind, + uint32_t iterations, uint32_t timeout_ms) +{ + struct usb4_rdma_dv_query_caps_resp caps; + struct usb4_rdma_dv_queue_create req = {}; + struct usb4_rdma_dv_queue_resp resp = {}; + struct ibv_qp_init_attr qp_attr = {}; + struct ibv_mr *doorbell_mr = NULL; + struct ibv_mr *cq_mr = NULL; + struct ibv_mr *sq_mr = NULL; + struct ibv_qp *qp = NULL; + struct ibv_cq *verbs_cq = NULL; + struct ibv_pd *pd = NULL; + QueueMem mem = {}; + uint32_t sq_entries; + uint32_t cq_entries; + uint64_t start_ns; + int ret = 1; + int err; + + err = query_caps(ctx, &caps); + if (err) { + fprintf(stderr, "QUERY_CAPS failed: %s (%d)\n", strerror(err), + err); + return 1; + } + if (caps.abi_version != USB4_RDMA_DV_ABI_VERSION) { + fprintf(stderr, "unsupported DV ABI kernel=%u userspace=%u\n", + caps.abi_version, USB4_RDMA_DV_ABI_VERSION); + return 1; + } + + sq_entries = caps.default_sq_entries; + cq_entries = caps.default_cq_entries; + if (iterations < cq_entries) + cq_entries = iterations < USB4_RDMA_DV_MIN_QUEUE_ENTRIES ? + USB4_RDMA_DV_MIN_QUEUE_ENTRIES : + iterations; + + hipError_t hret = alloc_queue_mem(kind, &mem, &caps, sq_entries, + cq_entries); + if (hret != hipSuccess) { + fprintf(stderr, "alloc kind=%s failed: %s\n", kind, + hipGetErrorString(hret)); + return 1; + } + + pd = ibv_alloc_pd(ctx); + if (!pd) { + fprintf(stderr, "ibv_alloc_pd: %s\n", strerror(errno)); + goto out; + } + sq_mr = ibv_reg_mr(pd, mem.sq, mem.sq_bytes, IBV_ACCESS_LOCAL_WRITE); + cq_mr = ibv_reg_mr(pd, mem.cq, mem.cq_bytes, IBV_ACCESS_LOCAL_WRITE); + doorbell_mr = ibv_reg_mr(pd, mem.doorbell, mem.doorbell_bytes, + IBV_ACCESS_LOCAL_WRITE); + if (!sq_mr || !cq_mr || !doorbell_mr) { + fprintf(stderr, + "ibv_reg_mr queue memory failed kind=%s errno=%d %s\n", + kind, errno, strerror(errno)); + goto out; + } + + verbs_cq = ibv_create_cq(ctx, cq_entries, NULL, NULL, 0); + if (!verbs_cq) { + fprintf(stderr, "ibv_create_cq: %s\n", strerror(errno)); + goto out; + } + + qp_attr.send_cq = verbs_cq; + qp_attr.recv_cq = verbs_cq; + qp_attr.qp_type = IBV_QPT_RC; + qp_attr.cap.max_send_wr = sq_entries; + qp_attr.cap.max_recv_wr = 1; + qp_attr.cap.max_send_sge = 1; + qp_attr.cap.max_recv_sge = 1; + qp = ibv_create_qp(pd, &qp_attr); + if (!qp) { + fprintf(stderr, "ibv_create_qp: %s\n", strerror(errno)); + goto out; + } + + req.abi_version = USB4_RDMA_DV_ABI_VERSION; + req.sq_addr = (uintptr_t)mem.sq; + req.cq_addr = (uintptr_t)mem.cq; + req.doorbell_addr = (uintptr_t)mem.doorbell; + req.sq_entries = sq_entries; + req.cq_entries = cq_entries; + req.sq_stride = caps.wqe_size; + req.cq_stride = caps.cqe_size; + err = create_queue(qp, &req, &resp); + if (err) { + fprintf(stderr, "CREATE_QUEUE failed: %s (%d)\n", + strerror(err), err); + goto out; + } + + host_store_release_u32(&mem.doorbell->producer.generation, + resp.generation); + host_store_release_u32(&mem.doorbell->producer.cq_head, + usb4_rdma_dv_tail_pack(0, resp.generation)); + host_store_release_u32(&mem.doorbell->producer.sq_tail, + usb4_rdma_dv_tail_pack(0, resp.generation)); + + hret = hipMemset(mem.state, 0, sizeof(*mem.state)); + if (hret == hipSuccess) + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "hipMemset state failed: %s\n", + hipGetErrorString(hret)); + goto out_destroy; + } + + hipLaunchKernelGGL(kernel_cqe_plain_load_kernel, dim3(1), dim3(1), 0, + 0, mem.doorbell, mem.state, iterations, + resp.generation, 100000000u); + hret = hipGetLastError(); + if (hret != hipSuccess) { + fprintf(stderr, "kernel launch failed: %s\n", + hipGetErrorString(hret)); + goto out_destroy; + } + + start_ns = now_ns(); + for (uint32_t iter = 1; iter <= iterations; iter++) { + uint32_t slot = (iter - 1) % sq_entries; + uint32_t tail = usb4_rdma_dv_tail_pack(iter, resp.generation); + struct usb4_rdma_dv_kick kick = { + .sq_tail = tail, + }; + struct usb4_rdma_dv_wqe *wqe = &mem.sq[slot]; + + memset(wqe, 0, sizeof(*wqe)); + wqe->opcode = USB4_RDMA_DV_WQE_NOP; + wqe->flags = USB4_RDMA_DV_WQE_F_SIGNALED; + wqe->wr_id = PROBE_WR_ID_BASE + iter; + wqe->generation = resp.generation; + __atomic_thread_fence(__ATOMIC_RELEASE); + host_store_release_u32(&mem.doorbell->producer.sq_tail, tail); + + err = kick_queue(qp, &kick); + if (err) { + fprintf(stderr, "KICK failed iter=%u: %s (%d)\n", iter, + strerror(err), err); + goto out_destroy; + } + if (wait_gpu_seen(mem.state, iter, timeout_ms)) + goto out_destroy; + } + + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "hipDeviceSynchronize failed: %s\n", + hipGetErrorString(hret)); + goto out_destroy; + } + + { + double secs = (double)(now_ns() - start_ns) / 1000000000.0; + + printf("kind=%s direction=kernel-cqe-to-gpu-plain-load iterations=%u status=OK elapsed_sec=%.6f avg_us=%.3f gpu_seen=%u gpu_error=%u gpu_spins=%" PRIu64 "\n", + kind, iterations, secs, secs * 1000000.0 / iterations, + host_load_acquire_u32(&mem.state->gpu_seen), + host_load_acquire_u32(&mem.state->gpu_error), + mem.state->gpu_spins); + } + ret = 0; + +out_destroy: + err = destroy_queue(qp); + if (err) + fprintf(stderr, "DESTROY_QUEUE failed: %s (%d)\n", + strerror(err), err); +out: + if (qp && ibv_destroy_qp(qp)) + fprintf(stderr, "ibv_destroy_qp: %s\n", strerror(errno)); + if (verbs_cq && ibv_destroy_cq(verbs_cq)) + fprintf(stderr, "ibv_destroy_cq: %s\n", strerror(errno)); + if (doorbell_mr && ibv_dereg_mr(doorbell_mr)) + fprintf(stderr, "ibv_dereg_mr(doorbell): %s\n", strerror(errno)); + if (cq_mr && ibv_dereg_mr(cq_mr)) + fprintf(stderr, "ibv_dereg_mr(CQ): %s\n", strerror(errno)); + if (sq_mr && ibv_dereg_mr(sq_mr)) + fprintf(stderr, "ibv_dereg_mr(SQ): %s\n", strerror(errno)); + if (pd && ibv_dealloc_pd(pd)) + fprintf(stderr, "ibv_dealloc_pd: %s\n", strerror(errno)); + free_queue_mem(&mem); + return ret; +} + +static void usage(const char *argv0) +{ + fprintf(stderr, + "usage: %s [kind] [iterations] [device] [timeout-ms]\n" + " kind: managed|host|host-coherent|host-noncoherent|host-uncached\n", + argv0); +} + +int main(int argc, char **argv) +{ + const char *kind = argc > 1 ? argv[1] : "host-coherent"; + uint32_t iterations = argc > 2 ? strtoul(argv[2], NULL, 0) : 10000; + const char *wanted = argc > 3 ? argv[3] : "usb4_rdma0"; + uint32_t timeout_ms = argc > 4 ? strtoul(argv[4], NULL, 0) : 5000; + struct ibv_device **list = NULL; + struct ibv_context *ctx = NULL; + struct ibv_device *dev; + int count = 0; + int ret = 1; + hipError_t hret; + + if (!iterations || !timeout_ms) { + usage(argv[0]); + return 2; + } + + hret = hipSetDevice(0); + if (hret != hipSuccess) { + fprintf(stderr, "hipSetDevice failed: %s\n", + hipGetErrorString(hret)); + return 1; + } + + list = ibv_get_device_list(&count); + if (!list) { + perror("ibv_get_device_list"); + return 1; + } + dev = find_device(list, count, wanted); + if (!dev) { + fprintf(stderr, "no USB4 RDMA device found wanted=%s\n", wanted); + goto out; + } + ctx = ibv_open_device(dev); + if (!ctx) { + fprintf(stderr, "ibv_open_device(%s): %s\n", + ibv_get_device_name(dev), strerror(errno)); + goto out; + } + + printf("device=%s kind=%s iterations=%u timeout_ms=%u\n", + ibv_get_device_name(dev), kind, iterations, timeout_ms); + ret = run_probe(ctx, kind, iterations, timeout_ms); + +out: + if (ctx) + ibv_close_device(ctx); + if (list) + ibv_free_device_list(list); + return ret; +} diff --git a/userspace/bench/hip_rdma_write_visibility_probe.cpp b/userspace/bench/hip_rdma_write_visibility_probe.cpp new file mode 100644 index 0000000..721ed51 --- /dev/null +++ b/userspace/bench/hip_rdma_write_visibility_probe.cpp @@ -0,0 +1,1744 @@ +// SPDX-License-Identifier: GPL-2.0 OR BSD-3-Clause +/* + * Two-host probe for the exact USB4 GDA visibility pattern used by rocSHMEM + * alltoall staging: + * + * peer kernel writes payload into a HIP host allocation via RDMA WRITE + * peer kernel writes a later signal word via RDMA WRITE + * GPU observes the signal and then reads the payload + * + * The GPU intentionally preloads the payload before each signal wait. If the + * allocation/load mode can return stale GPU-cached payload after the kernel's + * CPU copy, this probe catches it without involving RCCL or rocSHMEM. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +#define VIS_MAGIC 0x56495332u /* VIS2 */ + +enum load_mode { + LOAD_NORMAL = 0, + LOAD_ATOMIC = 1, +}; + +enum source_fill { + SOURCE_FILL_CPU = 0, + SOURCE_FILL_GPU = 1, + SOURCE_FILL_GPU_HDP = 2, + SOURCE_FILL_GPU_SYNC = 3, + SOURCE_FILL_GPU_HDP_SYNC = 4, + SOURCE_FILL_GPU_HOST_HDP = 5, +}; + +enum source_reg { + SOURCE_REG_MR = 0, + SOURCE_REG_DMABUF = 1, +}; + +enum recv_reg { + RECV_REG_MR = 0, + RECV_REG_DMABUF = 1, +}; + +struct opts { + const char *role; + const char *dev; + const char *connect_host; + const char *kind; + const char *mode_name; + const char *source_kind; + const char *source_fill_name; + const char *source_reg_name; + const char *recv_reg_name; + int port; + int gid_index; + int ib_port; + int mtu; + int timeout_ms; + uint32_t count; + size_t size; + enum load_mode mode; + enum source_fill source_fill; + enum source_reg source_reg; + enum recv_reg recv_reg; +}; + +struct peer_info { + uint32_t magic; + uint32_t qpn; + uint32_t psn; + uint32_t lid; + uint32_t rkey; + uint64_t addr; + uint8_t gid[16]; +}; + +struct alignas(64) ProbeState { + uint32_t gpu_seen; + uint32_t gpu_error_seq; + uint32_t gpu_error_off; + uint32_t gpu_error_mode; + uint32_t gpu_got; + uint32_t gpu_want; + uint64_t gpu_spins; + uint32_t prefetch_sink; + uint8_t pad[64 - 6 * sizeof(uint32_t) - sizeof(uint64_t) - + sizeof(uint32_t)]; +}; + +struct HipRegion { + const char *kind; + void *base; + size_t bytes; + uint8_t *payload; + uint64_t *signal; + ProbeState *state; + size_t signal_offset; + size_t state_offset; + bool hip_alloc; + bool host_alloc; +}; + +struct alignas(64) SourceFillState { + uint32_t ready_seq; + uint32_t error_seq; + uint32_t error_off; + uint32_t pad[13]; +}; + +struct SenderSource { + const char *kind; + const char *fill_name; + const char *reg_name; + void *payload_base; + uint8_t *payload; + size_t payload_bytes; + bool payload_hip_alloc; + bool payload_hip_vmm_alloc; + bool payload_hsa_alloc; + bool payload_host_alloc; + bool payload_malloc_alloc; + hipMemGenericAllocationHandle_t payload_vmm_handle; + size_t payload_vmm_bytes; + struct ibv_mr *payload_mr; + int dmabuf_fd; + bool dmabuf_fd_is_hsa; + uint64_t dmabuf_offset; + void *signal_base; + uint64_t *signal; + struct ibv_mr *signal_mr; + SourceFillState *fill_state; + uint32_t *hdp_flush_ptr; +}; + +struct ReceiverMr { + const char *reg_name; + struct ibv_mr *mr; + int dmabuf_fd = -1; + uint64_t dmabuf_offset; + size_t length; +}; + +static uint64_t now_ns(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (uint64_t)ts.tv_sec * 1000000000ull + (uint64_t)ts.tv_nsec; +} + +static uintptr_t align_up(uintptr_t value, size_t alignment) +{ + return (value + alignment - 1) & ~(uintptr_t)(alignment - 1); +} + +static uint32_t host_load_acquire_u32(uint32_t *ptr) +{ + return __atomic_load_n(ptr, __ATOMIC_ACQUIRE); +} + +static uint64_t host_load_acquire_u64(uint64_t *ptr) +{ + return __atomic_load_n(ptr, __ATOMIC_ACQUIRE); +} + +static __host__ __device__ inline uint32_t pattern_word(uint32_t seq, + uint32_t word) +{ + return 0x9e3779b9u ^ (seq * 0x10001u) ^ (word * 0x45d9f3bu) ^ + (word >> 3); +} + +static int send_all(int fd, const void *buf, size_t len) +{ + const char *p = (const char *)buf; + + while (len) { + ssize_t n = send(fd, p, len, 0); + + if (n < 0) { + if (errno == EINTR) + continue; + return -1; + } + if (!n) + return -1; + p += n; + len -= (size_t)n; + } + return 0; +} + +static int recv_all(int fd, void *buf, size_t len) +{ + char *p = (char *)buf; + + while (len) { + ssize_t n = recv(fd, p, len, MSG_WAITALL); + + if (n < 0) { + if (errno == EINTR) + continue; + return -1; + } + if (!n) + return -1; + p += n; + len -= (size_t)n; + } + return 0; +} + +static int tcp_listen(int port) +{ + struct sockaddr_in addr; + int fd; + int one = 1; + + fd = socket(AF_INET, SOCK_STREAM, 0); + if (fd < 0) + return -1; + setsockopt(fd, SOL_SOCKET, SO_REUSEADDR, &one, sizeof(one)); + memset(&addr, 0, sizeof(addr)); + addr.sin_family = AF_INET; + addr.sin_addr.s_addr = htonl(INADDR_ANY); + addr.sin_port = htons((uint16_t)port); + if (bind(fd, (struct sockaddr *)&addr, sizeof(addr)) || + listen(fd, 1)) { + close(fd); + return -1; + } + return fd; +} + +static int tcp_connect(const char *host, int port) +{ + struct addrinfo hints; + struct addrinfo *res = NULL; + struct addrinfo *ai; + char portbuf[16]; + int fd = -1; + + memset(&hints, 0, sizeof(hints)); + hints.ai_socktype = SOCK_STREAM; + hints.ai_family = AF_UNSPEC; + snprintf(portbuf, sizeof(portbuf), "%d", port); + if (getaddrinfo(host, portbuf, &hints, &res)) + return -1; + for (ai = res; ai; ai = ai->ai_next) { + fd = socket(ai->ai_family, ai->ai_socktype, ai->ai_protocol); + if (fd < 0) + continue; + if (!connect(fd, ai->ai_addr, ai->ai_addrlen)) + break; + close(fd); + fd = -1; + } + freeaddrinfo(res); + return fd; +} + +static struct ibv_context *open_dev(const char *name) +{ + struct ibv_device **list; + struct ibv_context *ctx = NULL; + int n = 0; + + list = ibv_get_device_list(&n); + if (!list) + return NULL; + for (int i = 0; i < n; i++) { + if (!strcmp(ibv_get_device_name(list[i]), name)) { + ctx = ibv_open_device(list[i]); + break; + } + } + ibv_free_device_list(list); + return ctx; +} + +static enum ibv_mtu mtu_enum(int mtu) +{ + switch (mtu) { + case 256: + return IBV_MTU_256; + case 512: + return IBV_MTU_512; + case 1024: + return IBV_MTU_1024; + case 2048: + return IBV_MTU_2048; + case 4096: + default: + return IBV_MTU_4096; + } +} + +static int qp_to_rts(struct ibv_qp *qp, int port, int sgid_index, + const union ibv_gid *dgid, uint32_t dlid, + uint32_t dest_qpn, uint32_t local_psn, + uint32_t remote_psn, enum ibv_mtu path_mtu, + int qp_access_flags) +{ + struct ibv_qp_attr a; + int ret; + + memset(&a, 0, sizeof(a)); + a.qp_state = IBV_QPS_INIT; + a.pkey_index = 0; + a.port_num = (uint8_t)port; + a.qp_access_flags = qp_access_flags; + ret = ibv_modify_qp(qp, &a, IBV_QP_STATE | IBV_QP_PKEY_INDEX | + IBV_QP_PORT | IBV_QP_ACCESS_FLAGS); + if (ret) { + perror("modify INIT"); + return ret; + } + + memset(&a, 0, sizeof(a)); + a.qp_state = IBV_QPS_RTR; + a.path_mtu = path_mtu; + a.rq_psn = remote_psn; + a.dest_qp_num = dest_qpn; + a.max_dest_rd_atomic = 1; + a.min_rnr_timer = 12; + a.ah_attr.dlid = (uint16_t)dlid; + a.ah_attr.sl = 0; + a.ah_attr.src_path_bits = 0; + a.ah_attr.port_num = (uint8_t)port; + if (dgid->global.interface_id) { + a.ah_attr.is_global = 1; + a.ah_attr.grh.dgid = *dgid; + a.ah_attr.grh.sgid_index = sgid_index; + a.ah_attr.grh.hop_limit = 1; + } + ret = ibv_modify_qp(qp, &a, IBV_QP_STATE | IBV_QP_AV | + IBV_QP_PATH_MTU | IBV_QP_DEST_QPN | + IBV_QP_RQ_PSN | IBV_QP_MAX_DEST_RD_ATOMIC | + IBV_QP_MIN_RNR_TIMER); + if (ret) { + perror("modify RTR"); + return ret; + } + + memset(&a, 0, sizeof(a)); + a.qp_state = IBV_QPS_RTS; + a.sq_psn = local_psn; + a.timeout = 14; + a.retry_cnt = 7; + a.rnr_retry = 7; + a.max_rd_atomic = 1; + ret = ibv_modify_qp(qp, &a, IBV_QP_STATE | IBV_QP_SQ_PSN | + IBV_QP_TIMEOUT | IBV_QP_RETRY_CNT | + IBV_QP_RNR_RETRY | IBV_QP_MAX_QP_RD_ATOMIC); + if (ret) + perror("modify RTS"); + return ret; +} + +static int exchange_info(int fd, const struct peer_info *local, + struct peer_info *remote) +{ + if (send_all(fd, local, sizeof(*local)) || + recv_all(fd, remote, sizeof(*remote))) { + perror("metadata exchange"); + return -1; + } + if (remote->magic != local->magic) { + fprintf(stderr, "bad remote magic 0x%x\n", remote->magic); + return -1; + } + return 0; +} + +static int parse_mode(const char *s, enum load_mode *mode) +{ + if (!strcmp(s, "normal")) + *mode = LOAD_NORMAL; + else if (!strcmp(s, "atomic")) + *mode = LOAD_ATOMIC; + else + return -1; + return 0; +} + +static int parse_source_fill(const char *s, enum source_fill *fill) +{ + if (!strcmp(s, "cpu")) + *fill = SOURCE_FILL_CPU; + else if (!strcmp(s, "gpu")) + *fill = SOURCE_FILL_GPU; + else if (!strcmp(s, "gpu-hdp")) + *fill = SOURCE_FILL_GPU_HDP; + else if (!strcmp(s, "gpu-sync")) + *fill = SOURCE_FILL_GPU_SYNC; + else if (!strcmp(s, "gpu-hdp-sync")) + *fill = SOURCE_FILL_GPU_HDP_SYNC; + else if (!strcmp(s, "gpu-host-hdp")) + *fill = SOURCE_FILL_GPU_HOST_HDP; + else + return -1; + return 0; +} + +static int parse_source_reg(const char *s, enum source_reg *reg) +{ + if (!strcmp(s, "reg_mr")) + *reg = SOURCE_REG_MR; + else if (!strcmp(s, "dmabuf")) + *reg = SOURCE_REG_DMABUF; + else + return -1; + return 0; +} + +static int parse_recv_reg(const char *s, enum recv_reg *reg) +{ + if (!strcmp(s, "reg_mr")) + *reg = RECV_REG_MR; + else if (!strcmp(s, "dmabuf")) + *reg = RECV_REG_DMABUF; + else + return -1; + return 0; +} + +static void usage(const char *argv0) +{ + fprintf(stderr, + "usage: %s --role recv|send --dev DEV --gid-index N --port P\n" + " [--connect HOST] [--kind device|host-coherent|host-uncached|host|managed]\n" + " [--recv-reg reg_mr|dmabuf]\n" + " [--mode normal|atomic]\n" + " [--source-kind malloc|device|managed|host|host-coherent|host-uncached|\n" + " vmm-device-pinned|vmm-device-uncached|vmm-host-pinned|vmm-host-uncached|\n" + " ext-finegrained|ext-uncached|hsa-gpu-extended]\n" + " [--source-fill cpu|gpu|gpu-hdp|gpu-sync|gpu-hdp-sync|gpu-host-hdp]\n" + " [--source-reg reg_mr|dmabuf]\n" + " [--size BYTES] [--count N] [--timeout-ms N]\n", + argv0); +} + +static int parse_opts(int argc, char **argv, struct opts *o) +{ + memset(o, 0, sizeof(*o)); + o->dev = "usb4_rdma0"; + o->kind = "host-coherent"; + o->mode_name = "normal"; + o->source_kind = "malloc"; + o->source_fill_name = "cpu"; + o->source_reg_name = "reg_mr"; + o->recv_reg_name = "reg_mr"; + o->mode = LOAD_NORMAL; + o->source_fill = SOURCE_FILL_CPU; + o->source_reg = SOURCE_REG_MR; + o->recv_reg = RECV_REG_MR; + o->port = 18518; + o->gid_index = 1; + o->ib_port = 1; + o->mtu = 4096; + o->timeout_ms = 5000; + o->count = 1000; + o->size = 32; + + for (int i = 1; i < argc; i++) { + if (!strcmp(argv[i], "--role") && i + 1 < argc) { + o->role = argv[++i]; + } else if (!strcmp(argv[i], "--dev") && i + 1 < argc) { + o->dev = argv[++i]; + } else if (!strcmp(argv[i], "--connect") && i + 1 < argc) { + o->connect_host = argv[++i]; + } else if (!strcmp(argv[i], "--kind") && i + 1 < argc) { + o->kind = argv[++i]; + } else if (!strcmp(argv[i], "--mode") && i + 1 < argc) { + o->mode_name = argv[++i]; + if (parse_mode(o->mode_name, &o->mode)) + return -1; + } else if (!strcmp(argv[i], "--source-kind") && i + 1 < argc) { + o->source_kind = argv[++i]; + } else if (!strcmp(argv[i], "--source-fill") && i + 1 < argc) { + o->source_fill_name = argv[++i]; + if (parse_source_fill(o->source_fill_name, + &o->source_fill)) + return -1; + } else if (!strcmp(argv[i], "--source-reg") && i + 1 < argc) { + o->source_reg_name = argv[++i]; + if (parse_source_reg(o->source_reg_name, + &o->source_reg)) + return -1; + } else if (!strcmp(argv[i], "--recv-reg") && i + 1 < argc) { + o->recv_reg_name = argv[++i]; + if (parse_recv_reg(o->recv_reg_name, &o->recv_reg)) + return -1; + } else if (!strcmp(argv[i], "--port") && i + 1 < argc) { + o->port = atoi(argv[++i]); + } else if (!strcmp(argv[i], "--gid-index") && i + 1 < argc) { + o->gid_index = atoi(argv[++i]); + } else if (!strcmp(argv[i], "--ib-port") && i + 1 < argc) { + o->ib_port = atoi(argv[++i]); + } else if (!strcmp(argv[i], "--mtu") && i + 1 < argc) { + o->mtu = atoi(argv[++i]); + } else if (!strcmp(argv[i], "--timeout-ms") && i + 1 < argc) { + o->timeout_ms = atoi(argv[++i]); + } else if (!strcmp(argv[i], "--count") && i + 1 < argc) { + o->count = (uint32_t)strtoul(argv[++i], NULL, 0); + } else if (!strcmp(argv[i], "--size") && i + 1 < argc) { + o->size = strtoull(argv[++i], NULL, 0); + } else { + return -1; + } + } + + if (!o->role || (strcmp(o->role, "recv") && strcmp(o->role, "send")) || + o->port <= 0 || o->gid_index < 0 || o->ib_port <= 0 || + o->timeout_ms <= 0 || !o->count || !o->size || + (o->size % sizeof(uint32_t)) != 0) + return -1; + if (o->source_fill == SOURCE_FILL_CPU && + strcmp(o->source_kind, "malloc") && + strcmp(o->source_kind, "host") && + strcmp(o->source_kind, "host-coherent") && + strcmp(o->source_kind, "host-uncached") && + strcmp(o->source_kind, "managed")) + return -1; + if (o->source_fill != SOURCE_FILL_CPU && + !strcmp(o->source_kind, "malloc")) + return -1; + if (o->source_reg == SOURCE_REG_DMABUF && + !strcmp(o->source_kind, "malloc")) + return -1; + return 0; +} + +static bool source_kind_is_vmm(const char *kind) +{ + return !strcmp(kind, "vmm-device-pinned") || + !strcmp(kind, "vmm-device-uncached") || + !strcmp(kind, "vmm-host-pinned") || + !strcmp(kind, "vmm-host-uncached"); +} + +static bool source_kind_is_hsa(const char *kind) +{ + return !strcmp(kind, "hsa-gpu-extended"); +} + +static hipError_t alloc_region(const char *kind, size_t payload_size, + HipRegion *region) +{ + unsigned int flags = hipHostMallocMapped; + uintptr_t cursor; + void *base = NULL; + size_t bytes = align_up(payload_size, 64) + 64 + sizeof(ProbeState) + + 4096; + hipError_t ret; + + memset(region, 0, sizeof(*region)); + region->kind = kind; + region->bytes = bytes; + if (!strcmp(kind, "managed")) { + ret = hipMallocManaged(&base, bytes); + region->hip_alloc = true; + } else if (!strcmp(kind, "device")) { + ret = hipMalloc(&base, bytes); + region->hip_alloc = true; + } else if (!strcmp(kind, "host")) { + ret = hipHostMalloc(&base, bytes, flags); + region->host_alloc = true; + } else if (!strcmp(kind, "host-coherent")) { + ret = hipHostMalloc(&base, bytes, + flags | hipHostMallocCoherent); + region->host_alloc = true; + } else if (!strcmp(kind, "host-uncached")) { + ret = hipHostMalloc(&base, bytes, + flags | hipHostMallocUncached); + region->host_alloc = true; + } else { + return hipErrorInvalidValue; + } + if (ret != hipSuccess) + return ret; + + cursor = align_up((uintptr_t)base, 64); + region->payload = (uint8_t *)cursor; + cursor += align_up(payload_size, 64); + region->signal_offset = cursor - (uintptr_t)region->payload; + region->signal = (uint64_t *)cursor; + cursor += 64; + cursor = align_up(cursor, 64); + region->state_offset = cursor - (uintptr_t)region->payload; + region->state = (ProbeState *)cursor; + region->base = base; + if (!strcmp(kind, "device")) { + ret = hipMemset(base, 0, bytes); + if (ret == hipSuccess) + ret = hipDeviceSynchronize(); + if (ret != hipSuccess) { + (void)hipFree(base); + memset(region, 0, sizeof(*region)); + return ret; + } + } else { + memset(base, 0, bytes); + } + return hipSuccess; +} + +static void free_region(HipRegion *region) +{ + if (!region->base) + return; + if (region->hip_alloc) + (void)hipFree(region->base); + else if (region->host_alloc) + (void)hipHostFree(region->base); + memset(region, 0, sizeof(*region)); +} + +__device__ static inline uint32_t gpu_atomic_load_u32(const uint32_t *ptr) +{ + return __hip_atomic_load(const_cast(ptr), __ATOMIC_ACQUIRE, + __HIP_MEMORY_SCOPE_SYSTEM); +} + +__device__ static inline uint64_t gpu_atomic_load_u64(const uint64_t *ptr) +{ + return __hip_atomic_load(const_cast(ptr), __ATOMIC_ACQUIRE, + __HIP_MEMORY_SCOPE_SYSTEM); +} + +__device__ static inline uint32_t gpu_payload_load_u32(const uint32_t *ptr, + enum load_mode mode) +{ + if (mode == LOAD_ATOMIC) + return gpu_atomic_load_u32(ptr); + return *reinterpret_cast(ptr); +} + +__device__ static inline void gpu_store_release_u32(uint32_t *ptr, + uint32_t value) +{ + __hip_atomic_store(ptr, value, __ATOMIC_RELEASE, + __HIP_MEMORY_SCOPE_SYSTEM); +} + +__device__ static inline bool gpu_cas_system_u32(uint32_t *ptr, + uint32_t expected, + uint32_t desired) +{ + return __hip_atomic_compare_exchange_strong( + ptr, &expected, desired, __ATOMIC_ACQ_REL, __ATOMIC_ACQUIRE, + __HIP_MEMORY_SCOPE_SYSTEM); +} + +__global__ void source_fill_kernel(uint8_t *payload, SourceFillState *state, + uint32_t words, uint32_t seq, + uint32_t *hdp_flush_ptr) +{ + uint32_t *payload_words = reinterpret_cast(payload); + uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x; + uint32_t nthreads = blockDim.x * gridDim.x; + + for (uint32_t word = tid; word < words; word += nthreads) + payload_words[word] = pattern_word(seq, word); + + __syncthreads(); + if (tid == 0) { + __threadfence_system(); + if (hdp_flush_ptr) { + __atomic_store_n(hdp_flush_ptr, 1u, __ATOMIC_SEQ_CST); + __threadfence_system(); + } + gpu_store_release_u32(&state->ready_seq, seq); + } +} + +__global__ void rdma_write_visibility_kernel(uint8_t *payload, + uint64_t *signal, + ProbeState *state, + uint32_t words, + uint32_t count, + enum load_mode mode, + uint32_t max_spins) +{ + __shared__ uint32_t block_error; + uint32_t *payload_words = reinterpret_cast(payload); + uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x; + uint32_t nthreads = blockDim.x * gridDim.x; + uint32_t sink = 0; + + for (uint32_t seq = 1; seq <= count; ++seq) { + if (tid == 0) + block_error = 0; + __syncthreads(); + + for (uint32_t word = tid; word < words; word += nthreads) + sink ^= *reinterpret_cast( + payload_words + word); + if (tid == 0) + state->prefetch_sink ^= sink; + __syncthreads(); + + uint32_t spins = 0; + while (gpu_atomic_load_u64(signal) != seq) { + if (++spins == max_spins) { + if (tid == 0) { + state->gpu_spins += spins; + state->gpu_error_mode = + static_cast(mode); + gpu_store_release_u32( + &state->gpu_error_seq, seq); + } + return; + } + } + + if (tid == 0) + state->gpu_spins += spins; + for (uint32_t word = tid; word < words; word += nthreads) { + uint32_t got = + gpu_payload_load_u32(payload_words + word, mode); + uint32_t want = pattern_word(seq, word); + + if (got != want) { + if (gpu_cas_system_u32(&state->gpu_error_seq, + 0u, seq)) { + state->gpu_error_off = word * 4; + state->gpu_error_mode = + static_cast(mode); + state->gpu_got = got; + state->gpu_want = want; + __threadfence_system(); + } + atomicExch(&block_error, 1u); + } + } + + __syncthreads(); + if (block_error) + return; + if (tid == 0) + gpu_store_release_u32(&state->gpu_seen, seq); + __syncthreads(); + } +} + +static int wait_gpu_seen(HipRegion *region, uint32_t seq, int timeout_ms) +{ + uint64_t deadline = now_ns() + (uint64_t)timeout_ms * 1000000ull; + + while (now_ns() < deadline) { + uint32_t err = host_load_acquire_u32(®ion->state->gpu_error_seq); + uint32_t seen = host_load_acquire_u32(®ion->state->gpu_seen); + + if (err) { + fprintf(stderr, + "gpu error seq=%u seen=%u off=%u mode=%u got=0x%08x want=0x%08x signal=%" PRIu64 "\n", + err, seen, region->state->gpu_error_off, + region->state->gpu_error_mode, + region->state->gpu_got, region->state->gpu_want, + host_load_acquire_u64(region->signal)); + return -1; + } + if (seen >= seq) + return 0; + usleep(100); + } + fprintf(stderr, "timeout waiting gpu_seen=%u got=%u signal=%" PRIu64 "\n", + seq, host_load_acquire_u32(®ion->state->gpu_seen), + host_load_acquire_u64(region->signal)); + return -1; +} + +static int poll_send_cq(struct ibv_cq *cq, int timeout_ms) +{ + uint64_t deadline = now_ns() + (uint64_t)timeout_ms * 1000000ull; + struct ibv_wc wc; + + while (now_ns() < deadline) { + int n = ibv_poll_cq(cq, 1, &wc); + + if (n < 0) { + fprintf(stderr, "ibv_poll_cq failed: %d\n", n); + return -1; + } + if (!n) { + usleep(100); + continue; + } + if (wc.status != IBV_WC_SUCCESS) { + fprintf(stderr, + "send wc error wr_id=%" PRIu64 " status=%u opcode=%u byte_len=%u\n", + wc.wr_id, wc.status, wc.opcode, wc.byte_len); + return -1; + } + return 0; + } + fprintf(stderr, "timeout waiting send CQE\n"); + return -1; +} + +static void fill_payload(uint32_t *payload, size_t bytes, uint32_t seq) +{ + size_t words = bytes / sizeof(uint32_t); + + for (size_t i = 0; i < words; ++i) + payload[i] = pattern_word(seq, (uint32_t)i); +} + +struct HsaAgentInfo { + hsa_agent_t agent; + hsa_device_type_t type; +}; + +struct HsaPoolInfo { + hsa_amd_memory_pool_t pool; + hsa_agent_t owner; + hsa_device_type_t owner_type; + hsa_amd_segment_t segment; + hsa_amd_memory_pool_location_t location; + uint32_t flags; + bool alloc_allowed; + size_t alloc_granule; +}; + +struct HsaPoolEnumCtx { + hsa_agent_t owner; + hsa_device_type_t owner_type; + std::vector *pools; +}; + +static const char *hsa_status_name(hsa_status_t status) +{ + const char *name = NULL; + + if (hsa_status_string(status, &name) != HSA_STATUS_SUCCESS || !name) + return "unknown"; + return name; +} + +static hsa_status_t collect_hsa_agent(hsa_agent_t agent, void *data) +{ + std::vector *agents = + static_cast *>(data); + HsaAgentInfo info = {}; + hsa_status_t status; + + info.agent = agent; + status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &info.type); + if (status != HSA_STATUS_SUCCESS) + info.type = static_cast(-1); + agents->push_back(info); + return HSA_STATUS_SUCCESS; +} + +static hsa_status_t collect_hsa_pool(hsa_amd_memory_pool_t pool, void *data) +{ + HsaPoolEnumCtx *ctx = static_cast(data); + HsaPoolInfo info = {}; + hsa_status_t status; + + info.pool = pool; + info.owner = ctx->owner; + info.owner_type = ctx->owner_type; + + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &info.segment); + if (status != HSA_STATUS_SUCCESS) + return HSA_STATUS_SUCCESS; + if (info.segment == HSA_AMD_SEGMENT_GLOBAL) { + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, + &info.flags); + } + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, + &info.alloc_allowed); + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_LOCATION, &info.location); + if (info.alloc_allowed) { + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, + &info.alloc_granule); + } + ctx->pools->push_back(info); + return HSA_STATUS_SUCCESS; +} + +static bool hsa_pool_matches_source_kind(const HsaPoolInfo &pool, + const char *kind) +{ + if (!strcmp(kind, "hsa-gpu-extended")) { + return pool.owner_type == HSA_DEVICE_TYPE_GPU && + pool.segment == HSA_AMD_SEGMENT_GLOBAL && + pool.location == HSA_AMD_MEMORY_POOL_LOCATION_GPU && + pool.alloc_allowed && + (pool.flags & + HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED); + } + return false; +} + +static size_t round_up_size(size_t value, size_t alignment) +{ + if (!alignment) + return value; + size_t rem = value % alignment; + return rem ? value + alignment - rem : value; +} + +static hipError_t alloc_source_payload_hsa(const char *kind, size_t bytes, + SenderSource *src) +{ + std::vector agents; + std::vector pools; + hsa_amd_memory_pool_t selected = {}; + size_t granule = 4096; + void *ptr = NULL; + hsa_status_t status; + + status = hsa_init(); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_init failed: %s\n", + hsa_status_name(status)); + return hipErrorInvalidValue; + } + + status = hsa_iterate_agents(collect_hsa_agent, &agents); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_iterate_agents failed: %s\n", + hsa_status_name(status)); + return hipErrorInvalidValue; + } + + for (const HsaAgentInfo &agent : agents) { + HsaPoolEnumCtx ctx = { agent.agent, agent.type, &pools }; + (void)hsa_amd_agent_iterate_memory_pools( + agent.agent, collect_hsa_pool, &ctx); + } + + for (const HsaPoolInfo &pool : pools) { + if (!hsa_pool_matches_source_kind(pool, kind)) + continue; + selected = pool.pool; + granule = pool.alloc_granule ? pool.alloc_granule : 4096; + break; + } + if (!selected.handle) { + fprintf(stderr, "no HSA pool matches source kind=%s\n", kind); + return hipErrorInvalidValue; + } + + size_t size = round_up_size(bytes, granule); + status = hsa_amd_memory_pool_allocate(selected, size, 0, &ptr); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_amd_memory_pool_allocate kind=%s failed: %s\n", + kind, hsa_status_name(status)); + return hipErrorOutOfMemory; + } + + std::vector allowed_agents; + for (const HsaAgentInfo &agent : agents) { + hsa_amd_memory_pool_access_t access = + HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED; + + if (agent.type != HSA_DEVICE_TYPE_CPU && + agent.type != HSA_DEVICE_TYPE_GPU) + continue; + status = hsa_amd_agent_memory_pool_get_info( + agent.agent, selected, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + if (status == HSA_STATUS_SUCCESS && + access != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) + allowed_agents.push_back(agent.agent); + } + + status = hsa_amd_agents_allow_access( + static_cast(allowed_agents.size()), + allowed_agents.data(), NULL, ptr); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_amd_agents_allow_access kind=%s failed: %s\n", + kind, hsa_status_name(status)); + (void)hsa_amd_memory_pool_free(ptr); + return hipErrorInvalidValue; + } + + src->payload_base = ptr; + src->payload = static_cast(ptr); + src->payload_bytes = size; + src->payload_hsa_alloc = true; + memset(ptr, 0, bytes); + return hipSuccess; +} + +static hipError_t alloc_source_payload_vmm(const char *kind, size_t bytes, + SenderSource *src) +{ + hipMemAllocationProp prop = {}; + hipMemAccessDesc access = {}; + hipMemGenericAllocationHandle_t handle = NULL; + void *ptr = NULL; + size_t granule = 0; + size_t map_bytes; + int device = 0; + hipError_t ret; + + ret = hipGetDevice(&device); + if (ret != hipSuccess) + return ret; + + prop.type = strstr(kind, "uncached") ? hipMemAllocationTypeUncached : + hipMemAllocationTypePinned; + prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor; + prop.location.type = strstr(kind, "device") ? + hipMemLocationTypeDevice : + hipMemLocationTypeHost; + prop.location.id = prop.location.type == hipMemLocationTypeDevice ? + device : + 0; + prop.allocFlags.gpuDirectRDMACapable = 1; + + ret = hipMemGetAllocationGranularity( + &granule, &prop, hipMemAllocationGranularityMinimum); + if (ret != hipSuccess) + return ret; + if (!granule) + granule = 4096; + map_bytes = (bytes + granule - 1) & ~(granule - 1); + + ret = hipMemAddressReserve(&ptr, map_bytes, granule, NULL, 0); + if (ret != hipSuccess) + return ret; + + ret = hipMemCreate(&handle, map_bytes, &prop, 0); + if (ret != hipSuccess) + goto err_addr; + + ret = hipMemMap(ptr, map_bytes, 0, handle, 0); + if (ret != hipSuccess) + goto err_handle; + + access.location.type = hipMemLocationTypeDevice; + access.location.id = device; + access.flags = hipMemAccessFlagsProtReadWrite; + ret = hipMemSetAccess(ptr, map_bytes, &access, 1); + if (ret != hipSuccess) + goto err_unmap; + + src->payload_base = ptr; + src->payload = static_cast(ptr); + src->payload_vmm_handle = handle; + src->payload_vmm_bytes = map_bytes; + src->payload_hip_vmm_alloc = true; + return hipMemset(ptr, 0, bytes); + +err_unmap: + (void)hipMemUnmap(ptr, map_bytes); +err_handle: + (void)hipMemRelease(handle); +err_addr: + (void)hipMemAddressFree(ptr, map_bytes); + return ret; +} + +static hipError_t alloc_source_payload(const char *kind, size_t bytes, + SenderSource *src) +{ + unsigned int flags = hipHostMallocMapped; + + if (!strcmp(kind, "malloc")) { + if (posix_memalign(&src->payload_base, 64, bytes)) + return hipErrorOutOfMemory; + src->payload_malloc_alloc = true; + memset(src->payload_base, 0, bytes); + src->payload = (uint8_t *)src->payload_base; + return hipSuccess; + } + + if (source_kind_is_hsa(kind)) { + return alloc_source_payload_hsa(kind, bytes, src); + } else if (source_kind_is_vmm(kind)) { + return alloc_source_payload_vmm(kind, bytes, src); + } else if (!strcmp(kind, "ext-finegrained")) { + hipError_t ret = hipExtMallocWithFlags( + &src->payload_base, bytes, hipDeviceMallocFinegrained); + if (ret != hipSuccess) + return ret; + src->payload_hip_alloc = true; + } else if (!strcmp(kind, "ext-uncached")) { + hipError_t ret = hipExtMallocWithFlags( + &src->payload_base, bytes, hipDeviceMallocUncached); + if (ret != hipSuccess) + return ret; + src->payload_hip_alloc = true; + } else if (!strcmp(kind, "device")) { + hipError_t ret = hipMalloc(&src->payload_base, bytes); + if (ret != hipSuccess) + return ret; + src->payload_hip_alloc = true; + } else if (!strcmp(kind, "managed")) { + hipError_t ret = hipMallocManaged(&src->payload_base, bytes); + if (ret != hipSuccess) + return ret; + src->payload_hip_alloc = true; + } else if (!strcmp(kind, "host")) { + hipError_t ret = hipHostMalloc(&src->payload_base, bytes, + flags); + if (ret != hipSuccess) + return ret; + src->payload_host_alloc = true; + } else if (!strcmp(kind, "host-coherent")) { + hipError_t ret = hipHostMalloc(&src->payload_base, bytes, + flags | + hipHostMallocCoherent); + if (ret != hipSuccess) + return ret; + src->payload_host_alloc = true; + } else if (!strcmp(kind, "host-uncached")) { + hipError_t ret = hipHostMalloc(&src->payload_base, bytes, + flags | + hipHostMallocUncached); + if (ret != hipSuccess) + return ret; + src->payload_host_alloc = true; + } else { + return hipErrorInvalidValue; + } + + src->payload = (uint8_t *)src->payload_base; + return hipMemset(src->payload_base, 0, bytes); +} + +static int init_sender_source(struct opts *o, struct ibv_pd *pd, + SenderSource *src) +{ + int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | + IBV_ACCESS_REMOTE_READ; + hipError_t hret; + + memset(src, 0, sizeof(*src)); + src->kind = o->source_kind; + src->fill_name = o->source_fill_name; + src->reg_name = o->source_reg_name; + src->dmabuf_fd = -1; + src->payload_bytes = align_up(o->size, 4096); + if (src->payload_bytes < 4096) + src->payload_bytes = 4096; + + hret = alloc_source_payload(o->source_kind, src->payload_bytes, src); + if (hret != hipSuccess) { + fprintf(stderr, "sender payload alloc kind=%s failed: %s\n", + o->source_kind, hipGetErrorString(hret)); + return -1; + } + if (src->payload_hip_alloc || src->payload_host_alloc) { + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "sender payload init sync failed: %s\n", + hipGetErrorString(hret)); + return -1; + } + } + + if (posix_memalign(&src->signal_base, 64, 64)) { + fprintf(stderr, "sender signal alloc failed\n"); + return -1; + } + memset(src->signal_base, 0, 64); + src->signal = (uint64_t *)src->signal_base; + + src->signal_mr = ibv_reg_mr(pd, src->signal_base, 64, + IBV_ACCESS_LOCAL_WRITE); + if (!src->signal_mr) { + fprintf(stderr, "sender signal ibv_reg_mr: %s\n", + strerror(errno)); + return -1; + } + + if (o->source_reg == SOURCE_REG_DMABUF) { + if (src->payload_hip_vmm_alloc) { + hipError_t export_ret; + + export_ret = hipMemExportToShareableHandle( + &src->dmabuf_fd, src->payload_vmm_handle, + hipMemHandleTypePosixFileDescriptor, 0); + if (export_ret != hipSuccess) { + fprintf(stderr, + "sender hipMemExportToShareableHandle kind=%s failed: %s\n", + o->source_kind, + hipGetErrorString(export_ret)); + return -1; + } + src->dmabuf_offset = 0; + src->dmabuf_fd_is_hsa = false; + } else { + hsa_status_t hsa_ret; + + hsa_ret = hsa_amd_portable_export_dmabuf( + src->payload, src->payload_bytes, + &src->dmabuf_fd, &src->dmabuf_offset); + if (hsa_ret != HSA_STATUS_SUCCESS) { + fprintf(stderr, + "sender dmabuf export kind=%s failed hsa_status=%d\n", + o->source_kind, hsa_ret); + return -1; + } + src->dmabuf_fd_is_hsa = true; + } + src->payload_mr = ibv_reg_dmabuf_mr( + pd, src->dmabuf_offset, src->payload_bytes, + (uintptr_t)src->payload, src->dmabuf_fd, access); + } else { + src->payload_mr = ibv_reg_mr(pd, src->payload, + src->payload_bytes, access); + } + if (!src->payload_mr) { + fprintf(stderr, "sender payload reg kind=%s method=%s: %s\n", + o->source_kind, o->source_reg_name, strerror(errno)); + return -1; + } + + if (o->source_fill != SOURCE_FILL_CPU) { + hret = hipHostMalloc(reinterpret_cast(&src->fill_state), + sizeof(*src->fill_state), + hipHostMallocMapped | + hipHostMallocCoherent); + if (hret != hipSuccess) { + fprintf(stderr, "sender fill state alloc failed: %s\n", + hipGetErrorString(hret)); + return -1; + } + memset(src->fill_state, 0, sizeof(*src->fill_state)); + if (o->source_fill == SOURCE_FILL_GPU_HDP || + o->source_fill == SOURCE_FILL_GPU_HDP_SYNC || + o->source_fill == SOURCE_FILL_GPU_HOST_HDP) { + hret = hipDeviceGetAttribute( + reinterpret_cast(&src->hdp_flush_ptr), + hipDeviceAttributeHdpMemFlushCntl, 0); + if (hret != hipSuccess || !src->hdp_flush_ptr) { + fprintf(stderr, + "hipDeviceAttributeHdpMemFlushCntl failed: %s ptr=%p\n", + hipGetErrorString(hret), + src->hdp_flush_ptr); + return -1; + } + } + } + + printf("send_source kind=%s fill=%s reg=%s payload=%p payload_bytes=%zu lkey=0x%x signal=%p signal_lkey=0x%x dmabuf_fd=%d dmabuf_offset=%" PRIu64 " hdp=%p\n", + src->kind, src->fill_name, src->reg_name, src->payload, + src->payload_bytes, src->payload_mr->lkey, src->signal, + src->signal_mr->lkey, src->dmabuf_fd, src->dmabuf_offset, + src->hdp_flush_ptr); + return 0; +} + +static void cleanup_sender_source(SenderSource *src) +{ + if (src->payload_mr && ibv_dereg_mr(src->payload_mr)) + fprintf(stderr, "sender payload ibv_dereg_mr: %s\n", + strerror(errno)); + if (src->signal_mr && ibv_dereg_mr(src->signal_mr)) + fprintf(stderr, "sender signal ibv_dereg_mr: %s\n", + strerror(errno)); + if (src->dmabuf_fd >= 0) { + if (src->dmabuf_fd_is_hsa) + (void)hsa_amd_portable_close_dmabuf(src->dmabuf_fd); + else + (void)close(src->dmabuf_fd); + } + if (src->fill_state) + (void)hipHostFree(src->fill_state); + if (src->payload_host_alloc) + (void)hipHostFree(src->payload_base); + else if (src->payload_hip_vmm_alloc) { + (void)hipMemUnmap(src->payload_base, src->payload_vmm_bytes); + (void)hipMemAddressFree(src->payload_base, + src->payload_vmm_bytes); + (void)hipMemRelease(src->payload_vmm_handle); + } + else if (src->payload_hsa_alloc) + (void)hsa_amd_memory_pool_free(src->payload_base); + else if (src->payload_hip_alloc) + (void)hipFree(src->payload_base); + else if (src->payload_malloc_alloc) + free(src->payload_base); + free(src->signal_base); +} + +static int init_receiver_mr(struct opts *o, struct ibv_pd *pd, + HipRegion *region, ReceiverMr *recv) +{ + int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE; + + memset(recv, 0, sizeof(*recv)); + recv->reg_name = o->recv_reg_name; + recv->dmabuf_fd = -1; + recv->length = region->state_offset + sizeof(*region->state); + + if (o->recv_reg == RECV_REG_DMABUF) { + hsa_status_t hsa_ret; + + hsa_ret = hsa_amd_portable_export_dmabuf( + region->payload, recv->length, &recv->dmabuf_fd, + &recv->dmabuf_offset); + if (hsa_ret != HSA_STATUS_SUCCESS) { + fprintf(stderr, + "receiver dmabuf export kind=%s failed hsa_status=%d\n", + o->kind, hsa_ret); + return -1; + } + recv->mr = ibv_reg_dmabuf_mr( + pd, recv->dmabuf_offset, recv->length, + (uintptr_t)region->payload, recv->dmabuf_fd, access); + } else { + recv->mr = ibv_reg_mr(pd, region->payload, recv->length, + access); + } + if (!recv->mr) { + fprintf(stderr, "receiver reg kind=%s method=%s: %s\n", + o->kind, o->recv_reg_name, strerror(errno)); + if (recv->dmabuf_fd >= 0) { + (void)hsa_amd_portable_close_dmabuf(recv->dmabuf_fd); + recv->dmabuf_fd = -1; + } + return -1; + } + + printf("recv_mr kind=%s reg=%s length=%zu rkey=0x%x dmabuf_fd=%d dmabuf_offset=%" PRIu64 "\n", + o->kind, recv->reg_name, recv->length, recv->mr->rkey, + recv->dmabuf_fd, recv->dmabuf_offset); + return 0; +} + +static void cleanup_receiver_mr(ReceiverMr *recv) +{ + if (recv->mr && ibv_dereg_mr(recv->mr)) + fprintf(stderr, "receiver ibv_dereg_mr: %s\n", + strerror(errno)); + if (recv->dmabuf_fd >= 0) + (void)hsa_amd_portable_close_dmabuf(recv->dmabuf_fd); +} + +static int wait_source_ready(SenderSource *src, uint32_t seq, int timeout_ms) +{ + uint64_t deadline = now_ns() + (uint64_t)timeout_ms * 1000000ull; + + while (now_ns() < deadline) { + uint32_t err = host_load_acquire_u32(&src->fill_state->error_seq); + uint32_t ready = + host_load_acquire_u32(&src->fill_state->ready_seq); + + if (err) { + fprintf(stderr, "source fill error seq=%u off=%u\n", + err, src->fill_state->error_off); + return -1; + } + if (ready >= seq) + return 0; + usleep(50); + } + + fprintf(stderr, "timeout waiting source_ready=%u got=%u\n", seq, + host_load_acquire_u32(&src->fill_state->ready_seq)); + return -1; +} + +static int prepare_source_payload(SenderSource *src, struct opts *o, + uint32_t seq) +{ + if (o->source_fill == SOURCE_FILL_CPU) { + fill_payload((uint32_t *)src->payload, o->size, seq); + return 0; + } + + __atomic_store_n(&src->fill_state->ready_seq, 0, __ATOMIC_RELEASE); + hipLaunchKernelGGL(source_fill_kernel, dim3(1), dim3(256), 0, 0, + src->payload, src->fill_state, + (uint32_t)(o->size / sizeof(uint32_t)), seq, + (o->source_fill == SOURCE_FILL_GPU_HDP || + o->source_fill == SOURCE_FILL_GPU_HDP_SYNC) ? + src->hdp_flush_ptr : + nullptr); + hipError_t hret = hipGetLastError(); + if (hret != hipSuccess) { + fprintf(stderr, "source fill launch failed: %s\n", + hipGetErrorString(hret)); + return -1; + } + if (wait_source_ready(src, seq, o->timeout_ms)) + return -1; + if (o->source_fill == SOURCE_FILL_GPU_HOST_HDP) { + __atomic_store_n(src->hdp_flush_ptr, 1u, __ATOMIC_SEQ_CST); + __sync_synchronize(); + } + if (o->source_fill == SOURCE_FILL_GPU_SYNC || + o->source_fill == SOURCE_FILL_GPU_HDP_SYNC) { + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "source fill sync failed: %s\n", + hipGetErrorString(hret)); + return -1; + } + } + return 0; +} + +static int post_write_pair(struct ibv_qp *qp, struct ibv_mr *payload_mr, + struct ibv_mr *signal_mr, uint8_t *payload, + uint64_t *signal, size_t size, + uint64_t remote_payload, uint32_t remote_rkey, + size_t remote_signal_offset, uint32_t seq) +{ + struct ibv_sge payload_sge; + struct ibv_sge signal_sge; + struct ibv_send_wr payload_wr; + struct ibv_send_wr signal_wr; + struct ibv_send_wr *bad = NULL; + + *signal = seq; + memset(&payload_sge, 0, sizeof(payload_sge)); + payload_sge.addr = (uintptr_t)payload; + payload_sge.length = (uint32_t)size; + payload_sge.lkey = payload_mr->lkey; + memset(&signal_sge, 0, sizeof(signal_sge)); + signal_sge.addr = (uintptr_t)signal; + signal_sge.length = sizeof(*signal); + signal_sge.lkey = signal_mr->lkey; + + memset(&payload_wr, 0, sizeof(payload_wr)); + payload_wr.wr_id = ((uint64_t)seq << 1); + payload_wr.sg_list = &payload_sge; + payload_wr.num_sge = 1; + payload_wr.opcode = IBV_WR_RDMA_WRITE; + payload_wr.wr.rdma.remote_addr = remote_payload; + payload_wr.wr.rdma.rkey = remote_rkey; + payload_wr.next = &signal_wr; + + memset(&signal_wr, 0, sizeof(signal_wr)); + signal_wr.wr_id = ((uint64_t)seq << 1) | 1u; + signal_wr.sg_list = &signal_sge; + signal_wr.num_sge = 1; + signal_wr.opcode = IBV_WR_RDMA_WRITE; + signal_wr.send_flags = IBV_SEND_SIGNALED | IBV_SEND_FENCE; + signal_wr.wr.rdma.remote_addr = remote_payload + remote_signal_offset; + signal_wr.wr.rdma.rkey = remote_rkey; + + return ibv_post_send(qp, &payload_wr, &bad); +} + +static int run_receiver(struct opts *o, int sock, HipRegion *region, + uint32_t rkey, uint32_t qpn) +{ + char ready = 'R'; + char ack; + hipError_t hret; + uint64_t start_ns; + int ret = 1; + + printf("recv_region kind=%s payload=%p signal=%p state=%p bytes=%zu signal_offset=%zu state_offset=%zu rkey=0x%x local_qpn=%u\n", + o->kind, region->payload, region->signal, region->state, + region->bytes, region->signal_offset, region->state_offset, + rkey, qpn); + + hret = hipMemset(region->state, 0, sizeof(*region->state)); + if (hret == hipSuccess) + hret = hipMemset(region->payload, 0, o->size); + if (hret == hipSuccess) + hret = hipMemset(region->signal, 0, sizeof(*region->signal)); + if (hret == hipSuccess) + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "receiver hip init failed: %s\n", + hipGetErrorString(hret)); + return 1; + } + + hipLaunchKernelGGL(rdma_write_visibility_kernel, dim3(1), dim3(256), + 0, 0, region->payload, region->signal, region->state, + (uint32_t)(o->size / sizeof(uint32_t)), o->count, + o->mode, 100000000u); + hret = hipGetLastError(); + if (hret != hipSuccess) { + fprintf(stderr, "kernel launch failed: %s\n", + hipGetErrorString(hret)); + return 1; + } + + if (send_all(sock, &ready, 1)) { + perror("ready send"); + goto out_sync; + } + + start_ns = now_ns(); + for (uint32_t seq = 1; seq <= o->count; ++seq) { + if (wait_gpu_seen(region, seq, o->timeout_ms)) + goto out_sync; + ack = 'A'; + if (send_all(sock, &ack, 1)) { + perror("ack send"); + goto out_sync; + } + } + + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "hipDeviceSynchronize failed: %s\n", + hipGetErrorString(hret)); + goto out_sync; + } + + { + double secs = (double)(now_ns() - start_ns) / 1000000000.0; + printf("recv_result kind=%s mode=%s size=%zu count=%u status=OK elapsed_sec=%.6f avg_us=%.3f gpu_spins=%" PRIu64 " prefetch_sink=0x%x\n", + o->kind, o->mode_name, o->size, o->count, secs, + secs * 1000000.0 / (double)o->count, + region->state->gpu_spins, region->state->prefetch_sink); + } + ret = 0; + +out_sync: + if (ret) + (void)hipDeviceSynchronize(); + return ret; +} + +static int run_sender(struct opts *o, int sock, struct ibv_pd *pd, + struct ibv_qp *qp, struct ibv_cq *cq, + const struct peer_info *remote) +{ + SenderSource src; + char ready; + char ack; + size_t signal_offset = align_up(o->size, 64); + uint64_t start_ns; + int ret = 1; + + if (init_sender_source(o, pd, &src)) + return 1; + + if (recv_all(sock, &ready, 1) || ready != 'R') { + fprintf(stderr, "receiver did not signal ready\n"); + goto out_src; + } + + start_ns = now_ns(); + for (uint32_t seq = 1; seq <= o->count; ++seq) { + if (prepare_source_payload(&src, o, seq)) + goto out_src; + if (post_write_pair(qp, src.payload_mr, src.signal_mr, + src.payload, src.signal, o->size, + remote->addr, remote->rkey, signal_offset, + seq)) { + perror("ibv_post_send RDMA_WRITE"); + goto out_src; + } + if (poll_send_cq(cq, o->timeout_ms)) + goto out_src; + if (recv_all(sock, &ack, 1) || ack != 'A') { + fprintf(stderr, "receiver did not ack seq=%u\n", seq); + goto out_src; + } + } + + { + double secs = (double)(now_ns() - start_ns) / 1000000000.0; + printf("send_result source_kind=%s source_fill=%s source_reg=%s size=%zu count=%u status=OK elapsed_sec=%.6f avg_us=%.3f remote_addr=0x%016" PRIx64 " remote_rkey=0x%x\n", + o->source_kind, o->source_fill_name, o->source_reg_name, + o->size, o->count, secs, + secs * 1000000.0 / (double)o->count, + remote->addr, remote->rkey); + } + ret = 0; + +out_src: + cleanup_sender_source(&src); + return ret; +} + +int main(int argc, char **argv) +{ + struct opts o; + struct ibv_context *ctx = NULL; + struct ibv_pd *pd = NULL; + struct ibv_cq *cq = NULL; + struct ibv_qp *qp = NULL; + struct ibv_qp_init_attr qp_attr; + ReceiverMr recv_mr = {}; + HipRegion recv_region = {}; + bool recv_region_allocated = false; + union ibv_gid local_gid, remote_gid; + struct ibv_port_attr port_attr; + struct peer_info local, remote; + uint32_t psn; + int listen_fd = -1; + int sock = -1; + int ret = 1; + + if (parse_opts(argc, argv, &o)) { + usage(argv[0]); + return 2; + } + if (!strcmp(o.role, "send") && !o.connect_host) { + usage(argv[0]); + return 2; + } + if (hipSetDevice(0) != hipSuccess) { + fprintf(stderr, "hipSetDevice(0) failed\n"); + return 1; + } + + ctx = open_dev(o.dev); + if (!ctx) { + fprintf(stderr, "ibv_open_device(%s): %s\n", o.dev, + strerror(errno)); + goto out; + } + if (ibv_query_port(ctx, o.ib_port, &port_attr)) { + perror("ibv_query_port"); + goto out; + } + if (ibv_query_gid(ctx, o.ib_port, o.gid_index, &local_gid)) { + perror("ibv_query_gid"); + goto out; + } + pd = ibv_alloc_pd(ctx); + if (!pd) { + perror("ibv_alloc_pd"); + goto out; + } + cq = ibv_create_cq(ctx, 32, NULL, NULL, 0); + if (!cq) { + perror("ibv_create_cq"); + goto out; + } + memset(&qp_attr, 0, sizeof(qp_attr)); + qp_attr.send_cq = cq; + qp_attr.recv_cq = cq; + qp_attr.qp_type = IBV_QPT_RC; + qp_attr.cap.max_send_wr = 32; + qp_attr.cap.max_recv_wr = 1; + qp_attr.cap.max_send_sge = 1; + qp_attr.cap.max_recv_sge = 1; + qp = ibv_create_qp(pd, &qp_attr); + if (!qp) { + perror("ibv_create_qp"); + goto out; + } + + psn = (uint32_t)(now_ns() & 0xffffffu); + memset(&local, 0, sizeof(local)); + local.magic = VIS_MAGIC; + local.qpn = qp->qp_num; + local.psn = psn; + local.lid = port_attr.lid; + memcpy(local.gid, local_gid.raw, sizeof(local.gid)); + + if (!strcmp(o.role, "recv")) { + hipError_t hret = alloc_region(o.kind, o.size, &recv_region); + + if (hret != hipSuccess) { + fprintf(stderr, "receiver region alloc failed: %s\n", + hipGetErrorString(hret)); + goto out; + } + recv_region_allocated = true; + if (init_receiver_mr(&o, pd, &recv_region, &recv_mr)) + goto out; + local.addr = (uintptr_t)recv_region.payload; + local.rkey = recv_mr.mr->rkey; + + listen_fd = tcp_listen(o.port); + if (listen_fd < 0) { + perror("tcp_listen"); + goto out; + } + printf("recv_listen port=%d kind=%s reg=%s mode=%s size=%zu count=%u qpn=%u\n", + o.port, o.kind, o.recv_reg_name, o.mode_name, o.size, + o.count, local.qpn); + sock = accept(listen_fd, NULL, NULL); + if (sock < 0) { + perror("accept"); + goto out; + } + } else { + sock = tcp_connect(o.connect_host, o.port); + if (sock < 0) { + perror("tcp_connect"); + goto out; + } + printf("send_connect host=%s port=%d size=%zu count=%u qpn=%u\n", + o.connect_host, o.port, o.size, o.count, local.qpn); + } + + if (exchange_info(sock, &local, &remote)) + goto out; + memcpy(remote_gid.raw, remote.gid, sizeof(remote_gid.raw)); + if (qp_to_rts(qp, o.ib_port, o.gid_index, &remote_gid, remote.lid, + remote.qpn, local.psn, remote.psn, mtu_enum(o.mtu), + IBV_ACCESS_REMOTE_WRITE)) + goto out; + + if (!strcmp(o.role, "recv")) + ret = run_receiver(&o, sock, &recv_region, recv_mr.mr->rkey, + qp->qp_num); + else + ret = run_sender(&o, sock, pd, qp, cq, &remote); + +out: + cleanup_receiver_mr(&recv_mr); + if (recv_region_allocated) + free_region(&recv_region); + if (sock >= 0) + close(sock); + if (listen_fd >= 0) + close(listen_fd); + if (qp && ibv_destroy_qp(qp)) + fprintf(stderr, "ibv_destroy_qp: %s\n", strerror(errno)); + if (cq && ibv_destroy_cq(cq)) + fprintf(stderr, "ibv_destroy_cq: %s\n", strerror(errno)); + if (pd && ibv_dealloc_pd(pd)) + fprintf(stderr, "ibv_dealloc_pd: %s\n", strerror(errno)); + if (ctx && ibv_close_device(ctx)) + fprintf(stderr, "ibv_close_device: %s\n", strerror(errno)); + return ret; +} diff --git a/userspace/bench/hip_reg_mr_probe.cpp b/userspace/bench/hip_reg_mr_probe.cpp new file mode 100644 index 0000000..74b60fd --- /dev/null +++ b/userspace/bench/hip_reg_mr_probe.cpp @@ -0,0 +1,175 @@ +// SPDX-License-Identifier: MIT +/* + * Probe whether libibverbs can register ROCm allocation types with usb4_rdma. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +static struct ibv_context *open_context(const char *name) +{ + struct ibv_device **list; + struct ibv_context *ctx = NULL; + int n = 0; + + list = ibv_get_device_list(&n); + if (!list) + return NULL; + for (int i = 0; i < n; i++) { + if (!strcmp(ibv_get_device_name(list[i]), name)) { + ctx = ibv_open_device(list[i]); + break; + } + } + ibv_free_device_list(list); + return ctx; +} + +int main(int argc, char **argv) +{ + const char *dev = "usb4_rdma0"; + const char *kind = argc > 1 ? argv[1] : "device"; + size_t size = argc > 2 ? strtoull(argv[2], NULL, 0) : 4096; + bool gpu_touch = false; + bool use_dmabuf = false; + unsigned int host_flags = hipHostMallocMapped; + struct ibv_context *ctx; + struct ibv_pd *pd; + struct ibv_mr *mr; + void *ptr = NULL; + hipError_t hret; + int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | + IBV_ACCESS_REMOTE_READ; + int ret = 1; + + for (int i = 3; i < argc; i++) { + if (!strcmp(argv[i], "gpu-touch")) { + gpu_touch = true; + } else if (!strcmp(argv[i], "dmabuf")) { + use_dmabuf = true; + } else { + fprintf(stderr, + "usage: %s device|managed|host|host-coherent|host-noncoherent|host-uncached [bytes] [gpu-touch] [dmabuf]\n", + argv[0]); + return 2; + } + } + + hret = hipSetDevice(0); + if (hret != hipSuccess) { + fprintf(stderr, "hipSetDevice failed: %s\n", + hipGetErrorString(hret)); + return 1; + } + + if (!strcmp(kind, "device")) { + hret = hipMalloc(&ptr, size); + } else if (!strcmp(kind, "managed")) { + hret = hipMallocManaged(&ptr, size); + } else if (!strcmp(kind, "host")) { + hret = hipHostMalloc(&ptr, size, host_flags); + } else if (!strcmp(kind, "host-coherent")) { + host_flags |= hipHostMallocCoherent; + hret = hipHostMalloc(&ptr, size, host_flags); + } else if (!strcmp(kind, "host-noncoherent")) { + host_flags |= hipHostMallocNonCoherent; + hret = hipHostMalloc(&ptr, size, host_flags); + } else if (!strcmp(kind, "host-uncached")) { + host_flags |= hipHostMallocUncached; + hret = hipHostMalloc(&ptr, size, host_flags); + } else { + fprintf(stderr, + "usage: %s device|managed|host|host-coherent|host-noncoherent|host-uncached [bytes] [gpu-touch] [dmabuf]\n", + argv[0]); + return 2; + } + if (hret != hipSuccess) { + fprintf(stderr, "hip allocation kind=%s failed: %s\n", kind, + hipGetErrorString(hret)); + return 1; + } + if (gpu_touch) { + hret = hipMemset(ptr, 0x5a, size); + if (hret == hipSuccess) + hret = hipDeviceSynchronize(); + if (hret != hipSuccess) { + fprintf(stderr, "gpu-touch kind=%s failed: %s\n", kind, + hipGetErrorString(hret)); + goto out_free; + } + } + + ctx = open_context(dev); + if (!ctx) { + fprintf(stderr, "open_context(%s) failed\n", dev); + goto out_free; + } + pd = ibv_alloc_pd(ctx); + if (!pd) { + perror("ibv_alloc_pd"); + goto out_close; + } + + errno = 0; + if (use_dmabuf) { + uint64_t offset = 0; + int fd = -1; + hsa_status_t hsa_ret; + + hsa_ret = hsa_amd_portable_export_dmabuf(ptr, size, &fd, + &offset); + if (hsa_ret != HSA_STATUS_SUCCESS) { + printf("kind=%s method=dmabuf gpu_touch=%d ptr=%p size=%zu export=FAIL hsa_status=%d\n", + kind, gpu_touch, ptr, size, hsa_ret); + goto out_pd; + } + + mr = ibv_reg_dmabuf_mr(pd, offset, size, (uint64_t)ptr, fd, + access); + if (!mr) { + printf("kind=%s method=dmabuf gpu_touch=%d ptr=%p size=%zu fd=%d offset=%llu reg_mr=FAIL errno=%d %s\n", + kind, gpu_touch, ptr, size, fd, + (unsigned long long)offset, errno, + strerror(errno)); + (void)hsa_amd_portable_close_dmabuf(fd); + goto out_pd; + } + + printf("kind=%s method=dmabuf gpu_touch=%d ptr=%p size=%zu fd=%d offset=%llu reg_mr=OK lkey=0x%x rkey=0x%x\n", + kind, gpu_touch, ptr, size, fd, + (unsigned long long)offset, mr->lkey, mr->rkey); + ibv_dereg_mr(mr); + (void)hsa_amd_portable_close_dmabuf(fd); + ret = 0; + goto out_pd; + } + + mr = ibv_reg_mr(pd, ptr, size, access); + if (!mr) { + printf("kind=%s method=reg_mr gpu_touch=%d ptr=%p size=%zu reg_mr=FAIL errno=%d %s\n", + kind, gpu_touch, ptr, size, errno, strerror(errno)); + } else { + printf("kind=%s method=reg_mr gpu_touch=%d ptr=%p size=%zu reg_mr=OK lkey=0x%x rkey=0x%x\n", + kind, gpu_touch, ptr, size, mr->lkey, mr->rkey); + ibv_dereg_mr(mr); + ret = 0; + } + +out_pd: + ibv_dealloc_pd(pd); +out_close: + ibv_close_device(ctx); +out_free: + if (!strncmp(kind, "host", 4)) + (void)hipHostFree(ptr); + else + (void)hipFree(ptr); + return ret; +} diff --git a/userspace/bench/hsa_fine_grain_pool_probe.cpp b/userspace/bench/hsa_fine_grain_pool_probe.cpp new file mode 100644 index 0000000..af21a4b --- /dev/null +++ b/userspace/bench/hsa_fine_grain_pool_probe.cpp @@ -0,0 +1,421 @@ +// SPDX-License-Identifier: MIT +/* + * Probe HSA fine-grained memory pools for the USB4 GDA queue-memory contract. + * + * The ABI can use HSA fine-grained allocations only if their runtime alignment + * is compatible with the SQ/CQ/doorbell layout and if usb4_rdma can register + * them as MRs. This program prints pool properties and attempts ibv_reg_mr() + * for candidate fine-grained allocations. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +struct AgentInfo { + hsa_agent_t agent; + char name[64]; + hsa_device_type_t type; +}; + +struct PoolInfo { + hsa_amd_memory_pool_t pool; + size_t owner_index; + hsa_amd_segment_t segment; + uint32_t flags; + bool alloc_allowed; + bool accessible_by_all; + hsa_amd_memory_pool_location_t location; + size_t alloc_granule; + size_t rec_granule; + size_t alloc_alignment; + size_t max_size; +}; + +struct PoolEnumCtx { + const std::vector *agents; + size_t owner_index; + std::vector *pools; +}; + +static const char *hsa_status_name(hsa_status_t status) +{ + const char *name = NULL; + + if (hsa_status_string(status, &name) != HSA_STATUS_SUCCESS || !name) + return "unknown"; + return name; +} + +static const char *agent_type_name(hsa_device_type_t type) +{ + switch (type) { + case HSA_DEVICE_TYPE_CPU: + return "cpu"; + case HSA_DEVICE_TYPE_GPU: + return "gpu"; + case HSA_DEVICE_TYPE_DSP: + return "dsp"; + case HSA_DEVICE_TYPE_AIE: + return "aie"; + default: + return "unknown"; + } +} + +static const char *segment_name(hsa_amd_segment_t segment) +{ + switch (segment) { + case HSA_AMD_SEGMENT_GLOBAL: + return "global"; + case HSA_AMD_SEGMENT_READONLY: + return "readonly"; + case HSA_AMD_SEGMENT_PRIVATE: + return "private"; + case HSA_AMD_SEGMENT_GROUP: + return "group"; + default: + return "unknown"; + } +} + +static const char *location_name(hsa_amd_memory_pool_location_t location) +{ + switch (location) { + case HSA_AMD_MEMORY_POOL_LOCATION_CPU: + return "cpu"; + case HSA_AMD_MEMORY_POOL_LOCATION_GPU: + return "gpu"; + default: + return "unknown"; + } +} + +static const char *access_name(hsa_amd_memory_pool_access_t access) +{ + switch (access) { + case HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED: + return "never"; + case HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT: + return "default"; + case HSA_AMD_MEMORY_POOL_ACCESS_DISALLOWED_BY_DEFAULT: + return "allow_access"; + default: + return "unknown"; + } +} + +static bool pool_is_candidate(const PoolInfo &pool) +{ + if (pool.segment != HSA_AMD_SEGMENT_GLOBAL || !pool.alloc_allowed) + return false; + return pool.flags & (HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED | + HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED); +} + +static size_t round_up(size_t value, size_t granule) +{ + if (!granule) + return value; + size_t rem = value % granule; + return rem ? value + granule - rem : value; +} + +static struct ibv_context *open_context(const char *name) +{ + struct ibv_device **list; + struct ibv_context *ctx = NULL; + int n = 0; + + list = ibv_get_device_list(&n); + if (!list) + return NULL; + + for (int i = 0; i < n; i++) { + if (!strcmp(ibv_get_device_name(list[i]), name)) { + ctx = ibv_open_device(list[i]); + break; + } + } + + ibv_free_device_list(list); + return ctx; +} + +static hsa_status_t collect_agent(hsa_agent_t agent, void *data) +{ + std::vector *agents = + static_cast *>(data); + AgentInfo info = {}; + hsa_status_t status; + + info.agent = agent; + status = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, info.name); + if (status != HSA_STATUS_SUCCESS) + snprintf(info.name, sizeof(info.name), "unknown"); + info.name[sizeof(info.name) - 1] = '\0'; + + status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &info.type); + if (status != HSA_STATUS_SUCCESS) + info.type = static_cast(-1); + + agents->push_back(info); + return HSA_STATUS_SUCCESS; +} + +static hsa_status_t collect_pool(hsa_amd_memory_pool_t pool, void *data) +{ + PoolEnumCtx *ctx = static_cast(data); + PoolInfo info = {}; + hsa_status_t status; + + info.pool = pool; + info.owner_index = ctx->owner_index; + + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &info.segment); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "pool=0x%016" PRIx64 " query=segment status=%s\n", + pool.handle, hsa_status_name(status)); + return HSA_STATUS_SUCCESS; + } + + if (info.segment == HSA_AMD_SEGMENT_GLOBAL) { + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, + &info.flags); + if (status != HSA_STATUS_SUCCESS) + info.flags = 0; + } + + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, + &info.alloc_allowed); + if (status != HSA_STATUS_SUCCESS) + info.alloc_allowed = false; + + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALL, + &info.accessible_by_all); + if (status != HSA_STATUS_SUCCESS) + info.accessible_by_all = false; + + status = hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_LOCATION, &info.location); + if (status != HSA_STATUS_SUCCESS) + info.location = static_cast(-1); + + if (info.alloc_allowed) { + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, + &info.alloc_granule); + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_REC_GRANULE, + &info.rec_granule); + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT, + &info.alloc_alignment); + (void)hsa_amd_memory_pool_get_info( + pool, HSA_AMD_MEMORY_POOL_INFO_ALLOC_MAX_SIZE, + &info.max_size); + } + + ctx->pools->push_back(info); + return HSA_STATUS_SUCCESS; +} + +static void print_pool_access(const PoolInfo &pool, + const std::vector &agents) +{ + for (size_t i = 0; i < agents.size(); i++) { + hsa_amd_memory_pool_access_t access = + HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED; + hsa_status_t status = hsa_amd_agent_memory_pool_get_info( + agents[i].agent, pool.pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + + printf("pool_access pool=0x%016" PRIx64 " agent=%zu name=\"%s\" access=%s status=%s\n", + pool.pool.handle, i, agents[i].name, + status == HSA_STATUS_SUCCESS ? access_name(access) : "query-fail", + hsa_status_name(status)); + } +} + +static void probe_allocation(const PoolInfo &pool, + const std::vector &agents, + struct ibv_pd *pd, size_t requested_size) +{ + std::vector allowed_agents; + void *ptr = NULL; + size_t size = round_up(requested_size, pool.alloc_granule); + hsa_status_t status; + + for (const AgentInfo &agent : agents) { + hsa_amd_memory_pool_access_t access = + HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED; + + if (agent.type != HSA_DEVICE_TYPE_CPU && + agent.type != HSA_DEVICE_TYPE_GPU) + continue; + + status = hsa_amd_agent_memory_pool_get_info( + agent.agent, pool.pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access); + if (status == HSA_STATUS_SUCCESS && + access != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) + allowed_agents.push_back(agent.agent); + } + + status = hsa_amd_memory_pool_allocate(pool.pool, size, 0, &ptr); + if (status != HSA_STATUS_SUCCESS) { + printf("alloc_probe pool=0x%016" PRIx64 " owner=%zu requested=%zu size=%zu alloc=FAIL status=%s\n", + pool.pool.handle, pool.owner_index, requested_size, size, + hsa_status_name(status)); + return; + } + + status = hsa_amd_agents_allow_access( + static_cast(allowed_agents.size()), + allowed_agents.data(), NULL, ptr); + + uintptr_t addr = reinterpret_cast(ptr); + printf("alloc_probe pool=0x%016" PRIx64 " owner=%zu grain=fine extended=%u ptr=%p requested=%zu size=%zu alloc_alignment=%zu ptr_mod_alignment=%zu ptr_mod_4096=%zu allow_scope=cpu_gpu allow_agents=%zu allow_access=%s", + pool.pool.handle, pool.owner_index, + !!(pool.flags & + HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED), + ptr, requested_size, size, pool.alloc_alignment, + pool.alloc_alignment ? addr % pool.alloc_alignment : 0, + addr % 4096, allowed_agents.size(), hsa_status_name(status)); + + if (pd) { + int access = IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | + IBV_ACCESS_REMOTE_READ; + errno = 0; + struct ibv_mr *mr = ibv_reg_mr(pd, ptr, size, access); + if (mr) { + printf(" reg_mr=OK lkey=0x%x rkey=0x%x\n", mr->lkey, + mr->rkey); + (void)ibv_dereg_mr(mr); + } else { + printf(" reg_mr=FAIL errno=%d %s\n", errno, + strerror(errno)); + } + + int dmabuf_fd = -1; + uint64_t dmabuf_offset = 0; + status = hsa_amd_portable_export_dmabuf(ptr, size, &dmabuf_fd, + &dmabuf_offset); + if (status != HSA_STATUS_SUCCESS) { + printf(" dmabuf_export=FAIL status=%s\n", + hsa_status_name(status)); + } else { + errno = 0; + mr = ibv_reg_dmabuf_mr(pd, dmabuf_offset, size, + reinterpret_cast(ptr), + dmabuf_fd, access); + if (mr) { + printf(" dmabuf_export=OK fd=%d offset=%" PRIu64 " reg_dmabuf=OK lkey=0x%x rkey=0x%x\n", + dmabuf_fd, dmabuf_offset, mr->lkey, + mr->rkey); + (void)ibv_dereg_mr(mr); + } else { + printf(" dmabuf_export=OK fd=%d offset=%" PRIu64 " reg_dmabuf=FAIL errno=%d %s\n", + dmabuf_fd, dmabuf_offset, errno, + strerror(errno)); + } + (void)hsa_amd_portable_close_dmabuf(dmabuf_fd); + } + } else { + printf(" reg_mr=SKIP\n"); + } + + (void)hsa_amd_memory_pool_free(ptr); +} + +int main(int argc, char **argv) +{ + const char *dev = argc > 1 ? argv[1] : "usb4_rdma0"; + size_t requested_size = argc > 2 ? strtoull(argv[2], NULL, 0) : 4096; + std::vector agents; + std::vector pools; + struct ibv_context *verbs = NULL; + struct ibv_pd *pd = NULL; + hsa_status_t status; + + status = hsa_init(); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_init failed: %s\n", hsa_status_name(status)); + return 1; + } + + status = hsa_iterate_agents(collect_agent, &agents); + if (status != HSA_STATUS_SUCCESS) { + fprintf(stderr, "hsa_iterate_agents failed: %s\n", + hsa_status_name(status)); + (void)hsa_shut_down(); + return 1; + } + + printf("requested_size=%zu rdma_device=%s agents=%zu\n", requested_size, + dev, agents.size()); + for (size_t i = 0; i < agents.size(); i++) { + printf("agent index=%zu handle=0x%016" PRIx64 " type=%s name=\"%s\"\n", + i, agents[i].agent.handle, agent_type_name(agents[i].type), + agents[i].name); + + PoolEnumCtx ctx = { &agents, i, &pools }; + status = hsa_amd_agent_iterate_memory_pools(agents[i].agent, + collect_pool, &ctx); + if (status != HSA_STATUS_SUCCESS) + fprintf(stderr, + "agent=%zu iterate_memory_pools failed: %s\n", + i, hsa_status_name(status)); + } + + if (strcmp(dev, "none")) { + verbs = open_context(dev); + if (!verbs) { + fprintf(stderr, "open_context(%s) failed; skipping reg_mr\n", + dev); + } else { + pd = ibv_alloc_pd(verbs); + if (!pd) + perror("ibv_alloc_pd"); + } + } + + for (const PoolInfo &pool : pools) { + const AgentInfo &owner = agents[pool.owner_index]; + bool fine = pool.flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED; + bool extended = pool.flags & + HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED; + + printf("pool owner=%zu owner_name=\"%s\" owner_type=%s handle=0x%016" PRIx64 " segment=%s location=%s flags=0x%08x fine=%u extended=%u alloc_allowed=%u alloc_granule=%zu rec_granule=%zu alloc_alignment=%zu accessible_by_all=%u max_size=%zu candidate=%u\n", + pool.owner_index, owner.name, agent_type_name(owner.type), + pool.pool.handle, segment_name(pool.segment), + location_name(pool.location), pool.flags, fine, extended, + pool.alloc_allowed, pool.alloc_granule, pool.rec_granule, + pool.alloc_alignment, pool.accessible_by_all, pool.max_size, + pool_is_candidate(pool)); + print_pool_access(pool, agents); + + if (pool_is_candidate(pool)) + probe_allocation(pool, agents, pd, requested_size); + } + + if (pd) + (void)ibv_dealloc_pd(pd); + if (verbs) + (void)ibv_close_device(verbs); + (void)hsa_shut_down(); + return 0; +} diff --git a/userspace/bench/rx_side_dma_validate.sh b/userspace/bench/rx_side_dma_validate.sh new file mode 100755 index 0000000..3971ffb --- /dev/null +++ b/userspace/bench/rx_side_dma_validate.sh @@ -0,0 +1,157 @@ +#!/usr/bin/env bash +# RX-side DMA + link-layer validation for codex/gda-v2-rebased-port @ 75cf776. +# +# Stages: +# 1. Copy the built module to strix-1, reload it. +# 2. Confirm usb4_rdma ports report IB_LINK_LAYER_INFINIBAND. +# 3. CPU regression: rc_write_verify (RC WRITE bounce-buffer path unchanged). +# 4. DMABUF receiver smoke: dmabuf_mr_probe + HIP receiver probe; verify +# data_rx_dmabuf_zcopy increments and data_rx_dmabuf_zcopy_error stays 0. +# +# Usage: ./userspace/bench/rx_side_dma_validate.sh +# HOSTS = "strix-1 strix-2" (server is HOSTS[0]) +# BENCH_ROOT = /mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase +# OUT_ROOT = /mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase/.bench-artifacts/rx-side-dma +# +# Exit codes: +# 0 = all stages pass +# 1 = usage / config error +# 2 = module load / link-layer / counter check failure +# 3 = regression test failure + +set -euo pipefail + +HOSTS=${HOSTS:-"strix-1 strix-2"} +BENCH_ROOT=${BENCH_ROOT:-/mnt/Home/src/thunderbolt-ibverbs-gda-v2-rebase} +OUT_ROOT=${OUT_ROOT:-${BENCH_ROOT}/.bench-artifacts/rx-side-dma} +RUN_ID=${RUN_ID:-$(date -u +%Y%m%dT%H%M%SZ)} +OUT_DIR=${OUT_DIR:-${OUT_ROOT}/${RUN_ID}} +SSH_OPTS=${SSH_OPTS:-"-o ConnectTimeout=5 -o BatchMode=yes -o ServerAliveInterval=5 -o ServerAliveCountMax=3"} +SERVER=${HOSTS%% *} +CLIENT=${HOSTS##* } + +if [ "$SERVER" = "$CLIENT" ]; then + echo "HOSTS must list at least 2 nodes (server, client). Got: $HOSTS" >&2 + exit 1 +fi + +stage_reload() { + log "stage=reload host=$SERVER" + "${BENCH_ROOT}/tools/tbv-target-module.sh" "$SERVER" --booted-kernel --reload --options 'profile=linux_perf tbnet=prefer_rdma bind_services=1' 2>&1 | tee "$OUT_DIR/reload-${SERVER}.log" + # Wait for the new usb4_rdma devices to register; on a clean boot the + # source-aware XDomain handler initialises in time, on a reload the kernel + # autoload path is required (which is why the doc recommends a reboot). + for _ in $(seq 1 30); do + if ssh $SSH_OPTS "$SERVER" 'test -n "$(ls /sys/class/infiniband/usb4_rdma* 2>/dev/null)"' 2>/dev/null; then + break + fi + sleep 1 + done +} + +############################################################################## +# Stage 1: copy + reload module on SERVER +############################################################################## +stage_reload() { + log "stage=reload host=$SERVER" + "${BENCH_ROOT}/tools/tbv-target-module.sh" "$SERVER" --booted-kernel --copy 2>&1 | tee "$OUT_DIR/reload-${SERVER}.log" + ssh $SSH_OPTS "$SERVER" 'sudo rmmod thunderbolt_ibverbs 2>/dev/null || true; sudo modprobe thunderbolt_ibverbs' 2>&1 | tee -a "$OUT_DIR/reload-${SERVER}.log" +} + +############################################################################## +# Stage 2: confirm link_layer is InfiniBand on SERVER (and CLIENT) +############################################################################## +stage_link_layer() { + log "stage=link_layer hosts=$HOSTS" + local fail=0 + for h in $HOSTS; do + local ll + ll=$(ssh $SSH_OPTS "$h" 'for d in /sys/class/infiniband/usb4_rdma*; do cat "$d/ports/1/link_layer"; done' 2>/dev/null) || { log "host=$h ssh failed"; fail=1; continue; } + log "host=$h link_layer=$ll" + if [ "$ll" != "InfiniBand" ]; then + log "FAIL host=$h expected InfiniBand, got '$ll'" + fail=1 + fi + done + return $fail +} + +############################################################################## +# Stage 3: CPU regression on SERVER/CLIENT (rc_write_verify) +############################################################################## +stage_cpu_regression() { + log "stage=cpu_regression server=$SERVER client=$CLIENT" + ssh $SSH_OPTS "$SERVER" "cd $BENCH_ROOT && result/bin/rc_write_verify -d usb4_rdma0 -s 4096 -c 1000" 2>&1 \ + | tee "$OUT_DIR/rc_write_verify-server.log" & + local server_pid=$! + sleep 1 + ssh $SSH_OPTS "$CLIENT" "cd $BENCH_ROOT && result/bin/rc_write_verify -d usb4_rdma0 -s 4096 -c 1000 \$(ssh $SSH_OPTS $SERVER 'cat /sys/class/infiniband/usb4_rdma0/ports/1/gid_attrs/* 2>/dev/null | head -1 || true' &1 \ + | tee "$OUT_DIR/rc_write_verify-client.log" & + local client_pid=$! + wait $server_pid + local server_rc=$? + wait $client_pid + local client_rc=$? + log "rc_write_verify server_rc=$server_rc client_rc=$client_rc" + grep -E 'data_tx_ack_error|data_rx_copy_error|data_rx_dmabuf_zcopy_error|cqe_error|completion_error' "$OUT_DIR/rc_write_verify-server.log" "$OUT_DIR/rc_write_verify-client.log" 2>/dev/null | head -20 + return $(( server_rc + client_rc )) +} + +############################################################################## +# Stage 4: DMABUF receiver smoke +############################################################################## +stage_dmabuf_rx() { + log "stage=dmabuf_rx server=$SERVER client=$CLIENT" + # 4a: import probe on the receiver + ssh $SSH_OPTS "$SERVER" "cd $BENCH_ROOT && result/bin/dmabuf_mr_probe -d usb4_rdma0 --length 4096 --fd 5" 2>&1 \ + | tee "$OUT_DIR/dmabuf_mr_probe.log" + # 4b: before/after counter snapshot + local before after + before=$(ssh $SSH_OPTS "$SERVER" 'cat /sys/kernel/debug/tbv/data_rx_dmabuf_zcopy 2>/dev/null || echo 0') + log "data_rx_dmabuf_zcopy(before)=$before" + # 4c: run HIP receiver probe (if available) + if ssh $SSH_OPTS "$SERVER" "test -x $BENCH_ROOT/result/bin/hip_rdma_write_visibility_probe" 2>/dev/null; then + log "running HIP RDMA write visibility probe on $SERVER" + ssh $SSH_OPTS "$SERVER" "cd $BENCH_ROOT && result/bin/hip_rdma_write_visibility_probe --recv-reg dmabuf -d usb4_rdma0 -s 4096 -c 10" 2>&1 \ + | tee "$OUT_DIR/hip_recv.log" + else + log "HIP probe not available at $BENCH_ROOT/result/bin/hip_rdma_write_visibility_probe; skipping live RDMA write into dmabuf" + fi + after=$(ssh $SSH_OPTS "$SERVER" 'cat /sys/kernel/debug/tbv/data_rx_dmabuf_zcopy 2>/dev/null || echo 0') + log "data_rx_dmabuf_zcopy(after)=$after" + if [ "$after" -le "$before" ]; then + log "WARN data_rx_dmabuf_zcopy did not increment ($before -> $after). Check that a dmabuf receiver ran." + fi + local errs + errs=$(ssh $SSH_OPTS "$SERVER" 'cat /sys/kernel/debug/tbv/data_rx_dmabuf_zcopy_error 2>/dev/null || echo 0') + log "data_rx_dmabuf_zcopy_error=$errs" + if [ "$errs" != "0" ]; then + log "FAIL data_rx_dmabuf_zcopy_error=$errs" + return 2 + fi + return 0 +} + +############################################################################## +# Driver +############################################################################## +log "out_dir=$OUT_DIR run_id=$RUN_ID" +trap 'log "aborted"; exit 130' INT TERM + +if ! stage_reload; then + log "FAIL stage=reload" + exit 2 +fi +if ! stage_link_layer; then + log "FAIL stage=link_layer" + exit 2 +fi +if ! stage_cpu_regression; then + log "FAIL stage=cpu_regression" + exit 3 +fi +if ! stage_dmabuf_rx; then + log "FAIL stage=dmabuf_rx" + exit 2 +fi +log "PASS all stages" diff --git a/userspace/bench/tbv_dv_caps_probe.c b/userspace/bench/tbv_dv_caps_probe.c new file mode 100644 index 0000000..d901f56 --- /dev/null +++ b/userspace/bench/tbv_dv_caps_probe.c @@ -0,0 +1,734 @@ +// SPDX-License-Identifier: GPL-2.0 OR BSD-3-Clause +#define _POSIX_C_SOURCE 200112L +/* + * USB4 RDMA Direct Verbs (DV) QUERY_CAPS and queue smoke probe. + * + * Opens a usb4_rdma* / usb4_apple* verbs device, issues the private + * USB4_RDMA_DV_METHOD_QUERY_CAPS method via the raw RDMA_VERBS_IOCTL ABI, + * and prints the reported capabilities and queue-memory layout. With -q it + * also attaches a DV queue, KICKs signaled NOP WQEs, and checks CQEs. + * + * Uses the raw ioctl rather than rdma-core's private execute_ioctl() helper + * so the probe can be built as an ordinary test binary outside the provider + * tree. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include "usb4_rdma_dv.h" + +static bool device_name_matches(const char *name) +{ + return strncmp(name, "usb4_rdma", strlen("usb4_rdma")) == 0 || + strncmp(name, "usb4_apple", strlen("usb4_apple")) == 0; +} + +static struct ibv_device *find_device(struct ibv_device **list, int count, + const char *wanted) +{ + struct ibv_device *fallback = NULL; + int i; + + for (i = 0; i < count; i++) { + const char *name = ibv_get_device_name(list[i]); + + if (wanted && strcmp(name, wanted) == 0) + return list[i]; + if (!fallback && device_name_matches(name)) + fallback = list[i]; + } + return fallback; +} + +static int query_caps(struct ibv_context *ctx, + struct usb4_rdma_dv_query_caps_resp *resp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[1]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_QUERY_CAPS, + .num_attrs = 1, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_QUERY_CAPS_RESP, + .len = sizeof(*resp), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)resp, + }, + }, + }; + + memset(resp, 0, sizeof(*resp)); + if (ioctl(ctx->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int create_queue(struct ibv_qp *qp, + const struct usb4_rdma_dv_queue_create *req, + struct usb4_rdma_dv_queue_resp *resp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[3]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_CREATE_QUEUE, + .num_attrs = 3, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_REQ, + .len = sizeof(*req), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)req, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_CREATE_QUEUE_RESP, + .len = sizeof(*resp), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)resp, + }, + }, + }; + + memset(resp, 0, sizeof(*resp)); + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int destroy_queue(struct ibv_qp *qp) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[1]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_DESTROY_QUEUE, + .num_attrs = 1, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_DESTROY_QUEUE_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + }, + }; + + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int kick_queue(struct ibv_qp *qp, const struct usb4_rdma_dv_kick *req) +{ + struct { + struct ib_uverbs_ioctl_hdr hdr; + struct ib_uverbs_attr attrs[2]; + } cmd = { + .hdr = { + .length = sizeof(cmd), + .object_id = USB4_RDMA_DV_OBJECT_DEVICE, + .method_id = USB4_RDMA_DV_METHOD_KICK, + .num_attrs = 2, + .driver_id = RDMA_DRIVER_UNKNOWN, + }, + .attrs = { + { + .attr_id = USB4_RDMA_DV_ATTR_KICK_QP, + .flags = UVERBS_ATTR_F_MANDATORY, + .data = qp->handle, + }, + { + .attr_id = USB4_RDMA_DV_ATTR_KICK_REQ, + .len = sizeof(*req), + .flags = UVERBS_ATTR_F_MANDATORY, + .data = (uintptr_t)req, + }, + }, + }; + + if (ioctl(qp->context->cmd_fd, RDMA_VERBS_IOCTL, &cmd) < 0) + return errno; + return 0; +} + +static int alloc_aligned(size_t alignment, size_t size, void **ptr) +{ + int ret; + + *ptr = NULL; + ret = posix_memalign(ptr, alignment, size); + if (ret) + return ret; + memset(*ptr, 0, size); + return 0; +} + +static void store_release_u32(uint32_t *ptr, uint32_t value) +{ + __atomic_store_n(ptr, value, __ATOMIC_RELEASE); +} + +static uint32_t load_acquire_u32(uint32_t *ptr) +{ + return __atomic_load_n(ptr, __ATOMIC_ACQUIRE); +} + +static uint8_t next_generation(uint32_t generation) +{ + uint32_t mask = (1u << USB4_RDMA_DV_TAIL_GENERATION_BITS) - 1; + uint32_t next = (generation + 1) & mask; + + return next ? next : 1; +} + +static int expect_cqe(const char *label, struct usb4_rdma_dv_cqe *cqe, + struct usb4_rdma_dv_doorbell *doorbell, uint32_t index, + uint32_t generation, uint64_t wr_id, uint32_t status, + uint32_t opcode) +{ + uint32_t cq_tail = load_acquire_u32(&doorbell->consumer.cq_tail); + uint32_t sq_head = load_acquire_u32(&doorbell->consumer.sq_head); + + printf("%s wr_id=0x%016" PRIx64 " status=%u opcode=%u sq_head=%u cq_tail=%u qp_state=%u\n", + label, (uint64_t)cqe->wr_id, cqe->status, cqe->opcode, + usb4_rdma_dv_tail_index(sq_head), + usb4_rdma_dv_tail_index(cq_tail), + load_acquire_u32(&doorbell->consumer.qp_state)); + + if (usb4_rdma_dv_tail_generation(cq_tail) != generation || + usb4_rdma_dv_tail_index(cq_tail) != index + 1) { + fprintf(stderr, "%s: unexpected cq_tail=0x%08x\n", label, + cq_tail); + return 1; + } + if (usb4_rdma_dv_tail_generation(sq_head) != generation || + usb4_rdma_dv_tail_index(sq_head) != index + 1) { + fprintf(stderr, "%s: unexpected sq_head=0x%08x\n", label, + sq_head); + return 1; + } + if (cqe->wr_id != wr_id || cqe->status != status || + cqe->opcode != opcode) { + fprintf(stderr, + "%s: unexpected CQE wr_id=0x%016" PRIx64 " status=%u opcode=%u\n", + label, (uint64_t)cqe->wr_id, cqe->status, cqe->opcode); + return 1; + } + + return 0; +} + +static int run_kick_smoke(struct ibv_qp *qp, + struct usb4_rdma_dv_wqe *sq, + struct usb4_rdma_dv_cqe *cq, + struct usb4_rdma_dv_doorbell *doorbell, + uint32_t generation) +{ + struct usb4_rdma_dv_kick kick = {}; + uint32_t tail; + int err; + + store_release_u32(&doorbell->producer.generation, generation); + store_release_u32(&doorbell->producer.cq_head, + usb4_rdma_dv_tail_pack(0, generation)); + store_release_u32(&doorbell->producer.sq_tail, + usb4_rdma_dv_tail_pack(0, generation)); + + memset(&sq[0], 0, sizeof(sq[0])); + sq[0].opcode = USB4_RDMA_DV_WQE_NOP; + sq[0].flags = USB4_RDMA_DV_WQE_F_SIGNALED; + sq[0].wr_id = 0x4b49434b4e4f5031ULL; + sq[0].generation = generation; + __atomic_thread_fence(__ATOMIC_RELEASE); + tail = usb4_rdma_dv_tail_pack(1, generation); + store_release_u32(&doorbell->producer.sq_tail, tail); + kick.sq_tail = tail; + + err = kick_queue(qp, &kick); + if (err) { + fprintf(stderr, "KICK nop failed: %s (%d)\n", strerror(err), + err); + return 1; + } + if (expect_cqe("kick_nop", &cq[0], doorbell, 0, generation, + sq[0].wr_id, USB4_RDMA_DV_CQE_SUCCESS, + USB4_RDMA_DV_WQE_NOP)) + return 1; + store_release_u32(&doorbell->producer.cq_head, + usb4_rdma_dv_tail_pack(1, generation)); + + memset(&sq[1], 0, sizeof(sq[1])); + sq[1].opcode = USB4_RDMA_DV_WQE_NOP; + sq[1].flags = USB4_RDMA_DV_WQE_F_SIGNALED; + sq[1].wr_id = 0x4b49434b5354414cULL; + sq[1].generation = next_generation(generation); + __atomic_thread_fence(__ATOMIC_RELEASE); + tail = usb4_rdma_dv_tail_pack(2, generation); + store_release_u32(&doorbell->producer.sq_tail, tail); + kick.sq_tail = tail; + + err = kick_queue(qp, &kick); + if (err) { + fprintf(stderr, "KICK stale-gen failed: %s (%d)\n", + strerror(err), err); + return 1; + } + if (expect_cqe("kick_stale_gen", &cq[1], doorbell, 1, generation, + sq[1].wr_id, USB4_RDMA_DV_CQE_STALE_GEN, + USB4_RDMA_DV_WQE_NOP)) + return 1; + store_release_u32(&doorbell->producer.cq_head, + usb4_rdma_dv_tail_pack(2, generation)); + + return 0; +} + +static void print_caps_bitmap(uint32_t caps) +{ + const struct { + uint32_t bit; + const char *name; + } bits[] = { + { USB4_RDMA_DV_CAP_SEND, "send" }, + { USB4_RDMA_DV_CAP_SEND_IMM, "send_imm" }, + { USB4_RDMA_DV_CAP_WRITE, "write" }, + { USB4_RDMA_DV_CAP_WRITE_IMM, "write_imm" }, + { USB4_RDMA_DV_CAP_FENCE, "fence" }, + { USB4_RDMA_DV_CAP_READ, "read" }, + { USB4_RDMA_DV_CAP_ATOMIC_FETCH_ADD, "atomic_fetch_add" }, + { USB4_RDMA_DV_CAP_ATOMIC_SWAP, "atomic_swap" }, + { USB4_RDMA_DV_CAP_ATOMIC_CMP_SWAP, "atomic_cmp_swap" }, + }; + bool first = true; + size_t i; + + printf("caps=0x%08" PRIx32, caps); + if (!caps) { + printf(" (none)\n"); + return; + } + printf(" ("); + for (i = 0; i < sizeof(bits) / sizeof(bits[0]); i++) { + if (!(caps & bits[i].bit)) + continue; + printf("%s%s", first ? "" : "|", bits[i].name); + first = false; + } + printf(")\n"); +} + +static void print_doorbell_field(const char *name, const char *writer, + const char *reader, size_t offset, size_t size) +{ + printf("doorbell_field name=%s offset=%zu size=%zu writer=%s reader=%s\n", + name, offset, size, writer, reader); +} + +static void print_doorbell_layout(void) +{ + size_t producer_off = offsetof(struct usb4_rdma_dv_doorbell, producer); + size_t consumer_off = offsetof(struct usb4_rdma_dv_doorbell, consumer); + + printf("doorbell_layout record_size=%zu producer_line_offset=%zu producer_line_size=%zu consumer_line_offset=%zu consumer_line_size=%zu\n", + sizeof(struct usb4_rdma_dv_doorbell), producer_off, + sizeof(struct usb4_rdma_dv_doorbell_producer_line), + consumer_off, sizeof(struct usb4_rdma_dv_doorbell_consumer_line)); + print_doorbell_field( + "producer.sq_tail", "gpu", "kernel", + producer_off + + offsetof(struct usb4_rdma_dv_doorbell_producer_line, sq_tail), + sizeof(((struct usb4_rdma_dv_doorbell_producer_line *)0)->sq_tail)); + print_doorbell_field( + "producer.cq_head", "gpu", "kernel", + producer_off + + offsetof(struct usb4_rdma_dv_doorbell_producer_line, cq_head), + sizeof(((struct usb4_rdma_dv_doorbell_producer_line *)0)->cq_head)); + print_doorbell_field( + "producer.generation", "gpu", "kernel", + producer_off + + offsetof(struct usb4_rdma_dv_doorbell_producer_line, generation), + sizeof(((struct usb4_rdma_dv_doorbell_producer_line *)0)->generation)); + print_doorbell_field( + "consumer.sq_head", "kernel", "gpu", + consumer_off + + offsetof(struct usb4_rdma_dv_doorbell_consumer_line, sq_head), + sizeof(((struct usb4_rdma_dv_doorbell_consumer_line *)0)->sq_head)); + print_doorbell_field( + "consumer.cq_tail", "kernel", "gpu", + consumer_off + + offsetof(struct usb4_rdma_dv_doorbell_consumer_line, cq_tail), + sizeof(((struct usb4_rdma_dv_doorbell_consumer_line *)0)->cq_tail)); + print_doorbell_field( + "consumer.qp_state", "kernel", "gpu", + consumer_off + + offsetof(struct usb4_rdma_dv_doorbell_consumer_line, qp_state), + sizeof(((struct usb4_rdma_dv_doorbell_consumer_line *)0)->qp_state)); + print_doorbell_field( + "consumer.generation", "kernel", "gpu", + consumer_off + + offsetof(struct usb4_rdma_dv_doorbell_consumer_line, generation), + sizeof(((struct usb4_rdma_dv_doorbell_consumer_line *)0)->generation)); + printf("wqe_field name=generation offset=%zu size=%zu\n", + offsetof(struct usb4_rdma_dv_wqe, generation), + sizeof(((struct usb4_rdma_dv_wqe *)0)->generation)); +} + +static int run_queue_test(struct ibv_context *ctx, + const struct usb4_rdma_dv_query_caps_resp *caps) +{ + struct usb4_rdma_dv_queue_create req = {}; + struct usb4_rdma_dv_queue_resp resp = {}; + struct ibv_qp_init_attr qp_attr = {}; + struct ibv_mr *doorbell_mr = NULL; + struct ibv_mr *cq_mr = NULL; + struct ibv_mr *sq_mr = NULL; + struct ibv_qp *qp = NULL; + struct ibv_cq *cq = NULL; + struct ibv_pd *pd = NULL; + size_t doorbell_bytes = caps->doorbell_page_size; + size_t cq_bytes = (size_t)caps->default_cq_entries * caps->cqe_size; + size_t sq_bytes = (size_t)caps->default_sq_entries * caps->wqe_size; + void *doorbell = NULL; + void *cq_buf = NULL; + void *sq_buf = NULL; + int access = IBV_ACCESS_LOCAL_WRITE; + int ret = 1; + int err; + + if (caps->abi_version != USB4_RDMA_DV_ABI_VERSION) { + fprintf(stderr, "queue test skipped: unsupported ABI %u\n", + caps->abi_version); + return 2; + } + + err = alloc_aligned(caps->wqe_size, sq_bytes, &sq_buf); + if (err) { + fprintf(stderr, "alloc SQ: %s\n", strerror(err)); + goto out; + } + err = alloc_aligned(caps->cqe_size, cq_bytes, &cq_buf); + if (err) { + fprintf(stderr, "alloc CQ: %s\n", strerror(err)); + goto out; + } + err = alloc_aligned(caps->doorbell_page_size, doorbell_bytes, &doorbell); + if (err) { + fprintf(stderr, "alloc doorbell: %s\n", strerror(err)); + goto out; + } + + pd = ibv_alloc_pd(ctx); + if (!pd) { + fprintf(stderr, "ibv_alloc_pd: %s\n", strerror(errno)); + goto out; + } + sq_mr = ibv_reg_mr(pd, sq_buf, sq_bytes, access); + if (!sq_mr) { + fprintf(stderr, "ibv_reg_mr(SQ): %s\n", strerror(errno)); + goto out; + } + cq_mr = ibv_reg_mr(pd, cq_buf, cq_bytes, access); + if (!cq_mr) { + fprintf(stderr, "ibv_reg_mr(CQ): %s\n", strerror(errno)); + goto out; + } + doorbell_mr = ibv_reg_mr(pd, doorbell, doorbell_bytes, access); + if (!doorbell_mr) { + fprintf(stderr, "ibv_reg_mr(doorbell): %s\n", strerror(errno)); + goto out; + } + + cq = ibv_create_cq(ctx, caps->default_cq_entries, NULL, NULL, 0); + if (!cq) { + fprintf(stderr, "ibv_create_cq: %s\n", strerror(errno)); + goto out; + } + + qp_attr.send_cq = cq; + qp_attr.recv_cq = cq; + qp_attr.qp_type = IBV_QPT_RC; + qp_attr.cap.max_send_wr = caps->default_sq_entries; + qp_attr.cap.max_recv_wr = 1; + qp_attr.cap.max_send_sge = 1; + qp_attr.cap.max_recv_sge = 1; + qp = ibv_create_qp(pd, &qp_attr); + if (!qp) { + fprintf(stderr, "ibv_create_qp: %s\n", strerror(errno)); + goto out; + } + + req.abi_version = USB4_RDMA_DV_ABI_VERSION; + req.sq_addr = (uintptr_t)sq_buf; + req.cq_addr = (uintptr_t)cq_buf; + req.doorbell_addr = (uintptr_t)doorbell; + req.sq_entries = caps->default_sq_entries; + req.cq_entries = caps->default_cq_entries; + req.sq_stride = caps->wqe_size; + req.cq_stride = caps->cqe_size; + + err = create_queue(qp, &req, &resp); + if (err) { + fprintf(stderr, "CREATE_QUEUE failed: %s (%d)\n", + strerror(err), err); + goto out; + } + + printf("create_queue qp_num=%u generation=%u sq_entries=%u cq_entries=%u\n", + resp.qp_num, resp.generation, req.sq_entries, req.cq_entries); + + if (resp.qp_num != qp->qp_num) { + fprintf(stderr, + "CREATE_QUEUE returned qp_num=%u (expected %u)\n", + resp.qp_num, qp->qp_num); + (void)destroy_queue(qp); + goto out; + } + if (!resp.generation) { + fprintf(stderr, "CREATE_QUEUE returned generation=0 (must be nonzero)\n"); + (void)destroy_queue(qp); + goto out; + } + + { + /* + * The kernel publishes the LIVE state into the consumer line + * of the doorbell. Verify the GPU side would read what the + * contract promises: packed (index=0, generation=X) tails + * and qp_state=LIVE. Without a fence the test buffer is + * already coherent on the CPU side after the ioctl returns. + */ + const struct usb4_rdma_dv_doorbell *db = doorbell; + + printf("doorbell_after_create sq_head=0x%08x cq_tail=0x%08x qp_state=%u generation=%u\n", + db->consumer.sq_head, db->consumer.cq_tail, + db->consumer.qp_state, db->consumer.generation); + + if (db->consumer.generation != resp.generation || + db->consumer.qp_state != USB4_RDMA_DV_QP_LIVE || + usb4_rdma_dv_tail_generation(db->consumer.sq_head) != resp.generation || + usb4_rdma_dv_tail_generation(db->consumer.cq_tail) != resp.generation || + usb4_rdma_dv_tail_index(db->consumer.sq_head) != 0 || + usb4_rdma_dv_tail_index(db->consumer.cq_tail) != 0) { + fprintf(stderr, + "doorbell consumer line did not match LIVE state\n"); + (void)destroy_queue(qp); + goto out; + } + } + + err = run_kick_smoke(qp, sq_buf, cq_buf, doorbell, resp.generation); + if (err) { + (void)destroy_queue(qp); + goto out; + } + + /* Second CREATE_QUEUE must be rejected with -EBUSY while the + * first is still attached. */ + err = create_queue(qp, &req, &resp); + if (err != EBUSY) { + fprintf(stderr, + "second CREATE_QUEUE should have failed with EBUSY, got %s (%d)\n", + err ? strerror(err) : "success", err); + (void)destroy_queue(qp); + goto out; + } + printf("create_queue_again rejected ebusy=ok\n"); + + err = destroy_queue(qp); + if (err) { + fprintf(stderr, "DESTROY_QUEUE failed: %s (%d)\n", + strerror(err), err); + goto out; + } + { + const struct usb4_rdma_dv_doorbell *db = doorbell; + + printf("doorbell_after_destroy qp_state=%u generation=%u\n", + db->consumer.qp_state, db->consumer.generation); + if (db->consumer.qp_state != USB4_RDMA_DV_QP_DEAD || + db->consumer.generation == resp.generation || + !db->consumer.generation) { + fprintf(stderr, + "doorbell consumer line did not reach DEAD with bumped generation\n"); + goto out; + } + } + printf("destroy_queue ok\n"); + + /* DESTROY_QUEUE on an already-destroyed queue must return ENOENT. */ + err = destroy_queue(qp); + if (err != ENOENT) { + fprintf(stderr, + "second DESTROY_QUEUE should have failed with ENOENT, got %s (%d)\n", + err ? strerror(err) : "success", err); + goto out; + } + printf("destroy_queue_again rejected enoent=ok\n"); + + ret = 0; +out: + if (qp && ibv_destroy_qp(qp)) + fprintf(stderr, "ibv_destroy_qp: %s\n", strerror(errno)); + if (cq && ibv_destroy_cq(cq)) + fprintf(stderr, "ibv_destroy_cq: %s\n", strerror(errno)); + if (doorbell_mr && ibv_dereg_mr(doorbell_mr)) + fprintf(stderr, "ibv_dereg_mr(doorbell): %s\n", strerror(errno)); + if (cq_mr && ibv_dereg_mr(cq_mr)) + fprintf(stderr, "ibv_dereg_mr(CQ): %s\n", strerror(errno)); + if (sq_mr && ibv_dereg_mr(sq_mr)) + fprintf(stderr, "ibv_dereg_mr(SQ): %s\n", strerror(errno)); + if (pd && ibv_dealloc_pd(pd)) + fprintf(stderr, "ibv_dealloc_pd: %s\n", strerror(errno)); + free(doorbell); + free(cq_buf); + free(sq_buf); + return ret; +} + +static void usage(const char *argv0) +{ + fprintf(stderr, + "usage: %s [-d ] [-l] [-q]\n" + "\n" + "Probe USB4 RDMA Direct Verbs QUERY_CAPS on a usb4_rdma* device.\n" + "\n" + "Options:\n" + " -d Match device by name (default: first usb4_rdma* / usb4_apple*)\n" + " -l Print doorbell/WQE layout only; do not open a device\n" + " -q After QUERY_CAPS, exercise CREATE_QUEUE/DESTROY_QUEUE on a fresh QP\n" + " -h Print this help\n", + argv0); +} + +int main(int argc, char **argv) +{ + struct usb4_rdma_dv_query_caps_resp resp; + struct ibv_device **list; + struct ibv_device *dev; + struct ibv_context *ctx; + const char *wanted = NULL; + bool layout_only = false; + bool queue_test = false; + int num_devices = 0; + int opt; + int err; + + while ((opt = getopt(argc, argv, "d:lqh")) != -1) { + switch (opt) { + case 'd': + wanted = optarg; + break; + case 'l': + layout_only = true; + break; + case 'q': + queue_test = true; + break; + case 'h': + usage(argv[0]); + return 0; + default: + usage(argv[0]); + return 2; + } + } + + if (layout_only) { + print_doorbell_layout(); + return 0; + } + + list = ibv_get_device_list(&num_devices); + if (!list) { + fprintf(stderr, "ibv_get_device_list failed: %s\n", + strerror(errno)); + return 1; + } + + dev = find_device(list, num_devices, wanted); + if (!dev) { + fprintf(stderr, + "no matching usb4_rdma/usb4_apple device found%s%s\n", + wanted ? " for name=" : "", wanted ? wanted : ""); + ibv_free_device_list(list); + return 1; + } + + ctx = ibv_open_device(dev); + if (!ctx) { + fprintf(stderr, "ibv_open_device(%s) failed: %s\n", + ibv_get_device_name(dev), strerror(errno)); + ibv_free_device_list(list); + return 1; + } + + err = query_caps(ctx, &resp); + if (err) { + fprintf(stderr, "QUERY_CAPS ioctl failed on %s: %s\n", + ibv_get_device_name(dev), strerror(err)); + ibv_close_device(ctx); + ibv_free_device_list(list); + return 1; + } + + printf("device=%s\n", ibv_get_device_name(dev)); + printf("abi_version=%" PRIu32 "\n", resp.abi_version); + print_caps_bitmap(resp.caps); + printf("max_sq_entries=%" PRIu32 " default_sq_entries=%" PRIu32 "\n", + resp.max_sq_entries, resp.default_sq_entries); + printf("max_cq_entries=%" PRIu32 " default_cq_entries=%" PRIu32 "\n", + resp.max_cq_entries, resp.default_cq_entries); + printf("wqe_size=%" PRIu32 " cqe_size=%" PRIu32 "\n", + resp.wqe_size, resp.cqe_size); + printf("doorbell_record_size=%" PRIu32 " doorbell_page_size=%" PRIu32 "\n", + resp.doorbell_record_size, resp.doorbell_page_size); + printf("tail_index_bits=%" PRIu32 " tail_generation_bits=%" PRIu32 "\n", + resp.tail_index_bits, resp.tail_generation_bits); + print_doorbell_layout(); + + if (queue_test) { + int rc = run_queue_test(ctx, &resp); + + ibv_close_device(ctx); + ibv_free_device_list(list); + return rc; + } + + ibv_close_device(ctx); + ibv_free_device_list(list); + return 0; +} diff --git a/userspace/bench/uc_oneway.c b/userspace/bench/uc_oneway.c index 60e2d89..25291db 100644 --- a/userspace/bench/uc_oneway.c +++ b/userspace/bench/uc_oneway.c @@ -1298,3 +1298,4 @@ int main(int argc, char **argv) ibv_close_device(ctx); return ret; } + diff --git a/userspace/usb4_rdma/usb4_rdma_dv.h b/userspace/usb4_rdma/usb4_rdma_dv.h new file mode 100644 index 0000000..38315b6 --- /dev/null +++ b/userspace/usb4_rdma/usb4_rdma_dv.h @@ -0,0 +1,402 @@ +/* SPDX-License-Identifier: GPL-2.0 OR BSD-3-Clause */ +#ifndef USB4_RDMA_DV_H +#define USB4_RDMA_DV_H + +#include + +/* + * USB4 RDMA Direct Verbs / GDA ABI, v2. + * + * ABI version policy: + * - abi_version 0 is invalid. + * - v1 is the first experimental queue ABI. + * - v2 adds RDMA_READ and software-backed 64-bit atomics for rocSHMEM + * signal/counter operations. Atomics are executed by the target kernel + * against the target MR under provider-side serialization; USB4 does not + * expose hardware RNIC atomics. + * - v2 also defines USB4_RDMA_DV_WQE_F_LOCAL_LOOPBACK for provider-known + * self-target atomics. When set on an atomic WQE, the kernel must validate + * that the destination QP and target MR are local and then execute the atomic + * locally instead of emitting a native-data packet. This is required for + * rocSHMEM self-polling barriers because the selected RoCE GID may be an + * IPv4/IPv6 netdev GID rather than the driver's synthetic link-local GID. + * - A method input carrying abi_version must exactly match a version returned + * by QUERY_CAPS unless that method explicitly documents compatible older + * input. + * - New struct fields may be appended behind existing reserved fields only if + * zero/default semantics keep old userspace working. Any reinterpretation of + * existing fields requires an ABI version bump. + * + * This is not a hardware RNIC ABI. The GPU produces software-RNIC WQEs into + * host-visible memory. The kernel owns the Thunderbolt/NHI rings, consumes + * these WQEs, performs the existing native SEND/WRITE data path, and produces + * CQEs back into host-visible memory. + * + * Queue-memory requirements: + * - SQ, CQ, and doorbell pages must be host-visible and coherent for CPU/GPU + * system-scope atomics. On ROCm the validated queue-memory allocators are + * hipMallocManaged, hipHostMalloc(..., hipHostMallocMapped), + * hipHostMallocCoherent, hipHostMallocUncached, and HSA CPU-location + * fine-grained or extended-scope fine-grained pools with CPU/GPU agent access + * granted. hipHostMallocNonCoherent is not valid queue memory; it can + * silently corrupt producer/consumer handshakes. HSA GPU-location pools are + * not valid queue memory for v1 unless they can also be registered through + * the RDMA MR path. + * - v1 queue memory is allocated by userspace and pinned by the kernel during + * CREATE_QUEUE. The kernel does not allocate or mmap queue pages in v1. + * - The 32-bit queue head/tail words must live in a dedicated doorbell page, + * separate from WQE/CQE arrays. The kernel poller watches SQ tail words, not + * descriptor cachelines. v1 uses one 128-byte doorbell record per QP inside a + * dedicated 4 KiB page-sized mapping. + * - The doorbell record is split into two 64-byte cachelines. The producer line + * contains GPU-written/kernel-read fields. The consumer line contains + * kernel-written/GPU-read fields. Do not move fields across those lines + * without an ABI bump. + * + * Producer rules for GPU -> kernel SQ: + * store WQE generation matching the current QP generation + * store WQE fields + * system-scope release fence/atomic + * system-scope atomic store of producer.sq_tail with the same generation + * + * Consumer rules for kernel reading SQ: + * smp_load_acquire(producer.sq_tail) + * reject if the packed generation or producer.generation is stale + * read WQE fields + * reject any WQE whose generation is stale + * smp_store_release(consumer.sq_head) after consuming WQEs + * + * Producer rules for kernel -> GPU CQ: + * store CQE fields + * smp_store_release(consumer.cq_tail) + * + * Consumer rules for GPU reading CQ: + * system-scope acquire load of consumer.cq_tail + * read CQE fields + * system-scope release store of producer.cq_head after consuming CQEs + * + * Generation protocol: + * - CREATE_QUEUE returns generation X and initializes consumer.generation to X. + * - Before producing WQEs, userspace must mirror X into producer.generation. + * - All SQ/CQ head and tail words use usb4_rdma_dv_tail_pack(). The index part + * is a monotonically increasing producer/consumer counter modulo 2^24; the + * ring slot is index % queue_depth. The generation part is X. + * - The kernel rejects producer.sq_tail, producer.cq_head, or any WQE whose + * embedded generation differs from X. + * - DESTROY_QUEUE bumps consumer.generation before teardown completes, so stale + * producer writes from an old QP lifetime are detectable by later drain/poll + * attempts. + * - DESTROY_QUEUE is a teardown operation, not a drain operation. v1 discards + * in-flight or not-yet-completed DV WQEs without producing FLUSH_ERR CQEs. + * Consumers that require completions must stop producing, wait for SQ head + * and CQ tail to catch up, and only then destroy the queue. + * + * v1 does not define a GPU-produced receive queue. DV owns only the + * GPU-produced send queue and DV completion queue surface. The receive queue + * stays kernel-owned, and the peer may post receive WQEs through standard + * verbs while the QP has an active DV queue. SEND/SEND_IMM and + * RDMA_WRITE_WITH_IMM consume peer receive WQEs and complete on the peer's + * normal verbs receive CQ. RDMA_WRITE, RDMA_READ, and atomics do not consume + * a peer receive WQE. + * + * CQE ordering and errors: + * - v2 produces CQEs in SQ WQE order. Internal native SEND/WRITE/READ/ATOMIC + * completions may arrive out of order, but the kernel buffers them until + * prior WQEs have completed or been skipped as unsignaled successes. + * - Transient transport admission pressure is not a WQE error. If the native + * path cannot currently reserve local TX queue/ring resources, the kernel + * leaves SQ head on the blocked WQE, produces no CQE, and retries on a later + * poll scan. Later WQEs on the same QP are not consumed until that WQE is + * admitted, preserving per-QP SEND/WRITE ordering. Other QPs may continue to + * drain in round-robin order. + * - USB4_RDMA_DV_WQE_F_FENCE is a software admission fence. A fenced WQE is + * not admitted to the native transport until all earlier DV WQEs on the same + * QP have completed and the ordered DV completion stream has advanced past + * them. This provides an explicit software ordering primitive for future + * compound protocols; it is not intended to make a high-rate producer spin on + * sender-side admission retries. + * - A WQE generation mismatch is observable: the kernel writes a CQE with + * USB4_RDMA_DV_CQE_STALE_GEN, advances SQ head past that WQE, and leaves the + * QP live. + * - CQ overflow is fatal for the DV queue. The kernel moves the QP to ERR, + * writes consumer.qp_state = USB4_RDMA_DV_QP_ERR if the doorbell page is + * still mapped, and stops posting DV CQEs. If there is no CQ slot available, + * no sentinel CQE is guaranteed. + */ + +#define USB4_RDMA_DV_ABI_VERSION 2 + +#define USB4_RDMA_DV_MIN_QUEUE_ENTRIES 2 +#define USB4_RDMA_DV_MAX_SQ_ENTRIES 1024 +#define USB4_RDMA_DV_MAX_CQ_ENTRIES 4096 +#define USB4_RDMA_DV_DEFAULT_SQ_ENTRIES 256 +#define USB4_RDMA_DV_DEFAULT_CQ_ENTRIES 512 +#define USB4_RDMA_DV_WQE_SIZE 64 /* bytes */ +#define USB4_RDMA_DV_CQE_SIZE 64 /* bytes */ +#define USB4_RDMA_DV_DOORBELL_LINE_SIZE 64 /* bytes */ +#define USB4_RDMA_DV_DOORBELL_RECORD_SIZE \ + (2 * USB4_RDMA_DV_DOORBELL_LINE_SIZE) /* bytes */ +#define USB4_RDMA_DV_DOORBELL_PAGE_SIZE 4096 /* bytes */ + +/* + * Private uverbs ids. These use the RDMA core driver namespace bit + * (UVERBS_ID_DRIVER_NS / UVERBS_API_NS_FLAG == 1 << 12), but keep the value + * local so this header can be included by small standalone probes without + * pulling in rdma-core's private ioctl helper headers. + */ +#define USB4_RDMA_DV_DRIVER_NS (1u << 12) + +enum usb4_rdma_dv_objects { + USB4_RDMA_DV_OBJECT_DEVICE = USB4_RDMA_DV_DRIVER_NS, +}; + +enum usb4_rdma_dv_methods { + USB4_RDMA_DV_METHOD_QUERY_CAPS = USB4_RDMA_DV_DRIVER_NS, + USB4_RDMA_DV_METHOD_CREATE_QUEUE, + USB4_RDMA_DV_METHOD_DESTROY_QUEUE, + USB4_RDMA_DV_METHOD_KICK, +}; + +enum usb4_rdma_dv_query_caps_attrs { + USB4_RDMA_DV_ATTR_QUERY_CAPS_RESP = USB4_RDMA_DV_DRIVER_NS, +}; + +enum usb4_rdma_dv_create_queue_attrs { + USB4_RDMA_DV_ATTR_CREATE_QUEUE_QP = USB4_RDMA_DV_DRIVER_NS, + USB4_RDMA_DV_ATTR_CREATE_QUEUE_REQ, + USB4_RDMA_DV_ATTR_CREATE_QUEUE_RESP, +}; + +enum usb4_rdma_dv_destroy_queue_attrs { + USB4_RDMA_DV_ATTR_DESTROY_QUEUE_QP = USB4_RDMA_DV_DRIVER_NS, +}; + +enum usb4_rdma_dv_kick_attrs { + USB4_RDMA_DV_ATTR_KICK_QP = USB4_RDMA_DV_DRIVER_NS, + USB4_RDMA_DV_ATTR_KICK_REQ, +}; + +#define USB4_RDMA_DV_TAIL_INDEX_BITS 24 +#define USB4_RDMA_DV_TAIL_INDEX_MASK ((1u << USB4_RDMA_DV_TAIL_INDEX_BITS) - 1) +#define USB4_RDMA_DV_TAIL_GEN_SHIFT USB4_RDMA_DV_TAIL_INDEX_BITS +#define USB4_RDMA_DV_TAIL_GENERATION_BITS \ + (32 - USB4_RDMA_DV_TAIL_INDEX_BITS) + +static inline __u32 usb4_rdma_dv_tail_pack(__u32 index, __u8 generation) +{ + return (index & USB4_RDMA_DV_TAIL_INDEX_MASK) | + ((__u32)generation << USB4_RDMA_DV_TAIL_GEN_SHIFT); +} + +static inline __u32 usb4_rdma_dv_tail_index(__u32 packed) +{ + return packed & USB4_RDMA_DV_TAIL_INDEX_MASK; +} + +static inline __u8 usb4_rdma_dv_tail_generation(__u32 packed) +{ + return packed >> USB4_RDMA_DV_TAIL_GEN_SHIFT; +} + +enum usb4_rdma_dv_caps { + USB4_RDMA_DV_CAP_SEND = 1u << 0, + USB4_RDMA_DV_CAP_SEND_IMM = 1u << 1, + USB4_RDMA_DV_CAP_WRITE = 1u << 2, + USB4_RDMA_DV_CAP_WRITE_IMM = 1u << 3, + USB4_RDMA_DV_CAP_FENCE = 1u << 4, + USB4_RDMA_DV_CAP_READ = 1u << 5, + USB4_RDMA_DV_CAP_ATOMIC_FETCH_ADD = 1u << 6, + USB4_RDMA_DV_CAP_ATOMIC_SWAP = 1u << 7, + USB4_RDMA_DV_CAP_ATOMIC_CMP_SWAP = 1u << 8, +}; + +#define USB4_RDMA_DV_CAPS_V2 \ + (USB4_RDMA_DV_CAP_SEND | \ + USB4_RDMA_DV_CAP_SEND_IMM | \ + USB4_RDMA_DV_CAP_WRITE | \ + USB4_RDMA_DV_CAP_WRITE_IMM | \ + USB4_RDMA_DV_CAP_FENCE | \ + USB4_RDMA_DV_CAP_READ | \ + USB4_RDMA_DV_CAP_ATOMIC_FETCH_ADD | \ + USB4_RDMA_DV_CAP_ATOMIC_SWAP | \ + USB4_RDMA_DV_CAP_ATOMIC_CMP_SWAP) + +#define USB4_RDMA_DV_CAPS_V1 \ + (USB4_RDMA_DV_CAP_SEND | \ + USB4_RDMA_DV_CAP_SEND_IMM | \ + USB4_RDMA_DV_CAP_WRITE | \ + USB4_RDMA_DV_CAP_WRITE_IMM | \ + USB4_RDMA_DV_CAP_FENCE) + +enum usb4_rdma_dv_wqe_opcode { + USB4_RDMA_DV_WQE_NOP = 0, + USB4_RDMA_DV_WQE_SEND = 1, + USB4_RDMA_DV_WQE_SEND_IMM = 2, + USB4_RDMA_DV_WQE_RDMA_WRITE = 3, + USB4_RDMA_DV_WQE_RDMA_WRITE_IMM = 4, + USB4_RDMA_DV_WQE_RDMA_READ = 5, + USB4_RDMA_DV_WQE_ATOMIC_FETCH_ADD = 6, + USB4_RDMA_DV_WQE_ATOMIC_SWAP = 7, + USB4_RDMA_DV_WQE_ATOMIC_CMP_SWAP = 8, +}; + +enum usb4_rdma_dv_wqe_flags { + USB4_RDMA_DV_WQE_F_SIGNALED = 1u << 0, + USB4_RDMA_DV_WQE_F_SOLICITED = 1u << 1, + USB4_RDMA_DV_WQE_F_FENCE = 1u << 2, + USB4_RDMA_DV_WQE_F_LOCAL_LOOPBACK = 1u << 3, +}; + +enum usb4_rdma_dv_cqe_status { + USB4_RDMA_DV_CQE_SUCCESS = 0, + USB4_RDMA_DV_CQE_WR_FLUSH_ERR = 1, + USB4_RDMA_DV_CQE_LOCAL_LEN_ERR = 2, + USB4_RDMA_DV_CQE_LOCAL_PROT_ERR = 3, + USB4_RDMA_DV_CQE_REMOTE_ACCESS_ERR = 4, + USB4_RDMA_DV_CQE_RETRY_EXC_ERR = 5, + USB4_RDMA_DV_CQE_STALE_GEN = 6, + USB4_RDMA_DV_CQE_GENERAL_ERR = 255, +}; + +enum usb4_rdma_dv_qp_state { + USB4_RDMA_DV_QP_LIVE = 0, + USB4_RDMA_DV_QP_ERR = 1, + USB4_RDMA_DV_QP_DEAD = 2, +}; + +/* + * v2 WQE format: + * - exactly one SGE + * - no inline data + * - SEND, SEND_WITH_IMM, RDMA_WRITE, RDMA_WRITE_WITH_IMM, RDMA_READ + * - 64-bit ATOMIC_FETCH_ADD, ATOMIC_SWAP, and ATOMIC_CMP_SWAP + * + * For atomics, length must be 8 and local_addr/lkey identify an 8-byte local + * result slot. FETCH_ADD and SWAP use reserved1[0] as the add/swap value. + * CMP_SWAP uses reserved1[0] as the swap value and reserved1[1] as the compare + * value. The returned value is the original remote value. Users that do not + * need the returned value must still provide a valid 8-byte local result slot; + * providers may reuse a scratch MR for that case. + */ +struct usb4_rdma_dv_wqe { + __u16 opcode; + __u16 flags; + __u32 length; + __u64 wr_id; + __u64 local_addr; + __u32 lkey; + __u32 rkey; + __u64 remote_addr; + __u32 imm_data; + __u32 generation; + __u64 reserved1[2]; +} __attribute__((packed, aligned(64))); + +struct usb4_rdma_dv_cqe { + __u64 wr_id; + __u32 status; + __u32 opcode; + __u32 byte_len; + __u32 imm_data; + __u32 vendor_err; + __u32 qp_state; + __u64 reserved[4]; +} __attribute__((packed, aligned(64))); + +/* + * Doorbell page layout for one queue pair. Each 32-bit word is independently + * polled with acquire semantics by its consumer. A stale GPU write from a torn + * down QP is detected by the explicit generation fields and by the generation + * bits in packed queue indexes. + */ +struct usb4_rdma_dv_doorbell_producer_line { + __u32 sq_tail; + __u32 cq_head; + __u32 generation; + __u32 reserved0; + __u64 reserved1[6]; +} __attribute__((packed, aligned(64))); + +struct usb4_rdma_dv_doorbell_consumer_line { + __u32 sq_head; + __u32 cq_tail; + __u32 qp_state; + __u32 generation; + __u64 reserved1[6]; +} __attribute__((packed, aligned(64))); + +struct usb4_rdma_dv_doorbell { + struct usb4_rdma_dv_doorbell_producer_line producer; + struct usb4_rdma_dv_doorbell_consumer_line consumer; +} __attribute__((packed, aligned(64))); + +/* + * Capability fields are ceilings and recommendations, not fixed allocations. + * Userspace requests actual depths in CREATE_QUEUE. The kernel rejects depths + * below USB4_RDMA_DV_MIN_QUEUE_ENTRIES or above max_*_entries. + * + * On the current two-rail USB4 software-RNIC path, 256 SQ entries and 512 CQ + * entries are the recommended starting point. Deeper queues are exposed as a + * cap for consumers that deliberately batch or multiplex more work, not + * because the Thunderbolt wire needs 1024 WQEs in flight to fill its BDP. + */ +struct usb4_rdma_dv_query_caps_resp { + __u32 abi_version; + __u32 caps; + __u32 max_sq_entries; + __u32 max_cq_entries; + __u32 default_sq_entries; + __u32 default_cq_entries; + __u32 wqe_size; + __u32 cqe_size; + __u32 doorbell_record_size; + __u32 doorbell_page_size; + __u32 tail_index_bits; + __u32 tail_generation_bits; + __u32 reserved[4]; +}; + +/* + * Queue addresses are userspace virtual addresses of host-visible coherent + * memory. The kernel pins these ranges during CREATE_QUEUE and releases them + * during DESTROY_QUEUE/QP destroy; v1 does not require separate MR keys for + * SQ/CQ/doorbell queue memory. + * sq_stride/cq_stride are in bytes and must equal the caps-reported WQE/CQE + * sizes for v1. doorbell_addr is page-aligned; the first record in that page + * is a struct usb4_rdma_dv_doorbell. + */ +struct usb4_rdma_dv_queue_create { + __u32 abi_version; + __u32 flags; + __u64 sq_addr; + __u64 cq_addr; + __u64 doorbell_addr; + __u32 sq_entries; + __u32 cq_entries; + __u32 sq_stride; + __u32 cq_stride; + __u32 reserved[4]; +}; + +struct usb4_rdma_dv_queue_resp { + __u32 qp_num; + __u32 generation; + __u32 flags; + __u32 reserved0; +}; + +/* + * Optional explicit-kick v1 path. KICK is a drain trigger: the handler + * consumes all WQEs between the current SQ head and sq_tail, then returns. If + * SQ head already equals sq_tail, KICK is a cheap no-op. Kernel-poll mode uses + * the same drain operation after observing producer.sq_tail directly; KICK is + * not required when polling is enabled. sq_tail is the packed tail value + * userspace already published to producer.sq_tail. + */ +struct usb4_rdma_dv_kick { + __u32 sq_tail; + __u32 flags; + __u32 reserved[2]; +}; + +#endif /* USB4_RDMA_DV_H */