266 lines · plain
1.. _libc_gpu_rpc:2 3======================4Remote Procedure Calls5======================6 7.. contents:: Table of Contents8 :depth: 49 :local:10 11Remote Procedure Call Implementation12====================================13 14Traditionally, the C library abstracts over several functions that interface15with the platform's operating system through system calls. The GPU however does16not provide an operating system that can handle target dependent operations.17Instead, we implemented remote procedure calls to interface with the host's18operating system while executing on a GPU.19 20We implemented remote procedure calls using unified virtual memory to create a21shared communicate channel between the two processes. This memory is often22pinned memory that can be accessed asynchronously and atomically by multiple23processes simultaneously. This support means that we can simply provide mutual24exclusion on a shared buffer to swap work back and forth between the host system25and the GPU. We can then use this to create a simple client-server protocol26using this shared memory.27 28This work treats the GPU as a client and the host as a server. The client29initiates a communication while the server listens for them. In order to30communicate between the host and the device, we simply maintain a buffer of31memory and two mailboxes. One mailbox is write-only while the other is32read-only. This exposes three primitive operations: using the buffer, giving33away ownership, and waiting for ownership. This is implemented as a half-duplex34transmission channel between the two sides. We decided to assign ownership of35the buffer to the client when the inbox and outbox bits are equal and to the36server when they are not.37 38In order to make this transmission channel thread-safe, we abstract ownership of39the given mailbox pair and buffer around a port, effectively acting as a lock40and an index into the allocated buffer slice. The server and device have41independent locks around the given port. In this scheme, the buffer can be used42to communicate intent and data generically with the server. We then simply43provide multiple copies of this protocol and expose them as multiple ports.44 45If this were simply a standard CPU system, this would be sufficient. However,46GPUs have my unique architectural challenges. First, GPU threads execute in47lock-step with each other in groups typically called warps or wavefronts. We48need to target the smallest unit of independent parallelism, so the RPC49interface needs to handle an entire group of threads at once. This is done by50increasing the size of the buffer and adding a thread mask argument so the51server knows which threads are active when it handles the communication. Second,52GPUs generally have no forward progress guarantees. In order to guarantee we do53not encounter deadlocks while executing it is required that the number of ports54matches the maximum amount of hardware parallelism on the device. It is also55very important that the thread mask remains consistent while interfacing with56the port.57 58.. image:: ./rpc-diagram.svg59 :width: 75%60 :align: center61 62The above diagram outlines the architecture of the RPC interface. For clarity63the following list will explain the operations done by the client and server64respectively when initiating a communication.65 66First, a communication from the perspective of the client:67 68* The client searches for an available port and claims the lock.69* The client checks that the port is still available to the current device and70 continues if so.71* The client writes its data to the fixed-size packet and toggles its outbox.72* The client waits until its inbox matches its outbox.73* The client reads the data from the fixed-size packet.74* The client closes the port and continues executing.75 76Now, the same communication from the perspective of the server:77 78* The server searches for an available port with pending work and claims the79 lock.80* The server checks that the port is still available to the current device.81* The server reads the opcode to perform the expected operation, in this82 case a receive and then send.83* The server reads the data from the fixed-size packet.84* The server writes its data to the fixed-size packet and toggles its outbox.85* The server closes the port and continues searching for ports that need to be86 serviced87 88This architecture currently requires that the host periodically checks the RPC89server's buffer for ports with pending work. Note that a port can be closed90without waiting for its submitted work to be completed. This allows us to model91asynchronous operations that do not need to wait until the server has completed92them. If an operation requires more data than the fixed size buffer, we simply93send multiple packets back and forth in a streaming fashion.94 95Client Example96--------------97 98The Client API is not currently exported by the LLVM C library. This is99primarily due to being written in C++ and relying on internal data structures.100It uses a simple send and receive interface with a fixed-size packet. The101following example uses the RPC interface to call a function pointer on the102server.103 104This code first opens a port with the given opcode to facilitate the105communication. It then copies over the argument struct to the server using the106``send_n`` interface to stream arbitrary bytes. The next send operation provides107the server with the function pointer that will be executed. The final receive108operation is a no-op and simply forces the client to wait until the server is109done. It can be omitted if asynchronous execution is desired.110 111.. code-block:: c++112 113 void rpc_host_call(void *fn, void *data, size_t size) {114 rpc::Client::Port port = rpc::client.open<RPC_HOST_CALL>();115 port.send_n(data, size);116 port.send([=](rpc::Buffer *buffer) {117 buffer->data[0] = reinterpret_cast<uintptr_t>(fn);118 });119 port.recv([](rpc::Buffer *) {});120 port.close();121 }122 123Server Example124--------------125 126This example shows the server-side handling of the previous client example. When127the server is checked, if there are any ports with pending work it will check128the opcode and perform the appropriate action. In this case, the action is to129call a function pointer provided by the client.130 131In this example, the server simply runs forever in a separate thread for132brevity's sake. Because the client is a GPU potentially handling several threads133at once, the server needs to loop over all the active threads on the GPU. We134abstract this into the ``lane_size`` variable, which is simply the device's warp135or wavefront size. The identifier is simply the threads index into the current136warp or wavefront. We allocate memory to copy the struct data into, and then137call the given function pointer with that copied data. The final send simply138signals completion and uses the implicit thread mask to delete the temporary139data.140 141.. code-block:: c++142 143 for(;;) {144 auto port = server.try_open(index);145 if (!port)146 return continue;147 148 switch(port->get_opcode()) {149 case RPC_HOST_CALL: {150 uint64_t sizes[LANE_SIZE];151 void *args[LANE_SIZE];152 port->recv_n(args, sizes, [&](uint64_t size) { return new char[size]; });153 port->recv([&](rpc::Buffer *buffer, uint32_t id) {154 reinterpret_cast<void (*)(void *)>(buffer->data[0])(args[id]);155 });156 port->send([&](rpc::Buffer *, uint32_t id) {157 delete[] reinterpret_cast<uint8_t *>(args[id]);158 });159 break;160 }161 default:162 port->recv([](rpc::Buffer *) {});163 break;164 }165 }166 167CUDA Server Example168-------------------169 170The following code shows an example of using the exported RPC interface along171with the C library to manually configure a working server using the CUDA172language. Other runtimes can use the presence of the ``__llvm_rpc_client``173in the GPU executable as an indicator for whether or not the server can be174checked. These details should ideally be handled by the GPU language runtime,175but the following example shows how it can be used by a standard user.176 177.. _libc_gpu_cuda_server:178 179.. code-block:: cuda180 181 #include <cstdio>182 #include <cstdlib>183 #include <cuda_runtime.h>184 185 #include <shared/rpc.h>186 #include <shared/rpc_opcodes.h>187 #include <shared/rpc_server.h>188 189 [[noreturn]] void handle_error(cudaError_t err) {190 fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));191 exit(EXIT_FAILURE);192 }193 194 // Routes the library symbol into the CUDA runtime interface.195 [[gnu::weak]] __device__ rpc::Client client asm("__llvm_rpc_client");196 197 // The device-side overload of the standard C function to call.198 extern "C" __device__ int puts(const char *);199 200 // Calls the C library function from the GPU C library.201 __global__ void hello() { puts("Hello world!"); }202 203 int main() {204 void *rpc_client = nullptr;205 if (cudaError_t err = cudaGetSymbolAddress(&rpc_client, client))206 handle_error(err);207 208 // Initialize the RPC client and server interface.209 uint32_t warp_size = 32;210 void *rpc_buffer = nullptr;211 if (cudaError_t err = cudaMallocHost(212 &rpc_buffer,213 rpc::Server::allocation_size(warp_size, rpc::MAX_PORT_COUNT)))214 handle_error(err);215 rpc::Server server(rpc::MAX_PORT_COUNT, rpc_buffer);216 rpc::Client client(rpc::MAX_PORT_COUNT, rpc_buffer);217 218 // Initialize the client on the device so it can communicate with the server.219 if (cudaError_t err = cudaMemcpy(rpc_client, &client, sizeof(rpc::Client),220 cudaMemcpyHostToDevice))221 handle_error(err);222 223 cudaStream_t stream;224 if (cudaError_t err = cudaStreamCreate(&stream))225 handle_error(err);226 227 // Execute the kernel.228 hello<<<1, 1, 0, stream>>>();229 230 // While the kernel is executing, check the RPC server for work to do.231 // Requires non-blocking CUDA kernels but avoids a separate thread.232 do {233 auto port = server.try_open(warp_size, /*index=*/0);234 if (!port)235 continue;236 237 // Only available in-tree from the 'libc' sources.238 handle_libc_opcodes(*port, warp_size);239 port->close();240 } while (cudaStreamQuery(stream) == cudaErrorNotReady);241 }242 243The above code must be compiled in CUDA's relocatable device code mode and with244the advanced offloading driver to link in the library. Currently this can be245done with the following invocation. Using LTO avoids the overhead normally246associated with relocatable device code linking. The C library for GPU's247handling is included through the ``shared/`` directory. This is not currently248installed as it does not use a stable interface.249 250 251.. code-block:: sh252 253 $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart \254 -I<install-path>include -L<install-path>/lib -Xoffload-linker -lc \255 -O3 -foffload-lto -o hello256 $> ./hello257 Hello world!258 259Extensions260----------261 262The opcode is a 32-bit integer that must be unique to the requested operation.263All opcodes used by ``libc`` internally have the character ``c`` in the most264significant byte. Any other opcode is available for use outside of the ``libc``265implementation.266