[libc-commits] [libc] 07102a1 - [libc] Implement the 'abort' function on the GPU

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2023-08-31T08:40:15-05:00
New Revision: 07102a11941d9287b58063e3d764694974205d53

URL: https://github.com/llvm/llvm-project/commit/07102a11941d9287b58063e3d764694974205d53
DIFF: https://github.com/llvm/llvm-project/commit/07102a11941d9287b58063e3d764694974205d53.diff

LOG: [libc] Implement the 'abort' function on the GPU

This function implements the `abort` function on the GPU. The
implementation here closely mirros the `exit` call where we first
synchornize with the RPC server to make sure it's listening and then we
exit on the GPU.

I was unsure if this should be a simple `__builtin_assert` on the GPU. I
elected to go with an RPC approach to make this a more "true" `abort`
call. That is, it should invoke some signal handlers and exit with the
proper code according to the implemented C library on the server.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D159210

Added: 
    libc/src/stdlib/gpu/abort.cpp

Modified: 
    libc/config/gpu/entrypoints.txt
    libc/docs/gpu/support.rst
    libc/include/llvm-libc-types/rpc_opcodes_t.h
    libc/src/stdlib/gpu/CMakeLists.txt
    libc/utils/gpu/server/rpc_server.cpp

Removed: 
    


################################################################################
diff  --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
index bc8f9649f49168..1d3047ab3c146b 100644
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -65,11 +65,10 @@ set(TARGET_LIBC_ENTRYPOINTS
     libc.src.stdlib.strtoll
     libc.src.stdlib.strtoul
     libc.src.stdlib.strtoull
-
-    # stdlib.h entrypoints
     libc.src.stdlib._Exit
     libc.src.stdlib.atexit
     libc.src.stdlib.exit
+    libc.src.stdlib.abort
 
     # Only implemented in the test suite
     libc.src.stdlib.malloc

diff  --git a/libc/docs/gpu/support.rst b/libc/docs/gpu/support.rst
index 3818b65d12863d..d1c27c7e8032d7 100644
--- a/libc/docs/gpu/support.rst
+++ b/libc/docs/gpu/support.rst
@@ -90,6 +90,7 @@ atof           |check|
 atol           |check|
 atoll          |check|
 exit           |check|    |check|
+abort          |check|    |check|
 labs           |check|
 llabs          |check|
 div            |check|

diff  --git a/libc/include/llvm-libc-types/rpc_opcodes_t.h b/libc/include/llvm-libc-types/rpc_opcodes_t.h
index f28fb5babf4fcd..3916c0ad953892 100644
--- a/libc/include/llvm-libc-types/rpc_opcodes_t.h
+++ b/libc/include/llvm-libc-types/rpc_opcodes_t.h
@@ -22,6 +22,7 @@ typedef enum : unsigned short {
   RPC_MALLOC = 9,
   RPC_FREE = 10,
   RPC_HOST_CALL = 11,
+  RPC_ABORT = 12,
 } rpc_opcode_t;
 
 #endif // __LLVM_LIBC_TYPES_RPC_OPCODE_H__

diff  --git a/libc/src/stdlib/gpu/CMakeLists.txt b/libc/src/stdlib/gpu/CMakeLists.txt
index 6a5d57623fd410..71ae10648d004e 100644
--- a/libc/src/stdlib/gpu/CMakeLists.txt
+++ b/libc/src/stdlib/gpu/CMakeLists.txt
@@ -19,3 +19,14 @@ add_entrypoint_object(
     libc.include.stdlib
     libc.src.__support.RPC.rpc_client
 )
+
+add_entrypoint_object(
+  abort
+  SRCS
+    abort.cpp
+  HDRS
+    ../abort.h
+  DEPENDS
+    libc.include.stdlib
+    libc.src.__support.RPC.rpc_client
+)

diff  --git a/libc/src/stdlib/gpu/abort.cpp b/libc/src/stdlib/gpu/abort.cpp
new file mode 100644
index 00000000000000..5e208dba88169d
--- /dev/null
+++ b/libc/src/stdlib/gpu/abort.cpp
@@ -0,0 +1,31 @@
+//===-- GPU implementation of abort ---------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/__support/RPC/rpc_client.h"
+#include "src/__support/common.h"
+
+#include "src/stdlib/abort.h"
+
+namespace __llvm_libc {
+
+LLVM_LIBC_FUNCTION(void, abort, ()) {
+  // We want to first make sure the server is listening before we abort.
+  rpc::Client::Port port = rpc::client.open<RPC_ABORT>();
+  port.send_and_recv([](rpc::Buffer *) {}, [](rpc::Buffer *) {});
+  port.send([&](rpc::Buffer *) {});
+  port.close();
+
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+  LIBC_INLINE_ASM("exit;" ::: "memory");
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  __builtin_amdgcn_endpgm();
+#endif
+  __builtin_unreachable();
+}
+
+} // namespace __llvm_libc

diff  --git a/libc/utils/gpu/server/rpc_server.cpp b/libc/utils/gpu/server/rpc_server.cpp
index 2a5494f7d3e978..22e937f0595c6a 100644
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/utils/gpu/server/rpc_server.cpp
@@ -145,6 +145,13 @@ struct Server {
       });
       break;
     }
+    case RPC_ABORT: {
+      // Send a response to the client to signal that we are ready to abort.
+      port->recv_and_send([](rpc::Buffer *) {});
+      port->recv([](rpc::Buffer *) {});
+      abort();
+      break;
+    }
     case RPC_HOST_CALL: {
       uint64_t sizes[lane_size] = {0};
       void *args[lane_size] = {nullptr};