Revert "[libc] Make RPC server handling header only (llvm#131205)"

This reverts commit 8437b7f.

Author: ronlieb

libc/docs/gpu/rpc.rst

@@ -184,7 +184,6 @@ 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));
@@ -231,10 +230,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);
@@ -243,16 +242,14 @@ 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 GPU's
-handling is included through the ``shared/`` directory. This is not currently
-installed as it does not use a stable interface.
-
+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.
 
 .. code-block:: sh
 
   $> clang++ -x cuda rpc.cpp --offload-arch=native -fgpu-rdc -lcudart \
-       -I<install-path>include -L<install-path>/lib -Xoffload-linker -lc \
-       -O3 -foffload-lto -o hello
+       -I<install-path>include -L<install-path>/lib -lllvmlibc_rpc_server \
+       -Xoffload-linker -lc -O3 -foffload-lto -o hello
   $> ./hello
   Hello world!
 

libc/shared/rpc_opcodes.h

@@ -50,4 +50,10 @@ 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/shared/rpc_server.h

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

libc/utils/gpu/CMakeLists.txt

@@ -1,8 +1,5 @@
 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/CMakeLists.txt

@@ -13,7 +13,6 @@
 
 #include "shared/rpc.h"
 #include "shared/rpc_opcodes.h"
-#include "shared/rpc_server.h"
 
 #include <cstddef>
 #include <cstdint>
@@ -182,7 +181,7 @@ inline uint32_t handle_server(rpc::Server &server, uint32_t index,
     break;
   }
   default:
-    status = LIBC_NAMESPACE::shared::handle_libc_opcodes(*port, num_lanes);
+    status = handle_libc_opcodes(*port, num_lanes);
     break;
   }
 

libc/utils/gpu/loader/Loader.h

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

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

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

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

@@ -0,0 +1,30 @@
+add_library(llvmlibc_rpc_server STATIC rpc_server.cpp)
+
+# Include the RPC implemenation from libc.
+target_include_directories(llvmlibc_rpc_server PRIVATE ${LIBC_SOURCE_DIR})
+target_include_directories(llvmlibc_rpc_server PUBLIC ${LIBC_SOURCE_DIR}/include)
+target_include_directories(llvmlibc_rpc_server PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
+
+# Ignore unsupported clang attributes if we're using GCC.
+target_compile_options(llvmlibc_rpc_server PUBLIC
+                       $<$<CXX_COMPILER_ID:Clang>:-Wno-c99-extensions>
+                       $<$<CXX_COMPILER_ID:GNU>:-Wno-attributes>)
+target_compile_definitions(llvmlibc_rpc_server PUBLIC
+                           LIBC_COPT_USE_C_ASSERT
+                           LIBC_COPT_MEMCPY_USE_EMBEDDED_TINY
+                           LIBC_TYPES_LONG_DOUBLE_IS_FLOAT64
+                           LIBC_COPT_ARRAY_ARG_LIST
+                           LIBC_COPT_PRINTF_DISABLE_WRITE_INT
+                           LIBC_COPT_PRINTF_DISABLE_INDEX_MODE
+                           LIBC_COPT_PRINTF_DISABLE_STRERROR
+                           LIBC_NAMESPACE=${LIBC_NAMESPACE})
+
+# Install the server and associated header.
+install(FILES ${LIBC_SOURCE_DIR}/shared/rpc.h
+              ${LIBC_SOURCE_DIR}/shared/rpc_util.h
+              ${LIBC_SOURCE_DIR}/shared/rpc_opcodes.h
+        DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/shared
+        COMPONENT libc-headers)
+install(TARGETS llvmlibc_rpc_server
+        ARCHIVE DESTINATION "lib${LLVM_LIBDIR_SUFFIX}"
+        COMPONENT libc)

libc/utils/gpu/server/CMakeLists.txt

@@ -5,45 +5,25 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-//
-// This file is intended to be used externally as part of the `shared/`
-// interface. For that purpose, we manually define a few options normally
-// handled by the libc build system.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef LLVM_LIBC_SRC___SUPPORT_RPC_RPC_SERVER_H
-#define LLVM_LIBC_SRC___SUPPORT_RPC_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"
 
+#include "src/__support/CPP/type_traits.h"
 #include "src/__support/arg_list.h"
 #include "src/stdio/printf_core/converter.h"
 #include "src/stdio/printf_core/parser.h"
 #include "src/stdio/printf_core/writer.h"
 
-#include "hdr/stdio_overlay.h"
-#include "hdr/stdlib_overlay.h"
+#include <stdio.h>
+#include <stdlib.h>
 
-namespace LIBC_NAMESPACE_DECL {
-namespace internal {
+namespace LIBC_NAMESPACE {
 
 // Minimal replacement for 'std::vector' that works for trivial types.
 template <typename T> class TempVector {
@@ -55,66 +35,68 @@ template <typename T> class TempVector {
   size_t capacity;
 
 public:
-  LIBC_INLINE TempVector() : data(nullptr), current(0), capacity(0) {}
+  TempVector() : data(nullptr), current(0), capacity(0) {}
 
-  LIBC_INLINE ~TempVector() { free(data); }
+  ~TempVector() { free(data); }
 
-  LIBC_INLINE void push_back(const T &value) {
+  void push_back(const T &value) {
     if (current == capacity)
       grow();
     data[current] = T(value);
     ++current;
   }
 
-  LIBC_INLINE void push_back(T &&value) {
+  void push_back(T &&value) {
     if (current == capacity)
       grow();
     data[current] = T(static_cast<T &&>(value));
     ++current;
   }
 
-  LIBC_INLINE void pop_back() { --current; }
+  void pop_back() { --current; }
 
-  LIBC_INLINE bool empty() { return current == 0; }
+  bool empty() { return current == 0; }
 
-  LIBC_INLINE size_t size() { return current; }
+  size_t size() { return current; }
 
-  LIBC_INLINE T &operator[](size_t index) { return data[index]; }
+  T &operator[](size_t index) { return data[index]; }
 
-  LIBC_INLINE T &back() { return data[current - 1]; }
+  T &back() { return data[current - 1]; }
 
 private:
-  LIBC_INLINE void grow() {
+  void grow() {
     size_t new_capacity = capacity ? capacity * 2 : 1;
     void *new_data = realloc(data, new_capacity * sizeof(T));
+    if (!new_data)
+      abort();
     data = static_cast<T *>(new_data);
     capacity = new_capacity;
   }
 };
 
 struct TempStorage {
-  LIBC_INLINE char *alloc(size_t size) {
+  char *alloc(size_t size) {
     storage.push_back(reinterpret_cast<char *>(malloc(size)));
     return storage.back();
   }
 
-  LIBC_INLINE ~TempStorage() {
+  ~TempStorage() {
     for (size_t i = 0; i < storage.size(); ++i)
       free(storage[i]);
   }
 
   TempVector<char *> storage;
 };
 
-// Get the associated stream out of an encoded number.
-LIBC_INLINE static ::FILE *to_stream(uintptr_t f) {
-  enum Stream {
-    File = 0,
-    Stdin = 1,
-    Stdout = 2,
-    Stderr = 3,
-  };
+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) {
   ::FILE *stream = reinterpret_cast<FILE *>(f & ~0x3ull);
   Stream type = static_cast<Stream>(f & 0x3ull);
   if (type == Stdin)
@@ -127,8 +109,7 @@ LIBC_INLINE static ::FILE *to_stream(uintptr_t f) {
 }
 
 template <bool packed, uint32_t num_lanes>
-LIBC_INLINE static void handle_printf(rpc::Server::Port &port,
-                                      TempStorage &temp_storage) {
+static 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 ||
@@ -287,8 +268,7 @@ LIBC_INLINE static void handle_printf(rpc::Server::Port &port,
       }
     }
 
-    results[lane] = static_cast<int>(
-        fwrite(buffer, 1, writer.get_chars_written(), files[lane]));
+    results[lane] = fwrite(buffer, 1, writer.get_chars_written(), files[lane]);
     if (results[lane] != writer.get_chars_written() || ret == -1)
       results[lane] = -1;
   }
@@ -302,7 +282,7 @@ LIBC_INLINE static void handle_printf(rpc::Server::Port &port,
 }
 
 template <uint32_t num_lanes>
-LIBC_INLINE static rpc::Status handle_port_impl(rpc::Server::Port &port) {
+rpc::Status handle_port_impl(rpc::Server::Port &port) {
   TempStorage temp_storage;
 
   switch (port.get_opcode()) {
@@ -353,9 +333,8 @@ LIBC_INLINE static rpc::Status handle_port_impl(rpc::Server::Port &port) {
     void *data[num_lanes] = {nullptr};
     port.recv([&](rpc::Buffer *buffer, uint32_t id) {
       data[id] = temp_storage.alloc(buffer->data[0]);
-      const char *str = ::fgets(reinterpret_cast<char *>(data[id]),
-                                static_cast<int>(buffer->data[0]),
-                                to_stream(buffer->data[1]));
+      const char *str = fgets(reinterpret_cast<char *>(data[id]),
+                              buffer->data[0], to_stream(buffer->data[1]));
       sizes[id] = !str ? 0 : __builtin_strlen(str) + 1;
     });
     port.send_n(data, sizes);
@@ -374,9 +353,9 @@ LIBC_INLINE static rpc::Status handle_port_impl(rpc::Server::Port &port) {
     break;
   }
   case LIBC_CLOSE_FILE: {
-    port.recv_and_send([&](rpc::Buffer *buffer, uint32_t) {
+    port.recv_and_send([&](rpc::Buffer *buffer, uint32_t id) {
       FILE *file = reinterpret_cast<FILE *>(buffer->data[0]);
-      buffer->data[0] = ::fclose(file);
+      buffer->data[0] = fclose(file);
     });
     break;
   }
@@ -519,28 +498,21 @@ LIBC_INLINE static rpc::Status handle_port_impl(rpc::Server::Port &port) {
   return rpc::RPC_SUCCESS;
 }
 
-} // namespace internal
-} // namespace LIBC_NAMESPACE_DECL
+} // namespace LIBC_NAMESPACE
 
-namespace LIBC_NAMESPACE_DECL {
 namespace rpc {
-
-// Handles any opcode generated from the 'libc' client code.
-LIBC_INLINE ::rpc::Status handle_libc_opcodes(::rpc::Server::Port &port,
-                                              uint32_t num_lanes) {
+// 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) {
   switch (num_lanes) {
   case 1:
-    return internal::handle_port_impl<1>(port);
+    return LIBC_NAMESPACE::handle_port_impl<1>(port);
   case 32:
-    return internal::handle_port_impl<32>(port);
+    return LIBC_NAMESPACE::handle_port_impl<32>(port);
   case 64:
-    return internal::handle_port_impl<64>(port);
+    return LIBC_NAMESPACE::handle_port_impl<64>(port);
   default:
-    return ::rpc::RPC_ERROR;
+    return rpc::RPC_ERROR;
   }
 }
-
 } // namespace rpc
-} // namespace LIBC_NAMESPACE_DECL
-
-#endif // LLVM_LIBC_SRC___SUPPORT_RPC_RPC_SERVER_H