[libc] Make RPC server handling header only #131205
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
jhuber6
requested review from
JonChesterfield,
lntue,
michaelrj-google and
shiltian
llvmbot
added
cmake
|
@llvm/pr-subscribers-backend-amdgpu @llvm/pr-subscribers-libc Author: Joseph Huber (jhuber6) ChangesSummary: Generally, this makes it easier to share code without weird Full diff: https://github.com/llvm/llvm-project/pull/131205.diff 12 Files Affected:
diff --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst
index dde4207df6e49..1d6aaea2adcfe 100644
--- a/libc/docs/gpu/rpc.rst
+++ b/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!
diff --git a/libc/shared/rpc_opcodes.h b/libc/shared/rpc_opcodes.h
index 270c35dec28b8..6de41cd1899e7 100644
--- a/libc/shared/rpc_opcodes.h
+++ b/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
diff --git a/libc/utils/gpu/server/rpc_server.cpp b/libc/shared/rpc_server.h
similarity index 92%
rename from libc/utils/gpu/server/rpc_server.cpp
rename to libc/shared/rpc_server.h
index caffc0aee772b..6d1f52b69c02f 100644
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/shared/rpc_server.h
@@ -1,4 +1,4 @@
-//===-- 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.
@@ -6,11 +6,25 @@
//
//===----------------------------------------------------------------------===//
+#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
diff --git a/libc/utils/gpu/CMakeLists.txt b/libc/utils/gpu/CMakeLists.txt
index 7c15f36052cf3..e529646a1206e 100644
--- a/libc/utils/gpu/CMakeLists.txt
+++ b/libc/utils/gpu/CMakeLists.txt
@@ -1,2 +1 @@
-add_subdirectory(server)
add_subdirectory(loader)
diff --git a/libc/utils/gpu/loader/CMakeLists.txt b/libc/utils/gpu/loader/CMakeLists.txt
index 60597a67ce57a..9b3bd009dc0f1 100644
--- a/libc/utils/gpu/loader/CMakeLists.txt
+++ b/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
diff --git a/libc/utils/gpu/loader/Loader.h b/libc/utils/gpu/loader/Loader.h
index 8e86f63969326..d30a7c0c9669f 100644
--- a/libc/utils/gpu/loader/Loader.h
+++ b/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>
diff --git a/libc/utils/gpu/loader/amdgpu/CMakeLists.txt b/libc/utils/gpu/loader/amdgpu/CMakeLists.txt
index 80c5ae357416a..17878daf0b6fe 100644
--- a/libc/utils/gpu/loader/amdgpu/CMakeLists.txt
+++ b/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)
diff --git a/libc/utils/gpu/loader/nvptx/CMakeLists.txt b/libc/utils/gpu/loader/nvptx/CMakeLists.txt
index 21453b9ca0348..42510ac31dad4 100644
--- a/libc/utils/gpu/loader/nvptx/CMakeLists.txt
+++ b/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)
diff --git a/libc/utils/gpu/server/CMakeLists.txt b/libc/utils/gpu/server/CMakeLists.txt
deleted file mode 100644
index 7ca101e42a0af..0000000000000
--- a/libc/utils/gpu/server/CMakeLists.txt
+++ /dev/null
@@ -1,30 +0,0 @@
-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)
diff --git a/llvm/cmake/modules/FindLibcCommonUtils.cmake b/llvm/cmake/modules/FindLibcCommonUtils.cmake
new file mode 100644
index 0000000000000..0e65fdff7c34b
--- /dev/null
+++ b/llvm/cmake/modules/FindLibcCommonUtils.cmake
@@ -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()
diff --git a/offload/plugins-nextgen/common/CMakeLists.txt b/offload/plugins-nextgen/common/CMakeLists.txt
index de219efc8f79c..ffc431f68dbc5 100644
--- a/offload/plugins-nextgen/common/CMakeLists.txt
+++ b/offload/plugins-nextgen/common/CMakeLists.txt
@@ -21,20 +21,9 @@ if (NOT LLVM_LINK_LLVM_DYLIB)
endforeach()
endif()
-# Include the RPC server from the `libc` project if available.
+# Include the RPC server from the `libc` project.
include(FindLibcCommonUtils)
target_link_libraries(PluginCommon PRIVATE llvm-libc-common-utilities)
-if(TARGET llvmlibc_rpc_server AND ${LIBOMPTARGET_GPU_LIBC_SUPPORT})
- target_link_libraries(PluginCommon PRIVATE llvmlibc_rpc_server)
- target_compile_definitions(PluginCommon PRIVATE LIBOMPTARGET_RPC_SUPPORT)
-elseif(${LIBOMPTARGET_GPU_LIBC_SUPPORT})
- find_library(llvmlibc_rpc_server NAMES llvmlibc_rpc_server
- PATHS ${LIBOMPTARGET_LLVM_LIBRARY_DIR} NO_DEFAULT_PATH)
- if(llvmlibc_rpc_server)
- target_link_libraries(PluginCommon PRIVATE ${llvmlibc_rpc_server})
- target_compile_definitions(PluginCommon PRIVATE LIBOMPTARGET_RPC_SUPPORT)
- endif()
-endif()
# Define the TARGET_NAME and DEBUG_PREFIX.
target_compile_definitions(PluginCommon PRIVATE
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 70f572923d4b1..670edb5f446ca 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -15,6 +15,7 @@
#include "shared/rpc.h"
#include "shared/rpc_opcodes.h"
+#include "shared/rpc_server.h"
using namespace llvm;
using namespace omp;
@@ -88,10 +89,8 @@ static rpc::Status runServer(plugin::GenericDeviceTy &Device, void *Buffer) {
handleOffloadOpcodes(Device, *Port, Device.getWarpSize());
// Let the `libc` library handle any other unhandled opcodes.
-#ifdef LIBOMPTARGET_RPC_SUPPORT
if (Status == rpc::RPC_UNHANDLED_OPCODE)
Status = handle_libc_opcodes(*Port, Device.getWarpSize());
-#endif
Port->close();
|
|
@llvm/pr-subscribers-offload Author: Joseph Huber (jhuber6) ChangesSummary: Generally, this makes it easier to share code without weird Full diff: https://github.com/llvm/llvm-project/pull/131205.diff 12 Files Affected:
diff --git a/libc/docs/gpu/rpc.rst b/libc/docs/gpu/rpc.rst
index dde4207df6e49..1d6aaea2adcfe 100644
--- a/libc/docs/gpu/rpc.rst
+++ b/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!
diff --git a/libc/shared/rpc_opcodes.h b/libc/shared/rpc_opcodes.h
index 270c35dec28b8..6de41cd1899e7 100644
--- a/libc/shared/rpc_opcodes.h
+++ b/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
diff --git a/libc/utils/gpu/server/rpc_server.cpp b/libc/shared/rpc_server.h
similarity index 92%
rename from libc/utils/gpu/server/rpc_server.cpp
rename to libc/shared/rpc_server.h
index caffc0aee772b..6d1f52b69c02f 100644
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/shared/rpc_server.h
@@ -1,4 +1,4 @@
-//===-- 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.
@@ -6,11 +6,25 @@
//
//===----------------------------------------------------------------------===//
+#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
diff --git a/libc/utils/gpu/CMakeLists.txt b/libc/utils/gpu/CMakeLists.txt
index 7c15f36052cf3..e529646a1206e 100644
--- a/libc/utils/gpu/CMakeLists.txt
+++ b/libc/utils/gpu/CMakeLists.txt
@@ -1,2 +1 @@
-add_subdirectory(server)
add_subdirectory(loader)
diff --git a/libc/utils/gpu/loader/CMakeLists.txt b/libc/utils/gpu/loader/CMakeLists.txt
index 60597a67ce57a..9b3bd009dc0f1 100644
--- a/libc/utils/gpu/loader/CMakeLists.txt
+++ b/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
diff --git a/libc/utils/gpu/loader/Loader.h b/libc/utils/gpu/loader/Loader.h
index 8e86f63969326..d30a7c0c9669f 100644
--- a/libc/utils/gpu/loader/Loader.h
+++ b/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>
diff --git a/libc/utils/gpu/loader/amdgpu/CMakeLists.txt b/libc/utils/gpu/loader/amdgpu/CMakeLists.txt
index 80c5ae357416a..17878daf0b6fe 100644
--- a/libc/utils/gpu/loader/amdgpu/CMakeLists.txt
+++ b/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)
diff --git a/libc/utils/gpu/loader/nvptx/CMakeLists.txt b/libc/utils/gpu/loader/nvptx/CMakeLists.txt
index 21453b9ca0348..42510ac31dad4 100644
--- a/libc/utils/gpu/loader/nvptx/CMakeLists.txt
+++ b/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)
diff --git a/libc/utils/gpu/server/CMakeLists.txt b/libc/utils/gpu/server/CMakeLists.txt
deleted file mode 100644
index 7ca101e42a0af..0000000000000
--- a/libc/utils/gpu/server/CMakeLists.txt
+++ /dev/null
@@ -1,30 +0,0 @@
-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)
diff --git a/llvm/cmake/modules/FindLibcCommonUtils.cmake b/llvm/cmake/modules/FindLibcCommonUtils.cmake
new file mode 100644
index 0000000000000..0e65fdff7c34b
--- /dev/null
+++ b/llvm/cmake/modules/FindLibcCommonUtils.cmake
@@ -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()
diff --git a/offload/plugins-nextgen/common/CMakeLists.txt b/offload/plugins-nextgen/common/CMakeLists.txt
index de219efc8f79c..ffc431f68dbc5 100644
--- a/offload/plugins-nextgen/common/CMakeLists.txt
+++ b/offload/plugins-nextgen/common/CMakeLists.txt
@@ -21,20 +21,9 @@ if (NOT LLVM_LINK_LLVM_DYLIB)
endforeach()
endif()
-# Include the RPC server from the `libc` project if available.
+# Include the RPC server from the `libc` project.
include(FindLibcCommonUtils)
target_link_libraries(PluginCommon PRIVATE llvm-libc-common-utilities)
-if(TARGET llvmlibc_rpc_server AND ${LIBOMPTARGET_GPU_LIBC_SUPPORT})
- target_link_libraries(PluginCommon PRIVATE llvmlibc_rpc_server)
- target_compile_definitions(PluginCommon PRIVATE LIBOMPTARGET_RPC_SUPPORT)
-elseif(${LIBOMPTARGET_GPU_LIBC_SUPPORT})
- find_library(llvmlibc_rpc_server NAMES llvmlibc_rpc_server
- PATHS ${LIBOMPTARGET_LLVM_LIBRARY_DIR} NO_DEFAULT_PATH)
- if(llvmlibc_rpc_server)
- target_link_libraries(PluginCommon PRIVATE ${llvmlibc_rpc_server})
- target_compile_definitions(PluginCommon PRIVATE LIBOMPTARGET_RPC_SUPPORT)
- endif()
-endif()
# Define the TARGET_NAME and DEBUG_PREFIX.
target_compile_definitions(PluginCommon PRIVATE
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index 70f572923d4b1..670edb5f446ca 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -15,6 +15,7 @@
#include "shared/rpc.h"
#include "shared/rpc_opcodes.h"
+#include "shared/rpc_server.h"
using namespace llvm;
using namespace omp;
@@ -88,10 +89,8 @@ static rpc::Status runServer(plugin::GenericDeviceTy &Device, void *Buffer) {
handleOffloadOpcodes(Device, *Port, Device.getWarpSize());
// Let the `libc` library handle any other unhandled opcodes.
-#ifdef LIBOMPTARGET_RPC_SUPPORT
if (Status == rpc::RPC_UNHANDLED_OPCODE)
Status = handle_libc_opcodes(*Port, Device.getWarpSize());
-#endif
Port->close();
|
michaelrj-google
requested changes
Mar 13, 2025
libc/shared/rpc_server.h
Outdated
| @@ -1,16 +1,30 @@ | |||
| //===-- Shared memory RPC server instantiation ------------------*- C++ -*-===// | |||
| //===-- RPC server handling -----------------------------------------------===// | |||
There was a problem hiding this comment.
this file doesn't go in /shared, since that's only for the proxy headers to expose this code externally. This should go in /src/__support/ then a proxy header similar to /shared/fpbits.h should expose a narrow interface from this.
Comment on lines
14
to
+16
| #include "shared/rpc.h" | ||
| #include "shared/rpc_opcodes.h" | ||
| #include "shared/rpc_server.h" |
There was a problem hiding this comment.
the libc code isn't intended to include from /shared, if you need to use code both inside of libc and externally it should go in /src/__support with /shared only exporting the parts you want to use in other projects.
There was a problem hiding this comment.
It's going to be really tough making that work but I'll try.
There was a problem hiding this comment.
Okay I think this works, PTAL.
There was a problem hiding this comment.
now that you've moved rpc_server.h to __support, this should be fine to change to #include "src/__support/RPC/rpc_server.h". Then in a followup you can do the same thing with the other rpc headers in /shared so that they're all consistent.
There was a problem hiding this comment.
But we want to use the shared interface since this is exported from libc right
There was a problem hiding this comment.
I.e. we need to use the shared interface since it's the one that defines the LIBC_NAMESpACE and other stuff.
There was a problem hiding this comment.
the shared interface is for external callers, but internally libc should use its own interface. Also I forgot about the previous discussion on placement of the other rpc headers in /shared, they're fine there for now but iirc the plan is to move them to /include/rpc eventually. Writing that down so I remember it.
There was a problem hiding this comment.
But this isn't internal, this will be moved outside of libc soon and the offload/ stuff is definitely outside of libc.
There was a problem hiding this comment.
ah, if this is getting moved outside of libc then this is correct. I was under the impression you weren't planning to move it.
jhuber6
force-pushed
the
RemoveRPCServer
branch
from
bc4046e to
f077fed
Compare
michaelrj-google
approved these changes
Mar 13, 2025
Comment on lines
14
to
+16
| #include "shared/rpc.h" | ||
| #include "shared/rpc_opcodes.h" | ||
| #include "shared/rpc_server.h" |
There was a problem hiding this comment.
ah, if this is getting moved outside of libc then this is correct. I was under the impression you weren't planning to move it.
|
Yeah, I'm considering installing this somehow. We can still provide a library that just defines a linkable entrypoint for thie header only function. What do you think is the reasonable way to install that? |
|
You could add the |
True, though we'd still need to export the C++ header. But it's completely freestanding and independent from the rest of |
you could also add a small installable library to |
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.
jhuber6
force-pushed
the
RemoveRPCServer
branch
from
f077fed to
8732a49
Compare
searlmc1
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Mar 14, 2025
This reverts commit 8437b7f.
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/146/builds/2485 Here is the relevant piece of the build log for the reference |
|
This commit broke the build on s390x because libc prevents building on this architecture: |
I see, realistically that could probably default to a simple version, no reason for it to error when there's a perfectly usable fallback. |
searlmc1
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Mar 14, 2025
needed to add sym link: temporarily ./runtimes/cmake/Modules/FindLibcCommonUtils.cmake -> ../../../llvm/cmake/modules/FindLibcCommonUtils.cmake 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.
searlmc1
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Mar 14, 2025
|
This commit broke standalone offload builds: Looks like it's missing an include directory. |
Uhh, do those go through the runtime build? I think that might be a side-effect of me moving where the |
If I understand the question correctly, then no — we build straight from
Ah, yes, that explains it. Since the file is part of installed LLVM,
I can only test it partially right now — after restoring the file from the preceding commit, it fails due to missing files from Also, if you wanted to move it out of |
All that file does is point to the
Good point. |
mgorny
added a commit
to mgorny/llvm-project
that referenced
this pull request
Mar 17, 2025
…uilds Move `FindLibcCommonUtils` from LLVM's CMake module directory to the shared top-level CMake directory, as the module is intended to be used from within the source tree rather than the installed LLVM version. This fixes standalone offload builds after llvm#131205.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
7 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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 withinLLVM 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
libcproject sothat we don't need to bootstrap those and can build them in LLVM.