[OpenMP] Unconditionally provide an RPC client interface for OpenMP #117933
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
carlobertolli,
dhruvachak,
JonChesterfield,
jplehr,
ronlieb,
saiislam and
shiltian
|
@llvm/pr-subscribers-backend-amdgpu @llvm/pr-subscribers-offload Author: Joseph Huber (jhuber6) ChangesSummary: The interface here uses a weak symbol for the RPC client by the same The presense of this symbol is what controls whether or not we set up Full diff: https://github.com/llvm/llvm-project/pull/117933.diff 6 Files Affected:
diff --git a/offload/DeviceRTL/CMakeLists.txt b/offload/DeviceRTL/CMakeLists.txt
index c76ad018ab4fe7..3da83e5c307132 100644
--- a/offload/DeviceRTL/CMakeLists.txt
+++ b/offload/DeviceRTL/CMakeLists.txt
@@ -131,6 +131,7 @@ set(bc_flags -c -foffload-lto -std=c++17 -fvisibility=hidden
-DOMPTARGET_DEVICE_RUNTIME
-I${include_directory}
-I${devicertl_base_directory}/../include
+ -I${LLVM_MAIN_SRC_DIR}/../libc
${LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL}
)
@@ -275,6 +276,7 @@ function(compileDeviceRTLLibrary target_cpu target_name target_triple)
target_compile_definitions(${ide_target_name} PRIVATE SHARED_SCRATCHPAD_SIZE=512)
target_include_directories(${ide_target_name} PRIVATE
${include_directory}
+ ${LLVM_MAIN_SRC_DIR}/../libc
${devicertl_base_directory}/../include
${LIBOMPTARGET_LLVM_INCLUDE_DIRS}
)
diff --git a/offload/DeviceRTL/src/Misc.cpp b/offload/DeviceRTL/src/Misc.cpp
index 8e690f6fd8e7ce..c1df477365bcb6 100644
--- a/offload/DeviceRTL/src/Misc.cpp
+++ b/offload/DeviceRTL/src/Misc.cpp
@@ -12,6 +12,8 @@
#include "Allocator.h"
#include "Configuration.h"
#include "DeviceTypes.h"
+#include "Shared/RPCOpcodes.h"
+#include "shared/rpc.h"
#include "Debug.h"
@@ -110,6 +112,12 @@ void *indirectCallLookup(void *HstPtr) {
return HstPtr;
}
+/// The openmp client instance used to communicate with the server.
+/// FIXME: This is marked as 'retain' so that it is not removed via
+/// `-mlink-builtin-bitcode`
+[[gnu::visibility("protected"), gnu::weak,
+ gnu::retain]] rpc::Client Client asm("__llvm_rpc_client");
+
} // namespace impl
} // namespace ompx
@@ -156,6 +164,20 @@ void omp_free(void *ptr, omp_allocator_handle_t allocator) {
return;
}
}
+
+unsigned long long __llvm_omp_host_call(void *fn, void *data, size_t size) {
+ rpc::Client::Port Port = ompx::impl::Client.open<OFFLOAD_HOST_CALL>();
+ Port.send_n(data, size);
+ Port.send([=](rpc::Buffer *buffer, uint32_t) {
+ buffer->data[0] = reinterpret_cast<uintptr_t>(fn);
+ });
+ unsigned long long Ret;
+ Port.recv([&](rpc::Buffer *Buffer, uint32_t) {
+ Ret = static_cast<unsigned long long>(Buffer->data[0]);
+ });
+ Port.close();
+ return Ret;
+}
}
///}
diff --git a/offload/DeviceRTL/src/exports b/offload/DeviceRTL/src/exports
index 288ddf90b4a9f2..01667e7aba827a 100644
--- a/offload/DeviceRTL/src/exports
+++ b/offload/DeviceRTL/src/exports
@@ -15,4 +15,5 @@ malloc
free
memcmp
printf
+__llvm_rpc_client
__assert_fail
diff --git a/offload/include/Shared/RPCOpcodes.h b/offload/include/Shared/RPCOpcodes.h
new file mode 100644
index 00000000000000..beee29df1f7076
--- /dev/null
+++ b/offload/include/Shared/RPCOpcodes.h
@@ -0,0 +1,25 @@
+//===-- Shared/RPCOpcodes.h - Offload specific RPC opcodes ----- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines RPC opcodes that are specifically used by the OpenMP device runtime.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OMPTARGET_SHARED_RPC_OPCODES_H
+#define OMPTARGET_SHARED_RPC_OPCODES_H
+
+#define LLVM_OFFLOAD_RPC_BASE 'o'
+#define LLVM_OFFLOAD_OPCODE(n) (LLVM_OFFLOAD_RPC_BASE << 24 | n)
+
+typedef enum {
+ OFFLOAD_HOST_CALL = LLVM_OFFLOAD_OPCODE(0),
+} offload_opcode_t;
+
+#undef LLVM_OFFLOAD_OPCODE
+
+#endif // OMPTARGET_SHARED_RPC_OPCODES_H
diff --git a/offload/plugins-nextgen/common/src/RPC.cpp b/offload/plugins-nextgen/common/src/RPC.cpp
index c35431da69eb65..38509b5d78e615 100644
--- a/offload/plugins-nextgen/common/src/RPC.cpp
+++ b/offload/plugins-nextgen/common/src/RPC.cpp
@@ -9,6 +9,7 @@
#include "RPC.h"
#include "Shared/Debug.h"
+#include "Shared/RPCOpcodes.h"
#include "PluginInterface.h"
@@ -93,6 +94,22 @@ Error RPCServerTy::runServer(plugin::GenericDeviceTy &Device) {
});
break;
}
+ case OFFLOAD_HOST_CALL: {
+ uint64_t Sizes[64] = {0};
+ unsigned long long Results[64] = {0};
+ void *Args[64] = {nullptr};
+ Port->recv_n(Args, Sizes, [&](uint64_t Size) { return new char[Size]; });
+ Port->recv([&](rpc::Buffer *buffer, uint32_t ID) {
+ using FuncPtrTy = unsigned long long (*)(void *);
+ auto Func = reinterpret_cast<FuncPtrTy>(buffer->data[0]);
+ Results[ID] = Func(Args[ID]);
+ });
+ Port->send([&](rpc::Buffer *Buffer, uint32_t ID) {
+ Buffer->data[0] = static_cast<uint64_t>(Results[ID]);
+ delete[] reinterpret_cast<char *>(Args[ID]);
+ });
+ break;
+ }
default:
// Let the `libc` library handle any other unhandled opcodes.
Status = handle_libc_opcodes(*Port, Device.getWarpSize());
diff --git a/offload/test/libc/host_call.c b/offload/test/api/omp_host_call.c
similarity index 82%
rename from offload/test/libc/host_call.c
rename to offload/test/api/omp_host_call.c
index 61c4e14d5b3881..a3d1a97822b31b 100644
--- a/offload/test/libc/host_call.c
+++ b/offload/test/api/omp_host_call.c
@@ -1,20 +1,18 @@
// RUN: %libomptarget-compile-run-and-check-generic
-// REQUIRES: libc
-
#include <assert.h>
#include <omp.h>
#include <stdio.h>
#pragma omp begin declare variant match(device = {kind(gpu)})
// Extension provided by the 'libc' project.
-unsigned long long rpc_host_call(void *fn, void *args, size_t size);
-#pragma omp declare target to(rpc_host_call) device_type(nohost)
+unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size);
+#pragma omp declare target to(__llvm_omp_host_call) device_type(nohost)
#pragma omp end declare variant
#pragma omp begin declare variant match(device = {kind(cpu)})
// Dummy host implementation to make this work for all targets.
-unsigned long long rpc_host_call(void *fn, void *args, size_t size) {
+unsigned long long __llvm_omp_host_call(void *fn, void *args, size_t size) {
return ((unsigned long long (*)(void *))fn)(args);
}
#pragma omp end declare variant
@@ -58,7 +56,8 @@ int main() {
#pragma omp parallel num_threads(2)
{
args_t args = {omp_get_thread_num(), omp_get_team_num()};
- unsigned long long res = rpc_host_call(fn_ptr, &args, sizeof(args_t));
+ unsigned long long res =
+ __llvm_omp_host_call(fn_ptr, &args, sizeof(args_t));
printf("Result: %d\n", (int)res);
}
}
|
jplehr
reviewed
Nov 28, 2024
jhuber6
force-pushed
the
OpenMPRPCClient
branch
from
913ddc8 to
08c8d80
Compare
|
Thank you. IMHO that's easier to understand. |
Summary: This patch adds an RPC interface that lives directly in the OpenMP device runtime. This allows OpenMP to implement custom opcodes. Currently this is only providing the host call interface, which is the raw version of reverse offloading. Previously this lived in `libc/` as an extension which is not the correct place. The interface here uses a weak symbol for the RPC client by the same name that the `libc` interface uses. This means that it will defer to the libc one if both are present so we don't need to set up multiple instances. The presense of this symbol is what controls whether or not we set up the RPC server. Because this is an external symbol it normally won't be optimized out, so there's a special pass in OpenMPOpt that deletes this symbol if it is unused during linking. That means at `O0` the RPC server will always be present now, but will be removed trivially if it's not used at O1 and higher.
jhuber6
force-pushed
the
OpenMPRPCClient
branch
from
08c8d80 to
b9a1a67
Compare
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
jhuber6
force-pushed
the
OpenMPRPCClient
branch
from
b9a1a67 to
3d5d76e
Compare
jplehr
approved these changes
Dec 2, 2024
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/73/builds/9461 Here is the relevant piece of the build log for the reference |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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 adds an RPC interface that lives directly in the OpenMP
device runtime. This allows OpenMP to implement custom opcodes.
Currently this is only providing the host call interface, which is the
raw version of reverse offloading. Previously this lived in
libc/asan extension which is not the correct place.
The interface here uses a weak symbol for the RPC client by the same
name that the
libcinterface uses. This means that it will defer tothe libc one if both are present so we don't need to set up multiple
instances.
The presense of this symbol is what controls whether or not we set up
the RPC server. Because this is an external symbol it normally won't be
optimized out, so there's a special pass in OpenMPOpt that deletes this
symbol if it is unused during linking. That means at
O0the RPC serverwill always be present now, but will be removed trivially if it's not
used at O1 and higher.