diff --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -65,11 +65,10 @@
     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
--- a/libc/docs/gpu/support.rst
+++ b/libc/docs/gpu/support.rst
@@ -90,6 +90,7 @@
 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
--- a/libc/include/llvm-libc-types/rpc_opcodes_t.h
+++ b/libc/include/llvm-libc-types/rpc_opcodes_t.h
@@ -22,6 +22,7 @@
   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
--- a/libc/src/stdlib/gpu/CMakeLists.txt
+++ b/libc/src/stdlib/gpu/CMakeLists.txt
@@ -19,3 +19,14 @@
     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
--- /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
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/utils/gpu/server/rpc_server.cpp
@@ -145,6 +145,13 @@
       });
       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};