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

[jhuber6]

Summary

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.

[jdoerfert]

LG

libc/config/gpu/entrypoints.txt

set(TARGET_LIBC_ENTRYPOINTS
     libc.src.stdlib.lldiv
     libc.src.stdlib.strtod
     libc.src.stdlib.strtof
     libc.src.stdlib.strtol
     libc.src.stdlib.strtold
     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.src.stdlib.aligned_alloc
     libc.src.stdlib.realloc
     libc.src.stdlib.free
 
     # errno.h entrypoints

libc/docs/gpu/support.rst

 Function Name  Available  RPC Required
 =============  =========  ============
 abs            |check|
 atoi           |check|
 atof           |check|
 atol           |check|
 atoll          |check|
 exit           |check|    |check|
+abort          |check|    |check|
 labs           |check|
 llabs          |check|
 div            |check|
 ldiv           |check|
 lldiv          |check|
 strtod         |check|
 strtof         |check|
 strtol         |check|

libc/include/llvm-libc-types/rpc_opcodes_t.h

typedef enum : unsigned short {
   RPC_WRITE_TO_STREAM = 4,
   RPC_READ_FROM_STDIN = 5,
   RPC_READ_FROM_STREAM = 6,
   RPC_OPEN_FILE = 7,
   RPC_CLOSE_FILE = 8,
   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

add_entrypoint_object(
   SRCS
     free.cpp
   HDRS
     ../free.h
   DEPENDS
     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

+//===-- 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

case RPC_EXIT: {
       port->recv_and_send([](rpc::Buffer *) {});
       port->recv([](rpc::Buffer *buffer) {
         int status = 0;
         std::memcpy(&status, buffer->data, sizeof(int));
         exit(status);
       });
       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};
       port->recv_n(args, sizes, [&](uint64_t size) { return new char[size]; });
       port->recv([&](rpc::Buffer *buffer, uint32_t id) {
         reinterpret_cast<void (*)(void *)>(buffer->data[0])(args[id]);
       });
       port->send([&](rpc::Buffer *, uint32_t id) {