[libc] Make RPC server handling header only

Summary:
This patch moves the RPC server handling to be a header only utility
stored in the `shared/` directory. This is intended to be shared within
LLVM for the loaders and `offload/` handling.

Generally, this makes it easier to share code without weird
cross-project binaries being plucked out of the build system. It also
allows us to soon move the loader interface out of the `libc` project so
that we don't need to bootstrap those and can build them in LLVM.

Author: jhuber6

libc/docs/gpu/rpc.rst

@@ -184,6 +184,7 @@ but the following example shows how it can be used by a standard user.
 
   #include <shared/rpc.h>
   #include <shared/rpc_opcodes.h>
+  #include <shared/rpc_server.h>
 
   [[noreturn]] void handle_error(cudaError_t err) {
     fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
@@ -230,10 +231,10 @@ but the following example shows how it can be used by a standard user.
     // Requires non-blocking CUDA kernels but avoids a separate thread.
     do {
       auto port = server.try_open(warp_size, /*index=*/0);
-      // From libllvmlibc_rpc_server.a in the installation.
       if (!port)
         continue;
 
+      // Only available in-tree from the 'libc' sources.
       handle_libc_opcodes(*port, warp_size);
       port->close();
     } while (cudaStreamQuery(stream) == cudaErrorNotReady);
@@ -242,14 +243,16 @@ but the following example shows how it can be used by a standard user.
 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. The C library for GPUs is
-linked in by forwarding the static library to the device-side link job.
+associated with relocatable device code linking. The C library for GPU's
+handling is included through the ``shared/`` directory. This is not currently
+installed as it does not use a stable interface.
+
 
 .. code-block:: sh
 
   $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart \
-       -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \
-       -Xoffload-linker -lc -O3 -foffload-lto -o hello
+       -I<install-path>include -L<install-path>/lib -Xoffload-linker -lc \
+       -O3 -foffload-lto -o hello
   $> ./hello
   Hello world!
 

libc/shared/rpc_opcodes.h

@@ -50,10 +50,4 @@ typedef enum {
 
 #undef LLVM_LIBC_OPCODE
 
-namespace rpc {
-// The implementation of this function currently lives in the utility directory
-// at 'utils/gpu/server/rpc_server.cpp'.
-rpc::Status handle_libc_opcodes(rpc::Server::Port &port, uint32_t num_lanes);
-} // namespace rpc
-
 #endif // LLVM_LIBC_SHARED_RPC_OPCODES_H

libc/utils/gpu/server/rpc_server.cpp renamed to libc/shared/rpc_server.h

@@ -1,16 +1,30 @@
-//===-- Shared memory RPC server instantiation ------------------*- C++ -*-===//
+//===-- RPC server handling -----------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
+#ifndef LLVM_LIBC_SHARED_RPC_SERVER_H
+#define LLVM_LIBC_SHARED_RPC_SERVER_H
+
 // Workaround for missing __has_builtin in < GCC 10.
 #ifndef __has_builtin
 #define __has_builtin(x) 0
 #endif
 
+// Configs for using the LLVM libc writer interface.
+#define LIBC_COPT_USE_C_ASSERT
+#define LIBC_COPT_MEMCPY_USE_EMBEDDED_TINY
+#define LIBC_COPT_ARRAY_ARG_LIST
+#define LIBC_COPT_PRINTF_DISABLE_WRITE_INT
+#define LIBC_COPT_PRINTF_DISABLE_INDEX_MODE
+#define LIBC_COPT_PRINTF_DISABLE_STRERROR
+
+// The 'long double' type is 8 byte
+#define LIBC_TYPES_LONG_DOUBLE_IS_FLOAT64
+
 #include "shared/rpc.h"
 #include "shared/rpc_opcodes.h"
 
@@ -24,6 +38,7 @@
 #include <stdlib.h>
 
 namespace LIBC_NAMESPACE {
+namespace internal {
 
 // Minimal replacement for 'std::vector' that works for trivial types.
 template <typename T> class TempVector {
@@ -88,15 +103,15 @@ struct TempStorage {
   TempVector<char *> storage;
 };
 
-enum Stream {
-  File = 0,
-  Stdin = 1,
-  Stdout = 2,
-  Stderr = 3,
-};
-
 // Get the associated stream out of an encoded number.
-LIBC_INLINE ::FILE *to_stream(uintptr_t f) {
+static inline ::FILE *to_stream(uintptr_t f) {
+  enum Stream {
+    File = 0,
+    Stdin = 1,
+    Stdout = 2,
+    Stderr = 3,
+  };
+
   ::FILE *stream = reinterpret_cast<FILE *>(f & ~0x3ull);
   Stream type = static_cast<Stream>(f & 0x3ull);
   if (type == Stdin)
@@ -109,7 +124,8 @@ LIBC_INLINE ::FILE *to_stream(uintptr_t f) {
 }
 
 template <bool packed, uint32_t num_lanes>
-static void handle_printf(rpc::Server::Port &port, TempStorage &temp_storage) {
+static inline void handle_printf(rpc::Server::Port &port,
+                                 TempStorage &temp_storage) {
   FILE *files[num_lanes] = {nullptr};
   // Get the appropriate output stream to use.
   if (port.get_opcode() == LIBC_PRINTF_TO_STREAM ||
@@ -282,7 +298,7 @@ static void handle_printf(rpc::Server::Port &port, TempStorage &temp_storage) {
 }
 
 template <uint32_t num_lanes>
-rpc::Status handle_port_impl(rpc::Server::Port &port) {
+static inline rpc::Status handle_port_impl(rpc::Server::Port &port) {
   TempStorage temp_storage;
 
   switch (port.get_opcode()) {
@@ -498,21 +514,24 @@ rpc::Status handle_port_impl(rpc::Server::Port &port) {
   return rpc::RPC_SUCCESS;
 }
 
+} // namespace internal
 } // namespace LIBC_NAMESPACE
 
 namespace rpc {
-// The implementation of this function currently lives in the utility directory
-// at 'utils/gpu/server/rpc_server.cpp'.
-rpc::Status handle_libc_opcodes(rpc::Server::Port &port, uint32_t num_lanes) {
+// Handles any opcode generated from the 'libc' client code.
+static inline rpc::Status handle_libc_opcodes(rpc::Server::Port &port,
+                                              uint32_t num_lanes) {
   switch (num_lanes) {
   case 1:
-    return LIBC_NAMESPACE::handle_port_impl<1>(port);
+    return LIBC_NAMESPACE::internal::handle_port_impl<1>(port);
   case 32:
-    return LIBC_NAMESPACE::handle_port_impl<32>(port);
+    return LIBC_NAMESPACE::internal::handle_port_impl<32>(port);
   case 64:
-    return LIBC_NAMESPACE::handle_port_impl<64>(port);
+    return LIBC_NAMESPACE::internal::handle_port_impl<64>(port);
   default:
     return rpc::RPC_ERROR;
   }
 }
 } // namespace rpc
+
+#endif // LLVM_LIBC_SHARED_RPC_SERVER_H

libc/utils/gpu/CMakeLists.txt

@@ -1,2 +1 @@
-add_subdirectory(server)
 add_subdirectory(loader)

libc/utils/gpu/loader/CMakeLists.txt

@@ -1,5 +1,8 @@
 add_library(gpu_loader OBJECT Main.cpp)
 
+include(FindLibcCommonUtils)
+target_link_libraries(gpu_loader PUBLIC llvm-libc-common-utilities)
+
 target_include_directories(gpu_loader PUBLIC
   ${CMAKE_CURRENT_SOURCE_DIR}
   ${LIBC_SOURCE_DIR}/include

libc/utils/gpu/loader/Loader.h

@@ -13,6 +13,7 @@
 
 #include "shared/rpc.h"
 #include "shared/rpc_opcodes.h"
+#include "shared/rpc_server.h"
 
 #include <cstddef>
 #include <cstdint>

libc/utils/gpu/loader/amdgpu/CMakeLists.txt

@@ -7,10 +7,4 @@ set(LLVM_LINK_COMPONENTS
   )
 
 add_llvm_executable(amdhsa-loader amdhsa-loader.cpp)
-
-target_link_libraries(amdhsa-loader
-  PRIVATE
-  hsa-runtime64::hsa-runtime64
-  gpu_loader
-  llvmlibc_rpc_server
-)
+target_link_libraries(amdhsa-loader PRIVATE hsa-runtime64::hsa-runtime64 gpu_loader)

libc/utils/gpu/loader/nvptx/CMakeLists.txt

@@ -6,10 +6,4 @@ set(LLVM_LINK_COMPONENTS
   )
 
 add_llvm_executable(nvptx-loader nvptx-loader.cpp)
-
-target_link_libraries(nvptx-loader
-  PRIVATE
-  gpu_loader
-  llvmlibc_rpc_server
-  CUDA::cuda_driver
-)
+target_link_libraries(nvptx-loader PRIVATE gpu_loader CUDA::cuda_driver)

libc/utils/gpu/server/CMakeLists.txt

@@ -0,0 +1,19 @@
+#===--------------------------------------------------------------------===//
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for details.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+#===--------------------------------------------------------------------===//
+
+if(NOT TARGET llvm-libc-common-utilities)
+  set(libc_path ${CMAKE_CURRENT_LIST_DIR}/../../../libc)
+  if (EXISTS ${libc_path} AND IS_DIRECTORY ${libc_path})
+    add_library(llvm-libc-common-utilities INTERFACE)
+    # TODO: Reorganize the libc shared section so that it can be included without
+    # adding the root "libc" directory to the include path.
+    target_include_directories(llvm-libc-common-utilities INTERFACE ${libc_path})
+    target_compile_definitions(llvm-libc-common-utilities INTERFACE LIBC_NAMESPACE=__llvm_libc_common_utils)
+    target_compile_features(llvm-libc-common-utilities INTERFACE cxx_std_17)
+  endif()
+endif()