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/assert/gpu | |
| parent | e831a32c93c1ab404785773cc7c08c01730d61e5 (diff) | |
[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/assert/gpu')
| -rw-r--r-- | libc/src/assert/gpu/__assert_fail.cpp | 10 |
1 files changed, 2 insertions, 8 deletions
diff --git a/libc/src/assert/gpu/__assert_fail.cpp b/libc/src/assert/gpu/__assert_fail.cpp index b8ee168b069d..30c5f7fff602 100644 --- a/libc/src/assert/gpu/__assert_fail.cpp +++ b/libc/src/assert/gpu/__assert_fail.cpp @@ -27,14 +27,8 @@ LLVM_LIBC_FUNCTION(void, __assert_fail, uint32_t claimed = gpu::is_first_lane(mask) ? !lock.fetch_or(1, cpp::MemoryOrder::ACQUIRE) : 0; - if (!gpu::broadcast_value(mask, claimed)) { -#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(); - } + if (!gpu::broadcast_value(mask, claimed)) + gpu::end_program(); // Only a single line should be printed if an assertion is hit. if (gpu::is_first_lane(mask)) |