[libc-commits] [libc] [libc][Docs] Update the GPU RPC documentation (PR #79069)
Nick Desaulniers via libc-commits
libc-commits at lists.llvm.org
Tue Jan 23 10:30:56 PST 2024
================
@@ -11,10 +11,298 @@ Remote Procedure Calls
Remote Procedure Call Implementation
====================================
-Certain features from the standard C library, such as allocation or printing,
-require support from the operating system. We instead implement a remote
-procedure call (RPC) interface to allow submitting work from the GPU to a host
-server that forwards it to the host system.
+Traditionally, the C library abstracts over several functions that interface
+with the platform's operating system through system calls. The GPU however does
+not provide an operating system that can handle target dependent operations.
+Instead, we implemented remote procedure calls to interface with the host's
+operating system while executing on a GPU.
+
+We implemented remote procedure calls using unified virtual memory to create a
+shared communicate channel between the two processes. This memory is often
+pinned memory that can be accessed asynchronously and atomically by multiple
+processes simultaneously. This supports means that we can simply provide mutual
+exclusion on a shared better to swap work back and forth between the host system
+and the GPU. We can then use this to create a simple client-server protocol
+using this shared memory.
+
+This work treats the GPU as a client and the host as a server. The client
+initiates a communication while the server listens for them. In order to
+communicate between the host and the device, we simply maintain a buffer of
+memory and two mailboxes. One mailbox is write-only while the other is
+read-only. This exposes three primitive operations: using the buffer, giving
+away ownership, and waiting for ownership. This is implemented as a half-duplex
+transmission channel between the two sides. We decided to assign ownership of
+the buffer to the client when the inbox and outbox bits are equal and to the
+server when they are not.
+
+In order to make this transmission channel thread-safe, we abstract ownership of
+the given mailbox pair and buffer around a port, effectively acting as a lock
+and an index into the allocated buffer slice. The server and device have
+independent locks around the given port. In this scheme, the buffer can be used
+to communicate intent and data generically with the server. We them simply
+provide multiple copies of this protocol and expose them as multiple ports.
+
+If this were simply a standard CPU system, this would be sufficient. However,
+GPUs have my unique architectural challenges. First, GPU threads execute in
+lock-step with each other in groups typically called warps or wavefronts. We
+need to target the smallest unit of independent parallelism, so the RPC
+interface needs to handle an entire group of threads at once. This is done by
+increasing the size of the buffer and adding a thread mask argument so the
+server knows which threads are active when it handles the communication. Second,
+GPUs generally have no forward progress guarantees. In order to guarantee we do
+not encounter deadlocks while executing it is required that the number of ports
+matches the maximum amount of hardware parallelism on the device. It is also
+very important that the thread mask remains consistent while interfacing with
+the port.
+
+.. image:: ./rpc-diagram.svg
+ :width: 75%
+ :align: center
+
+The above diagram outlines the architecture of the RPC interface. For clarity
+the following list will explain the operations done by the client and server
+respectively when initiating a communication.
+
+First, a communication from the perspective of the client:
+
+* The client searches for an available port and claims the lock.
+* The client checks that the port is still available to the current device and
+ continues if so.
+* The client writes its data to the fixed-size packet and toggles its outbox.
+* The client waits until its inbox matches its outbox.
+* The client reads the data from the fixed-size packet.
+* The client closes the port and continues executing.
+
+Now, the same communication from the perspective of the server:
+
+* The server searches for an available port with pending work and claims the
+ lock.
+* The server checks that the port is still available to the current device.
+* The server reads the opcode to perform the expected operation, in this
+ case a receive and then send.
+* The server reads the data from the fixed-size packet.
+* The server writes its data to the fixed-size packet and toggles its outbox.
+* The server closes the port and continues searching for ports that need to be
+ serviced
+
+This architecture currently requires that the host periodically checks the RPC
+server's buffer for ports with pending work. Note that a port can be closed
+without waiting for its submitted work to be completed. This allows us to model
+asynchronous operations that do not need to wait until the server has completed
+them. If an operation requires more data than the fixed size buffer, we simply
+send multiple packets back and forth in a streaming fashion.
+
+Server Library
+--------------
+
+The RPC server's basic functionality is provided by the LLVM C library. A static
+library called ``libllvmlibc_rpc_server.a`` includes handling for the basic
+operations, such as printing or exiting. This has a small API that handles
+setting up the unified buffer and an interface to check the opcodes.
+
+Some operations are too divergent to provide generic implementations for, such
+as allocating device accessible memory. For these cases, we provide a callback
+registration scheme to add a custom handler for any given opcode through the
+port API. More information can be found in the installed header
+``<install>/include/gpu-none-llvm/rpc_server.h``.
+
+Client Example
+--------------
+
+The Client API is not currently exported by the LLVM C library. This is
+primarily due to being written in C++ and relying on internal data structures.
+It uses a simple send and receive interface with a fixed-size packet. The
+following example uses the RPC interface to call a function pointer on the
+server.
+
+This code first opens a port with the given opcode to facilitate the
+communication. It then copies over the argument struct to the server using the
+``send_n`` interface to stream arbitrary bytes. The next send operation provides
+the server with the function pointer that will be executed. The final receive
+operation is a no-op and simply forces the client to wait until the server is
+done. It can be omitted if asynchronous execution is desired.
+
+.. code-block:: c++
+
+ void rpc_host_call(void *fn, void *data, size_t size) {
+ rpc::Client::Port port = rpc::client.open<RPC_HOST_CALL>();
+ port.send_n(data, size);
+ port.send([=](rpc::Buffer *buffer) {
+ buffer->data[0] = reinterpret_cast<uintptr_t>(fn);
+ });
+ port.recv([](rpc::Buffer *) {});
+ port.close();
+ }
+
+Server Example
+--------------
+
+This example shows the server-side handling of the previous client example. When
+the server is checked, if there are any ports with pending work it will check
+the opcode and perform the appropriate action. In this case, the action is to
+call a function pointer provided by the client.
+
+In this example, the server simply runs forever in a separate thread for
+brevity's sake. Because the client is a GPU potentially handling several threads
+at once, the server needs to loop over all the active threads on the GPU. We
+abstract this into the ``lane_size`` variable, which is simply the device's warp
+or wavefront size. The identifier is simply the threads index into the current
+warp or wavefront. We allocate memory to copy the struct data into, and then
+call the given function pointer with that copied data. The final send simply
+signals completion and uses the implicit thread mask to delete the temporary
+data.
+
+.. code-block:: c++
+
+ for(;;) {
+ auto port = server.try_open(index);
+ if (!port)
+ return continue;
+
+ switch(port->get_opcode()) {
+ case RPC_HOST_CALL: {
+ uint64_t sizes[LANE_SIZE];
+ void *args[LANE_SIZE];
+ port->recv_n(args, sizes, [&](uint64_t size) { return new char[size]; });
+ port->recv([&](rpc::Buffer *buffer, uint32_t id) {
+ reinterpret_cast<void (*)(void *)>(buffer->data[0])(args[id]);
+ });
+ port->send([&](rpc::Buffer *, uint32_t id) {
+ delete[] reinterpret_cast<uint8_t *>(args[id]);
+ });
+ break;
+ }
+ default:
+ port->recv([](rpc::Buffer *) {});
+ break;
+ }
+ }
+
+CUDA Server Example
+-------------------
+
+The following code shows an example of using the exported RPC interface along
+with the C library to manually configure a working server using the CUDA
+language. Other runtimes can use the presence of the ``__llvm_libc_rpc_client``
+in the GPU executable as an indicator for whether or not the server can be
+checked. These details should ideally be handled by the GPU language runtime,
+but the following example shows how it can be used by a standard user.
+
+.. code-block:: cuda
+
+ #include <cstdio>
+ #include <cstdlib>
+ #include <cuda_runtime.h>
+
+ #include <gpu-none-llvm/rpc_server.h>
+
+ [[noreturn]] void handle_error(cudaError_t err) {
+ fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
+ exit(EXIT_FAILURE);
+ }
+
+ [[noreturn]] void handle_error(rpc_status_t err) {
+ fprintf(stderr, "RPC error: %d\n", err);
+ exit(EXIT_FAILURE);
+ }
+
+ // The handle to the RPC client provided by the C library.
+ extern "C" __device__ void *__llvm_libc_rpc_client;
+
+ __global__ void get_client_ptr(void **ptr) { *ptr = __llvm_libc_rpc_client; }
+
+ // Obtain the RPC client's handle from the device. The CUDA language cannot look
+ // up the symbol directly like the driver API, so we launch a kernel to read it.
+ void *get_rpc_client() {
+ void *rpc_client = nullptr;
+ void **rpc_client_d = nullptr;
+
+ if (cudaError_t err = cudaMalloc(&rpc_client_d, sizeof(void *)))
+ handle_error(err);
+ get_client_ptr<<<1, 1>>>(rpc_client_d);
+ if (cudaError_t err = cudaDeviceSynchronize())
+ handle_error(err);
+ if (cudaError_t err = cudaMemcpy(&rpc_client, rpc_client_d, sizeof(void *),
+ cudaMemcpyDeviceToHost))
+ handle_error(err);
+ return rpc_client;
+ }
+
+ // Routines to allocate mapped memory that both the host and the device can
+ // access asychonrously to communicate with eachother.
+ void *alloc_host(size_t size, void *) {
+ void *sharable_ptr;
+ if (cudaError_t err = cudaMallocHost(&sharable_ptr, sizeof(void *)))
+ handle_error(err);
+ return sharable_ptr;
+ };
+
+ void free_host(void *ptr, void *) {
+ if (cudaError_t err = cudaFreeHost(ptr))
+ handle_error(err);
+ }
+
+ // The device-side overload of the standard C function to call.
+ extern "C" __device__ int puts(const char *);
+
+ // Calls the C library function from the GPU C library.
+ __global__ void hello() { puts("Hello world!"); }
+
+ int main() {
+ int device = 0;
+ // Initialize the RPC server to run on a single device.
+ if (rpc_status_t err = rpc_init(/*num_device=*/1))
+ handle_error(err);
+
+ // Initialize the RPC server to run on the given device.
+ if (rpc_status_t err =
+ rpc_server_init(device, RPC_MAXIMUM_PORT_COUNT,
+ /*warp_size=*/32, alloc_host, /*data=*/nullptr))
+ handle_error(err);
+
+ // Initialize the RPC client by copying the buffer to the device's handle.
+ void *rpc_client = get_rpc_client();
+ if (cudaError_t err =
+ cudaMemcpy(rpc_client, rpc_get_client_buffer(device),
+ rpc_get_client_size(), cudaMemcpyHostToDevice))
+ handle_error(err);
+
+ cudaStream_t stream;
+ if (cudaError_t err = cudaStreamCreate(&stream))
+ handle_error(err);
+
+ // Execute the kernel.
+ hello<<<1, 1, 0, stream>>>();
+
+ // While the kernel is executing, check the RPC server for work to do.
+ while (cudaStreamQuery(stream) == cudaErrorNotReady)
+ if (rpc_status_t err = rpc_handle_server(device))
+ handle_error(err);
+
+ // Shut down the server running on the given device.
+ if (rpc_status_t err =
+ rpc_server_shutdown(device, free_host, /*data=*/nullptr))
+ handle_error(err);
+
+ // Shut down the entire RPC server interface.
+ if (rpc_status_t err = rpc_shutdown())
+ handle_error(err);
+
+ return EXIT_SUCCESS;
+ }
+
+The above code must be compiled in CUDA's relocatable device code mode and with
+the advanced offloading driver to link in the library. Currently this can be
+done with the following invocation. Using LTO avoids the overhead normally
+associated with relocatable device code linking.
+
+.. code-block:: sh
+
+ $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart -lcgpu \
+ -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \
----------------
nickdesaulniers wrote:
add `/` before `include` for the `-I` to be consistent with `-L`
https://github.com/llvm/llvm-project/pull/79069
More information about the libc-commits
mailing list