API Reference

API Reference#

This page documents the hipObject public C API, extracted from annotated source headers by Doxygen and rendered via Breathe.

Core Functionality#

typedef void *hipObjHandle_t#: Opaque handle to an S3 object for RDMA I/O.

hipObjError_t hipObjGetVersion(unsigned *major, unsigned *minor, unsigned *patch)#

Get the library’s version as numbers.

Writes the major, minor, and patch version of the hipObject library that’s running. These can differ from the HIPOBJ_VERSION_MAJOR, HIPOBJ_VERSION_MINOR, and HIPOBJ_VERSION_PATCH macros of the header that the application was compiled with, if it’s linked against a different hipObject.

A NULL pointer skips that part of the version, and the others are still written. It always returns hipObjSuccess.

Parameters:

major – Where to write the major version, or NULL
minor – Where to write the minor version, or NULL
patch – Where to write the patch version, or NULL

Returns:

hipObjError_t

const char *hipObjGetVersionString(void)#: Return the library version as a string.

hipObjError_t hipObjInit(hipObjConfig_t *config)#

Initialize the hipObject library.

Selects the closest NIC to the GPU (see hipObjConfig_t), opens it, and creates the RDMA connections that transfers use (see Thread Safety).

hipObjGetErrorString(), hipObjGetVersionString(), hipObjGetVersion(), hipObjParseRdmaReply(), hipObjTokenClientNic(), and hipObjPutRdmaToken() work without it, and hipObjShutdown() does nothing without it. The other functions return hipObjNotInitialized until it has succeeded.

It reads hipObject’s environment variables, so don’t change the environment on another thread while it runs (see Thread Safety).

Returns hipObjInvalidValue if config is NULL, flags isn’t 0, or gpuDevice is below -1 or names a GPU that doesn’t exist; hipObjAlreadyInitialized if hipObject is already initialized; hipObjNicNotFound if there’s no NIC to use; and hipObjRdmaError if libibverbs doesn’t load, a HIP call fails (hipError says which error), or the NIC can’t be opened or the connections created.

Parameters:: config – Pointer to configuration struct
Returns:: hipObjError_t

hipObjError_t hipObjShutdown(void)#

Shut down the hipObject library.

Releases all RDMA resources. All registered buffers are implicitly deregistered.

It doesn’t wait for calls in progress on other threads, or for the transfer whose callback calls it. They finish normally, and the resources they use are released when the last of them returns (see Thread Safety).

A staging buffer that a copy that timed out may still use (see hipObjBufSync()) isn’t freed.

Returns:: hipObjError_t

HIPOBJ_VERSION_MAJOR#: hipObject major version number

HIPOBJ_VERSION_MINOR#: hipObject minor version number

HIPOBJ_VERSION_PATCH#: hipObject patch version number

struct hipObjConfig_t#

#include <hipobj.h>

Configuration for hipObject initialization.

hipObject doesn’t make S3 requests, so it doesn’t use endpoint, region, accessKey, or secretKey: the application’s S3 client does.

gpuDevice is the HIP device index of the GPU that hipObject chooses the closest RDMA NIC to, or -1 for the calling thread’s current device. On a host without a GPU, set it to -1, and set nicHint.

nicHint lists the RDMA devices that hipObject can choose from, by name (as ibv_devices lists them), separated by commas; NULL or “” allows them all. If hipObject can’t choose a NIC by where the GPU is, as on a host without one, it opens the device that nicHint names, which must then be a single name.

Set the fields you don’t use to 0 or NULL.

Public Members

const char *accessKey#: Not used

const char *endpoint#: Not used

uint32_t flags#: Reserved, set to 0

int gpuDevice#: HIP device index, or -1

const char *nicHint#: RDMA NIC names, or NULL

const char *region#: Not used

const char *secretKey#: Not used

struct hipObjOps_t#

#include <hipobj.h>

Callback struct for S3 SDK integration.

The application provides these callbacks so that hipObject can inject RDMA tokens into S3 requests and receive RDMA reply tags from the server. This mirrors cuObject’s CUObjOps_t pattern.

The callbacks can call any hipObject function; see Thread Safety.

Public Members

int (*recvReply)(void *ctx, char *reply, size_t *replyLen)#

Called by hipObject to receive the RDMA reply. The application should extract x-amz-rdma-reply from the S3 response headers.

Param ctx:: User-supplied context pointer
Param reply:: Buffer to receive the reply tag
Param replyLen:: In: buffer size; out: actual length
Return:: 0 on success, negative on failure

int (*sendRequest)(void *ctx, const char *token, size_t tokenLen)#

Called by hipObject to send the RDMA token. The application should embed the token in the x-amz-rdma-token S3 request header.

Param ctx:: User-supplied context pointer
Param token:: Hex-encoded RDMA token
Param tokenLen:: Length of token in bytes
Return:: 0 on success, negative on failure

Thread Safety#

Every hipObject function can be called from any thread at any time, including from the hipObjOps_t callbacks, and calls on different threads run at the same time. A call waits for another only briefly, while the other changes something that both use, such as the list of registered buffers, with two exceptions: a transfer can wait for an RDMA connection (see below), and calls to hipObjInit() and hipObjShutdown() run one at a time. hipObjInit() reads hipObject’s environment variables, so it mustn’t run while another thread changes the environment, for example with setenv(), putenv(), or unsetenv().

Transfers. hipObjGet() and hipObjPut() calls run at the same time as each other, and as every other call. Each transfer uses an RDMA connection of its own, from a set that hipObjInit() creates. The HIPOBJ_MAX_TRANSFERS environment variable sets how many it creates, and so how many transfers can be in progress at once: a whole number from 1 to 1024. A value that isn’t one is ignored, leaving the default of 16. A transfer that starts while every connection is in use waits for one to be free.

Callbacks. A callback can call any function, including hipObjGet() and hipObjPut(). A transfer that a callback starts doesn’t wait for a connection, since the one it would wait for could be the one the callback’s own transfer holds: if none is free, it returns hipObjBusy. A callback that waits for another thread’s transfer can wait forever for the same reason, and hipObject can’t detect it. hipObject tells which transfers a callback starts by the thread it runs on, so a callback must return on the thread that called it. This matters if it runs in a fiber or coroutine that can move to another thread.

Buffers. hipObjBufDeregister() returns hipObjBusy, and leaves the buffer registered, while another call is using the buffer: a transfer, including the one whose callback deregisters it, hipObjBufSync(), or hipObjGetRdmaToken(). Calls on other buffers don’t wait for each other. As with any C function, calls that run at the same time mustn’t share memory that one of them writes to: overlapping parts of a registered buffer, or a token that one of them releases. For a buffer that’s staged through host memory (see hipObjBufSync()), hipObjPut() and HIPOBJ_SYNC_TO_HOST write the staging buffer, so they count as writing the part of the buffer they use.

Shutdown. hipObjShutdown() doesn’t wait for the calls in progress. They finish normally, and hipObject releases its RDMA resources and buffer registrations when the last of them returns. A transfer that’s waiting for a connection returns hipObjNotInitialized. Calls that start after hipObjShutdown(), on any thread, and need an initialized library fail with hipObjNotInitialized until hipObjInit() is called again, which can be before the earlier calls have finished.

Transfers made of several calls. A transfer that the application drives itself, with hipObjGetRdmaToken(), its own S3 request, and hipObjBufSync(), is several calls, and hipObject doesn’t know that the transfer is in progress between them. Other threads’ calls can run between them, and nothing keeps those from deregistering the buffer or shutting hipObject down: the application must keep the buffer registered, and hipObject initialized, until the transfer ends. Every token from hipObjGetRdmaToken() names the same RDMA connection, so the application must also make sure that such transfers don’t overlap each other. They can overlap hipObjGet() and hipObjPut() transfers, which use connections of their own.

Errors and Error Handling#

enum hipObjOpError_t#

Operation-level error codes.

Values:

enumerator hipObjSuccess#

enumerator hipObjInvalidValue#

enumerator hipObjNotInitialized#

enumerator hipObjAlreadyInitialized#

enumerator hipObjRdmaError#

enumerator hipObjS3Error#

enumerator hipObjBufNotRegistered#

enumerator hipObjBufAlreadyRegistered#

enumerator hipObjNicNotFound#

enumerator hipObjDmabufNotSupported#

enumerator hipObjSizeTooLarge#

enumerator hipObjInternalError#

enumerator hipObjBusy#: The buffer or connection is in use. A staging copy of the buffer that timed out is still running (see hipObjBufSync()); another call is using a buffer that hipObjBufDeregister() was given; or a transfer started from a callback found every RDMA connection in use (see Thread Safety). Try again later. hipObject added this code; cuObject doesn’t have it.

const char *hipObjGetErrorString(hipObjOpError_t err)#: Return a human-readable string for an operation error code.

HIPOBJ_BASE_ERR#: Base value for hipObject error codes.

struct hipObjError_t#: #include <hipobj.h>

Compound error type carrying both an operation error and an optional HIP error.

Buffer Registration#

hipObjError_t hipObjBufDeregister(void *devPtr)#

Deregister a previously registered buffer.

While another call is using the buffer (see Thread Safety), or a staging copy of the buffer that timed out is still running (see hipObjBufSync()), this returns hipObjBusy and leaves the buffer registered. Call it again once the call or the copy has finished, and don’t free the buffer until then.

Parameters:: devPtr – Pointer previously passed to hipObjBufRegister or hipObjBufRegisterHost
Returns:: hipObjError_t

hipObjError_t hipObjBufRegister(void *devPtr, size_t size)#

Exports the GPU buffer via dmabuf and registers it with the RDMA NIC. Maximum 4 GiB per registration. The buffer must lie within a single HIP allocation: returns hipObjInvalidValue if it runs past the end of the allocation that holds devPtr.

Parameters:

devPtr – Pointer returned by hipMalloc
size – Size of the buffer in bytes

Returns:

hipObjError_t

hipObjError_t hipObjBufRegisterHost(void *hostPtr, size_t size)#

Registers user-owned CPU RAM, such as a buffer allocated with malloc() or hipHostMalloc(), with the RDMA NIC. hipObject does not allocate or free the host buffer. Maximum 4 GiB per registration.

Parameters:

hostPtr – Pointer to CPU-accessible memory
size – Size of the buffer in bytes

Returns:

hipObjError_t

Data Transfer (GET / PUT)#

hipObjError_t hipObjBufSync(void *devPtr, size_t size, off_t offset, int direction)#

Stage a registered buffer between device and host memory.

When a device pointer cannot be registered with the NIC directly (no dmabuf support, or a fabric with no peer-to-peer path to the GPU), hipObjBufRegister() registers a host staging buffer instead and the peer’s RDMA reads and writes land there. Callers driving a transfer through hipObjGetRdmaToken() must therefore call this with HIPOBJ_SYNC_TO_HOST before a PUT and HIPOBJ_SYNC_TO_DEVICE after a GET. It is a no-op, reporting success, for a directly registered buffer; hipObjGet() and hipObjPut() do it themselves.

When this returns hipObjSuccess, the copy has finished. If the copy doesn’t finish within the staging timeout (the HIPOBJ_STAGE_TIMEOUT_MS environment variable), this returns hipObjInternalError and leaves the copy running. It can still read or write the buffer after this returns. Until it finishes, hipObject keeps the staging buffer it uses, and calls that start after the copy times out can’t use the buffer: hipObjGet(), hipObjPut(), hipObjBufSync(), and hipObjGetRdmaToken() return hipObjBusy for it, and hipObjBufDeregister() returns hipObjBusy and leaves it registered. Calls that had already started aren’t stopped. Once the copy has finished, the buffer can be used and deregistered again. If it never finishes, the staging buffer is never freed.

hipObject doesn’t wait for GPU work that uses the buffer, on any stream. Before HIPOBJ_SYNC_TO_HOST, make sure that the GPU work that writes the buffer has finished, for example with hipStreamSynchronize(). Before HIPOBJ_SYNC_TO_DEVICE, make sure that no GPU work that reads or writes the buffer is still running.

hipObjError_t hipObjGet(hipObjHandle_t handle, void *devPtr, size_t size, off_t offset, hipObjOps_t *ops, void *ctx)#

GET: fetch an S3 object into a registered buffer via RDMA.

The server performs an RDMA WRITE to push data into the registered buffer, starting offset bytes into it. The bytes from offset to offset + size must lie inside the buffer. hipObject doesn’t send an object offset; the request the callbacks make selects the part of the object to fetch.

Make sure that no GPU work that reads or writes the buffer is still running when you call this. When it returns hipObjSuccess, the data is in the buffer.

If every RDMA connection is in use by other transfers, this waits for one to be free, or, if it’s called from a callback, returns hipObjBusy (see Thread Safety).

For a buffer that’s staged through host memory (see hipObjBufSync()), if the copy into the buffer doesn’t finish within the staging timeout (the HIPOBJ_STAGE_TIMEOUT_MS environment variable), this returns hipObjInternalError and leaves the copy running. It can still write the buffer after this returns, and until it finishes, calls that start after this returns can’t use or deregister the buffer (see hipObjBufSync()).

Parameters:

handle – S3 object handle (from application)
devPtr – Registered GPU or host buffer
size – Number of bytes to transfer
offset – Byte offset into the registered buffer
ops – S3 SDK callbacks
ctx – User context passed to callbacks

Returns:

hipObjError_t

hipObjError_t hipObjGetRdmaToken(const void *devPtr, size_t size, int op, char **outToken)#

Mint a hex-encoded RC RDMA token for a registered buffer.

The caller must release *outToken with hipObjPutRdmaToken(). op is HIPOBJ_RDMA_OP_PUT or HIPOBJ_RDMA_OP_GET (reserved).

Returns hipObjBusy while a staging copy of the buffer that timed out is still running (see hipObjBufSync()).

hipObjError_t hipObjParseRdmaReply(const char *reply, size_t replyLen, int *httpCode)#

Parse an x-amz-rdma-reply header value.

On success writes the HTTP-style reply code (200, 204, 206, 501) to httpCode.

hipObjError_t hipObjPut(hipObjHandle_t handle, const void *devPtr, size_t size, off_t offset, hipObjOps_t *ops, void *ctx)#

PUT: store a registered buffer to an S3 object via RDMA.

The server performs an RDMA READ to pull data from the registered buffer, starting offset bytes into it. The bytes from offset to offset + size must lie inside the buffer. hipObject doesn’t send an object offset; the request the callbacks make selects where the data goes in the object.

hipObject doesn’t wait for GPU work that writes the buffer, on any stream. Make sure that it has finished, for example with hipStreamSynchronize(), before you call this.

If every RDMA connection is in use by other transfers, this waits for one to be free, or, if it’s called from a callback, returns hipObjBusy (see Thread Safety).

For a buffer that’s staged through host memory (see hipObjBufSync()), if the copy out of the buffer doesn’t finish within the staging timeout (the HIPOBJ_STAGE_TIMEOUT_MS environment variable), this returns hipObjInternalError and leaves the copy running. It can still read the buffer after this returns, and until it finishes, calls that start after this returns can’t use or deregister the buffer (see hipObjBufSync()).

Parameters:

handle – S3 object handle (from application)
devPtr – Registered GPU or host buffer
size – Number of bytes to transfer
offset – Byte offset into the registered buffer
ops – S3 SDK callbacks
ctx – User context passed to callbacks

Returns:

hipObjError_t

hipObjError_t hipObjPutRdmaToken(char *token)#: Release a token allocated by hipObjGetRdmaToken()

hipObjError_t hipObjTokenClientNic(const char *token, char *nicIp, size_t nicIpLen)#

Extract client NIC IPv4 from a minted RDMA token.

Writes a dotted-quad address into nicIp when the token GID carries an IPv4-mapped RoCEv2 suffix. Returns hipObjSuccess with an empty string when no address is encoded.

HIPOBJ_RDMA_OP_GET#: RDMA token operation: server RDMA WRITE (client GET)

HIPOBJ_RDMA_OP_PUT#: RDMA token operation: server RDMA READ (client PUT)

HIPOBJ_RDMA_REPLY_NOT_IMPLEMENTED#: x-amz-rdma-reply: server declined RDMA (HTTP fallback)

HIPOBJ_SYNC_TO_DEVICE#: hipObjBufSync direction: staging buffer -> device

HIPOBJ_SYNC_TO_HOST#: hipObjBufSync direction: device -> staging buffer

API Reference

Contents

API Reference#

Core Functionality#

Thread Safety#

Errors and Error Handling#

Buffer Registration#

Data Transfer (GET / PUT)#