diff options
| author | Joseph Huber <35342157+jhuber6@users.noreply.github.com> | 2023-09-12 14:59:02 -0500 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2023-09-12 14:59:02 -0500 |
| commit | 688019851e696f13ce4701fc9a8ec3e116221288 (patch) | |
| tree | b760cb7732d26cd84ff5bce5f892e9af12db8019 /libc/src/stdlib | |
| parent | e831a32c93c1ab404785773cc7c08c01730d61e5 (diff) | |
| download | llvm-688019851e696f13ce4701fc9a8ec3e116221288.zip llvm-688019851e696f13ce4701fc9a8ec3e116221288.tar.gz llvm-688019851e696f13ce4701fc9a8ec3e116221288.tar.bz2 | |
[libc][NFC] Factor GPU exiting into a common function (#66093)
Summary:
We currently call the GPU routine to terminate the current thread in
three separate locations .This should be wrapped into a helper function
to simplify the implementation.
Diffstat (limited to 'libc/src/stdlib')
| -rw-r--r-- | libc/src/stdlib/gpu/abort.cpp | 7 |
1 files changed, 1 insertions, 6 deletions
diff --git a/libc/src/stdlib/gpu/abort.cpp b/libc/src/stdlib/gpu/abort.cpp index 5e208db..441e776 100644 --- a/libc/src/stdlib/gpu/abort.cpp +++ b/libc/src/stdlib/gpu/abort.cpp @@ -20,12 +20,7 @@ LLVM_LIBC_FUNCTION(void, abort, ()) { 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(); + gpu::end_program(); } } // namespace __llvm_libc |