[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

Author: jhuber6

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

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|

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__

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
+)

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

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};