[libc] [llvm] [libc] Make RPC server handling header only (PR #131205)
Joseph Huber via llvm-commits
llvm-commits at lists.llvm.org
Thu Mar 13 14:33:16 PDT 2025
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/131205
>From f077fed8019feca55215e73784f76d07c7fbac52 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Thu, 13 Mar 2025 14:41:27 -0500
Subject: [PATCH] [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.
---
libc/docs/gpu/rpc.rst | 13 ++-
libc/shared/rpc_opcodes.h | 6 -
libc/shared/rpc_server.h | 22 ++++
.../__support/RPC/rpc_server.h} | 110 +++++++++++-------
libc/utils/gpu/CMakeLists.txt | 1 -
libc/utils/gpu/loader/CMakeLists.txt | 3 +
libc/utils/gpu/loader/Loader.h | 3 +-
libc/utils/gpu/loader/amdgpu/CMakeLists.txt | 8 +-
libc/utils/gpu/loader/nvptx/CMakeLists.txt | 8 +-
libc/utils/gpu/server/CMakeLists.txt | 30 -----
llvm/cmake/modules/FindLibcCommonUtils.cmake | 19 +++
offload/plugins-nextgen/common/CMakeLists.txt | 13 +--
offload/plugins-nextgen/common/src/RPC.cpp | 6 +-
13 files changed, 129 insertions(+), 113 deletions(-)
create mode 100644 libc/shared/rpc_server.h
rename libc/{utils/gpu/server/rpc_server.cpp => src/__support/RPC/rpc_server.h} (84%)
delete mode 100644 libc/utils/gpu/server/CMakeLists.txt
create mode 100644 llvm/cmake/modules/FindLibcCommonUtils.cmake
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/shared/rpc_server.h b/libc/shared/rpc_server.h
new file mode 100644
index 0000000000000..5509094b944ad
--- /dev/null
+++ b/libc/shared/rpc_server.h
@@ -0,0 +1,22 @@
+//===-- Shared RPC server interface -----------------------------*- C++ -*-===//
+//
+// 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
+
+#include "src/__support/RPC/rpc_server.h"
+
+namespace LIBC_NAMESPACE_DECL {
+namespace shared {
+
+using LIBC_NAMESPACE::rpc::handle_libc_opcodes;
+
+} // namespace shared
+} // namespace LIBC_NAMESPACE_DECL
+
+#endif // LLVM_LIBC_SHARED_RPC_SERVER_H
diff --git a/libc/utils/gpu/server/rpc_server.cpp b/libc/src/__support/RPC/rpc_server.h
similarity index 84%
rename from libc/utils/gpu/server/rpc_server.cpp
rename to libc/src/__support/RPC/rpc_server.h
index caffc0aee772b..7387eba9ceb26 100644
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/src/__support/RPC/rpc_server.h
@@ -5,25 +5,45 @@
// 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 <stdio.h>
-#include <stdlib.h>
+#include "hdr/stdio_overlay.h"
+#include "hdr/stdlib_overlay.h"
-namespace LIBC_NAMESPACE {
+namespace LIBC_NAMESPACE_DECL {
+namespace internal {
// Minimal replacement for 'std::vector' that works for trivial types.
template <typename T> class TempVector {
@@ -35,52 +55,50 @@ template <typename T> class TempVector {
size_t capacity;
public:
- TempVector() : data(nullptr), current(0), capacity(0) {}
+ LIBC_INLINE TempVector() : data(nullptr), current(0), capacity(0) {}
- ~TempVector() { free(data); }
+ LIBC_INLINE ~TempVector() { free(data); }
- void push_back(const T &value) {
+ LIBC_INLINE void push_back(const T &value) {
if (current == capacity)
grow();
data[current] = T(value);
++current;
}
- void push_back(T &&value) {
+ LIBC_INLINE void push_back(T &&value) {
if (current == capacity)
grow();
data[current] = T(static_cast<T &&>(value));
++current;
}
- void pop_back() { --current; }
+ LIBC_INLINE void pop_back() { --current; }
- bool empty() { return current == 0; }
+ LIBC_INLINE bool empty() { return current == 0; }
- size_t size() { return current; }
+ LIBC_INLINE size_t size() { return current; }
- T &operator[](size_t index) { return data[index]; }
+ LIBC_INLINE T &operator[](size_t index) { return data[index]; }
- T &back() { return data[current - 1]; }
+ LIBC_INLINE T &back() { return data[current - 1]; }
private:
- void grow() {
+ LIBC_INLINE 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 {
- char *alloc(size_t size) {
+ LIBC_INLINE char *alloc(size_t size) {
storage.push_back(reinterpret_cast<char *>(malloc(size)));
return storage.back();
}
- ~TempStorage() {
+ LIBC_INLINE ~TempStorage() {
for (size_t i = 0; i < storage.size(); ++i)
free(storage[i]);
}
@@ -88,15 +106,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) {
+LIBC_INLINE static ::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 +127,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) {
+LIBC_INLINE 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 ||
@@ -268,7 +287,8 @@ static void handle_printf(rpc::Server::Port &port, TempStorage &temp_storage) {
}
}
- results[lane] = fwrite(buffer, 1, writer.get_chars_written(), files[lane]);
+ results[lane] = static_cast<int>(
+ fwrite(buffer, 1, writer.get_chars_written(), files[lane]));
if (results[lane] != writer.get_chars_written() || ret == -1)
results[lane] = -1;
}
@@ -282,7 +302,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) {
+LIBC_INLINE static rpc::Status handle_port_impl(rpc::Server::Port &port) {
TempStorage temp_storage;
switch (port.get_opcode()) {
@@ -333,8 +353,9 @@ 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]),
- buffer->data[0], to_stream(buffer->data[1]));
+ const char *str = ::fgets(reinterpret_cast<char *>(data[id]),
+ static_cast<int>(buffer->data[0]),
+ to_stream(buffer->data[1]));
sizes[id] = !str ? 0 : __builtin_strlen(str) + 1;
});
port.send_n(data, sizes);
@@ -353,9 +374,9 @@ rpc::Status handle_port_impl(rpc::Server::Port &port) {
break;
}
case LIBC_CLOSE_FILE: {
- port.recv_and_send([&](rpc::Buffer *buffer, uint32_t id) {
+ port.recv_and_send([&](rpc::Buffer *buffer, uint32_t) {
FILE *file = reinterpret_cast<FILE *>(buffer->data[0]);
- buffer->data[0] = fclose(file);
+ buffer->data[0] = ::fclose(file);
});
break;
}
@@ -498,21 +519,28 @@ rpc::Status handle_port_impl(rpc::Server::Port &port) {
return rpc::RPC_SUCCESS;
}
-} // namespace LIBC_NAMESPACE
+} // namespace internal
+} // namespace LIBC_NAMESPACE_DECL
+namespace LIBC_NAMESPACE_DECL {
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.
+LIBC_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 internal::handle_port_impl<1>(port);
case 32:
- return LIBC_NAMESPACE::handle_port_impl<32>(port);
+ return internal::handle_port_impl<32>(port);
case 64:
- return LIBC_NAMESPACE::handle_port_impl<64>(port);
+ return internal::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
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..ec05117a041ab 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>
@@ -181,7 +182,7 @@ inline uint32_t handle_server(rpc::Server &server, uint32_t index,
break;
}
default:
- status = handle_libc_opcodes(*port, num_lanes);
+ status = LIBC_NAMESPACE::shared::handle_libc_opcodes(*port, num_lanes);
break;
}
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..fc90bb2e032f2 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,9 @@ 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
+ Status = LIBC_NAMESPACE::shared::handle_libc_opcodes(*Port,
+ Device.getWarpSize());
Port->close();
More information about the llvm-commits
mailing list