Architecture

Architecture#

This page describes how the hipObject library does its work: the code in src/, the helpers in shared/, and the public header include/hipobj.h. It doesn’t cover the build system, the tests, the examples, or the MinIO bridge in integrations/. The wire formats are in hipObject cuObject Interoperability, and the rules for changing the code are in AGENTS.md in the repository root.

What the Library Does#

hipObject moves the data of an S3 object between an S3 server and a buffer in GPU or host memory by RDMA, so that the payload doesn’t pass through the client’s CPU. It does not speak HTTP or S3. The work is split three ways:

  • The application owns the S3 conversation. It sends the S3 request, with an RDMA token from hipObject in the x-amz-rdma-token header, and hands the x-amz-rdma-reply header from the response back to hipObject. It does this either through the two hipObjOps_t callbacks that hipObjGet() and hipObjPut() call, or by driving the transfer itself with hipObjGetRdmaToken(), hipObjBufSync(), and hipObjParseRdmaReply().

  • The S3 server moves the payload with one-sided RDMA: an RDMA WRITE into the client’s buffer for a GET, an RDMA READ from it for a PUT. The client posts no work requests of its own.

  • hipObject prepares the client side: it chooses the NIC closest to the GPU, opens it, registers buffers with it (as dma-bufs for GPU memory), mints the tokens that describe a buffer and a queue pair (QP), parses the server’s reply, connects the QP to the peer the reply names, waits for the transfer’s completion, and copies through host memory when the NIC can’t reach the GPU directly.

The headers are wire-compatible with NVIDIA cuObject, but hipObject uses the Reliable Connection (RC) transport, not cuObject’s Dynamic Connection (DC), on Broadcom (bnxt) and AMD Pensando (ionic) NICs.

Source Map#

File

Role

include/hipobj.h

The public C API, including the thread-safety contract (the threads group)

src/hipobj.cpp

The C API’s implementation: argument checks, the transfer sequence, staging copies

src/state.h, state.cpp

Platform, Options, Instance, the current instance, init and shutdown

src/buffer.h, buffer.cpp

BufferMap: registered buffers, dma-buf and staging-buffer registration

src/abandoned-copies.h, .cpp

AbandonedCopies: staging copies that timed out but may still run

src/rdma/transport.h, .cpp

RdmaDevice, RcConnection, QP state transitions, completion polling

src/rdma/connection-pool.h, .cpp

ConnectionPool and Lease: one connection per transfer

src/rdma/token.h, .cpp

Token encoding and decoding, x-amz-rdma-reply parsing

src/rdma/vendor-ops.h, bnxt/, ionic/

Vendor QP attributes, compiled in by HIPOBJ_BNXT and HIPOBJ_IONIC

src/s3/control.h, .cpp

Calling the hipObjOps_t callbacks

src/common/hip-seam.*

HipRuntime interface and SystemHip (HIP, plus HSA’s dma-buf export)

src/common/ibv-wrapper.*

Verbs interface and LoadedVerbs (libibverbs through dlopen())

src/common/ibv-core.h

hipObject’s own copy of the libibverbs ABI types

src/common/ibv-ptr.h

std::unique_ptr types for verbs objects, and own()

src/common/nic-seam.h, rdma-topology.*

NicTopology interface, SysfsTopology, NIC and GID selection

src/common/thread-safety.h

Mutexes and locks annotated for Clang’s thread safety analysis

src/common/no-destructor.h, dl-handle.h, gid.h, ibv-gid.h

Small helpers

shared/

Helpers shared with tools and tests: span, address, parse, sockaddr, warnings

Layers#

              application, with its own S3 client
                  |                         ^
      hipobj.h C API calls          hipObjOps_t callbacks
                  v                         |
+-----------------------------------------------------------------+
| hipobj.cpp: argument checks, transfer sequence, staging copies  |
+-----------------------------------------------------------------+
| state: the current Instance                                     |
|   options | device | tokenConn | pool | buffers | abandoned     |
+----------------+-------------------+--------------+-------------+
| buffer,        | rdma/transport,   | rdma/token   | s3/control  |
| abandoned-     | connection-pool,  |              |             |
| copies         | vendor attributes |              |             |
+----------------+-------------------+--------------+-------------+
| Platform:   HipRuntime    |    Verbs      |    NicTopology      |
+---------------------------+---------------+---------------------+
             |                      |                  |
   HIP runtime, and HSA      libibverbs.so.1      /sys/class/
   for dma-buf export        (dlopen)             infiniband

Calls go down the diagram, or across a row, never up. Nothing below hipobj.cpp calls a public function, and only the three Platform implementations call the HIP runtime, libibverbs, or sysfs.

The Instance#

All of the library’s state lives in one Instance (src/state.h), built by hipObjInit() and released after hipObjShutdown(). It holds:

  • platform: the outside world (see The Platform).

  • options: settings read from the environment once, at init (see Configuration).

  • config: the GPU and NIC that init chose. Nothing in the library reads it after init; the tests do.

  • device: the open RdmaDevice: its verbs context, protection domain (PD), port, GID, and vendor ID.

  • tokenConn: the one connection whose QP number hipObjGetRdmaToken() puts in its tokens.

  • pool: the ConnectionPool that hipObjGet() and hipObjPut() lease connections from.

  • buffers: the BufferMap of registered buffers.

  • abandoned: the AbandonedCopies list.

Everything but the last three is fixed once the instance is built, and those three each have a lock of their own, so the instance needs no lock.

Finding the Instance#

The process-wide state is a Current (src/state.cpp): the instance in a std::shared_ptr, a shared mutex that guards it, and an initMutex that init and shutdown hold. It’s held in a NoDestructor, so it’s never destroyed: destroying an instance releases verbs and HIP resources, which isn’t safe during static destruction.

Every public function that needs the library starts with currentInstance(), which copies the shared_ptr under a shared lock and releases the lock. The copy keeps the instance alive for the rest of the call, even if another thread shuts the library down meanwhile, and the instance found is passed down to everything the call does. This is why nothing inside the library calls a public function: it could find a different instance.

Init#

hipObjInit() checks its hipObjConfig_t, returns hipObjAlreadyInitialized if there’s an instance, gets the system platform (see The Platform), reads the options (see Configuration), and calls initWith(), which holds initMutex throughout:

  1. Resolves the GPU. A gpuDevice of -1 means the calling thread’s current HIP device. Having no GPU at all is only acceptable with a nicHint.

  2. Chooses the NIC (see NIC and GID Selection).

  3. Opens the device, which also chooses the port and GID.

  4. Creates the token connection.

  5. Builds the Instance, whose ConnectionPool creates HIPOBJ_MAX_TRANSFERS connections. Creating fewer is a failure.

  6. Publishes the instance under the exclusive lock.

Everything is built in locals first, so a failure at any step releases what was built, and nothing is published.

initWith() takes the Platform and Options as parameters; the tests call it directly with a fake platform.

Shutdown#

hipObjShutdown() (shutdownInstance()) moves the instance out of Current, so later calls find the library uninitialized, and closes its connection pool, which wakes any transfer waiting for a connection with hipObjNotInitialized. It doesn’t wait for calls in progress: the instance is destroyed when the last call holding it returns. A new hipObjInit() can run before then, and the old calls keep using the old instance.

The instance’s members are destroyed in reverse order, which gives the teardown order the verbs objects need:

  1. abandoned: copies that are over are released; the staging buffers of those still running are leaked on purpose.

  2. buffers: each MR is deregistered, then its staging buffer freed.

  3. pool, then tokenConn: each QP is destroyed, then its CQ.

  4. device: the PD is deallocated, then the context closed.

  5. platform outlives all of them.

The Platform#

Platform (src/state.h) bundles references to three interfaces, which are the library’s only route to the outside world:

Interface

Production class

Reaches

HipRuntime

SystemHip

HIP: devices, PCI bus ids, pinned host memory, async copies, events, pointer attributes; HSA’s hsa_amd_portable_export_dmabuf()

Verbs

LoadedVerbs

libibverbs: devices, PDs, MRs (plain and dma-buf), CQs, QPs, polling

NicTopology

SysfsTopology

/sys/class/infiniband: each NIC’s PCI bus id and NUMA node, and the RoCE version of each GID

Their methods are const, so one object serves every thread. Objects below the instance hold a reference to the narrowest interface they need, given to their constructors; there are no global accessors. The tests substitute a FakePlatform and run the rest of the library unchanged.

systemPlatform() builds the production platform on first use and never destroys it, so an application that never calls hipObjInit() never loads libibverbs, and once loaded, libibverbs stays loaded.

libibverbs#

The library doesn’t link to libibverbs or include its headers. loadLibibverbs() opens libibverbs.so.1 with dlopen(), by its soname, which the runtime package installs, and looks up each function with dlsym(). ibv-core.h declares the ABI’s structures itself (derived from rocSHMEM by way of rocm-xio), which ties hipObject to that soname’s ABI. Two parts are optional:

  • ibv_reg_dmabuf_mr() (rdma-core 34 and later). Without it, supportsDmabuf() is false, and GPU memory can only be registered with a plain ibv_reg_mr() (see Buffer Registration).

  • ibv_poll_cq() and ibv_post_send(), which are inline functions in the real header. Without the symbols, LoadedVerbs calls the provider’s entries in the context’s ops table, as the inline functions do.

Every verbs object is owned by a std::unique_ptr whose deleter holds the Verbs it came from (ibv-ptr.h), created with own<Ptr>(verbs, raw). HostBufPtr (pinned host memory) and HipEventPtr do the same with the HipRuntime. dropOwnership() gives up an object without releasing it, for the deliberate leaks described below.

NIC and GID Selection#

At init (src/common/rdma-topology.cpp):

NIC

GetClosestNicToGpu() asks HIP for the GPU’s PCI bus id and the topology for every NIC with an active port. nicHint here is a comma-separated allow-list of device names. The distance between two bus ids is the number of characters after their common prefix, ignoring the PCI function; the first of the closest NICs wins. It’s a heuristic over the bus-id strings, not a walk of the PCIe tree. The topology also reports each NIC’s NUMA node, but the selection doesn’t use it.

Fallback

If there’s no GPU, or no NIC with a known bus id, and nicHint is set, the hint is opened as a device name directly. This is how GPU-less hosts, such as the emulated-NIC CI lanes, initialize. Without a hint, init fails with hipObjNicNotFound.

Port

The device is always opened on port 1, although the topology records each NIC’s first active port.

GID

PickBestGid() reads the port’s GID table, skips unconfigured entries and those whose RoCE version is unknown, and prefers, from best to worst: RoCE v2 IPv4-mapped, RoCE v2 global, RoCE v2 link-local, RoCE v1 global, RoCE v1 link-local. Ties go to the lowest index; with no candidate, index 0 is used.

Vendor

The device’s PCI vendor ID is read once, so that connecting a QP can apply the bnxt or ionic attributes without querying it.

Connections#

RcConnection (src/rdma/transport.h) is an RC QP and the CQ it completes on, both owned, on an RdmaDevice that has to outlive it. createRcConnection() creates a 16-entry CQ and a QP with 8 send and 8 receive work requests of one SGE each, and moves the QP to INIT with remote read and write access.

QP state changes:

  • connectRcPeer() moves INIT to RTR, aimed at the peer token’s QP number and GID with a 4096-byte MTU and the device’s GID index, then to RTS (timeout 14, 7 retries, 7 RNR retries, PSNs 0). If the device is a Broadcom (0x14E4) or Pensando (0x1DD8) NIC and that backend was compiled in, its attributes are applied; today both set the same values the generic code does.

  • resetRcConnection() moves the QP to RESET, drains the CQ, and moves it back to INIT, ready for a different peer. An RC QP can’t reconnect from RTS any other way.

  • pollCompletion() polls the CQ, sleeping 1 ms between empty polls, until a completion arrives, a completion reports an error, or the timeout passes. An error completion’s status and vendor error are printed to stderr, since nothing else keeps them.

The Connection Pool#

Each hipObjGet() and hipObjPut() needs a QP of its own, since a connected RC QP is bound to one peer. ConnectionPool (src/rdma/connection-pool.h) creates a fixed set of connections at init and leases them out:

  • acquire() returns a Lease, an RAII handle to one connection. If none is free it waits on a condition variable, unless the calling thread already holds a lease (a thread-local count): then it returns Busy, since the connection it would wait for could be its own, as when a transfer’s callback starts another transfer.

  • When a lease ends, a connection that was connected is reset. If it can’t be reset it’s replaced with a new one, and if that fails too it’s dropped and the pool shrinks; an empty pool fails transfers with hipObjRdmaError.

  • close(), from shutdown, makes acquire() fail with Closed, including calls already waiting.

The connection slots are a vector whose size never changes, so a lease can use its slot without the pool’s lock; the lock only guards the free list and counters.

The Token Connection#

hipObjGetRdmaToken() serves transfers the application drives itself. hipObject can’t tell when such a transfer ends, so it can’t lease a connection for it. Instead, every such token names the instance’s single tokenConn, which is never leased, never connected, and stays in INIT. This is why the API documentation says that application-driven transfers mustn’t overlap each other.

Buffer Registration#

BufferMap (src/buffer.h) is a std::map from the address a buffer was registered at to a std::shared_ptr<const BufEntry>. Lookups are by that exact address: a transfer passes the registered address and selects part of the buffer with its offset. The map holds at most 256 buffers, each up to 4 GiB.

A buffer is registered in one of three ways:

Registered with

How

Address in tokens

Extra work per transfer

hipObjBufRegister(), directly

dma-buf MR, or a plain MR if that succeeds

The caller’s address

hipDeviceSynchronize() after the RDMA

hipObjBufRegister(), staged

Plain MR over a pinned host buffer of the same size

The staging buffer’s

A copy to the staging buffer before a PUT, or from it after a GET

hipObjBufRegisterHost()

Plain MR over the caller’s memory

The caller’s address

None

For memory that HIP reports as device memory, hipObjBufRegister() first checks that the range lies within one HIP allocation. registerDeviceMemory() then tries, in order:

  1. If dma-buf registration is enabled (HIPOBJ_DMABUF) and libibverbs supports it: export the range as a dma-buf through HSA, and register it with ibv_reg_dmabuf_mr(), using the caller’s address as the I/O virtual address. The dma-buf fd is closed as soon as the call returns, since the MR keeps its own reference. Whether this works is found out by trying, not predicted from kernel configuration.

  2. A plain ibv_reg_mr(), which works for host memory and generally fails for device memory.

  3. If both fail for device memory: allocate a pinned host staging buffer with hipHostMalloc() and register that, after a warning on stderr that the buffer is no longer GPU-direct. HIPOBJ_REQUIRE_GPU_DIRECT makes this a failure instead, so that CI can’t pass on the fallback without noticing.

Every MR gets local write, remote read, and remote write access, whatever the direction of the transfers that will use it.

Registration runs without the map’s lock: the map is checked first under a shared lock, so a duplicate or a full map doesn’t register for nothing, and checked again when the entry is inserted under the exclusive lock. If two threads register the same address, the loser’s registration is released after the lock.

A BufEntry declares the staging buffer before the MR, so the MR is deregistered before the memory it covers is freed.

deregisterBuffer() refuses, with hipObjBusy, to remove an entry that anything else holds a copy of: a call that found it with find(), or an abandoned copy. Since copies are only made under the map’s lock, the check can’t miss one. The MR and staging buffer are released after the lock.

A Transfer#

hipObjGet() and hipObjPut() (src/hipobj.cpp) share runRdmaTransfer(). With the instance found and the arguments checked:

  1. Look up the buffer with buffers.find(), check that [offset, offset + size) lies inside it without computing the sum, and refuse a staged buffer with hipObjBusy while an abandoned copy of it is still running (see Staging Copies).

  2. PUT, staged: copy the range from the device to the staging buffer.

  3. Lease a connection from the pool.

  4. Mint the token: the lease’s QP number, the device’s GID, the MR’s rkey, the buffer’s token address plus the offset, the size, RC, the port number, and a LID of 0. It’s encoded as 88 lowercase hex digits (44 little-endian bytes, laid out in hipObject cuObject Interoperability) into a NUL-terminated array on the stack.

  5. sendRequest(ctx, token, 88): the application embeds the token in its S3 request.

  6. recvReply(ctx, buffer, &length): the application returns the x-amz-rdma-reply value, into a 512-byte buffer on the stack. Anything but a success reply (ok, 200, 204, or 206), including 501, which means the server declined RDMA, fails the call with hipObjS3Error. Falling back to plain HTTP is the application’s job; it has the reply.

  7. Connect, if the reply has the form 200:<88 hex digits>, to the peer QP that token names. connectRcPeer() uses only the token’s transport (which must be RC), QP number, GID, and LID; the other fields are reserved in a reply (see hipObject cuObject Interoperability).

  8. Wait for a completion on the lease’s CQ, for up to 5 seconds.

  9. Synchronize the device, for a directly registered GPU buffer.

  10. End the lease, which resets the QP if it was connected and returns it to the pool.

  11. GET, staged: copy the range from the staging buffer to the device.

The hipObjHandle_t argument is unused, and neither the token nor the transfer records whether it’s a GET or a PUT: the server knows that from the S3 request.

No lock is held during the callbacks, so they can call any public function. They run on the calling thread.

caller          hipObject                 app's S3 client       S3 server
  |                 |                            |                   |
  |-- hipObjGet --->| find buffer, lease conn,   |                   |
  |                 | mint token                 |                   |
  |                 |-- sendRequest(token) ----->|-- GET + token --->|
  |                 |                            |   RDMA WRITE into |
  |                 |                            |   the buffer      |
  |                 |<-- recvReply(reply) -------|<-- reply ---------|
  |                 | parse reply, connect QP,   |                   |
  |                 | poll CQ, sync device,      |                   |
  |                 | return lease               |                   |
  |                 | (staged: staging -> GPU)   |                   |
  |<-- result ------|                            |                   |

The diagram shows the server doing its RDMA before it replies, which is what the in-repo test server (hipobj-rdma-test-server) does; see Open Questions about that ordering.

Transfers the Application Drives#

The MinIO bridge, and any application that doesn’t fit the callback model, uses the lower-level calls instead:

  • hipObjGetRdmaToken() mints a token for the start of a registered buffer, on the token connection (see The Token Connection), and returns a malloc()ed copy that hipObjPutRdmaToken() frees.

  • hipObjBufSync() performs the staging copy that hipObjPut() and hipObjGet() would (HIPOBJ_SYNC_TO_HOST before a PUT, HIPOBJ_SYNC_TO_DEVICE after a GET), and does nothing for a buffer that isn’t staged.

  • hipObjParseRdmaReply() turns an x-amz-rdma-reply value into its HTTP status.

  • hipObjTokenClientNic() extracts the IPv4 address from a token’s GID, when the GID is IPv4-mapped.

These calls don’t connect a QP or wait for a completion. The last two, like hipObjGetErrorString(), hipObjGetVersionString(), and hipObjGetVersion(), don’t need an initialized library.

Tokens and Replies#

src/rdma/token.cpp serializes a token field by field, with explicit little-endian byte order, never by copying a struct, and a static_assert fails the build if a field’s size changes the wire format. Decoding accepts exactly 88 hex digits in either case and a transport byte of DC or RC, and nothing else. A reply is ok, err, a three-digit status from 100 to 599, or 200: followed by a token; trailing NUL, CR, and LF characters are ignored, and anything else is rejected rather than guessed at.

Staging Copies#

A staging copy (stageCopyWithDeadline() in src/hipobj.cpp) is a hipMemcpyAsync() on the calling thread’s per-thread default stream (so copies from different threads don’t serialize on the null stream), with an event recorded after it. The event is polled every millisecond until the copy finishes or HIPOBJ_STAGE_TIMEOUT_MS (default 30 s) passes.

The timeout exists because ROCm waits on a copy by spinning in user space: a DMA that never completes would otherwise hang the caller forever. But a copy that timed out can’t be stopped, and can still read or write the staging buffer whenever the GPU gets to it. So the copy’s event and the buffer’s BufEntry are handed to AbandonedCopies (src/abandoned-copies.h):

  • Holding the entry keeps the staging buffer allocated, and makes deregistration report hipObjBusy.

  • Transfers, syncs, and token requests on that buffer return hipObjBusy while the copy is running, since they’d mix their data with its.

  • Copies that have finished are released when they’re next checked: by a call on the same buffer, or by any hipObjBufDeregister().

  • If the instance is destroyed while a copy is still running, its event and staging buffer are leaked rather than freed under it.

Handing a copy over must not fail, or a running copy would be left untracked. So each copy reserves room in the list (reserve(), which returns a Slot) before it starts, and add() uses that room without allocating.

Concurrency#

The caller-facing contract is in the header’s threads group (see API Reference), and the rules for keeping the code in line with it are in AGENTS.md. In short:

Lock

Guards

Held

Current::initMutex

Init and shutdown, one at a time

Throughout initWith() and shutdownInstance()

Current::mutex (shared)

The current instance pointer

Shared to copy it, exclusive to change it

BufferMap::mutex_ (shared)

The map of entries

Shared for lookups, exclusive to insert or remove

ConnectionPool::mutex_

The free list, live count, and closed flag

To lease or return a connection. A thread waiting for one releases it while it waits, and takes it back to check the pool again.

AbandonedCopies::mutex_

The list of copies and reservations

To reserve room or add a copy, and while querying the copies’ events, which doesn’t wait

Locks guard lookups, not work: a lookup copies a std::shared_ptr (or takes a Lease) and releases the lock, and no lock is held across a callback, a verbs call, or a HIP call, apart from the exceptions in the table. They’re the annotated wrappers in src/common/thread-safety.h, so Clang’s -Wthread-safety checks what each one guards and the order Current’s two are taken in.

The one wait in a transfer is for a connection, and the thread-local lease count turns the wait that could deadlock, a callback waiting for its own connection, into hipObjBusy.

Errors#

Every public function but hipObjGetErrorString() and hipObjGetVersionString(), which return strings, returns a hipObjError_t, built with makeError() (src/hipobj-private.h): an operation error, and the HIP error behind it where there is one. Each public function body is a function try block, so no exception crosses the C API. The functions that return a hipObjError_t turn any exception, whatever its type, into hipObjInternalError; the two that return strings return a fallback string instead ("Unknown error" and "0.0.0").

Internal functions report failure with bool, std::optional, or a result enum class, and hipobj.cpp maps those to error codes. Details that an error code can’t carry, such as a failed work completion’s status or a fallback to a staging buffer, go to stderr; the hot path prints nothing on success.

Configuration#

parseOptions() reads these environment variables once, in hipObjInit(). An unset or invalid value leaves the default.

Variable

Default

Effect

HIPOBJ_MAX_TRANSFERS

16

Connections in the pool (1 to 1024), so how many transfers run at once

HIPOBJ_STAGE_TIMEOUT_MS

30000

How long a staging copy can take (up to 24 hours)

HIPOBJ_DMABUF

on

Values starting 0, n, or N turn off dma-buf registration

HIPOBJ_REQUIRE_GPU_DIRECT

off

Any value but one starting 0, n, or N makes the staging fallback a failure

HIPOBJ_RDMA_DEBUG

off

Any value not starting 0 reports failed dma-buf exports

The completion timeout, 5 seconds, isn’t configurable.

Open Questions#

The code doesn’t show how these are meant to work:

When the client’s QP is connected

A transfer connects its QP to the peer only after recvReply() returns, that is, after the S3 response has arrived, and the token connection is never connected at all. An RC QP in INIT doesn’t respond to incoming RDMA, yet the test server connects its own QP and does its RDMA before it sends the reply. Issue #22 reported this ordering as a defect; it was closed by the “v2” two-phase protocol, which was later removed (#203). The emulated-hardware CI lanes assert that RDMA transfers succeed, including through the MinIO bridge, which uses the never-connected token connection, so it isn’t clear whether real NICs work with this ordering, or how they’re expected to.

Which completion a transfer waits for

The client posts no work requests, and on an RC QP, one-sided RDMA reads and writes by the peer don’t generate completions on the responder’s CQ. pollCompletion() waits for any completion and treats a timeout as failure. The comment in finishTransferAfterReply() says the check reports “no evidence of failure” rather than a verified transfer, but it isn’t clear which completion is expected to arrive in a successful transfer.